ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-02 04:31:25 +00:00
Code Issues Packages Projects Releases Wiki Activity

[CK_TILE] Fix the Wrong Output Generated by Gemm Examples on GFX11/12 (#2713)

Browse Source 
* Introduce macro CK_TILE_USE_WMMA

Signed-off-by: Tianyuan Wu <Tianyuan.Wu@amd.com>

* Make CK_TILE_USE_WMMA global for all examples

Signed-off-by: Tianyuan Wu <Tianyuan.Wu@amd.com>

* Remove CK_TILE_USE_WMMA from config.hpp

Signed-off-by: Tianyuan Wu <Tianyuan.Wu@amd.com>

---------

Signed-off-by: Tianyuan Wu <Tianyuan.Wu@amd.com>
This commit is contained in:
Tianyuan Wu
2025-08-26 03:55:35 +08:00
committed by GitHub
parent 1d4a3341f0
commit e9605ed36d
4 changed files with 18 additions and 11 deletions
Download Patch File Download Diff File Expand all files Collapse all files

10
example/ck_tile/03_gemm/gemm_basic.cpp
View File

@@ -26,6 +26,15 @@ float gemm(const ck_tile::GemmHostArgs& args, const ck_tile::stream_config& s)
constexpr ck_tile::index_t N_Tile = 256;
constexpr ck_tile::index_t K_Tile = 64;
#if CK_TILE_USE_WMMA
constexpr ck_tile::index_t M_Warp = 4;
constexpr ck_tile::index_t N_Warp = 2;
constexpr ck_tile::index_t K_Warp = 1;
constexpr ck_tile::index_t M_Warp_Tile = 16;
constexpr ck_tile::index_t N_Warp_Tile = 16;
constexpr ck_tile::index_t K_Warp_Tile = 16;
#else
constexpr ck_tile::index_t M_Warp = 2;
constexpr ck_tile::index_t N_Warp = 2;
constexpr ck_tile::index_t K_Warp = 1;
@@ -33,6 +42,7 @@ float gemm(const ck_tile::GemmHostArgs& args, const ck_tile::stream_config& s)
constexpr ck_tile::index_t M_Warp_Tile = 32;
constexpr ck_tile::index_t N_Warp_Tile = 32;
constexpr ck_tile::index_t K_Warp_Tile = 16;
#endif
using CodegenGemmShape =
ck_tile::TileGemmShape<ck_tile::sequence<M_Tile, N_Tile, K_Tile>,