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

[yaxunl]

Summary

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.

[arichardson]

Merging this change broke our out-of-tree CHERI targets (and Arm Morello). We use addrspace(200) for *all* globals (including vtables). It would have been nice if I had been added to this review considering that I added line you are changing here.

If vtables are not in the default globals address space, I think we need another way of expressing this. I think CGM.getContext().getTargetAddressSpace(LangAS::Default)) should also be correct for your use-case?

[yaxunl]

In D103835#2999545, @arichardson wrote:

Merging this change broke our out-of-tree CHERI targets (and Arm Morello). We use addrspace(200) for *all* globals (including vtables). It would have been nice if I had been added to this review considering that I added line you are changing here.

If vtables are not in the default globals address space, I think we need another way of expressing this. I think CGM.getContext().getTargetAddressSpace(LangAS::Default)) should also be correct for your use-case?

vtbl addr space should be the same as this pointer. If I use addr space of this pointer and not assuming it is default addr space, will it work for you? Thanks.

[arichardson]

In D103835#2999731, @yaxunl wrote:

In D103835#2999545, @arichardson wrote:

Merging this change broke our out-of-tree CHERI targets (and Arm Morello). We use addrspace(200) for *all* globals (including vtables). It would have been nice if I had been added to this review considering that I added line you are changing here.

If vtables are not in the default globals address space, I think we need another way of expressing this. I think CGM.getContext().getTargetAddressSpace(LangAS::Default)) should also be correct for your use-case?

vtbl addr space should be the same as this pointer. If I use addr space of this pointer and not assuming it is default addr space, will it work for you? Thanks.

Yes that would also work. Thanks!

[rjmccall]

Hmm. I think "v-tables are in the address space of the object pointer" is not a good assumption. Probably this ought to be determined by the C++ ABI for the target. In principle it could even be class-specific, but I think we can start by assuming it's universal.

It should be decided by the AST-level ABI abstraction so that it properly affects record layout, since different address spaces can have different pointer sizes.

[yaxunl]

In D103835#3001011, @rjmccall wrote:

Hmm. I think "v-tables are in the address space of the object pointer" is not a good assumption. Probably this ought to be determined by the C++ ABI for the target. In principle it could even be class-specific, but I think we can start by assuming it's universal.

It should be decided by the AST-level ABI abstraction so that it properly affects record layout, since different address spaces can have different pointer sizes.

Sorry my previous description was not accurate.

Currently vtbl addr space is assumed to be the default global addr space of LLVM IR, which is determined by the data layout of the LLVM IR. This patch did not change that.

The vtbl field of a class is a pointer to default addr space. When the vtbl field gets initialized, it is casted to a pointer to default global addr space so that the vtbl can be stored to it. The addr space of the vtbl field itself should be the same as this pointer. So we were talking about a pointer to the vtbl field of an object, not the addr space of vtbl.

[rjmccall]

I see. Yes, in that case, that sounds right.

[yaxunl]

https://reviews.llvm.org/D109841

clang/lib/CodeGen/CGClass.cpp

void CodeGenFunction::InitializeVTablePointer(const VPtr &Vptr) {
   // Finally, store the address point. Use the same LLVM types as the field to
   // support optimization.
   unsigned GlobalsAS = CGM.getDataLayout().getDefaultGlobalsAddressSpace();
   unsigned ProgAS = CGM.getDataLayout().getProgramAddressSpace();
   llvm::Type *VTablePtrTy =
       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);
 
   llvm::StoreInst *Store = Builder.CreateStore(VTableAddressPoint, VTableField);
   TBAAAccessInfo TBAAInfo = CGM.getTBAAVTablePtrAccessInfo(VTablePtrTy);
   CGM.DecorateInstructionWithTBAA(Store, TBAAInfo);
   if (CGM.getCodeGenOpts().OptimizationLevel > 0 &&
       CGM.getCodeGenOpts().StrictVTablePointers)

clang/test/CodeGenCUDA/vtbl.cu

+// 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;
+}