Skip to content

feat: support dimension sizes less than 8 #11776

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
illsilin merged 33 commits into from
Sep 22, 2026
Merged

feat: support dimension sizes less than 8 #11776

Show file tree
Hide file tree
Changes from 13 commits
Commits
Show all changes
33 commits
Select commit Hold shift + click to select a range
3c6c293
Add A,B,C vector sizes as parameters to GemmConfig
Sep 4, 2026
798bd15
support for m,n,k < 8
bidlekm Sep 8, 2026
6e23498
run_gemm_example simplify with dispatching
bidlekm Sep 8, 2026
b90cb7c
check the fastest dimensions
bidlekm Sep 8, 2026
fd25e9f
remove gfx1250 ifdefs
bidlekm Sep 8, 2026
9b9aa50
batched gemm workaround
bidlekm Sep 8, 2026
d245f7a
remove fp32 path added
bidlekm Sep 8, 2026
c75cbc4
templated lambda instead of tag dispatch
bidlekm Sep 9, 2026
7456663
fix: fall back to padded, natural-vector-size GemmConfig when shape i…
bidlekm Sep 10, 2026
ca1df2f
review changes
bidlekm Sep 11, 2026
d923e2b
clang format
bidlekm Sep 13, 2026
dff9e9c
Merge branch 'develop' into users/bidlekm/ck/lower-dim-gemm-enable
bidlekm Sep 14, 2026
2d0fee9
CI fix
bidlekm Sep 14, 2026
528a8e2
adding different batch gemm layouts and a compatibility fix for fp32 …
bidlekm Sep 15, 2026
3e45105
clang format
bidlekm Sep 15, 2026
dab60c3
CI fix
bidlekm Sep 16, 2026
95a2b58
CI fix
bidlekm Sep 16, 2026
cde934a
test fix
bidlekm Sep 17, 2026
bec445b
Merge branch 'develop' into users/bidlekm/ck/lower-dim-gemm-enable
bidlekm Sep 17, 2026
eb8a10e
Merge branch 'develop' into users/bidlekm/ck/lower-dim-gemm-enable
bidlekm Sep 18, 2026
3b5dfeb
Merge branch 'develop' into users/bidlekm/ck/lower-dim-gemm-enable
bidlekm Sep 18, 2026
7b6a9b3
Merge branch 'develop' into users/bidlekm/ck/lower-dim-gemm-enable
illsilin Sep 18, 2026
e6bc051
Merge branch 'develop' into users/bidlekm/ck/lower-dim-gemm-enable
bidlekm Sep 20, 2026
c4532cd
Merge branch 'develop' into users/bidlekm/ck/lower-dim-gemm-enable
bidlekm Sep 21, 2026
d64d4d1
Merge branch 'develop' into users/bidlekm/ck/lower-dim-gemm-enable
bidlekm Sep 21, 2026
d2ecd74
Merge branch 'develop' into users/bidlekm/ck/lower-dim-gemm-enable
bidlekm Sep 21, 2026
142fc3a
Merge branch 'develop' into users/bidlekm/ck/lower-dim-gemm-enable
bidlekm Sep 22, 2026
4cae78b
resolve merge conflict
bidlekm Sep 22, 2026
8bccbd5
Merge branch 'develop' into users/bidlekm/ck/lower-dim-gemm-enable
bidlekm Sep 22, 2026
db14be1
ci fix
bidlekm Sep 22, 2026
2326976
Merge branch 'users/bidlekm/ck/lower-dim-gemm-enable' of github.com:R…
bidlekm Sep 22, 2026
4332ab1
Merge branch 'develop' into users/bidlekm/ck/lower-dim-gemm-enable
bidlekm Sep 22, 2026
51d399c
Merge branch 'develop' into users/bidlekm/ck/lower-dim-gemm-enable
bidlekm Sep 22, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion projects/composablekernel/example/ck_tile/03_gemm/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,7 +1,7 @@
# Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
# SPDX-License-Identifier: MIT

if(GPU_TARGETS MATCHES "gfx94|gfx95|gfx90a|gfx125")
if(GPU_TARGETS MATCHES "gfx94|gfx95|gfx90a|gfx125|gfx11|gfx12")
add_executable(tile_example_gemm_basic gemm_basic.cpp)
add_executable(tile_example_gemm_universal universal_gemm.cpp)
add_executable(tile_example_gemm_weight_preshuffle gemm_weight_preshuffle.cpp)
Expand Down
5 changes: 4 additions & 1 deletion projects/composablekernel/example/ck_tile/03_gemm/gemm_splitk_two_stage_invoker.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -79,7 +79,10 @@ struct SplitKTwoStageInvoker
ck_tile::element_wise::PassThrough,
ck_tile::element_wise::PassThrough,
AComputeDataType,
BComputeDataType>;
BComputeDataType,
GemmConfig::FixedVectorSize,
GemmConfig::VectorSizeA,
GemmConfig::VectorSizeB>;
using WorkspaceType = ck_tile::remove_cvref_t<typename GemmConfig::WorkspaceType>;

