NVIDIA / NVIDIA/cutlass

[QST][CUTEDSL] Address Misalignment in FP8 Gemm

Open
#2,902 6 comments 0 reactions 0 assignees View on GitHub

Nobody has claimed this yet.

? - Needs Triage inactive-90d question
Dominant language
C++
Stars
10.5k
Forks
2.1k
Avg merge
3d 11h
Merged PRs (30d)
7

Description

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))>
--------------------------------

Contributor guide

No contributing guide indexed for this repository

First steps

  1. Read the whole issue, then the project's contributing guide.
  2. Comment on the issue to say you are picking it up — it saves two people doing the same work.
  3. Fork the repository and make your change on a branch.
  4. Open a pull request that references the issue number.

Research direction

Start at thr_copy_ldmatrix_A.partition_S(sA) and inspect the reported shared-memory tensor and copy-view layouts, especially their alignment values. Reproduce the FP8 GEMM load and determine what change prevents the 16-byte misaligned-address error when loading from shared memory to registers.

Written by the indexing model from the issue text.

Assessment

Tech stack
cpp
Domain
hpc, performance
Issue type
Bug
Difficulty
4/5
Estimated time
3-5 days
Activity status
Quiet
Clarity
Needs clarification
Newbie friendliness
35/100

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.