This revision was automatically updated to reflect the committed changes. Closed by commit rG054cc3b1b469: [CUDA][HIP] Fix store of vtbl in ctor (authored by yaxunl). Herald added a project: clang.
Repository: rG LLVM Github Monorepo CHANGES SINCE LAST ACTION https://reviews.llvm.org/D103835/new/ https://reviews.llvm.org/D103835 Files: clang/lib/CodeGen/CGClass.cpp clang/test/CodeGenCUDA/vtbl.cu Index: clang/test/CodeGenCUDA/vtbl.cu =================================================================== --- /dev/null +++ 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; +} Index: clang/lib/CodeGen/CGClass.cpp =================================================================== --- clang/lib/CodeGen/CGClass.cpp +++ clang/lib/CodeGen/CGClass.cpp @@ -2518,8 +2518,10 @@ 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);
Index: clang/test/CodeGenCUDA/vtbl.cu =================================================================== --- /dev/null +++ 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; +} Index: clang/lib/CodeGen/CGClass.cpp =================================================================== --- clang/lib/CodeGen/CGClass.cpp +++ clang/lib/CodeGen/CGClass.cpp @@ -2518,8 +2518,10 @@ 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);
_______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits