NVIDIA / NVIDIA/warp

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

Open Beginner friendly
#1,968 0 comments 0 reactions 0 assignees View on GitHub

Nobody has claimed this yet.

Dominant language
Python
Stars
7.1k
Forks
624
Avg merge
3d 17h
Merged PRs (30d)
5

Description

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

Contributor guide

Open the contributing guide

First steps

  1. Read the whole issue, then the project's contributing guide.
  2. Comment on the issue to say you are picking it up — it saves two people doing the same work.
  3. Fork the repository and make your change on a branch.
  4. Open a pull request that references the issue number.

Research direction

Open warp/native/cuda_crt.h at the replacement CUDA attribute macros, then run the supplied Python reproducer with wp.config.llvm_cuda enabled. Add a Clang CUDA regression test covering the two launch-bounds forms and verify compilation, expected values, and preserved PTX bounds.

Written by the indexing model from the issue text.

Assessment

Tech stack
cpp, python
Domain
compilers
Issue type
Bug
Difficulty
2/5
Estimated time
1-3 hours
Activity status
Active
Clarity
Clearly specified
Newbie friendliness
88/100

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.