ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-06-29 19:28:33 +00:00
Code Issues Packages Projects Releases Wiki Activity

Add asm for no-loop v3_128x128x128

Browse Source
This commit is contained in:
OscarXu
2025-05-06 10:47:29 +08:00
parent efbb85be2a
commit e7fe8587f6
2 changed files with 24 additions and 24 deletions
Download Patch File Download Diff File Expand all files Collapse all files

BIN
example/65_gemm_multiply_multiply/hsa/moe_bs_stage2_v3_128x128x128.co
View File

Binary file not shown.

48
include/ck/tensor_operation/gpu/device/impl/device_moe_gemm_blockscale.hpp
View File

@@ -491,30 +491,30 @@ struct DeviceMoeGemmBlockScale
RunKernel(kernel);
}
}
// else if constexpr(BlkGemmPipelineVer == BlockGemmPipelineVersion::v2 ||
// BlkGemmPipelineVer == BlockGemmPipelineVersion::v3)
// {
// if(GridwiseGemm::CalculateKBlockLoopTailNum(K_split) == TailNumber::Odd)
// {
// const auto kernel = kernel_moe_gemm_2lds<GridwiseGemm,
// false,
// MemoryDataOp,
// minimum_occupancy,
// IsInputGemm,
// TailNumber::Odd>;
// RunKernel(kernel);
// }
// else
// {
// const auto kernel = kernel_moe_gemm_2lds<GridwiseGemm,
// false,
// MemoryDataOp,
// minimum_occupancy,
// IsInputGemm,
// TailNumber::Even>;
// RunKernel(kernel);
// }
// }
else if constexpr(BlkGemmPipelineVer == BlockGemmPipelineVersion::v2 ||
BlkGemmPipelineVer == BlockGemmPipelineVersion::v3)
{
if(GridwiseGemm::CalculateKBlockLoopTailNum(K_split) == TailNumber::Odd)
{
const auto kernel = kernel_moe_gemm_2lds<GridwiseGemm,
false,
MemoryDataOp,
minimum_occupancy,
IsInputGemm,
TailNumber::Odd>;
RunKernel(kernel);
}
else
{
const auto kernel = kernel_moe_gemm_2lds<GridwiseGemm,
false,
MemoryDataOp,
minimum_occupancy,
IsInputGemm,
TailNumber::Even>;
RunKernel(kernel);
}
}
}
#endif
#endif