[clang] [CIR] Keep the address space in thunk this and return adjustments (PR #234183)

[Arseniy Obolenskiy via cfe-commits]

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 <arseniy.obolenskiy at amd.com>
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 <arseniy.obolenskiy at amd.com>
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());