Skip to content

[BUG][CuTeDSL] cute.copy(..., tma_desc_ptr=p) with a gmem-space p passes the Python checks, then fails in libNVVM #3671

Open

Description

@zkyue

Which component has the problem?

CuTe DSL

Bug Report

Describe the bug

For TensorMapManager(..., 128), get_tensormap_ptr() defaults to !cute.ptr<!cute_nvgpu.tma_descriptor_tiled, gmem, align<128>>; copy_tensormap accepts smem or gmem destinations, so a gmem workspace needs the gmem form. Passing the same pointer to cute.copy(atom, src, dst, tma_bar_ptr=bar, tma_desc_ptr=p) gets through the isinstance(tma_desc_ptr, Pointer) check in cute/nvgpu/cpasync/copy.py, is lowered as ptr addrspace(1), and is handed to llvm.nvvm.cp.async.bulk.tensor.g2s.*, whose descriptor operand is a generic ptr. libNVVM then fails with a signature mismatch, and the printed suggestion points at CUDA_TOOLKIT_PATH. Converting first with get_tensormap_ptr(p, cute.AddressSpace.generic) compiles and loads the right data.

Steps/Code to reproduce bug

One CTA loads a 64x64 fp32 tile with TMA through a global-memory copy of the atom's descriptor, first via a generic pointer, then via the default gmem one. 4.7.1 and 4.8.0 behave the same.

import torch
import cutlass
import cutlass.cute as cute
from cutlass.cute.nvgpu import cpasync
from cutlass.cute.runtime import from_dlpack
from cutlass.utils import SmemAllocator, TensorMapManager, TensorMapUpdateMode

M, N = 64, 64


@cute.kernel
def load_tile(atom: cute.CopyAtom, mX: cute.Tensor, y: cute.Tensor, ws: cute.Tensor,
              space: cutlass.Constexpr):
    tidx, _, _ = cute.arch.thread_idx()
    smem = SmemAllocator()
    sX = smem.allocate_tensor(cutlass.Float32, cute.make_layout((M, N, 1), stride=(N, 1, M * N)), 128)
    bar = smem.allocate_array(cutlass.Int64, 1)
    gX = cute.local_tile(mX, (M, N), (None, None))
    tXsX, tXgX = cpasync.tma_partition(atom, 0, cute.make_layout(1),
                                       cute.group_modes(sX, 0, 2), cute.group_modes(gX, 0, 2))
    tm = TensorMapManager(TensorMapUpdateMode.GMEM, 128)
    desc = tm.get_tensormap_ptr(ws.iterator)  # default address_space=AddressSpace.gmem
    copy_desc = desc
    if cutlass.const_expr(space == "generic"):
        copy_desc = tm.get_tensormap_ptr(desc, cute.AddressSpace.generic)
    if tidx == 0:
        cpasync.copy_tensormap(atom, desc)
        cpasync.fence_tma_desc_release()
        cute.arch.mbarrier_init(bar, 1)
        cute.arch.mbarrier_init_fence()
    cute.arch.sync_threads()
    if tidx == 0:
        cute.arch.mbarrier_arrive_and_expect_tx(bar, M * N * 4)
    if tidx < 32:  # whole warp: the TMA copy elects one lane itself
        cpasync.fence_tma_desc_acquire(desc)
        cute.copy(atom, tXgX[(None, 0, 0)], tXsX[(None, 0)], tma_bar_ptr=bar, tma_desc_ptr=copy_desc)
    cute.arch.mbarrier_wait(bar, 0)
    for m in cutlass.range_constexpr(M):
        y[m, tidx] = sX[m, tidx, 0]


@cute.jit
def run(x: cute.Tensor, y: cute.Tensor, ws: cute.Tensor, space: cutlass.Constexpr):
    atom, mX = cpasync.make_tiled_tma_atom(
        cpasync.CopyBulkTensorTileG2SOp(), x, cute.make_layout((M, N), stride=(N, 1)), (M, N))
    load_tile(atom, mX, y, ws, space).launch(grid=(1, 1, 1), block=(N, 1, 1), smem=M * N * 4 + 1024)


