Report Clang CUDA compile errors (GH-1325)

[snico432]

Description

A failed Clang CUDA compilation now raises an exception immediately instead of silently continuing and causing a downstream module-load exception that doesn't indicate the root cause. Closes #1325.

Checklist

• I am familiar with the Contributing Guidelines.
• New or existing tests cover these changes.
• The documentation is up to date with these changes.

Test plan

Added test_invalid_native_func_compile_error in warp/tests/cuda/test_clang_cuda.py.

Ran:

uv run warp/tests/cuda/test_clang_cuda.py -k invalid_native_func_compile_error

Result: test collected and skipped locally because I don't have a CUDA device available.

Bug fix

Repro from #1325:

import sys
import warp as wp

mode = sys.argv[1] if len(sys.argv) > 1 else "nvrtc"
wp.config.llvm_cuda = (mode == "clang")

@wp.func_native("""
    INVALID SOURCE;
""")
def bad_func(x: float) -> float: ...

@wp.kernel
def test_kernel(a: wp.array(dtype=float)):
    i = wp.tid()
    a[i] = bad_func(1.0)

a = wp.zeros(10, dtype=float, device="cuda:0")
wp.launch(test_kernel, dim=10, inputs=[a], device="cuda:0")

NVRTC (uv run repro.py nvrtc):

Exception: CUDA kernel build failed with error code 6

Clang CUDA (uv run repro.py clang):

Exception: Failed to load CUDA module '__main__'

Summary by CodeRabbit

• Bug Fixes

  • Clang/LLVM CUDA compilation failures are now reported immediately during compilation rather than surfacing later as module load errors, improving error clarity and faster failure detection.

• Tests

  • Added test coverage that verifies native-function/CUDA compile failures are detected and reported at build time to prevent silent downstream errors.

• Documentation

  • Changelog updated to note the improved immediate reporting of CUDA compilation failures.

[copy-pr-bot]

This pull request requires additional validation before any workflows can run on NVIDIA's runners.

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.

[coderabbitai]

No actionable comments were generated in the recent review. 🎉

ℹ️ Recent review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yml

Review profile: CHILL

Plan: Enterprise

Run ID: 89f3ad0c-0773-44d4-bdd3-9781ef0af37a

📥 Commits

Reviewing files that changed from the base of the PR and between 5ad2363 and cd99b5b.

📒 Files selected for processing (3)

• CHANGELOG.md
• warp/_src/build.py
• warp/tests/cuda/test_clang_cuda.py

✅ Files skipped from review due to trivial changes (1)

• CHANGELOG.md

🚧 Files skipped from review as they are similar to previous changes (2)

• warp/_src/build.py
• warp/tests/cuda/test_clang_cuda.py

---

📝 Walkthrough

Walkthrough

The PR consolidates error checking in build_cuda() so LLVM/Clang CUDA compile errors set and check err like the NVRTC path, adds a unit test that triggers and asserts an immediate compile failure for a broken native function, and updates CHANGELOG.md documenting the fix.

Changes

Unified CUDA Compilation Error Handling

Layer / File(s)	Summary
Unified error handling in build_cuda() warp/_src/build.py	The build_cuda() function assigns the LLVM compilation result to err and moves the if err != 0 exception check to a shared location after both compilation branches, ensuring both Clang and NVRTC failures raise immediately.
Test for invalid native function compilation error warp/tests/cuda/test_clang_cuda.py	Adds test_invalid_native_func_compile_error() which defines a native-function stub with INVALID SOURCE and a uniquely-scoped kernel that calls it; the test asserts wp.launch raises with a message matching CUDA kernel build failed with error code.
Changelog documentation CHANGELOG.md	Documents that Clang CUDA compilation failures are now reported immediately, referencing GH-1325.

sequenceDiagram
  participant build_cuda
  participant wp_compile_cuda
  participant wp_cuda_compile_program
  alt llvm_cuda enabled
    build_cuda->>wp_compile_cuda: compile CUDA source (Clang/LLVM)
    wp_compile_cuda-->>build_cuda: err
  else llvm_cuda disabled
    build_cuda->>wp_cuda_compile_program: compile CUDA source (NVRTC)
    wp_cuda_compile_program-->>build_cuda: err
  end
  build_cuda->>build_cuda: if err != 0 raise Exception("CUDA kernel build failed with error code ...")

Loading

🎯 2 (Simple) | ⏱️ ~12 minutes

🚥 Pre-merge checks | ✅ 4 | ❌ 1

❌ Failed checks (1 warning)

Check name	Status	Explanation	Resolution
Docstring Coverage	⚠️ Warning	Docstring coverage is 0.00% which is insufficient. The required threshold is 80.00%.	Write docstrings for the functions missing them to satisfy the coverage threshold.

✅ Passed checks (4 passed)

Check name	Status	Explanation
Description Check	✅ Passed	Check skipped - CodeRabbit’s high-level summary is enabled.
Title check	✅ Passed	The title accurately summarizes the main change: reporting Clang CUDA compile errors immediately instead of silently failing, directly addressing issue GH-1325.
Linked Issues check	✅ Passed	The PR fully implements all requirements from issue #1325: checks wp_compile_cuda return value, raises immediate exception on failure, aligns Clang path with NVRTC behavior, and includes a test verifying the error is caught.
Out of Scope Changes check	✅ Passed	All changes are directly related to issue #1325: CHANGELOG entry documents the fix, build.py implements error checking, and test validates the behavior. No extraneous modifications detected.

_{✏️ Tip: You can configure your own custom pre-merge checks in the settings.}

✨ Finishing Touches

🧪 Generate unit tests (beta)

• Create PR with unit tests

---

_{Comment @coderabbitai help to get the list of available commands and usage tips.}

[greptile-apps]

Greptile Summary

This PR fixes a silent error-swallowing bug in the Clang/LLVM CUDA compilation path where the return value of wp_compile_cuda was discarded, causing downstream module-load failures instead of immediate, informative exceptions.

• warp/_src/build.py: Captures the return value of wp_compile_cuda and unifies the error check (if err != 0) to cover both the llvm and NVRTC compilation paths after the if/else block.
• warp/tests/cuda/test_clang_cuda.py: Adds test_invalid_native_func_compile_error which defines a kernel using an invalid native snippet and asserts that wp.launch raises the expected exception with the "CUDA kernel build failed with error code" message.

Confidence Score: 5/5

Safe to merge — the change is a minimal, targeted fix that captures and checks a previously discarded return value.

The fix is small and well-scoped: wp_compile_cuda has restype=ctypes.c_int explicitly set in context.py, so the captured err is always a Python int. Moving the unified if err != 0 check outside the if/else correctly handles both code paths without introducing any new risk. The test follows patterns used extensively across the test suite.

No files require special attention.

Important Files Changed

Filename	Overview
warp/_src/build.py	Captures return value of wp_compile_cuda (restype=c_int, confirmed in context.py) and moves error check outside the if/else so both llvm and NVRTC paths raise on non-zero codes.
warp/tests/cuda/test_clang_cuda.py	Adds test_invalid_native_func_compile_error using the established module=unique pattern; assertRaisesRegex matches the exact error string produced by build.py.
CHANGELOG.md	Adds changelog entry for the fix with correct issue reference.

Flowchart

%%{init: {'theme': 'neutral'}}%%
flowchart TD
    A[build_cuda called] --> B{llvm_cuda?}
    B -- Yes --> C["err = wp_compile_cuda(...)"]
    B -- No --> D["err = wp_cuda_compile_program(...)"]
    C --> E{err != 0?}
    D --> E
    E -- Yes --> F["raise Exception CUDA kernel build failed"]
    E -- No --> G[Return / Continue]

Loading

_{Reviews (4): Last reviewed commit: "Report Clang CUDA compile errors (GH-132..." | Re-trigger Greptile}

[c0d1f1ed]

I think this looks good. @snico432 what prompted wanting to fix this? The Clang compilation path for CUDA is just in a prototype state in case we have a use for it later or need another compiler to compare to. Just curious if you have a specific use for it.

[snico432]

@c0d1f1ed I don't have a specific use case. An engineer from your team came to one of my lectures and gave a presentation about the project and I found it interesting. I'm just trying to learn more about compilers while hopefully making some open-source contributions.

[shi-eric]

/ok to test cd99b5b

[shi-eric]

Thanks for the contribution, @snico432! I appreciate that you kept the branch consistently at one commit that was up-to-date with main which made it easy to merge this on our internal GitLab repository. The addition of tests, a changelog entry, and DCO signoff also saved some time.

Some minor feedback for future pull requests in the repo:

• Communicate in the original GitHub issue that you would like to work on it
• Allow the original commenter of a thread to decide whether to resolve it rather than resolving it yourself

There is overhead in reviewing pull requests, and our team is pretty small, so we would prefer to spend our time reviewing more substantive pull requests. But everyone has to start somewhere, and your pull request was already in a great shape 😃

[snico432]

Thanks for the feedback @shi-eric! I'll make those adjustments in future PRs.