Repository navigation
[BUG] TVM-FFI exported functions hang under multithreads #3706
Description
Which component has the problem?
CuTe DSL
Bug Report
Describe the bug
If you compile a cute.jit function with tvm-ffi, and launch it from multithreads, it would cause a deadlock
My agent debugged a bit using LLDB and looks like it's something related to cuda_dialect_init_library_once in libcute_dsl_runtime.so which is not open sourced. The issue seems to be some lock is not properly released when it returns. Because there is no source my agent can only inspect the assembly so I don't really know what is going on exactly.
This is what agent interprets from the assembly:
int cuda_dialect_init_library_once(int64_t *state, load_fn, per_dev_fn, err_fn) {
if (*state) return 0; // fast path, no lock
while (!CAS(&g_lock, 0, 1)) ; // 0x1d260e: acquire
if (*state) return 0; // 0x1d261d -> 0x1d261f: BUG 1, returns holding g_lock
void *lib = 0;
if (load_fn(&lib)) { err_fn(...); return -1; } // BUG 2: returns holding g_lock
if (per_dev_fn) { /* loop devices; if all fail: */ err_fn(...); return -1; } // BUG 2
*state = (int64_t)lib;
g_lock = 0; // 0x1d26b2: the only release
return 0;
}
Steps/Code to reproduce bug
Note: I can only reproduce this on a HGX 8 GPU node with a x64 CPU. If I run these on a GB200 node with ARM CPU they wouldn't hang somehow.
It can be reproduced easily from a simple multi-gpu JAX program:
Note: it only works with jax_tvm_ffi.register_ffi_target where the function is exported using tvm-ffi. If you use jax.ffi.ffi_call it won't hang
# JAX repro: a CuTe DSL kernel exported with --enable-tvm-ffi and registered as an XLA FFI
# target via jax_tvm_ffi.register_ffi_target. In a single-process multi-GPU jit, XLA's
# per-device threads make the first call of the kernel concurrently, which leaks the init
# lock in libcute_dsl_runtime.so and the step never finishes. Needs >= 3 local GPUs.
# Usage: python3 tvmffi_jax_repro.py (all local GPUs -> hangs; kill it)
# CUDA_VISIBLE_DEVICES=0 python3 ... (control -> OK)
import os
os.environ.setdefault("XLA_PYTHON_CLIENT_PREALLOCATE", "false")
import jax
import jax.numpy as jnp
import jax_tvm_ffi
from jax.sharding import Mesh, PartitionSpec as P
import cutlass.cute as cute
from cutlass import Float32
N = 256 # elements per device; one CTA
@cute.kernel
def add_one_kernel(x: cute.Tensor, y: cute.Tensor):
i, _, _ = cute.arch.thread_idx()
y[i] = x[i] + 1.0
@cute.jit
def add_one(x: cute.Tensor, y: cute.Tensor, stream):
add_one_kernel(x, y).launch(grid=(1, 1, 1), block=(N, 1, 1), stream=stream)
fake = lambda: cute.runtime.make_fake_compact_tensor(
Float32, (N,), stride_order=(0,), memspace=cute.AddressSpace.gmem, assumed_align=16
)
compiled = cute.compile(add_one, fake(), fake(), cute.runtime.make_fake_stream(use_tvm_ffi_env_stream=True),
options="--enable-tvm-ffi")
jax_tvm_ffi.register_ffi_target("cute_add_one", compiled, ["args", "rets"], platform="gpu")
def per_device(x):
return jax.ffi.ffi_call("cute_add_one", jax.ShapeDtypeStruct(x.shape, x.dtype))(x)
mesh = Mesh(jax.local_devices(), ("x",))
step = jax.jit(jax.shard_map(per_device, mesh=mesh, in_specs=P("x"), out_specs=P("x")))
x = jnp.arange(mesh.size * N, dtype=jnp.float32)
out = step(x).block_until_ready() # hangs here on >= 3 GPUs
print(f"OK: {mesh.size}-device step finished; result correct={bool(jnp.array_equal(out, x + 1))}")However, the bug is not related to JAX, and you can also reproduce by using multithreads:
# Minimal repro: libcute_dsl_runtime.so leaks its init lock when threads make the first
# call of a --enable-tvm-ffi function concurrently. With 3 threads, one loads the module,
# one returns still holding the lock, and the third spins forever: the script never exits.
import threading
import cutlass.cute as cute
@cute.kernel
def noop():
pass
@cute.jit
def launch():
noop().launch(grid=(1, 1, 1), block=(1, 1, 1))
f = cute.compile(launch, options="--enable-tvm-ffi")
barrier = threading.Barrier(3)
def first_call_of_f():
barrier.wait()
f()
threads = [threading.Thread(target=first_call_of_f) for _ in range(3)]
for t in threads:
t.start()
for t in threads:
t.join()
print("OK: no hang")Expected behavior
No hang
Environment details (please complete the following information):
➜ TransformerEngine git:(debug) ✗ pip show nvidia-cutlass-dsl
Name: nvidia-cutlass-dsl
Version: 4.7.0
Summary: NVIDIA CUTLASS Python DSL
Home-page:
Author: NVIDIA Corporation
Author-email:
License:
Location: /usr/local/lib/python3.12/dist-packages
Requires: nvidia-cutlass-dsl-libs-base, nvidia-cutlass-dsl-libs-cu12
Required-by: nvidia-cudnn-frontend
Additional context
Add any other context about the problem here.
Activity
[-][BUG] TVM-FFI export functions hang under multithreads[/-][+][BUG] TVM-FFI exported functions hang under multithreads[/+]aetherxeg-source commentedon Oct 2, 2026
The concurrency invariant here is stronger than “initialization happens once”:
"once-initialized state requires all exit paths to release initialization authority"
If the slow path holds a global initialization lock, every success, already-initialized race, and failure path must terminate with the lock in a valid state.
A useful state model is:
"uninitialized -> initializing(owner) -> initialized | failed-retryable"
where leaving "initializing(owner)" requires an explicit release transition.
The adversarial test should synchronize 3+ threads at first use and independently exercise:
Then assert bounded completion and exactly one committed initialized state.
This catches the class of bug where the logical state is correct but the authority/lock protecting its transition is leaked.
Related reference implementation / design work: AETHER X GLOBAL ExecSurface — https://github.com/AETHERXGLOBAL/execsurface
kainzhong commentedon Oct 2, 2026
Confirm this is fixed in 4.8.0