Skip to content

[BUG] CuTe DSL CopyUniversalOp ignores l1c_evict_priority again (regression of #3199 since 4.6.0, loads and stores) #3711

Open

Description

@gau-nernst

Which component has the problem?

CuTe DSL

Bug Report

Describe the bug
A clear and concise description of what the bug is.

A gmem <-> rmem cute.copy through a CopyUniversalOp atom built with l1c_evict_priority=... emits plain ld.global / st.global without the .L1:: qualifier, for both loads and stores and for every priority I tried (NO_ALLOCATE, EVICT_FIRST, EVICT_LAST).

This is the same symptom as #3199 (loads), which was resolved in 4.5.0. It works in 4.5.3 and regressed in 4.6.0. Releases checked: 4.5.3 (good), 4.6.0, 4.6.3, 4.7.0, 4.7.1, 4.8.0 (bad). The PTX operand type also changes from .u64 to .b64 at 4.6.0, which suggests the copy now goes through a different lowering path that drops the attribute.

cute.arch.load / cute.arch.store with level1_eviction_priority= still emit the qualifier on 4.8.0, so the problem looks limited to the copy-atom path.

Steps/Code to reproduce bug
Follow this guide http://matthewrocklin.com/blog/work/2018/02/28/minimal-bug-reports to craft a minimal bug report. This helps us reproduce the issue you're having and resolve the issue more quickly.

Compile-only (no GPU run needed). It prints the global loads / stores in the generated PTX:

# CuTeDSL MRE: CopyUniversalOp drops l1c_evict_priority (regressed in 4.6.0).
#
# A gmem <-> rmem cute.copy through
#   make_copy_atom(CopyUniversalOp(), ..., l1c_evict_priority=NO_ALLOCATE)
# should emit ld/st.global.L1::no_allocate. 4.5.x does; 4.6.0 - 4.8.0 emit a
# plain ld/st.global for both loads and stores (EVICT_FIRST / EVICT_LAST are
# dropped the same way). cute.arch.load / cute.arch.store with
# level1_eviction_priority still emit the qualifier.
#
# Compile-only; prints the global loads / stores in the generated PTX.
#   python atom_evict_hint.py [NO_ALLOCATE|EVICT_FIRST|EVICT_LAST]
import glob
import importlib.metadata
import os
import re
import sys
import tempfile

DUMP = tempfile.mkdtemp(prefix="evict_hint_")
os.environ["CUTE_DSL_KEEP_PTX"] = "1"
os.environ["CUTE_DSL_DUMP_DIR"] = DUMP
os.environ["CUTE_DSL_DISABLE_FILE_CACHING"] = "1"

import cutlass.cute as cute  # noqa: E402
from cutlass import Int32  # noqa: E402

PRIO = getattr(
    cute.nvgpu.CacheEvictionPriority, sys.argv[1] if len(sys.argv) > 1 else "NO_ALLOCATE"
)


def make(hinted):
    @cute.kernel
    def kernel(src: cute.Tensor, dst: cute.Tensor):
        regs = cute.make_rmem_tensor(8, Int32)
        atom = cute.make_copy_atom(
            cute.nvgpu.CopyUniversalOp(), Int32, num_bits_per_copy=256,
            l1c_evict_priority=PRIO,
        )
        plain = cute.make_copy_atom(
            cute.nvgpu.CopyUniversalOp(), Int32, num_bits_per_copy=256
        )
        cute.copy(atom if hinted == "load" else plain, src, regs)
        cute.copy(atom if hinted == "store" else plain, regs, dst)

    @cute.jit
    def run(src: cute.Tensor, dst: cute.Tensor):
        kernel(src, dst).launch(grid=(1, 1, 1), block=(1, 1, 1))

    return run


t = cute.runtime.make_fake_tensor(Int32, (8,), stride=(1,), assumed_align=32)
print("nvidia-cutlass-dsl", importlib.metadata.version("nvidia-cutlass-dsl"), PRIO)
for hinted in ("load", "store"):
    cute.compile(make(hinted), t, t)
    ptx = max(glob.glob(f"{DUMP}/**/*.ptx", recursive=True), key=os.path.getmtime)
    ops = sorted(set(re.findall(r"(?:ld|st)\.global[\w.:]*", open(ptx).read())))
    print(f"  hinted {hinted:5s}: {ops}")

Output on 4.8.0 (same on 4.6.0 - 4.7.1):

nvidia-cutlass-dsl 4.8.0 CacheEvictionPriority.NO_ALLOCATE
  hinted load : ['ld.global.v4.b64', 'st.global.v4.b64']
  hinted store: ['ld.global.v4.b64', 'st.global.v4.b64']
nvidia-cutlass-dsl 4.8.0 CacheEvictionPriority.EVICT_FIRST
  hinted load : ['ld.global.v4.b64', 'st.global.v4.b64']
  hinted store: ['ld.global.v4.b64', 'st.global.v4.b64']

Output on 4.5.3:

nvidia-cutlass-dsl 4.5.3 CacheEvictionPriority.NO_ALLOCATE
  hinted load : ['ld.global.L1::no_allocate.v4.u64', 'st.global.v4.u64']
  hinted store: ['ld.global.v4.u64', 'st.global.L1::no_allocate.v4.u64']
nvidia-cutlass-dsl 4.5.3 CacheEvictionPriority.EVICT_FIRST
  hinted load : ['ld.global.L1::evict_first.v4.u64', 'st.global.v4.u64']
  hinted store: ['ld.global.v4.u64', 'st.global.L1::evict_first.v4.u64']

Expected behavior
A clear and concise description of what you expected to happen.

The hinted direction carries the requested qualifier, as in 4.5.3: ld.global.L1::no_allocate... / st.global.L1::no_allocate... (SASS LDG.E.NA / STG.E.NA). If the attribute is no longer supported on this atom, make_copy_atom should reject it instead of silently dropping it.

Environment details (please complete the following information):

  • Environment location: Bare-metal
  • GPU: NVIDIA GB300 (sm_103a), driver 610.43.02
  • OS: Ubuntu 24.04.4 LTS (aarch64), Python 3.12.13
  • nvidia-cutlass-dsl: 4.8.0 (also checked 4.5.3, 4.6.0, 4.6.3, 4.7.0, 4.7.1 from PyPI)

Environment details (please complete the following information):

  • Environment location: [Bare-metal, Docker, Cloud(specify cloud provider)]

Additional context
Add any other context about the problem here.

Activity

  1. yunweili3 commented on Oct 5, 2026

    @yunweili3

    Reproduced on B200 (sm_100a) with 4.8.0 vs 4.5.3, same output as above.

    Root cause: the 4.6 update removed the memory attributes from the universal copy atom type itself. CopyAtomSIMTSyncCopyType.get now only takes (value type, bits per copy) (4.5.3 took mem order, mem scope, l1 cache eviction priority, invariant as well), but CopyUniversalOp._make_trait kept **kwargs and never looks at it, so the attribute is swallowed with no warning. The IR has no slot for it any more, so it cannot be restored on this atom. The specialized ops do carry it: the same MRE with CopyG2ROp for the load and CopyR2GOp for the store emits ld.global.L1::no_allocate.v4.b64 / st.global.L1::no_allocate.v4.b64 on 4.8.0, so that is the supported path.

    Side effect of the same regression: the deprecated cutlass.utils.distributed.ld_bypass asked CopyUniversalOp for a VOLATILE/SYS load and has been emitting a plain ld.global since 4.6.

    PR #3713 makes CopyUniversalOp reject these attributes with a TypeError that points at the specialized ops (the second behaviour you asked for), and switches ld_bypass to CopyG2ROp.

  2. brandon-yujie-sun commented on Oct 19, 2026

    @brandon-yujie-sun
    Collaborator

    @gau-nernst thanks for reporting the issue. It's now fixed in 4.8.1.

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