ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-20 21:09:08 +00:00
Code Issues Packages Projects Releases Wiki Activity

refactor

Browse Source 
[ROCm/composable_kernel commit: 42f4c7fd56]
This commit is contained in:
Chao Liu
2019-02-05 18:04:23 -06:00
parent ebf3d27947
commit 008027fef6
7 changed files with 63 additions and 31 deletions
Download Patch File Download Diff File Expand all files Collapse all files

10
driver/device_implicit_gemm_convolution_2_cnhw_csrk_knhw.cuh
View File

@@ -77,8 +77,8 @@ void device_implicit_gemm_convolution_2_cnhw_csrk_knhw(InDesc,
constexpr unsigned KPerThread = 16;
constexpr unsigned CPerThread = 1;
constexpr unsigned GemmRowThreadPerCluster = 4;
constexpr unsigned GemmColumnThreadPerCluster = 8;
constexpr unsigned GemmThreadPerColumnPerCluster = 4;
constexpr unsigned GemmThreadPerRowPerCluster = 8;
constexpr unsigned InBlockCopyThreadPerDim0 = 4;
constexpr unsigned InBlockCopyThreadPerDim1 = 16;
@@ -120,7 +120,7 @@ void device_implicit_gemm_convolution_2_cnhw_csrk_knhw(InDesc,
#if 1
gridwise_implicit_gemm_convolution_2_cnhw_csrk_knhw
#else
#elif 0
gridwise_implicit_gemm_convolution_2_cnhw_csrk_knhw_lds_pipeline
#endif
<GridSize,
@@ -135,8 +135,8 @@ void device_implicit_gemm_convolution_2_cnhw_csrk_knhw(InDesc,
BPerThread,
KPerThread,
CPerThread,
GemmRowThreadPerCluster,
GemmColumnThreadPerCluster,
GemmThreadPerColumnPerCluster,
GemmThreadPerRowPerCluster,
InBlockCopyThreadPerDim0,
InBlockCopyThreadPerDim1,
WeiBlockCopyThreadPerDim0,

20
driver/device_implicit_gemm_convolution_2_cnhw_srck_knhw.cuh
View File

@@ -76,8 +76,8 @@ void device_implicit_gemm_convolution_2_cnhw_srck_knhw(InDesc,
constexpr unsigned KPerThread = 1;
constexpr unsigned CPerThread = 1;
constexpr unsigned GemmThreadPerClusterRow = 1;
constexpr unsigned GemmThreadPerClusterColumn = 4;
constexpr unsigned GemmThreadPerColumnPerCluster = 1;
constexpr unsigned GemmThreadPerRowPerCluster = 4;
constexpr unsigned BlockSize = 32;
#elif 0
@@ -89,8 +89,8 @@ void device_implicit_gemm_convolution_2_cnhw_srck_knhw(InDesc,
constexpr unsigned KPerThread = 8;
constexpr unsigned CPerThread = 1;
constexpr unsigned GemmThreadPerClusterRow = 4;
constexpr unsigned GemmThreadPerClusterColumn = 4;
constexpr unsigned GemmThreadPerColumnPerCluster = 4;
constexpr unsigned GemmThreadPerRowPerCluster = 4;
constexpr unsigned BlockSize = 128;
#elif 0
@@ -102,8 +102,8 @@ void device_implicit_gemm_convolution_2_cnhw_srck_knhw(InDesc,
constexpr unsigned KPerThread = 8;
constexpr unsigned CPerThread = 1;
constexpr unsigned GemmRowThreadPerCluster = 4;
constexpr unsigned GemmColumnThreadPerCluster = 4;
constexpr unsigned GemmThreadPerColumnPerCluster = 4;
constexpr unsigned GemmThreadPerRowPerCluster = 4;
constexpr unsigned InBlockCopyThreadPerDim0 = 2;
constexpr unsigned InBlockCopyThreadPerDim1 = 64;
@@ -119,8 +119,8 @@ void device_implicit_gemm_convolution_2_cnhw_srck_knhw(InDesc,
constexpr unsigned KPerThread = 16;
constexpr unsigned CPerThread = 2;
constexpr unsigned GemmRowThreadPerCluster = 8;
constexpr unsigned GemmColumnThreadPerCluster = 8;
constexpr unsigned GemmThreadPerColumnPerCluster = 8;
constexpr unsigned GemmThreadPerRowPerCluster = 8;
constexpr unsigned InBlockCopyThreadPerDim0 = 8;
constexpr unsigned InBlockCopyThreadPerDim1 = 16;
@@ -171,8 +171,8 @@ void device_implicit_gemm_convolution_2_cnhw_srck_knhw(InDesc,
BPerThread,
KPerThread,
CPerThread,
GemmRowThreadPerCluster,
GemmColumnThreadPerCluster,
GemmThreadPerColumnPerCluster,
GemmThreadPerRowPerCluster,
InBlockCopyThreadPerDim0,
InBlockCopyThreadPerDim1>
<<<grid_dim, block_dim>>>(in_cnhw_desc,