[clang] 054cc3b - [CUDA][HIP] Fix store of vtbl in ctor

[Yaxun Liu via cfe-commits]

Author: Yaxun (Sam) Liu
Date: 2021-06-08T10:24:44-04:00
New Revision: 054cc3b1b469de4b0cb25d1dc3af43c679c5dc44

URL: https://github.com/llvm/llvm-project/commit/054cc3b1b469de4b0cb25d1dc3af43c679c5dc44
DIFF: https://github.com/llvm/llvm-project/commit/054cc3b1b469de4b0cb25d1dc3af43c679c5dc44.diff

LOG: [CUDA][HIP] Fix store of vtbl in ctor

vtbl itself is in default global address space. When clang emits
ctor, it gets a pointer to the vtbl field based on the this pointer,
then stores vtbl to the pointer.

Since this pointer can point to any address space (e.g. an object
created in stack), this pointer points to default address space, therefore
the pointer to vtbl field in this object should also be in default
address space.

Currently, clang incorrectly casts the pointer to vtbl field in this object
to global address space. This caused assertions in backend.

This patch fixes that by removing the incorrect addr space cast.

Reviewed by: Artem Belevich

Differential Revision: https://reviews.llvm.org/D103835

Added: 
    clang/test/CodeGenCUDA/vtbl.cu

Modified: 
    clang/lib/CodeGen/CGClass.cpp

Removed: 
    


################################################################################
diff  --git a/clang/lib/CodeGen/CGClass.cpp b/clang/lib/CodeGen/CGClass.cpp
index 50681da6608d4..3551c5e51f3a1 100644
--- a/clang/lib/CodeGen/CGClass.cpp
+++ b/clang/lib/CodeGen/CGClass.cpp
@@ -2518,8 +2518,10 @@ void CodeGenFunction::InitializeVTablePointer(const VPtr &Vptr) {
       llvm::FunctionType::get(CGM.Int32Ty, /*isVarArg=*/true)
           ->getPointerTo(ProgAS)
           ->getPointerTo(GlobalsAS);
+  // vtable field is is derived from `this` pointer, therefore it should be in
+  // default address space.
   VTableField = Builder.CreatePointerBitCastOrAddrSpaceCast(
-      VTableField, VTablePtrTy->getPointerTo(GlobalsAS));
+      VTableField, VTablePtrTy->getPointerTo());
   VTableAddressPoint = Builder.CreatePointerBitCastOrAddrSpaceCast(
       VTableAddressPoint, VTablePtrTy);
 

diff  --git a/clang/test/CodeGenCUDA/vtbl.cu b/clang/test/CodeGenCUDA/vtbl.cu
new file mode 100644
index 0000000000000..e09d904c2aade
--- /dev/null
+++ b/clang/test/CodeGenCUDA/vtbl.cu
@@ -0,0 +1,17 @@
+// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -target-cpu gfx906 \
+// RUN:   -emit-llvm -o - %s | FileCheck %s
+
+#include "Inputs/cuda.h"
+
+// CHECK-LABEL: define {{.*}}@_ZN1AC2Ev(%struct.A* nonnull align 8 dereferenceable(8) %this)
+// CHECK: store %struct.A* %this, %struct.A** %this.addr.ascast
+// CHECK: %this1 = load %struct.A*, %struct.A** %this.addr.ascast
+// CHECK: %[[VTFIELD:.*]] = bitcast %struct.A* %this1 to i32 (...)* addrspace(1)**
+// CHECK: store i32 (...)* addrspace(1)* bitcast{{.*}} @_ZTV1A{{.*}}, i32 (...)* addrspace(1)** %[[VTFIELD]]
+struct A {
+  __device__ virtual void vf() {}
+};
+
+__global__ void kern() {
+  A a;
+}