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; +} _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits