@wp.kernel(launch_bounds=...) fails under wp.config.llvm_cuda

Bug Description

Kernels using @wp.kernel(launch_bounds=...) compile and run through NVRTC, but fail when wp.config.llvm_cuda selects the bundled Clang/LLVM CUDA compiler. I encountered this while evaluating the experimental Clang CUDA path with kernels using the documented launch-bounds option.

This prevents comparing the two compiler paths on these kernels without removing the launch-bounds option or working around the missing macro. Removing the option changes the compilation constraints, so it is not an equivalent compiler comparison.

Save the following as a Python file and run it with a CUDA-enabled Warp build that includes warp-clang:

import numpy as np
import warp as wp

wp.config.llvm_cuda = True


@wp.kernel(launch_bounds=(64, 1), enable_backward=False, module="unique")
def bounded_double(values: wp.array[float]):
    i = wp.tid()
    values[i] *= 2.0


values = wp.array(np.arange(8, dtype=np.float32), device="cuda:0")
wp.launch(bounded_double, dim=8, inputs=[values], block_dim=64, device="cuda:0")
print(values.numpy())

With wp.config.llvm_cuda = False, this prints:

[ 0.  2.  4.  6.  8. 10. 12. 14.]

With True, the first diagnostic is:

error: no template named '__launch_bounds__'; did you mean 'wp::launch_bounds_t'?
extern "C" __launch_bounds__(64, 1) __global__ void ...

The generated kernel uses CUDA's __launch_bounds__ spelling. warp/native/cuda_crt.h provides replacement CUDA attribute macros under the comment Attributes otherwise defined by CUDA's crt/host_defines.h. It defines __global__, __device__, and other attributes, but omits __launch_bounds__. Injecting the equivalent definition into the generated translation unit makes the kernel compile and return the expected values:

#define __launch_bounds__(...) __attribute__((launch_bounds(__VA_ARGS__)))

The resulting PTX contains .maxntid 64 and .minnctapersm 1. The single-argument form also compiles with this definition and emits .maxntid 64. Separate fresh kernel caches were used for the backend comparison.

Expected: Clang CUDA compiles the kernel and preserves the requested launch bounds. Actual: it rejects the entry-point declaration before emitting PTX.

Related: #1026 tracks broader Clang CUDA compatibility work; its linked analysis does not mention this launch-bounds failure. I can submit a focused fix in cuda_crt.h with a Clang CUDA regression test.

System Information

• Warp: 1.18.0.dev5, built from source at 1bf652059b9e
• Bundled LLVM: 22.1.8
• CUDA Toolkit / NVRTC: 13.0; NVIDIA driver: 595.84
• GPU: NVIDIA GeForce RTX 4050 Laptop GPU
• Ubuntu 24.04, Python 3.12.3