Skip to content

Commit 5d78f3f

Browse files
henrylhtsangfacebook-github-bot
authored andcommitted
Switch alignemtn to 8 for cutlass 4 upgrade (#2491)
Summary: As titled. static_assert(detail::is_aligned<ElementC_, AlignmentC, ElementD_, AlignmentD>(), ^ Differential Revision: D77745963
1 parent 2a24a00 commit 5d78f3f

File tree

1 file changed

+4
-4
lines changed

1 file changed

+4
-4
lines changed

torchao/csrc/cuda/activation24/sparse_gemm.cu

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -95,10 +95,10 @@ struct SparseRowwiseKernel<cutlass::float_e4m3_t> {
9595
float,
9696
ElementOut,
9797
cutlass::layout::RowMajor,
98-
1,
98+
8,
9999
ElementOut,
100100
cutlass::layout::RowMajor,
101-
1,
101+
8,
102102
cutlass::epilogue::TmaWarpSpecializedCooperative,
103103
EpilogueEVT>::CollectiveOp;
104104

@@ -172,10 +172,10 @@ struct SparseRowwiseKernel<cutlass::bfloat16_t> {
172172
float,
173173
ElementOut,
174174
cutlass::layout::RowMajor,
175-
1,
175+
8,
176176
ElementOut,
177177
cutlass::layout::RowMajor,
178-
1,
178+
8,
179179
cutlass::epilogue::TmaWarpSpecializedCooperative,
180180
EpilogueEVT>::CollectiveOp;
181181

0 commit comments

Comments
 (0)