[clang] [Clang] [CUDA] Provide device-side definitions of `__cxa_[pure|deleted]_virtual()` (PR #217427)
Ambrose Leeb via cfe-commits
cfe-commits at lists.llvm.org
Fri Sep 11 09:48:33 PDT 2026
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 <aleeb at nvidia.com>
Date: Wed, 19 Aug 2026 20:29:26 +0200
Subject: [PATCH 1/8] [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 <aleeb at nvidia.com>
Date: Wed, 19 Aug 2026 20:42:22 +0200
Subject: [PATCH 2/8] 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"
>From 6019ffad4b4e5b24304867f650aee61dbcf76373 Mon Sep 17 00:00:00 2001
From: Ambrose Leeb <aleeb at nvidia.com>
Date: Thu, 3 Sep 2026 20:50:09 +0200
Subject: [PATCH 3/8] Emit functions in codegen
---
clang/lib/CodeGen/CGVTables.cpp | 33 ++++++++++++++-----
.../test/CodeGenCUDA/pure_deleted_virtual.cu | 22 +++++++++----
.../nvptx_target_pure_deleted_codegen.cpp | 3 +-
3 files changed, 42 insertions(+), 16 deletions(-)
diff --git a/clang/lib/CodeGen/CGVTables.cpp b/clang/lib/CodeGen/CGVTables.cpp
index 2d09ec90c013f..80014d9c89b82 100644
--- a/clang/lib/CodeGen/CGVTables.cpp
+++ b/clang/lib/CodeGen/CGVTables.cpp
@@ -829,18 +829,33 @@ void CodeGenVTables::addVTableComponent(ConstantArrayBuilder &builder,
if (RelativeCXXABIVTables)
return llvm::ConstantPointerNull::get(CGM.GlobalsInt8PtrTy);
- // For NVPTX devices in OpenMP emit special functon as null pointers,
- // otherwise linking ends up with unresolved references.
- if (CGM.getLangOpts().OpenMP && CGM.getLangOpts().OpenMPIsTargetDevice &&
- CGM.getTriple().isNVPTX())
- return llvm::ConstantPointerNull::get(CGM.GlobalsInt8PtrTy);
llvm::FunctionType *fnTy =
llvm::FunctionType::get(CGM.VoidTy, /*isVarArg=*/false);
- llvm::Constant *fn = cast<llvm::Constant>(
+ auto *F = cast<llvm::Function>(
CGM.CreateRuntimeFunction(fnTy, name).getCallee());
- if (auto f = dyn_cast<llvm::Function>(fn))
- f->setUnnamedAddr(llvm::GlobalValue::UnnamedAddr::Global);
- return fn;
+ F->setUnnamedAddr(llvm::GlobalValue::UnnamedAddr::Global);
+
+ // The Microsoft ABI uses the same function name for pure and deleted
+ // virtual functions.
+ if (!F->empty())
+ return F;
+
+ // For NVPTX devices in OpenMP and CUDA, provide a weak definition that
+ // traps, otherwise linking ends up with unresolved references.
+ if ((CGM.getLangOpts().OpenMP &&
+ CGM.getLangOpts().OpenMPIsTargetDevice &&
+ CGM.getTriple().isNVPTX()) ||
+ (CGM.getLangOpts().CUDA && CGM.getLangOpts().CUDAIsDevice)) {
+ F->setLinkage(llvm::GlobalValue::WeakAnyLinkage);
+ CodeGenFunction CGF(CGM);
+ const CGFunctionInfo &FI = CGM.getTypes().arrangeNullaryFunction();
+ CGF.StartFunction(GlobalDecl(), CGM.getContext().VoidTy, F, FI,
+ FunctionArgList{});
+ CGF.EmitTrapCallAndMakeUnreachable();
+ CGF.FinishFunction();
+ }
+
+ return F;
};
llvm::Constant *fnPtr;
diff --git a/clang/test/CodeGenCUDA/pure_deleted_virtual.cu b/clang/test/CodeGenCUDA/pure_deleted_virtual.cu
index 18a12d84c797a..b269555268b91 100644
--- a/clang/test/CodeGenCUDA/pure_deleted_virtual.cu
+++ b/clang/test/CodeGenCUDA/pure_deleted_virtual.cu
@@ -1,9 +1,19 @@
-// RUN: %clang --cuda-device-only -S -emit-llvm -o - %s 2>&1 | FileCheck %s
+// RUN: %clang_cc1 -emit-llvm %s -o - -fcuda-is-device -triple nvptx64 | 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.
+// available in device code.
-// CHECK-DAG: define weak {{.*}} void @__cxa_pure_virtual()
-// CHECK-DAG: define weak {{.*}} void @__cxa_deleted_virtual()
+#define __device__ __attribute__((__device__))
+
+struct S {
+ __device__ virtual void anchor();
+ __device__ virtual void pure() = 0;
+ __device__ virtual void deleted() = delete;
+};
+
+// Anchor function to force vtable emission.
+__device__ void S::anchor() {}
+
+// CHECK-DAG: @_ZTV1S = {{.*}} constant { [5 x ptr] } { [5 x ptr] [ptr null, ptr null, ptr @_ZN1S6anchorEv, ptr @__cxa_pure_virtual, ptr @__cxa_deleted_virtual] }
+// CHECK-DAG: define weak{{.*}} void @__cxa_pure_virtual()
+// CHECK-DAG: define weak{{.*}} void @__cxa_deleted_virtual()
diff --git a/clang/test/OpenMP/nvptx_target_pure_deleted_codegen.cpp b/clang/test/OpenMP/nvptx_target_pure_deleted_codegen.cpp
index 071a501b66398..b5187aeaa26e0 100644
--- a/clang/test/OpenMP/nvptx_target_pure_deleted_codegen.cpp
+++ b/clang/test/OpenMP/nvptx_target_pure_deleted_codegen.cpp
@@ -11,7 +11,8 @@
// CHECK-NOT: class_type_info
// CHECK-DAG: @_ZTV7Derived = linkonce_odr protected unnamed_addr constant { [3 x ptr] } { [3 x ptr] [ptr null, ptr null, ptr @_ZN7Derived3fooEv] }
-// CHECK-DAG: @_ZTV4Base = linkonce_odr protected unnamed_addr constant { [3 x ptr] } zeroinitializer
+// CHECK-DAG: @_ZTV4Base = linkonce_odr protected unnamed_addr constant { [3 x ptr] } { [3 x ptr] [ptr null, ptr null, ptr @__cxa_pure_virtual] }
+// CHECK-DAG: define weak{{.*}} void @__cxa_pure_virtual()
// CHECK-NOT: class_type_info
class Base {
public:
>From 61e8f0f68b0e1d0b0f8247c51138e8b50b9f78be Mon Sep 17 00:00:00 2001
From: Ambrose Leeb <aleeb at nvidia.com>
Date: Thu, 3 Sep 2026 20:53:35 +0200
Subject: [PATCH 4/8] remove headers
---
clang/lib/Headers/CMakeLists.txt | 1 -
.../Headers/__clang_cuda_runtime_wrapper.h | 1 -
.../Headers/__clang_gpu_device_virtual_trap.h | 37 -------------------
.../lib/Headers/__clang_hip_runtime_wrapper.h | 17 ++++++++-
4 files changed, 16 insertions(+), 40 deletions(-)
delete mode 100644 clang/lib/Headers/__clang_gpu_device_virtual_trap.h
diff --git a/clang/lib/Headers/CMakeLists.txt b/clang/lib/Headers/CMakeLists.txt
index 67376781bb8ba..3d845423759ac 100644
--- a/clang/lib/Headers/CMakeLists.txt
+++ b/clang/lib/Headers/CMakeLists.txt
@@ -311,7 +311,6 @@ 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 82cb1f6d50537..3871958e00ea1 100644
--- a/clang/lib/Headers/__clang_cuda_runtime_wrapper.h
+++ b/clang/lib/Headers/__clang_cuda_runtime_wrapper.h
@@ -491,7 +491,6 @@ __device__ inline __cuda_builtin_gridDim_t::operator uint3() const {
#include <__clang_cuda_cmath.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
// 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
deleted file mode 100644
index f7c92e30551b1..0000000000000
--- a/clang/lib/Headers/__clang_gpu_device_virtual_trap.h
+++ /dev/null
@@ -1,37 +0,0 @@
-//===---- __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 a94acf375b20e..4b8cffd86f044 100644
--- a/clang/lib/Headers/__clang_hip_runtime_wrapper.h
+++ b/clang/lib/Headers/__clang_hip_runtime_wrapper.h
@@ -32,7 +32,22 @@
#define nullptr NULL;
#endif
-#include <__clang_gpu_device_virtual_trap.h>
+#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
#if !defined(__HIPCC_RTC__)
#if __has_include("hip/hip_version.h")
>From 2616b3fa58b1a51ed47e794cead2533349ec00a8 Mon Sep 17 00:00:00 2001
From: Ambrose Leeb <aleeb at nvidia.com>
Date: Thu, 3 Sep 2026 21:04:20 +0200
Subject: [PATCH 5/8] reorder headers
---
clang/lib/Headers/__clang_cuda_runtime_wrapper.h | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/clang/lib/Headers/__clang_cuda_runtime_wrapper.h b/clang/lib/Headers/__clang_cuda_runtime_wrapper.h
index 3871958e00ea1..29178bf0db8a1 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_complex_builtins.h>
#include <__clang_cuda_intrinsics.h>
+#include <__clang_cuda_complex_builtins.h>
// curand_mtgp32_kernel helpfully redeclares blockDim and threadIdx in host
// mode, giving them their "proper" types of dim3 and uint3. This is
>From 689d47e3dca43028fdf7de4dbc4074f152e3098b Mon Sep 17 00:00:00 2001
From: Ambrose Leeb <aleeb at nvidia.com>
Date: Thu, 3 Sep 2026 23:47:33 +0200
Subject: [PATCH 6/8] Apply patch to all device targets
---
clang/lib/CodeGen/CGVTables.cpp | 7 ++-----
clang/lib/Headers/__clang_hip_runtime_wrapper.h | 17 -----------------
clang/test/CodeGenCUDA/pure_deleted_virtual.cu | 7 +++++--
3 files changed, 7 insertions(+), 24 deletions(-)
diff --git a/clang/lib/CodeGen/CGVTables.cpp b/clang/lib/CodeGen/CGVTables.cpp
index 80014d9c89b82..ba4e90c0f95ba 100644
--- a/clang/lib/CodeGen/CGVTables.cpp
+++ b/clang/lib/CodeGen/CGVTables.cpp
@@ -840,12 +840,9 @@ void CodeGenVTables::addVTableComponent(ConstantArrayBuilder &builder,
if (!F->empty())
return F;
- // For NVPTX devices in OpenMP and CUDA, provide a weak definition that
+ // For device compilation, provide a weak definition that
// traps, otherwise linking ends up with unresolved references.
- if ((CGM.getLangOpts().OpenMP &&
- CGM.getLangOpts().OpenMPIsTargetDevice &&
- CGM.getTriple().isNVPTX()) ||
- (CGM.getLangOpts().CUDA && CGM.getLangOpts().CUDAIsDevice)) {
+ if (CGM.getLangOpts().isTargetDevice()) {
F->setLinkage(llvm::GlobalValue::WeakAnyLinkage);
CodeGenFunction CGF(CGM);
const CGFunctionInfo &FI = CGM.getTypes().arrangeNullaryFunction();
diff --git a/clang/lib/Headers/__clang_hip_runtime_wrapper.h b/clang/lib/Headers/__clang_hip_runtime_wrapper.h
index 4b8cffd86f044..75c61b3319094 100644
--- a/clang/lib/Headers/__clang_hip_runtime_wrapper.h
+++ b/clang/lib/Headers/__clang_hip_runtime_wrapper.h
@@ -32,23 +32,6 @@
#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
-
#if !defined(__HIPCC_RTC__)
#if __has_include("hip/hip_version.h")
#include "hip/hip_version.h"
diff --git a/clang/test/CodeGenCUDA/pure_deleted_virtual.cu b/clang/test/CodeGenCUDA/pure_deleted_virtual.cu
index b269555268b91..110a0a349dfa0 100644
--- a/clang/test/CodeGenCUDA/pure_deleted_virtual.cu
+++ b/clang/test/CodeGenCUDA/pure_deleted_virtual.cu
@@ -1,4 +1,5 @@
-// RUN: %clang_cc1 -emit-llvm %s -o - -fcuda-is-device -triple nvptx64 | FileCheck %s
+// RUN: %clang_cc1 -emit-llvm %s -o - -fcuda-is-device -triple nvptx64 | FileCheck %s --check-prefix=CHECK,CUDA
+// RUN: %clang_cc1 -emit-llvm %s -o - -fcuda-is-device -triple amdgpu-amd-amdhsa | FileCheck %s --check-prefix=CHECK,HIP
// Check that __cxa_pure_virtual() and __cxa_deleted_virtual() are always
// available in device code.
@@ -14,6 +15,8 @@ struct S {
// Anchor function to force vtable emission.
__device__ void S::anchor() {}
-// CHECK-DAG: @_ZTV1S = {{.*}} constant { [5 x ptr] } { [5 x ptr] [ptr null, ptr null, ptr @_ZN1S6anchorEv, ptr @__cxa_pure_virtual, ptr @__cxa_deleted_virtual] }
+// CUDA-DAG: @_ZTV1S = {{.*}} constant { [5 x ptr] } { [5 x ptr] [ptr null, ptr null, ptr @_ZN1S6anchorEv, ptr @__cxa_pure_virtual, ptr @__cxa_deleted_virtual] }
+// HIP-DAG: @_ZTV1S = {{.*}} constant { [5 x ptr addrspace(1)] } { [5 x ptr addrspace(1)] [ptr addrspace(1) null, ptr addrspace(1) null, ptr addrspace(1) addrspacecast (ptr @_ZN1S6anchorEv to ptr addrspace(1)), ptr addrspace(1) addrspacecast (ptr @__cxa_pure_virtual to ptr addrspace(1)), ptr addrspace(1) addrspacecast (ptr @__cxa_deleted_virtual to ptr addrspace(1))] }, align 8
+
// CHECK-DAG: define weak{{.*}} void @__cxa_pure_virtual()
// CHECK-DAG: define weak{{.*}} void @__cxa_deleted_virtual()
>From 7aea1cab2daf1c29a3f6263d4f72596100379ac6 Mon Sep 17 00:00:00 2001
From: Ambrose Leeb <aleeb at nvidia.com>
Date: Fri, 4 Sep 2026 20:20:26 +0200
Subject: [PATCH 7/8] fix HIP test
---
clang/test/Headers/hip-header.hip | 5 +++--
1 file changed, 3 insertions(+), 2 deletions(-)
diff --git a/clang/test/Headers/hip-header.hip b/clang/test/Headers/hip-header.hip
index e5693b5c3aa35..2471c15b6c626 100644
--- a/clang/test/Headers/hip-header.hip
+++ b/clang/test/Headers/hip-header.hip
@@ -92,8 +92,6 @@ __device__ void test_vf() {
}
// CHECK: @_ZTV7derived = linkonce_odr unnamed_addr addrspace(1) constant { [4 x ptr addrspace(1)] } { [4 x ptr addrspace(1)] [ptr addrspace(1) null, ptr addrspace(1) null, ptr addrspace(1) addrspacecast (ptr @_ZN7derived2pvEv to ptr addrspace(1)), ptr addrspace(1) addrspacecast (ptr @__cxa_deleted_virtual to ptr addrspace(1))] }, comdat, align 8
// CHECK: @_ZTV4base = linkonce_odr unnamed_addr addrspace(1) constant { [4 x ptr addrspace(1)] } { [4 x ptr addrspace(1)] [ptr addrspace(1) null, ptr addrspace(1) null, ptr addrspace(1) addrspacecast (ptr @__cxa_pure_virtual to ptr addrspace(1)), ptr addrspace(1) addrspacecast (ptr @__cxa_deleted_virtual to ptr addrspace(1))] }, comdat, align 8
-// CHECK: define{{.*}}void @__cxa_pure_virtual()
-// CHECK: define{{.*}}void @__cxa_deleted_virtual()
struct Number {
__device__ Number(float _x) : x(_x) {}
@@ -201,3 +199,6 @@ void test_malloc_host(void *a) {
a = std::malloc(42);
std::free(a);
}
+
+// CHECK-DAG: define weak{{.*}} void @__cxa_pure_virtual()
+// CHECK-DAG: define weak{{.*}} void @__cxa_deleted_virtual()
>From e0734c77049c2a87608f371784714056be3f23ab Mon Sep 17 00:00:00 2001
From: Ambrose Leeb <aleeb at nvidia.com>
Date: Fri, 11 Sep 2026 18:48:16 +0200
Subject: [PATCH 8/8] add spirv64 test
---
clang/test/CodeGenCUDA/pure_deleted_virtual.cu | 2 ++
1 file changed, 2 insertions(+)
diff --git a/clang/test/CodeGenCUDA/pure_deleted_virtual.cu b/clang/test/CodeGenCUDA/pure_deleted_virtual.cu
index 110a0a349dfa0..f158a91a5e3ef 100644
--- a/clang/test/CodeGenCUDA/pure_deleted_virtual.cu
+++ b/clang/test/CodeGenCUDA/pure_deleted_virtual.cu
@@ -1,5 +1,6 @@
// RUN: %clang_cc1 -emit-llvm %s -o - -fcuda-is-device -triple nvptx64 | FileCheck %s --check-prefix=CHECK,CUDA
// RUN: %clang_cc1 -emit-llvm %s -o - -fcuda-is-device -triple amdgpu-amd-amdhsa | FileCheck %s --check-prefix=CHECK,HIP
+// RUN: %clang_cc1 -emit-llvm %s -o - -fcuda-is-device -triple spirv64-amd-amdhsa | FileCheck %s --check-prefix=CHECK,SPIRV64
// Check that __cxa_pure_virtual() and __cxa_deleted_virtual() are always
// available in device code.
@@ -17,6 +18,7 @@ __device__ void S::anchor() {}
// CUDA-DAG: @_ZTV1S = {{.*}} constant { [5 x ptr] } { [5 x ptr] [ptr null, ptr null, ptr @_ZN1S6anchorEv, ptr @__cxa_pure_virtual, ptr @__cxa_deleted_virtual] }
// HIP-DAG: @_ZTV1S = {{.*}} constant { [5 x ptr addrspace(1)] } { [5 x ptr addrspace(1)] [ptr addrspace(1) null, ptr addrspace(1) null, ptr addrspace(1) addrspacecast (ptr @_ZN1S6anchorEv to ptr addrspace(1)), ptr addrspace(1) addrspacecast (ptr @__cxa_pure_virtual to ptr addrspace(1)), ptr addrspace(1) addrspacecast (ptr @__cxa_deleted_virtual to ptr addrspace(1))] }, align 8
+// SPIRV64-DAG: @_ZTV1S = {{.*}} constant { [5 x ptr addrspace(1)] } { [5 x ptr addrspace(1)] [ptr addrspace(1) null, ptr addrspace(1) null, ptr addrspace(1) addrspacecast (ptr addrspace(4) @_ZN1S6anchorEv to ptr addrspace(1)), ptr addrspace(1) addrspacecast (ptr addrspace(4) @__cxa_pure_virtual to ptr addrspace(1)), ptr addrspace(1) addrspacecast (ptr addrspace(4) @__cxa_deleted_virtual to ptr addrspace(1))] }, align 8
// CHECK-DAG: define weak{{.*}} void @__cxa_pure_virtual()
// CHECK-DAG: define weak{{.*}} void @__cxa_deleted_virtual()
More information about the cfe-commits
mailing list