leavelet/cutlass
1
0
Fork 0
Code Issues Pull Requests Actions 1 Packages Projects Releases Wiki Activity

[Bug Fix]Set NumSplitsM to 1 when TileShapeM < 128 in sm90 fp8 blockwise scaling CollectiveMma (#2965)

Browse Source 
* Fix NumSplitsM when TileShapeM < 128.

* Use cute::conditional_t to replace std::conditional_t.
This commit is contained in:
Qi Yuhang
2026-01-23 15:56:52 +08:00
committed by GitHub
parent 0edaa6e47d
commit 2fafefb7b9
1 changed files with 2 additions and 1 deletions
Download Patch File Download Diff File Expand all files Collapse all files

3
include/cutlass/gemm/collective/sm90_mma_array_tma_gmma_ss_warpspecialized_fp8_blockwise_scaling.hpp
View File

@@ -212,7 +212,8 @@ struct CollectiveMma<
static_assert(cute::is_same_v<ElementAccumulator, ElementBlockScale>,
"ElementAccumulator and ElementBlockScale should be same datatype");
using NumSplitsM = cute::C<get<0>(TileShape_{}) / 128>;
// For TileShapeM < 128, NumSplitsM should be 1
using NumSplitsM = cute::conditional_t<get<0>(TileShape_{}) < _128{}, _1, cute::C<get<0>(TileShape_{}) / 128>>;
static_assert(NumSplitsM{} == 1 || NumSplitsM{} == 2);
struct SharedStorage {