[clang] [CIR][CUDA][HIP] Exclude wrong-side virtual functions from vtables (PR #228433)
Steffen Larsen via cfe-commits
cfe-commits at lists.llvm.org
Sun Oct 4 23:24:32 PDT 2026
https://github.com/steffenlarsen updated https://github.com/llvm/llvm-project/pull/228433
>From 120620e8ca12ba77345f14b18f7ac869c1e9e5c1 Mon Sep 17 00:00:00 2001
From: Steffen Holst Larsen <sholstla at amd.com>
Date: Fri, 2 Oct 2026 07:49:09 -0500
Subject: [PATCH] [CIR][CUDA][HIP] Exclude wrong-side virtual functions from
vtables
OGCG builds the vtables of a CUDA/HIP compilation for the side
being compiled, that is a slot whose virtual function cannot be emitted
on that side is null.
Port both parts of CodeGenVTables::addVTableComponent to
CIRGenVTables::getVTableComponent. As in OGCG, a null slot that holds a
thunk still advances the thunk index, so later thunks keep their slots.
Assisted-by: Claude Opus 5.5
Signed-off-by: Steffen Holst Larsen <sholstla at amd.com>
---
clang/lib/CIR/CodeGen/CIRGenVTables.cpp | 44 +++++++++-
.../CIR/CodeGenCUDA/vtable-host-device.cu | 81 +++++++++++++++++++
2 files changed, 121 insertions(+), 4 deletions(-)
create mode 100644 clang/test/CIR/CodeGenCUDA/vtable-host-device.cu
diff --git a/clang/lib/CIR/CodeGen/CIRGenVTables.cpp b/clang/lib/CIR/CodeGen/CIRGenVTables.cpp
index 123246f3b3d3f..7c51751016044 100644
--- a/clang/lib/CIR/CodeGen/CIRGenVTables.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenVTables.cpp
@@ -174,7 +174,28 @@ mlir::Attribute CIRGenVTables::getVTableComponent(
cgm.getASTContext().getTargetInfo().emitVectorDeletingDtors(
cgm.getASTContext().getLangOpts()));
- assert(!cir::MissingFeatures::cudaSupport());
+ const bool isThunk =
+ nextVTableThunkIndex < layout.vtable_thunks().size() &&
+ layout.vtable_thunks()[nextVTableThunkIndex].first == componentIndex;
+
+ if (cgm.getLangOpts().CUDA) {
+ // Emit NULL for methods we can't codegen on this
+ // side. Otherwise we'd end up with vtable with unresolved
+ // references.
+ const CXXMethodDecl *md = cast<CXXMethodDecl>(gd.getDecl());
+ // OK on device side: functions w/ __device__ attribute
+ // OK on host side: anything except __device__-only functions.
+ bool canEmitMethod =
+ cgm.getLangOpts().CUDAIsDevice
+ ? md->hasAttr<CUDADeviceAttr>()
+ : (md->hasAttr<CUDAHostAttr>() || !md->hasAttr<CUDADeviceAttr>());
+ if (!canEmitMethod) {
+ if (isThunk)
+ nextVTableThunkIndex++;
+ return builder.getConstNullPtrAttr(builder.getUInt8PtrTy());
+ }
+ // Method is acceptable, continue processing as usual.
+ }
auto getSpecialVirtFn = [&](StringRef name) -> cir::FuncOp {
assert(!cir::MissingFeatures::vtableRelativeLayout());
@@ -189,6 +210,23 @@ mlir::Attribute CIRGenVTables::getVTableComponent(
cir::FuncOp fnPtr = cgm.createRuntimeFunction(fnTy, name);
assert(!cir::MissingFeatures::opGlobalUnnamedAddr());
+
+ // The Microsoft ABI uses the same function name for pure and deleted
+ // virtual functions.
+ if (!fnPtr.isDeclaration())
+ return fnPtr;
+
+ // For device compilation, provide a weak definition that traps,
+ // otherwise linking ends up with unresolved references.
+ if (cgm.getLangOpts().isTargetDevice()) {
+ fnPtr.setLinkage(cir::GlobalLinkageKind::WeakAnyLinkage);
+ mlir::SymbolTable::setSymbolVisibility(
+ fnPtr, cgm.getMLIRVisibilityFromCIRLinkage(fnPtr.getLinkage()));
+ mlir::OpBuilder::InsertionGuard guard(builder);
+ builder.setInsertionPointToStart(fnPtr.addEntryBlock());
+ cir::TrapOp::create(builder, fnPtr.getLoc());
+ }
+
return fnPtr;
};
@@ -203,9 +241,7 @@ mlir::Attribute CIRGenVTables::getVTableComponent(
deletedVirtualFn =
getSpecialVirtFn(cgm.getCXXABI().getDeletedVirtualCallName());
fnPtr = deletedVirtualFn;
- } else if (nextVTableThunkIndex < layout.vtable_thunks().size() &&
- layout.vtable_thunks()[nextVTableThunkIndex].first ==
- componentIndex) {
+ } else if (isThunk) {
const ThunkInfo &thunkInfo =
layout.vtable_thunks()[nextVTableThunkIndex].second;
nextVTableThunkIndex++;
diff --git a/clang/test/CIR/CodeGenCUDA/vtable-host-device.cu b/clang/test/CIR/CodeGenCUDA/vtable-host-device.cu
new file mode 100644
index 0000000000000..fa2dc88704a56
--- /dev/null
+++ b/clang/test/CIR/CodeGenCUDA/vtable-host-device.cu
@@ -0,0 +1,81 @@
+#include "Inputs/cuda.h"
+
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fcuda-is-device -fclangir -emit-cir %s -o - \
+// RUN: | FileCheck %s --check-prefix=CIR-DEV --implicit-check-not=@_ZNK6Square4nameEv --implicit-check-not=@_ZN1C9host_onlyEv
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fcuda-is-device -fclangir -emit-llvm %s -o - \
+// RUN: | FileCheck %s --check-prefix=DEV --implicit-check-not=@_ZNK6Square4nameEv --implicit-check-not=@_ZN1C9host_onlyEv
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fcuda-is-device -emit-llvm %s -o - \
+// RUN: | FileCheck %s --check-prefix=DEV --implicit-check-not=@_ZNK6Square4nameEv --implicit-check-not=@_ZN1C9host_onlyEv
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -x cuda -fcuda-is-device -fclangir -emit-llvm %s -o - \
+// RUN: | FileCheck %s --check-prefix=DEV --implicit-check-not=@_ZNK6Square4nameEv --implicit-check-not=@_ZN1C9host_onlyEv
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -x cuda -fcuda-is-device -emit-llvm %s -o - \
+// RUN: | FileCheck %s --check-prefix=DEV --implicit-check-not=@_ZNK6Square4nameEv --implicit-check-not=@_ZN1C9host_onlyEv
+
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -x hip -fhip-new-launch-api -fclangir -emit-llvm %s -o - \
+// RUN: | FileCheck %s --check-prefix=HOST --implicit-check-not=@_ZNK6Square4areaEv
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -x hip -fhip-new-launch-api -emit-llvm %s -o - \
+// RUN: | FileCheck %s --check-prefix=HOST --implicit-check-not=@_ZNK6Square4areaEv
+
+// A vtable has null in the slots of virtual functions that can't be compiled
+// on the current side, so those functions aren't referenced or emitted there.
+// On the device, pure and deleted virtual functions are weak functions that
+// trap, since no device library provides them.
+
+struct Shape {
+ __host__ __device__ virtual ~Shape() {}
+ __device__ virtual float area() const = 0;
+ virtual const char *name() const = 0;
+ __host__ __device__ virtual void removed() = delete;
+};
+
+struct Square : Shape {
+ float s;
+ __host__ __device__ Square(float s) : s(s) {}
+ __device__ float area() const override { return s * s; }
+ const char *name() const override { return "square"; }
+};
+
+__global__ void kernel(float *out) {
+ Square sq(3.0f);
+ *out = sq.area();
+}
+
+const char *host_name() { return Square(3.0f).name(); }
+
+// CIR-DEV: cir.global {{.*}} @_ZTV6Square = #cir.vtable<{#cir.const_array<[#cir.ptr<null> : !cir.ptr<!u8i>, #cir.ptr<null> : !cir.ptr<!u8i>, #cir.global_view<@_ZN6SquareD1Ev> : !cir.ptr<!u8i>, #cir.global_view<@_ZN6SquareD0Ev> : !cir.ptr<!u8i>, #cir.global_view<@_ZNK6Square4areaEv> : !cir.ptr<!u8i>, #cir.ptr<null> : !cir.ptr<!u8i>, #cir.global_view<@__cxa_deleted_virtual> : !cir.ptr<!u8i>]>
+// CIR-DEV: cir.global {{.*}} @_ZTV5Shape = #cir.vtable<{#cir.const_array<[#cir.ptr<null> : !cir.ptr<!u8i>, #cir.ptr<null> : !cir.ptr<!u8i>, #cir.global_view<@_ZN5ShapeD1Ev> : !cir.ptr<!u8i>, #cir.global_view<@_ZN5ShapeD0Ev> : !cir.ptr<!u8i>, #cir.global_view<@__cxa_pure_virtual> : !cir.ptr<!u8i>, #cir.ptr<null> : !cir.ptr<!u8i>, #cir.global_view<@__cxa_deleted_virtual> : !cir.ptr<!u8i>]>
+
+// DEV-DAG: @_ZTV6Square = {{.*}}[ptr {{[^,]*}}null, ptr {{[^,]*}}null, ptr {{[^,]*}}@_ZN6SquareD1Ev{{[^,]*}}, ptr {{[^,]*}}@_ZN6SquareD0Ev{{[^,]*}}, ptr {{[^,]*}}@_ZNK6Square4areaEv{{[^,]*}}, ptr {{[^,]*}}null, ptr {{[^,]*}}@__cxa_deleted_virtual{{[^,]*}}]
+// DEV-DAG: @_ZTV5Shape = {{.*}}[ptr {{[^,]*}}null, ptr {{[^,]*}}null, ptr {{[^,]*}}@_ZN5ShapeD1Ev{{[^,]*}}, ptr {{[^,]*}}@_ZN5ShapeD0Ev{{[^,]*}}, ptr {{[^,]*}}@__cxa_pure_virtual{{[^,]*}}, ptr {{[^,]*}}null, ptr {{[^,]*}}@__cxa_deleted_virtual{{[^,]*}}]
+
+// HOST-DAG: @_ZTV6Square = {{.*}}[ptr null, ptr @_ZTI6Square, ptr @_ZN6SquareD1Ev, ptr @_ZN6SquareD0Ev, ptr null, ptr @_ZNK6Square4nameEv, ptr @__cxa_deleted_virtual]
+// HOST-DAG: @_ZTV5Shape = {{.*}}[ptr null, ptr @_ZTI5Shape, ptr @_ZN5ShapeD1Ev, ptr @_ZN5ShapeD0Ev, ptr null, ptr @__cxa_pure_virtual, ptr @__cxa_deleted_virtual]
+// HOST-DAG: declare {{.*}}void @__cxa_pure_virtual()
+// HOST-DAG: declare {{.*}}void @__cxa_deleted_virtual()
+
+// CIR-DEV: cir.func weak {{.*}}@__cxa_deleted_virtual() {
+// CIR-DEV-NEXT: cir.trap
+// CIR-DEV-NEXT: }
+// CIR-DEV: cir.func weak {{.*}}@__cxa_pure_virtual() {
+// CIR-DEV-NEXT: cir.trap
+// CIR-DEV-NEXT: }
+
+// DEV-DAG: define weak {{.*}}void @__cxa_pure_virtual()
+// DEV-DAG: define weak {{.*}}void @__cxa_deleted_virtual()
+
+// A host-only override that needs a thunk since its slot is null and the thunk
+// index still advances, so the next thunk lands in its own slot.
+struct A {
+ virtual void host_only() = 0;
+ __device__ virtual void f1() = 0;
+};
+struct B {
+ __device__ virtual void f2() {}
+};
+struct C : B, A {
+ void host_only() override {}
+ __device__ void f1() override {}
+};
+__device__ void use_c() { C c; }
+
+// DEV-DAG: @_ZTV1C = {{.*}}[ptr {{[^,]*}}null, ptr {{[^,]*}}null, ptr {{[^,]*}}@_ZN1B2f2Ev{{[^,]*}}, ptr {{[^,]*}}null, ptr {{[^,]*}}@_ZN1C2f1Ev{{[^,]*}}], [4 x ptr{{[^]]*}}] [ptr {{[^,]*}}inttoptr (i64 -8 to ptr{{[^,]*}}), ptr {{[^,]*}}null, ptr {{[^,]*}}null, ptr {{[^,]*}}@_ZThn8_N1C2f1Ev{{[^,]*}}]
More information about the cfe-commits
mailing list