ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-26 16:04:58 +00:00
Code Issues Packages Projects Releases Wiki Activity

GEMM/Conv+BiasAdd+ReLU+Add (#55)

Browse Source 
* gemm+activation

* move C pointwise operation into threadwise copy

* add pointwise operation to A/B matrix

* update ckProfiler

* adding bias add

* adding bias add

* adding bias add

* added bias add; worked around compiler issues

* clean up

* clean up

* Update README.md

* Update README.md

* Update README.md

* clean up

* add conv_xdl example

* adding conv_xdl_bias_relu_add example

* add conv+bias+relu+add, but has register spill issue

* tweak

* tweak

* refactor

* Update README.md

update readme for example/2_gemm_xdl_bias_relu_add

* clean up

* Update README.md

update readme for example/3_conv_xdl

* Update README.md

[ROCm/composable_kernel commit: 41cdd3801a]
This commit is contained in:
Chao Liu
2021-12-02 20:07:37 -06:00
committed by GitHub
parent 0572b85298
commit 9c85245412
40 changed files with 4336 additions and 250 deletions
Download Patch File Download Diff File Expand all files Collapse all files

16
device_operation/include/device_conv_instance.hpp
View File

@@ -2,6 +2,7 @@
#define DEVICE_CONV_INSTANTCE_HPP
#include "device_conv.hpp"
#include "element_wise_operation.hpp"
namespace ck {
namespace tensor_operation {
@@ -15,7 +16,10 @@ template <ck::index_t NDimSpatial,
typename InLayout,
typename WeiLayout,
typename OutLayout>
void add_device_conv_fwd_instance(std::vector<DeviceConvFwdPtr>&);
void add_device_conv_fwd_instance(
std::vector<DeviceConvFwdPtr<ck::tensor_operation::element_wise::PassThrough,
ck::tensor_operation::element_wise::PassThrough,
ck::tensor_operation::element_wise::PassThrough>>&);
template <ck::index_t NDimSpatial,
typename InDataType,
@@ -24,7 +28,10 @@ template <ck::index_t NDimSpatial,
typename InLayout,
typename WeiLayout,
typename OutLayout>
void add_device_conv_bwd_instance(std::vector<DeviceConvBwdPtr>&);
void add_device_conv_bwd_instance(
std::vector<DeviceConvBwdPtr<ck::tensor_operation::element_wise::PassThrough,
ck::tensor_operation::element_wise::PassThrough,
ck::tensor_operation::element_wise::PassThrough>>&);
template <ck::index_t NDimSpatial,
typename InDataType,
@@ -33,7 +40,10 @@ template <ck::index_t NDimSpatial,
typename InLayout,
typename WeiLayout,
typename OutLayout>
void add_device_conv_wrw_instance(std::vector<DeviceConvWrwPtr>&);
void add_device_conv_wrw_instance(
std::vector<DeviceConvWrwPtr<ck::tensor_operation::element_wise::PassThrough,
ck::tensor_operation::element_wise::PassThrough,
ck::tensor_operation::element_wise::PassThrough>>&);
} // namespace device_conv_instance
} // namespace device