[clang] [llvm] [mlir] [AMDGPU] Introduce builtin, intrinsic and ROCDL instruction for buffer_inv (PR #228497)
Dan Zimmerman via llvm-commits
llvm-commits at lists.llvm.org
Fri Oct 2 08:56:36 PDT 2026
https://github.com/danzimm created https://github.com/llvm/llvm-project/pull/228497
Depends on https://github.com/llvm/llvm-project/pull/226254: The first 5 commits are from that branch, this PR is the last 2 commits.
This PR was written with assistance by codex. I've audited and manually edited to make sure the PR is ready for review.
This PR is a second follow up to https://github.com/llvm/llvm-project/pull/221115, it exposes:
- `__builtin_amdgcn_buffer_inv(_Constant int)`; I added validation in SemaAMDGPU too
- `@llvm.amdgcn.buffer.inv(i32 immarg)`
- `rocdl.buffer.inv`
The lowering of `buffer_inv` was already implemented. Consequently defining a pattern on the pseudo was the only thing needed, outside of exposing the interface.
## Notes
- I implemented both the builtin and the intrinsic as there was ambiguity in https://github.com/llvm/llvm-project/pull/221115:
- https://github.com/llvm/llvm-project/pull/221115#issuecomment-5599993703 specifies `intrinsics like __builtin_amdgcn_buffer_wbinvl1` when suggesting a pattern to follow
- As I originally wanted to emit these instructions from triton, I also implemented the rocdl interface
>From a7f0eb06a85448adfc2292a3c67604279a359dce Mon Sep 17 00:00:00 2001
From: Dan Zimmerman <danzimm at meta.com>
Date: Wed, 23 Sep 2026 09:12:41 -0700
Subject: [PATCH 1/7] [AMDGPU] Add buffer_inv subtarget feature
---
llvm/lib/Target/AMDGPU/AMDGPU.td | 8 +++++++-
llvm/unittests/TargetParser/TargetParserTest.cpp | 12 ++++++++++++
2 files changed, 19 insertions(+), 1 deletion(-)
diff --git a/llvm/lib/Target/AMDGPU/AMDGPU.td b/llvm/lib/Target/AMDGPU/AMDGPU.td
index f96934e0b84349..e04bc53394577b 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPU.td
+++ b/llvm/lib/Target/AMDGPU/AMDGPU.td
@@ -510,6 +510,10 @@ defm GFX940Insts : AMDGPUSubtargetFeature<"gfx940-insts",
/*GenPredicate=*/0
>;
+defm BufferInvInst : AMDGPUSubtargetFeature<"buffer-inv-inst",
+ "Has buffer_inv instruction"
+>;
+
defm Permlane16Insts : AMDGPUSubtargetFeature<"permlane16-insts",
"Has v_permlane16_b32/v_permlanex16_b32 instructions"
>;
@@ -2079,6 +2083,7 @@ def FeatureISAVersion9_4_Common : FeatureSet<
FeatureAGPRAlloc,
FeatureTgSplitSupport,
FeatureGFX940Insts,
+ FeatureBufferInvInst,
FeatureRequiresAlignedVGPRs,
FeatureFmaMixInsts,
FeatureDLInsts,
@@ -3324,7 +3329,8 @@ def AMDGPUFrontendVisibleFeatures {
FeatureAtomicFaddRtnInsts, FeatureAtomicFlatPkAdd16Insts, FeatureAtomicGlobalPkAddBF16Inst,
FeatureBF16ConversionInsts,
FeatureBF16PackedInsts, FeatureBF16TransInsts, FeatureBF8ConversionScaleInsts,
- FeatureBVHRayTracingInsts, FeatureBitOp3Insts, FeatureCIInsts,
+ FeatureBVHRayTracingInsts, FeatureBitOp3Insts, FeatureBufferInvInst,
+ FeatureCIInsts,
FeatureClusters, FeatureCubeInsts, FeatureCvtPkNormVOP2Insts,
FeatureCvtSrPkBF16F32Inst, FeatureDLInsts, FeatureDPP,
FeatureDot10Insts, FeatureDot11Insts, FeatureDot12Insts,
diff --git a/llvm/unittests/TargetParser/TargetParserTest.cpp b/llvm/unittests/TargetParser/TargetParserTest.cpp
index 62d6f3a4c955b4..2847942227df33 100644
--- a/llvm/unittests/TargetParser/TargetParserTest.cpp
+++ b/llvm/unittests/TargetParser/TargetParserTest.cpp
@@ -2875,6 +2875,18 @@ TEST(TargetParserTest, testAMDGPUgetFeatureBitset) {
EXPECT_TRUE(Empty.empty());
}
+TEST(TargetParserTest, testAMDGPUBufferInvInstFeature) {
+ auto Has = [](AMDGPU::GPUKind AK) {
+ return AMDGPU::getFeatureBitset(AK).test(AMDGPU::FEAT_BUFFER_INV_INST);
+ };
+
+ EXPECT_FALSE(Has(AMDGPU::GK_GFX90A));
+ EXPECT_TRUE(Has(AMDGPU::GK_GFX942));
+ EXPECT_TRUE(Has(AMDGPU::GK_GFX950));
+ EXPECT_TRUE(Has(AMDGPU::GK_GFX9_4_GENERIC));
+ EXPECT_FALSE(Has(AMDGPU::GK_GFX1250));
+}
+
TEST(TargetParserTest, testAMDGPUHalfAddressableLDSFeature) {
auto Has = [](AMDGPU::GPUKind AK) {
return AMDGPU::getFeatureBitset(AK).test(
>From 5a4808272dbfb15e7b38b88cee9f85614f9f6e73 Mon Sep 17 00:00:00 2001
From: Dan Zimmerman <danzimm at meta.com>
Date: Thu, 24 Sep 2026 10:20:47 -0700
Subject: [PATCH 2/7] [AMDGPU] Use BufferInvInst for buffer_inv predicates
---
llvm/lib/Target/AMDGPU/BUFInstructions.td | 8 +++++---
llvm/test/MC/AMDGPU/buffer-inv-subtarget-feature.s | 10 ++++++++++
2 files changed, 15 insertions(+), 3 deletions(-)
create mode 100644 llvm/test/MC/AMDGPU/buffer-inv-subtarget-feature.s
diff --git a/llvm/lib/Target/AMDGPU/BUFInstructions.td b/llvm/lib/Target/AMDGPU/BUFInstructions.td
index 057c0f103ce800..2b5d20c6849f78 100644
--- a/llvm/lib/Target/AMDGPU/BUFInstructions.td
+++ b/llvm/lib/Target/AMDGPU/BUFInstructions.td
@@ -1429,7 +1429,7 @@ let SubtargetPredicate = HasAtomicFMinFMaxF64GlobalInsts in {
}
def BUFFER_INV : MUBUF_Invalidate<"buffer_inv"> {
- let SubtargetPredicate = isGFX940Plus;
+ let SubtargetPredicate = HasBufferInvInst;
let has_glc = 1;
let has_sccb = 1;
let InOperandList = (ins CPol_0:$cpol);
@@ -3660,10 +3660,12 @@ let AsmString = BUFFER_WBL2.Mnemonic, // drop flags
defm BUFFER_WBL2 : MUBUF_Real_gfx90a<0x28>;
defm BUFFER_INVL2 : MUBUF_Real_gfx90a<0x29>;
-let SubtargetPredicate = isGFX940Plus in {
+let SubtargetPredicate = isGFX940Plus in
def BUFFER_WBL2_gfx940 : MUBUF_Real_gfx940<0x28, BUFFER_WBL2>;
+
+let SubtargetPredicate = HasBufferInvInst,
+ AssemblerPredicate = HasBufferInvInst in
def BUFFER_INV_gfx940 : MUBUF_Real_gfx940<0x29, BUFFER_INV>;
-}
class MTBUF_Real_Base_vi <bits<4> op, MTBUF_Pseudo ps, int Enc> :
MTBUF_Real<ps>,
diff --git a/llvm/test/MC/AMDGPU/buffer-inv-subtarget-feature.s b/llvm/test/MC/AMDGPU/buffer-inv-subtarget-feature.s
new file mode 100644
index 00000000000000..52d7c4e5eed493
--- /dev/null
+++ b/llvm/test/MC/AMDGPU/buffer-inv-subtarget-feature.s
@@ -0,0 +1,10 @@
+// RUN: llvm-mc -triple=amdgpu9.42 -show-encoding %s | FileCheck --check-prefix=SUPPORTED %s
+// RUN: llvm-mc -triple=amdgpu9.50 -show-encoding %s | FileCheck --check-prefix=SUPPORTED %s
+// RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx9-4-generic -show-encoding %s | FileCheck --check-prefix=SUPPORTED %s
+// RUN: not llvm-mc -triple=amdgpu9.0a -filetype=null %s 2>&1 | FileCheck --check-prefix=GFX90A %s
+// RUN: not llvm-mc -triple=amdgpu9.42 -mattr=-buffer-inv-inst -filetype=null %s 2>&1 | FileCheck --check-prefix=DISABLED %s
+
+buffer_inv sc0 sc1
+// SUPPORTED: buffer_inv sc0 sc1 ; encoding: [0x00,0xc0,0xa4,0xe0,0x00,0x00,0x00,0x00]
+// GFX90A: error: instruction not supported on this GPU (gfx90a): buffer_inv
+// DISABLED: error: instruction not supported on this GPU (gfx942): buffer_inv
>From 2930ba5d493802bb79b073c83c201c77998af8f0 Mon Sep 17 00:00:00 2001
From: Dan Zimmerman <danzimm at meta.com>
Date: Thu, 24 Sep 2026 10:21:38 -0700
Subject: [PATCH 3/7] [AMDGPU] Use BufferInvInst in memory legalizer
---
llvm/lib/Target/AMDGPU/SIMemoryLegalizer.cpp | 6 +++---
1 file changed, 3 insertions(+), 3 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/SIMemoryLegalizer.cpp b/llvm/lib/Target/AMDGPU/SIMemoryLegalizer.cpp
index e818e06e3bb048..10804d65b82b0f 100644
--- a/llvm/lib/Target/AMDGPU/SIMemoryLegalizer.cpp
+++ b/llvm/lib/Target/AMDGPU/SIMemoryLegalizer.cpp
@@ -1396,7 +1396,7 @@ bool SIGfx6CacheControl::insertAcquire(MachineBasicBlock::iterator &MI,
if (canAffectGlobalAddrSpace(AddrSpace)) {
switch (Scope) {
case SIAtomicScope::SYSTEM:
- if (ST.hasGFX940Insts()) {
+ if (ST.hasBufferInvInst()) {
// Ensures that following loads will not see stale remote VMEM data or
// stale local VMEM data with MTYPE NC. Local VMEM data with MTYPE RW
// and CC will never be stale due to the local memory probes.
@@ -1428,7 +1428,7 @@ bool SIGfx6CacheControl::insertAcquire(MachineBasicBlock::iterator &MI,
}
[[fallthrough]];
case SIAtomicScope::AGENT:
- if (ST.hasGFX940Insts()) {
+ if (ST.hasBufferInvInst()) {
// Ensures that following loads will not see stale remote date or local
// MTYPE NC global data. Local MTYPE RW and CC memory will never be
// stale due to the memory probes.
@@ -1445,7 +1445,7 @@ bool SIGfx6CacheControl::insertAcquire(MachineBasicBlock::iterator &MI,
break;
case SIAtomicScope::WORKGROUP:
if (TgSplitEnabled) {
- if (ST.hasGFX940Insts()) {
+ if (ST.hasBufferInvInst()) {
// In threadgroup split mode the waves of a work-group can be
// executing on different CUs. Therefore need to invalidate the L1
// which is per CU. Otherwise in non-threadgroup split mode all waves
>From ad9c59817dce63bd49ab8e1bc197919125cb6b62 Mon Sep 17 00:00:00 2001
From: Dan Zimmerman <danzimm at meta.com>
Date: Thu, 24 Sep 2026 14:44:56 -0700
Subject: [PATCH 4/7] [AMDGPU] Remove redundant buffer_inv feature tests
---
llvm/test/MC/AMDGPU/buffer-inv-subtarget-feature.s | 10 ----------
llvm/unittests/TargetParser/TargetParserTest.cpp | 12 ------------
2 files changed, 22 deletions(-)
delete mode 100644 llvm/test/MC/AMDGPU/buffer-inv-subtarget-feature.s
diff --git a/llvm/test/MC/AMDGPU/buffer-inv-subtarget-feature.s b/llvm/test/MC/AMDGPU/buffer-inv-subtarget-feature.s
deleted file mode 100644
index 52d7c4e5eed493..00000000000000
--- a/llvm/test/MC/AMDGPU/buffer-inv-subtarget-feature.s
+++ /dev/null
@@ -1,10 +0,0 @@
-// RUN: llvm-mc -triple=amdgpu9.42 -show-encoding %s | FileCheck --check-prefix=SUPPORTED %s
-// RUN: llvm-mc -triple=amdgpu9.50 -show-encoding %s | FileCheck --check-prefix=SUPPORTED %s
-// RUN: llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx9-4-generic -show-encoding %s | FileCheck --check-prefix=SUPPORTED %s
-// RUN: not llvm-mc -triple=amdgpu9.0a -filetype=null %s 2>&1 | FileCheck --check-prefix=GFX90A %s
-// RUN: not llvm-mc -triple=amdgpu9.42 -mattr=-buffer-inv-inst -filetype=null %s 2>&1 | FileCheck --check-prefix=DISABLED %s
-
-buffer_inv sc0 sc1
-// SUPPORTED: buffer_inv sc0 sc1 ; encoding: [0x00,0xc0,0xa4,0xe0,0x00,0x00,0x00,0x00]
-// GFX90A: error: instruction not supported on this GPU (gfx90a): buffer_inv
-// DISABLED: error: instruction not supported on this GPU (gfx942): buffer_inv
diff --git a/llvm/unittests/TargetParser/TargetParserTest.cpp b/llvm/unittests/TargetParser/TargetParserTest.cpp
index 2847942227df33..62d6f3a4c955b4 100644
--- a/llvm/unittests/TargetParser/TargetParserTest.cpp
+++ b/llvm/unittests/TargetParser/TargetParserTest.cpp
@@ -2875,18 +2875,6 @@ TEST(TargetParserTest, testAMDGPUgetFeatureBitset) {
EXPECT_TRUE(Empty.empty());
}
-TEST(TargetParserTest, testAMDGPUBufferInvInstFeature) {
- auto Has = [](AMDGPU::GPUKind AK) {
- return AMDGPU::getFeatureBitset(AK).test(AMDGPU::FEAT_BUFFER_INV_INST);
- };
-
- EXPECT_FALSE(Has(AMDGPU::GK_GFX90A));
- EXPECT_TRUE(Has(AMDGPU::GK_GFX942));
- EXPECT_TRUE(Has(AMDGPU::GK_GFX950));
- EXPECT_TRUE(Has(AMDGPU::GK_GFX9_4_GENERIC));
- EXPECT_FALSE(Has(AMDGPU::GK_GFX1250));
-}
-
TEST(TargetParserTest, testAMDGPUHalfAddressableLDSFeature) {
auto Has = [](AMDGPU::GPUKind AK) {
return AMDGPU::getFeatureBitset(AK).test(
>From 28aebcc0fd4e7a338f06537b0681183ed2b27347 Mon Sep 17 00:00:00 2001
From: Dan Zimmerman <danzimm at meta.com>
Date: Thu, 24 Sep 2026 14:48:12 -0700
Subject: [PATCH 5/7] [AMDGPU] Inherit predicates for GFX940 buffer
invalidation instructions
---
llvm/lib/Target/AMDGPU/BUFInstructions.td | 4 ----
1 file changed, 4 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/BUFInstructions.td b/llvm/lib/Target/AMDGPU/BUFInstructions.td
index 2b5d20c6849f78..6cb0f77e0dbeb3 100644
--- a/llvm/lib/Target/AMDGPU/BUFInstructions.td
+++ b/llvm/lib/Target/AMDGPU/BUFInstructions.td
@@ -3660,11 +3660,7 @@ let AsmString = BUFFER_WBL2.Mnemonic, // drop flags
defm BUFFER_WBL2 : MUBUF_Real_gfx90a<0x28>;
defm BUFFER_INVL2 : MUBUF_Real_gfx90a<0x29>;
-let SubtargetPredicate = isGFX940Plus in
def BUFFER_WBL2_gfx940 : MUBUF_Real_gfx940<0x28, BUFFER_WBL2>;
-
-let SubtargetPredicate = HasBufferInvInst,
- AssemblerPredicate = HasBufferInvInst in
def BUFFER_INV_gfx940 : MUBUF_Real_gfx940<0x29, BUFFER_INV>;
class MTBUF_Real_Base_vi <bits<4> op, MTBUF_Pseudo ps, int Enc> :
>From eb220861ac5a12f5e53106e16dbdb717610e1d13 Mon Sep 17 00:00:00 2001
From: Dan Zimmerman <danzimm at meta.com>
Date: Tue, 8 Sep 2026 13:55:51 -0700
Subject: [PATCH 6/7] [AMDGPU] Add buffer_inv intrinsic and builtin
---
clang/include/clang/Basic/BuiltinsAMDGPU.td | 1 +
.../clang/Basic/DiagnosticSemaKinds.td | 3 ++
clang/lib/Sema/SemaAMDGPU.cpp | 13 +++++++
.../builtins-amdgcn-buffer-inv.cl | 35 +++++++++++++++++++
llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 8 +++++
.../AMDGPU/AMDGPURegBankLegalizeRules.cpp | 1 +
llvm/lib/Target/AMDGPU/BUFInstructions.td | 1 +
.../CodeGen/AMDGPU/llvm.amdgcn.buffer.inv.ll | 22 ++++++++++++
.../Verifier/AMDGPU/llvm.amdgcn.buffer.inv.ll | 18 ++++++++++
9 files changed, 102 insertions(+)
create mode 100644 clang/test/CodeGenOpenCL/builtins-amdgcn-buffer-inv.cl
create mode 100644 llvm/test/CodeGen/AMDGPU/llvm.amdgcn.buffer.inv.ll
create mode 100644 llvm/test/Verifier/AMDGPU/llvm.amdgcn.buffer.inv.ll
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.td b/clang/include/clang/Basic/BuiltinsAMDGPU.td
index 4fd604390d3ce3..79b54b366e1d56 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.td
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.td
@@ -157,6 +157,7 @@ def __builtin_amdgcn_sched_group_barrier : AMDGPUBuiltin<"void(_Constant int, _C
def __builtin_amdgcn_iglp_opt : AMDGPUBuiltin<"void(_Constant int)">;
def __builtin_amdgcn_s_dcache_inv : AMDGPUBuiltin<"void()">;
def __builtin_amdgcn_buffer_wbinvl1 : AMDGPUBuiltin<"void()">;
+def __builtin_amdgcn_buffer_inv : AMDGPUBuiltin<"void(_Constant int)", [], "buffer-inv-inst">;
def __builtin_amdgcn_fence : AMDGPUBuiltin<"void(unsigned int, char const *, ...)">;
def __builtin_amdgcn_groupstaticsize : AMDGPUBuiltin<"unsigned int()">;
def __builtin_amdgcn_wavefrontsize : AMDGPUBuiltin<"unsigned int()", [Const]>;
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 01775af5d301ca..3d0861b1051cac 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -14336,6 +14336,9 @@ def note_amdgpu_named_barrier_reason_inherited : Note<
// AMDGCN builtins diagnostics
def err_amdgcn_load_lds_size_invalid_value : Error<"invalid size value">;
def note_amdgcn_load_lds_size_valid_value : Note<"size must be %select{1, 2, or 4|1, 2, 4, 12 or 16}0">;
+def err_amdgcn_buffer_inv_invalid_cpol
+ : Error<"argument to '__builtin_amdgcn_buffer_inv' must be a combination "
+ "of the sc0 (1) and sc1 (16) cache-policy bits">;
def err_invalid_sync_scope
: Error<"unsupported atomic synchronization scope '%0'">;
def err_amdgcn_processor_is_arg_not_literal
diff --git a/clang/lib/Sema/SemaAMDGPU.cpp b/clang/lib/Sema/SemaAMDGPU.cpp
index ee6d989d2b107a..44abf0f688aec6 100644
--- a/clang/lib/Sema/SemaAMDGPU.cpp
+++ b/clang/lib/Sema/SemaAMDGPU.cpp
@@ -21,6 +21,7 @@
#include "clang/Sema/Ownership.h"
#include "clang/Sema/Scope.h"
#include "clang/Sema/Sema.h"
+#include "llvm/ADT/STLExtras.h"
#include "llvm/ADT/SmallVector.h"
#include "llvm/ADT/StringExtras.h"
#include "llvm/ADT/StringMap.h"
@@ -168,6 +169,18 @@ bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(const TargetInfo &TI,
case AMDGPU::BI__builtin_amdgcn_s_setreg:
return SemaRef.BuiltinConstantArgRange(TheCall, /*ArgNum=*/0, /*Low=*/0,
/*High=*/UINT16_MAX);
+ case AMDGPU::BI__builtin_amdgcn_buffer_inv: {
+ llvm::APSInt CPol;
+ if (SemaRef.BuiltinConstantArg(TheCall, /*ArgNum=*/0, CPol))
+ return true;
+
+ // BUFFER_INV only supports the SC0 (1) and SC1 (16) cache-policy bits.
+ if (!llvm::is_contained({0u, 1u, 16u, 17u}, CPol.getZExtValue()))
+ return Diag(TheCall->getArg(0)->getExprLoc(),
+ diag::err_amdgcn_buffer_inv_invalid_cpol)
+ << TheCall->getArg(0)->getSourceRange();
+ return false;
+ }
case AMDGPU::BI__builtin_amdgcn_s_wait_event: {
llvm::APSInt Result;
if (SemaRef.BuiltinConstantArg(TheCall, 0, Result))
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-buffer-inv.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-buffer-inv.cl
new file mode 100644
index 00000000000000..257a3649c751a4
--- /dev/null
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-buffer-inv.cl
@@ -0,0 +1,35 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgpu9.4-amd-amdhsa -DTEST_VALID -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgpu9.4-amd-amdhsa -DTEST_INVALID -fsyntax-only -verify %s
+// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgpu9.0a-amd-amdhsa -DTEST_UNSUPPORTED -emit-llvm -o /dev/null -verify %s
+
+#ifdef TEST_VALID
+// CHECK-LABEL: @test_buffer_inv(
+// CHECK: call void @llvm.amdgcn.buffer.inv(i32 0)
+// CHECK: call void @llvm.amdgcn.buffer.inv(i32 1)
+// CHECK: call void @llvm.amdgcn.buffer.inv(i32 16)
+// CHECK: call void @llvm.amdgcn.buffer.inv(i32 17)
+void test_buffer_inv(void) {
+ __builtin_amdgcn_buffer_inv(0);
+ __builtin_amdgcn_buffer_inv(1);
+ __builtin_amdgcn_buffer_inv(16);
+ __builtin_amdgcn_buffer_inv(17);
+}
+#endif
+
+#ifdef TEST_INVALID
+void test_nonconstant(int cpol) {
+ __builtin_amdgcn_buffer_inv(cpol); // expected-error {{argument to '__builtin_amdgcn_buffer_inv' must be a constant integer}}
+}
+
+void test_invalid_policy(void) {
+ // expected-error at +1 {{must be a combination of the sc0 (1) and sc1 (16) cache-policy bits}}
+ __builtin_amdgcn_buffer_inv(2);
+}
+#endif
+
+#ifdef TEST_UNSUPPORTED
+void test_unsupported(void) {
+ __builtin_amdgcn_buffer_inv(0); // expected-error {{'__builtin_amdgcn_buffer_inv' needs target feature buffer-inv-inst}}
+}
+#endif
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index e7829eb29ba34f..374f3b9c3fdc35 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -2236,6 +2236,14 @@ def int_amdgcn_buffer_wbinvl1 :
ClangBuiltin<"__builtin_amdgcn_buffer_wbinvl1">,
DefaultAttrsIntrinsic<[], [], [IntrNoMem, IntrHasSideEffects]>;
+let TargetFeatures = "buffer-inv-inst" in
+def int_amdgcn_buffer_inv :
+ ClangBuiltin<"__builtin_amdgcn_buffer_inv">,
+ DefaultAttrsIntrinsic<[], [llvm_i32_ty],
+ [ImmArg<ArgIndex<0>>,
+ RangeSet<ArgIndex<0>, [[0, 1], [16, 17]]>,
+ IntrNoMem, IntrHasSideEffects]>;
+
def int_amdgcn_s_dcache_inv :
ClangBuiltin<"__builtin_amdgcn_s_dcache_inv">,
DefaultAttrsIntrinsic<[], [], [IntrNoMem, IntrHasSideEffects]>;
diff --git a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp
index 1d92bca3e12abd..4b48c05ccfb2d0 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp
@@ -1801,6 +1801,7 @@ RegBankLegalizeRules::RegBankLegalizeRules(const GCNSubtarget &_ST,
// Intrinsics with no register operands.
addRulesForIOpcs({amdgcn_asyncmark,
+ amdgcn_buffer_inv,
amdgcn_endpgm,
amdgcn_iglp_opt,
amdgcn_init_exec,
diff --git a/llvm/lib/Target/AMDGPU/BUFInstructions.td b/llvm/lib/Target/AMDGPU/BUFInstructions.td
index 6cb0f77e0dbeb3..b8776084b31ed8 100644
--- a/llvm/lib/Target/AMDGPU/BUFInstructions.td
+++ b/llvm/lib/Target/AMDGPU/BUFInstructions.td
@@ -1434,6 +1434,7 @@ def BUFFER_INV : MUBUF_Invalidate<"buffer_inv"> {
let has_sccb = 1;
let InOperandList = (ins CPol_0:$cpol);
let AsmOperands = "$cpol";
+ let Pattern = [(int_amdgcn_buffer_inv timm:$cpol)];
}
def BUFFER_GL0_INV : MUBUF_Invalidate<"buffer_gl0_inv">;
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.buffer.inv.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.buffer.inv.ll
new file mode 100644
index 00000000000000..e4717be9a36dc5
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.buffer.inv.ll
@@ -0,0 +1,22 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
+; RUN: llc -mtriple=amdgpu9.4-amd-amdhsa -verify-machineinstrs < %s | FileCheck --check-prefix=GFX94 %s
+; RUN: llc -mtriple=amdgpu9.4-amd-amdhsa -global-isel -global-isel-abort=1 -verify-machineinstrs < %s | FileCheck --check-prefix=GFX94 %s
+; RUN: not llc -mtriple=amdgpu9.0a-amd-amdhsa -filetype=null < %s 2>&1 | FileCheck --check-prefix=GFX90A-ERR %s
+
+; GFX90A-ERR: llvm.amdgcn.buffer.inv requires target feature 'buffer-inv-inst'
+
+declare void @llvm.amdgcn.buffer.inv(i32 immarg)
+
+define amdgpu_kernel void @buffer_inv_cache_policies() {
+; GFX94-LABEL: buffer_inv_cache_policies:
+; GFX94: ; %bb.0:
+; GFX94-NEXT: buffer_inv
+; GFX94-NEXT: buffer_inv sc0
+; GFX94-NEXT: buffer_inv sc1
+; GFX94-NEXT: buffer_inv sc0 sc1
+ call void @llvm.amdgcn.buffer.inv(i32 0)
+ call void @llvm.amdgcn.buffer.inv(i32 1)
+ call void @llvm.amdgcn.buffer.inv(i32 16)
+ call void @llvm.amdgcn.buffer.inv(i32 17)
+ ret void
+}
diff --git a/llvm/test/Verifier/AMDGPU/llvm.amdgcn.buffer.inv.ll b/llvm/test/Verifier/AMDGPU/llvm.amdgcn.buffer.inv.ll
new file mode 100644
index 00000000000000..26d2d8929cd2a7
--- /dev/null
+++ b/llvm/test/Verifier/AMDGPU/llvm.amdgcn.buffer.inv.ll
@@ -0,0 +1,18 @@
+; RUN: not llvm-as %s -disable-output 2>&1 | FileCheck %s
+
+declare void @llvm.amdgcn.buffer.inv(i32)
+
+define void @nonconstant(i32 %cpol) {
+ ; CHECK: immarg operand has non-immediate parameter
+ ; CHECK-NEXT: i32 %cpol
+ ; CHECK-NEXT: call void @llvm.amdgcn.buffer.inv(i32 %cpol)
+ call void @llvm.amdgcn.buffer.inv(i32 %cpol)
+ ret void
+}
+
+define void @invalid_cache_policy() {
+ ; CHECK: immarg value 2 for arg 0 out of range set
+ ; CHECK-NEXT: call void @llvm.amdgcn.buffer.inv(i32 2)
+ call void @llvm.amdgcn.buffer.inv(i32 2)
+ ret void
+}
>From 16de7f0e56b29d3366655d09568288be17a703b6 Mon Sep 17 00:00:00 2001
From: Dan Zimmerman <danzimm at meta.com>
Date: Thu, 10 Sep 2026 10:00:43 -0700
Subject: [PATCH 7/7] [MLIR][ROCDL] Expose buffer.inv intrinsic
---
mlir/include/mlir/Dialect/LLVMIR/ROCDLAttrs.td | 4 ++++
mlir/include/mlir/Dialect/LLVMIR/ROCDLEnums.td | 10 ++++++++++
mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td | 16 ++++++++++++++++
.../LLVMIR/rocdl-cache-policy-invalid.mlir | 8 ++++++++
mlir/test/Dialect/LLVMIR/rocdl.mlir | 13 +++++++++++++
mlir/test/Target/LLVMIR/rocdl.mlir | 13 +++++++++++++
6 files changed, 64 insertions(+)
diff --git a/mlir/include/mlir/Dialect/LLVMIR/ROCDLAttrs.td b/mlir/include/mlir/Dialect/LLVMIR/ROCDLAttrs.td
index 075fa75ab5ad87..013c451f4b0a0c 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/ROCDLAttrs.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/ROCDLAttrs.td
@@ -93,6 +93,10 @@ def ROCDL_Gfx942CachePolicyAttr
: ROCDL_IntrinsicIntegerEnumAttr<ROCDL_Gfx942CachePolicy,
"gfx942_cache_policy">;
+def ROCDL_BufferInvCachePolicyAttr
+ : ROCDL_IntrinsicIntegerEnumAttr<ROCDL_BufferInvCachePolicy,
+ "buffer_inv_cache_policy">;
+
def ROCDL_Gfx12CachePolicyAttr
: ROCDL_IntrinsicIntegerEnumAttr<ROCDL_Gfx12CachePolicy,
"gfx12_cache_policy">;
diff --git a/mlir/include/mlir/Dialect/LLVMIR/ROCDLEnums.td b/mlir/include/mlir/Dialect/LLVMIR/ROCDLEnums.td
index fb704ecf4c0de8..e6116be1f8d1a7 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/ROCDLEnums.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/ROCDLEnums.td
@@ -208,6 +208,16 @@ def ROCDL_Gfx942CachePolicy : ROCDL_I32BitEnum<"Gfx942CachePolicy",
let printBitEnumQuoted = 0;
}
+def ROCDL_BufferInvCachePolicy : ROCDL_I32BitEnum<"BufferInvCachePolicy",
+ "buffer invalidate cache policy bits",
+ [
+ ROCDL_Gfx942CachePolicyNone,
+ ROCDL_Gfx942CachePolicySC0,
+ ROCDL_Gfx942CachePolicySC1
+ ]> {
+ let printBitEnumQuoted = 0;
+}
+
def ROCDL_Gfx12CachePolicyNone : I32BitEnumCaseNone<"none">;
def ROCDL_Gfx12CachePolicyNT : I32BitEnumCaseBit<"nt", 0>;
def ROCDL_Gfx12CachePolicyHT : I32BitEnumCaseBit<"ht", 1>;
diff --git a/mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td b/mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td
index 716fd92e810ace..92cc04f91c6ec1 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td
@@ -322,6 +322,22 @@ def ROCDL_SWaitcntOp : ROCDL_ConcreteNonMemIntrOp<"s.waitcnt", [], 0, [0], ["bit
}];
}
+def ROCDL_BufferInvOp :
+ ROCDL_ConcreteNonMemIntrOp<"buffer.inv", [], 0, [0], ["cpol"]>,
+ Arguments<(ins ROCDL_BufferInvCachePolicyAttr:$cpol)> {
+ let summary = "Invalidate vector caches";
+ let description = [{
+ Issues `buffer_inv` with the selected `sc0` and `sc1` cache-policy flags.
+ The flags may be combined; `none` selects neither.
+
+ Example:
+ ```mlir
+ rocdl.buffer.inv sc0|sc1
+ ```
+ }];
+ let assemblyFormat = "enum($cpol) attr-dict";
+}
+
def ROCDL_SSleepOp : ROCDL_ConcreteNonMemIntrOp<"s.sleep", [], 0, [0], ["count"]>,
Arguments<(ins I32Attr:$count)> {
let assemblyFormat = "attr-dict $count";
diff --git a/mlir/test/Dialect/LLVMIR/rocdl-cache-policy-invalid.mlir b/mlir/test/Dialect/LLVMIR/rocdl-cache-policy-invalid.mlir
index 354d1ef661b6e9..769feacb5e69f0 100644
--- a/mlir/test/Dialect/LLVMIR/rocdl-cache-policy-invalid.mlir
+++ b/mlir/test/Dialect/LLVMIR/rocdl-cache-policy-invalid.mlir
@@ -36,3 +36,11 @@ llvm.func @atomic_buffer_rejects_gfx12(%rsrc : vector<4xi32>,
%0 = rocdl.raw.buffer.atomic.smax %vdata, %rsrc, %offset, %soffset, gfx12<nt> : i32
llvm.return
}
+
+// -----
+
+llvm.func @buffer_inv_rejects_unsupported_policy() {
+ // expected-error at +1 {{expected string or keyword containing one of the following enum values for attribute 'cpol'}}
+ rocdl.buffer.inv nt
+ llvm.return
+}
diff --git a/mlir/test/Dialect/LLVMIR/rocdl.mlir b/mlir/test/Dialect/LLVMIR/rocdl.mlir
index 1b5b6419368819..dd9bbe46e44860 100644
--- a/mlir/test/Dialect/LLVMIR/rocdl.mlir
+++ b/mlir/test/Dialect/LLVMIR/rocdl.mlir
@@ -1206,6 +1206,19 @@ llvm.func @rocdl.s.waitcnt() {
llvm.return
}
+llvm.func @rocdl.buffer.inv() {
+ // CHECK-LABEL: rocdl.buffer.inv
+ // CHECK-NEXT: rocdl.buffer.inv none
+ // CHECK-NEXT: rocdl.buffer.inv sc0
+ // CHECK-NEXT: rocdl.buffer.inv sc1
+ // CHECK-NEXT: rocdl.buffer.inv sc0|sc1
+ rocdl.buffer.inv none
+ rocdl.buffer.inv sc0
+ rocdl.buffer.inv sc1
+ rocdl.buffer.inv sc0|sc1
+ llvm.return
+}
+
llvm.func @rocdl.s.sleep() {
// CHECK-LABEL: rocdl.s.sleep
// CHECK: rocdl.s.sleep 0
diff --git a/mlir/test/Target/LLVMIR/rocdl.mlir b/mlir/test/Target/LLVMIR/rocdl.mlir
index b4a94cef770fac..ff4dea36024af2 100644
--- a/mlir/test/Target/LLVMIR/rocdl.mlir
+++ b/mlir/test/Target/LLVMIR/rocdl.mlir
@@ -240,6 +240,19 @@ llvm.func @rocdl.s.waitcnt() {
llvm.return
}
+llvm.func @rocdl.buffer.inv() {
+ // CHECK-LABEL: rocdl.buffer.inv
+ // CHECK-NEXT: call void @llvm.amdgcn.buffer.inv(i32 0)
+ // CHECK-NEXT: call void @llvm.amdgcn.buffer.inv(i32 1)
+ // CHECK-NEXT: call void @llvm.amdgcn.buffer.inv(i32 16)
+ // CHECK-NEXT: call void @llvm.amdgcn.buffer.inv(i32 17)
+ rocdl.buffer.inv none
+ rocdl.buffer.inv sc0
+ rocdl.buffer.inv sc1
+ rocdl.buffer.inv sc0|sc1
+ llvm.return
+}
+
llvm.func @rocdl.s.sleep() {
// CHECK-LABEL: rocdl.s.sleep
// CHECK-NEXT: call void @llvm.amdgcn.s.sleep(i32 0)
More information about the llvm-commits
mailing list