[FeathBow]

Why

When the device refuses a kernel launch, the call returns as if the kernel had run. The launcher issues kernel<<<grid, block, smem, nullptr>>>(...) and never reads the launch status, so the outputs keep their old contents and no error reaches the caller.

On a single sm_89 GPU, tests built for the bundled sm_90 target fail only on their values, and compute-sanitizer shows cudaLaunchKernel returning cudaErrorNoKernelImageForDevice.

What

The CUDA launcher reads cudaGetLastError() after the launch and returns the CUDA error text when the device refused it, or null otherwise. The host entry checks that result and raises it with TVM_FFI_THROW(RuntimeError), naming the launched function. No exception crosses the C ABI, and Python receives a RuntimeError.

A new test in tests/ops/tir/test_launch.py asks for 1 MiB of dynamic shared memory, which every CUDA device refuses, and expects that error. The two codegen tests that check the launcher declaration now expect const char*.

Contract

docs/spec/codegen.md requires the launcher to return the CUDA error text of a refused launch and the host entry to raise it. docs/spec/runtime.md states that a refused launch raises RuntimeError from the call.

Risk

Calls whose launch was refused now raise. On a single sm_89 GPU this turns four tests/scaffold cases from passed to failed, because their sm_90 kernels never ran there. On a single sm_90 GPU they pass with and without this change.

Each launch adds one cudaGetLastError() call and one check on the host. The comment below shows no difference in time per call beyond the spread between runs.

[FeathBow]

Test plan

The new test in tests/ops/tir/test_launch.py fails on main with DID NOT RAISE RuntimeError and passes with this change on both GPUs below. It is the one test node this change adds.

pytest tests ran on main and on this change on two machines:

• a single sm_89 GPU with Python 3.12, PyTorch 2.11 for CUDA 12.9 and nvcc 12.8
• a single sm_90 GPU with Python 3.12, PyTorch 2.13 for CUDA 13.0 and nvcc 13.1

pre-commit run --all-files passes.

A/B on a single sm_90 GPU

pytest tests	main	this change
collected	2153	2154
passed	2136	2137
failed	4	4
skipped	13	13

The new test is the only one that changes outcome. The four failures on both sides are two tests that need a second GPU and two that cannot find a CUDA header in this machine's toolkit layout. These sm_90 results come from the first commit of this PR; CI covers the current version on sm_90.

A/B on a single sm_89 GPU

pytest tests	main	this change
collected	2153	2154
passed	2123	2120
failed	17	21
skipped	13	13

The new test passes. Four test_scaffold::test_skeleton_builds_and_every_export_is_callable cases pass on main and fail here with no kernel image is available for execution on the device.

The 15 tests that fail on both sides now report that error in place of a value mismatch. The bundled target is sm_90, so those launches never ran on this GPU.

With a small kernel built for sm_89:

	main	this change
launch with 1 MiB of dynamic shared memory	returns, output untouched	RuntimeError: tilefoundry: launch of ... failed: invalid argument
time per call, two runs	2.664 µs, 2.471 µs	2.532 µs, 2.496 µs

Each time is the median of 7 rounds of 5000 calls. Across all runs the rounds range from 2.32 to 2.97 µs.

[FeathBow]

I would also like to add sm_89 support to TileFoundry. The plan is an sm_89 architecture document, a device document for the card we test on, and the existing MMA, ldmatrix and cp.async paths running on that card.

The device schema calls global memory memory.hbm.*, which does not fit GDDR cards, so the first step would rename those facts.

If this fits your plans for TileFoundry, I would be glad to hear what you think.

[zhen8838]

Thanks. The fix is right: a refused launch no longer returns as though the kernel ran. The test is precise: 1 MiB of dynamic shared memory is refused on every device, so it hits exactly the new path.

Two inline comments: one on how the error crosses the C ABI, one on what the launch status can attribute.

[zhen8838]

This throws a C++ exception out of an extern "C" function. The host unit declares the launcher (src/tilefoundry/codegen/cpu/module.py:211-218, _declare_callee) and calls it (:207), and the host entry is what tvm-ffi registers (src/tilefoundry/codegen/cpu/templates/cpu_entry.cpp.j2:8). This works today only because every caller is C++ inside one library. A C or ctypes caller would terminate.

