Skip to content

[BUG] CuTe DSL: cute.compile stalls for minutes when no GPU is visible #3681

Open

Description

@hwanseoc

Which component has the problem?

CuTe DSL

Bug Report

Describe the bug

With no GPU visible and CUTE_DSL_ARCH unset, cute.compile stalls for minutes with no output. It then compiles for the default arch. On 4.8.0 a trivial kernel took 773 s; on 4.7.0 it took 325 s.

Arch detection recurses through its own error message:

  1. env_manager.arch calls detect_gpu_arch, which calls cuInit.
  2. cuInit fails with CUDA_ERROR_NO_DEVICE, so checkCudaErrors raises DSLCudaRuntimeError.
  3. That error's constructor calls _get_friendly_cuda_error_message, which reads env_manager.arch again. _arch is still unset, so this is step 1 again.

Each round pays a failing cuInit, about 2 s here, until Python's recursion limit. Then the except in detect_gpu_arch unwinds it. A faulthandler dump shows these frames repeating:

base_dsl/runtime/cuda.py:163  get_compute_capability_major_minor
base_dsl/env_manager.py:422   detect_gpu_arch
base_dsl/env_manager.py:999   arch
base_dsl/common.py:302        _get_friendly_cuda_error_message
base_dsl/common.py:459        __init__
base_dsl/runtime/cuda.py:359  checkCudaErrors

Steps/Code to reproduce bug

import time
import cuda.bindings.driver as cuda
import cutlass
import cutlass.cute as cute
from cutlass.cute.runtime import make_fake_stream


@cute.kernel
def noop_kernel(x: cutlass.Int32):
    pass


@cute.jit
def noop(x: cutlass.Int32, stream: cuda.CUstream):
    noop_kernel(x).launch(grid=(1, 1, 1), block=(32, 1, 1), stream=stream)


start = time.perf_counter()
cute.compile(noop, cutlass.Int32(0), make_fake_stream(), options="--enable-tvm-ffi")
print(f"cutlass-dsl {cutlass.__version__}: compiled in {time.perf_counter() - start:.0f} s", flush=True)
$ CUDA_VISIBLE_DEVICES= python repro.py
cutlass-dsl 4.8.0: compiled in 773 s
$ CUDA_VISIBLE_DEVICES= CUTE_DSL_ARCH=sm_100a python repro.py
cutlass-dsl 4.8.0: compiled in 0 s

Expected behavior

With no device, compile quickly for the default arch, or raise an error that points at CUTE_DSL_ARCH. Building an error message shouldn't re-enter device detection.

Environment details (please complete the following information):

  • Environment location: Bare-metal, run with CUDA_VISIBLE_DEVICES= to hide the GPU
  • nvidia-cutlass-dsl 4.8.0 and 4.7.0 (cu13), cuda-python 13.4.1
  • B200 host, driver 580.159.03, Ubuntu 22.04, Python 3.12

Additional context

This hits CPU-only build machines that compile kernels ahead of time. Workaround: set CUTE_DSL_ARCH.

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

Metadata

Metadata

Assignees

No one assigned

    Labels

    Type

    No type

    Projects

    No projects

      Milestone

      No milestone

      Relationships

      None yet

      Development

      No branches or pull requests

      Issue actions