[QST][CUTEDSL] Address Misalignment in FP8 Gemm

[Willie-Qu]

I am seeking help from Cutlass community. I encountered misaligned address (16 bytes requested) when load matrix from smem to register. I found out the problem started at "thr_copy_ldmatrix_A.partition_S(sA)".
Does anyone know how to debug and solve this issue?

--------------------------------
DEBUG: sA: 
tensor<ptr<i8, smem, align<1024>> o ((64,1),(8,8),(1,4)):((1,0),(64,512),(0,4096))>

DEBUG: thr_copy_ldmatrix_A: 
Tiled Copy
  Tiler MN:        (32:1,32:1)
  TV Layout tiled: ((4,8,2,2),((4,2,2),(1,1))):((128,1,16,0),((32,8,512),(0,0)))
Copy Atom
  ThrID:           32:1
  TV Layout Src:   ((2,2,4,2),16):((16,128,32,0),1)
  TV Layout Dst:   ((4,8),(1,2,2,2)):((32,1),(1,16,8,128))
  Value type:      i8
--------------------------------
tCsA_copy_view = thr_copy_ldmatrix_A.partition_S(sA)

DEBUG: tCsA_copy_view: 
tensor<ptr<i8, smem, align<8>> o (((8,2),2),2,2,(1,4)):(((1,128),1024),32,2048,(0,4096))>
--------------------------------

[Willie-Qu]

code to reproduce the problem. Just a simple FP8 gemm demo modified from cutlass/examples cute
File: reproduce_code.py

[m96-chan]

Additional Findings from SM120 (RTX 5090) Testing

We've been trying to get FP8 GEMM working on SM120 (Blackwell GeForce / RTX 5090) and encountered the same issue. Here are our additional observations that may help with diagnosis:

Environment

• GPU: NVIDIA GeForce RTX 5090 (SM 120)
• CUDA: 13.1
• CUTLASS: 4.3.3 (git submodule)
• Drivers tested: 576.40 and 576.52 (both reproduce the issue)
• OS: Windows 11

Reproduction

We reproduced the issue using CUTLASS example 87a (FP8 Blackwell GEMM):

nvcc -std=c++17 -O3 -arch=sm_120a \
  -I cutlass/include -I cutlass/tools/util/include \
  cutlass/examples/87a_blackwell_fp8_gemm/87a_blackwell_fp8_gemm.cu \
  -o example87a.exe

Result: LDSM address misaligned: requested 16 byte alignment but address X has 8 bytes alignment

Alternative Approaches Tested

We attempted several fallback strategies:

Approach	Kernel	Result
SM120 native (example 87a)	KernelScheduleSm120	FAIL: misaligned address (this bug)
SM100 fallback (example 81)	KernelScheduleSm100Blockwise	FAIL: initialize failed: 7 (kErrorInternal)
SM90 fallback (example 54)	KernelTmaWarpSpecializedCooperative	FAIL: initialize failed: 7 (kErrorInternal)
CUTLASS 2.x (SM80 ArchTag)	TF32/FP16/BF16 standard GEMM	WORKS

Key Observations

1. SM100/SM90 kernels do not work on SM120: Even though SM120 is "Blackwell architecture", the SM100 (datacenter) and SM90 (Hopper) CUTLASS 3.x/4.x kernels fail at initialize() with status 7 (kErrorInternal). This suggests architectural differences in tensor core generations between B100/B200 (datacenter) and RTX 5090 (consumer).

2. CUTLASS 2.x works fine on SM120: Standard GEMM kernels using cutlass::arch::Sm80 ArchTag (TF32, FP16, BF16) work correctly on SM120. This confirms the GPU itself is functional.

3. FP8 is currently blocked on SM120: There is no working FP8 GEMM path for RTX 5090:

   • Native SM120 FP8: blocked by this alignment bug
   • SM100 FP8: incompatible with SM120
   • SM90 FP8: incompatible with SM120

4. Driver version does not affect the issue: Tested on both 576.40 and 576.52 with identical results.

Workaround (non-FP8)

For those needing GEMM on SM120 while waiting for a fix, CUTLASS 2.x kernels with SM80 ArchTag work:

using Gemm = cutlass::gemm::device::Gemm<
    float, cutlass::layout::ColumnMajor,  // A
    float, cutlass::layout::ColumnMajor,  // B
    float, cutlass::layout::ColumnMajor,  // C
    float,                                 // Accumulator
    cutlass::arch::OpClassTensorOp,
    cutlass::arch::Sm80,                   // Works on SM120
    cutlass::gemm::GemmShape<128, 128, 16>,
    cutlass::gemm::GemmShape<64, 64, 16>,
    cutlass::gemm::GemmShape<16, 8, 8>,
    // ... epilogue, swizzle, stages
>;

This provides TensorCore TF32 computation at least, though not FP8.

Request

Could the CUTLASS team clarify:

1. Is SM120 FP8 GEMM intended to work with example 87a, or is a different configuration needed?
2. Is there a timeline for fixing the partition_S() alignment issue for SM120?

Thank you for your work on CUTLASS!

[m96-chan]

Workaround: Runtime Alignment Check in copy_sm75.hpp

We encountered this issue on SM120 (RTX 5090) and found a workaround by patching the LDSM copy operations to check alignment at runtime and fall back to scalar loads when misaligned.

Patch for cute/arch/copy_sm75.hpp

In SM75_U32x4_LDSM_N::copy():

#if defined(CUTE_ARCH_LDSM_SM75_ACTIVATED)
    uint32_t smem_int_ptr = cast_smem_ptr_to_uint(&smem_src);
    
    // Workaround for Issue #2902: partition_S() drops alignment
    if ((smem_int_ptr & 0xF) == 0) {
      // 16-byte aligned: use ldmatrix (fast path)
      asm volatile ("ldmatrix.sync.aligned.x4.m8n8.shared.b16 {%0, %1, %2, %3}, [%4];\n"
          : "=r"(dst0), "=r"(dst1), "=r"(dst2), "=r"(dst3)
          :  "r"(smem_int_ptr));
    } else {
      // Misaligned: use scalar loads (slow but correct)
      asm volatile (
          "ld.shared.u32 %0, [%4];\n"
          "ld.shared.u32 %1, [%5];\n"
          "ld.shared.u32 %2, [%6];\n"
          "ld.shared.u32 %3, [%7];\n"
          : "=r"(dst0), "=r"(dst1), "=r"(dst2), "=r"(dst3)
          : "r"(smem_int_ptr),
            "r"(smem_int_ptr + 4u),
            "r"(smem_int_ptr + 8u),
            "r"(smem_int_ptr + 12u));
    }
#else
    // ...
#endif

Same pattern applies to SM75_U32x2_LDSM_N, SM75_U16x4_LDSM_T, SM75_U16x8_LDSM_T, and the equivalents in copy_sm100.hpp and cutlass/arch/memory_sm75.h.

Caveats

• This is a workaround, not a proper fix. The scalar load path is slower than ldmatrix.
• The root cause (alignment being dropped in partition_S()) should still be fixed properly.
• We only tested on SM120; not sure if this affects other architectures.

Hope this helps others who hit this issue before a proper fix lands!

[github-actions]

This issue has been labeled inactive-30d due to no recent activity in the past 30 days. Please close this issue if no further response or action is needed. Otherwise, please respond with a comment indicating any updates or changes to the original issue and/or confirm this issue still needs to be addressed. This issue will be labeled inactive-90d if there is no activity in the next 60 days.