[BUG] CuTe DSL 4.7.1: the same MLIR compiles to different PTX from run to run (address-dependent codegen in the embedded libNVVM 13.3); not reproduced in 4.8.1

Open
  • Is something wrong?
#3,797 1 comment 2 reactions 0 assignees View on GitHub

Maintainers usually reply within 4 days

Nobody has claimed this yet.

Assessment

Difficulty
4/5
Estimated time
3-5 days
Newbie friendliness
35/100
Issue type
Bug
Clarity
Clearly specified
Activity status
Active
Tech stack
cpp, python
Domain
backend, compilers

Research direction

Start with the supplied repro.py and run it against CuTe DSL 4.7.1 as described; it saves the input MLIR and pipeline and reports PTX variants. Compare results with 4.8.1 and inspect the recorded environment and libNVVM version details. Done means establishing whether libNVVM 13.3 nondeterminism is known or guaranteed, and identifying whether 13.4 fixes it or merely avoids this reproduction.

Written by the indexing model from the issue text.

Description

bug CuTe DSL
Which component has the problem?

CuTe DSL

Bug Report

Describe the bug

Compiling the same kernel with CuTe DSL 4.7.1 does not always give the same PTX/SASS. In our builds about 2.5% of the kernels (161 of 6491, every target from sm_80 to sm_121a) came out in two or more variants over three cold builds of the same tree. Which kernels vary changes from build to build. The variants differ in instruction order and register numbering, and sometimes in register count or a folded constant.

