[clang] 78ae840 - [Clang] [CUDA] Provide device-side definitions of `__cxa_[pure|deleted]_virtual()` (#217427)
via cfe-commits
cfe-commits at lists.llvm.org
Fri Sep 11 12:12:28 PDT 2026
Author: Ambrose Leeb
Date: 2026-09-11T21:12:23+02:00
New Revision: 78ae84025be3bc2e711e803b79b8515075c97a80
URL: https://github.com/llvm/llvm-project/commit/78ae84025be3bc2e711e803b79b8515075c97a80
DIFF: https://github.com/llvm/llvm-project/commit/78ae84025be3bc2e711e803b79b8515075c97a80.diff
LOG: [Clang] [CUDA] Provide device-side definitions of `__cxa_[pure|deleted]_virtual()` (#217427)
Currently, we don't provide definitions of these on the device side, and
since there is also no C++ runtime on the device side, if a call to a
pure or deleted virtual function makes it into the final IR, the program
will fail to assemble because `ptxas` can’t find a definition of
`__cxa_pure_virtual()` or `__cxa_deleted_virtual()`.
This patch updates Clang to emit weak definitions of these functions
that trap.
Fixes #49183, fixes #67533.
Added:
clang/test/CodeGenCUDA/pure_deleted_virtual.cu
Modified:
clang/docs/ReleaseNotes.md
clang/lib/CodeGen/CGVTables.cpp
clang/lib/Headers/__clang_hip_runtime_wrapper.h
clang/test/Headers/hip-header.hip
clang/test/OpenMP/nvptx_target_pure_deleted_codegen.cpp
Removed:
################################################################################
diff --git a/clang/docs/ReleaseNotes.md b/clang/docs/ReleaseNotes.md
index ca29a33a2abba..26124b330a432 100644
--- a/clang/docs/ReleaseNotes.md
+++ b/clang/docs/ReleaseNotes.md
@@ -795,6 +795,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/CodeGen/CGVTables.cpp b/clang/lib/CodeGen/CGVTables.cpp
index 2d09ec90c013f..ba4e90c0f95ba 100644
--- a/clang/lib/CodeGen/CGVTables.cpp
+++ b/clang/lib/CodeGen/CGVTables.cpp
@@ -829,18 +829,30 @@ 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 device compilation, provide a weak definition that
+ // traps, otherwise linking ends up with unresolved references.
+ if (CGM.getLangOpts().isTargetDevice()) {
+ 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/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
new file mode 100644
index 0000000000000..f158a91a5e3ef
--- /dev/null
+++ b/clang/test/CodeGenCUDA/pure_deleted_virtual.cu
@@ -0,0 +1,24 @@
+// 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.
+
+#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() {}
+
+// 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()
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()
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:
More information about the cfe-commits
mailing list