ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-11 17:51:40 +00:00
Code Issues Packages Projects Releases Wiki Activity

Clean up where skiplds params is exposed

Browse Source
This commit is contained in:
Damien Lejeune
2025-07-30 20:07:08 +00:00
parent c1548bf08c
commit 96b26b2a3c
5 changed files with 13 additions and 30 deletions
Download Patch File Download Diff File Expand all files Collapse all files

12
include/ck_tile/ops/gemm/kernel/batched_gemm_kernel.hpp
View File

@@ -31,9 +31,7 @@ struct BatchedGemmHostArgs : public ck_tile::UniversalGemmHostArgs<>
ck_tile::index_t batch_stride_A_,
ck_tile::index_t batch_stride_B_,
ck_tile::index_t batch_stride_C_,
ck_tile::index_t batch_count_,
bool skip_a_lds_,
bool skip_b_lds_)
ck_tile::index_t batch_count_)
: UniversalGemmHostArgs<>({a_ptr_},
{b_ptr_},
{/*ds_ptr*/},
@@ -45,9 +43,7 @@ struct BatchedGemmHostArgs : public ck_tile::UniversalGemmHostArgs<>
{stride_A_},
{stride_B_},
{/*stride_Ds_*/},
stride_C_,
skip_a_lds_,
skip_b_lds_),
stride_C_),
batch_stride_A(batch_stride_A_),
batch_stride_B(batch_stride_B_),
batch_stride_E(batch_stride_C_),
@@ -144,9 +140,7 @@ struct BatchedGemmKernel
hostArgs.stride_Bs,
hostArgs.stride_Ds,
hostArgs.stride_E,
hostArgs.k_batch,
hostArgs.skip_a_lds,
hostArgs.skip_b_lds},
hostArgs.k_batch},
hostArgs.batch_stride_A,
hostArgs.batch_stride_B,
hostArgs.batch_stride_E,

4
include/ck_tile/ops/gemm/kernel/gemm_kernel.hpp
View File

@@ -158,9 +158,7 @@ struct GemmKernel
{hostArgs.stride_A},
{hostArgs.stride_B},
{/*hostArgs.stride_Ds*/},
hostArgs.stride_E,
hostArgs.skip_a_lds,
hostArgs.skip_b_lds));
hostArgs.stride_E));
}
CK_TILE_HOST static auto

18
include/ck_tile/ops/gemm/kernel/universal_gemm_kernel.hpp
View File

@@ -41,9 +41,7 @@ struct UniversalGemmHostArgs
const std::array<index_t, NumATensor>& stride_As_,
const std::array<index_t, NumBTensor>& stride_Bs_,
const std::array<index_t, NumDTensor>& stride_Ds_,
index_t stride_E_,
bool skip_a_lds_,
bool skip_b_lds_)
index_t stride_E_)
: as_ptr(as_ptr_),
bs_ptr(bs_ptr_),
ds_ptr(ds_ptr_),
@@ -55,9 +53,7 @@ struct UniversalGemmHostArgs
stride_Bs(stride_Bs_),
stride_Ds(stride_Ds_),
stride_E(stride_E_),
k_batch(k_batch_),
skip_a_lds(skip_a_lds_),
skip_b_lds(skip_b_lds_)
k_batch(k_batch_)
{
}
@@ -82,8 +78,6 @@ struct UniversalGemmHostArgs
};
index_t k_batch;
bool skip_a_lds;
bool skip_b_lds;
};
/// @brief The GEMM kernel device arguments.
@@ -117,10 +111,6 @@ struct UniversalGemmKernelArgs
/// (in memory) of E tensor.
index_t stride_E;
index_t k_batch;
/// @brief Flag to skip loading A tensor tile into LDS.
bool skip_a_lds;
/// @brief Flag to skip loading B tensor tile into LDS.
bool skip_b_lds;
};
/// @brief The Universal GEMM kernel template.
@@ -292,9 +282,7 @@ struct UniversalGemmKernel
hostArgs.stride_Bs,
hostArgs.stride_Ds,
hostArgs.stride_E,
hostArgs.k_batch,
hostArgs.skip_a_lds,
hostArgs.skip_b_lds};
hostArgs.k_batch};
}
CK_TILE_HOST_DEVICE static constexpr index_t GetSmemSize()

3
include/ck_tile/ops/gemm/pipeline/tile_gemm_traits.hpp
View File

@@ -67,6 +67,9 @@ struct TileGemmUniversalTraits
static constexpr bool UsePersistentKernel = UsePersistentKernel_;
static constexpr index_t NumWaveGroups = NumWaveGroups_;
static constexpr bool Preshuffle = Preshuffle_;
static constexpr bool kSkipALds = SkipALds_;
static constexpr bool kSkipBLds = SkipBLds_;
};
template <bool kPadM_,

6
test/ck_tile/batched_gemm/test_batched_gemm_util.hpp
View File

@@ -70,12 +70,12 @@ class TestCkTileBatchedGemm : public ::testing::Test
kPadN,
kPadK,
DoubleSmemBuffer,
kSkipALds,
kSkipBLds,
ALayout,
BLayout,
CLayout,
TransposeC>;
TransposeC,
kSkipALds,
kSkipBLds>;
using GemmPipelineProblem =
ck_tile::GemmPipelineProblem<ADataType, BDataType, AccDataType, GemmShape, Traits>;