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/6] [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/6] 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 <[email protected]> Date: Thu, 3 Sep 2026 20:50:09 +0200 Subject: [PATCH 3/6] 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 <[email protected]> Date: Thu, 3 Sep 2026 20:53:35 +0200 Subject: [PATCH 4/6] 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 <[email protected]> Date: Thu, 3 Sep 2026 21:04:20 +0200 Subject: [PATCH 5/6] 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 <[email protected]> Date: Thu, 3 Sep 2026 23:47:33 +0200 Subject: [PATCH 6/6] 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() _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
