[clang] [llvm] AMDGPU: Set the addrspacecast nonnull flag instead of the intrinsic (PR #220927)
Matt Arsenault via cfe-commits
cfe-commits at lists.llvm.org
Tue Sep 8 04:22:44 PDT 2026
https://github.com/arsenm updated https://github.com/llvm/llvm-project/pull/220927
>From 0efd43b5056066d89586182d17eb78a0cadb2411 Mon Sep 17 00:00:00 2001
From: Matt Arsenault <Matthew.Arsenault at amd.com>
Date: Thu, 20 Aug 2026 22:33:36 +0200
Subject: [PATCH] AMDGPU: Set the addrspacecast nonnull flag instead of the
intrinsic
AMDGPUCodeGenPrepare proved the source of certain flat<->local/private
addrspacecasts non-null and rewrote them to
llvm.amdgcn.addrspacecast.nonnull. Now that the flag is honored in
codegen, set it in place on the existing instruction instead.
Co-authored-by: Claude (Claude-Opus-4.8) <noreply at anthropic.com>
---
clang/test/CodeGenCUDA/fp-contract.cu | 1 +
.../Target/AMDGPU/AMDGPUCodeGenPrepare.cpp | 25 +++++++++++--------
.../codegen-prepare-addrspacecast-non-null.ll | 24 +++++++++---------
3 files changed, 28 insertions(+), 22 deletions(-)
diff --git a/clang/test/CodeGenCUDA/fp-contract.cu b/clang/test/CodeGenCUDA/fp-contract.cu
index d98d5ca3f462b..112a443045d1c 100644
--- a/clang/test/CodeGenCUDA/fp-contract.cu
+++ b/clang/test/CodeGenCUDA/fp-contract.cu
@@ -181,6 +181,7 @@ __host__ __device__ float func(float a, float b, float c) { return a + b * c; }
// NV-ON: fma.rn.f32
// NV-ON-NEXT: st.param.b32
// AMD-ON: v_fmac_f32_e32
+// AMD-ON-NEXT: s_waitcnt vmcnt(0) lgkmcnt(0)
// AMD-ON-NEXT: s_setpc_b64
// NV-OFF: mul.rn.f32
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUCodeGenPrepare.cpp b/llvm/lib/Target/AMDGPU/AMDGPUCodeGenPrepare.cpp
index 8caf6e77309b9..86deb93d11ce8 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUCodeGenPrepare.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUCodeGenPrepare.cpp
@@ -1999,14 +1999,23 @@ static bool isPtrKnownNeverNull(const Value *V, const DataLayout &DL,
}
bool AMDGPUCodeGenPrepareImpl::visitAddrSpaceCastInst(AddrSpaceCastInst &I) {
- // Intrinsic doesn't support vectors, also it seems that it's often difficult
- // to prove that a vector cannot have any nulls in it so it's unclear if it's
- // worth supporting.
+ // TODO: This is target-independent reasoning about the source pointer being
+ // non-null, and would fit better in a generic pass such as
+ // AggressiveInstCombine. It lives here for now because proving the source is
+ // not the null value requires knowing the numeric null pointer value of the
+ // source address space, which is not yet a first-class IR concept.
+
+ // If the flag is already set there is nothing to do.
+ if (I.hasNonNull())
+ return false;
+
+ // It is often difficult to prove that a vector of pointers cannot have any
+ // nulls in it, so it's unclear if it's worth supporting.
if (I.getType()->isVectorTy())
return false;
- // Check if this can be lowered to a amdgcn.addrspacecast.nonnull.
- // This is only worthwhile for casts from/to priv/local to flat.
+ // The nonnull flag only affects the lowering of casts from/to priv/local to
+ // flat, so only bother proving non-null for those.
const unsigned SrcAS = I.getSrcAddressSpace();
const unsigned DstAS = I.getDestAddressSpace();
@@ -2027,11 +2036,7 @@ bool AMDGPUCodeGenPrepareImpl::visitAddrSpaceCastInst(AddrSpaceCastInst &I) {
}))
return false;
- IRBuilder<> B(&I);
- auto *Intrin = B.CreateIntrinsic(
- I.getType(), Intrinsic::amdgcn_addrspacecast_nonnull, {I.getOperand(0)});
- I.replaceAllUsesWith(Intrin);
- DeadVals.push_back(&I);
+ I.setNonNull();
return true;
}
diff --git a/llvm/test/CodeGen/AMDGPU/codegen-prepare-addrspacecast-non-null.ll b/llvm/test/CodeGen/AMDGPU/codegen-prepare-addrspacecast-non-null.ll
index c87a7fba94ab3..dbed069faef1a 100644
--- a/llvm/test/CodeGen/AMDGPU/codegen-prepare-addrspacecast-non-null.ll
+++ b/llvm/test/CodeGen/AMDGPU/codegen-prepare-addrspacecast-non-null.ll
@@ -11,7 +11,7 @@
define void @local_to_flat_nonnull_arg(ptr addrspace(3) nonnull %ptr) {
; OPT-LABEL: define void @local_to_flat_nonnull_arg(
; OPT-SAME: ptr addrspace(3) nonnull [[PTR:%.*]]) {
-; OPT-NEXT: [[TMP1:%.*]] = call ptr @llvm.amdgcn.addrspacecast.nonnull.p0.p3(ptr addrspace(3) [[PTR]])
+; OPT-NEXT: [[TMP1:%.*]] = addrspacecast nonnull ptr addrspace(3) [[PTR]] to ptr
; OPT-NEXT: store volatile i32 7, ptr [[TMP1]], align 4
; OPT-NEXT: ret void
;
@@ -32,7 +32,7 @@ define void @local_to_flat_nonnull_arg(ptr addrspace(3) nonnull %ptr) {
define void @private_to_flat_nonnull_arg(ptr addrspace(5) nonnull %ptr) {
; OPT-LABEL: define void @private_to_flat_nonnull_arg(
; OPT-SAME: ptr addrspace(5) nonnull [[PTR:%.*]]) {
-; OPT-NEXT: [[TMP1:%.*]] = call ptr @llvm.amdgcn.addrspacecast.nonnull.p0.p5(ptr addrspace(5) [[PTR]])
+; OPT-NEXT: [[TMP1:%.*]] = addrspacecast nonnull ptr addrspace(5) [[PTR]] to ptr
; OPT-NEXT: store volatile i32 7, ptr [[TMP1]], align 4
; OPT-NEXT: ret void
;
@@ -53,7 +53,7 @@ define void @private_to_flat_nonnull_arg(ptr addrspace(5) nonnull %ptr) {
define void @flat_to_local_nonnull_arg(ptr nonnull %ptr) {
; OPT-LABEL: define void @flat_to_local_nonnull_arg(
; OPT-SAME: ptr nonnull [[PTR:%.*]]) {
-; OPT-NEXT: [[TMP1:%.*]] = call ptr addrspace(3) @llvm.amdgcn.addrspacecast.nonnull.p3.p0(ptr [[PTR]])
+; OPT-NEXT: [[TMP1:%.*]] = addrspacecast nonnull ptr [[PTR]] to ptr addrspace(3)
; OPT-NEXT: store volatile i32 7, ptr addrspace(3) [[TMP1]], align 4
; OPT-NEXT: ret void
;
@@ -72,7 +72,7 @@ define void @flat_to_local_nonnull_arg(ptr nonnull %ptr) {
define void @flat_to_private_nonnull_arg(ptr nonnull %ptr) {
; OPT-LABEL: define void @flat_to_private_nonnull_arg(
; OPT-SAME: ptr nonnull [[PTR:%.*]]) {
-; OPT-NEXT: [[TMP1:%.*]] = call ptr addrspace(5) @llvm.amdgcn.addrspacecast.nonnull.p5.p0(ptr [[PTR]])
+; OPT-NEXT: [[TMP1:%.*]] = addrspacecast nonnull ptr [[PTR]] to ptr addrspace(5)
; OPT-NEXT: store volatile i32 7, ptr addrspace(5) [[TMP1]], align 4
; OPT-NEXT: ret void
;
@@ -92,7 +92,7 @@ define void @local_to_flat_nonnull_load(ptr %p) {
; OPT-LABEL: define void @local_to_flat_nonnull_load(
; OPT-SAME: ptr [[P:%.*]]) {
; OPT-NEXT: [[PTR:%.*]] = load ptr addrspace(3), ptr [[P]], align 4, !nonnull [[META0:![0-9]+]]
-; OPT-NEXT: [[X:%.*]] = call ptr @llvm.amdgcn.addrspacecast.nonnull.p0.p3(ptr addrspace(3) [[PTR]])
+; OPT-NEXT: [[X:%.*]] = addrspacecast nonnull ptr addrspace(3) [[PTR]] to ptr
; OPT-NEXT: store volatile i32 7, ptr [[X]], align 4
; OPT-NEXT: ret void
;
@@ -117,7 +117,7 @@ define void @private_to_flat_nonnull_load(ptr %p) {
; OPT-LABEL: define void @private_to_flat_nonnull_load(
; OPT-SAME: ptr [[P:%.*]]) {
; OPT-NEXT: [[PTR:%.*]] = load ptr addrspace(5), ptr [[P]], align 4, !nonnull [[META0]]
-; OPT-NEXT: [[X:%.*]] = call ptr @llvm.amdgcn.addrspacecast.nonnull.p0.p5(ptr addrspace(5) [[PTR]])
+; OPT-NEXT: [[X:%.*]] = addrspacecast nonnull ptr addrspace(5) [[PTR]] to ptr
; OPT-NEXT: store volatile i32 7, ptr [[X]], align 4
; OPT-NEXT: ret void
;
@@ -142,7 +142,7 @@ define void @flat_to_local_nonnull_load(ptr %p) {
; OPT-LABEL: define void @flat_to_local_nonnull_load(
; OPT-SAME: ptr [[P:%.*]]) {
; OPT-NEXT: [[PTR:%.*]] = load ptr, ptr [[P]], align 8, !nonnull [[META0]]
-; OPT-NEXT: [[X:%.*]] = call ptr addrspace(3) @llvm.amdgcn.addrspacecast.nonnull.p3.p0(ptr [[PTR]])
+; OPT-NEXT: [[X:%.*]] = addrspacecast nonnull ptr [[PTR]] to ptr addrspace(3)
; OPT-NEXT: store volatile i32 7, ptr addrspace(3) [[X]], align 4
; OPT-NEXT: ret void
;
@@ -175,7 +175,7 @@ define void @flat_to_private_nonnull_load(ptr %p) {
; OPT-LABEL: define void @flat_to_private_nonnull_load(
; OPT-SAME: ptr [[P:%.*]]) {
; OPT-NEXT: [[PTR:%.*]] = load ptr, ptr [[P]], align 8, !nonnull [[META0]]
-; OPT-NEXT: [[X:%.*]] = call ptr addrspace(5) @llvm.amdgcn.addrspacecast.nonnull.p5.p0(ptr [[PTR]])
+; OPT-NEXT: [[X:%.*]] = addrspacecast nonnull ptr [[PTR]] to ptr addrspace(5)
; OPT-NEXT: store volatile i32 7, ptr addrspace(5) [[X]], align 4
; OPT-NEXT: ret void
;
@@ -208,7 +208,7 @@ define void @private_alloca_to_flat(ptr %ptr) {
; OPT-LABEL: define void @private_alloca_to_flat(
; OPT-SAME: ptr [[PTR:%.*]]) {
; OPT-NEXT: [[ALLOCA:%.*]] = alloca i8, align 1, addrspace(5)
-; OPT-NEXT: [[TMP1:%.*]] = call ptr @llvm.amdgcn.addrspacecast.nonnull.p0.p5(ptr addrspace(5) [[ALLOCA]])
+; OPT-NEXT: [[TMP1:%.*]] = addrspacecast nonnull ptr addrspace(5) [[ALLOCA]] to ptr
; OPT-NEXT: store volatile i32 7, ptr [[TMP1]], align 4
; OPT-NEXT: ret void
;
@@ -248,7 +248,7 @@ define void @knownbits_on_flat_to_priv(ptr %ptr) {
; OPT-NEXT: [[PTR_INT:%.*]] = ptrtoint ptr [[PTR]] to i64
; OPT-NEXT: [[PTR_OR:%.*]] = or i64 [[PTR_INT]], 15
; OPT-NEXT: [[KB_PTR:%.*]] = inttoptr i64 [[PTR_OR]] to ptr
-; OPT-NEXT: [[TMP1:%.*]] = call ptr addrspace(5) @llvm.amdgcn.addrspacecast.nonnull.p5.p0(ptr [[KB_PTR]])
+; OPT-NEXT: [[TMP1:%.*]] = addrspacecast nonnull ptr [[KB_PTR]] to ptr addrspace(5)
; OPT-NEXT: store volatile i32 7, ptr addrspace(5) [[TMP1]], align 4
; OPT-NEXT: ret void
;
@@ -274,7 +274,7 @@ define void @knownbits_on_priv_to_flat(ptr addrspace(5) %ptr) {
; OPT-NEXT: [[PTR_INT:%.*]] = ptrtoint ptr addrspace(5) [[PTR]] to i32
; OPT-NEXT: [[PTR_OR:%.*]] = and i32 [[PTR_INT]], 65535
; OPT-NEXT: [[KB_PTR:%.*]] = inttoptr i32 [[PTR_OR]] to ptr addrspace(5)
-; OPT-NEXT: [[TMP1:%.*]] = call ptr @llvm.amdgcn.addrspacecast.nonnull.p0.p5(ptr addrspace(5) [[KB_PTR]])
+; OPT-NEXT: [[TMP1:%.*]] = addrspacecast nonnull ptr addrspace(5) [[KB_PTR]] to ptr
; OPT-NEXT: store volatile i32 7, ptr [[TMP1]], align 4
; OPT-NEXT: ret void
;
@@ -312,7 +312,7 @@ define void @recursive_phis(i1 %cond, ptr addrspace(5) %ptr) {
; OPT-NEXT: br label %[[FINALLY]]
; OPT: [[FINALLY]]:
; OPT-NEXT: [[PHI_PTR]] = phi ptr addrspace(5) [ [[KB_PTR]], %[[THEN]] ], [ [[OTHER_PHI]], %[[ELSE]] ]
-; OPT-NEXT: [[TMP0:%.*]] = call ptr @llvm.amdgcn.addrspacecast.nonnull.p0.p5(ptr addrspace(5) [[PHI_PTR]])
+; OPT-NEXT: [[TMP0:%.*]] = addrspacecast nonnull ptr addrspace(5) [[PHI_PTR]] to ptr
; OPT-NEXT: store volatile i32 7, ptr [[TMP0]], align 4
; OPT-NEXT: br i1 [[COND]], label %[[ELSE]], label %[[END:.*]]
; OPT: [[END]]:
More information about the cfe-commits
mailing list