https://github.com/koparasy created https://github.com/llvm/llvm-project/pull/228652
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. >From ae0573f2f79079b30a408c60825347d2080aab46 Mon Sep 17 00:00:00 2001 From: Konstantinos Parasyris <[email protected]> Date: Fri, 2 Oct 2026 20:54:12 -0700 Subject: [PATCH] [CIR] Preserve address spaces in emitPointerWithAlignment casts 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. Co-Authored-By: Claude Opus 5.5 (1M context) <[email protected]> --- clang/lib/CIR/CodeGen/CIRGenBuilder.h | 3 +- clang/lib/CIR/CodeGen/CIRGenExpr.cpp | 4 +- .../CIR/CodeGen/subscript-addrspace-cast.cpp | 69 +++++++++++++++++++ .../CodeGenBuiltins/builtin-stdc-bit-c2y.c | 18 +++++ 4 files changed, 92 insertions(+), 2 deletions(-) create mode 100644 clang/test/CIR/CodeGen/subscript-addrspace-cast.cpp diff --git a/clang/lib/CIR/CodeGen/CIRGenBuilder.h b/clang/lib/CIR/CodeGen/CIRGenBuilder.h index d224feb83b03df..2e9bae5e7d2c80 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 62740dc246827c..f5eafaf393c029 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 00000000000000..c75d615138961f --- /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 3921ab6926d680..f698cf039142e6 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); } _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
