Skip to content

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

Description

@and2049

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

Activity

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Metadata

Metadata

Assignees

No one assigned

    Labels

    No labels
    No labels

    Type

    No type

    Projects

    No projects

      Milestone

      No milestone

      Relationships

      None yet

      Development

      No branches or pull requests

      Issue actions