[BUG] CuTe DSL TMEM allocator triggers internal struct.scalar pointer deprecation warnings
Nobody has claimed this yet.
- Dominant language
- C++
- Stars
- 10.5k
- Forks
- 2.1k
- Avg merge
- 3d 11h
- Merged PRs (30d)
- 7
Description
Which component has the problem?
CuTe DSL
Bug Report
[
Describe the bug
Not high pri but im tyring to drop my deprecation warnings to 0
Compiling a Blackwell CuTe DSL kernel through the public cutlass.utils.TmemAllocator API emits deprecation warnings from inside CuTe DSL:
nvidia_cutlass_dsl/python_packages/cutlass/cute/core.py:5811: DeprecationWarning:
Use explicit `struct.scalar.ptr` for pointer instead.
The user kernel does not access struct.scalar.value directly. The warning stacks lead to CuTe DSL's own TMEM helpers:
cutlass/cute/arch/tmem.py::alloc_tmemaccessessmem_ptr_to_write_address.valuecutlass/cute/arch/tmem.py::retrieve_tmem_ptraccessesptr_to_buffer_holding_addr.value
Both values can be struct.scalar pointer wrappers, whose .value property is deprecated in favor of .ptr.
Steps/Code to reproduce bug
- Compile a Blackwell CuTe DSL kernel that allocates TMEM with
cutlass.utils.TmemAllocatorand callsretrieve_ptr(). - Force compilation rather than loading a cached artifact and enable deprecation warnings:
CUTE_DSL_NO_CACHE=1 PYTHONWARNINGS=always::DeprecationWarning python kernel.py
The relevant warning stacks are:
cutlass/utils/tmem_allocator.py:444 in then_block_1
cute.arch.alloc_tmem(...)
cutlass/cute/arch/tmem.py:155 in alloc_tmem
smem_ptr_to_write_address.value
cutlass/cute/core.py:5811 in value
warnings.warn("Use explicit `struct.scalar.ptr` for pointer instead.")
and:
cutlass/utils/tmem_allocator.py:480 in retrieve_ptr
return cute.arch.retrieve_tmem_ptr(...)
cutlass/cute/arch/tmem.py:110 in retrieve_tmem_ptr
ptr_to_buffer_holding_addr.value
cutlass/cute/core.py:5811 in value
warnings.warn("Use explicit `struct.scalar.ptr` for pointer instead.")
The same internal .value accesses are present in the published 4.7.0 wheel and on main:
- https://github.com/NVIDIA/cutlass/blob/main/python/CuTeDSL/cutlass/cute/arch/tmem.py#L100-L111
- https://github.com/NVIDIA/cutlass/blob/main/python/CuTeDSL/cutlass/cute/arch/tmem.py#L151-L160
- https://github.com/NVIDIA/cutlass/blob/main/python/CuTeDSL/cutlass/cute/core.py#L5886-L5894
Expected behavior
Using the public TMEM allocator API should not trigger CuTe DSL's own pointer deprecation warnings. The internal TMEM helpers should extract the explicit pointer (struct.scalar.ptr) when passed a scalar pointer wrapper while continuing to support ordinary Pointer arguments.
Environment details
- Environment location: bare metal
- GPU: NVIDIA GB300
- Driver: 580.126.20
- CUDA toolkit: 13.1
- Architecture: aarch64
- Python: 3.13.12
- Observed with
nvidia-cutlass-dsl==4.6.0.dev0 - Confirmed the same internal accesses remain in the published
nvidia-cutlass-dsl==4.7.0wheel
Additional context
In one training workload, four identical warnings appear because two separately compiled backward kernels each allocate TMEM and retrieve its pointer. Capturing full warning stacks confirms all four warnings originate in the two CuTe DSL internal helper paths above, rather than direct use of the deprecated property by the user kernels.
Contributor guide
No contributing guide indexed for this repository
First steps
- Read the whole issue, then the project's contributing guide.
- Comment on the issue to say you are picking it up — it saves two people doing the same work.
- Fork the repository and make your change on a branch.
- Open a pull request that references the issue number.
Research direction
Start in cutlass/cute/arch/tmem.py at retrieve_tmem_ptr and alloc_tmem, then inspect the related calls from cutlass/utils/tmem_allocator.py. Run the supplied CUTE_DSL_NO_CACHE=1 PYTHONWARNINGS=always::DeprecationWarning reproduction with a Blackwell TMEM allocator kernel. Done means the helpers support scalar pointer wrappers and ordinary Pointer arguments without emitting the internal deprecation warnings.
Written by the indexing model from the issue text.
Assessment
- Tech stack
- python
- Domain
- compilers
- Issue type
- Bug
- Difficulty
- 2/5
- Estimated time
- 1-3 hours
- Activity status
- Active
- Clarity
- Clearly specified
- Newbie friendliness
- 78/100