[clang] [CIR][AMDGPU] Fix address space of __builtin_frame_address result (PR #233724)
Arseniy Obolenskiy via cfe-commits
cfe-commits at lists.llvm.org
Tue Oct 20 02:08:44 PDT 2026
https://github.com/aobolensk updated https://github.com/llvm/llvm-project/pull/233724
>From e1c5471003c7b1b5fde6b1468a75359ddb6692c9 Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy at amd.com>
Date: Tue, 20 Oct 2026 09:19:27 +0200
Subject: [PATCH 1/2] [CIR][AMDGPU] Fix address space of
__builtin_frame_address result
Use an address space cast instead of a bitcast, which failed CIR verification
---
clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp | 5 +++--
clang/test/CIR/CodeGenBuiltins/builtins.cpp | 14 ++++++++++++++
2 files changed, 17 insertions(+), 2 deletions(-)
diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp
index a94139cff20ff..160578037066c 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp
@@ -2538,8 +2538,9 @@ RValue CIRGenFunction::emitBuiltinExpr(const GlobalDecl &gd, unsigned builtinID,
mlir::Value addr = cir::FrameAddrOp::create(
builder, loc, allocaInt8PtrTy,
builder.getConstAPInt(loc, builder.getUInt32Ty(), level));
- return RValue::get(
- builder.createCast(loc, cir::CastKind::bitcast, addr, voidPtrTy));
+ addr = builder.createBitcast(
+ addr, builder.getVoidPtrTy(getCIRAllocaAddressSpace()));
+ return RValue::get(performAddrSpaceCast(addr, convertType(e->getType())));
}
case Builtin::BI__builtin_extract_return_addr:
case Builtin::BI__builtin_frob_return_addr:
diff --git a/clang/test/CIR/CodeGenBuiltins/builtins.cpp b/clang/test/CIR/CodeGenBuiltins/builtins.cpp
index 920f0696dba2d..8e4be00faaba0 100644
--- a/clang/test/CIR/CodeGenBuiltins/builtins.cpp
+++ b/clang/test/CIR/CodeGenBuiltins/builtins.cpp
@@ -4,6 +4,10 @@
// RUN: FileCheck --input-file=%t-cir.ll %s -check-prefix=LLVM
// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -Wno-unused-value -emit-llvm %s -o %t.ll
// RUN: FileCheck --input-file=%t.ll %s -check-prefix=LLVM
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -Wno-unused-value -fclangir -emit-cir %s -o %t-amdgpu.cir
+// RUN: FileCheck --input-file=%t-amdgpu.cir %s -check-prefix=CIR-AMDGPU
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -Wno-unused-value -fclangir -emit-llvm %s -o %t-amdgpu.ll
+// RUN: FileCheck --input-file=%t-amdgpu.ll %s -check-prefix=LLVM-AMDGPU
double fabs(double x) {
return __builtin_fabs(x);
@@ -35,4 +39,14 @@ extern "C" void *test_frame_address(void) {
// LLVM-LABEL: @test_frame_address
// LLVM: {{%.*}} = call ptr @llvm.frameaddress.p0(i32 1)
+ // CIR-AMDGPU-LABEL: test_frame_address
+ // CIR-AMDGPU: [[ARG:%.*]] = cir.const #cir.int<1> : !u32i
+ // CIR-AMDGPU: %[[ADDR:.*]] = cir.frame_address([[ARG]]) : !cir.ptr<!u8i, target_address_space(5)>
+ // CIR-AMDGPU: %[[CAST:.*]] = cir.cast bitcast %[[ADDR]] : !cir.ptr<!u8i, target_address_space(5)> -> !cir.ptr<!void, target_address_space(5)>
+ // CIR-AMDGPU: {{%.*}} = cir.cast address_space %[[CAST]] : !cir.ptr<!void, target_address_space(5)> -> !cir.ptr<!void>
+
+ // LLVM-AMDGPU-LABEL: @test_frame_address
+ // LLVM-AMDGPU: %[[ADDR:.*]] = call ptr addrspace(5) @llvm.frameaddress.p5(i32 1)
+ // LLVM-AMDGPU: {{%.*}} = addrspacecast ptr addrspace(5) %[[ADDR]] to ptr
+
}
>From 2dbdbcd9553b6a1a6ec2448ce17783bad8d736de Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy at amd.com>
Date: Tue, 20 Oct 2026 11:08:03 +0200
Subject: [PATCH 2/2] [CIR] Take function addresses in the program address
space
CIR emitted function addresses in address space 0, but some targets put functions in other, like address space 4 on spirv64-amd-amdhsa
---
clang/lib/CIR/CodeGen/CIRGenExpr.cpp | 6 ++++--
clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp | 10 +++++++++-
clang/lib/CIR/CodeGen/CIRGenModule.cpp | 4 +++-
clang/lib/CIR/Dialect/IR/CIRDialect.cpp | 11 ++++++++++-
clang/test/CIR/CodeGen/address-space.c | 7 +++++++
clang/test/CIR/CodeGen/no-proto-then-def.c | 7 +++++++
.../test/CIR/CodeGenHIP/spirv-constant-addrspace.hip | 9 +++++++++
7 files changed, 49 insertions(+), 5 deletions(-)
diff --git a/clang/lib/CIR/CodeGen/CIRGenExpr.cpp b/clang/lib/CIR/CodeGen/CIRGenExpr.cpp
index dfd6fa725e77f..84d30de7803c9 100644
--- a/clang/lib/CIR/CodeGen/CIRGenExpr.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenExpr.cpp
@@ -1009,7 +1009,9 @@ static LValue emitFunctionDeclLValue(CIRGenFunction &cgf, const Expr *e,
assert(!cir::MissingFeatures::sanitizers());
mlir::Type fnTy = funcOp.getFunctionType();
- mlir::Type ptrTy = cir::PointerType::get(fnTy);
+ mlir::ptr::MemorySpaceAttrInterface fnAddrSpace =
+ cgf.cgm.getTypes().getPointerAddressSpace(fd->getType());
+ mlir::Type ptrTy = cir::PointerType::get(fnTy, fnAddrSpace);
mlir::Value addr;
// On the HIP host, a reference to a __global__ kernel must resolve to the
@@ -1033,7 +1035,7 @@ static LValue emitFunctionDeclLValue(CIRGenFunction &cgf, const Expr *e,
if (funcOp.getFunctionType() != cgf.convertType(fd->getType())) {
fnTy = cgf.convertType(fd->getType());
- ptrTy = cir::PointerType::get(fnTy);
+ ptrTy = cir::PointerType::get(fnTy, fnAddrSpace);
addr = cir::CastOp::create(cgf.getBuilder(), addr.getLoc(), ptrTy,
cir::CastKind::bitcast, addr);
diff --git a/clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp b/clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp
index 2082963025238..633fa8e193307 100644
--- a/clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenExprScalar.cpp
@@ -2653,7 +2653,15 @@ mlir::Value ScalarExprEmitter::VisitCastExpr(CastExpr *ce) {
mlir::Value src = Visit(const_cast<Expr *>(subExpr));
mlir::Type dstTy = cgf.convertType(destTy);
- assert(!cir::MissingFeatures::addressSpace());
+ // Function pointers use the program address space and object pointers the
+ // default one, so a bitcast between them may change the address space.
+ auto srcPtrTy = mlir::dyn_cast<cir::PointerType>(src.getType());
+ auto dstPtrTy = mlir::dyn_cast<cir::PointerType>(dstTy);
+ if (srcPtrTy && dstPtrTy &&
+ srcPtrTy.getAddrSpace() != dstPtrTy.getAddrSpace())
+ src = cgf.performAddrSpaceCast(
+ src, cir::PointerType::get(srcPtrTy.getPointee(),
+ dstPtrTy.getAddrSpace()));
if (cgf.sanOpts.has(SanitizerKind::CFIUnrelatedCast))
cgf.getCIRGenModule().errorNYI(subExpr->getSourceRange(),
diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
index 614a731158e45..0b964ed2b7516 100644
--- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp
@@ -2139,7 +2139,9 @@ void CIRGenModule::replaceUsesOfNonProtoTypeWithRealFunction(
// canonicalizer.
mlir::Value res = getGlobalOp.getAddr();
const mlir::Type oldResTy = res.getType();
- const auto newPtrTy = cir::PointerType::get(newFn.getFunctionType());
+ const auto newPtrTy = cir::PointerType::get(
+ newFn.getFunctionType(),
+ mlir::cast<cir::PointerType>(oldResTy).getAddrSpace());
if (oldResTy != newPtrTy) {
res.setType(newPtrTy);
builder.setInsertionPointAfter(getGlobalOp.getOperation());
diff --git a/clang/lib/CIR/Dialect/IR/CIRDialect.cpp b/clang/lib/CIR/Dialect/IR/CIRDialect.cpp
index 4dfbfbb195ca4..c09a60624d505 100644
--- a/clang/lib/CIR/Dialect/IR/CIRDialect.cpp
+++ b/clang/lib/CIR/Dialect/IR/CIRDialect.cpp
@@ -13,6 +13,7 @@
#include "clang/CIR/Dialect/IR/CIRDialect.h"
#include "clang/CIR/Dialect/IR/CIRAttrs.h"
+#include "clang/CIR/Dialect/IR/CIRDataLayout.h"
#include "clang/CIR/Dialect/IR/CIROpsEnums.h"
#include "clang/CIR/Dialect/IR/CIRTypes.h"
@@ -2759,7 +2760,15 @@ cir::GetGlobalOp::verifySymbolUses(SymbolTableCollection &symbolTable) {
<< resultType.getPointee() << "' does not match type " << symTy
<< " of the global @" << getName();
- if (symAddrSpaceAttr != resultType.getAddrSpace()) {
+ // A function address may also be taken in the program address space.
+ auto targetAS =
+ dyn_cast_if_present<TargetAddressSpaceAttr>(resultType.getAddrSpace());
+ bool isProgramAS =
+ isa<FuncOp>(op) && targetAS &&
+ targetAS.getValue() ==
+ CIRDataLayout(getOperation()->getParentOfType<ModuleOp>())
+ .getProgramAddressSpace();
+ if (symAddrSpaceAttr != resultType.getAddrSpace() && !isProgramAS) {
return emitOpError()
<< "result type address space does not match the address "
"space of the global @"
diff --git a/clang/test/CIR/CodeGen/address-space.c b/clang/test/CIR/CodeGen/address-space.c
index e806e08bca19f..13d52e5201da4 100644
--- a/clang/test/CIR/CodeGen/address-space.c
+++ b/clang/test/CIR/CodeGen/address-space.c
@@ -4,6 +4,8 @@
// RUN: FileCheck --input-file=%t.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
+// RUN: %clang_cc1 -triple avr -fclangir -emit-llvm %s -o %t-avr.ll
+// RUN: FileCheck --input-file=%t-avr.ll %s -check-prefix=AVR
// Test global variable with address space
// CIR: cir.global external @gvar = #cir.ptr<null> : !cir.ptr<!s32i, target_address_space(1)>
@@ -58,3 +60,8 @@ struct S { int a[4]; };
void copy_to_as(struct S __attribute__((address_space(1))) *p, struct S *q) {
*p = *q;
}
+
+// A function pointer in the program address space cast to an object pointer.
+// AVR-LABEL: define{{.*}} ptr @fn_to_void_ptr
+// AVR: store ptr addrspacecast (ptr addrspace(1) @copy_to_as to ptr), ptr
+void *fn_to_void_ptr(void) { return (void *)copy_to_as; }
diff --git a/clang/test/CIR/CodeGen/no-proto-then-def.c b/clang/test/CIR/CodeGen/no-proto-then-def.c
index c142c4f0a16f5..4eb3c8f753569 100644
--- a/clang/test/CIR/CodeGen/no-proto-then-def.c
+++ b/clang/test/CIR/CodeGen/no-proto-then-def.c
@@ -4,6 +4,8 @@
// RUN: FileCheck --input-file=%t-cir.ll %s -check-prefix=LLVM
// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm -Wno-deprecated-non-prototype %s -o %t.ll
// RUN: FileCheck --input-file=%t.ll %s -check-prefix=LLVM
+// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -fclangir -emit-cir -Wno-deprecated-non-prototype %s -o %t-spirv.cir
+// RUN: FileCheck --input-file=%t-spirv.cir %s -check-prefix=SPIRV
extern int default_proc();
@@ -22,6 +24,11 @@ int default_proc(int a) { return a; }
// CIR: cir.get_global @default_proc : !cir.ptr<!cir.func<(!s32i) -> !s32i>>
// CIR: cir.cast bitcast %{{.*}} : !cir.ptr<!cir.func<(!s32i) -> !s32i>> -> !cir.ptr<!cir.func<(...) -> !s32i>>
+// The retyped address keeps the program address space.
+// SPIRV: cir.func{{.*}}@test_proc_ptr
+// SPIRV: cir.get_global @default_proc : !cir.ptr<!cir.func<(!s32i) -> !s32i>, target_address_space(4)>
+// SPIRV: cir.cast bitcast %{{.*}} : !cir.ptr<!cir.func<(!s32i) -> !s32i>, target_address_space(4)> -> !cir.ptr<!cir.func<(...) -> !s32i>, target_address_space(4)>
+
// LLVM: define{{.*}} @test_proc_ptr
// LLVM: store ptr @default_proc, ptr
// LLVM: define{{.*}} @default_proc
diff --git a/clang/test/CIR/CodeGenHIP/spirv-constant-addrspace.hip b/clang/test/CIR/CodeGenHIP/spirv-constant-addrspace.hip
index 1ddc16332f65c..a93895fbf1a9e 100644
--- a/clang/test/CIR/CodeGenHIP/spirv-constant-addrspace.hip
+++ b/clang/test/CIR/CodeGenHIP/spirv-constant-addrspace.hip
@@ -60,3 +60,12 @@ __device__ int get_shared() {
// An explicit opencl_constant qualifier keeps UniformConstant.
// LLVM-LABEL: define{{.*}} i32 @_Z6get_ocPU3AS2i(ptr addrspace(2)
__device__ int get_oc(__attribute__((opencl_constant)) int *p) { return *p; }
+
+// Function addresses are in the program address space 4.
+// CIR-LABEL: cir.func{{.*}} @_Z7take_fnv
+// CIR: cir.get_global @_Z6fn_refv : !cir.ptr<!cir.func<()>, target_address_space(4)>
+// LLVM-LABEL: define{{.*}} void @_Z7take_fnv
+// LLVM: call{{.*}} @_Z6use_fnPFvvE(ptr addrspace(4) noundef @_Z6fn_refv)
+__device__ void fn_ref();
+__device__ void use_fn(void (*)());
+__device__ void take_fn() { use_fn(fn_ref); }
More information about the cfe-commits
mailing list