[GH-ISSUE #291] [question] sm100_fp8_gemm_1d1d_impl: UTCCP transposer for scale factor is strange #147

New issue

Closed

opened 2026-10-26 21:13:40 +00:00 by bsakaguchi · 0 comments

bsakaguchi commented

2026-10-26 21:13:40 +00:00

Owner

Originally created by @songna2026 on GitHub (Mar 7, 2026).
Original GitHub issue: https://github.com/deepseek-ai/DeepGEMM/issues/291

Hello,

I see in kernel sm100_fp8_gemm_1d1d_impl, there is a warp 2 to do rearrangement for scale factor. I did not see such instruction in cutlass。
I think tcgen05.cp can correctly copy scale factor from SMEM to TMEM with SMEM layout as below:
4 means that four scale factor packed as unit32 in dim K.

tcgen05.cp.cta_group::1.32x128b.warpx4 will copy above data from SMEM to TMEM with below TMEM layout：

Is there any other deeper meaning behind adding warp 2?

Looking forward for your reply. Thanks a lot!

Originally created by @songna2026 on GitHub (Mar 7, 2026). Original GitHub issue: https://github.com/deepseek-ai/DeepGEMM/issues/291 Hello, I see in kernel sm100_fp8_gemm_1d1d_impl, there is a warp 2 to do rearrangement for scale factor. I did not see such instruction in cutlass。 I think tcgen05.cp can correctly copy scale factor from SMEM to TMEM with SMEM layout as below: 4 means that four scale factor packed as unit32 in dim K. <img width="344" height="568" alt="Image" src="https://github.com/user-attachments/assets/95523f81-c298-4fc2-8d68-5d928cb2c975" /> tcgen05.cp.cta_group::1.32x128b.warpx4 will copy above data from SMEM to TMEM with below TMEM layout： <img width="1570" height="934" alt="Image" src="https://github.com/user-attachments/assets/e4d43909-ba99-4ced-a1c2-9ab43413161f" /> Is there any other deeper meaning behind adding warp 2? Looking forward for your reply. Thanks a lot!

bsakaguchi closed this issue

2026-10-26 21:13:40 +00:00

bsakaguchi referenced this issue

2026-10-26 21:15:31 +00:00

[PR #149] [MERGED] Fix SM90 GEMM #265