using GemmPipeline = typename PipelineTypeTraits<
Expand Down
14 changes: 14 additions & 0 deletions projects/composablekernel/example/ck_tile/03_gemm/gemm_utils.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -8,10 +8,13 @@
#include "ck_tile/ops/epilogue.hpp"
#include "ck_tile/ops/gemm.hpp"
#include "ck_tile/utility/json_dump.hpp"
#include "gemm_common/vector_size_fallback_dispatch.hpp"

#include <string>
#include <variant>

using ck_tile_example::GemmConfigVectorSizeFallback;

struct GemmConfigBase
{
static constexpr bool kPadM = false;
Expand Down Expand Up @@ -41,6 +44,14 @@ struct GemmConfigBase
static constexpr ck_tile::DataCachePrefetchKind DataCachePrefetchB =
ck_tile::DataCachePrefetchKind::None;
static constexpr bool Async = false;

static constexpr bool FixedVectorSize = false;
// If FixedVectorSize==true: use these vector sizes for A/B loads and C store
static constexpr ck_tile::index_t VectorSizeA = 1;
static constexpr ck_tile::index_t VectorSizeB = 1;
static constexpr ck_tile::index_t VectorSizeC = 1;

static constexpr bool EnableSmallerVectorLoadFallback = false;
};

template <typename PrecType>
Expand Down Expand Up @@ -176,6 +187,7 @@ struct GemmConfigComputeV3_3 : public GemmConfigBase
template <typename PrecType>
struct GemmConfigComputeV3_WMMA : public GemmConfigBase
{

static constexpr ck_tile::index_t M_Tile = 128;
static constexpr ck_tile::index_t N_Tile = 128;
static constexpr ck_tile::index_t K_Tile = 64 / sizeof(PrecType);
Expand All @@ -193,6 +205,8 @@ struct GemmConfigComputeV3_WMMA : public GemmConfigBase
static constexpr ck_tile::GemmPipeline Pipeline = ck_tile::GemmPipeline::COMPUTE_V3;

static constexpr int kBlockPerCu = 2;

static constexpr bool EnableSmallerVectorLoadFallback = true;
};

template <typename PrecType>
Expand Down
62 changes: 37 additions & 25 deletions projects/composablekernel/example/ck_tile/03_gemm/run_gemm_example.inc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5,6 +5,7 @@
#include "ck_tile/host/permute_pk_int4.hpp"
#include "ck_tile/host/tensor_shuffle_utils.hpp"
#include "ck_tile/ops/common/utils.hpp"
#include "gemm_common/vector_size_fallback_dispatch.hpp"

template <typename Layout>
static constexpr inline auto is_row_major(Layout layout_)
Expand Down Expand Up @@ -347,31 +348,42 @@ int run_gemm_example_with_layouts(ck_tile::ArgParser& arg_parser,
c_m_n_dev_buf.SetZero();
c_m_n_dev_result.SetZero();

float ave_time = invoke_gemm<GemmConfig,
Invoker,
ADataTypeBuf,
BDataTypeBuf,
ck_tile::tuple<>,
AccDataType,
CDataType,
ALayout,
BLayout,
ck_tile::tuple<>,
CLayout>(a_m_k_dev_buf,
b_k_n_dev_buf,
c_m_n_dev_buf,
M,
N,
K,
stride_A,
stride_B,
stride_C,
kbatch,
n_warmup,
n_repeat,
persistent,
flush_cache,
rotating_count);
auto invoke = [&]<typename Cfg>() {
Comment thread
bidlekm marked this conversation as resolved.
Resolved
return invoke_gemm<Cfg,
Invoker,
ADataTypeBuf,
BDataTypeBuf,
ck_tile::tuple<>,
AccDataType,
CDataType,
ALayout,
BLayout,
ck_tile::tuple<>,
CLayout>(a_m_k_dev_buf,
b_k_n_dev_buf,
c_m_n_dev_buf,
M,
N,
K,
stride_A,
stride_B,
stride_C,
kbatch,
n_warmup,
n_repeat,
persistent,
flush_cache,
rotating_count);
};

float ave_time = ck_tile_example::dispatch_vector_size_fallback<GemmConfig,
ADataTypeBuf,
BDataTypeBuf,
CDataType,
ALayout,
BLayout,
CLayout>(
M, N, K, stride_A, stride_B, stride_C, kbatch, invoke);

c_m_n_dev_buf.FromDevice(c_m_n_dev_result.data());

Expand Down
23 changes: 15 additions & 8 deletions projects/composablekernel/example/ck_tile/03_gemm/universal_gemm_invoker.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -74,7 +74,8 @@ struct UniversalInvoker
VectorSize,
GemmConfig::DataCachePrefetchA,
GemmConfig::DataCachePrefetchB,
GemmConfig::Async>;
GemmConfig::Async,
false>; // LargeTensors

constexpr auto scheduler = GemmConfig::Scheduler;

Expand Down Expand Up @@ -103,7 +104,10 @@ struct UniversalInvoker
ck_tile::element_wise::PassThrough,
ck_tile::element_wise::PassThrough,
AComputeDataType,
BComputeDataType>;
BComputeDataType,
GemmConfig::FixedVectorSize,
GemmConfig::VectorSizeA,
GemmConfig::VectorSizeB>;

using GemmPipeline = typename PipelineTypeTraits<
GemmConfig::Pipeline>::template GemmPipeline<UniversalGemmProblem>;
Expand All @@ -127,8 +131,8 @@ struct UniversalInvoker
GemmConfig::K_Warp_Tile,
UniversalGemmProblem::TransposeC,
GemmConfig::NumWaveGroups,
false, /*FixedVectorSize_*/
1, /*VectorSizeC_*/
GemmConfig::FixedVectorSize,
Comment thread
bidlekm marked this conversation as resolved.
Resolved
GemmConfig::VectorSizeC,
1, /*BlockedXDLN_PerWarp_*/
GemmConfig::DoubleSmemBuffer, /*DoubleSmemBuffer*/
AComputeDataType,
Expand Down Expand Up @@ -294,7 +298,10 @@ struct UniversalInvoker
ck_tile::element_wise::PassThrough,
ck_tile::element_wise::PassThrough,
AComputeDataType,
BComputeDataType>;
BComputeDataType,
GemmConfig::FixedVectorSize,
GemmConfig::VectorSizeA,
GemmConfig::VectorSizeB>;

using GemmPipeline = typename PipelineTypeTraits<
GemmConfig::Pipeline>::template GemmPipeline<UniversalGemmProblem>;
Expand All @@ -317,9 +324,9 @@ struct UniversalInvoker
GemmConfig::K_Warp_Tile,
UniversalGemmProblem::TransposeC,
GemmConfig::NumWaveGroups,
false, /*FixedVectorSize_*/
1, /*VectorSizeC_*/
1, /*BlockedXDLN_PerWarp_*/
GemmConfig::FixedVectorSize,
GemmConfig::VectorSizeC,
1, /*BlockedXDLN_PerWarp_*/
GemmConfig::DoubleSmemBuffer,
AComputeDataType,
BComputeDataType>>;
Expand Down
63 changes: 37 additions & 26 deletions projects/composablekernel/example/ck_tile/16_batched_gemm/batched_gemm.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -45,9 +45,9 @@ float batched_gemm(const ck_tile::BatchedGemmHostArgs& args, const ck_tile::stre
#endif
constexpr bool DoubleSmemBuffer = GemmConfig::DoubleSmemBuffer;

constexpr bool kPadM = false;
constexpr bool kPadN = false;
constexpr bool kPadK = false;
constexpr bool kPadM = GemmConfig::kPadM;
constexpr bool kPadN = GemmConfig::kPadN;
constexpr bool kPadK = GemmConfig::kPadK;

constexpr bool TransposeC = false;

Expand All @@ -72,33 +72,44 @@ float batched_gemm(const ck_tile::BatchedGemmHostArgs& args, const ck_tile::stre
TransposeC>;
constexpr auto scheduler = GemmConfig::Scheduler;

using UniversalGemmProblem = ck_tile::UniversalGemmPipelineProblem<ADataType,
BDataType,
AccDataType,
GemmShape,
GemmUniversalTraits,
scheduler>;
using UniversalGemmProblem =
ck_tile::UniversalGemmPipelineProblem<ADataType,
BDataType,
AccDataType,
GemmShape,
GemmUniversalTraits,
scheduler,
ck_tile::element_wise::PassThrough,
ck_tile::element_wise::PassThrough,
ADataType,
BDataType,
GemmConfig::FixedVectorSize,
GemmConfig::VectorSizeA,
GemmConfig::VectorSizeB>;

using GemmPipeline = typename PipelineTypeTraits<GemmConfig::Pipeline>::template GemmPipeline<
UniversalGemmProblem>;

using GemmEpilogue = ck_tile::CShuffleEpilogue<
ck_tile::CShuffleEpilogueProblem<ADataType,
BDataType,
DsDataType,
AccDataType,
CDataType,
DsLayout,
CLayout,
CDEElementWise,
TilePartitioner::MPerBlock,
TilePartitioner::NPerBlock,
M_Warp,
N_Warp,
M_Warp_Tile,
N_Warp_Tile,
K_Warp_Tile,
UniversalGemmProblem::TransposeC>>;
using GemmEpilogue =
ck_tile::CShuffleEpilogue<ck_tile::CShuffleEpilogueProblem<ADataType,
BDataType,
DsDataType,
AccDataType,
CDataType,
DsLayout,
CLayout,
CDEElementWise,
TilePartitioner::MPerBlock,
TilePartitioner::NPerBlock,
M_Warp,
N_Warp,
M_Warp_Tile,
N_Warp_Tile,
K_Warp_Tile,
UniversalGemmProblem::TransposeC,
1, /*kNumWaveGroups_*/
GemmConfig::FixedVectorSize,
GemmConfig::VectorSizeC>>;

using Kernel = ck_tile::BatchedGemmKernel<TilePartitioner, GemmPipeline, GemmEpilogue>;
auto kargs = Kernel::MakeKernelArgs(args);
Expand Down
27 changes: 23 additions & 4 deletions projects/composablekernel/example/ck_tile/16_batched_gemm/batched_gemm.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -10,8 +10,25 @@
#include "ck_tile/ops/gemm/kernel/batched_gemm_kernel.hpp"
#include "ck_tile/ops/elementwise/unary_element_wise_operation.hpp"
#include "ck_tile/utility/json_dump.hpp"
#include "gemm_common/vector_size_fallback_dispatch.hpp"

struct GemmConfigMemory
using ck_tile_example::GemmConfigVectorSizeFallback;

struct GemmConfigBase
{
static constexpr bool kPadM = false;
static constexpr bool kPadN = false;
static constexpr bool kPadK = false;

static constexpr bool FixedVectorSize = false;
static constexpr ck_tile::index_t VectorSizeA = 1;
static constexpr ck_tile::index_t VectorSizeB = 1;
static constexpr ck_tile::index_t VectorSizeC = 1;

static constexpr bool EnableSmallerVectorLoadFallback = false;
};

struct GemmConfigMemory : public GemmConfigBase
{
// Memory friendly for Interwave scheduler
static constexpr ck_tile::index_t M_Tile = 128;
Expand All @@ -31,7 +48,7 @@ struct GemmConfigMemory
static constexpr auto Scheduler = ck_tile::GemmPipelineScheduler::Interwave;
};

struct GemmConfigV3
struct GemmConfigV3 : public GemmConfigBase
{
// Compute friendly for Intrawave scheduler
static constexpr ck_tile::index_t M_Tile = 256;
Expand All @@ -51,7 +68,7 @@ struct GemmConfigV3
static constexpr auto Scheduler = ck_tile::GemmPipelineScheduler::Intrawave;
};

struct GemmConfigV4
struct GemmConfigV4 : public GemmConfigBase
{
// Compute friendly for Intrawave scheduler
// Using the ping pong reader in the lds level
Expand All @@ -72,7 +89,7 @@ struct GemmConfigV4
static constexpr auto Scheduler = ck_tile::GemmPipelineScheduler::Intrawave;
};

struct GemmConfigV3_Wmma
struct GemmConfigV3_Wmma : public GemmConfigBase
{
// Compute friendly for Intrawave scheduler
static constexpr ck_tile::index_t M_Tile = 128;
Expand All @@ -90,6 +107,8 @@ struct GemmConfigV3_Wmma
static constexpr bool DoubleSmemBuffer = false;
static constexpr ck_tile::GemmPipeline Pipeline = ck_tile::GemmPipeline::COMPUTE_V3;
static constexpr auto Scheduler = ck_tile::GemmPipelineScheduler::Intrawave;

static constexpr bool EnableSmallerVectorLoadFallback = true;
};

template <ck_tile::GemmPipeline PipelineId>
Expand Down
Loading
Loading