[Mlir-commits] [clang] [llvm] [mlir] [AMDGPU] Introduce builtin, intrinsic and ROCDL instruction for buffer_inv (PR #228497)
Dan Zimmerman
llvmlistbot at llvm.org
Tue Oct 20 11:53:55 PDT 2026
https://github.com/danzimm updated https://github.com/llvm/llvm-project/pull/228497
>From 1886b11f71735dcfe0ed742ac2f3b51427c38170 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 1/5] [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 34a1bd2d301a37..770f3af8d0e0fa 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 e6bdf2e7c41842..001689e849b7c1 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -14437,6 +14437,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 b1f6fa33649bd1..11472b1689f6c9 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 ed80bc4628b6fa..d8020748a93888 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -2331,6 +2331,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 acc501dc2bcef7..672fc772534f9c 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp
@@ -1805,6 +1805,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 3964b477f6672f..25e06698f48065 100644
--- a/llvm/lib/Target/AMDGPU/BUFInstructions.td
+++ b/llvm/lib/Target/AMDGPU/BUFInstructions.td
@@ -1432,6 +1432,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 75f3ede914ccf689fb77b1c34a33746bea2b4f96 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 2/5] [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 2f067b8ba98fda..6e6bb2c13dcbf4 100644
--- a/mlir/test/Target/LLVMIR/rocdl.mlir
+++ b/mlir/test/Target/LLVMIR/rocdl.mlir
@@ -247,6 +247,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)
>From 3010fd0cff00c35a70a653d85ad903de336de50c Mon Sep 17 00:00:00 2001
From: Dan Zimmerman <danzimm at meta.com>
Date: Tue, 6 Oct 2026 11:20:13 -0700
Subject: [PATCH 3/5] [AMDGPU] Move buffer_inv diagnostics to SemaOpenCL
---
.../test/CodeGenOpenCL/builtins-amdgcn-buffer-inv.cl | 12 ------------
clang/test/SemaOpenCL/builtins-amdgcn-buffer-inv.cl | 11 +++++++++++
2 files changed, 11 insertions(+), 12 deletions(-)
create mode 100644 clang/test/SemaOpenCL/builtins-amdgcn-buffer-inv.cl
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-buffer-inv.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-buffer-inv.cl
index 257a3649c751a4..b74e4f0d5129e0 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-buffer-inv.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-buffer-inv.cl
@@ -1,6 +1,5 @@
// 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
@@ -17,17 +16,6 @@ void test_buffer_inv(void) {
}
#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}}
diff --git a/clang/test/SemaOpenCL/builtins-amdgcn-buffer-inv.cl b/clang/test/SemaOpenCL/builtins-amdgcn-buffer-inv.cl
new file mode 100644
index 00000000000000..cc944d2b4d6624
--- /dev/null
+++ b/clang/test/SemaOpenCL/builtins-amdgcn-buffer-inv.cl
@@ -0,0 +1,11 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgpu9.4-amd-amdhsa -fsyntax-only -verify %s
+
+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);
+}
>From 53f3cf75cdeb9dc68614713c859c999d87425a5a Mon Sep 17 00:00:00 2001
From: Dan Zimmerman <danzimm at meta.com>
Date: Tue, 6 Oct 2026 11:21:32 -0700
Subject: [PATCH 4/5] [AMDGPU] Refresh buffer_inv intrinsic codegen test
---
llvm/test/CodeGen/AMDGPU/llvm.amdgcn.buffer.inv.ll | 7 ++++---
1 file changed, 4 insertions(+), 3 deletions(-)
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.buffer.inv.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.buffer.inv.ll
index e4717be9a36dc5..20fecd79569b08 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.buffer.inv.ll
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.buffer.inv.ll
@@ -1,6 +1,6 @@
-; 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
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
+; RUN: llc -mtriple=amdgpu9.4-amd-amdhsa < %s | FileCheck --check-prefix=GFX94 %s
+; RUN: llc -mtriple=amdgpu9.4-amd-amdhsa -global-isel < %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'
@@ -14,6 +14,7 @@ define amdgpu_kernel void @buffer_inv_cache_policies() {
; GFX94-NEXT: buffer_inv sc0
; GFX94-NEXT: buffer_inv sc1
; GFX94-NEXT: buffer_inv sc0 sc1
+; GFX94-NEXT: s_endpgm
call void @llvm.amdgcn.buffer.inv(i32 0)
call void @llvm.amdgcn.buffer.inv(i32 1)
call void @llvm.amdgcn.buffer.inv(i32 16)
>From 99ed7e716652b907d46ce7761476c4877b889082 Mon Sep 17 00:00:00 2001
From: Dan Zimmerman <danzimm at meta.com>
Date: Tue, 13 Oct 2026 06:38:51 -0700
Subject: [PATCH 5/5] Use ConfinedAttr in ROCDL changes
---
.../include/mlir/Dialect/LLVMIR/ROCDLAttrs.td | 18 ++++++++++--
.../include/mlir/Dialect/LLVMIR/ROCDLEnums.td | 10 -------
.../LLVMIR/rocdl-cache-policy-invalid.mlir | 28 +++++++++++++++++--
3 files changed, 41 insertions(+), 15 deletions(-)
diff --git a/mlir/include/mlir/Dialect/LLVMIR/ROCDLAttrs.td b/mlir/include/mlir/Dialect/LLVMIR/ROCDLAttrs.td
index 013c451f4b0a0c..a2304e7144fe64 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/ROCDLAttrs.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/ROCDLAttrs.td
@@ -93,9 +93,21 @@ def ROCDL_Gfx942CachePolicyAttr
: ROCDL_IntrinsicIntegerEnumAttr<ROCDL_Gfx942CachePolicy,
"gfx942_cache_policy">;
-def ROCDL_BufferInvCachePolicyAttr
- : ROCDL_IntrinsicIntegerEnumAttr<ROCDL_BufferInvCachePolicy,
- "buffer_inv_cache_policy">;
+// BUFFER_INV only supports SC0 (bit 0) and SC1 (bit 4).
+def ROCDL_BufferInvCachePolicyBits : AttrConstraint<
+ CPred<[{
+ (static_cast<uint32_t>(
+ ::llvm::cast<::mlir::ROCDL::Gfx942CachePolicyAttr>($_self)
+ .getValue()) &
+ ~static_cast<uint32_t>(
+ ::mlir::ROCDL::Gfx942CachePolicy::sc0 |
+ ::mlir::ROCDL::Gfx942CachePolicy::sc1)) == 0
+ }]>,
+ "with no bits other than sc0 and sc1 set">;
+
+def ROCDL_BufferInvCachePolicyAttr :
+ ConfinedAttr<ROCDL_Gfx942CachePolicyAttr,
+ [ROCDL_BufferInvCachePolicyBits]>;
def ROCDL_Gfx12CachePolicyAttr
: ROCDL_IntrinsicIntegerEnumAttr<ROCDL_Gfx12CachePolicy,
diff --git a/mlir/include/mlir/Dialect/LLVMIR/ROCDLEnums.td b/mlir/include/mlir/Dialect/LLVMIR/ROCDLEnums.td
index e6116be1f8d1a7..fb704ecf4c0de8 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/ROCDLEnums.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/ROCDLEnums.td
@@ -208,16 +208,6 @@ 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/test/Dialect/LLVMIR/rocdl-cache-policy-invalid.mlir b/mlir/test/Dialect/LLVMIR/rocdl-cache-policy-invalid.mlir
index 769feacb5e69f0..318482a47ae64f 100644
--- a/mlir/test/Dialect/LLVMIR/rocdl-cache-policy-invalid.mlir
+++ b/mlir/test/Dialect/LLVMIR/rocdl-cache-policy-invalid.mlir
@@ -39,8 +39,32 @@ llvm.func @atomic_buffer_rejects_gfx12(%rsrc : vector<4xi32>,
// -----
-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'}}
+llvm.func @buffer_inv_rejects_nt() {
+ // expected-error at +1 {{failed to satisfy constraint: gfx942 buffer cache policy bits with no bits other than sc0 and sc1 set}}
rocdl.buffer.inv nt
llvm.return
}
+
+// -----
+
+llvm.func @buffer_inv_rejects_swz() {
+ // expected-error at +1 {{failed to satisfy constraint: gfx942 buffer cache policy bits with no bits other than sc0 and sc1 set}}
+ rocdl.buffer.inv swz
+ llvm.return
+}
+
+// -----
+
+llvm.func @buffer_inv_rejects_volatile() {
+ // expected-error at +1 {{failed to satisfy constraint: gfx942 buffer cache policy bits with no bits other than sc0 and sc1 set}}
+ rocdl.buffer.inv volatile
+ llvm.return
+}
+
+// -----
+
+llvm.func @buffer_inv_rejects_mixed_policy() {
+ // expected-error at +1 {{failed to satisfy constraint: gfx942 buffer cache policy bits with no bits other than sc0 and sc1 set}}
+ rocdl.buffer.inv sc0|nt
+ llvm.return
+}
More information about the Mlir-commits
mailing list