The input to the backend does not vary. The MLIR the DSL hands to its pass pipeline is byte-identical in every run, and so is the LLVM-dialect module after the DSL's MLIR passes (cute-to-nvvm{... llvm-only=true}). The variation is in the libNVVM step, and it depends on where the compiler's allocations land:

  • The same input gives different PTX across processes, and also when the same module goes through the pipeline again in one process.
  • With ASLR disabled (setarch -R) and no extra threads in the process (OPENBLAS_NUM_THREADS=1; numpy's OpenBLAS starts its pool at import), 32 of 32 runs gave the same PTX. Either one alone was not enough.
  • Our kernels' LLVM IR, fed as text to the public nvidia-nvvm==13.3.73 libNVVM through the C API (nvvmCompileProgram, no DSL involved), gives the same PTX variants. It varies at -opt=2/-opt=3 and not at -opt=0/-opt=1. With nvidia-nvvm 13.4.59 or 13.4.92, every run gives the same PTX. We cannot share that IR; the text IR of the CUTLASS example below is rejected by libNVVM 13.3's parser ("parse expected type" in the kernel's parameter list), so the repro below goes through the DSL's pipeline.

With CuTe DSL 4.8.1 (embedded libNVVM 13.4.59) we see no variation: 0 of 6491 kernels over five cold builds, and 180 of 180 runs of the repro below.

Steps/Code to reproduce bug

Public code only: the Blackwell FMHA example of this repository at v4.7.1 (causal, head_dim 128, non-persistent), compiled for sm_100a with fake tensors. No GPU and no PyTorch needed.

pip install "nvidia-cutlass-dsl[cu13]==4.7.1"
git clone --depth 1 -b v4.7.1 https://github.com/NVIDIA/cutlass
python repro.py cutlass 120

The script:

  1. Calls cute.compile once and saves the module the DSL hands to its pass pipeline (generic form) and the pipeline string.
  2. Runs that module through that pipeline N times in one process: a new ir.Context, ir.Module.parse, PassManager.parse(pipeline).run(...), then a hash of the dumped PTX.
  3. Between runs, mallocs blocks of seeded random sizes and frees every other one, so each run sees another heap layout.

It prints the variants and exits 1 if there is more than one.

repro.py 
"""CuTe DSL 4.7.1: the same MLIR, run through the DSL's own cute-to-nvvm pipeline again in the same process, gives
different PTX. The kernel is the Blackwell FMHA example of the CUTLASS repository (v4.7.1; causal, head_dim 128,
non-persistent), compiled for sm_100a with fake tensors: no GPU and no PyTorch needed.

    pip install "nvidia-cutlass-dsl[cu13]==4.7.1"
    git clone --depth 1 -b v4.7.1 https://github.com/NVIDIA/cutlass
    python repro.py cutlass [N]

1. cute.compile of the kernel, once, saving the module the DSL hands to its pass pipeline (generic form) and the
   pipeline string: ./input.mlir, ./pipeline.txt.
2. N times (default 120): a new ir.Context, parse ./input.mlir, PassManager.parse(pipeline).run(...), hash the PTX the
   pipeline dumps. Before each run the script mallocs blocks of seeded random sizes and frees every other one, so the
   runs see different heap layouts.
It prints how many runs gave each PTX, keeps one PTX per variant in ./variants/ and exits 1 if there is more than one.
"""
import collections, ctypes, hashlib, os, random, re, shutil, sys, tempfile, types
from pathlib import Path

os.environ.setdefault("CUDA_VISIBLE_DEVICES", "")
os.environ.setdefault("CUTE_DSL_ARCH", "sm_100a")
cutlass_dir = Path(sys.argv[1])
count = int(sys.argv[2]) if len(sys.argv) > 2 else 120

# the example imports torch at module level but uses it only in its run(): a stand-in is enough to compile
class _Stub(types.ModuleType):
    def __getattr__(self, name):
        if name.startswith("__"):
            raise AttributeError(name)
        return _Stub(f"{self.__name__}.{name}")
    def __call__(self, *args, **kwargs):
        return _Stub("call")
try:
    import torch  # noqa: F401
except ImportError:
    sys.modules["torch"] = _Stub("torch")

examples = cutlass_dir / "examples/python/CuTeDSL"
sys.path[:0] = [str(examples), str(examples / "cute/blackwell/kernel/attention/fmha")]
import cutlass
import cutlass.cute as cute
from cutlass._mlir import ir, passmanager
from cutlass.base_dsl import compiler
from cutlass.cute.typing import Float32, Int32
import fmha
from helpers import fmha_helpers

# 1. the module and pipeline of one compile
compile_module = compiler.Compiler.compile
def keep_input(self, module, pipeline, *args, **kwargs):
    if "cute-to-nvvm" in pipeline:
        Path("input.mlir").write_text(module.operation.get_asm(print_generic_op_form=True))
        Path("pipeline.txt").write_text(re.sub(r"dump-ptx-path='[^']*'", "", pipeline))
    return compile_module(self, module, pipeline, *args, **kwargs)
compiler.Compiler.compile = keep_input

def fake(rank):
    shape = tuple(cute.sym_int() for _ in range(rank - 1)) + (128,)
    stride = tuple(cute.sym_int64(divisibility=8) for _ in range(rank - 1)) + (1,)
    return cute.runtime.make_fake_tensor(cutlass.Float16, shape, stride=stride, assumed_align=32)

kernel = fmha.BlackwellFusedMultiHeadAttentionForward(
    Float32, Float32, (128, 128), 128, False, fmha_helpers.MaskEnum.WINDOW_MASK_INFERENCE, False, False
)
q, k, v, o = (fake(5) for _ in range(4))
cute.compile(
    kernel, q, k, v, o, tuple(Int32(1) for _ in range(8)), None, None, None, None,
    Float32(1.0), Float32(1.0), Float32(1.0), None, None, Int32(0), None, None, cute.runtime.make_fake_stream(), False,
    options="--gpu-arch=sm_100a --opt-level=2",
)
compiler.Compiler.compile = compile_module
text = Path("input.mlir").read_text()
pipeline = Path("pipeline.txt").read_text()

# 2. the same module through the same pipeline, again and again
libc = ctypes.CDLL(None)
libc.malloc.restype = ctypes.c_void_p
libc.free.argtypes = [ctypes.c_void_p]
rng = random.Random(0)
work = Path(tempfile.mkdtemp(prefix="cutedsl_repro_"))
out = Path("variants")
shutil.rmtree(out, ignore_errors=True)
out.mkdir()
seen = collections.Counter()
for i in range(count):
    blocks = [libc.malloc(rng.choice((16, 24, 40, 72, 136, 264, 520, 1032, 4104, 65544))) for _ in range(rng.randrange(1, 3000))]
    for block in blocks[::2]:
        libc.free(block)
    dump = work / f"run{i}"
    with ir.Context() as context:
        context.enable_multithreading(False)
        module = ir.Module.parse(text)
        passmanager.PassManager.parse(pipeline.rstrip("})") + f" dump-ptx-path='{dump}'}})").run(module.operation)
    (ptx,) = work.glob(f"run{i}.*ptx")
    digest = hashlib.sha256(ptx.read_bytes()).hexdigest()[:12]
    seen[digest] += 1
    if not (out / f"{digest}.ptx").exists():
        shutil.copy(ptx, out / f"{digest}.ptx")
    print(f"run {i}: {digest}", flush=True)
shutil.rmtree(work)
print(f"cutlass-dsl {cutlass.__version__}: {count} runs of one module, {len(seen)} PTX variants: {dict(seen)}")
sys.exit(1 if len(seen) > 1 else 0)

Observed with 4.7.1: two PTX variants. Six processes of 60 runs each gave 55/5, 60/0, 53/7, 58/2, 55/5 and 57/3, so 22 of 360 runs produced the second variant. Diff excerpt (the rest is register renumbering):

-	.reg .b64 	%rd<3000>;
+	.reg .b64 	%rd<2999>;
@@
-	cvta.param.u64 	%rd399, %rd4;
+	cvta.param.u64 	%rd398, %rd4;
+	prefetch.tensormap 	[%rd398];
+	cvta.param.u64 	%rd399, %rd3;
 ...
-	cvta.param.u64 	%rd402, %rd1;
-	prefetch.tensormap 	[%rd402];

With nvidia-cutlass-dsl[cu13]==4.8.1 the same script gives one PTX in every run.

Expected behavior

The same input, compiler and options always give the same PTX and cubin.

Environment details (please complete the following information):

  • Environment location: Docker container on bare metal, Linux x86_64 (Ubuntu 24.04.3, glibc 2.39, kernel 5.10). Compile only (CUTE_DSL_ARCH=sm_100a, CUDA_VISIBLE_DEVICES=), no GPU used.
  • nvidia-cutlass-dsl 4.7.1 (nvidia-cutlass-dsl-libs-cu13 4.7.1; embedded libNVVM "CL-37800683, Cuda compilation tools, release 13.3, V13.3.27"). Compared with 4.8.1 (embedded libNVVM 13.4, V13.4.59).
  • Standalone libNVVM checks: nvidia-nvvm 13.3.73 varies; 13.4.59 and 13.4.92 do not.
  • Python 3.13.14.

Additional context

Ruled out:

  • The Python frontend: the pipeline input is byte-identical in every run.
  • PYTHONHASHSEED: fixing it does not help.
  • Dump paths and the pid: with ASLR off and one thread, runs with different paths and pids gave the same PTX.
  • MLIR multithreading: disabled on every context in the repro.
  • #3564 (link-libraries order): our pipelines link no libraries.

Questions:

  1. Is this a known libNVVM 13.3 issue that was fixed in 13.4? The 4.8.x and CUDA 13.4 release notes do not mention it, and we could not tell a fix from a change that no longer triggers on these kernels.
  2. Is deterministic code generation (same input and options give the same PTX and cubin) something CuTe DSL and libNVVM guarantee and test? We rely on it to cache and compare compiled kernels.
Dominant language
C++
Stars
10.7k
Forks
2.2k
Avg merge
3d 20h
Merged PRs (30d)
4

Getting set up

This project ships no dev container, Dockerfile or contributing guide, so setting up is up to you: start from its README, and see our first-contribution guide for the general steps.

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.

More from NVIDIA/cutlass

All issues in NVIDIA/cutlass

Similar issues

More C++ issues

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.