[GH-ISSUE #479] [Bug] SM100 HC-prenorm TMEM deallocation fault #200

New issue

Open

opened 2026-10-26 21:14:39 +00:00 by bsakaguchi · 0 comments

bsakaguchi commented

2026-10-26 21:14:39 +00:00

Owner

Originally created by @henrylhtsang on GitHub (Oct 21, 2026).
Original GitHub issue: https://github.com/deepseek-ai/DeepGEMM/issues/479

Environment

B200; DeepSeek V4.1 Flash.

Observed behavior

CUDA fault in sm100_tf32_hc_prenorm_gemm_impl:

CUDBG_EXCEPTION_WARP_TMEM_ACCESS_CHECK (20)
Tensor Memory column 3136 being dealloced by instruction
tcgen05.dealloc is not allocated. Columns allocated are: 192-223.

The kernel ignores the returned TMEM base and calls free(0, kNumTmemCols). This assumes the allocated TMEM base is always zero. Is that assumption guaranteed when kernels run concurrently on separate CUDA streams?

CUTLASS uses the returned base:

tmem_allocator.allocate(cute::TMEM::Sm100TmemCapacityColumns, &shared_storage.tmem_base_ptr);
// After __syncthreads():
tCtAcc.data() = shared_storage.tmem_base_ptr;

Then frees that same base:

tmem_allocator.free(shared_storage.tmem_base_ptr, cute::TMEM::Sm100TmemCapacityColumns);

Reproduction

No standalone reproducer or run with an upstream DeepGEMM build yet. Causality is unconfirmed.

Originally created by @henrylhtsang on GitHub (Oct 21, 2026). Original GitHub issue: https://github.com/deepseek-ai/DeepGEMM/issues/479 ## Environment B200; DeepSeek V4.1 Flash. ## Observed behavior CUDA fault in `sm100_tf32_hc_prenorm_gemm_impl`: ```text CUDBG_EXCEPTION_WARP_TMEM_ACCESS_CHECK (20) Tensor Memory column 3136 being dealloced by instruction tcgen05.dealloc is not allocated. Columns allocated are: 192-223. ``` The [kernel](https://github.com/deepseek-ai/DeepGEMM/blob/5081f8513af57e61696de4f4f04b0e707ffb503e/deep_gemm/include/deep_gemm/impls/sm100_tf32_hc_prenorm_gemm.cuh#L267) ignores the returned TMEM base and calls `free(0, kNumTmemCols)`. This assumes the allocated TMEM base is always zero. Is that assumption guaranteed when kernels run concurrently on separate CUDA streams? CUTLASS [uses the returned base](https://github.com/NVIDIA/cutlass/blob/91843b71d90737e0a9b86f7d9df213cff5a3526f/examples/cute/tutorial/blackwell/01_mma_sm100.cu#L238-L241): ```cpp tmem_allocator.allocate(cute::TMEM::Sm100TmemCapacityColumns, &shared_storage.tmem_base_ptr); // After __syncthreads(): tCtAcc.data() = shared_storage.tmem_base_ptr; ``` Then [frees that same base](https://github.com/NVIDIA/cutlass/blob/91843b71d90737e0a9b86f7d9df213cff5a3526f/examples/cute/tutorial/blackwell/01_mma_sm100.cu#L329): ```cpp tmem_allocator.free(shared_storage.tmem_base_ptr, cute::TMEM::Sm100TmemCapacityColumns); ``` ## Reproduction No standalone reproducer or run with an upstream DeepGEMM build yet. Causality is unconfirmed.

bsakaguchi referenced this issue

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

[PR #200] [MERGED] Add DeepSeek V3.2 indexer logit kernels for H800 (SM90) #290