From 260efca649a18c110aeda9e618c0c6772c77ceeb Mon Sep 17 00:00:00 2001
From: Anshil Gandhi <Anshil.Gandhi@amd.com>
Date: Mon, 27 Jan 2025 23:37:08 +0000
Subject: [PATCH] [CUDA] Emit NULL in VTable in the last resort

Before adding a pointer to a deleting destructor
in the VTable, check if it is either a pure
virtual function, deleted virtual member function or a
function which desires an addition of thunk
functions.

Change-Id: I791acb42011d253a2a6b06e5f33090abdaa90872
---
 clang/lib/CodeGen/CGVTables.cpp | 35 +++++++++++++++++----------------
 1 file changed, 18 insertions(+), 17 deletions(-)

diff --git a/clang/lib/CodeGen/CGVTables.cpp b/clang/lib/CodeGen/CGVTables.cpp
index 7f729d359b82b..f42ab580e6897 100644
--- a/clang/lib/CodeGen/CGVTables.cpp
+++ b/clang/lib/CodeGen/CGVTables.cpp
@@ -772,23 +772,6 @@ void CodeGenVTables::addVTableComponent(ConstantArrayBuilder &builder,
   case VTableComponent::CK_DeletingDtorPointer: {
     GlobalDecl GD = component.getGlobalDecl();
 
-    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)
-        return builder.add(
-            llvm::ConstantExpr::getNullValue(CGM.GlobalsInt8PtrTy));
-      // Method is acceptable, continue processing as usual.
-    }
-
     auto getSpecialVirtualFn = [&](StringRef name) -> llvm::Constant * {
       // FIXME(PR43094): When merging comdat groups, lld can select a local
       // symbol as the signature symbol even though it cannot be accessed
@@ -845,6 +828,24 @@ void CodeGenVTables::addVTableComponent(ConstantArrayBuilder &builder,
 
     // Otherwise we can use the method definition directly.
     } else {
+
+      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)
+          return builder.add(
+              llvm::ConstantExpr::getNullValue(CGM.GlobalsInt8PtrTy));
+        // Method is acceptable, continue processing as usual.
+      }
+
       llvm::Type *fnTy = CGM.getTypes().GetFunctionTypeForVTable(GD);
       fnPtr = CGM.GetAddrOfFunction(GD, fnTy, /*ForVTable=*/true);
       if (CGM.getCodeGenOpts().PointerAuth.CXXVirtualFunctionPointers)