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 8 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
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
37 changes: 37 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 @@ -41,6 +41,39 @@ 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;

// Enable for gfx1250 RCR
static constexpr bool EnableKPadFallback = false;
static constexpr bool EnableMNPadFallback = false;
Comment thread
bidlekm marked this conversation as resolved.
Resolved
};

// A,B vector sizes must be multiples of K
template <typename GemmConfig,
ck_tile::index_t VectorSizeA_,
ck_tile::index_t VectorSizeB_,
ck_tile::index_t VectorSizeC_>
struct GemmConfigFixedVectorSize : public GemmConfig
{
// Split-K partitions are aligned to K_Warp_Tile: preserve split-K remainder
static_assert(GemmConfig::K_Warp_Tile % VectorSizeA_ == 0 &&
GemmConfig::K_Warp_Tile % VectorSizeB_ == 0,
"A/B vector width must divide K_Warp_Tile");

// Enable K/M/N padding
static constexpr bool kPadM = true;
static constexpr bool kPadN = true;
static constexpr bool kPadK = true;

static constexpr bool FixedVectorSize = true;
static constexpr ck_tile::index_t VectorSizeA = VectorSizeA_;
static constexpr ck_tile::index_t VectorSizeB = VectorSizeB_;
static constexpr ck_tile::index_t VectorSizeC = VectorSizeC_;
};

template <typename PrecType>
Expand Down Expand Up @@ -176,6 +209,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 +227,9 @@ struct GemmConfigComputeV3_WMMA : public GemmConfigBase
static constexpr ck_tile::GemmPipeline Pipeline = ck_tile::GemmPipeline::COMPUTE_V3;

static constexpr int kBlockPerCu = 2;

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

template <typename PrecType>
Expand Down
117 changes: 90 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 @@ -222,7 +222,7 @@ int run_gemm_example_with_layouts(ck_tile::ArgParser& arg_parser,
const BLayout b_layout = BLayout{},
[[maybe_unused]] const CLayout c_layout = CLayout{})
{
// Use GemmTypeConfig to get actual data types for tensor operations

Comment thread
bidlekm marked this conversation as resolved.
Resolved
using TypeConfig = GemmTypeConfig<ADataType_, BDataType_, CDataType_>;
using ADataTypeBuf = typename TypeConfig::ADataType;
using BDataTypeBuf = typename TypeConfig::BDataType;
Expand Down Expand Up @@ -304,7 +304,7 @@ int run_gemm_example_with_layouts(ck_tile::ArgParser& arg_parser,
return ck_tile::shuffle_b_v0<GemmConfig>(b_k_n);
}
}();
// shuffled buffer B for device implementation

if constexpr(std::is_same_v<BDataTypeBuf, ck_tile::pk_int4_t>)
{
ck_tile::permute_vectors_i4x4_b(b_shuffle_host);
Expand Down Expand Up @@ -347,31 +347,94 @@ 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);
constexpr bool APacksK = std::is_same_v<ALayout, ck_tile::tensor_layout::gemm::RowMajor>;
constexpr bool BPacksN = std::is_same_v<BLayout, ck_tile::tensor_layout::gemm::RowMajor>;
constexpr bool CPacksN = std::is_same_v<CLayout, ck_tile::tensor_layout::gemm::RowMajor>;
Comment thread
bidlekm marked this conversation as resolved.
Resolved

constexpr bool k_pad_fallback_eligible =
GemmConfig::EnableKPadFallback && !GemmConfig::FixedVectorSize &&
std::is_same_v<ADataTypeBuf, BDataTypeBuf> &&
ck_tile::numeric_traits<ADataTypeBuf>::PackedSize == 1 &&
ck_tile::numeric_traits<BDataTypeBuf>::PackedSize == 1;

constexpr bool mn_pad_fallback_eligible =
GemmConfig::EnableMNPadFallback && !GemmConfig::FixedVectorSize &&
std::is_same_v<ADataTypeBuf, BDataTypeBuf> &&
ck_tile::numeric_traits<ADataTypeBuf>::PackedSize == 1 &&
ck_tile::numeric_traits<BDataTypeBuf>::PackedSize == 1 &&
ck_tile::numeric_traits<CDataType>::PackedSize == 1;

constexpr ck_tile::index_t ABVectorSize = 16 / sizeof(ADataTypeBuf);
constexpr ck_tile::index_t CVectorSize = 16 / sizeof(CDataType);
Comment thread
bidlekm marked this conversation as resolved.
Resolved

