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

[CK-Tile] fix block scale example for gfx1201 (#3283)

Browse Source
This commit is contained in:
Khushbu Agarwal
2025-11-25 13:10:28 -08:00
committed by GitHub
parent 9ac2666d5b
commit 37ea160088
3 changed files with 28 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files

5
example/ck_tile/38_block_scale_gemm/gemm_bquant_quantgrouped_preshuffleb.cpp
View File

@@ -3,8 +3,13 @@
#include "run_gemm_quant_example.inc"
#if CK_TILE_USE_WMMA
template <typename T>
using GemmConfig = GemmConfigPreshuffleB_BQuant_Prefill_Wmma<T>;
#else
template <typename T>
using GemmConfig = GemmConfigPreshuffleB_BQuant_Prefill<T>;
#endif
void bquant_quantgrouped_preshuffleb_instance_factory(
std::unordered_map<size_t, std::function<int(const ck_tile::ArgParser&)>>& lut)

5
example/ck_tile/38_block_scale_gemm/gemm_bquant_quantgrouped_preshuffleb_preshufflequant.cpp
View File

@@ -3,8 +3,13 @@
#include "run_gemm_quant_example.inc"
#if CK_TILE_USE_WMMA
template <typename T>
using GemmConfig = GemmConfigPreshuffleB_PreshuffleBQuant_Prefill_Wmma<T>;
#else
template <typename T>
using GemmConfig = GemmConfigPreshuffleB_PreshuffleBQuant_Prefill<T>;
#endif
void bquant_quantgrouped_preshuffleb_preshufflequant_instance_factory(
std::unordered_map<size_t, std::function<int(const ck_tile::ArgParser&)>>& lut)

18
example/ck_tile/38_block_scale_gemm/gemm_utils.hpp
View File

@@ -250,6 +250,24 @@ struct GemmConfigBQuantPrefill_Wmma : public GemmConfigBQuantPrefill<PrecType>
static constexpr ck_tile::index_t K_Warp_Tile = 16;
};
template <typename PrecType>
struct GemmConfigPreshuffleB_BQuant_Prefill_Wmma
: public GemmConfigPreshuffleB_BQuant_Prefill<PrecType>
{
static constexpr ck_tile::index_t M_Warp_Tile = 16;
static constexpr ck_tile::index_t N_Warp_Tile = 16;
static constexpr ck_tile::index_t K_Warp_Tile = 16;
};
template <typename PrecType>
struct GemmConfigPreshuffleB_PreshuffleBQuant_Prefill_Wmma
: public GemmConfigPreshuffleB_PreshuffleBQuant_Prefill<PrecType>
{
static constexpr ck_tile::index_t M_Warp_Tile = 16;
static constexpr ck_tile::index_t N_Warp_Tile = 16;
static constexpr ck_tile::index_t K_Warp_Tile = 16;
};
template <typename ADataType_,
typename BDataType_ = ADataType_,
typename CDataType_ = ADataType_,