Repository navigation
[BUG] Cache policy from cute.copy is dropped when the copy atom is created with invariant=True #3726
Description
Which component has the problem?
CuTe DSL
Bug Report
Describe the bug
When you create a copy atom with invariant=True and pass a L2 cache policy to cute.copy's cache_policy param, it will be silently dropped.
Steps/Code to reproduce bug
"""Minimal repro: CopyG2ROp(invariant=True) + cute.copy(cache_policy=...) drops the L2 policy.
Usage: python repro.py {inv,noinv} (profile with ncu and read the SASS of `k`)
"""
import sys
import torch
import cutlass
from cutlass import cute, BFloat16, Int64, Float32
from cutlass.cute.runtime import from_dlpack
INVARIANT = sys.argv[1] == "inv"
@cute.kernel
def k(mX: cute.Tensor, mY: cute.Tensor):
tidx, _, _ = cute.arch.thread_idx()
policy = cute.arch.inline_ptx(
"createpolicy.fractional.L2::evict_first.b64 {$w0}, {$r0};",
write_only_types=[Int64], read_only_args=[Float32(1.0)])
load = cute.make_copy_atom(cute.nvgpu.CopyG2ROp(), BFloat16, num_bits_per_copy=256,
invariant=INVARIANT)
src = cute.make_tensor((mX.iterator + tidx * 16).align(32), cute.make_layout(16))
r = cute.make_rmem_tensor((16,), BFloat16)
cute.copy(load, src, r, cache_policy=policy)
dst = cute.make_tensor(mY.iterator + tidx * 16, cute.make_layout(16))
cute.autovec_copy(r, dst)
@cute.jit
def launch(mX: cute.Tensor, mY: cute.Tensor):
k(mX, mY).launch(grid=[1, 1, 1], block=[32, 1, 1])
x = torch.randn(512, dtype=torch.bfloat16, device="cuda"); y = torch.empty_like(x)
launch(from_dlpack(x, assumed_align=32), from_dlpack(y, assumed_align=32))
torch.cuda.synchronize()
assert torch.equal(x, y)
print("ok", "invariant=" + str(INVARIANT))Run CUTE_DSL_KEEP_PTX=1 python cutedsl_invariant_cache_policy_repro.py and examine the PTX, you can see ld.global.nc.v8.b32 doesn't have the cache hint despite we did pass one in CuTeDSL.
Run CUTE_DSL_KEEP_PTX=1 python cutedsl_invariant_cache_policy_repro.py --no-invariant which will use invariant=False, and examine the PTX, this time the cache hint presents: ld.global.L2::cache_hint.v4.b64
Expected behavior
I can use both cache policy and invariant hint
Environment details (please complete the following information):
- Environment location: GB200, nvidia-cutlass-dsl==4.8.0
Additional context
Add any other context about the problem here.
Activity
yunweili3 commentedon Oct 10, 2026
Reproduced on B200 with nvidia-cutlass-dsl 4.8.0, and the drop is below the Python layer.
With
CUTE_DSL_KEEP=ir-debug,ptx, the raw IR for theinvariant=Truecase still carries the policy:but the PTX comes out as
ld.global.nc.v8.b32with noL2::cache_hint, while the same copy withinvariant=Falselowers told.global.L2::cache_hint.v4.b64 {...}, [%rd7], %rd1. So the lowering ofcute.copyfor an invariant G2R atom takes the non-coherent-load path and never looks at thepolicy field.
Two related observations:
invariant=Truetogether with any other non-default loadattribute (
l1c_evict_priority,load_cache_mode, ...):invariant_load is not allowed when the other non-default memory features (e.g., mem_order) are set to non-default values. The L2 policyescapes that check because it is a runtime value attached later through
atom.set_value.ld.global.nc.L2::cache_hint.v8.b32 {...}, [addr], policy;assembles fine with
ptxas -arch=sm_100a(CUDA 13.1), so the lowering could emit it.Until the lowering does, #3727 makes the Python side fail loudly instead of silently dropping
the policy:
cute.copy(..., cache_policy=...)with aninvariant=TrueG2R atom raises aValueErrorpointing at the two options. Workaround meanwhile:invariant=Falsewith the policy(you lose the
.ncpath), or an inline-PTXld.global.nc.L2::cache_hintfor the hot loop.