Skip to content

[BUG] CuTeDSL 4.8.0 nvpu-optimization fails without a diagnostic on a sum of 10 dynamic extents #3699

Open

Description

@allispaul

Which component has the problem?

CuTe DSL

Bug Report

Describe the bug
In nvidia-cutlass-dsl 4.8.0, compilation of a host function containing a linear chain of 10 or more cute.tuple_add ops on dynamic int tuples fails with a generic "compiler pass failed error". Compilation succeeds for <= 9 tuple_adds. The same code compiles in 4.7.1 with a chain of length up to 64.

The failing pass is nvgpu-optimization; disabling it removes the compile error on 4.8.0; for a chain of length 10, you can also successfully compile the code by saving the partially transformed module after failure and rerunning the identical full pipeline, with the pass still enabled.

Steps/Code to reproduce bug
Save the following as repro_tuple_add_chain.py:

import argparse                                                                              
                                                                                                 
import cutlass                                                                               
import torch                                                                                 
from cutlass import cute                                                                     
from cutlass.cute.runtime import from_dlpack                                                                                                                                             
                                                                                                 
@cute.kernel                                                                                 
def kernel(out: cute.Tensor):                                                                
    tidx, _, _ = cute.arch.thread_idx()                                                      
    if tidx == 0:                                                                            
        out[0] = cutlass.Int32(1)                                                            
                                                                                            
@cute.jit                                                                                    
def host(out: cute.Tensor, tensors: tuple):                                                  
    total = sum(cute.size(t, mode=[0]) for t in tensors)                                     
    kernel(out).launch(grid=(total, 1, 1), block=(32, 1, 1))                                 
                                                                                                                                                                                               
def main():                                                                                  
    parser = argparse.ArgumentParser()                                                       
    parser.add_argument("--n", type=int, default=10, help="number of summed sizes")          
    args = parser.parse_args()                                                               
                                                                                           
    out = from_dlpack(torch.zeros(1, dtype=torch.int32, device="cuda"))                      
    tensors = tuple(                                                                         
        from_dlpack(torch.zeros(8, dtype=torch.int32, device="cuda")).mark_layout_dynamic()  
        for _ in range(args.n)                                                               
    )                                                                                        
    cute.compile(host, out, tensors)                                                         
    print(f"n={args.n}: compiled OK")                                                        
                                                                                                 
                                                                                                 
if __name__ == "__main__":                                                                   
    main()

For example, run the failing case in an isolated 4.8.0 environment:

CUDA_VISIBLE_DEVICES=3 CUTE_DSL_NO_CACHE=1 \                      
    uv run --no-project --with 'nvidia-cutlass-dsl[cu13]==4.8.0' \  
    --with torch==2.13.0 python repro_tuple_add_chain.py --n 10

Change the version to 4.7.1 for the comparison. The abbreviated commands below assume the selected environment, CUDA_VISIBLE_DEVICES=3, and CUTE_DSL_NO_CACHE=1.

Command 4.7.1 4.8.0
python repro_tuple_add_chain.py --n 9 OK OK
python repro_tuple_add_chain.py --n 10 OK compiler pass failed
python repro_tuple_add_chain.py --n 64 OK compiler pass failed
CUTE_DSL_COMPILER_OPT="use-nvgpu-optimization-pass=false" python repro_tuple_add_chain.py --n 64 n/a OK

Error output with 4.8.0:

cutlass.base_dsl.compiler.CompilerDiagnosticError:  
error: compiler pass failed                         
    error: Failure while executing pass pipeline                                                                                  

Additional reproduction: save the failed pass output, then retry

Save this as repro_tuple_add_chain_retry.py beside
repro_tuple_add_chain.py:

"""Save a failed compilation's partially rewritten IR and retry the full pipeline."""

import json
import os
from pathlib import Path

# Disable caching before importing the DSL so the compiler hook always runs.
os.environ["CUTE_DSL_NO_CACHE"] = "1"

from cutlass._mlir import ir, passmanager
from cutlass.base_dsl.compiler import Compiler

from repro_tuple_add_chain import main

original_compile = Compiler.compile
output = Path(os.environ.get("REPRO_DUMP_DIR", "tuple-retry-dump"))


