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 <[email protected]> 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 <[email protected]> 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 <[email protected]> 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 <[email protected]> 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 <[email protected]> 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 <[email protected]> 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@+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 <[email protected]> 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@+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) _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
