ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-17 00:58:44 +00:00
Code Issues Packages Projects Releases Wiki Activity

[GEMM] Add new macor options check

Browse Source
This commit is contained in:
Clement Lin
2025-03-30 10:07:21 +08:00
committed by Philip Maybank
parent 1099762267
commit a95665a6af
3 changed files with 18 additions and 3 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
example/ck_tile/99_toy_example/02_gemm/block_gemm_pipeline_agmem_bgmem_creg_default_policy.hpp
View File

@@ -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<kNPerBlock>{}, number<kKPerBlock / kKPack>{}, number<kKPack>{}),

11
example/ck_tile/99_toy_example/02_gemm/config.h
View File

@@ -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

8
example/ck_tile/99_toy_example/02_gemm/gemm.cpp
View File

@@ -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");