ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-17 03:19:48 +00:00
Code Issues Packages Projects Releases Wiki Activity

tune

Browse Source 
[ROCm/composable_kernel commit: e9ac4855f8]
This commit is contained in:
Chao Liu
2019-01-21 16:38:13 -06:00
parent 41c65ce2db
commit a8ef877996
4 changed files with 18 additions and 15 deletions
Download Patch File Download Diff File Expand all files Collapse all files

8
src/include/gemm.cuh
View File

@@ -388,9 +388,9 @@ struct blockwise_gemm_block_a_block_b_thread_c
const unsigned thread_work_cluster_id =
thread_id - cluster_work_block_id * (MThreadPerCluster * NThreadPerCluster);
const unsigned m_cluster_work_block_id = cluster_work_block_id / NThreadPerCluster;
const unsigned m_cluster_work_block_id = cluster_work_block_id / NClusterWork;
const unsigned n_cluster_work_block_id =
cluster_work_block_id - m_cluster_work_block_id * NThreadPerCluster;
cluster_work_block_id - m_cluster_work_block_id * NClusterWork;
const unsigned m_thread_work_cluster_id =
thread_work_cluster_id / NThreadPerCluster;
@@ -401,12 +401,12 @@ struct blockwise_gemm_block_a_block_b_thread_c
if(get_block_1d_id() == 0)
{
printf("%u %u, \t"
//"MClusterWork %u MThreadPerCluster %u NClusterWork %u NThreadPerCluster %u \t"
"MClusterWork %u MThreadPerCluster %u NClusterWork %u NThreadPerCluster %u \t"
"m_cluster_work_block_id %u n_cluster_work_block_id %u \t"
"m_thread_work_cluster_id %u n_thread_work_cluster_id %u \t"
"\n",
get_block_1d_id(), get_thread_local_1d_id(),
//MClusterWork, MThreadPerCluster, NClusterWork, NThreadPerCluster,
MClusterWork, MThreadPerCluster, NClusterWork, NThreadPerCluster,
m_cluster_work_block_id, n_cluster_work_block_id,
m_thread_work_cluster_id, n_thread_work_cluster_id);
}

5
src/include/gridwise_implicit_gemm_convolution_2_cnhw_srck_knhw.cuh
View File

@@ -239,10 +239,13 @@ gridwise_implicit_gemm_convolution_2_cnhw_srck_knhw(InGlobalDesc,
p_out_thread[out_kb_thread_desc.Get1dIndex(k, b)]);
}
#endif
if(k_data < K && n_data < N && h_data < Ho && w_data < Wo)
if(n_data < N && h_data < Ho && w_data < Wo)
{
#if 1
p_out_global[out_knhw_global_desc.Get1dIndex(k_data, n_data, h_data, w_data)] =
p_out_thread[out_kb_thread_desc.Get1dIndex(k, b)];
#endif
#if 0
if(get_block_1d_id() == 0)
{