NVIDIA / NVIDIA/cutlass

[BUG] CuTe DSL TMEM allocator triggers internal struct.scalar pointer deprecation warnings

Open Beginner friendly
#3,451 1 comment 0 reactions 0 assignees View on GitHub

Nobody has claimed this yet.

CuTe DSL
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_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

  1. Compile a Blackwell CuTe DSL kernel that allocates TMEM with cutlass.utils.TmemAllocator and calls retrieve_ptr().
  2. 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.

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

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.