const bool needsScalarAB =
k_pad_fallback_eligible && (((APacksK ? K : M) % ABVectorSize != 0) ||
((BPacksN ? N : K) % ABVectorSize != 0));
const bool needsScalarC =
mn_pad_fallback_eligible && ((CPacksN ? N : M) % CVectorSize != 0);
Comment thread
bidlekm marked this conversation as resolved.
Resolved

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;
Comment thread
bidlekm marked this conversation as resolved.
Resolved
if constexpr(k_pad_fallback_eligible && mn_pad_fallback_eligible)
{
if(needsScalarAB && needsScalarC)
ave_time = invoke.template operator()<GemmConfigFixedVectorSize<GemmConfig, 1, 1, 1>>();
else if(needsScalarAB)
ave_time = invoke.template
operator()<GemmConfigFixedVectorSize<GemmConfig, 1, 1, CVectorSize>>();
else if(needsScalarC)
ave_time = invoke.template operator()<
GemmConfigFixedVectorSize<GemmConfig, ABVectorSize, ABVectorSize, 1>>();
else
ave_time = invoke.template operator()<GemmConfig>();
}
Comment thread
bidlekm marked this conversation as resolved.
Resolved
else if constexpr(k_pad_fallback_eligible)
{
if(needsScalarAB)
ave_time = invoke.template
operator()<GemmConfigFixedVectorSize<GemmConfig, 1, 1, CVectorSize>>();
else
ave_time = invoke.template operator()<GemmConfig>();
}
else if constexpr(mn_pad_fallback_eligible)
{
if(needsScalarC)
ave_time = invoke.template operator()<
GemmConfigFixedVectorSize<GemmConfig, ABVectorSize, ABVectorSize, 1>>();
else
ave_time = invoke.template operator()<GemmConfig>();
}
Comment thread
bidlekm marked this conversation as resolved.
Resolved
else
{
ave_time = invoke.template operator()<GemmConfig>();
}

c_m_n_dev_buf.FromDevice(c_m_n_dev_result.data());

Expand Down
12 changes: 8 additions & 4 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
21 changes: 16 additions & 5 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,10 @@ 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;
using VectorSizeDefaults = GemmConfigVectorSizeDefaults<GemmConfig>;
constexpr bool kPadM = VectorSizeDefaults::kPadM;
constexpr bool kPadN = VectorSizeDefaults::kPadN;
constexpr bool kPadK = VectorSizeDefaults::kPadK;

constexpr bool TransposeC = false;

Expand Down Expand Up @@ -77,7 +78,14 @@ float batched_gemm(const ck_tile::BatchedGemmHostArgs& args, const ck_tile::stre
AccDataType,
GemmShape,
GemmUniversalTraits,
scheduler>;
scheduler,
ck_tile::element_wise::PassThrough,
ck_tile::element_wise::PassThrough,
ADataType,
BDataType,
VectorSizeDefaults::FixedVectorSize,
VectorSizeDefaults::VectorSizeA,
VectorSizeDefaults::VectorSizeB>;

using GemmPipeline = typename PipelineTypeTraits<GemmConfig::Pipeline>::template GemmPipeline<
UniversalGemmProblem>;
Expand All @@ -98,7 +106,10 @@ float batched_gemm(const ck_tile::BatchedGemmHostArgs& args, const ck_tile::stre
M_Warp_Tile,
N_Warp_Tile,
K_Warp_Tile,
UniversalGemmProblem::TransposeC>>;
UniversalGemmProblem::TransposeC,
1, /*kNumWaveGroups_*/
VectorSizeDefaults::FixedVectorSize,
VectorSizeDefaults::VectorSizeC>>;

using Kernel = ck_tile::BatchedGemmKernel<TilePartitioner, GemmPipeline, GemmEpilogue>;
auto kargs = Kernel::MakeKernelArgs(args);
Expand Down
40 changes: 40 additions & 0 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 @@ -11,6 +11,46 @@
#include "ck_tile/ops/elementwise/unary_element_wise_operation.hpp"
#include "ck_tile/utility/json_dump.hpp"

template <typename GemmConfig,
ck_tile::index_t VectorSizeA_,
ck_tile::index_t VectorSizeB_,
ck_tile::index_t VectorSizeC_>
struct GemmConfigFixedVectorSize : public GemmConfig
Comment thread
bidlekm marked this conversation as resolved.
Resolved
{
static_assert(GemmConfig::K_Warp_Tile % VectorSizeA_ == 0 &&
GemmConfig::K_Warp_Tile % VectorSizeB_ == 0,
"A/B vector width must divide K_Warp_Tile");
Comment thread
bidlekm marked this conversation as resolved.
Resolved

static constexpr bool kPadM = true;
static constexpr bool kPadN = true;
static constexpr bool kPadK = true;

static constexpr bool FixedVectorSize = true;
static constexpr ck_tile::index_t VectorSizeA = VectorSizeA_;
static constexpr ck_tile::index_t VectorSizeB = VectorSizeB_;
static constexpr ck_tile::index_t VectorSizeC = VectorSizeC_;
};

struct GemmConfigNoFixedVectorSize
{
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;
};

template <typename T>
using has_fixed_vector_size_t = decltype(T::FixedVectorSize);

template <typename GemmConfig>
using GemmConfigVectorSizeDefaults =
std::conditional_t<ck_tile::is_detected<has_fixed_vector_size_t, GemmConfig>::value,
Comment thread
bidlekm marked this conversation as resolved.
Resolved
GemmConfig,
GemmConfigNoFixedVectorSize>;

struct GemmConfigMemory
{
// Memory friendly for Interwave scheduler
Expand Down
Loading
Loading