https://github.com/Sirraide updated https://github.com/llvm/llvm-project/pull/217427
>From 9bfbb8f727a0457c83821d453a2500371cd68377 Mon Sep 17 00:00:00 2001 From: Ambrose Leeb <[email protected]> Date: Wed, 19 Aug 2026 20:29:26 +0200 Subject: [PATCH 1/2] [Clang] [CUDA] Provide device-side definitions of __cxa_[pure|deleted]_virtual() Update our `__clang_*` headers to provide weak symbols for these like we're already doing for HIP (in fact, I just moved those exact definitions into a separate header that is now used for both CUDA and HIP). The issues that this fixes reported a `ptxas` error, but I don't think our lit infrastructure supports checking if `ptxas` is even available in the `PATH`, so instead, I just added a test that checks that the definitions make it into the device-side IR. Fixes #49183, fixes #67533. --- clang/docs/ReleaseNotes.md | 5 +++ clang/lib/Headers/CMakeLists.txt | 1 + .../Headers/__clang_cuda_runtime_wrapper.h | 1 + .../Headers/__clang_gpu_device_virtual_trap.h | 39 +++++++++++++++++++ .../lib/Headers/__clang_hip_runtime_wrapper.h | 17 +------- .../test/CodeGenCUDA/pure_deleted_virtual.cu | 9 +++++ 6 files changed, 56 insertions(+), 16 deletions(-) create mode 100644 clang/lib/Headers/__clang_gpu_device_virtual_trap.h create mode 100644 clang/test/CodeGenCUDA/pure_deleted_virtual.cu diff --git a/clang/docs/ReleaseNotes.md b/clang/docs/ReleaseNotes.md index e4a6f72f8fec5..924fe8b373808 100644 --- a/clang/docs/ReleaseNotes.md +++ b/clang/docs/ReleaseNotes.md @@ -548,6 +548,11 @@ features cannot lower the translation-unit ABI level; - Added `--cuda-emit-nvcc-abi` to emit the NVCC-compatible host registration ABI (`__cudaRegisterLinkedBinary`). +- Clang now provides device-side definitions of `__cxa_pure_virtual()` and + `__cxa_deleted_virtual()`; previously, any (potential) call to a pure/deleted + virtual function that could not be optimised out would cause the program to + fail to assemble. This is now fixed. (#GH49183) (#GH67533) + #### AIX Support #### NetBSD Support diff --git a/clang/lib/Headers/CMakeLists.txt b/clang/lib/Headers/CMakeLists.txt index 21b6eb5e38052..cb6567c22f0cd 100644 --- a/clang/lib/Headers/CMakeLists.txt +++ b/clang/lib/Headers/CMakeLists.txt @@ -311,6 +311,7 @@ set(gpu_files __clang_gpu_device_functions.h __clang_gpu_intrinsics.h __clang_gpu_runtime_wrapper.h + __clang_gpu_device_virtual_trap.h ) set(gpu_hip_wrapper_files diff --git a/clang/lib/Headers/__clang_cuda_runtime_wrapper.h b/clang/lib/Headers/__clang_cuda_runtime_wrapper.h index 29178bf0db8a1..efd5016c2847d 100644 --- a/clang/lib/Headers/__clang_cuda_runtime_wrapper.h +++ b/clang/lib/Headers/__clang_cuda_runtime_wrapper.h @@ -491,6 +491,7 @@ __device__ inline __cuda_builtin_gridDim_t::operator uint3() const { #include <__clang_cuda_cmath.h> #include <__clang_cuda_intrinsics.h> #include <__clang_cuda_complex_builtins.h> +#include <__clang_gpu_device_virtual_trap.h> // curand_mtgp32_kernel helpfully redeclares blockDim and threadIdx in host // mode, giving them their "proper" types of dim3 and uint3. This is diff --git a/clang/lib/Headers/__clang_gpu_device_virtual_trap.h b/clang/lib/Headers/__clang_gpu_device_virtual_trap.h new file mode 100644 index 0000000000000..0f212d7bc3994 --- /dev/null +++ b/clang/lib/Headers/__clang_gpu_device_virtual_trap.h @@ -0,0 +1,39 @@ +//===---- __clang_gpu_device_virtual_trap.h - Virtual Trap Functions --------=== +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// +// This file provides device-side definitions of __cxa_pure_virtual() and +// __cxa_deleted_virtual(). +// +//===----------------------------------------------------------------------===// + +#ifndef __CLANG_GPU_DEVICE_VIRTUAL_TRAP_H__ +#define __CLANG_GPU_DEVICE_VIRTUAL_TRAP_H__ + +#if defined(__CUDA__) || defined(__HIP__) + +#ifdef __cplusplus +extern "C" { +__attribute__((__visibility__("default"))) +__attribute__((weak)) +__attribute__((noreturn)) +__device__ void __cxa_pure_virtual(void) { + __builtin_trap(); +} + +__attribute__((__visibility__("default"))) +__attribute__((weak)) +__attribute__((noreturn)) +__device__ void __cxa_deleted_virtual(void) { + __builtin_trap(); +} +} // extern "C" +#endif //__cplusplus + +#endif // defined(__HIP__) || defined(__CUDA__) + +#endif // __CLANG_GPU_DEVICE_VIRTUAL_TRAP_H__ diff --git a/clang/lib/Headers/__clang_hip_runtime_wrapper.h b/clang/lib/Headers/__clang_hip_runtime_wrapper.h index 4b8cffd86f044..a94acf375b20e 100644 --- a/clang/lib/Headers/__clang_hip_runtime_wrapper.h +++ b/clang/lib/Headers/__clang_hip_runtime_wrapper.h @@ -32,22 +32,7 @@ #define nullptr NULL; #endif -#ifdef __cplusplus -extern "C" { - __attribute__((__visibility__("default"))) - __attribute__((weak)) - __attribute__((noreturn)) - __device__ void __cxa_pure_virtual(void) { - __builtin_trap(); - } - __attribute__((__visibility__("default"))) - __attribute__((weak)) - __attribute__((noreturn)) - __device__ void __cxa_deleted_virtual(void) { - __builtin_trap(); - } -} -#endif //__cplusplus +#include <__clang_gpu_device_virtual_trap.h> #if !defined(__HIPCC_RTC__) #if __has_include("hip/hip_version.h") diff --git a/clang/test/CodeGenCUDA/pure_deleted_virtual.cu b/clang/test/CodeGenCUDA/pure_deleted_virtual.cu new file mode 100644 index 0000000000000..18a12d84c797a --- /dev/null +++ b/clang/test/CodeGenCUDA/pure_deleted_virtual.cu @@ -0,0 +1,9 @@ +// RUN: %clang --cuda-device-only -S -emit-llvm -o - %s 2>&1 | FileCheck %s + +// Check that __cxa_pure_virtual() and __cxa_deleted_virtual() are always +// available in device code. These functions are defined in a header included +// by __clang_cuda_runtime_wrapper.h, so use the driver here rather than +// invoking the frontend directly to make sure they are pulled in. + +// CHECK-DAG: define weak {{.*}} void @__cxa_pure_virtual() +// CHECK-DAG: define weak {{.*}} void @__cxa_deleted_virtual() >From 3fd7b4f03ada901185897ae27d455b3ff1e33644 Mon Sep 17 00:00:00 2001 From: Ambrose Leeb <[email protected]> Date: Wed, 19 Aug 2026 20:42:22 +0200 Subject: [PATCH 2/2] clang-format --- clang/lib/Headers/__clang_cuda_runtime_wrapper.h | 2 +- .../lib/Headers/__clang_gpu_device_virtual_trap.h | 14 ++++++-------- 2 files changed, 7 insertions(+), 9 deletions(-) diff --git a/clang/lib/Headers/__clang_cuda_runtime_wrapper.h b/clang/lib/Headers/__clang_cuda_runtime_wrapper.h index efd5016c2847d..82cb1f6d50537 100644 --- a/clang/lib/Headers/__clang_cuda_runtime_wrapper.h +++ b/clang/lib/Headers/__clang_cuda_runtime_wrapper.h @@ -489,8 +489,8 @@ __device__ inline __cuda_builtin_gridDim_t::operator uint3() const { } #include <__clang_cuda_cmath.h> -#include <__clang_cuda_intrinsics.h> #include <__clang_cuda_complex_builtins.h> +#include <__clang_cuda_intrinsics.h> #include <__clang_gpu_device_virtual_trap.h> // curand_mtgp32_kernel helpfully redeclares blockDim and threadIdx in host diff --git a/clang/lib/Headers/__clang_gpu_device_virtual_trap.h b/clang/lib/Headers/__clang_gpu_device_virtual_trap.h index 0f212d7bc3994..f7c92e30551b1 100644 --- a/clang/lib/Headers/__clang_gpu_device_virtual_trap.h +++ b/clang/lib/Headers/__clang_gpu_device_virtual_trap.h @@ -18,17 +18,15 @@ #ifdef __cplusplus extern "C" { -__attribute__((__visibility__("default"))) -__attribute__((weak)) -__attribute__((noreturn)) -__device__ void __cxa_pure_virtual(void) { +__attribute__((__visibility__("default"))) __attribute__((weak)) +__attribute__((noreturn)) __device__ void +__cxa_pure_virtual(void) { __builtin_trap(); } -__attribute__((__visibility__("default"))) -__attribute__((weak)) -__attribute__((noreturn)) -__device__ void __cxa_deleted_virtual(void) { +__attribute__((__visibility__("default"))) __attribute__((weak)) +__attribute__((noreturn)) __device__ void +__cxa_deleted_virtual(void) { __builtin_trap(); } } // extern "C" _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