def compile_with_retry(self, module, pipeline, *args, **kwargs):
    """Retry the identical full pipeline on serialized failure-state MLIR."""
    output.mkdir(parents=True, exist_ok=False)
    (output / "pipeline.txt").write_text(pipeline)
    (output / "before.mlir").write_text(str(module))
    pm = passmanager.PassManager.parse(pipeline)
    (output / "expanded-pipeline.txt").write_text(str(pm))
    pm.enable_verifier(kwargs.get("enable_verifier", False))
    pm.enable_timing()
    diagnostics = []

    def record_diagnostic(diagnostic):
        """Retain native diagnostics without changing pass behavior."""
        diagnostics.append(str(diagnostic))
        return True

    handler = module.context.attach_diagnostic_handler(record_diagnostic)
    try:
        pm.run(module.operation)
    except ir.MLIRError as error:
        (output / "first-error.txt").write_text(str(error))
        (output / "after-failure.mlir").write_text(str(module))
        print(f"First pipeline failed: {error}")
    else:
        raise RuntimeError("Expected the first pipeline to fail; use 4.8.0 and --n 10")
    finally:
        handler.detach()
        (output / "first-diagnostics.json").write_text(json.dumps(diagnostics, indent=2))
        # Destruction prints native pass timings, including the last pass reached.
        del pm

    replay = ir.Module.parse((output / "after-failure.mlir").read_text())
    replay.operation.verify()
    result = original_compile(self, replay, pipeline, *args, **kwargs)
    (output / "after-retry.mlir").write_text(str(result))
    print("Full pipeline succeeded on reparsed after-failure.mlir")
    return result


if __name__ == "__main__":
    Compiler.compile = compile_with_retry
    try:
        main()
    finally:
        Compiler.compile = original_compile

This runs the same Python reproducer, intercepting
the native compilation entry point rather than supplying hand-written IR.

CUDA_VISIBLE_DEVICES=3 REPRO_DUMP_DIR=/tmp/tuple-retry-example \
  uv run --no-project --with 'nvidia-cutlass-dsl[cu13]==4.8.0' \
  --with torch==2.13.0 python repro_tuple_add_chain_retry.py --n 10

Use a fresh output directory. Do not set the pass-disable option; the retry
script disables the compilation cache itself. It performs these steps:

  1. Save the compiler's original pipeline string and its input module.
  2. Run that pipeline with native pass timing enabled. It executes the prefix
    through NvgpuOptimizationPass, which fails and stops the pipeline.
  3. Catch the failure and emit the partially rewritten module left by the
    failed pass
    to after-failure.mlir.
  4. Parse that file into a new MLIR module and verify it.
  5. Run the original compiler entry point with the same full pipeline string
    on that module. The pipeline succeeds, including nvgpu-optimization and
    the downstream compilation/JIT path; cute.compile returns normally.

The important distinction is that this is not a clean stop immediately
before nvgpu-optimization, nor a retry from the original untouched input.
The first, failing invocation has already modified the module.

Output, omitting the earlier entries in the native timing table:

... EliminateUnnecessarySyncPass
... NvgpuOptimizationPass
... Rest
... Total
First pipeline failed: Failure while executing pass pipeline
Full pipeline succeeded on reparsed after-failure.mlir
n=10: compiled OK

The output directory contains:

  • pipeline.txt: original full pipeline, reused without alteration.
  • expanded-pipeline.txt: constituent pass names for inspection, not the
    string reparsed for the retry.
  • before.mlir: original pipeline input.
  • first-error.txt and first-diagnostics.json: generic failure and the
    empty native diagnostic list ([]).
  • after-failure.mlir: the saved, verified input to the second attempt.
  • after-retry.mlir: successfully lowered module.

This succeeded repeatedly for --n 10.

Expected behavior
The valid Python example should compile, as it does on 4.7.1.

Environment details (please complete the following information):

  • nvidia-cutlass-dsl 4.8.0 (nvidia-cutlass-dsl-libs-cu13 4.8.0)
  • Python 3.12.3, torch 2.13.0
  • GPU: NVIDIA B300 SXM6 AC (target sm_103a), driver 580.126.09
  • Linux 6.8.0-100-generic
  • Comparison: nvidia-cutlass-dsl[cu13] 4.7.1 in a separate environment

Additional context
Thanks to @reubenconducts for helping track this down!

Activity

  1. brandon-yujie-sun commented on Oct 8, 2026

    @brandon-yujie-sun
    Collaborator
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