[clang] [CIR] Preserve address spaces in emitPointerWithAlignment casts (PR #228652)
via cfe-commits
cfe-commits at lists.llvm.org
Fri Oct 2 20:56:37 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-clang
Author: Konstantinos Parasyris (koparasy)
<details>
<summary>Changes</summary>
emitPointerWithAlignment handled CK_AddressSpaceConversion like CK_BitCast and never emitted the address-space cast, so the result kept the source address space. When the address escaped (returned reference, stored pointer), the store path bitcast the destination slot to the wrong address space. This miscompiled sycl::multi_ptr::operator[] on SPIR-V: a local-memory offset was used as a generic address. Emit the cast after the element bitcast, as classic CodeGen does.
createElementBitCast also built the new pointer type in the default address space, so changing the element type of a non-default-AS pointer produced an AS-changing bitcast that the verifier rejects (for example ((int *)p)[i] or __builtin_stdc_memreverse8 on an AS3 pointer). Keep the source address space, which matches classic's withElementType.
---
Full diff: https://github.com/llvm/llvm-project/pull/228652.diff
4 Files Affected:
- (modified) clang/lib/CIR/CodeGen/CIRGenBuilder.h (+2-1)
- (modified) clang/lib/CIR/CodeGen/CIRGenExpr.cpp (+3-1)
- (added) clang/test/CIR/CodeGen/subscript-addrspace-cast.cpp (+69)
- (modified) clang/test/CIR/CodeGenBuiltins/builtin-stdc-bit-c2y.c (+18)
``````````diff
diff --git a/clang/lib/CIR/CodeGen/CIRGenBuilder.h b/clang/lib/CIR/CodeGen/CIRGenBuilder.h
index d224feb83b03d..2e9bae5e7d2c8 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuilder.h
+++ b/clang/lib/CIR/CodeGen/CIRGenBuilder.h
@@ -539,7 +539,8 @@ class CIRGenBuilderTy : public cir::CIRBaseBuilderTy {
if (destType == addr.getElementType())
return addr;
- auto ptrTy = getPointerTo(destType);
+ auto srcPtrTy = mlir::cast<cir::PointerType>(addr.getPointer().getType());
+ auto ptrTy = getPointerTo(destType, srcPtrTy.getAddrSpace());
return Address(createBitcast(loc, addr.getPointer(), ptrTy), destType,
addr.getAlignment());
}
diff --git a/clang/lib/CIR/CodeGen/CIRGenExpr.cpp b/clang/lib/CIR/CodeGen/CIRGenExpr.cpp
index 62740dc246827..f5eafaf393c02 100644
--- a/clang/lib/CIR/CodeGen/CIRGenExpr.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenExpr.cpp
@@ -173,7 +173,9 @@ Address CIRGenFunction::emitPointerWithAlignment(const Expr *expr,
convertTypeForMem(expr->getType()->getPointeeType());
addr = getBuilder().createElementBitCast(getLoc(expr->getSourceRange()),
addr, eltTy);
- assert(!cir::MissingFeatures::addressSpace());
+ if (ce->getCastKind() == CK_AddressSpaceConversion)
+ addr = addr.withPointer(performAddrSpaceCast(
+ addr.getPointer(), convertType(expr->getType())));
return addr;
}
diff --git a/clang/test/CIR/CodeGen/subscript-addrspace-cast.cpp b/clang/test/CIR/CodeGen/subscript-addrspace-cast.cpp
new file mode 100644
index 0000000000000..c75d615138961
--- /dev/null
+++ b/clang/test/CIR/CodeGen/subscript-addrspace-cast.cpp
@@ -0,0 +1,69 @@
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fclangir -emit-cir %s -o %t.cir
+// RUN: FileCheck --input-file=%t.cir %s --check-prefix=CIR
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fclangir -emit-llvm %s -o %t-cir.ll
+// RUN: FileCheck --input-file=%t-cir.ll %s --check-prefix=LLVM
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm %s -o %t.ll
+// RUN: FileCheck --input-file=%t.ll %s --check-prefix=OGCG
+
+#define AS3 __attribute__((address_space(3)))
+
+float &at(AS3 float *p, long i) {
+ return ((float *)p)[i];
+}
+
+// CIR-LABEL: cir.func{{.*}} @_Z2atPU3AS3fl(
+// CIR: %[[RETVAL:.*]] = cir.alloca "__retval" {{.*}} : !cir.ptr<!cir.ptr<!cir.float>>
+// CIR: %[[P:.*]] = cir.load{{.*}} : !cir.ptr<!cir.ptr<!cir.float, target_address_space(3)>>, !cir.ptr<!cir.float, target_address_space(3)>
+// CIR: %[[G:.*]] = cir.cast address_space %[[P]] : !cir.ptr<!cir.float, target_address_space(3)> -> !cir.ptr<!cir.float>
+// CIR: %[[ELT:.*]] = cir.ptr_stride %[[G]], %{{.*}} : (!cir.ptr<!cir.float>, !s64i) -> !cir.ptr<!cir.float>
+// CIR-NOT: cir.cast bitcast
+// CIR: cir.store %[[ELT]], %[[RETVAL]] : !cir.ptr<!cir.float>, !cir.ptr<!cir.ptr<!cir.float>>
+
+// LLVM-LABEL: define {{.*}}@_Z2atPU3AS3fl(
+// LLVM: %[[G:.*]] = addrspacecast ptr addrspace(3) %{{.*}} to ptr
+// LLVM: getelementptr float, ptr %[[G]]
+
+// OGCG-LABEL: define {{.*}}@_Z2atPU3AS3fl(
+// OGCG: %[[G:.*]] = addrspacecast ptr addrspace(3) %{{.*}} to ptr
+// OGCG: getelementptr inbounds float, ptr %[[G]]
+
+int &at_int(AS3 float *p, long i) {
+ return ((int *)p)[i];
+}
+
+// CIR-LABEL: cir.func{{.*}} @_Z6at_intPU3AS3fl(
+// CIR: %[[P:.*]] = cir.load{{.*}} : !cir.ptr<!cir.ptr<!cir.float, target_address_space(3)>>, !cir.ptr<!cir.float, target_address_space(3)>
+// CIR: %[[BC:.*]] = cir.cast bitcast %[[P]] : !cir.ptr<!cir.float, target_address_space(3)> -> !cir.ptr<!s32i, target_address_space(3)>
+// CIR: %[[G:.*]] = cir.cast address_space %[[BC]] : !cir.ptr<!s32i, target_address_space(3)> -> !cir.ptr<!s32i>
+// CIR: cir.ptr_stride %[[G]], %{{.*}} : (!cir.ptr<!s32i>, !s64i) -> !cir.ptr<!s32i>
+
+// LLVM-LABEL: define {{.*}}@_Z6at_intPU3AS3fl(
+// LLVM: %[[G:.*]] = addrspacecast ptr addrspace(3) %{{.*}} to ptr
+// LLVM: getelementptr i32, ptr %[[G]]
+
+// OGCG-LABEL: define {{.*}}@_Z6at_intPU3AS3fl(
+// OGCG: %[[G:.*]] = addrspacecast ptr addrspace(3) %{{.*}} to ptr
+// OGCG: getelementptr inbounds i32, ptr %[[G]]
+
+float *g;
+void store_addr(AS3 float *p, long i) {
+ g = &((float *)p)[i];
+}
+
+// CIR-LABEL: cir.func{{.*}} @_Z10store_addrPU3AS3fl(
+// CIR: %[[P:.*]] = cir.load{{.*}} : !cir.ptr<!cir.ptr<!cir.float, target_address_space(3)>>, !cir.ptr<!cir.float, target_address_space(3)>
+// CIR: %[[G:.*]] = cir.cast address_space %[[P]] : !cir.ptr<!cir.float, target_address_space(3)> -> !cir.ptr<!cir.float>
+// CIR: %[[ELT:.*]] = cir.ptr_stride %[[G]], %{{.*}} : (!cir.ptr<!cir.float>, !s64i) -> !cir.ptr<!cir.float>
+// CIR: %[[GADDR:.*]] = cir.get_global @g : !cir.ptr<!cir.ptr<!cir.float>>
+// CIR-NOT: cir.cast bitcast
+// CIR: cir.store{{.*}} %[[ELT]], %[[GADDR]] : !cir.ptr<!cir.float>, !cir.ptr<!cir.ptr<!cir.float>>
+
+// LLVM-LABEL: define {{.*}}@_Z10store_addrPU3AS3fl(
+// LLVM: %[[G:.*]] = addrspacecast ptr addrspace(3) %{{.*}} to ptr
+// LLVM: %[[ELT:.*]] = getelementptr float, ptr %[[G]]
+// LLVM: store ptr %[[ELT]], ptr @g
+
+// OGCG-LABEL: define {{.*}}@_Z10store_addrPU3AS3fl(
+// OGCG: %[[G:.*]] = addrspacecast ptr addrspace(3) %{{.*}} to ptr
+// OGCG: %[[ELT:.*]] = getelementptr inbounds float, ptr %[[G]]
+// OGCG: store ptr %[[ELT]], ptr @g
diff --git a/clang/test/CIR/CodeGenBuiltins/builtin-stdc-bit-c2y.c b/clang/test/CIR/CodeGenBuiltins/builtin-stdc-bit-c2y.c
index 3921ab6926d68..f698cf039142e 100644
--- a/clang/test/CIR/CodeGenBuiltins/builtin-stdc-bit-c2y.c
+++ b/clang/test/CIR/CodeGenBuiltins/builtin-stdc-bit-c2y.c
@@ -95,6 +95,24 @@ void test_builtin_stdc_memreverse8_u64(unsigned char *p) {
// LLVM: call i64 @llvm.bswap.i64(
// LLVM: store i64
+void test_builtin_stdc_memreverse8_as3(
+ __attribute__((address_space(3))) unsigned char *p) {
+ __builtin_stdc_memreverse8(4, p);
+}
+
+// CIR-LABEL: test_builtin_stdc_memreverse8_as3
+// CIR: %[[P:.*]] = cir.load {{.*}} : !cir.ptr<!cir.ptr<!u8i, target_address_space(3)>>, !cir.ptr<!u8i, target_address_space(3)>
+// CIR: %[[CAST:.*]] = cir.cast bitcast %[[P]] : !cir.ptr<!u8i, target_address_space(3)> -> !cir.ptr<!u32i, target_address_space(3)>
+// CIR: %[[VAL:.*]] = cir.load {{.*}} %[[CAST]] : !cir.ptr<!u32i, target_address_space(3)>, !u32i
+// CIR: %[[SWAP:.*]] = cir.byte_swap %[[VAL]] : !u32i
+// CIR: cir.store {{.*}} %[[SWAP]], %[[CAST]] : !u32i, !cir.ptr<!u32i, target_address_space(3)>
+
+// LLVM-LABEL: test_builtin_stdc_memreverse8_as3
+// LLVM: %[[P:.*]] = load ptr addrspace(3), ptr
+// LLVM: %[[VAL:.*]] = load i32, ptr addrspace(3) %[[P]]
+// LLVM: %[[SWAP:.*]] = call i32 @llvm.bswap.i32(i32 %[[VAL]])
+// LLVM: store i32 %[[SWAP]], ptr addrspace(3) %[[P]]
+
void test_builtin_stdc_memreverse8_size3(unsigned char *p) {
__builtin_stdc_memreverse8(3, p);
}
``````````
</details>
https://github.com/llvm/llvm-project/pull/228652
More information about the cfe-commits
mailing list