print("cutlass", cutlass.__version__, "|", torch.cuda.get_device_name(), "| torch", torch.__version__)
x = torch.arange(M * N, dtype=torch.float32, device="cuda").reshape(M, N)
ws = torch.zeros(16, dtype=torch.int64, device="cuda")  # 128 B, one descriptor
fx, fws = from_dlpack(x, assumed_align=16), from_dlpack(ws, assumed_align=128)
for space in ("generic", "gmem"):
    print("tma_desc_ptr address space:", space, flush=True)
    y = torch.zeros_like(x)
    fy = from_dlpack(y, assumed_align=16)
    compiled = cute.compile(run, fx, fy, fws, space)
    compiled(fx, fy, fws)
    torch.cuda.synchronize()
    print("  y == x:", torch.equal(y, x), flush=True)

Output on 4.8.0, minus three cutlass.utils DeprecationWarnings:

cutlass 4.8.0 | NVIDIA H100 80GB HBM3 | torch 2.12.0a0+5aff3928d8.nv26.05
tma_desc_ptr address space: generic
  y == x: True
tma_desc_ptr address space: gmem
Call parameter type does not match function signature!
  %26 = getelementptr i64, ptr addrspace(1) %15, i32 0
 ptr  call void @llvm.nvvm.cp.async.bulk.tensor.g2s.cta.tile.2d(ptr addrspace(3) %8, ptr addrspace(3) %11, ptr addrspace(1) %26, i32 0, i32 0, i64 0, i1 true)
error: NVVM backend compilation failed
  error: libNVVM failed while compiling generated device IR.
  note: target architecture: sm_90a
  note: backend log:
        Failed creating the llvm::Module.
  suggestion: check that CUDA_TOOLKIT_PATH is set correctly and that the CUDA toolkit supports
              target architecture sm_90a; rerun with IR dump flags if backend context is needed

Expected behavior

Either the lowering casts a gmem descriptor pointer to generic, or cute.copy checks tma_desc_ptr.memspace and raises a ValueError that says to pass AddressSpace.generic. Every tma_desc_ptr site in copy.py has the same type-only check, in 4.8.0 and on main (0b55a2f).

Environment details

  • Environment location: container (Kubernetes pod)
  • nvidia-cutlass-dsl 4.7.1 and 4.8.0, H100 80GB HBM3 (sm_90a), host driver 590.44.01, CUDA 13.2 compat libraries loaded (libcuda.so.595.58.03, libnvidia-nvvm70.so.4), Python 3.12, torch 2.12

Additional context

So a kernel that keeps its own descriptor copy holds two pointers to the same 128 bytes. NVIDIA's grouped GEMM example converts to generic by hand before cute.copy; neither get_tensormap_ptr nor cute.copy documents it.

Activity

  1. kvnloo commented on Sep 27, 2026

    @kvnloo

    I think the most useful fix is to reject or normalize the descriptor pointer before it reaches NVVM.

    Right now the API accepts:

    tma_desc_ptr=<gmem pointer>

    through the Python-side Pointer check, but the NVVM intrinsic expects the descriptor operand in the generic address space.

    That gives us a bad contract:

    Python accepts -> IR lowers -> libNVVM fails with a misleading toolkit error

    A cleaner contract would be one of:

    • automatically apply the supported generic-address conversion for a valid TensorMap descriptor pointer; or
    • reject the non-generic pointer at the CuTe DSL boundary with an actionable error.

    Regression matrix:

    • manager-produced generic pointer -> compiles and loads correct data;
    • manager-produced gmem pointer -> either normalized or rejected before NVVM;
    • ordinary baked descriptor path -> unchanged.

    I would prefer fixing the API boundary rather than relying on every caller to know that get_tensormap_ptr(..., AddressSpace.generic) is required.

  2. github-actions commented on Oct 27, 2026

    @github-actions

    This issue has been labeled inactive-30d due to no recent activity in the past 30 days. Please close this issue if no further response or action is needed. Otherwise, please respond with a comment indicating any updates or changes to the original issue and/or confirm this issue still needs to be addressed. This issue will be labeled inactive-90d if there is no activity in the next 60 days.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Metadata

Metadata

Assignees

No one assigned

    Type

    No type

    Projects

    No projects

      Milestone

      No milestone

      Relationships

      None yet

      Development

      No branches or pull requests

      Issue actions