Fix misaligned local memory access in 02_dump_reg_shmem - #3633
Open
random25160765-collab wants to merge 1 commit into
Open
Fix misaligned local memory access in 02_dump_reg_shmem#3633random25160765-collab wants to merge 1 commit into
random25160765-collab wants to merge 1 commit into
Conversation
Fragment is a plain cutlass::Array (2-byte aligned for half_t) but the iterator accesses it through a 16-byte-aligned AlignedArray. When the fragment is homed in local memory, the frame slot is allocated at +0xc while the stores are emitted as 16-byte STL.128, so the effective address is 12 (mod 16) and the kernel faults with cudaErrorMisalignedAddress on sm_120. Aligning the fragment moves the slot to +0x10 and the example completes. Fixes NVIDIA#3581 Co-Authored-By: Claude Code <noreply@anthropic.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Fixes #3581
examples/02_dump_reg_shmemfaults onsm_120withcudaErrorMisalignedAddress(error 716); the example printsFailedand exits255.Why
The iterator accesses its
Fragmentthrough a type that requires more alignment thanFragmentitself declares:include/cutlass/transform/threadblock/predicated_tile_iterator.h:181—AccessType = AlignedArray<Element, AccessSize, (AccessSize * sizeof_bits<Element>::value / 8)>→ 16-byte aligned forhalf_tinclude/cutlass/transform/threadblock/predicated_tile_iterator.h:191—Fragment = cutlass::Array<Element, ...>→ 2-byte aligned (cutlass::Arraydeclares noalignas; onlyAlignedArraydoes)include/cutlass/transform/threadblock/predicated_tile_iterator.h:327—AccessType *frag_ptr = reinterpret_cast<AccessType *>(&frag);When the compiler homes the fragment in local memory, it allocates the frame slot from the declared alignment (2 → offset
+0xc) but emits the stores from the cast-implied alignment (16 →STL.128). The slot lands at12 (mod 16), so the 16-byte store is misaligned and the kernel faults.Fix
examples/02_dump_reg_shmem/dump_reg_shmem.cu:85:Verification
RTX 5060 (
sm_120), CUDA 13.3,main@147295a3, unmodified apart from this change:STL.128offsets+0xc0xc, 0x1c, … 0x7c(all≡ 12 mod 16)alignas(16)+0x100x10, 0x20, … 0x80(all 16-byte aligned)Registers and frame size are unchanged (40 registers, 144-byte frame) — only the slot offset moves, so the placement is not constrained by the frame itself.
Built with
-DCUTLASS_NVCC_ARCHS=120 -DCUTLASS_ENABLE_EXAMPLES=ON; the example prints all of its dumps and exits0.Scope
This is the minimal fix for the example. The underlying question — whether the
Fragmenttypedef itself should declare the alignment its own accessors require — affects every caller that declarestypename Iterator::Fragment frag;and is left to the maintainers. Details, SASS and a CUTLASS-free reproducer are in #3581.🤖 Generated with Claude Code