https://github.com/aobolensk updated https://github.com/llvm/llvm-project/pull/234183
>From 282f5e49d9400be4ccf8b3345927ed9fa7bb1aa5 Mon Sep 17 00:00:00 2001 From: Arseniy Obolenskiy <[email protected]> Date: Wed, 21 Oct 2026 14:21:00 +0200 Subject: [PATCH 1/2] [CIR] Keep the address space in thunk this and return adjustments Thunks cast this to a plain i8 pointer, which broke on spirv64-amd where this lives in address space 4 --- clang/lib/CIR/CodeGen/CIRGenItaniumCXXABI.cpp | 10 ++++---- clang/test/CIR/CodeGenHIP/thunk-addrspace.hip | 23 +++++++++++++++++++ 2 files changed, 29 insertions(+), 4 deletions(-) create mode 100644 clang/test/CIR/CodeGenHIP/thunk-addrspace.hip diff --git a/clang/lib/CIR/CodeGen/CIRGenItaniumCXXABI.cpp b/clang/lib/CIR/CodeGen/CIRGenItaniumCXXABI.cpp index 9c952dd6054686..1f27561f4208f3 100644 --- a/clang/lib/CIR/CodeGen/CIRGenItaniumCXXABI.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenItaniumCXXABI.cpp @@ -2257,8 +2257,9 @@ static mlir::Value performTypeAdjustment(CIRGenFunction &cgf, CIRGenBuilderTy &builder = cgf.getBuilder(); mlir::Location loc = builder.getUnknownLoc(); - cir::PointerType i8PtrTy = builder.getUInt8PtrTy(); - mlir::Value v = builder.createBitcast(initialPtr.getPointer(), i8PtrTy); + mlir::Value v = + builder.createPtrBitcast(initialPtr.getPointer(), builder.getUInt8Ty()); + auto i8PtrTy = mlir::cast<cir::PointerType>(v.getType()); // In a base-to-derived cast, the non-virtual adjustment is applied first. if (nonVirtualAdjustment && !isReturnAdjustment) { @@ -2271,11 +2272,12 @@ static mlir::Value performTypeAdjustment(CIRGenFunction &cgf, if (virtualAdjustment) { mlir::Value vtablePtr = cgf.getVTablePtr( loc, Address(v, clang::CharUnits::One()), unadjustedClass); - vtablePtr = builder.createBitcast(vtablePtr, i8PtrTy); + cir::PointerType vtableI8PtrTy = builder.getUInt8PtrTy(); + vtablePtr = builder.createBitcast(vtablePtr, vtableI8PtrTy); mlir::Value offset; mlir::Value offsetPtr = - cir::PtrStrideOp::create(builder, loc, i8PtrTy, vtablePtr, + cir::PtrStrideOp::create(builder, loc, vtableI8PtrTy, vtablePtr, builder.getSInt64(virtualAdjustment, loc)); if (cgf.cgm.getLangOpts().RelativeCXXABIVTables) { assert(!cir::MissingFeatures::vtableRelativeLayout()); diff --git a/clang/test/CIR/CodeGenHIP/thunk-addrspace.hip b/clang/test/CIR/CodeGenHIP/thunk-addrspace.hip new file mode 100644 index 00000000000000..6326ce485aeea0 --- /dev/null +++ b/clang/test/CIR/CodeGenHIP/thunk-addrspace.hip @@ -0,0 +1,23 @@ +// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -x hip -fcuda-is-device -fclangir -emit-cir %s -o %t.cir +// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s +// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -x hip -fcuda-is-device -fclangir -emit-llvm %s -o - \ +// RUN: | FileCheck --check-prefix=LLVM %s +// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -x hip -fcuda-is-device -emit-llvm %s -o - \ +// RUN: | FileCheck --check-prefix=LLVM %s + +#define __device__ __attribute__((device)) + +struct A { __device__ virtual void f(); }; +struct B { __device__ virtual void g(); int b; }; +struct C : A, B { __device__ void g() override; }; +__device__ void C::g() {} + +// CIR-LABEL: cir.func {{.*}} @_ZThn8_N1C1gEv +// CIR: %[[THIS:.*]] = cir.load {{.*}}, !cir.ptr<!rec_C, target_address_space(4)> +// CIR: %[[RAW:.*]] = cir.cast bitcast %[[THIS]] : !cir.ptr<!rec_C, target_address_space(4)> -> !cir.ptr<!u8i, target_address_space(4)> +// CIR: %[[ADJ:.*]] = cir.ptr_stride %[[RAW]], %{{.*}} : (!cir.ptr<!u8i, target_address_space(4)>, !s64i) -> !cir.ptr<!u8i, target_address_space(4)> +// CIR: cir.cast bitcast %[[ADJ]] : !cir.ptr<!u8i, target_address_space(4)> -> !cir.ptr<!rec_C, target_address_space(4)> + +// LLVM-LABEL: define spir_func void @_ZThn8_N1C1gEv( +// LLVM: %[[THIS:.*]] = load ptr addrspace(4), ptr addrspace(4) +// LLVM: getelementptr {{.*}}i8, ptr addrspace(4) %[[THIS]], i64 -8 >From 1834c9131f757b5be44dfb3c7ea44234bbaf5eae Mon Sep 17 00:00:00 2001 From: Arseniy Obolenskiy <[email protected]> Date: Fri, 23 Oct 2026 13:44:49 +0200 Subject: [PATCH 2/2] suggestion --- clang/lib/CIR/CodeGen/CIRGenItaniumCXXABI.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenItaniumCXXABI.cpp b/clang/lib/CIR/CodeGen/CIRGenItaniumCXXABI.cpp index 1f27561f4208f3..55f4dd547a971e 100644 --- a/clang/lib/CIR/CodeGen/CIRGenItaniumCXXABI.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenItaniumCXXABI.cpp @@ -2257,6 +2257,7 @@ static mlir::Value performTypeAdjustment(CIRGenFunction &cgf, CIRGenBuilderTy &builder = cgf.getBuilder(); mlir::Location loc = builder.getUnknownLoc(); + cir::PointerType ui8DefaultASPtrTy = builder.getUInt8PtrTy(); mlir::Value v = builder.createPtrBitcast(initialPtr.getPointer(), builder.getUInt8Ty()); auto i8PtrTy = mlir::cast<cir::PointerType>(v.getType()); @@ -2272,12 +2273,11 @@ static mlir::Value performTypeAdjustment(CIRGenFunction &cgf, if (virtualAdjustment) { mlir::Value vtablePtr = cgf.getVTablePtr( loc, Address(v, clang::CharUnits::One()), unadjustedClass); - cir::PointerType vtableI8PtrTy = builder.getUInt8PtrTy(); - vtablePtr = builder.createBitcast(vtablePtr, vtableI8PtrTy); + vtablePtr = builder.createBitcast(vtablePtr, ui8DefaultASPtrTy); mlir::Value offset; mlir::Value offsetPtr = - cir::PtrStrideOp::create(builder, loc, vtableI8PtrTy, vtablePtr, + cir::PtrStrideOp::create(builder, loc, ui8DefaultASPtrTy, vtablePtr, builder.getSInt64(virtualAdjustment, loc)); if (cgf.cgm.getLangOpts().RelativeCXXABIVTables) { assert(!cir::MissingFeatures::vtableRelativeLayout()); _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
