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 all 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
2 changes: 1 addition & 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 @@ -15,7 +15,7 @@ struct GemmConfigTwoStage_Wmma : public GemmConfigComputeV3_WMMA<PrecType_>
{
using WorkspaceType = ck_tile::remove_cvref_t<WorkspaceType_>;

static constexpr bool EnableKPadFallback = false; // TODO invoker not supported yet
static constexpr bool EnableSmallerVectorLoadFallback = false; // TODO invoker not supported yet
};

struct SplitKTwoStageInvoker
Expand Down
15 changes: 9 additions & 6 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;

// Max. vectorized global memory access in bytes.
static constexpr ck_tile::index_t kMaxVectorBytes = 16;

Expand Down Expand Up @@ -56,11 +59,11 @@ struct GemmConfigBase
static constexpr ck_tile::index_t VectorSizeB = 1;
static constexpr ck_tile::index_t VectorSizeC = 1;

// Enable for RCR
static constexpr bool EnableKPadFallback = false;
static constexpr bool EnableSmallerVectorLoadFallback = false;
};

// A,B vector sizes must divide K_Warp_Tile.
// A,B vector sizes must divide K_Warp_Tile. Used directly by test_gemm_unaligned_k.cpp; the
// general run_gemm_example.inc dispatch path uses GemmConfigVectorSizeFallback instead.
template <typename GemmConfig,
ck_tile::index_t VectorSizeA_,
ck_tile::index_t VectorSizeB_,
Expand Down Expand Up @@ -236,7 +239,7 @@ struct GemmConfigComputeV3_WMMA : public GemmConfigBase

static constexpr int kBlockPerCu = 2;

static constexpr bool EnableKPadFallback = true;
static constexpr bool EnableSmallerVectorLoadFallback = true;
};

template <typename PrecType>
Expand All @@ -245,8 +248,8 @@ struct GemmConfigComputeV3_WMMA_ClusterLaunch : public GemmConfigComputeV3_WMMA<
static constexpr ck_tile::index_t kClusterSizeM = 2;
static constexpr ck_tile::index_t kClusterSizeN = 2;

// KPad fallback not validated on the cluster-launch path yet, disable for now.
static constexpr bool EnableKPadFallback = false;
// Vector-size fallback not validated on the cluster-launch path yet, disable for now.
static constexpr bool EnableSmallerVectorLoadFallback = false;
};

template <typename PrecType>
Expand Down
37 changes: 10 additions & 27 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"

#include <type_traits>

Expand Down Expand Up @@ -349,19 +350,7 @@ int run_gemm_example_with_layouts(ck_tile::ArgParser& arg_parser,
c_m_n_dev_buf.SetZero();
c_m_n_dev_result.SetZero();

// Enable fallback: RCR, PackedSize=1
constexpr bool k_pad_fallback_eligible =
GemmConfig::EnableKPadFallback && !GemmConfig::FixedVectorSize &&
std::is_same_v<ADataTypeBuf, BDataTypeBuf> &&
std::is_same_v<ALayout, ck_tile::tensor_layout::gemm::RowMajor> &&
std::is_same_v<BLayout, ck_tile::tensor_layout::gemm::ColumnMajor> &&
std::is_same_v<CLayout, ck_tile::tensor_layout::gemm::RowMajor> &&
ck_tile::numeric_traits<ADataTypeBuf>::PackedSize == 1 &&
ck_tile::numeric_traits<BDataTypeBuf>::PackedSize == 1;

// Invoke GEMM
auto run_invoke = [&](auto cfg_tag) {
using Cfg = typename decltype(cfg_tag)::type;
auto invoke = [&]<typename Cfg>() {
Comment thread
bidlekm marked this conversation as resolved.
Resolved
return invoke_gemm<Cfg,
Invoker,
ADataTypeBuf,
Expand Down Expand Up @@ -389,20 +378,14 @@ int run_gemm_example_with_layouts(ck_tile::ArgParser& arg_parser,
rotating_count);
};

float ave_time;
if constexpr(k_pad_fallback_eligible)
{
using KPadConfig = GemmConfigFixedVectorSize<GemmConfig, 1, 1, kMaxVectorElems<CDataType>>;

ave_time = (K % kMaxVectorElems<ADataTypeBuf> != 0)
? run_invoke(std::type_identity<KPadConfig>{}) // unaligned K
: run_invoke(std::type_identity<GemmConfig>{}); // aligned K
}
else
{
// Do not use KPad fallback
ave_time = run_invoke(std::type_identity<GemmConfig>{});
}
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
14 changes: 9 additions & 5 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 @@ -297,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 @@ -320,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