Skip to content

Commit 09aa536

Browse files
committed
refactor: remove dead code from gemm universal kernel
1 parent b06f6f5 commit 09aa536

File tree

2 files changed

+2
-5
lines changed

2 files changed

+2
-5
lines changed

include/ck_tile/ops/gemm/kernel/universal_gemm_kernel.hpp

Lines changed: 1 addition & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -1035,7 +1035,6 @@ struct UniversalGemmKernel
10351035
* @param block_idx_n The GEMM's output N dimension tile index processed by this workgroup.
10361036
*
10371037
*/
1038-
template <bool UseDefaultScheduler = true>
10391038
CK_TILE_DEVICE static void RunGemm(const std::array<const ADataType*, NumATensor>& as_ptr,
10401039
const std::array<const BDataType*, NumBTensor>& bs_ptr,
10411040
const std::array<const void*, NumDTensor>& ds_ptr,
@@ -1161,9 +1160,7 @@ struct UniversalGemmKernel
11611160
// allocate LDS
11621161
__shared__ char smem_ptr[GetSmemSize()];
11631162

1164-
constexpr auto scheduler_type =
1165-
GemmPipeline::DoubleSmemBuffer || (GemmPipeline::NumWaveGroups == 1);
1166-
RunGemm<scheduler_type>(
1163+
RunGemm(
11671164
as_ptr, bs_ptr, kargs.ds_ptr, e_ptr, smem_ptr, kargs, splitk_batch_offset, i_m, i_n);
11681165
}
11691166

include/ck_tile/ops/gemm/pipeline/gemm_pipeline_problem.hpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -80,7 +80,7 @@ struct GemmPipelineProblemBase
8080
static constexpr bool kPadK = Traits::kPadK;
8181

8282
static constexpr bool DoubleSmemBuffer = Traits::DoubleSmemBuffer;
83-
static constexpr auto Scheduler = GemmPipelineScheduler::Default;
83+
static constexpr auto Scheduler = GemmPipelineScheduler::Intrawave;
8484
static constexpr index_t VectorLoadSize = Traits::_VectorSize;
8585

8686
// In the base situation, the Preshuffle setting should be false.

0 commit comments

Comments
 (0)