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 b8faa00d85..989ea29a02 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 @@ -122,7 +122,7 @@ struct BlockGemmPipelineAGmemBGmemCRegDefaultPolicy constexpr index_t kKPerBlock = Problem::BlockGemmShape::kK; constexpr index_t kKPack = 8; -#if defined(NAIVE_IMPLEMENTATION) +#if defined(PADDING_K_FIRST) || defined(NAIVE_IMPLEMENTATION) #pragma message ("BANK_CONFLICT: K_FIRST") constexpr auto b_lds_block_desc_0 = make_naive_tensor_descriptor( make_tuple(number{}, number{}, number{}), diff --git a/example/ck_tile/99_toy_example/02_gemm/config.h b/example/ck_tile/99_toy_example/02_gemm/config.h index dfff4a6e12..fa531be0ff 100644 --- a/example/ck_tile/99_toy_example/02_gemm/config.h +++ b/example/ck_tile/99_toy_example/02_gemm/config.h @@ -1,11 +1,15 @@ #if defined(KERNEL_A) + #define PADDING_K_FIRST #define USING_MFMA_32x32x_8x2 #elif defined(KERNEL_B) + #define PADDING_K_FIRST #define USING_MFMA_16x16x16 #elif defined(KERNEL_C) + #define PADDING_K_FIRST #define USING_MFMA_16x16x_16x2 #elif defined(KERNEL_D) + #define PADDING_K_FIRST #define USING_MFMA_16x16x_16x2 #define USING_XOR_BASED_BANK_CONFLICT_FREE #elif defined(KERNEL_E) @@ -17,13 +21,18 @@ #define USING_XOR_BASED_BANK_CONFLICT_FREE #define ADJUST_BLOCK_TILE_SHAPE #define ENABLE_PREFETCH - #define ENABLE_INSTRUCTION_SCH #elif defined(KERNEL_G) #define USING_MFMA_16x16x_16x2 #define USING_XOR_BASED_BANK_CONFLICT_FREE #define ADJUST_BLOCK_TILE_SHAPE #define ENABLE_PREFETCH #define ENABLE_INSTRUCTION_SCH +#elif defined(KERNEL_H) + #define USING_MFMA_16x16x_16x2 + #define USING_XOR_BASED_BANK_CONFLICT_FREE + #define ADJUST_BLOCK_TILE_SHAPE + #define ENABLE_PREFETCH + #define ENABLE_INSTRUCTION_SCH #define ENABLE_CACHE_AWARE_WG_SCH #else #define NAIVE_IMPLEMENTATION diff --git a/example/ck_tile/99_toy_example/02_gemm/gemm.cpp b/example/ck_tile/99_toy_example/02_gemm/gemm.cpp index 6fb5c58ed8..56b9f4c965 100644 --- a/example/ck_tile/99_toy_example/02_gemm/gemm.cpp +++ b/example/ck_tile/99_toy_example/02_gemm/gemm.cpp @@ -72,8 +72,14 @@ int main(int argc, char* argv[]) printf(" --> XOR-based banck conflict-free\n"); printf(" --> Adjust block tile shape\n"); printf(" --> Enable prefetch\n"); - printf(" --> Enable instruction schedule\n"); #elif defined(KERNEL_G) + printf("*** KernelF test ***\n"); + printf(" --> Using mfma_16x16x(16x2)\n"); + printf(" --> XOR-based banck conflict-free\n"); + printf(" --> Adjust block tile shape\n"); + printf(" --> Enable prefetch\n"); + printf(" --> Enable instruction schedule\n"); +#elif defined(KERNEL_H) printf("*** KernelG test ***\n"); printf(" --> Using mfma_16x16x(16x2)\n"); printf(" --> XOR-based banck conflict-free\n");