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: _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