Returning cudaError_t to the host is not an option either: docs/spec/codegen.md says the host module must not reference CUDA types.

Suggested shape:

• The launcher returns const char*: nullptr on success, cudaGetErrorString(status) otherwise (a static string, so no lifetime issue).

• The host emits:

  if (const char* error = <launcher>(...)) {
    TVM_FFI_THROW(RuntimeError) << "tilefoundry: launch of <name> failed: " << error;
  }

  TVM_FFI_THROW comes with <tvm/ffi/function.h>, which the host preamble already includes.

• Update both launcher declarations (cuda/module.py:46 and :82), _declare_callee, _one_launch, and the codegen.md sentence.

The error then crosses tvm-ffi as an ffi error of kind RuntimeError, and no exception crosses the C ABI. Today the std::runtime_error reaches Python as RuntimeError only by fallback: tvm-ffi's safe-call wrapper re-kinds any std::exception as InternalError, and Python maps unknown kinds to RuntimeError.

On the name: the host does not know the __global__ name. That is kernel_symbol() in cuda/abi.py (<mangled>_kernel). The host has the launcher name (ctx.signature_of(callee).name, <mangled>_launch) and the callee's FunctionSymbol (ctx.codegen_context.symbols.by_function[id(callee)]), whose export_name is the readable dotted name. export_name reads better in an error. The two spec sentences would then say "naming the launched function" instead of "the kernel".

[FeathBow]

Thanks, this is a cleaner boundary. The launcher now returns the CUDA error text, or null on success, and the host entry raises it with TVM_FFI_THROW(RuntimeError), so no exception crosses the C ABI.

Both launcher declarations, _declare_callee, _one_launch and the codegen.md sentence are updated, and the two codegen tests that check the declaration now expect const char*. The message names the launched function. On a single sm_89 GPU the full suite gives the same results as with the first commit.

[zhen8838]

cudaGetLastError() right after the launch also returns a sticky asynchronous error left by an earlier kernel, and any stale error an earlier runtime call left on this thread. So "launch of X failed" can name a kernel that did not cause the fault.

Suggest saying so here, for example: "The status read after the launch can also carry an asynchronous error from earlier device work, so the named kernel is where the error was observed, not necessarily where it occurred."

And let the message hint at it, the way PyTorch does ("CUDA kernel errors might be asynchronously reported at some other API call ... consider passing CUDA_LAUNCH_BLOCKING=1"). For example: tilefoundry: launch of <name> failed: <error> (this can be an asynchronous error from an earlier kernel; set CUDA_LAUNCH_BLOCKING=1 to locate it).

[zhen8838]

Thanks for this fix, and for offering sm_89 support: it is welcome. Please send it as a separate PR from this one. It is mostly data and checks; the code paths it needs already exist.

Item	Change	Notes
Architecture document	New src/tilefoundry/target/cuda/hardware/nvidia_sm89.toml, modelled on nvidia_sm90.toml	[facts.capabilities] lists only mma.sync, ldmatrix and cp.async, so T.cuda.sm80.Mma, T.ldmatrix and T.copy_async are selectable and WGMMA / TMA drop out. Documents in that directory are discovered automatically (src/tilefoundry/target/base.py:121).
Device document	New nvidia_<card>.toml, modelled on nvidia_h200_sxm.toml	SM count, clock, memory, L2, dense peaks and service rates of your card.
Runtime headers	Compile them once for sm_89	ops/copy_async_bulk holds sm_90-only PTX inside templates, so it should not be instantiated on sm_89; a compile confirms it.
Tests	A pytest marker that skips unless the local GPU is sm_89	Our CI has only H200 GPUs. It runs the compiler-side tests (the document checks in tests/target/test_hardware_resources.py, schedule facts / candidates, lowering, an nvcc compile for sm_89) but cannot execute sm_89 kernels, so please post your local results for those.

The memory.hbm. * you can only modify this name in the toml of sm_89. When reading it in the compiler, you only need to treat it as gmem-related facts.
The FP8 mma.sync atom that sm_89 adds is a new instruction declaration. It can go in the same PR or in a second one, as you prefer.

The contract for both documents is docs/spec/target.md §4.1 and §4.2.