[CUDA] Emit NULL in VTable in the last resort

[gandhi56]

For device functions which cannot be emitted into a VTable, emit a null value only if the corresponding declaration is none of the following:

• pure virtual function
• deleted virtual member function
• a function for which thunks may be emitted.

[llvmbot]

@llvm/pr-subscribers-clang-codegen

@llvm/pr-subscribers-clang

Author: Anshil Gandhi (gandhi56)

Changes

For device functions which cannot be emitted into a VTable, emit a null value only if the corresponding declaration is none of the following:

• pure virtual function
• deleted virtual member function
• a function for which thunks may be emitted.

---

Full diff: https://github.com/llvm/llvm-project/pull/124687.diff

1 Files Affected:

• (modified) clang/lib/CodeGen/CGVTables.cpp (+18-17)

diff --git a/clang/lib/CodeGen/CGVTables.cpp b/clang/lib/CodeGen/CGVTables.cpp
index 7f729d359b82b3..f42ab580e6897b 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)