diff --git a/example/ck_tile/99_toy_example/02_gemm/block_gemm_pipeline_agmem_bgmem_creg_default_policy.hpp b/example/ck_tile/99_toy_example/02_gemm/block_gemm_pipeline_agmem_bgmem_creg_default_policy.hpp index 9a9fa5a435..940c659031 100644 --- a/example/ck_tile/99_toy_example/02_gemm/block_gemm_pipeline_agmem_bgmem_creg_default_policy.hpp +++ b/example/ck_tile/99_toy_example/02_gemm/block_gemm_pipeline_agmem_bgmem_creg_default_policy.hpp @@ -27,6 +27,7 @@ struct BlockGemmPipelineAGmemBGmemCRegDefaultPolicy constexpr index_t kKPack = 8; #if BANK_CONFLICT_K_FIRST +#pragma message ("BANK_CONFLICT: K_FIRST") constexpr auto a_lds_block_desc_0 = make_naive_tensor_descriptor( make_tuple(number{}, number{}, number{}), make_tuple(number{}, number{}, number<1>{}), @@ -41,6 +42,7 @@ struct BlockGemmPipelineAGmemBGmemCRegDefaultPolicy make_tuple(sequence<0>{}, sequence<1>{})); #elif PADDING_K_FIRST +#pragma message ("BANK_CONFLICT: PADDING_K_FIRST") constexpr auto a_lds_block_desc_0 = make_naive_tensor_descriptor( make_tuple(number{}, number{}, number{}), make_tuple(number<(kKPerBlock / kKPack + 1) * kKPack>{}, number{}, number<1>{}), @@ -55,6 +57,7 @@ struct BlockGemmPipelineAGmemBGmemCRegDefaultPolicy make_tuple(sequence<0>{}, sequence<1>{})); #elif PADDING_MN_FIRST +#pragma message ("BANK_CONFLICT: PADDING_MN_FIRST") constexpr auto a_lds_block_desc_0 = make_naive_tensor_descriptor( make_tuple(number{}, number{}, number{}), make_tuple(number<(kMPerBlock + 1) * kKPack>{}, number{}, number<1>{}), @@ -69,6 +72,7 @@ struct BlockGemmPipelineAGmemBGmemCRegDefaultPolicy make_tuple(sequence<0>{}, sequence<1>{})); #elif XOR +#pragma message ("BANK_CONFLICT: XOR") using ADataType = remove_cvref_t; constexpr auto DataTypeSize = sizeof(ADataType); @@ -122,6 +126,7 @@ struct BlockGemmPipelineAGmemBGmemCRegDefaultPolicy constexpr index_t kKPack = 8; #if BANK_CONFLICT_K_FIRST +#pragma message ("BANK_CONFLICT: K_FIRST") constexpr auto b_lds_block_desc_0 = make_naive_tensor_descriptor( make_tuple(number{}, number{}, number{}), make_tuple(number{}, number{}, number<1>{}), @@ -136,6 +141,7 @@ struct BlockGemmPipelineAGmemBGmemCRegDefaultPolicy make_tuple(sequence<0>{}, sequence<1>{})); #elif PADDING_K_FIRST +#pragma message ("BANK_CONFLICT: PADDING_K_FIRST") constexpr auto b_lds_block_desc_0 = make_naive_tensor_descriptor( make_tuple(number{}, number{}, number{}), make_tuple(number<(kKPerBlock / kKPack + 1) * kKPack>{}, number{}, number<1>{}), @@ -150,6 +156,7 @@ struct BlockGemmPipelineAGmemBGmemCRegDefaultPolicy make_tuple(sequence<0>{}, sequence<1>{})); #elif PADDING_MN_FIRST +#pragma message ("BANK_CONFLICT: PADDING_MN_FIRST") constexpr auto b_lds_block_desc_0 = make_naive_tensor_descriptor( make_tuple(number{}, number{}, number{}), make_tuple(number<(kNPerBlock + 1) * kKPack>{}, number{}, number<1>{}), @@ -164,6 +171,7 @@ struct BlockGemmPipelineAGmemBGmemCRegDefaultPolicy make_tuple(sequence<0>{}, sequence<1>{})); #elif XOR +#pragma message ("BANK_CONFLICT: XOR") using BDataType = remove_cvref_t; constexpr auto DataTypeSize = sizeof(BDataType);