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_tmem accesses smem_ptr_to_write_address.value
cutlass/cute/arch/tmem.py::retrieve_tmem_ptr accesses ptr_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.TmemAllocator and calls retrieve_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:
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.0 wheel
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.
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.TmemAllocatorAPI emits deprecation warnings from inside CuTe DSL:The user kernel does not access
struct.scalar.valuedirectly. 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.valueBoth values can be
struct.scalarpointer wrappers, whose.valueproperty is deprecated in favor of.ptr.Steps/Code to reproduce bug
cutlass.utils.TmemAllocatorand callsretrieve_ptr().The relevant warning stacks are:
and:
The same internal
.valueaccesses are present in the published 4.7.0 wheel and onmain: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 ordinaryPointerarguments.Environment details
nvidia-cutlass-dsl==4.6.0.dev0nvidia-cutlass-dsl==4.7.0wheelAdditional 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.