[clang] [llvm] [AMDGPU] Add llvm.amdgcn.s.setprio.relaxed intrinsic (PR #204850)
Frederik Harwath via cfe-commits
cfe-commits at lists.llvm.org
Thu Aug 27 08:00:40 PDT 2026
https://github.com/frederik-h updated https://github.com/llvm/llvm-project/pull/204850
>From 0ebed862ba7f75c767876e33f325438d6b858f49 Mon Sep 17 00:00:00 2001
From: Frederik Harwath <fharwath at amd.com>
Date: Thu, 18 Jun 2026 09:58:18 -0400
Subject: [PATCH 01/10] [AMDGPU] Add scheduling mask operand to S_SETPRIO
Add a mask operand to S_SETPRIO that controls which instruction types
can be scheduled across it, using the same semantics as SCHED_BARRIER.
The default mask 0 creates a full scheduling boundary, preserving
existing behavior.
Add new llvm.amdgcn.s.setprio.mask(i16, i32) intrinsic with explicit
mask parameter.
---
llvm/docs/AMDGPUUsage.rst | 8 +++++
llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 6 ++++
llvm/lib/Target/AMDGPU/AMDGPUIGroupLP.cpp | 18 ++++++----
.../AMDGPU/AMDGPURegBankLegalizeRules.cpp | 1 +
.../Target/AMDGPU/AMDGPUSetWavePriority.cpp | 3 +-
.../lib/Target/AMDGPU/GCNHazardRecognizer.cpp | 9 +++--
llvm/lib/Target/AMDGPU/SIInstrInfo.cpp | 19 +++++++++--
llvm/lib/Target/AMDGPU/SIInstrInfo.h | 3 ++
llvm/lib/Target/AMDGPU/SOPInstructions.td | 18 ++++++++--
.../CodeGen/AMDGPU/llvm.amdgcn.s.setprio.ll | 33 +++++++++++++++++++
.../AMDGPU/required-export-priority.mir | 12 +++----
llvm/test/CodeGen/AMDGPU/sched-setprio.ll | 26 +++++++++++++++
12 files changed, 136 insertions(+), 20 deletions(-)
diff --git a/llvm/docs/AMDGPUUsage.rst b/llvm/docs/AMDGPUUsage.rst
index dba0997e4f099..57e51dc4bc1af 100644
--- a/llvm/docs/AMDGPUUsage.rst
+++ b/llvm/docs/AMDGPUUsage.rst
@@ -1724,6 +1724,14 @@ The AMDGPU backend implements the following LLVM IR intrinsics.
- 0x0400: All Transcendental (e.g. V_EXP) instructions may be scheduled across sched_barrier.
- 0x0800: All LDSDMA instructions may be scheduled across sched_barrier.
+ llvm.amdgcn.s.setprio.mask Sets the hardware wave priority and controls instruction scheduling across
+ the intrinsic in the same way as llvm.amdgcn.sched.barrier. Takes two parameters:
+
+ - Priority (i16): Hardware wave priority level. 0 = lowest, 3 = highest.
+ - Mask (i32): Scheduling barrier mask.
+
+ llvm.amdgcn.s.setprio Like llvm.amdgcn.s.setprio.mask with Mask = 0.
+
llvm.amdgcn.sched.group.barrier Creates schedule groups with specific properties to create custom scheduling
pipelines. The ordering between groups is enforced by the instruction scheduler.
The intrinsic applies to the code that precedes the intrinsic. The intrinsic
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index f1659f0cd803a..d478077321d67 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -2224,6 +2224,12 @@ def int_amdgcn_s_setprio :
DefaultAttrsIntrinsic<[], [llvm_i16_ty], [ImmArg<ArgIndex<0>>, IntrNoMem,
IntrHasSideEffects]>;
+def int_amdgcn_s_setprio_mask :
+ ClangBuiltin<"__builtin_amdgcn_s_setprio_mask">,
+ DefaultAttrsIntrinsic<[], [llvm_i16_ty, llvm_i32_ty],
+ [ImmArg<ArgIndex<0>>, ImmArg<ArgIndex<1>>,
+ IntrNoMem, IntrHasSideEffects]>;
+
def int_amdgcn_s_setprio_inc_wg :
ClangBuiltin<"__builtin_amdgcn_s_setprio_inc_wg">,
DefaultAttrsIntrinsic<[], [llvm_i16_ty], [ImmArg<ArgIndex<0>>, IntrNoMem,
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUIGroupLP.cpp b/llvm/lib/Target/AMDGPU/AMDGPUIGroupLP.cpp
index 6c95dc76f4f6f..930330a38715b 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUIGroupLP.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUIGroupLP.cpp
@@ -2604,7 +2604,7 @@ void IGroupLPDAGMutation::apply(ScheduleDAGInstrs *DAGInstrs) {
for (auto R = DAG->SUnits.rbegin(), E = DAG->SUnits.rend(); R != E; ++R) {
unsigned Opc = R->getInstr()->getOpcode();
// SCHED_[GROUP_]BARRIER and IGLP are mutually exclusive.
- if (Opc == AMDGPU::SCHED_BARRIER) {
+ if (TII->isSchedBarrierLike(*R->getInstr())) {
addSchedBarrierEdges(*R);
FoundSB = true;
} else if (Opc == AMDGPU::SCHED_GROUP_BARRIER) {
@@ -2629,11 +2629,17 @@ void IGroupLPDAGMutation::apply(ScheduleDAGInstrs *DAGInstrs) {
void IGroupLPDAGMutation::addSchedBarrierEdges(SUnit &SchedBarrier) {
MachineInstr &MI = *SchedBarrier.getInstr();
- assert(MI.getOpcode() == AMDGPU::SCHED_BARRIER);
- LLVM_DEBUG(dbgs() << "Building SchedGroup for SchedBarrier with Mask: "
- << MI.getOperand(0).getImm() << "\n");
- auto InvertedMask =
- invertSchedBarrierMask((SchedGroupMask)MI.getOperand(0).getImm());
+ assert(TII->isSchedBarrierLike(MI));
+
+ SchedGroupMask Mask =
+ static_cast<SchedGroupMask>(TII->getSchedBarrierLikeMask(MI));
+
+ LLVM_DEBUG(
+ dbgs()
+ << "Building SchedGroup for SchedBarrier-like instruction with Mask: "
+ << (unsigned)Mask << "\n");
+
+ auto InvertedMask = invertSchedBarrierMask(Mask);
SchedGroup SG(InvertedMask, std::nullopt, DAG, TII);
for (SUnit &SU : DAG->SUnits)
diff --git a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp
index b3dbc3fb5c97c..98fe88d6c48e0 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp
@@ -1790,6 +1790,7 @@ RegBankLegalizeRules::RegBankLegalizeRules(const GCNSubtarget &_ST,
amdgcn_s_nop,
amdgcn_s_sethalt,
amdgcn_s_setprio,
+ amdgcn_s_setprio_mask,
amdgcn_s_setprio_inc_wg,
amdgcn_s_sleep,
amdgcn_s_ttracedata_imm,
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUSetWavePriority.cpp b/llvm/lib/Target/AMDGPU/AMDGPUSetWavePriority.cpp
index 56aa3f6db83ad..6916a2561b2ef 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUSetWavePriority.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUSetWavePriority.cpp
@@ -85,7 +85,8 @@ AMDGPUSetWavePriority::BuildSetprioMI(MachineBasicBlock &MBB,
MachineBasicBlock::iterator I,
unsigned priority) const {
return BuildMI(MBB, I, DebugLoc(), TII->get(AMDGPU::S_SETPRIO))
- .addImm(priority);
+ .addImm(priority)
+ .addImm(0);
}
// Checks that for every predecessor Pred that can reach a VMEM load,
diff --git a/llvm/lib/Target/AMDGPU/GCNHazardRecognizer.cpp b/llvm/lib/Target/AMDGPU/GCNHazardRecognizer.cpp
index 2f2c05329c0d5..bf671c019225e 100644
--- a/llvm/lib/Target/AMDGPU/GCNHazardRecognizer.cpp
+++ b/llvm/lib/Target/AMDGPU/GCNHazardRecognizer.cpp
@@ -3736,7 +3736,8 @@ static bool ensureEntrySetPrio(MachineFunction *MF, int Priority,
}
BuildMI(EntryMBB, EntryMBB.begin(), DebugLoc(), TII.get(AMDGPU::S_SETPRIO))
- .addImm(Priority);
+ .addImm(Priority)
+ .addImm(0);
return true;
}
@@ -3814,7 +3815,8 @@ bool GCNHazardRecognizer::fixRequiredExportPriority(MachineInstr *MI) {
// Lower priority.
BuildMI(*MBB, NextMI, DL, TII.get(AMDGPU::S_SETPRIO))
- .addImm(PostExportPriority);
+ .addImm(PostExportPriority)
+ .addImm(0);
if (!EndOfShader) {
// Wait for exports to complete.
@@ -3829,7 +3831,8 @@ bool GCNHazardRecognizer::fixRequiredExportPriority(MachineInstr *MI) {
if (!EndOfShader) {
// Return to normal (higher) priority.
BuildMI(*MBB, NextMI, DL, TII.get(AMDGPU::S_SETPRIO))
- .addImm(NormalPriority);
+ .addImm(NormalPriority)
+ .addImm(0);
}
return true;
diff --git a/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp b/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp
index ba1f1b2effe9f..0d7eb1fcc61ff 100644
--- a/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp
@@ -4559,6 +4559,22 @@ static bool changesVGPRIndexingMode(const MachineInstr &MI) {
}
}
+bool SIInstrInfo::isSchedBarrierLike(const MachineInstr &MI) const {
+ unsigned Opc = MI.getOpcode();
+ return Opc == AMDGPU::SCHED_BARRIER || Opc == AMDGPU::S_SETPRIO;
+}
+
+unsigned SIInstrInfo::getSchedBarrierLikeMask(const MachineInstr &MI) const {
+ switch (MI.getOpcode()) {
+ case AMDGPU::SCHED_BARRIER:
+ return MI.getOperand(0).getImm();
+ case AMDGPU::S_SETPRIO:
+ return MI.getOperand(1).getImm();
+ default:
+ llvm_unreachable("Expected a SchedBarrier-like instruction");
+ }
+}
+
bool SIInstrInfo::isSchedulingBoundary(const MachineInstr &MI,
const MachineBasicBlock *MBB,
const MachineFunction &MF) const {
@@ -4577,7 +4593,7 @@ bool SIInstrInfo::isSchedulingBoundary(const MachineInstr &MI,
if (MI.getOpcode() == TargetOpcode::INLINEASM_BR)
return true;
- if (MI.getOpcode() == AMDGPU::SCHED_BARRIER && MI.getOperand(0).getImm() == 0)
+ if (isSchedBarrierLike(MI) && getSchedBarrierLikeMask(MI) == 0)
return true;
// Target-independent instructions do not have an implicit-use of EXEC, even
@@ -4586,7 +4602,6 @@ bool SIInstrInfo::isSchedulingBoundary(const MachineInstr &MI,
return MI.modifiesRegister(AMDGPU::EXEC, &RI) ||
MI.getOpcode() == AMDGPU::S_SETREG_IMM32_B32 ||
MI.getOpcode() == AMDGPU::S_SETREG_B32 ||
- MI.getOpcode() == AMDGPU::S_SETPRIO ||
MI.getOpcode() == AMDGPU::S_SETPRIO_INC_WG ||
changesVGPRIndexingMode(MI);
}
diff --git a/llvm/lib/Target/AMDGPU/SIInstrInfo.h b/llvm/lib/Target/AMDGPU/SIInstrInfo.h
index 228549b7e383f..b1627a5a8c4b5 100644
--- a/llvm/lib/Target/AMDGPU/SIInstrInfo.h
+++ b/llvm/lib/Target/AMDGPU/SIInstrInfo.h
@@ -1172,6 +1172,9 @@ class SIInstrInfo final : public AMDGPUGenInstrInfo {
Opcode == AMDGPU::DS_GWS_BARRIER;
}
+ bool isSchedBarrierLike(const MachineInstr &MI) const;
+ unsigned getSchedBarrierLikeMask(const MachineInstr &MI) const;
+
static bool isLoadMonitor(unsigned Opc) {
switch (Opc) {
case AMDGPU::GLOBAL_LOAD_MONITOR_B32:
diff --git a/llvm/lib/Target/AMDGPU/SOPInstructions.td b/llvm/lib/Target/AMDGPU/SOPInstructions.td
index 01dd444c477e2..4ed247086e27e 100644
--- a/llvm/lib/Target/AMDGPU/SOPInstructions.td
+++ b/llvm/lib/Target/AMDGPU/SOPInstructions.td
@@ -1753,10 +1753,24 @@ def S_SLEEP_VAR : SOP1_0_32 <"s_sleep_var", [(int_amdgcn_s_sleep_var SSrc_b32:$s
let hasSideEffects = 1;
}
-def S_SETPRIO : SOPP_Pseudo <"s_setprio", (ins i16imm:$simm16), "$simm16",
- [(int_amdgcn_s_setprio timm:$simm16)]> {
+def S_SETPRIO : SOPP_Pseudo <"s_setprio", (ins i16imm:$simm16, i32imm:$mask), "$simm16",
+ []> {
}
+} // End hasSideEffects = 1
+
+def : GCNPat<
+ (int_amdgcn_s_setprio timm:$priority),
+ (S_SETPRIO timm:$priority, (i32 0))
+>;
+
+def : GCNPat<
+ (int_amdgcn_s_setprio_mask timm:$priority, timm:$mask),
+ (S_SETPRIO timm:$priority, timm:$mask)
+>;
+
+let hasSideEffects = 1 in {
+
def S_SETPRIO_INC_WG : SOPP_Pseudo <"s_setprio_inc_wg", (ins i16imm:$simm16), "$simm16",
[(int_amdgcn_s_setprio_inc_wg timm:$simm16)]> {
let SubtargetPredicate = HasSetPrioIncWgInst;
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.setprio.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.setprio.ll
index 79d4f28254da6..0078f33782bf1 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.setprio.ll
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.setprio.ll
@@ -5,6 +5,7 @@
; RUN: llc -global-isel -new-reg-bank-select -mtriple=amdgcn -mcpu=gfx90a -show-mc-encoding < %s | FileCheck -check-prefix=GFX9 %s
declare void @llvm.amdgcn.s.setprio(i16) #0
+declare void @llvm.amdgcn.s.setprio.mask(i16, i32) #0
define void @test_llvm_amdgcn_s_setprio() #0 {
; GFX9-LABEL: test_llvm_amdgcn_s_setprio:
@@ -46,4 +47,36 @@ define void @test_llvm_amdgcn_s_setprio() #0 {
ret void
}
+define void @test_llvm_amdgcn_s_setprio_mask_full() #0 {
+; GFX9-LABEL: test_llvm_amdgcn_s_setprio_mask_full:
+; GFX9: ; %bb.0:
+; GFX9-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0) ; encoding: [0x00,0x00,0x8c,0xbf]
+; GFX9-NEXT: s_setprio 3 ; encoding: [0x03,0x00,0x8f,0xbf]
+; GFX9-NEXT: s_setpc_b64 s[30:31] ; encoding: [0x1e,0x1d,0x80,0xbe]
+;
+; SI-LABEL: test_llvm_amdgcn_s_setprio_mask_full:
+; SI: ; %bb.0:
+; SI-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0) ; encoding: [0x00,0x00,0x8c,0xbf]
+; SI-NEXT: s_setprio 3 ; encoding: [0x03,0x00,0x8f,0xbf]
+; SI-NEXT: s_setpc_b64 s[30:31] ; encoding: [0x1e,0x20,0x80,0xbe]
+ call void @llvm.amdgcn.s.setprio.mask(i16 3, i32 0)
+ ret void
+}
+
+define void @test_llvm_amdgcn_s_setprio_mask_selective() #0 {
+; GFX9-LABEL: test_llvm_amdgcn_s_setprio_mask_selective:
+; GFX9: ; %bb.0:
+; GFX9-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0) ; encoding: [0x00,0x00,0x8c,0xbf]
+; GFX9-NEXT: s_setprio 1 ; encoding: [0x01,0x00,0x8f,0xbf]
+; GFX9-NEXT: s_setpc_b64 s[30:31] ; encoding: [0x1e,0x1d,0x80,0xbe]
+;
+; SI-LABEL: test_llvm_amdgcn_s_setprio_mask_selective:
+; SI: ; %bb.0:
+; SI-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0) ; encoding: [0x00,0x00,0x8c,0xbf]
+; SI-NEXT: s_setprio 1 ; encoding: [0x01,0x00,0x8f,0xbf]
+; SI-NEXT: s_setpc_b64 s[30:31] ; encoding: [0x1e,0x20,0x80,0xbe]
+ call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 15)
+ ret void
+}
+
attributes #0 = { nounwind }
diff --git a/llvm/test/CodeGen/AMDGPU/required-export-priority.mir b/llvm/test/CodeGen/AMDGPU/required-export-priority.mir
index 70241033f35a8..9039ee14715ca 100644
--- a/llvm/test/CodeGen/AMDGPU/required-export-priority.mir
+++ b/llvm/test/CodeGen/AMDGPU/required-export-priority.mir
@@ -334,15 +334,15 @@ body: |
bb.1:
liveins: $vgpr0
- S_SETPRIO 3
+ S_SETPRIO 3, 0
$vgpr0 = V_OR_B32_e32 2, $vgpr0, implicit $exec
- S_SETPRIO 0
+ S_SETPRIO 0, 0
bb.2:
liveins: $vgpr0
- S_SETPRIO 1
+ S_SETPRIO 1, 0
$vgpr0 = V_OR_B32_e32 3, $vgpr0, implicit $exec
- S_SETPRIO 0
+ S_SETPRIO 0, 0
bb.3:
liveins: $vgpr0
@@ -381,9 +381,9 @@ body: |
; GFX1150-NEXT: S_SETPRIO 2
; GFX1150-NEXT: S_SETPRIO 3
; GFX1150-NEXT: S_ENDPGM 0
- S_SETPRIO 3
+ S_SETPRIO 3, 0
EXP 1, $vgpr0, $vgpr0, $vgpr0, $vgpr0, -1, -1, 15, implicit $exec
EXP_DONE 0, $vgpr0, $vgpr0, $vgpr0, $vgpr0, -1, -1, 15, implicit $exec
- S_SETPRIO 3
+ S_SETPRIO 3, 0
S_ENDPGM 0
...
diff --git a/llvm/test/CodeGen/AMDGPU/sched-setprio.ll b/llvm/test/CodeGen/AMDGPU/sched-setprio.ll
index 78a14719c48da..53f1fa1f4002e 100644
--- a/llvm/test/CodeGen/AMDGPU/sched-setprio.ll
+++ b/llvm/test/CodeGen/AMDGPU/sched-setprio.ll
@@ -1,6 +1,7 @@
; RUN: llc -mtriple=amdgcn -mcpu=gfx908 < %s | FileCheck --check-prefix=GCN %s
declare void @llvm.amdgcn.s.setprio(i16)
+declare void @llvm.amdgcn.s.setprio.mask(i16, i32)
declare <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float, float, <4 x float>, i32, i32, i32)
; GCN-LABEL: {{^}}test_mfma_f32_4x4x1f32:
@@ -18,3 +19,28 @@ bb:
store <4 x float> %mai.2, ptr addrspace(1) %arg
ret void
}
+
+; GCN-LABEL: {{^}}test_setprio_mask0_blocks_salu:
+; GCN: s_setprio 1
+; GCN-NEXT: s_add_i32
+define amdgpu_cs void @test_setprio_mask0_blocks_salu(ptr addrspace(1) %out, i32 inreg %x, i32 inreg %y) {
+ %add1 = add i32 %x, 1
+ call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 0)
+ %add2 = add i32 %y, 2
+ %sum = add i32 %add1, %add2
+ store i32 %sum, ptr addrspace(1) %out
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_setprio_mask4_allows_salu:
+; GCN: s_add_i32
+; GCN-NEXT: s_add_i32
+; GCN-NEXT: s_setprio 1
+define amdgpu_cs void @test_setprio_mask4_allows_salu(ptr addrspace(1) %out, i32 inreg %x, i32 inreg %y) {
+ %add1 = add i32 %x, 1
+ call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 4)
+ %add2 = add i32 %y, 2
+ %sum = add i32 %add1, %add2
+ store i32 %sum, ptr addrspace(1) %out
+ ret void
+}
>From c0c07f49451058bfd0671dbc0bc5a602196156c6 Mon Sep 17 00:00:00 2001
From: Frederik Harwath <fharwath at amd.com>
Date: Mon, 22 Jun 2026 04:10:41 -0400
Subject: [PATCH 02/10] [AMDGPU] Add more tests for S_SETPRIO scheduling mask
Add tests demonstrating the scheduling behavior of each mask bit type.
ALU, VALU, SALU, MFMA, TRANS can move, memory ops cannot.
---
llvm/test/CodeGen/AMDGPU/sched-setprio.ll | 261 ++++++++++++++++++++--
1 file changed, 241 insertions(+), 20 deletions(-)
diff --git a/llvm/test/CodeGen/AMDGPU/sched-setprio.ll b/llvm/test/CodeGen/AMDGPU/sched-setprio.ll
index 53f1fa1f4002e..f8dceed6cc950 100644
--- a/llvm/test/CodeGen/AMDGPU/sched-setprio.ll
+++ b/llvm/test/CodeGen/AMDGPU/sched-setprio.ll
@@ -1,46 +1,267 @@
; RUN: llc -mtriple=amdgcn -mcpu=gfx908 < %s | FileCheck --check-prefix=GCN %s
+; Test scheduling mask behavior for llvm.amdgcn.s.setprio{,.mask} intrinsics.
+
declare void @llvm.amdgcn.s.setprio(i16)
declare void @llvm.amdgcn.s.setprio.mask(i16, i32)
+declare float @llvm.amdgcn.rcp.f32(float)
declare <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float, float, <4 x float>, i32, i32, i32)
-; GCN-LABEL: {{^}}test_mfma_f32_4x4x1f32:
-; GCN: s_setprio 1
-; GCN: v_mfma
-; GCN: v_mfma
-; GCN: s_setprio 0
-define amdgpu_kernel void @test_mfma_f32_4x4x1f32(ptr addrspace(1) %arg) #0 {
-bb:
- %in.1 = load <4 x float>, ptr addrspace(1) %arg
- call void @llvm.amdgcn.s.setprio(i16 1)
- %mai.1 = tail call <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float 1.0, float 2.0, <4 x float> %in.1, i32 0, i32 0, i32 0)
- %mai.2 = tail call <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float 3.0, float 4.0, <4 x float> %mai.1, i32 0, i32 0, i32 0)
- call void @llvm.amdgcn.s.setprio(i16 0)
- store <4 x float> %mai.2, ptr addrspace(1) %arg
+; GCN-LABEL: {{^}}test_mask0_blocks_salu:
+; GCN: s_setprio 1
+; GCN-NEXT: s_add_i32
+define amdgpu_cs void @test_mask0_blocks_salu(ptr addrspace(1) %out, i32 inreg %x, i32 inreg %y) {
+ %add1 = add i32 %x, 1
+ call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 0)
+ %add2 = add i32 %y, 2
+ %sum = add i32 %add1, %add2
+ store i32 %sum, ptr addrspace(1) %out
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_mask0_blocks_valu:
+; GCN: v_add_f32_e32 v{{[0-9]+}}, 1.0
+; GCN-NEXT: s_setprio 1
+; GCN-NEXT: v_add_f32_e32 v{{[0-9]+}}, 2.0
+define amdgpu_cs void @test_mask0_blocks_valu(ptr addrspace(1) %out, float %x, float %y) {
+ %add1 = fadd float %x, 1.0
+ call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 0)
+ %add2 = fadd float %y, 2.0
+ %sum = fadd float %add1, %add2
+ store float %sum, ptr addrspace(1) %out
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_mask0_blocks_trans:
+; GCN: v_rcp_f32_e32 v{{[0-9]+}}, v{{[0-9]+}}
+; GCN-NEXT: s_setprio 1
+; GCN-NEXT: v_rcp_f32_e32 v{{[0-9]+}}, v{{[0-9]+}}
+define amdgpu_cs void @test_mask0_blocks_trans(ptr addrspace(1) %out, float %x, float %y) {
+ %rcp1 = call float @llvm.amdgcn.rcp.f32(float %x)
+ call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 0)
+ %rcp2 = call float @llvm.amdgcn.rcp.f32(float %y)
+ %sum = fadd float %rcp1, %rcp2
+ store float %sum, ptr addrspace(1) %out
ret void
}
-; GCN-LABEL: {{^}}test_setprio_mask0_blocks_salu:
+; GCN-LABEL: {{^}}test_mask4_allows_salu:
+; GCN: s_add_i32
+; GCN-NEXT: s_add_i32
+; GCN: s_setprio 1
+define amdgpu_cs void @test_mask4_allows_salu(ptr addrspace(1) %out, i32 inreg %x, i32 inreg %y) {
+ %add1 = add i32 %x, 1
+ call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 4)
+ %add2 = add i32 %y, 2
+ %sum = add i32 %add1, %add2
+ store i32 %sum, ptr addrspace(1) %out
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_mask4_blocks_valu:
+; GCN: v_add_f32_e32 v{{[0-9]+}}, 1.0
+; GCN-NEXT: s_setprio 1
+; GCN-NEXT: v_add_f32_e32 v{{[0-9]+}}, 2.0
+define amdgpu_cs void @test_mask4_blocks_valu(ptr addrspace(1) %out, float %x, float %y) {
+ %add1 = fadd float %x, 1.0
+ call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 4)
+ %add2 = fadd float %y, 2.0
+ %sum = fadd float %add1, %add2
+ store float %sum, ptr addrspace(1) %out
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_mask2_allows_valu:
+; GCN: v_add_f32_e32
+; GCN-NEXT: v_add_f32_e32
+; GCN-NEXT: v_add_f32_e32
+; GCN: s_setprio 1
+define amdgpu_cs void @test_mask2_allows_valu(ptr addrspace(1) %out, float %x, float %y) {
+ %add1 = fadd float %x, 1.0
+ call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 2)
+ %add2 = fadd float %y, 2.0
+ %sum = fadd float %add1, %add2
+ store float %sum, ptr addrspace(1) %out
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_mask2_blocks_salu:
; GCN: s_setprio 1
; GCN-NEXT: s_add_i32
-define amdgpu_cs void @test_setprio_mask0_blocks_salu(ptr addrspace(1) %out, i32 inreg %x, i32 inreg %y) {
+define amdgpu_cs void @test_mask2_blocks_salu(ptr addrspace(1) %out, i32 inreg %x, i32 inreg %y) {
%add1 = add i32 %x, 1
- call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 0)
+ call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 2)
%add2 = add i32 %y, 2
%sum = add i32 %add1, %add2
store i32 %sum, ptr addrspace(1) %out
ret void
}
-; GCN-LABEL: {{^}}test_setprio_mask4_allows_salu:
+
+; GCN-LABEL: {{^}}test_mask1_allows_salu:
; GCN: s_add_i32
; GCN-NEXT: s_add_i32
-; GCN-NEXT: s_setprio 1
-define amdgpu_cs void @test_setprio_mask4_allows_salu(ptr addrspace(1) %out, i32 inreg %x, i32 inreg %y) {
+; GCN: s_setprio 1
+define amdgpu_cs void @test_mask1_allows_salu(ptr addrspace(1) %out, i32 inreg %x, i32 inreg %y) {
%add1 = add i32 %x, 1
- call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 4)
+ call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 1)
%add2 = add i32 %y, 2
%sum = add i32 %add1, %add2
store i32 %sum, ptr addrspace(1) %out
ret void
}
+
+; GCN-LABEL: {{^}}test_mask1_allows_valu:
+; GCN: v_add_f32_e32
+; GCN-NEXT: v_add_f32_e32
+; GCN-NEXT: v_add_f32_e32
+; GCN: s_setprio 1
+define amdgpu_cs void @test_mask1_allows_valu(ptr addrspace(1) %out, float %x, float %y) {
+ %add1 = fadd float %x, 1.0
+ call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 1)
+ %add2 = fadd float %y, 2.0
+ %sum = fadd float %add1, %add2
+ store float %sum, ptr addrspace(1) %out
+ ret void
+}
+
+
+; GCN-LABEL: {{^}}test_mask1024_allows_trans:
+; GCN: v_rcp_f32_e32
+; GCN-NEXT: v_rcp_f32_e32
+; GCN: s_setprio 1
+define amdgpu_cs void @test_mask1024_allows_trans(ptr addrspace(1) %out, float %x, float %y) {
+ %rcp1 = call float @llvm.amdgcn.rcp.f32(float %x)
+ call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 1024)
+ %rcp2 = call float @llvm.amdgcn.rcp.f32(float %y)
+ %sum = fadd float %rcp1, %rcp2
+ store float %sum, ptr addrspace(1) %out
+ ret void
+}
+
+; GCN-LABEL: {{^}}test_mask1024_blocks_salu:
+; GCN: s_setprio 1
+; GCN-NEXT: s_add_i32
+define amdgpu_cs void @test_mask1024_blocks_salu(ptr addrspace(1) %out, i32 inreg %x, i32 inreg %y) {
+ %add1 = add i32 %x, 1
+ call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 1024)
+ %add2 = add i32 %y, 2
+ %sum = add i32 %add1, %add2
+ store i32 %sum, ptr addrspace(1) %out
+ ret void
+}
+
+
+; Test SALU+VALU (0x0006 = 6) allows both
+; GCN-LABEL: {{^}}test_mask6_allows_salu_and_valu:
+; GCN: s_add_i32
+; GCN-NEXT: s_add_i32
+; GCN: s_setprio 1
+define amdgpu_cs void @test_mask6_allows_salu_and_valu(ptr addrspace(1) %out, i32 inreg %x, i32 inreg %y) {
+ %add1 = add i32 %x, 1
+ call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 6)
+ %add2 = add i32 %y, 2
+ %sum = add i32 %add1, %add2
+ store i32 %sum, ptr addrspace(1) %out
+ ret void
+}
+
+
+; GCN-LABEL: {{^}}test_mask8_allows_mfma:
+; GCN: s_setprio 1
+; GCN: v_mfma_f32_4x4x1f32
+; GCN: v_mfma_f32_4x4x1f32
+define amdgpu_cs void @test_mask8_allows_mfma(ptr addrspace(1) %out, <4 x float> %in) {
+ %mfma1 = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float 1.0, float 2.0, <4 x float> %in, i32 0, i32 0, i32 0)
+ call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 8)
+ %mfma2 = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float 3.0, float 4.0, <4 x float> %mfma1, i32 0, i32 0, i32 0)
+ store <4 x float> %mfma2, ptr addrspace(1) %out
+ ret void
+}
+
+
+; GCN-LABEL: {{^}}test_mask0_blocks_mfma:
+; GCN: v_mfma_f32_4x4x1f32
+; GCN: s_setprio 1
+; GCN: v_mfma_f32_4x4x1f32
+define amdgpu_cs void @test_mask0_blocks_mfma(ptr addrspace(1) %out, <4 x float> %in) {
+ %mfma1 = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float 1.0, float 2.0, <4 x float> %in, i32 0, i32 0, i32 0)
+ call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 0)
+ %mfma2 = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float 3.0, float 4.0, <4 x float> %mfma1, i32 0, i32 0, i32 0)
+ store <4 x float> %mfma2, ptr addrspace(1) %out
+ ret void
+}
+
+
+; GCN-LABEL: {{^}}test_mask32_vmem_read:
+; GCN: s_setprio 1
+; GCN: global_load_dword
+define amdgpu_cs void @test_mask32_vmem_read(ptr addrspace(1) %in, ptr addrspace(1) %out, i32 inreg %x) {
+ %val1 = add i32 %x, 1
+ call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 32)
+ %load = load i32, ptr addrspace(1) %in
+ %sum = add i32 %val1, %load
+ store i32 %sum, ptr addrspace(1) %out
+ ret void
+}
+
+
+; GCN-LABEL: {{^}}test_mask64_vmem_write:
+; GCN: s_setprio 1
+; GCN: global_store_dword
+define amdgpu_cs void @test_mask64_vmem_write(ptr addrspace(1) %out, i32 inreg %x) {
+ %val = add i32 %x, 1
+ call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 64)
+ store i32 %val, ptr addrspace(1) %out
+ ret void
+}
+
+
+; GCN-LABEL: {{^}}test_mask16_vmem:
+; GCN: s_setprio 1
+; GCN: global_load_dword
+define amdgpu_cs void @test_mask16_vmem(ptr addrspace(1) %in, ptr addrspace(1) %out, i32 inreg %x) {
+ %val1 = add i32 %x, 1
+ call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 16)
+ %load = load i32, ptr addrspace(1) %in
+ %sum = add i32 %val1, %load
+ store i32 %sum, ptr addrspace(1) %out
+ ret void
+}
+
+
+; GCN-LABEL: {{^}}test_mask256_ds_read:
+; GCN: s_setprio 1
+; GCN: ds_read_b32
+define amdgpu_cs void @test_mask256_ds_read(ptr addrspace(3) %in, ptr addrspace(1) %out, i32 inreg %x) {
+ %val1 = add i32 %x, 1
+ call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 256)
+ %load = load i32, ptr addrspace(3) %in
+ %sum = add i32 %val1, %load
+ store i32 %sum, ptr addrspace(1) %out
+ ret void
+}
+
+
+; GCN-LABEL: {{^}}test_mask512_ds_write:
+; GCN: s_setprio 1
+; GCN: ds_write_b32
+define amdgpu_cs void @test_mask512_ds_write(ptr addrspace(3) %out, i32 inreg %x) {
+ %val = add i32 %x, 1
+ call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 512)
+ store i32 %val, ptr addrspace(3) %out
+ ret void
+}
+
+
+; GCN-LABEL: {{^}}test_mask128_ds:
+; GCN: s_setprio 1
+; GCN: ds_read_b32
+define amdgpu_cs void @test_mask128_ds(ptr addrspace(3) %in, ptr addrspace(1) %out, i32 inreg %x) {
+ %val1 = add i32 %x, 1
+ call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 128)
+ %load = load i32, ptr addrspace(3) %in
+ %sum = add i32 %val1, %load
+ store i32 %sum, ptr addrspace(1) %out
+ ret void
+}
>From 96d8e5c353747c9fde54376a9999d780de7115c4 Mon Sep 17 00:00:00 2001
From: Frederik Harwath <fharwath at amd.com>
Date: Mon, 22 Jun 2026 05:47:25 -0400
Subject: [PATCH 03/10] [AMDGPU] Reject incompatible mask bits
---
llvm/docs/AMDGPUUsage.rst | 4 +-
llvm/lib/Target/AMDGPU/SIInstrInfo.cpp | 12 ++
.../CodeGen/AMDGPU/llvm.amdgcn.s.setprio.ll | 12 +-
llvm/test/CodeGen/AMDGPU/sched-setprio.ll | 140 +++++++++---------
4 files changed, 91 insertions(+), 77 deletions(-)
diff --git a/llvm/docs/AMDGPUUsage.rst b/llvm/docs/AMDGPUUsage.rst
index 57e51dc4bc1af..0ca39c1db29da 100644
--- a/llvm/docs/AMDGPUUsage.rst
+++ b/llvm/docs/AMDGPUUsage.rst
@@ -1728,7 +1728,9 @@ The AMDGPU backend implements the following LLVM IR intrinsics.
the intrinsic in the same way as llvm.amdgcn.sched.barrier. Takes two parameters:
- Priority (i16): Hardware wave priority level. 0 = lowest, 3 = highest.
- - Mask (i32): Scheduling barrier mask.
+ - Mask (i32): Scheduling barrier mask. Memory mask bits (VMEM, VMEM_READ,
+ VMEM_WRITE, DS, DS_READ, DS_WRITE, LDSDMA) may not be set since they
+ contradict the S_SETPRIO instruction's inherent scheduling dependencies.
llvm.amdgcn.s.setprio Like llvm.amdgcn.s.setprio.mask with Mask = 0.
diff --git a/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp b/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp
index 0d7eb1fcc61ff..1594f4d679e6a 100644
--- a/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp
@@ -6059,6 +6059,18 @@ bool SIInstrInfo::verifyInstruction(const MachineInstr &MI,
}
}
+ if (Opcode == AMDGPU::S_SETPRIO) {
+ constexpr unsigned BarrierMemMaskBits =
+ (1u << 4) | (1u << 5) | (1u << 6) | // VMEM, VMEM_READ, VMEM_WRITE
+ (1u << 7) | (1u << 8) | (1u << 9) | // DS, DS_READ, DS_WRITE
+ (1u << 11); // LDSDMA
+ unsigned Mask = MI.getOperand(1).getImm();
+ if (Mask & BarrierMemMaskBits) {
+ ErrInfo = "S_SETPRIO mask contains invalid memory operation bits";
+ return false;
+ }
+ }
+
return true;
}
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.setprio.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.setprio.ll
index 0078f33782bf1..996ac40943f03 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.setprio.ll
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.setprio.ll
@@ -47,14 +47,14 @@ define void @test_llvm_amdgcn_s_setprio() #0 {
ret void
}
-define void @test_llvm_amdgcn_s_setprio_mask_full() #0 {
-; GFX9-LABEL: test_llvm_amdgcn_s_setprio_mask_full:
+define void @test_llvm_amdgcn_s_setprio_mask0() #0 {
+; GFX9-LABEL: test_llvm_amdgcn_s_setprio_mask0:
; GFX9: ; %bb.0:
; GFX9-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0) ; encoding: [0x00,0x00,0x8c,0xbf]
; GFX9-NEXT: s_setprio 3 ; encoding: [0x03,0x00,0x8f,0xbf]
; GFX9-NEXT: s_setpc_b64 s[30:31] ; encoding: [0x1e,0x1d,0x80,0xbe]
;
-; SI-LABEL: test_llvm_amdgcn_s_setprio_mask_full:
+; SI-LABEL: test_llvm_amdgcn_s_setprio_mask0:
; SI: ; %bb.0:
; SI-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0) ; encoding: [0x00,0x00,0x8c,0xbf]
; SI-NEXT: s_setprio 3 ; encoding: [0x03,0x00,0x8f,0xbf]
@@ -63,14 +63,14 @@ define void @test_llvm_amdgcn_s_setprio_mask_full() #0 {
ret void
}
-define void @test_llvm_amdgcn_s_setprio_mask_selective() #0 {
-; GFX9-LABEL: test_llvm_amdgcn_s_setprio_mask_selective:
+define void @test_llvm_amdgcn_s_setprio_mask15() #0 {
+; GFX9-LABEL: test_llvm_amdgcn_s_setprio_mask15:
; GFX9: ; %bb.0:
; GFX9-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0) ; encoding: [0x00,0x00,0x8c,0xbf]
; GFX9-NEXT: s_setprio 1 ; encoding: [0x01,0x00,0x8f,0xbf]
; GFX9-NEXT: s_setpc_b64 s[30:31] ; encoding: [0x1e,0x1d,0x80,0xbe]
;
-; SI-LABEL: test_llvm_amdgcn_s_setprio_mask_selective:
+; SI-LABEL: test_llvm_amdgcn_s_setprio_mask15:
; SI: ; %bb.0:
; SI-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0) ; encoding: [0x00,0x00,0x8c,0xbf]
; SI-NEXT: s_setprio 1 ; encoding: [0x01,0x00,0x8f,0xbf]
diff --git a/llvm/test/CodeGen/AMDGPU/sched-setprio.ll b/llvm/test/CodeGen/AMDGPU/sched-setprio.ll
index f8dceed6cc950..d2b74c5d33bcb 100644
--- a/llvm/test/CodeGen/AMDGPU/sched-setprio.ll
+++ b/llvm/test/CodeGen/AMDGPU/sched-setprio.ll
@@ -1,6 +1,21 @@
-; RUN: llc -mtriple=amdgcn -mcpu=gfx908 < %s | FileCheck --check-prefix=GCN %s
+; RUN: split-file %s %t
+; RUN: llc -mtriple=amdgcn -mcpu=gfx908 < %t/valid.ll | FileCheck --check-prefix=GCN %s
+; RUN: not --crash llc -mtriple=amdgcn -mcpu=gfx908 -verify-machineinstrs < %t/vmem.ll 2>&1 | FileCheck --check-prefix=CHECK-ERR %s
+; RUN: not --crash llc -mtriple=amdgcn -mcpu=gfx908 -verify-machineinstrs < %t/vmem-read.ll 2>&1 | FileCheck --check-prefix=CHECK-ERR %s
+; RUN: not --crash llc -mtriple=amdgcn -mcpu=gfx908 -verify-machineinstrs < %t/vmem-write.ll 2>&1 | FileCheck --check-prefix=CHECK-ERR %s
+; RUN: not --crash llc -mtriple=amdgcn -mcpu=gfx908 -verify-machineinstrs < %t/ds.ll 2>&1 | FileCheck --check-prefix=CHECK-ERR %s
+; RUN: not --crash llc -mtriple=amdgcn -mcpu=gfx908 -verify-machineinstrs < %t/ds-read.ll 2>&1 | FileCheck --check-prefix=CHECK-ERR %s
+; RUN: not --crash llc -mtriple=amdgcn -mcpu=gfx908 -verify-machineinstrs < %t/ds-write.ll 2>&1 | FileCheck --check-prefix=CHECK-ERR %s
+; RUN: not --crash llc -mtriple=amdgcn -mcpu=gfx908 -verify-machineinstrs < %t/ldsdma.ll 2>&1 | FileCheck --check-prefix=CHECK-ERR %s
-; Test scheduling mask behavior for llvm.amdgcn.s.setprio{,.mask} intrinsics.
+; CHECK-ERR: S_SETPRIO mask contains invalid memory operation bits
+
+; Demonstrate that the maks bits that are compatible with s_set_prio
+; allow corresponding instruction types to be scheduled across the s_setprio
+; instruction and that incompatible mask bits are rejected by the
+; verifier.
+
+;--- valid.ll
declare void @llvm.amdgcn.s.setprio(i16)
declare void @llvm.amdgcn.s.setprio.mask(i16, i32)
@@ -8,12 +23,14 @@ declare float @llvm.amdgcn.rcp.f32(float)
declare <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float, float, <4 x float>, i32, i32, i32)
; GCN-LABEL: {{^}}test_mask0_blocks_salu:
-; GCN: s_setprio 1
+; GCN: s_add_i32
+; GCN-NEXT: s_setprio 1
+; GCN-NEXT: s_add_i32
; GCN-NEXT: s_add_i32
define amdgpu_cs void @test_mask0_blocks_salu(ptr addrspace(1) %out, i32 inreg %x, i32 inreg %y) {
- %add1 = add i32 %x, 1
+ %add1 = add i32 %x, %y
call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 0)
- %add2 = add i32 %y, 2
+ %add2 = add i32 %y, %y
%sum = add i32 %add1, %add2
store i32 %sum, ptr addrspace(1) %out
ret void
@@ -48,11 +65,12 @@ define amdgpu_cs void @test_mask0_blocks_trans(ptr addrspace(1) %out, float %x,
; GCN-LABEL: {{^}}test_mask4_allows_salu:
; GCN: s_add_i32
; GCN-NEXT: s_add_i32
-; GCN: s_setprio 1
+; GCN-NEXT: s_add_i32
+; GCN-NEXT: s_setprio 1
define amdgpu_cs void @test_mask4_allows_salu(ptr addrspace(1) %out, i32 inreg %x, i32 inreg %y) {
- %add1 = add i32 %x, 1
+ %add1 = add i32 %x, %y
call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 4)
- %add2 = add i32 %y, 2
+ %add2 = add i32 %y, %y
%sum = add i32 %add1, %add2
store i32 %sum, ptr addrspace(1) %out
ret void
@@ -86,26 +104,28 @@ define amdgpu_cs void @test_mask2_allows_valu(ptr addrspace(1) %out, float %x, f
}
; GCN-LABEL: {{^}}test_mask2_blocks_salu:
-; GCN: s_setprio 1
+; GCN: s_add_i32
+; GCN-NEXT: s_setprio 1
+; GCN-NEXT: s_add_i32
; GCN-NEXT: s_add_i32
define amdgpu_cs void @test_mask2_blocks_salu(ptr addrspace(1) %out, i32 inreg %x, i32 inreg %y) {
- %add1 = add i32 %x, 1
+ %add1 = add i32 %x, %y
call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 2)
- %add2 = add i32 %y, 2
+ %add2 = add i32 %y, %y
%sum = add i32 %add1, %add2
store i32 %sum, ptr addrspace(1) %out
ret void
}
-
; GCN-LABEL: {{^}}test_mask1_allows_salu:
; GCN: s_add_i32
; GCN-NEXT: s_add_i32
+; GCN-NEXT: s_add_i32
; GCN: s_setprio 1
define amdgpu_cs void @test_mask1_allows_salu(ptr addrspace(1) %out, i32 inreg %x, i32 inreg %y) {
- %add1 = add i32 %x, 1
+ %add1 = add i32 %x, %y
call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 1)
- %add2 = add i32 %y, 2
+ %add2 = add i32 %y, %y
%sum = add i32 %add1, %add2
store i32 %sum, ptr addrspace(1) %out
ret void
@@ -125,7 +145,6 @@ define amdgpu_cs void @test_mask1_allows_valu(ptr addrspace(1) %out, float %x, f
ret void
}
-
; GCN-LABEL: {{^}}test_mask1024_allows_trans:
; GCN: v_rcp_f32_e32
; GCN-NEXT: v_rcp_f32_e32
@@ -140,33 +159,33 @@ define amdgpu_cs void @test_mask1024_allows_trans(ptr addrspace(1) %out, float %
}
; GCN-LABEL: {{^}}test_mask1024_blocks_salu:
-; GCN: s_setprio 1
+; GCN: s_add_i32
+; GCN-NEXT: s_setprio 1
+; GCN-NEXT: s_add_i32
; GCN-NEXT: s_add_i32
define amdgpu_cs void @test_mask1024_blocks_salu(ptr addrspace(1) %out, i32 inreg %x, i32 inreg %y) {
- %add1 = add i32 %x, 1
+ %add1 = add i32 %x, %y
call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 1024)
- %add2 = add i32 %y, 2
+ %add2 = add i32 %y, %y
%sum = add i32 %add1, %add2
store i32 %sum, ptr addrspace(1) %out
ret void
}
-
-; Test SALU+VALU (0x0006 = 6) allows both
; GCN-LABEL: {{^}}test_mask6_allows_salu_and_valu:
; GCN: s_add_i32
; GCN-NEXT: s_add_i32
+; GCN-NEXT: s_add_i32
; GCN: s_setprio 1
define amdgpu_cs void @test_mask6_allows_salu_and_valu(ptr addrspace(1) %out, i32 inreg %x, i32 inreg %y) {
- %add1 = add i32 %x, 1
+ %add1 = add i32 %x, %y
call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 6)
- %add2 = add i32 %y, 2
+ %add2 = add i32 %y, %y
%sum = add i32 %add1, %add2
store i32 %sum, ptr addrspace(1) %out
ret void
}
-
; GCN-LABEL: {{^}}test_mask8_allows_mfma:
; GCN: s_setprio 1
; GCN: v_mfma_f32_4x4x1f32
@@ -179,7 +198,6 @@ define amdgpu_cs void @test_mask8_allows_mfma(ptr addrspace(1) %out, <4 x float>
ret void
}
-
; GCN-LABEL: {{^}}test_mask0_blocks_mfma:
; GCN: v_mfma_f32_4x4x1f32
; GCN: s_setprio 1
@@ -192,76 +210,58 @@ define amdgpu_cs void @test_mask0_blocks_mfma(ptr addrspace(1) %out, <4 x float>
ret void
}
+;--- vmem.ll
+declare void @llvm.amdgcn.s.setprio.mask(i16, i32)
-; GCN-LABEL: {{^}}test_mask32_vmem_read:
-; GCN: s_setprio 1
-; GCN: global_load_dword
-define amdgpu_cs void @test_mask32_vmem_read(ptr addrspace(1) %in, ptr addrspace(1) %out, i32 inreg %x) {
- %val1 = add i32 %x, 1
+define amdgpu_cs void @test_invalid_vmem() {
+ call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 16)
+ ret void
+}
+
+;--- vmem-read.ll
+declare void @llvm.amdgcn.s.setprio.mask(i16, i32)
+
+define amdgpu_cs void @test_invalid_vmem_read() {
call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 32)
- %load = load i32, ptr addrspace(1) %in
- %sum = add i32 %val1, %load
- store i32 %sum, ptr addrspace(1) %out
ret void
}
+;--- vmem-write.ll
+declare void @llvm.amdgcn.s.setprio.mask(i16, i32)
-; GCN-LABEL: {{^}}test_mask64_vmem_write:
-; GCN: s_setprio 1
-; GCN: global_store_dword
-define amdgpu_cs void @test_mask64_vmem_write(ptr addrspace(1) %out, i32 inreg %x) {
- %val = add i32 %x, 1
+define amdgpu_cs void @test_invalid_vmem_write() {
call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 64)
- store i32 %val, ptr addrspace(1) %out
ret void
}
+;--- ds.ll
+declare void @llvm.amdgcn.s.setprio.mask(i16, i32)
-; GCN-LABEL: {{^}}test_mask16_vmem:
-; GCN: s_setprio 1
-; GCN: global_load_dword
-define amdgpu_cs void @test_mask16_vmem(ptr addrspace(1) %in, ptr addrspace(1) %out, i32 inreg %x) {
- %val1 = add i32 %x, 1
- call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 16)
- %load = load i32, ptr addrspace(1) %in
- %sum = add i32 %val1, %load
- store i32 %sum, ptr addrspace(1) %out
+define amdgpu_cs void @test_invalid_ds() {
+ call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 128)
ret void
}
+;--- ds-read.ll
+declare void @llvm.amdgcn.s.setprio.mask(i16, i32)
-; GCN-LABEL: {{^}}test_mask256_ds_read:
-; GCN: s_setprio 1
-; GCN: ds_read_b32
-define amdgpu_cs void @test_mask256_ds_read(ptr addrspace(3) %in, ptr addrspace(1) %out, i32 inreg %x) {
- %val1 = add i32 %x, 1
+define amdgpu_cs void @test_invalid_ds_read() {
call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 256)
- %load = load i32, ptr addrspace(3) %in
- %sum = add i32 %val1, %load
- store i32 %sum, ptr addrspace(1) %out
ret void
}
+;--- ds-write.ll
+declare void @llvm.amdgcn.s.setprio.mask(i16, i32)
-; GCN-LABEL: {{^}}test_mask512_ds_write:
-; GCN: s_setprio 1
-; GCN: ds_write_b32
-define amdgpu_cs void @test_mask512_ds_write(ptr addrspace(3) %out, i32 inreg %x) {
- %val = add i32 %x, 1
+define amdgpu_cs void @test_invalid_ds_write() {
call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 512)
- store i32 %val, ptr addrspace(3) %out
ret void
}
+;--- ldsdma.ll
+declare void @llvm.amdgcn.s.setprio.mask(i16, i32)
-; GCN-LABEL: {{^}}test_mask128_ds:
-; GCN: s_setprio 1
-; GCN: ds_read_b32
-define amdgpu_cs void @test_mask128_ds(ptr addrspace(3) %in, ptr addrspace(1) %out, i32 inreg %x) {
- %val1 = add i32 %x, 1
- call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 128)
- %load = load i32, ptr addrspace(3) %in
- %sum = add i32 %val1, %load
- store i32 %sum, ptr addrspace(1) %out
+define amdgpu_cs void @test_invalid_ldsdma() {
+ call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 2048)
ret void
}
>From 4f9512b574dae58eefad1ecff9fe98da4adc9ddd Mon Sep 17 00:00:00 2001
From: Frederik Harwath <fharwath at amd.com>
Date: Tue, 28 Jul 2026 12:16:35 -0400
Subject: [PATCH 04/10] Review changes
Address changes from copilot review
---
.../AMDGPU/required-export-priority.mir | 38 +++++++++----------
llvm/test/CodeGen/AMDGPU/sched-setprio.ll | 2 +-
2 files changed, 20 insertions(+), 20 deletions(-)
diff --git a/llvm/test/CodeGen/AMDGPU/required-export-priority.mir b/llvm/test/CodeGen/AMDGPU/required-export-priority.mir
index 9039ee14715ca..ed75c30b3311e 100644
--- a/llvm/test/CodeGen/AMDGPU/required-export-priority.mir
+++ b/llvm/test/CodeGen/AMDGPU/required-export-priority.mir
@@ -48,9 +48,9 @@ body: |
; GFX1150-LABEL: name: end_of_shader
; GFX1150: liveins: $vgpr0
; GFX1150-NEXT: {{ $}}
- ; GFX1150-NEXT: S_SETPRIO 2
+ ; GFX1150-NEXT: S_SETPRIO 2, 0
; GFX1150-NEXT: EXP_DONE 0, $vgpr0, $vgpr0, $vgpr0, $vgpr0, -1, -1, 15, implicit $exec
- ; GFX1150-NEXT: S_SETPRIO 0
+ ; GFX1150-NEXT: S_SETPRIO 0, 0
; GFX1150-NEXT: S_NOP 0
; GFX1150-NEXT: S_NOP 0
; GFX1150-NEXT: S_ENDPGM 0
@@ -75,13 +75,13 @@ body: |
; GFX1150-LABEL: name: end_of_shader_return_to_epilogue
; GFX1150: liveins: $vgpr0
; GFX1150-NEXT: {{ $}}
- ; GFX1150-NEXT: S_SETPRIO 2
+ ; GFX1150-NEXT: S_SETPRIO 2, 0
; GFX1150-NEXT: EXP_DONE 0, $vgpr0, $vgpr0, $vgpr0, $vgpr0, -1, -1, 15, implicit $exec
- ; GFX1150-NEXT: S_SETPRIO 0
+ ; GFX1150-NEXT: S_SETPRIO 0, 0
; GFX1150-NEXT: S_WAITCNT_EXPCNT $sgpr_null, 0
; GFX1150-NEXT: S_NOP 0
; GFX1150-NEXT: S_NOP 0
- ; GFX1150-NEXT: S_SETPRIO 2
+ ; GFX1150-NEXT: S_SETPRIO 2, 0
; GFX1150-NEXT: SI_RETURN_TO_EPILOG $vgpr0
EXP_DONE 0, $vgpr0, $vgpr0, $vgpr0, $vgpr0, -1, -1, 15, implicit $exec
SI_RETURN_TO_EPILOG $vgpr0
@@ -196,11 +196,11 @@ body: |
; GFX1150-LABEL: name: block_of_exports
; GFX1150: liveins: $vgpr0
; GFX1150-NEXT: {{ $}}
- ; GFX1150-NEXT: S_SETPRIO 2
+ ; GFX1150-NEXT: S_SETPRIO 2, 0
; GFX1150-NEXT: EXP 2, $vgpr0, $vgpr0, $vgpr0, $vgpr0, -1, -1, 15, implicit $exec
; GFX1150-NEXT: EXP 1, $vgpr0, $vgpr0, $vgpr0, $vgpr0, -1, -1, 15, implicit $exec
; GFX1150-NEXT: EXP_DONE 0, $vgpr0, $vgpr0, $vgpr0, $vgpr0, -1, -1, 15, implicit $exec
- ; GFX1150-NEXT: S_SETPRIO 0
+ ; GFX1150-NEXT: S_SETPRIO 0, 0
; GFX1150-NEXT: S_NOP 0
; GFX1150-NEXT: S_NOP 0
; GFX1150-NEXT: S_ENDPGM 0
@@ -231,23 +231,23 @@ body: |
; GFX1150-LABEL: name: sparse_exports
; GFX1150: liveins: $vgpr0
; GFX1150-NEXT: {{ $}}
- ; GFX1150-NEXT: S_SETPRIO 2
+ ; GFX1150-NEXT: S_SETPRIO 2, 0
; GFX1150-NEXT: EXP 2, $vgpr0, $vgpr0, $vgpr0, $vgpr0, -1, -1, 15, implicit $exec
- ; GFX1150-NEXT: S_SETPRIO 0
+ ; GFX1150-NEXT: S_SETPRIO 0, 0
; GFX1150-NEXT: S_WAITCNT_EXPCNT $sgpr_null, 0
; GFX1150-NEXT: S_NOP 0
; GFX1150-NEXT: S_NOP 0
- ; GFX1150-NEXT: S_SETPRIO 2
+ ; GFX1150-NEXT: S_SETPRIO 2, 0
; GFX1150-NEXT: $vgpr0 = V_AND_B32_e32 1, $vgpr0, implicit $exec
; GFX1150-NEXT: EXP 1, $vgpr0, $vgpr0, $vgpr0, $vgpr0, -1, -1, 15, implicit $exec
- ; GFX1150-NEXT: S_SETPRIO 0
+ ; GFX1150-NEXT: S_SETPRIO 0, 0
; GFX1150-NEXT: S_WAITCNT_EXPCNT $sgpr_null, 0
; GFX1150-NEXT: S_NOP 0
; GFX1150-NEXT: S_NOP 0
- ; GFX1150-NEXT: S_SETPRIO 2
+ ; GFX1150-NEXT: S_SETPRIO 2, 0
; GFX1150-NEXT: $vgpr0 = V_OR_B32_e32 2, $vgpr0, implicit $exec
; GFX1150-NEXT: EXP_DONE 0, $vgpr0, $vgpr0, $vgpr0, $vgpr0, -1, -1, 15, implicit $exec
- ; GFX1150-NEXT: S_SETPRIO 0
+ ; GFX1150-NEXT: S_SETPRIO 0, 0
; GFX1150-NEXT: S_NOP 0
; GFX1150-NEXT: S_NOP 0
; GFX1150-NEXT: S_ENDPGM 0
@@ -362,24 +362,24 @@ body: |
; GFX11-LABEL: name: existing_setprio_2
; GFX11: liveins: $vgpr0
; GFX11-NEXT: {{ $}}
- ; GFX11-NEXT: S_SETPRIO 3
+ ; GFX11-NEXT: S_SETPRIO 3, 0
; GFX11-NEXT: EXP 1, $vgpr0, $vgpr0, $vgpr0, $vgpr0, -1, -1, 15, implicit $exec
; GFX11-NEXT: EXP_DONE 0, $vgpr0, $vgpr0, $vgpr0, $vgpr0, -1, -1, 15, implicit $exec
- ; GFX11-NEXT: S_SETPRIO 3
+ ; GFX11-NEXT: S_SETPRIO 3, 0
; GFX11-NEXT: S_ENDPGM 0
;
; GFX1150-LABEL: name: existing_setprio_2
; GFX1150: liveins: $vgpr0
; GFX1150-NEXT: {{ $}}
- ; GFX1150-NEXT: S_SETPRIO 3
+ ; GFX1150-NEXT: S_SETPRIO 3, 0
; GFX1150-NEXT: EXP 1, $vgpr0, $vgpr0, $vgpr0, $vgpr0, -1, -1, 15, implicit $exec
; GFX1150-NEXT: EXP_DONE 0, $vgpr0, $vgpr0, $vgpr0, $vgpr0, -1, -1, 15, implicit $exec
- ; GFX1150-NEXT: S_SETPRIO 0
+ ; GFX1150-NEXT: S_SETPRIO 0, 0
; GFX1150-NEXT: S_WAITCNT_EXPCNT $sgpr_null, 0
; GFX1150-NEXT: S_NOP 0
; GFX1150-NEXT: S_NOP 0
- ; GFX1150-NEXT: S_SETPRIO 2
- ; GFX1150-NEXT: S_SETPRIO 3
+ ; GFX1150-NEXT: S_SETPRIO 2, 0
+ ; GFX1150-NEXT: S_SETPRIO 3, 0
; GFX1150-NEXT: S_ENDPGM 0
S_SETPRIO 3, 0
EXP 1, $vgpr0, $vgpr0, $vgpr0, $vgpr0, -1, -1, 15, implicit $exec
diff --git a/llvm/test/CodeGen/AMDGPU/sched-setprio.ll b/llvm/test/CodeGen/AMDGPU/sched-setprio.ll
index d2b74c5d33bcb..5439a38ce497a 100644
--- a/llvm/test/CodeGen/AMDGPU/sched-setprio.ll
+++ b/llvm/test/CodeGen/AMDGPU/sched-setprio.ll
@@ -10,7 +10,7 @@
; CHECK-ERR: S_SETPRIO mask contains invalid memory operation bits
-; Demonstrate that the maks bits that are compatible with s_set_prio
+; Demonstrate that the mask bits that are compatible with s_set_prio
; allow corresponding instruction types to be scheduled across the s_setprio
; instruction and that incompatible mask bits are rejected by the
; verifier.
>From 44ccb566c2e0801936da0e5843c0aa4acdcc67c4 Mon Sep 17 00:00:00 2001
From: Frederik Harwath <fharwath at amd.com>
Date: Wed, 26 Aug 2026 03:15:02 -0400
Subject: [PATCH 05/10] Revert "Review changes"
This reverts commit 4f9512b574dae58eefad1ecff9fe98da4adc9ddd.
---
.../AMDGPU/required-export-priority.mir | 38 +++++++++----------
llvm/test/CodeGen/AMDGPU/sched-setprio.ll | 2 +-
2 files changed, 20 insertions(+), 20 deletions(-)
diff --git a/llvm/test/CodeGen/AMDGPU/required-export-priority.mir b/llvm/test/CodeGen/AMDGPU/required-export-priority.mir
index 34924e1e62381..fc65dad5cdd5e 100644
--- a/llvm/test/CodeGen/AMDGPU/required-export-priority.mir
+++ b/llvm/test/CodeGen/AMDGPU/required-export-priority.mir
@@ -48,9 +48,9 @@ body: |
; GFX1150-LABEL: name: end_of_shader
; GFX1150: liveins: $vgpr0
; GFX1150-NEXT: {{ $}}
- ; GFX1150-NEXT: S_SETPRIO 2, 0
+ ; GFX1150-NEXT: S_SETPRIO 2
; GFX1150-NEXT: EXP_DONE 0, $vgpr0, $vgpr0, $vgpr0, $vgpr0, -1, -1, 15, implicit $exec
- ; GFX1150-NEXT: S_SETPRIO 0, 0
+ ; GFX1150-NEXT: S_SETPRIO 0
; GFX1150-NEXT: S_NOP 0
; GFX1150-NEXT: S_NOP 0
; GFX1150-NEXT: S_ENDPGM 0
@@ -75,13 +75,13 @@ body: |
; GFX1150-LABEL: name: end_of_shader_return_to_epilogue
; GFX1150: liveins: $vgpr0
; GFX1150-NEXT: {{ $}}
- ; GFX1150-NEXT: S_SETPRIO 2, 0
+ ; GFX1150-NEXT: S_SETPRIO 2
; GFX1150-NEXT: EXP_DONE 0, $vgpr0, $vgpr0, $vgpr0, $vgpr0, -1, -1, 15, implicit $exec
- ; GFX1150-NEXT: S_SETPRIO 0, 0
+ ; GFX1150-NEXT: S_SETPRIO 0
; GFX1150-NEXT: S_WAITCNT_EXPCNT $sgpr_null, 0
; GFX1150-NEXT: S_NOP 0
; GFX1150-NEXT: S_NOP 0
- ; GFX1150-NEXT: S_SETPRIO 2, 0
+ ; GFX1150-NEXT: S_SETPRIO 2
; GFX1150-NEXT: SI_RETURN_TO_EPILOG $vgpr0
EXP_DONE 0, $vgpr0, $vgpr0, $vgpr0, $vgpr0, -1, -1, 15, implicit $exec
SI_RETURN_TO_EPILOG $vgpr0
@@ -196,11 +196,11 @@ body: |
; GFX1150-LABEL: name: block_of_exports
; GFX1150: liveins: $vgpr0
; GFX1150-NEXT: {{ $}}
- ; GFX1150-NEXT: S_SETPRIO 2, 0
+ ; GFX1150-NEXT: S_SETPRIO 2
; GFX1150-NEXT: EXP 2, $vgpr0, $vgpr0, $vgpr0, $vgpr0, -1, -1, 15, implicit $exec
; GFX1150-NEXT: EXP 1, $vgpr0, $vgpr0, $vgpr0, $vgpr0, -1, -1, 15, implicit $exec
; GFX1150-NEXT: EXP_DONE 0, $vgpr0, $vgpr0, $vgpr0, $vgpr0, -1, -1, 15, implicit $exec
- ; GFX1150-NEXT: S_SETPRIO 0, 0
+ ; GFX1150-NEXT: S_SETPRIO 0
; GFX1150-NEXT: S_NOP 0
; GFX1150-NEXT: S_NOP 0
; GFX1150-NEXT: S_ENDPGM 0
@@ -231,23 +231,23 @@ body: |
; GFX1150-LABEL: name: sparse_exports
; GFX1150: liveins: $vgpr0
; GFX1150-NEXT: {{ $}}
- ; GFX1150-NEXT: S_SETPRIO 2, 0
+ ; GFX1150-NEXT: S_SETPRIO 2
; GFX1150-NEXT: EXP 2, $vgpr0, $vgpr0, $vgpr0, $vgpr0, -1, -1, 15, implicit $exec
- ; GFX1150-NEXT: S_SETPRIO 0, 0
+ ; GFX1150-NEXT: S_SETPRIO 0
; GFX1150-NEXT: S_WAITCNT_EXPCNT $sgpr_null, 0
; GFX1150-NEXT: S_NOP 0
; GFX1150-NEXT: S_NOP 0
- ; GFX1150-NEXT: S_SETPRIO 2, 0
+ ; GFX1150-NEXT: S_SETPRIO 2
; GFX1150-NEXT: $vgpr0 = V_AND_B32_e32 1, $vgpr0, implicit $exec
; GFX1150-NEXT: EXP 1, $vgpr0, $vgpr0, $vgpr0, $vgpr0, -1, -1, 15, implicit $exec
- ; GFX1150-NEXT: S_SETPRIO 0, 0
+ ; GFX1150-NEXT: S_SETPRIO 0
; GFX1150-NEXT: S_WAITCNT_EXPCNT $sgpr_null, 0
; GFX1150-NEXT: S_NOP 0
; GFX1150-NEXT: S_NOP 0
- ; GFX1150-NEXT: S_SETPRIO 2, 0
+ ; GFX1150-NEXT: S_SETPRIO 2
; GFX1150-NEXT: $vgpr0 = V_OR_B32_e32 2, $vgpr0, implicit $exec
; GFX1150-NEXT: EXP_DONE 0, $vgpr0, $vgpr0, $vgpr0, $vgpr0, -1, -1, 15, implicit $exec
- ; GFX1150-NEXT: S_SETPRIO 0, 0
+ ; GFX1150-NEXT: S_SETPRIO 0
; GFX1150-NEXT: S_NOP 0
; GFX1150-NEXT: S_NOP 0
; GFX1150-NEXT: S_ENDPGM 0
@@ -362,24 +362,24 @@ body: |
; GFX11-LABEL: name: existing_setprio_2
; GFX11: liveins: $vgpr0
; GFX11-NEXT: {{ $}}
- ; GFX11-NEXT: S_SETPRIO 3, 0
+ ; GFX11-NEXT: S_SETPRIO 3
; GFX11-NEXT: EXP 1, $vgpr0, $vgpr0, $vgpr0, $vgpr0, -1, -1, 15, implicit $exec
; GFX11-NEXT: EXP_DONE 0, $vgpr0, $vgpr0, $vgpr0, $vgpr0, -1, -1, 15, implicit $exec
- ; GFX11-NEXT: S_SETPRIO 3, 0
+ ; GFX11-NEXT: S_SETPRIO 3
; GFX11-NEXT: S_ENDPGM 0
;
; GFX1150-LABEL: name: existing_setprio_2
; GFX1150: liveins: $vgpr0
; GFX1150-NEXT: {{ $}}
- ; GFX1150-NEXT: S_SETPRIO 3, 0
+ ; GFX1150-NEXT: S_SETPRIO 3
; GFX1150-NEXT: EXP 1, $vgpr0, $vgpr0, $vgpr0, $vgpr0, -1, -1, 15, implicit $exec
; GFX1150-NEXT: EXP_DONE 0, $vgpr0, $vgpr0, $vgpr0, $vgpr0, -1, -1, 15, implicit $exec
- ; GFX1150-NEXT: S_SETPRIO 0, 0
+ ; GFX1150-NEXT: S_SETPRIO 0
; GFX1150-NEXT: S_WAITCNT_EXPCNT $sgpr_null, 0
; GFX1150-NEXT: S_NOP 0
; GFX1150-NEXT: S_NOP 0
- ; GFX1150-NEXT: S_SETPRIO 2, 0
- ; GFX1150-NEXT: S_SETPRIO 3, 0
+ ; GFX1150-NEXT: S_SETPRIO 2
+ ; GFX1150-NEXT: S_SETPRIO 3
; GFX1150-NEXT: S_ENDPGM 0
S_SETPRIO 3, 0
EXP 1, $vgpr0, $vgpr0, $vgpr0, $vgpr0, -1, -1, 15, implicit $exec
diff --git a/llvm/test/CodeGen/AMDGPU/sched-setprio.ll b/llvm/test/CodeGen/AMDGPU/sched-setprio.ll
index c0db14447e11b..da20dfb37dd36 100644
--- a/llvm/test/CodeGen/AMDGPU/sched-setprio.ll
+++ b/llvm/test/CodeGen/AMDGPU/sched-setprio.ll
@@ -10,7 +10,7 @@
; CHECK-ERR: S_SETPRIO mask contains invalid memory operation bits
-; Demonstrate that the mask bits that are compatible with s_set_prio
+; Demonstrate that the maks bits that are compatible with s_set_prio
; allow corresponding instruction types to be scheduled across the s_setprio
; instruction and that incompatible mask bits are rejected by the
; verifier.
>From 58214deb312e269c95de71fff825a65d1fd20dfb Mon Sep 17 00:00:00 2001
From: Frederik Harwath <fharwath at amd.com>
Date: Wed, 26 Aug 2026 03:16:43 -0400
Subject: [PATCH 06/10] Revert "[AMDGPU] Reject incompatible mask bits"
This reverts commit 96d8e5c353747c9fde54376a9999d780de7115c4.
---
llvm/docs/AMDGPUUsage.rst | 4 +-
llvm/lib/Target/AMDGPU/SIInstrInfo.cpp | 12 -
.../CodeGen/AMDGPU/llvm.amdgcn.s.setprio.ll | 12 +-
llvm/test/CodeGen/AMDGPU/sched-setprio.ll | 275 +-----------------
4 files changed, 21 insertions(+), 282 deletions(-)
diff --git a/llvm/docs/AMDGPUUsage.rst b/llvm/docs/AMDGPUUsage.rst
index e2af2257e8402..6bd1f17962e16 100644
--- a/llvm/docs/AMDGPUUsage.rst
+++ b/llvm/docs/AMDGPUUsage.rst
@@ -1919,9 +1919,7 @@ The AMDGPU backend implements the following LLVM IR intrinsics.
the intrinsic in the same way as llvm.amdgcn.sched.barrier. Takes two parameters:
- Priority (i16): Hardware wave priority level. 0 = lowest, 3 = highest.
- - Mask (i32): Scheduling barrier mask. Memory mask bits (VMEM, VMEM_READ,
- VMEM_WRITE, DS, DS_READ, DS_WRITE, LDSDMA) may not be set since they
- contradict the S_SETPRIO instruction's inherent scheduling dependencies.
+ - Mask (i32): Scheduling barrier mask.
llvm.amdgcn.s.setprio Like llvm.amdgcn.s.setprio.mask with Mask = 0.
diff --git a/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp b/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp
index 42a8d52f0a5b5..d0f7cb0981d8c 100644
--- a/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp
@@ -5973,18 +5973,6 @@ bool SIInstrInfo::verifyInstruction(const MachineInstr &MI,
}
}
- if (Opcode == AMDGPU::S_SETPRIO) {
- constexpr unsigned BarrierMemMaskBits =
- (1u << 4) | (1u << 5) | (1u << 6) | // VMEM, VMEM_READ, VMEM_WRITE
- (1u << 7) | (1u << 8) | (1u << 9) | // DS, DS_READ, DS_WRITE
- (1u << 11); // LDSDMA
- unsigned Mask = MI.getOperand(1).getImm();
- if (Mask & BarrierMemMaskBits) {
- ErrInfo = "S_SETPRIO mask contains invalid memory operation bits";
- return false;
- }
- }
-
return true;
}
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.setprio.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.setprio.ll
index 98dea66d3d99e..0592ca3bbfb14 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.setprio.ll
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.setprio.ll
@@ -47,14 +47,14 @@ define void @test_llvm_amdgcn_s_setprio() #0 {
ret void
}
-define void @test_llvm_amdgcn_s_setprio_mask0() #0 {
-; GFX9-LABEL: test_llvm_amdgcn_s_setprio_mask0:
+define void @test_llvm_amdgcn_s_setprio_mask_full() #0 {
+; GFX9-LABEL: test_llvm_amdgcn_s_setprio_mask_full:
; GFX9: ; %bb.0:
; GFX9-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0) ; encoding: [0x00,0x00,0x8c,0xbf]
; GFX9-NEXT: s_setprio 3 ; encoding: [0x03,0x00,0x8f,0xbf]
; GFX9-NEXT: s_setpc_b64 s[30:31] ; encoding: [0x1e,0x1d,0x80,0xbe]
;
-; SI-LABEL: test_llvm_amdgcn_s_setprio_mask0:
+; SI-LABEL: test_llvm_amdgcn_s_setprio_mask_full:
; SI: ; %bb.0:
; SI-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0) ; encoding: [0x00,0x00,0x8c,0xbf]
; SI-NEXT: s_setprio 3 ; encoding: [0x03,0x00,0x8f,0xbf]
@@ -63,14 +63,14 @@ define void @test_llvm_amdgcn_s_setprio_mask0() #0 {
ret void
}
-define void @test_llvm_amdgcn_s_setprio_mask15() #0 {
-; GFX9-LABEL: test_llvm_amdgcn_s_setprio_mask15:
+define void @test_llvm_amdgcn_s_setprio_mask_selective() #0 {
+; GFX9-LABEL: test_llvm_amdgcn_s_setprio_mask_selective:
; GFX9: ; %bb.0:
; GFX9-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0) ; encoding: [0x00,0x00,0x8c,0xbf]
; GFX9-NEXT: s_setprio 1 ; encoding: [0x01,0x00,0x8f,0xbf]
; GFX9-NEXT: s_setpc_b64 s[30:31] ; encoding: [0x1e,0x1d,0x80,0xbe]
;
-; SI-LABEL: test_llvm_amdgcn_s_setprio_mask15:
+; SI-LABEL: test_llvm_amdgcn_s_setprio_mask_selective:
; SI: ; %bb.0:
; SI-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0) ; encoding: [0x00,0x00,0x8c,0xbf]
; SI-NEXT: s_setprio 1 ; encoding: [0x01,0x00,0x8f,0xbf]
diff --git a/llvm/test/CodeGen/AMDGPU/sched-setprio.ll b/llvm/test/CodeGen/AMDGPU/sched-setprio.ll
index da20dfb37dd36..c6b83fd715782 100644
--- a/llvm/test/CodeGen/AMDGPU/sched-setprio.ll
+++ b/llvm/test/CodeGen/AMDGPU/sched-setprio.ll
@@ -1,267 +1,20 @@
-; RUN: split-file %s %t
-; RUN: llc -mtriple=amdgpu9.08 < %t/valid.ll | FileCheck --check-prefix=GCN %s
-; RUN: not --crash llc -mtriple=amdgpu9.08 -verify-machineinstrs < %t/vmem.ll 2>&1 | FileCheck --check-prefix=CHECK-ERR %s
-; RUN: not --crash llc -mtriple=amdgpu9.08 -verify-machineinstrs < %t/vmem-read.ll 2>&1 | FileCheck --check-prefix=CHECK-ERR %s
-; RUN: not --crash llc -mtriple=amdgpu9.08 -verify-machineinstrs < %t/vmem-write.ll 2>&1 | FileCheck --check-prefix=CHECK-ERR %s
-; RUN: not --crash llc -mtriple=amdgpu9.08 -verify-machineinstrs < %t/ds.ll 2>&1 | FileCheck --check-prefix=CHECK-ERR %s
-; RUN: not --crash llc -mtriple=amdgpu9.08 -verify-machineinstrs < %t/ds-read.ll 2>&1 | FileCheck --check-prefix=CHECK-ERR %s
-; RUN: not --crash llc -mtriple=amdgpu9.08 -verify-machineinstrs < %t/ds-write.ll 2>&1 | FileCheck --check-prefix=CHECK-ERR %s
-; RUN: not --crash llc -mtriple=amdgpu9.08 -verify-machineinstrs < %t/ldsdma.ll 2>&1 | FileCheck --check-prefix=CHECK-ERR %s
-
-; CHECK-ERR: S_SETPRIO mask contains invalid memory operation bits
-
-; Demonstrate that the maks bits that are compatible with s_set_prio
-; allow corresponding instruction types to be scheduled across the s_setprio
-; instruction and that incompatible mask bits are rejected by the
-; verifier.
-
-;--- valid.ll
+; RUN: llc -mtriple=amdgpu9.08 < %s | FileCheck --check-prefix=GCN %s
declare void @llvm.amdgcn.s.setprio(i16)
-declare void @llvm.amdgcn.s.setprio.mask(i16, i32)
-declare float @llvm.amdgcn.rcp.f32(float)
declare <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float, float, <4 x float>, i32, i32, i32)
-; GCN-LABEL: {{^}}test_mask0_blocks_salu:
-; GCN: s_add_i32
-; GCN-NEXT: s_setprio 1
-; GCN-NEXT: s_add_i32
-; GCN-NEXT: s_add_i32
-define amdgpu_cs void @test_mask0_blocks_salu(ptr addrspace(1) %out, i32 inreg %x, i32 inreg %y) {
- %add1 = add i32 %x, %y
- call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 0)
- %add2 = add i32 %y, %y
- %sum = add i32 %add1, %add2
- store i32 %sum, ptr addrspace(1) %out
- ret void
-}
-
-; GCN-LABEL: {{^}}test_mask0_blocks_valu:
-; GCN: v_add_f32_e32 v{{[0-9]+}}, 1.0
-; GCN-NEXT: s_setprio 1
-; GCN-NEXT: v_add_f32_e32 v{{[0-9]+}}, 2.0
-define amdgpu_cs void @test_mask0_blocks_valu(ptr addrspace(1) %out, float %x, float %y) {
- %add1 = fadd float %x, 1.0
- call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 0)
- %add2 = fadd float %y, 2.0
- %sum = fadd float %add1, %add2
- store float %sum, ptr addrspace(1) %out
- ret void
-}
-
-; GCN-LABEL: {{^}}test_mask0_blocks_trans:
-; GCN: v_rcp_f32_e32 v{{[0-9]+}}, v{{[0-9]+}}
-; GCN-NEXT: s_setprio 1
-; GCN-NEXT: v_rcp_f32_e32 v{{[0-9]+}}, v{{[0-9]+}}
-define amdgpu_cs void @test_mask0_blocks_trans(ptr addrspace(1) %out, float %x, float %y) {
- %rcp1 = call float @llvm.amdgcn.rcp.f32(float %x)
- call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 0)
- %rcp2 = call float @llvm.amdgcn.rcp.f32(float %y)
- %sum = fadd float %rcp1, %rcp2
- store float %sum, ptr addrspace(1) %out
- ret void
-}
-
-; GCN-LABEL: {{^}}test_mask4_allows_salu:
-; GCN: s_add_i32
-; GCN-NEXT: s_add_i32
-; GCN-NEXT: s_add_i32
-; GCN-NEXT: s_setprio 1
-define amdgpu_cs void @test_mask4_allows_salu(ptr addrspace(1) %out, i32 inreg %x, i32 inreg %y) {
- %add1 = add i32 %x, %y
- call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 4)
- %add2 = add i32 %y, %y
- %sum = add i32 %add1, %add2
- store i32 %sum, ptr addrspace(1) %out
- ret void
-}
-
-; GCN-LABEL: {{^}}test_mask4_blocks_valu:
-; GCN: v_add_f32_e32 v{{[0-9]+}}, 1.0
-; GCN-NEXT: s_setprio 1
-; GCN-NEXT: v_add_f32_e32 v{{[0-9]+}}, 2.0
-define amdgpu_cs void @test_mask4_blocks_valu(ptr addrspace(1) %out, float %x, float %y) {
- %add1 = fadd float %x, 1.0
- call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 4)
- %add2 = fadd float %y, 2.0
- %sum = fadd float %add1, %add2
- store float %sum, ptr addrspace(1) %out
- ret void
-}
-
-; GCN-LABEL: {{^}}test_mask2_allows_valu:
-; GCN: v_add_f32_e32
-; GCN-NEXT: v_add_f32_e32
-; GCN-NEXT: v_add_f32_e32
-; GCN: s_setprio 1
-define amdgpu_cs void @test_mask2_allows_valu(ptr addrspace(1) %out, float %x, float %y) {
- %add1 = fadd float %x, 1.0
- call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 2)
- %add2 = fadd float %y, 2.0
- %sum = fadd float %add1, %add2
- store float %sum, ptr addrspace(1) %out
- ret void
-}
-
-; GCN-LABEL: {{^}}test_mask2_blocks_salu:
-; GCN: s_add_i32
-; GCN-NEXT: s_setprio 1
-; GCN-NEXT: s_add_i32
-; GCN-NEXT: s_add_i32
-define amdgpu_cs void @test_mask2_blocks_salu(ptr addrspace(1) %out, i32 inreg %x, i32 inreg %y) {
- %add1 = add i32 %x, %y
- call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 2)
- %add2 = add i32 %y, %y
- %sum = add i32 %add1, %add2
- store i32 %sum, ptr addrspace(1) %out
- ret void
-}
-
-; GCN-LABEL: {{^}}test_mask1_allows_salu:
-; GCN: s_add_i32
-; GCN-NEXT: s_add_i32
-; GCN-NEXT: s_add_i32
-; GCN: s_setprio 1
-define amdgpu_cs void @test_mask1_allows_salu(ptr addrspace(1) %out, i32 inreg %x, i32 inreg %y) {
- %add1 = add i32 %x, %y
- call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 1)
- %add2 = add i32 %y, %y
- %sum = add i32 %add1, %add2
- store i32 %sum, ptr addrspace(1) %out
- ret void
-}
-
-; GCN-LABEL: {{^}}test_mask1_allows_valu:
-; GCN: v_add_f32_e32
-; GCN-NEXT: v_add_f32_e32
-; GCN-NEXT: v_add_f32_e32
-; GCN: s_setprio 1
-define amdgpu_cs void @test_mask1_allows_valu(ptr addrspace(1) %out, float %x, float %y) {
- %add1 = fadd float %x, 1.0
- call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 1)
- %add2 = fadd float %y, 2.0
- %sum = fadd float %add1, %add2
- store float %sum, ptr addrspace(1) %out
- ret void
-}
-
-; GCN-LABEL: {{^}}test_mask1024_allows_trans:
-; GCN: v_rcp_f32_e32
-; GCN-NEXT: v_rcp_f32_e32
-; GCN: s_setprio 1
-define amdgpu_cs void @test_mask1024_allows_trans(ptr addrspace(1) %out, float %x, float %y) {
- %rcp1 = call float @llvm.amdgcn.rcp.f32(float %x)
- call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 1024)
- %rcp2 = call float @llvm.amdgcn.rcp.f32(float %y)
- %sum = fadd float %rcp1, %rcp2
- store float %sum, ptr addrspace(1) %out
- ret void
-}
-
-; GCN-LABEL: {{^}}test_mask1024_blocks_salu:
-; GCN: s_add_i32
-; GCN-NEXT: s_setprio 1
-; GCN-NEXT: s_add_i32
-; GCN-NEXT: s_add_i32
-define amdgpu_cs void @test_mask1024_blocks_salu(ptr addrspace(1) %out, i32 inreg %x, i32 inreg %y) {
- %add1 = add i32 %x, %y
- call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 1024)
- %add2 = add i32 %y, %y
- %sum = add i32 %add1, %add2
- store i32 %sum, ptr addrspace(1) %out
- ret void
-}
-
-; GCN-LABEL: {{^}}test_mask6_allows_salu_and_valu:
-; GCN: s_add_i32
-; GCN-NEXT: s_add_i32
-; GCN-NEXT: s_add_i32
-; GCN: s_setprio 1
-define amdgpu_cs void @test_mask6_allows_salu_and_valu(ptr addrspace(1) %out, i32 inreg %x, i32 inreg %y) {
- %add1 = add i32 %x, %y
- call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 6)
- %add2 = add i32 %y, %y
- %sum = add i32 %add1, %add2
- store i32 %sum, ptr addrspace(1) %out
- ret void
-}
-
-; GCN-LABEL: {{^}}test_mask8_allows_mfma:
-; GCN: s_setprio 1
-; GCN: v_mfma_f32_4x4x1f32
-; GCN: v_mfma_f32_4x4x1f32
-define amdgpu_cs void @test_mask8_allows_mfma(ptr addrspace(1) %out, <4 x float> %in) {
- %mfma1 = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float 1.0, float 2.0, <4 x float> %in, i32 0, i32 0, i32 0)
- call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 8)
- %mfma2 = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float 3.0, float 4.0, <4 x float> %mfma1, i32 0, i32 0, i32 0)
- store <4 x float> %mfma2, ptr addrspace(1) %out
- ret void
-}
-
-; GCN-LABEL: {{^}}test_mask0_blocks_mfma:
-; GCN: v_mfma_f32_4x4x1f32
-; GCN: s_setprio 1
-; GCN: v_mfma_f32_4x4x1f32
-define amdgpu_cs void @test_mask0_blocks_mfma(ptr addrspace(1) %out, <4 x float> %in) {
- %mfma1 = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float 1.0, float 2.0, <4 x float> %in, i32 0, i32 0, i32 0)
- call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 0)
- %mfma2 = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float 3.0, float 4.0, <4 x float> %mfma1, i32 0, i32 0, i32 0)
- store <4 x float> %mfma2, ptr addrspace(1) %out
- ret void
-}
-
-;--- vmem.ll
-declare void @llvm.amdgcn.s.setprio.mask(i16, i32)
-
-define amdgpu_cs void @test_invalid_vmem() {
- call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 16)
- ret void
-}
-
-;--- vmem-read.ll
-declare void @llvm.amdgcn.s.setprio.mask(i16, i32)
-
-define amdgpu_cs void @test_invalid_vmem_read() {
- call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 32)
- ret void
-}
-
-;--- vmem-write.ll
-declare void @llvm.amdgcn.s.setprio.mask(i16, i32)
-
-define amdgpu_cs void @test_invalid_vmem_write() {
- call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 64)
- ret void
-}
-
-;--- ds.ll
-declare void @llvm.amdgcn.s.setprio.mask(i16, i32)
-
-define amdgpu_cs void @test_invalid_ds() {
- call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 128)
- ret void
-}
-
-;--- ds-read.ll
-declare void @llvm.amdgcn.s.setprio.mask(i16, i32)
-
-define amdgpu_cs void @test_invalid_ds_read() {
- call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 256)
- ret void
-}
-
-;--- ds-write.ll
-declare void @llvm.amdgcn.s.setprio.mask(i16, i32)
-
-define amdgpu_cs void @test_invalid_ds_write() {
- call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 512)
- ret void
-}
-
-;--- ldsdma.ll
-declare void @llvm.amdgcn.s.setprio.mask(i16, i32)
-
-define amdgpu_cs void @test_invalid_ldsdma() {
- call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 2048)
+; GCN-LABEL: {{^}}test_mfma_f32_4x4x1f32:
+; GCN: s_setprio 1
+; GCN: v_mfma
+; GCN: v_mfma
+; GCN: s_setprio 0
+define amdgpu_kernel void @test_mfma_f32_4x4x1f32(ptr addrspace(1) %arg) #0 {
+bb:
+ %in.1 = load <4 x float>, ptr addrspace(1) %arg
+ call void @llvm.amdgcn.s.setprio(i16 1)
+ %mai.1 = tail call <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float 1.0, float 2.0, <4 x float> %in.1, i32 0, i32 0, i32 0)
+ %mai.2 = tail call <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float 3.0, float 4.0, <4 x float> %mai.1, i32 0, i32 0, i32 0)
+ call void @llvm.amdgcn.s.setprio(i16 0)
+ store <4 x float> %mai.2, ptr addrspace(1) %arg
ret void
}
>From b1f69a193712311dc27f140ae1d53bf9d3651a66 Mon Sep 17 00:00:00 2001
From: Frederik Harwath <fharwath at amd.com>
Date: Wed, 26 Aug 2026 03:17:33 -0400
Subject: [PATCH 07/10] Revert "[AMDGPU] Add scheduling mask operand to
S_SETPRIO"
This reverts commit 0ebed862ba7f75c767876e33f325438d6b858f49.
---
llvm/docs/AMDGPUUsage.rst | 8 -----
llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 6 ----
llvm/lib/Target/AMDGPU/AMDGPUIGroupLP.cpp | 18 ++++------
.../AMDGPU/AMDGPURegBankLegalizeRules.cpp | 1 -
.../Target/AMDGPU/AMDGPUSetWavePriority.cpp | 3 +-
.../lib/Target/AMDGPU/GCNHazardRecognizer.cpp | 9 ++---
llvm/lib/Target/AMDGPU/SIInstrInfo.cpp | 19 ++---------
llvm/lib/Target/AMDGPU/SIInstrInfo.h | 3 --
llvm/lib/Target/AMDGPU/SOPInstructions.td | 18 ++--------
.../CodeGen/AMDGPU/llvm.amdgcn.s.setprio.ll | 33 -------------------
.../AMDGPU/required-export-priority.mir | 12 +++----
11 files changed, 20 insertions(+), 110 deletions(-)
diff --git a/llvm/docs/AMDGPUUsage.rst b/llvm/docs/AMDGPUUsage.rst
index 6bd1f17962e16..443dd7f9772d7 100644
--- a/llvm/docs/AMDGPUUsage.rst
+++ b/llvm/docs/AMDGPUUsage.rst
@@ -1915,14 +1915,6 @@ The AMDGPU backend implements the following LLVM IR intrinsics.
- 0x0400: All Transcendental (e.g. V_EXP) instructions may be scheduled across sched_barrier.
- 0x0800: All LDSDMA instructions may be scheduled across sched_barrier.
- llvm.amdgcn.s.setprio.mask Sets the hardware wave priority and controls instruction scheduling across
- the intrinsic in the same way as llvm.amdgcn.sched.barrier. Takes two parameters:
-
- - Priority (i16): Hardware wave priority level. 0 = lowest, 3 = highest.
- - Mask (i32): Scheduling barrier mask.
-
- llvm.amdgcn.s.setprio Like llvm.amdgcn.s.setprio.mask with Mask = 0.
-
llvm.amdgcn.sched.group.barrier Creates schedule groups with specific properties to create custom scheduling
pipelines. The ordering between groups is enforced by the instruction scheduler.
The intrinsic applies to the code that precedes the intrinsic. The intrinsic
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index d6d5b69b28168..565637b36131c 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -2229,12 +2229,6 @@ def int_amdgcn_s_setprio :
DefaultAttrsIntrinsic<[], [llvm_i16_ty], [ImmArg<ArgIndex<0>>, IntrNoMem,
IntrHasSideEffects]>;
-def int_amdgcn_s_setprio_mask :
- ClangBuiltin<"__builtin_amdgcn_s_setprio_mask">,
- DefaultAttrsIntrinsic<[], [llvm_i16_ty, llvm_i32_ty],
- [ImmArg<ArgIndex<0>>, ImmArg<ArgIndex<1>>,
- IntrNoMem, IntrHasSideEffects]>;
-
def int_amdgcn_s_setprio_inc_wg :
ClangBuiltin<"__builtin_amdgcn_s_setprio_inc_wg">,
DefaultAttrsIntrinsic<[], [llvm_i16_ty], [ImmArg<ArgIndex<0>>, IntrNoMem,
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUIGroupLP.cpp b/llvm/lib/Target/AMDGPU/AMDGPUIGroupLP.cpp
index 04bf0525eeb52..df7e80d62e065 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUIGroupLP.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUIGroupLP.cpp
@@ -2746,7 +2746,7 @@ void IGroupLPDAGMutation::apply(ScheduleDAGInstrs *DAGInstrs) {
for (auto R = DAG->SUnits.rbegin(), E = DAG->SUnits.rend(); R != E; ++R) {
unsigned Opc = R->getInstr()->getOpcode();
// SCHED_[GROUP_]BARRIER and IGLP are mutually exclusive.
- if (TII->isSchedBarrierLike(*R->getInstr())) {
+ if (Opc == AMDGPU::SCHED_BARRIER) {
addSchedBarrierEdges(*R);
FoundSB = true;
} else if (Opc == AMDGPU::SCHED_GROUP_BARRIER) {
@@ -2771,17 +2771,11 @@ void IGroupLPDAGMutation::apply(ScheduleDAGInstrs *DAGInstrs) {
void IGroupLPDAGMutation::addSchedBarrierEdges(SUnit &SchedBarrier) {
MachineInstr &MI = *SchedBarrier.getInstr();
- assert(TII->isSchedBarrierLike(MI));
-
- SchedGroupMask Mask =
- static_cast<SchedGroupMask>(TII->getSchedBarrierLikeMask(MI));
-
- LLVM_DEBUG(
- dbgs()
- << "Building SchedGroup for SchedBarrier-like instruction with Mask: "
- << (unsigned)Mask << "\n");
-
- auto InvertedMask = invertSchedBarrierMask(Mask);
+ assert(MI.getOpcode() == AMDGPU::SCHED_BARRIER);
+ LLVM_DEBUG(dbgs() << "Building SchedGroup for SchedBarrier with Mask: "
+ << MI.getOperand(0).getImm() << "\n");
+ auto InvertedMask =
+ invertSchedBarrierMask((SchedGroupMask)MI.getOperand(0).getImm());
SchedGroup SG(InvertedMask, std::nullopt, DAG, TII);
for (SUnit &SU : DAG->SUnits)
diff --git a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp
index 3eab1bc2c4c4a..881afaf920037 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp
@@ -1821,7 +1821,6 @@ RegBankLegalizeRules::RegBankLegalizeRules(const GCNSubtarget &_ST,
amdgcn_s_nop,
amdgcn_s_sethalt,
amdgcn_s_setprio,
- amdgcn_s_setprio_mask,
amdgcn_s_setprio_inc_wg,
amdgcn_s_sleep,
amdgcn_s_ttracedata_imm,
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUSetWavePriority.cpp b/llvm/lib/Target/AMDGPU/AMDGPUSetWavePriority.cpp
index df3fe8c9562f3..e3448be9f44f5 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUSetWavePriority.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUSetWavePriority.cpp
@@ -85,8 +85,7 @@ AMDGPUSetWavePriority::BuildSetprioMI(MachineBasicBlock &MBB,
MachineBasicBlock::iterator I,
unsigned priority) const {
return BuildMI(MBB, I, DebugLoc(), TII->get(AMDGPU::S_SETPRIO))
- .addImm(priority)
- .addImm(0);
+ .addImm(priority);
}
// Checks that for every predecessor Pred that can reach a VMEM load,
diff --git a/llvm/lib/Target/AMDGPU/GCNHazardRecognizer.cpp b/llvm/lib/Target/AMDGPU/GCNHazardRecognizer.cpp
index 3970f701fe164..af326d60aacb4 100644
--- a/llvm/lib/Target/AMDGPU/GCNHazardRecognizer.cpp
+++ b/llvm/lib/Target/AMDGPU/GCNHazardRecognizer.cpp
@@ -3699,8 +3699,7 @@ static bool ensureEntrySetPrio(MachineFunction *MF, int Priority,
}
BuildMI(EntryMBB, EntryMBB.begin(), DebugLoc(), TII.get(AMDGPU::S_SETPRIO))
- .addImm(Priority)
- .addImm(0);
+ .addImm(Priority);
return true;
}
@@ -3778,8 +3777,7 @@ bool GCNHazardRecognizer::fixRequiredExportPriority(MachineInstr *MI) {
// Lower priority.
BuildMI(*MBB, NextMI, DL, TII.get(AMDGPU::S_SETPRIO))
- .addImm(PostExportPriority)
- .addImm(0);
+ .addImm(PostExportPriority);
if (!EndOfShader) {
// Wait for exports to complete.
@@ -3794,8 +3792,7 @@ bool GCNHazardRecognizer::fixRequiredExportPriority(MachineInstr *MI) {
if (!EndOfShader) {
// Return to normal (higher) priority.
BuildMI(*MBB, NextMI, DL, TII.get(AMDGPU::S_SETPRIO))
- .addImm(NormalPriority)
- .addImm(0);
+ .addImm(NormalPriority);
}
return true;
diff --git a/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp b/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp
index d0f7cb0981d8c..b2c3772cb0041 100644
--- a/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp
@@ -4462,22 +4462,6 @@ static bool changesVGPRIndexingMode(const MachineInstr &MI) {
}
}
-bool SIInstrInfo::isSchedBarrierLike(const MachineInstr &MI) const {
- unsigned Opc = MI.getOpcode();
- return Opc == AMDGPU::SCHED_BARRIER || Opc == AMDGPU::S_SETPRIO;
-}
-
-unsigned SIInstrInfo::getSchedBarrierLikeMask(const MachineInstr &MI) const {
- switch (MI.getOpcode()) {
- case AMDGPU::SCHED_BARRIER:
- return MI.getOperand(0).getImm();
- case AMDGPU::S_SETPRIO:
- return MI.getOperand(1).getImm();
- default:
- llvm_unreachable("Expected a SchedBarrier-like instruction");
- }
-}
-
bool SIInstrInfo::isSchedulingBoundary(const MachineInstr &MI,
const MachineBasicBlock *MBB,
const MachineFunction &MF) const {
@@ -4496,7 +4480,7 @@ bool SIInstrInfo::isSchedulingBoundary(const MachineInstr &MI,
if (MI.getOpcode() == TargetOpcode::INLINEASM_BR)
return true;
- if (isSchedBarrierLike(MI) && getSchedBarrierLikeMask(MI) == 0)
+ if (MI.getOpcode() == AMDGPU::SCHED_BARRIER && MI.getOperand(0).getImm() == 0)
return true;
// Target-independent instructions do not have an implicit-use of EXEC, even
@@ -4505,6 +4489,7 @@ bool SIInstrInfo::isSchedulingBoundary(const MachineInstr &MI,
return MI.modifiesRegister(AMDGPU::EXEC, &RI) ||
MI.getOpcode() == AMDGPU::S_SETREG_IMM32_B32 ||
MI.getOpcode() == AMDGPU::S_SETREG_B32 ||
+ MI.getOpcode() == AMDGPU::S_SETPRIO ||
MI.getOpcode() == AMDGPU::S_SETPRIO_INC_WG ||
changesVGPRIndexingMode(MI);
}
diff --git a/llvm/lib/Target/AMDGPU/SIInstrInfo.h b/llvm/lib/Target/AMDGPU/SIInstrInfo.h
index 75c4eab1be937..4c8641a6091d7 100644
--- a/llvm/lib/Target/AMDGPU/SIInstrInfo.h
+++ b/llvm/lib/Target/AMDGPU/SIInstrInfo.h
@@ -1149,9 +1149,6 @@ class SIInstrInfo final : public AMDGPUGenInstrInfo {
Opcode == AMDGPU::DS_GWS_BARRIER;
}
- bool isSchedBarrierLike(const MachineInstr &MI) const;
- unsigned getSchedBarrierLikeMask(const MachineInstr &MI) const;
-
static bool isLoadMonitor(unsigned Opc) {
switch (Opc) {
case AMDGPU::GLOBAL_LOAD_MONITOR_B32:
diff --git a/llvm/lib/Target/AMDGPU/SOPInstructions.td b/llvm/lib/Target/AMDGPU/SOPInstructions.td
index b6b8ba1720994..8fa54de570e67 100644
--- a/llvm/lib/Target/AMDGPU/SOPInstructions.td
+++ b/llvm/lib/Target/AMDGPU/SOPInstructions.td
@@ -1754,24 +1754,10 @@ def S_SLEEP_VAR : SOP1_0_32 <"s_sleep_var", [(int_amdgcn_s_sleep_var SSrc_b32:$s
let hasSideEffects = 1;
}
-def S_SETPRIO : SOPP_Pseudo <"s_setprio", (ins i16imm:$simm16, i32imm:$mask), "$simm16",
- []> {
+def S_SETPRIO : SOPP_Pseudo <"s_setprio", (ins i16imm:$simm16), "$simm16",
+ [(int_amdgcn_s_setprio timm:$simm16)]> {
}
-} // End hasSideEffects = 1
-
-def : GCNPat<
- (int_amdgcn_s_setprio timm:$priority),
- (S_SETPRIO timm:$priority, (i32 0))
->;
-
-def : GCNPat<
- (int_amdgcn_s_setprio_mask timm:$priority, timm:$mask),
- (S_SETPRIO timm:$priority, timm:$mask)
->;
-
-let hasSideEffects = 1 in {
-
def S_SETPRIO_INC_WG : SOPP_Pseudo <"s_setprio_inc_wg", (ins i16imm:$simm16), "$simm16",
[(int_amdgcn_s_setprio_inc_wg timm:$simm16)]> {
let SubtargetPredicate = HasSetPrioIncWgInst;
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.setprio.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.setprio.ll
index 0592ca3bbfb14..87df1e152bef5 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.setprio.ll
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.setprio.ll
@@ -5,7 +5,6 @@
; RUN: llc -global-isel -mtriple=amdgpu9.0a -show-mc-encoding < %s | FileCheck -check-prefix=GFX9 %s
declare void @llvm.amdgcn.s.setprio(i16) #0
-declare void @llvm.amdgcn.s.setprio.mask(i16, i32) #0
define void @test_llvm_amdgcn_s_setprio() #0 {
; GFX9-LABEL: test_llvm_amdgcn_s_setprio:
@@ -47,36 +46,4 @@ define void @test_llvm_amdgcn_s_setprio() #0 {
ret void
}
-define void @test_llvm_amdgcn_s_setprio_mask_full() #0 {
-; GFX9-LABEL: test_llvm_amdgcn_s_setprio_mask_full:
-; GFX9: ; %bb.0:
-; GFX9-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0) ; encoding: [0x00,0x00,0x8c,0xbf]
-; GFX9-NEXT: s_setprio 3 ; encoding: [0x03,0x00,0x8f,0xbf]
-; GFX9-NEXT: s_setpc_b64 s[30:31] ; encoding: [0x1e,0x1d,0x80,0xbe]
-;
-; SI-LABEL: test_llvm_amdgcn_s_setprio_mask_full:
-; SI: ; %bb.0:
-; SI-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0) ; encoding: [0x00,0x00,0x8c,0xbf]
-; SI-NEXT: s_setprio 3 ; encoding: [0x03,0x00,0x8f,0xbf]
-; SI-NEXT: s_setpc_b64 s[30:31] ; encoding: [0x1e,0x20,0x80,0xbe]
- call void @llvm.amdgcn.s.setprio.mask(i16 3, i32 0)
- ret void
-}
-
-define void @test_llvm_amdgcn_s_setprio_mask_selective() #0 {
-; GFX9-LABEL: test_llvm_amdgcn_s_setprio_mask_selective:
-; GFX9: ; %bb.0:
-; GFX9-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0) ; encoding: [0x00,0x00,0x8c,0xbf]
-; GFX9-NEXT: s_setprio 1 ; encoding: [0x01,0x00,0x8f,0xbf]
-; GFX9-NEXT: s_setpc_b64 s[30:31] ; encoding: [0x1e,0x1d,0x80,0xbe]
-;
-; SI-LABEL: test_llvm_amdgcn_s_setprio_mask_selective:
-; SI: ; %bb.0:
-; SI-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0) ; encoding: [0x00,0x00,0x8c,0xbf]
-; SI-NEXT: s_setprio 1 ; encoding: [0x01,0x00,0x8f,0xbf]
-; SI-NEXT: s_setpc_b64 s[30:31] ; encoding: [0x1e,0x20,0x80,0xbe]
- call void @llvm.amdgcn.s.setprio.mask(i16 1, i32 15)
- ret void
-}
-
attributes #0 = { nounwind }
diff --git a/llvm/test/CodeGen/AMDGPU/required-export-priority.mir b/llvm/test/CodeGen/AMDGPU/required-export-priority.mir
index fc65dad5cdd5e..cf12eff2c0527 100644
--- a/llvm/test/CodeGen/AMDGPU/required-export-priority.mir
+++ b/llvm/test/CodeGen/AMDGPU/required-export-priority.mir
@@ -334,15 +334,15 @@ body: |
bb.1:
liveins: $vgpr0
- S_SETPRIO 3, 0
+ S_SETPRIO 3
$vgpr0 = V_OR_B32_e32 2, $vgpr0, implicit $exec
- S_SETPRIO 0, 0
+ S_SETPRIO 0
bb.2:
liveins: $vgpr0
- S_SETPRIO 1, 0
+ S_SETPRIO 1
$vgpr0 = V_OR_B32_e32 3, $vgpr0, implicit $exec
- S_SETPRIO 0, 0
+ S_SETPRIO 0
bb.3:
liveins: $vgpr0
@@ -381,9 +381,9 @@ body: |
; GFX1150-NEXT: S_SETPRIO 2
; GFX1150-NEXT: S_SETPRIO 3
; GFX1150-NEXT: S_ENDPGM 0
- S_SETPRIO 3, 0
+ S_SETPRIO 3
EXP 1, $vgpr0, $vgpr0, $vgpr0, $vgpr0, -1, -1, 15, implicit $exec
EXP_DONE 0, $vgpr0, $vgpr0, $vgpr0, $vgpr0, -1, -1, 15, implicit $exec
- S_SETPRIO 3, 0
+ S_SETPRIO 3
S_ENDPGM 0
...
>From 82a26ab036dfedba5f6d0acfb6c787c55434d807 Mon Sep 17 00:00:00 2001
From: Frederik Harwath <fharwath at amd.com>
Date: Wed, 26 Aug 2026 04:49:17 -0400
Subject: [PATCH 08/10] [AMDGPU] Add llvm.amdgcn.s.setprio.relaxed intrinsic
The S_SETPRIO instruction is treated as a scheduling boundary to prevent
unintended scheduling of VALU instructions across S_SETPRIO.
This restrics the scheduler more than necessary.
Stop treating S_SETPRIO as a scheduling boundary. Change the
llvm.amdgcn.s.setprio intrinsic to wrap S_SETPRIO in SCHED_BARRIER 0
instructions which effectively achieves the same scheduling behavior
as before. Add new llvm.amdgcn.s.setprio.relaxed intrinsic which
lowers into a single S_SETPRIO instruction. This can can be used to
pair S_SETPRIO with other less restrictive scheduling barriers.
---
clang/include/clang/Basic/BuiltinsAMDGPU.td | 1 +
clang/test/CodeGenOpenCL/builtins-amdgcn.cl | 9 ++
.../test/SemaOpenCL/builtins-amdgcn-error.cl | 6 +
llvm/docs/AMDGPUUsage.rst | 6 +
llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 6 +
.../AMDGPU/AMDGPURegBankLegalizeRules.cpp | 1 +
llvm/lib/Target/AMDGPU/SIISelLowering.cpp | 8 ++
llvm/lib/Target/AMDGPU/SIInstrInfo.cpp | 1 -
llvm/lib/Target/AMDGPU/SIInstructions.td | 10 ++
llvm/lib/Target/AMDGPU/SOPInstructions.td | 2 +-
.../CodeGen/AMDGPU/llvm.amdgcn.s.setprio.ll | 36 +++++
.../AMDGPU/llvm.amdgcn.s.setprio.relaxed.ll | 49 +++++++
.../AMDGPU/required-export-priority.ll | 4 +
.../CodeGen/AMDGPU/sched-setprio-relaxed.ll | 134 ++++++++++++++++++
llvm/test/CodeGen/AMDGPU/sched-setprio.ll | 4 +-
15 files changed, 274 insertions(+), 3 deletions(-)
create mode 100644 llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.setprio.relaxed.ll
create mode 100644 llvm/test/CodeGen/AMDGPU/sched-setprio-relaxed.ll
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.td b/clang/include/clang/Basic/BuiltinsAMDGPU.td
index 73b27a2b5c5cd..bd3f7ae08eb0c 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.td
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.td
@@ -208,6 +208,7 @@ def __builtin_amdgcn_s_sleep : AMDGPUBuiltin<"void(_Constant int)">;
def __builtin_amdgcn_s_incperflevel : AMDGPUBuiltin<"void(_Constant int)">;
def __builtin_amdgcn_s_decperflevel : AMDGPUBuiltin<"void(_Constant int)">;
def __builtin_amdgcn_s_setprio : AMDGPUBuiltin<"void(_Constant short)">;
+def __builtin_amdgcn_s_setprio_relaxed : AMDGPUBuiltin<"void(_Constant short)">;
def __builtin_amdgcn_ds_swizzle : AMDGPUBuiltin<"int(int, _Constant int)", [Const]>;
def __builtin_amdgcn_ds_permute : AMDGPUBuiltin<"int(int, int)", [Const]>;
def __builtin_amdgcn_ds_bpermute : AMDGPUBuiltin<"int(int, int)", [Const]>;
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
index 28c420a5760f4..1375806125034 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
@@ -1041,6 +1041,15 @@ void test_s_setprio()
__builtin_amdgcn_s_setprio(3);
}
+// CHECK-LABEL: @test_s_setprio_relaxed
+// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.s.setprio.relaxed(i16 0)
+// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.s.setprio.relaxed(i16 3)
+void test_s_setprio_relaxed()
+{
+ __builtin_amdgcn_s_setprio_relaxed(0);
+ __builtin_amdgcn_s_setprio_relaxed(3);
+}
+
// CHECK-LABEL: @test_read_exec(
// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.ballot.i64(i1 true)
void test_read_exec(global ulong* out) {
diff --git a/clang/test/SemaOpenCL/builtins-amdgcn-error.cl b/clang/test/SemaOpenCL/builtins-amdgcn-error.cl
index 10f00dd7da5f9..536de3704274f 100644
--- a/clang/test/SemaOpenCL/builtins-amdgcn-error.cl
+++ b/clang/test/SemaOpenCL/builtins-amdgcn-error.cl
@@ -60,6 +60,12 @@ void test_s_setprio(int x)
__builtin_amdgcn_s_setprio(65536); // expected-warning {{implicit conversion from 'int' to 'short' changes value from 65536 to 0}}
}
+void test_s_setprio_relaxed(int x)
+{
+ __builtin_amdgcn_s_setprio_relaxed(x); // expected-error {{argument to '__builtin_amdgcn_s_setprio_relaxed' must be a constant integer}}
+ __builtin_amdgcn_s_setprio_relaxed(65536); // expected-warning {{implicit conversion from 'int' to 'short' changes value from 65536 to 0}}
+}
+
void test_sched_barrier(int x)
{
__builtin_amdgcn_sched_barrier(x); // expected-error {{argument to '__builtin_amdgcn_sched_barrier' must be a constant integer}}
diff --git a/llvm/docs/AMDGPUUsage.rst b/llvm/docs/AMDGPUUsage.rst
index 443dd7f9772d7..0d1a551e0725a 100644
--- a/llvm/docs/AMDGPUUsage.rst
+++ b/llvm/docs/AMDGPUUsage.rst
@@ -1954,6 +1954,12 @@ The AMDGPU backend implements the following LLVM IR intrinsics.
The iglp_opt strategy implementations are subject to change.
+ llvm.amdgcn.s.setprio Provide access to the ``s_setprio`` instruction for setting the wave priority to the given value.
+ The intrinsic ensures that the instruction acts as a full barrier for instruction scheduling, i.e.
+ no instruction may be scheduled across it.
+
+ llvm.amdgcn.s.setprio.relaxed Like ``llvm.amdgcn.s.setprio``, but does not act as a scheduling barrier.
+
llvm.amdgcn.s.getpc Provides access to the s_getpc_b64 instruction, but with the return value
sign-extended from the width of the underlying PC hardware register even on
processors where the s_getpc_b64 instruction returns a zero-extended value.
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index 565637b36131c..2fad48178e436 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -2229,6 +2229,12 @@ def int_amdgcn_s_setprio :
DefaultAttrsIntrinsic<[], [llvm_i16_ty], [ImmArg<ArgIndex<0>>, IntrNoMem,
IntrHasSideEffects]>;
+// Like int_amdgcn_s_setprio but without scheduling boundary behavior.
+def int_amdgcn_s_setprio_relaxed :
+ ClangBuiltin<"__builtin_amdgcn_s_setprio_relaxed">,
+ DefaultAttrsIntrinsic<[], [llvm_i16_ty], [ImmArg<ArgIndex<0>>, IntrNoMem,
+ IntrHasSideEffects]>;
+
def int_amdgcn_s_setprio_inc_wg :
ClangBuiltin<"__builtin_amdgcn_s_setprio_inc_wg">,
DefaultAttrsIntrinsic<[], [llvm_i16_ty], [ImmArg<ArgIndex<0>>, IntrNoMem,
diff --git a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp
index 881afaf920037..dca4617ece0a7 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp
@@ -1822,6 +1822,7 @@ RegBankLegalizeRules::RegBankLegalizeRules(const GCNSubtarget &_ST,
amdgcn_s_sethalt,
amdgcn_s_setprio,
amdgcn_s_setprio_inc_wg,
+ amdgcn_s_setprio_relaxed,
amdgcn_s_sleep,
amdgcn_s_ttracedata_imm,
amdgcn_s_wait_asynccnt,
diff --git a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
index 78bbb7f2d6146..a1bd173af3ed8 100644
--- a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
+++ b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
@@ -7361,6 +7361,14 @@ SITargetLowering::EmitInstrWithCustomInserter(MachineInstr &MI,
MRI.setSimpleHint(MI.getOperand(0).getReg(), MI.getOperand(6).getReg());
return BB;
}
+ case AMDGPU::S_SETPRIO_BARRIER: {
+ int64_t Priority = MI.getOperand(0).getImm();
+ BuildMI(*BB, MI, DL, TII->get(AMDGPU::SCHED_BARRIER)).addImm(0);
+ BuildMI(*BB, MI, DL, TII->get(AMDGPU::S_SETPRIO)).addImm(Priority);
+ BuildMI(*BB, MI, DL, TII->get(AMDGPU::SCHED_BARRIER)).addImm(0);
+ MI.eraseFromParent();
+ return BB;
+ }
default:
if (TII->isImage(MI) || TII->isMUBUF(MI)) {
if (!MI.mayStore())
diff --git a/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp b/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp
index b2c3772cb0041..793c30d8369cb 100644
--- a/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp
@@ -4489,7 +4489,6 @@ bool SIInstrInfo::isSchedulingBoundary(const MachineInstr &MI,
return MI.modifiesRegister(AMDGPU::EXEC, &RI) ||
MI.getOpcode() == AMDGPU::S_SETREG_IMM32_B32 ||
MI.getOpcode() == AMDGPU::S_SETREG_B32 ||
- MI.getOpcode() == AMDGPU::S_SETPRIO ||
MI.getOpcode() == AMDGPU::S_SETPRIO_INC_WG ||
changesVGPRIndexingMode(MI);
}
diff --git a/llvm/lib/Target/AMDGPU/SIInstructions.td b/llvm/lib/Target/AMDGPU/SIInstructions.td
index 39df5dfd4d01d..7a897a20ed707 100644
--- a/llvm/lib/Target/AMDGPU/SIInstructions.td
+++ b/llvm/lib/Target/AMDGPU/SIInstructions.td
@@ -553,6 +553,16 @@ def SCHED_BARRIER : SPseudoInstSI<(outs), (ins i32imm:$mask),
let isMeta = 1;
}
+// Expands to SCHED_BARRIER 0 + S_SETPRIO + SCHED_BARRIER 0 to provide
+// backwards compatible behavior for setprio intrinsic since S_SETPRIO
+// used to be a scheduling boundary. It no longer is to allow for more
+// control over the scheduling behavior.
+def S_SETPRIO_BARRIER : SPseudoInstSI<(outs), (ins i16imm:$priority),
+ [(int_amdgcn_s_setprio timm:$priority)]> {
+ let usesCustomInserter = 1;
+ let hasSideEffects = 1;
+}
+
def SCHED_GROUP_BARRIER : SPseudoInstSI<
(outs),
(ins i32imm:$mask, i32imm:$size, i32imm:$syncid),
diff --git a/llvm/lib/Target/AMDGPU/SOPInstructions.td b/llvm/lib/Target/AMDGPU/SOPInstructions.td
index 8fa54de570e67..118091210d33d 100644
--- a/llvm/lib/Target/AMDGPU/SOPInstructions.td
+++ b/llvm/lib/Target/AMDGPU/SOPInstructions.td
@@ -1755,7 +1755,7 @@ def S_SLEEP_VAR : SOP1_0_32 <"s_sleep_var", [(int_amdgcn_s_sleep_var SSrc_b32:$s
}
def S_SETPRIO : SOPP_Pseudo <"s_setprio", (ins i16imm:$simm16), "$simm16",
- [(int_amdgcn_s_setprio timm:$simm16)]> {
+ [(int_amdgcn_s_setprio_relaxed timm:$simm16)]> {
}
def S_SETPRIO_INC_WG : SOPP_Pseudo <"s_setprio_inc_wg", (ins i16imm:$simm16), "$simm16",
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.setprio.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.setprio.ll
index 87df1e152bef5..2604361f1bd06 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.setprio.ll
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.setprio.ll
@@ -9,30 +9,66 @@ declare void @llvm.amdgcn.s.setprio(i16) #0
define void @test_llvm_amdgcn_s_setprio() #0 {
; GFX9-LABEL: test_llvm_amdgcn_s_setprio:
; GFX9: ; %bb.0:
+; GFX9-NEXT: ; sched_barrier mask(0x00000000)
; GFX9-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0) ; encoding: [0x00,0x00,0x8c,0xbf]
; GFX9-NEXT: s_setprio 0 ; encoding: [0x00,0x00,0x8f,0xbf]
+; GFX9-NEXT: ; sched_barrier mask(0x00000000)
+; GFX9-NEXT: ; sched_barrier mask(0x00000000)
; GFX9-NEXT: s_setprio 1 ; encoding: [0x01,0x00,0x8f,0xbf]
+; GFX9-NEXT: ; sched_barrier mask(0x00000000)
+; GFX9-NEXT: ; sched_barrier mask(0x00000000)
; GFX9-NEXT: s_setprio 2 ; encoding: [0x02,0x00,0x8f,0xbf]
+; GFX9-NEXT: ; sched_barrier mask(0x00000000)
+; GFX9-NEXT: ; sched_barrier mask(0x00000000)
; GFX9-NEXT: s_setprio 3 ; encoding: [0x03,0x00,0x8f,0xbf]
+; GFX9-NEXT: ; sched_barrier mask(0x00000000)
+; GFX9-NEXT: ; sched_barrier mask(0x00000000)
; GFX9-NEXT: s_setprio 10 ; encoding: [0x0a,0x00,0x8f,0xbf]
+; GFX9-NEXT: ; sched_barrier mask(0x00000000)
+; GFX9-NEXT: ; sched_barrier mask(0x00000000)
; GFX9-NEXT: s_setprio -1 ; encoding: [0xff,0xff,0x8f,0xbf]
+; GFX9-NEXT: ; sched_barrier mask(0x00000000)
+; GFX9-NEXT: ; sched_barrier mask(0x00000000)
; GFX9-NEXT: s_setprio 0 ; encoding: [0x00,0x00,0x8f,0xbf]
+; GFX9-NEXT: ; sched_barrier mask(0x00000000)
+; GFX9-NEXT: ; sched_barrier mask(0x00000000)
; GFX9-NEXT: s_setprio 1 ; encoding: [0x01,0x00,0x8f,0xbf]
+; GFX9-NEXT: ; sched_barrier mask(0x00000000)
+; GFX9-NEXT: ; sched_barrier mask(0x00000000)
; GFX9-NEXT: s_setprio -1 ; encoding: [0xff,0xff,0x8f,0xbf]
+; GFX9-NEXT: ; sched_barrier mask(0x00000000)
; GFX9-NEXT: s_setpc_b64 s[30:31] ; encoding: [0x1e,0x1d,0x80,0xbe]
;
; SI-LABEL: test_llvm_amdgcn_s_setprio:
; SI: ; %bb.0:
+; SI-NEXT: ; sched_barrier mask(0x00000000)
; SI-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0) ; encoding: [0x00,0x00,0x8c,0xbf]
; SI-NEXT: s_setprio 0 ; encoding: [0x00,0x00,0x8f,0xbf]
+; SI-NEXT: ; sched_barrier mask(0x00000000)
+; SI-NEXT: ; sched_barrier mask(0x00000000)
; SI-NEXT: s_setprio 1 ; encoding: [0x01,0x00,0x8f,0xbf]
+; SI-NEXT: ; sched_barrier mask(0x00000000)
+; SI-NEXT: ; sched_barrier mask(0x00000000)
; SI-NEXT: s_setprio 2 ; encoding: [0x02,0x00,0x8f,0xbf]
+; SI-NEXT: ; sched_barrier mask(0x00000000)
+; SI-NEXT: ; sched_barrier mask(0x00000000)
; SI-NEXT: s_setprio 3 ; encoding: [0x03,0x00,0x8f,0xbf]
+; SI-NEXT: ; sched_barrier mask(0x00000000)
+; SI-NEXT: ; sched_barrier mask(0x00000000)
; SI-NEXT: s_setprio 10 ; encoding: [0x0a,0x00,0x8f,0xbf]
+; SI-NEXT: ; sched_barrier mask(0x00000000)
+; SI-NEXT: ; sched_barrier mask(0x00000000)
; SI-NEXT: s_setprio -1 ; encoding: [0xff,0xff,0x8f,0xbf]
+; SI-NEXT: ; sched_barrier mask(0x00000000)
+; SI-NEXT: ; sched_barrier mask(0x00000000)
; SI-NEXT: s_setprio 0 ; encoding: [0x00,0x00,0x8f,0xbf]
+; SI-NEXT: ; sched_barrier mask(0x00000000)
+; SI-NEXT: ; sched_barrier mask(0x00000000)
; SI-NEXT: s_setprio 1 ; encoding: [0x01,0x00,0x8f,0xbf]
+; SI-NEXT: ; sched_barrier mask(0x00000000)
+; SI-NEXT: ; sched_barrier mask(0x00000000)
; SI-NEXT: s_setprio -1 ; encoding: [0xff,0xff,0x8f,0xbf]
+; SI-NEXT: ; sched_barrier mask(0x00000000)
; SI-NEXT: s_setpc_b64 s[30:31] ; encoding: [0x1e,0x20,0x80,0xbe]
call void @llvm.amdgcn.s.setprio(i16 0)
call void @llvm.amdgcn.s.setprio(i16 1)
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.setprio.relaxed.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.setprio.relaxed.ll
new file mode 100644
index 0000000000000..4701d4ad8dcff
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.setprio.relaxed.ll
@@ -0,0 +1,49 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py
+; RUN: llc -mtriple=amdgpu9.0a -show-mc-encoding < %s | FileCheck -check-prefix=GFX9 %s
+; RUN: llc -mtriple=amdgpu6.00 -show-mc-encoding < %s | FileCheck -check-prefix=SI %s
+; RUN: llc -global-isel -mtriple=amdgpu6.00 -show-mc-encoding < %s | FileCheck -check-prefix=SI %s
+; RUN: llc -global-isel -mtriple=amdgpu9.0a -show-mc-encoding < %s | FileCheck -check-prefix=GFX9 %s
+
+declare void @llvm.amdgcn.s.setprio.relaxed(i16) #0
+
+define void @test_llvm_amdgcn_s_setprio.relaxed() #0 {
+; GFX9-LABEL: test_llvm_amdgcn_s_setprio.relaxed:
+; GFX9: ; %bb.0:
+; GFX9-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0) ; encoding: [0x00,0x00,0x8c,0xbf]
+; GFX9-NEXT: s_setprio 0 ; encoding: [0x00,0x00,0x8f,0xbf]
+; GFX9-NEXT: s_setprio 1 ; encoding: [0x01,0x00,0x8f,0xbf]
+; GFX9-NEXT: s_setprio 2 ; encoding: [0x02,0x00,0x8f,0xbf]
+; GFX9-NEXT: s_setprio 3 ; encoding: [0x03,0x00,0x8f,0xbf]
+; GFX9-NEXT: s_setprio 10 ; encoding: [0x0a,0x00,0x8f,0xbf]
+; GFX9-NEXT: s_setprio -1 ; encoding: [0xff,0xff,0x8f,0xbf]
+; GFX9-NEXT: s_setprio 0 ; encoding: [0x00,0x00,0x8f,0xbf]
+; GFX9-NEXT: s_setprio 1 ; encoding: [0x01,0x00,0x8f,0xbf]
+; GFX9-NEXT: s_setprio -1 ; encoding: [0xff,0xff,0x8f,0xbf]
+; GFX9-NEXT: s_setpc_b64 s[30:31] ; encoding: [0x1e,0x1d,0x80,0xbe]
+;
+; SI-LABEL: test_llvm_amdgcn_s_setprio.relaxed:
+; SI: ; %bb.0:
+; SI-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0) ; encoding: [0x00,0x00,0x8c,0xbf]
+; SI-NEXT: s_setprio 0 ; encoding: [0x00,0x00,0x8f,0xbf]
+; SI-NEXT: s_setprio 1 ; encoding: [0x01,0x00,0x8f,0xbf]
+; SI-NEXT: s_setprio 2 ; encoding: [0x02,0x00,0x8f,0xbf]
+; SI-NEXT: s_setprio 3 ; encoding: [0x03,0x00,0x8f,0xbf]
+; SI-NEXT: s_setprio 10 ; encoding: [0x0a,0x00,0x8f,0xbf]
+; SI-NEXT: s_setprio -1 ; encoding: [0xff,0xff,0x8f,0xbf]
+; SI-NEXT: s_setprio 0 ; encoding: [0x00,0x00,0x8f,0xbf]
+; SI-NEXT: s_setprio 1 ; encoding: [0x01,0x00,0x8f,0xbf]
+; SI-NEXT: s_setprio -1 ; encoding: [0xff,0xff,0x8f,0xbf]
+; SI-NEXT: s_setpc_b64 s[30:31] ; encoding: [0x1e,0x20,0x80,0xbe]
+ call void @llvm.amdgcn.s.setprio.relaxed(i16 0)
+ call void @llvm.amdgcn.s.setprio.relaxed(i16 1)
+ call void @llvm.amdgcn.s.setprio.relaxed(i16 2)
+ call void @llvm.amdgcn.s.setprio.relaxed(i16 3)
+ call void @llvm.amdgcn.s.setprio.relaxed(i16 10)
+ call void @llvm.amdgcn.s.setprio.relaxed(i16 65535)
+ call void @llvm.amdgcn.s.setprio.relaxed(i16 65536)
+ call void @llvm.amdgcn.s.setprio.relaxed(i16 65537)
+ call void @llvm.amdgcn.s.setprio.relaxed(i16 -1)
+ ret void
+}
+
+attributes #0 = { nounwind }
diff --git a/llvm/test/CodeGen/AMDGPU/required-export-priority.ll b/llvm/test/CodeGen/AMDGPU/required-export-priority.ll
index b6e2e75c9d28c..5d464a4fd314d 100644
--- a/llvm/test/CodeGen/AMDGPU/required-export-priority.ll
+++ b/llvm/test/CodeGen/AMDGPU/required-export-priority.ll
@@ -445,7 +445,9 @@ define amdgpu_ps void @test_export_in_callee_prio(float %v) #0 {
; GFX11: ; %bb.0:
; GFX11-NEXT: s_mov_b32 s32, 0
; GFX11-NEXT: v_add_f32_e32 v0, 1.0, v0
+; GFX11-NEXT: ; sched_barrier mask(0x00000000)
; GFX11-NEXT: s_setprio 0
+; GFX11-NEXT: ; sched_barrier mask(0x00000000)
; GFX11-NEXT: s_getpc_b64 s[0:1]
; GFX11-NEXT: s_add_u32 s0, s0, test_export_gfx at gotpcrel32@lo+4
; GFX11-NEXT: s_addc_u32 s1, s1, test_export_gfx at gotpcrel32@hi+12
@@ -459,7 +461,9 @@ define amdgpu_ps void @test_export_in_callee_prio(float %v) #0 {
; GFX1150-NEXT: s_setprio 2
; GFX1150-NEXT: s_mov_b32 s32, 0
; GFX1150-NEXT: v_add_f32_e32 v0, 1.0, v0
+; GFX1150-NEXT: ; sched_barrier mask(0x00000000)
; GFX1150-NEXT: s_setprio 2
+; GFX1150-NEXT: ; sched_barrier mask(0x00000000)
; GFX1150-NEXT: s_getpc_b64 s[0:1]
; GFX1150-NEXT: s_add_u32 s0, s0, test_export_gfx at gotpcrel32@lo+4
; GFX1150-NEXT: s_addc_u32 s1, s1, test_export_gfx at gotpcrel32@hi+12
diff --git a/llvm/test/CodeGen/AMDGPU/sched-setprio-relaxed.ll b/llvm/test/CodeGen/AMDGPU/sched-setprio-relaxed.ll
new file mode 100644
index 0000000000000..63a51aa4420c8
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/sched-setprio-relaxed.ll
@@ -0,0 +1,134 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
+; RUN: llc -mtriple=amdgpu9.08-amd-amdhsa < %s | FileCheck --check-prefix=GCN %s
+
+; Test that llvm.amdgcn.s.setprio.relaxed allows instructions to be scheduled across it,
+; unlike llvm.amdgcn.s.setprio which is as a scheduling boundary.
+
+declare void @llvm.amdgcn.s.setprio.relaxed(i16)
+declare float @llvm.amdgcn.rcp.f32(float)
+
+define amdgpu_kernel void @test_relaxed_allows_salu(ptr addrspace(1) %out, i32 inreg %x, i32 inreg %y) {
+; GCN-LABEL: test_relaxed_allows_salu:
+; GCN: ; %bb.0:
+; GCN-NEXT: s_load_dwordx4 s[0:3], s[8:9], 0x0
+; GCN-NEXT: v_mov_b32_e32 v0, 0
+; GCN-NEXT: s_setprio 1
+; GCN-NEXT: s_waitcnt lgkmcnt(0)
+; GCN-NEXT: s_add_i32 s2, s2, s3
+; GCN-NEXT: s_add_i32 s3, s3, s3
+; GCN-NEXT: s_add_i32 s2, s2, s3
+; GCN-NEXT: v_mov_b32_e32 v1, s2
+; GCN-NEXT: global_store_dword v0, v1, s[0:1]
+; GCN-NEXT: s_endpgm
+ %add1 = add i32 %x, %y
+ call void @llvm.amdgcn.s.setprio.relaxed(i16 1)
+ %add2 = add i32 %y, %y
+ %sum = add i32 %add1, %add2
+ store i32 %sum, ptr addrspace(1) %out
+ ret void
+}
+
+define amdgpu_kernel void @test_relaxed_allows_valu(ptr addrspace(1) %out, float %x, float %y) {
+; GCN-LABEL: test_relaxed_allows_valu:
+; GCN: ; %bb.0:
+; GCN-NEXT: s_load_dwordx4 s[0:3], s[8:9], 0x0
+; GCN-NEXT: v_mov_b32_e32 v0, 0
+; GCN-NEXT: s_setprio 1
+; GCN-NEXT: s_waitcnt lgkmcnt(0)
+; GCN-NEXT: v_add_f32_e64 v1, s2, 1.0
+; GCN-NEXT: v_add_f32_e64 v2, s3, 2.0
+; GCN-NEXT: v_add_f32_e32 v1, v1, v2
+; GCN-NEXT: global_store_dword v0, v1, s[0:1]
+; GCN-NEXT: s_endpgm
+ %add1 = fadd float %x, 1.0
+ call void @llvm.amdgcn.s.setprio.relaxed(i16 1)
+ %add2 = fadd float %y, 2.0
+ %sum = fadd float %add1, %add2
+ store float %sum, ptr addrspace(1) %out
+ ret void
+}
+
+define amdgpu_kernel void @test_relaxed_allows_trans(ptr addrspace(1) %out, float %x, float %y) {
+; GCN-LABEL: test_relaxed_allows_trans:
+; GCN: ; %bb.0:
+; GCN-NEXT: s_load_dwordx4 s[0:3], s[8:9], 0x0
+; GCN-NEXT: v_mov_b32_e32 v2, 0
+; GCN-NEXT: s_setprio 1
+; GCN-NEXT: s_waitcnt lgkmcnt(0)
+; GCN-NEXT: v_rcp_f32_e32 v0, s2
+; GCN-NEXT: v_rcp_f32_e32 v1, s3
+; GCN-NEXT: v_add_f32_e32 v0, v0, v1
+; GCN-NEXT: global_store_dword v2, v0, s[0:1]
+; GCN-NEXT: s_endpgm
+ %rcp1 = call float @llvm.amdgcn.rcp.f32(float %x)
+ call void @llvm.amdgcn.s.setprio.relaxed(i16 1)
+ %rcp2 = call float @llvm.amdgcn.rcp.f32(float %y)
+ %sum = fadd float %rcp1, %rcp2
+ store float %sum, ptr addrspace(1) %out
+ ret void
+}
+
+define amdgpu_kernel void @test_relaxed_allows_mfma(ptr addrspace(1) %out, <4 x float> %in) {
+; GCN-LABEL: test_relaxed_allows_mfma:
+; GCN: ; %bb.0:
+; GCN-NEXT: s_load_dwordx2 s[4:5], s[8:9], 0x0
+; GCN-NEXT: s_load_dwordx4 s[0:3], s[8:9], 0x10
+; GCN-NEXT: v_mov_b32_e32 v0, 1.0
+; GCN-NEXT: v_mov_b32_e32 v2, 0x40400000
+; GCN-NEXT: v_mov_b32_e32 v1, 2.0
+; GCN-NEXT: v_mov_b32_e32 v4, 0
+; GCN-NEXT: s_waitcnt lgkmcnt(0)
+; GCN-NEXT: v_mov_b32_e32 v5, s0
+; GCN-NEXT: v_mov_b32_e32 v3, s1
+; GCN-NEXT: v_mov_b32_e32 v6, s2
+; GCN-NEXT: v_accvgpr_write_b32 a0, v5
+; GCN-NEXT: v_mov_b32_e32 v5, s3
+; GCN-NEXT: v_accvgpr_write_b32 a1, v3
+; GCN-NEXT: v_accvgpr_write_b32 a2, v6
+; GCN-NEXT: v_accvgpr_write_b32 a3, v5
+; GCN-NEXT: s_setprio 1
+; GCN-NEXT: v_mfma_f32_4x4x1f32 a[0:3], v0, v1, a[0:3]
+; GCN-NEXT: v_mov_b32_e32 v0, 4.0
+; GCN-NEXT: s_nop 1
+; GCN-NEXT: v_mfma_f32_4x4x1f32 a[0:3], v2, v0, a[0:3]
+; GCN-NEXT: s_nop 3
+; GCN-NEXT: v_accvgpr_read_b32 v0, a0
+; GCN-NEXT: v_accvgpr_read_b32 v1, a1
+; GCN-NEXT: v_accvgpr_read_b32 v2, a2
+; GCN-NEXT: v_accvgpr_read_b32 v3, a3
+; GCN-NEXT: s_nop 1
+; GCN-NEXT: global_store_dwordx4 v4, v[0:3], s[4:5]
+; GCN-NEXT: s_endpgm
+ %mfma1 = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float 1.0, float 2.0, <4 x float> %in, i32 0, i32 0, i32 0)
+ call void @llvm.amdgcn.s.setprio.relaxed(i16 1)
+ %mfma2 = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float 3.0, float 4.0, <4 x float> %mfma1, i32 0, i32 0, i32 0)
+ store <4 x float> %mfma2, ptr addrspace(1) %out
+ ret void
+}
+
+declare <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float, float, <4 x float>, i32 immarg, i32 immarg, i32 immarg)
+
+; Memory operations cannot move across s_setprio due to its side effects.
+define amdgpu_kernel void @test_relaxed_blocks_memory(ptr addrspace(1) %out1, ptr addrspace(1) %out2, ptr addrspace(1) %out3, float %x) {
+; GCN-LABEL: test_relaxed_blocks_memory:
+; GCN: ; %bb.0:
+; GCN-NEXT: s_load_dwordx4 s[0:3], s[8:9], 0x0
+; GCN-NEXT: s_load_dword s6, s[8:9], 0x18
+; GCN-NEXT: v_mov_b32_e32 v0, 0
+; GCN-NEXT: v_mov_b32_e32 v1, 1.0
+; GCN-NEXT: s_load_dwordx2 s[4:5], s[8:9], 0x10
+; GCN-NEXT: s_waitcnt lgkmcnt(0)
+; GCN-NEXT: global_store_dword v0, v1, s[0:1]
+; GCN-NEXT: v_rcp_f32_e32 v1, s6
+; GCN-NEXT: v_mov_b32_e32 v2, 2.0
+; GCN-NEXT: s_setprio 1
+; GCN-NEXT: global_store_dword v0, v2, s[2:3]
+; GCN-NEXT: global_store_dword v0, v1, s[4:5]
+; GCN-NEXT: s_endpgm
+ store float 1.0, ptr addrspace(1) %out1
+ %rcp = call float @llvm.amdgcn.rcp.f32(float %x)
+ call void @llvm.amdgcn.s.setprio.relaxed(i16 1)
+ store float 2.0, ptr addrspace(1) %out2
+ store float %rcp, ptr addrspace(1) %out3
+ ret void
+}
diff --git a/llvm/test/CodeGen/AMDGPU/sched-setprio.ll b/llvm/test/CodeGen/AMDGPU/sched-setprio.ll
index c6b83fd715782..3cb868734dc5b 100644
--- a/llvm/test/CodeGen/AMDGPU/sched-setprio.ll
+++ b/llvm/test/CodeGen/AMDGPU/sched-setprio.ll
@@ -4,7 +4,9 @@ declare void @llvm.amdgcn.s.setprio(i16)
declare <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float, float, <4 x float>, i32, i32, i32)
; GCN-LABEL: {{^}}test_mfma_f32_4x4x1f32:
-; GCN: s_setprio 1
+; GCN: sched_barrier mask(0x00000000)
+; GCN-NEXT: s_setprio 1
+; GCN-NEXT: sched_barrier mask(0x00000000)
; GCN: v_mfma
; GCN: v_mfma
; GCN: s_setprio 0
>From eadb83ff0881b275a2edfdeaa58a180cffcaeb20 Mon Sep 17 00:00:00 2001
From: Frederik Harwath <fharwath at amd.com>
Date: Thu, 27 Aug 2026 09:57:12 -0400
Subject: [PATCH 09/10] Revert "[AMDGPU] Introduce
llvm.amdgcn.s.setprio.relaxed"
This reverts commit c79c5a95ca8b00c694e3e84fac9b372ac144728c.
---
clang/include/clang/Basic/BuiltinsAMDGPU.td | 1 -
clang/test/CodeGenOpenCL/builtins-amdgcn.cl | 9 --
.../test/SemaOpenCL/builtins-amdgcn-error.cl | 6 -
llvm/docs/AMDGPUUsage.rst | 6 -
llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 6 -
.../AMDGPU/AMDGPURegBankLegalizeRules.cpp | 1 -
llvm/lib/Target/AMDGPU/SIISelLowering.cpp | 8 --
llvm/lib/Target/AMDGPU/SIInstrInfo.cpp | 1 +
llvm/lib/Target/AMDGPU/SIInstructions.td | 10 --
llvm/lib/Target/AMDGPU/SOPInstructions.td | 2 +-
.../CodeGen/AMDGPU/llvm.amdgcn.s.setprio.ll | 36 -----
.../AMDGPU/llvm.amdgcn.s.setprio.relaxed.ll | 49 -------
.../AMDGPU/required-export-priority.ll | 4 -
.../CodeGen/AMDGPU/sched-setprio-relaxed.ll | 134 ------------------
llvm/test/CodeGen/AMDGPU/sched-setprio.ll | 4 +-
15 files changed, 3 insertions(+), 274 deletions(-)
delete mode 100644 llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.setprio.relaxed.ll
delete mode 100644 llvm/test/CodeGen/AMDGPU/sched-setprio-relaxed.ll
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.td b/clang/include/clang/Basic/BuiltinsAMDGPU.td
index bd3f7ae08eb0c..73b27a2b5c5cd 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.td
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.td
@@ -208,7 +208,6 @@ def __builtin_amdgcn_s_sleep : AMDGPUBuiltin<"void(_Constant int)">;
def __builtin_amdgcn_s_incperflevel : AMDGPUBuiltin<"void(_Constant int)">;
def __builtin_amdgcn_s_decperflevel : AMDGPUBuiltin<"void(_Constant int)">;
def __builtin_amdgcn_s_setprio : AMDGPUBuiltin<"void(_Constant short)">;
-def __builtin_amdgcn_s_setprio_relaxed : AMDGPUBuiltin<"void(_Constant short)">;
def __builtin_amdgcn_ds_swizzle : AMDGPUBuiltin<"int(int, _Constant int)", [Const]>;
def __builtin_amdgcn_ds_permute : AMDGPUBuiltin<"int(int, int)", [Const]>;
def __builtin_amdgcn_ds_bpermute : AMDGPUBuiltin<"int(int, int)", [Const]>;
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
index 1375806125034..28c420a5760f4 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
@@ -1041,15 +1041,6 @@ void test_s_setprio()
__builtin_amdgcn_s_setprio(3);
}
-// CHECK-LABEL: @test_s_setprio_relaxed
-// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.s.setprio.relaxed(i16 0)
-// CHECK: {{.*}}call{{.*}} void @llvm.amdgcn.s.setprio.relaxed(i16 3)
-void test_s_setprio_relaxed()
-{
- __builtin_amdgcn_s_setprio_relaxed(0);
- __builtin_amdgcn_s_setprio_relaxed(3);
-}
-
// CHECK-LABEL: @test_read_exec(
// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.ballot.i64(i1 true)
void test_read_exec(global ulong* out) {
diff --git a/clang/test/SemaOpenCL/builtins-amdgcn-error.cl b/clang/test/SemaOpenCL/builtins-amdgcn-error.cl
index 536de3704274f..10f00dd7da5f9 100644
--- a/clang/test/SemaOpenCL/builtins-amdgcn-error.cl
+++ b/clang/test/SemaOpenCL/builtins-amdgcn-error.cl
@@ -60,12 +60,6 @@ void test_s_setprio(int x)
__builtin_amdgcn_s_setprio(65536); // expected-warning {{implicit conversion from 'int' to 'short' changes value from 65536 to 0}}
}
-void test_s_setprio_relaxed(int x)
-{
- __builtin_amdgcn_s_setprio_relaxed(x); // expected-error {{argument to '__builtin_amdgcn_s_setprio_relaxed' must be a constant integer}}
- __builtin_amdgcn_s_setprio_relaxed(65536); // expected-warning {{implicit conversion from 'int' to 'short' changes value from 65536 to 0}}
-}
-
void test_sched_barrier(int x)
{
__builtin_amdgcn_sched_barrier(x); // expected-error {{argument to '__builtin_amdgcn_sched_barrier' must be a constant integer}}
diff --git a/llvm/docs/AMDGPUUsage.rst b/llvm/docs/AMDGPUUsage.rst
index 0d1a551e0725a..443dd7f9772d7 100644
--- a/llvm/docs/AMDGPUUsage.rst
+++ b/llvm/docs/AMDGPUUsage.rst
@@ -1954,12 +1954,6 @@ The AMDGPU backend implements the following LLVM IR intrinsics.
The iglp_opt strategy implementations are subject to change.
- llvm.amdgcn.s.setprio Provide access to the ``s_setprio`` instruction for setting the wave priority to the given value.
- The intrinsic ensures that the instruction acts as a full barrier for instruction scheduling, i.e.
- no instruction may be scheduled across it.
-
- llvm.amdgcn.s.setprio.relaxed Like ``llvm.amdgcn.s.setprio``, but does not act as a scheduling barrier.
-
llvm.amdgcn.s.getpc Provides access to the s_getpc_b64 instruction, but with the return value
sign-extended from the width of the underlying PC hardware register even on
processors where the s_getpc_b64 instruction returns a zero-extended value.
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index 2fad48178e436..565637b36131c 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -2229,12 +2229,6 @@ def int_amdgcn_s_setprio :
DefaultAttrsIntrinsic<[], [llvm_i16_ty], [ImmArg<ArgIndex<0>>, IntrNoMem,
IntrHasSideEffects]>;
-// Like int_amdgcn_s_setprio but without scheduling boundary behavior.
-def int_amdgcn_s_setprio_relaxed :
- ClangBuiltin<"__builtin_amdgcn_s_setprio_relaxed">,
- DefaultAttrsIntrinsic<[], [llvm_i16_ty], [ImmArg<ArgIndex<0>>, IntrNoMem,
- IntrHasSideEffects]>;
-
def int_amdgcn_s_setprio_inc_wg :
ClangBuiltin<"__builtin_amdgcn_s_setprio_inc_wg">,
DefaultAttrsIntrinsic<[], [llvm_i16_ty], [ImmArg<ArgIndex<0>>, IntrNoMem,
diff --git a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp
index dca4617ece0a7..881afaf920037 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp
@@ -1822,7 +1822,6 @@ RegBankLegalizeRules::RegBankLegalizeRules(const GCNSubtarget &_ST,
amdgcn_s_sethalt,
amdgcn_s_setprio,
amdgcn_s_setprio_inc_wg,
- amdgcn_s_setprio_relaxed,
amdgcn_s_sleep,
amdgcn_s_ttracedata_imm,
amdgcn_s_wait_asynccnt,
diff --git a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
index a1bd173af3ed8..78bbb7f2d6146 100644
--- a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
+++ b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp
@@ -7361,14 +7361,6 @@ SITargetLowering::EmitInstrWithCustomInserter(MachineInstr &MI,
MRI.setSimpleHint(MI.getOperand(0).getReg(), MI.getOperand(6).getReg());
return BB;
}
- case AMDGPU::S_SETPRIO_BARRIER: {
- int64_t Priority = MI.getOperand(0).getImm();
- BuildMI(*BB, MI, DL, TII->get(AMDGPU::SCHED_BARRIER)).addImm(0);
- BuildMI(*BB, MI, DL, TII->get(AMDGPU::S_SETPRIO)).addImm(Priority);
- BuildMI(*BB, MI, DL, TII->get(AMDGPU::SCHED_BARRIER)).addImm(0);
- MI.eraseFromParent();
- return BB;
- }
default:
if (TII->isImage(MI) || TII->isMUBUF(MI)) {
if (!MI.mayStore())
diff --git a/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp b/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp
index 793c30d8369cb..b2c3772cb0041 100644
--- a/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp
@@ -4489,6 +4489,7 @@ bool SIInstrInfo::isSchedulingBoundary(const MachineInstr &MI,
return MI.modifiesRegister(AMDGPU::EXEC, &RI) ||
MI.getOpcode() == AMDGPU::S_SETREG_IMM32_B32 ||
MI.getOpcode() == AMDGPU::S_SETREG_B32 ||
+ MI.getOpcode() == AMDGPU::S_SETPRIO ||
MI.getOpcode() == AMDGPU::S_SETPRIO_INC_WG ||
changesVGPRIndexingMode(MI);
}
diff --git a/llvm/lib/Target/AMDGPU/SIInstructions.td b/llvm/lib/Target/AMDGPU/SIInstructions.td
index 7a897a20ed707..39df5dfd4d01d 100644
--- a/llvm/lib/Target/AMDGPU/SIInstructions.td
+++ b/llvm/lib/Target/AMDGPU/SIInstructions.td
@@ -553,16 +553,6 @@ def SCHED_BARRIER : SPseudoInstSI<(outs), (ins i32imm:$mask),
let isMeta = 1;
}
-// Expands to SCHED_BARRIER 0 + S_SETPRIO + SCHED_BARRIER 0 to provide
-// backwards compatible behavior for setprio intrinsic since S_SETPRIO
-// used to be a scheduling boundary. It no longer is to allow for more
-// control over the scheduling behavior.
-def S_SETPRIO_BARRIER : SPseudoInstSI<(outs), (ins i16imm:$priority),
- [(int_amdgcn_s_setprio timm:$priority)]> {
- let usesCustomInserter = 1;
- let hasSideEffects = 1;
-}
-
def SCHED_GROUP_BARRIER : SPseudoInstSI<
(outs),
(ins i32imm:$mask, i32imm:$size, i32imm:$syncid),
diff --git a/llvm/lib/Target/AMDGPU/SOPInstructions.td b/llvm/lib/Target/AMDGPU/SOPInstructions.td
index 118091210d33d..8fa54de570e67 100644
--- a/llvm/lib/Target/AMDGPU/SOPInstructions.td
+++ b/llvm/lib/Target/AMDGPU/SOPInstructions.td
@@ -1755,7 +1755,7 @@ def S_SLEEP_VAR : SOP1_0_32 <"s_sleep_var", [(int_amdgcn_s_sleep_var SSrc_b32:$s
}
def S_SETPRIO : SOPP_Pseudo <"s_setprio", (ins i16imm:$simm16), "$simm16",
- [(int_amdgcn_s_setprio_relaxed timm:$simm16)]> {
+ [(int_amdgcn_s_setprio timm:$simm16)]> {
}
def S_SETPRIO_INC_WG : SOPP_Pseudo <"s_setprio_inc_wg", (ins i16imm:$simm16), "$simm16",
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.setprio.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.setprio.ll
index 2604361f1bd06..87df1e152bef5 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.setprio.ll
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.setprio.ll
@@ -9,66 +9,30 @@ declare void @llvm.amdgcn.s.setprio(i16) #0
define void @test_llvm_amdgcn_s_setprio() #0 {
; GFX9-LABEL: test_llvm_amdgcn_s_setprio:
; GFX9: ; %bb.0:
-; GFX9-NEXT: ; sched_barrier mask(0x00000000)
; GFX9-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0) ; encoding: [0x00,0x00,0x8c,0xbf]
; GFX9-NEXT: s_setprio 0 ; encoding: [0x00,0x00,0x8f,0xbf]
-; GFX9-NEXT: ; sched_barrier mask(0x00000000)
-; GFX9-NEXT: ; sched_barrier mask(0x00000000)
; GFX9-NEXT: s_setprio 1 ; encoding: [0x01,0x00,0x8f,0xbf]
-; GFX9-NEXT: ; sched_barrier mask(0x00000000)
-; GFX9-NEXT: ; sched_barrier mask(0x00000000)
; GFX9-NEXT: s_setprio 2 ; encoding: [0x02,0x00,0x8f,0xbf]
-; GFX9-NEXT: ; sched_barrier mask(0x00000000)
-; GFX9-NEXT: ; sched_barrier mask(0x00000000)
; GFX9-NEXT: s_setprio 3 ; encoding: [0x03,0x00,0x8f,0xbf]
-; GFX9-NEXT: ; sched_barrier mask(0x00000000)
-; GFX9-NEXT: ; sched_barrier mask(0x00000000)
; GFX9-NEXT: s_setprio 10 ; encoding: [0x0a,0x00,0x8f,0xbf]
-; GFX9-NEXT: ; sched_barrier mask(0x00000000)
-; GFX9-NEXT: ; sched_barrier mask(0x00000000)
; GFX9-NEXT: s_setprio -1 ; encoding: [0xff,0xff,0x8f,0xbf]
-; GFX9-NEXT: ; sched_barrier mask(0x00000000)
-; GFX9-NEXT: ; sched_barrier mask(0x00000000)
; GFX9-NEXT: s_setprio 0 ; encoding: [0x00,0x00,0x8f,0xbf]
-; GFX9-NEXT: ; sched_barrier mask(0x00000000)
-; GFX9-NEXT: ; sched_barrier mask(0x00000000)
; GFX9-NEXT: s_setprio 1 ; encoding: [0x01,0x00,0x8f,0xbf]
-; GFX9-NEXT: ; sched_barrier mask(0x00000000)
-; GFX9-NEXT: ; sched_barrier mask(0x00000000)
; GFX9-NEXT: s_setprio -1 ; encoding: [0xff,0xff,0x8f,0xbf]
-; GFX9-NEXT: ; sched_barrier mask(0x00000000)
; GFX9-NEXT: s_setpc_b64 s[30:31] ; encoding: [0x1e,0x1d,0x80,0xbe]
;
; SI-LABEL: test_llvm_amdgcn_s_setprio:
; SI: ; %bb.0:
-; SI-NEXT: ; sched_barrier mask(0x00000000)
; SI-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0) ; encoding: [0x00,0x00,0x8c,0xbf]
; SI-NEXT: s_setprio 0 ; encoding: [0x00,0x00,0x8f,0xbf]
-; SI-NEXT: ; sched_barrier mask(0x00000000)
-; SI-NEXT: ; sched_barrier mask(0x00000000)
; SI-NEXT: s_setprio 1 ; encoding: [0x01,0x00,0x8f,0xbf]
-; SI-NEXT: ; sched_barrier mask(0x00000000)
-; SI-NEXT: ; sched_barrier mask(0x00000000)
; SI-NEXT: s_setprio 2 ; encoding: [0x02,0x00,0x8f,0xbf]
-; SI-NEXT: ; sched_barrier mask(0x00000000)
-; SI-NEXT: ; sched_barrier mask(0x00000000)
; SI-NEXT: s_setprio 3 ; encoding: [0x03,0x00,0x8f,0xbf]
-; SI-NEXT: ; sched_barrier mask(0x00000000)
-; SI-NEXT: ; sched_barrier mask(0x00000000)
; SI-NEXT: s_setprio 10 ; encoding: [0x0a,0x00,0x8f,0xbf]
-; SI-NEXT: ; sched_barrier mask(0x00000000)
-; SI-NEXT: ; sched_barrier mask(0x00000000)
; SI-NEXT: s_setprio -1 ; encoding: [0xff,0xff,0x8f,0xbf]
-; SI-NEXT: ; sched_barrier mask(0x00000000)
-; SI-NEXT: ; sched_barrier mask(0x00000000)
; SI-NEXT: s_setprio 0 ; encoding: [0x00,0x00,0x8f,0xbf]
-; SI-NEXT: ; sched_barrier mask(0x00000000)
-; SI-NEXT: ; sched_barrier mask(0x00000000)
; SI-NEXT: s_setprio 1 ; encoding: [0x01,0x00,0x8f,0xbf]
-; SI-NEXT: ; sched_barrier mask(0x00000000)
-; SI-NEXT: ; sched_barrier mask(0x00000000)
; SI-NEXT: s_setprio -1 ; encoding: [0xff,0xff,0x8f,0xbf]
-; SI-NEXT: ; sched_barrier mask(0x00000000)
; SI-NEXT: s_setpc_b64 s[30:31] ; encoding: [0x1e,0x20,0x80,0xbe]
call void @llvm.amdgcn.s.setprio(i16 0)
call void @llvm.amdgcn.s.setprio(i16 1)
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.setprio.relaxed.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.setprio.relaxed.ll
deleted file mode 100644
index 4701d4ad8dcff..0000000000000
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.setprio.relaxed.ll
+++ /dev/null
@@ -1,49 +0,0 @@
-; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py
-; RUN: llc -mtriple=amdgpu9.0a -show-mc-encoding < %s | FileCheck -check-prefix=GFX9 %s
-; RUN: llc -mtriple=amdgpu6.00 -show-mc-encoding < %s | FileCheck -check-prefix=SI %s
-; RUN: llc -global-isel -mtriple=amdgpu6.00 -show-mc-encoding < %s | FileCheck -check-prefix=SI %s
-; RUN: llc -global-isel -mtriple=amdgpu9.0a -show-mc-encoding < %s | FileCheck -check-prefix=GFX9 %s
-
-declare void @llvm.amdgcn.s.setprio.relaxed(i16) #0
-
-define void @test_llvm_amdgcn_s_setprio.relaxed() #0 {
-; GFX9-LABEL: test_llvm_amdgcn_s_setprio.relaxed:
-; GFX9: ; %bb.0:
-; GFX9-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0) ; encoding: [0x00,0x00,0x8c,0xbf]
-; GFX9-NEXT: s_setprio 0 ; encoding: [0x00,0x00,0x8f,0xbf]
-; GFX9-NEXT: s_setprio 1 ; encoding: [0x01,0x00,0x8f,0xbf]
-; GFX9-NEXT: s_setprio 2 ; encoding: [0x02,0x00,0x8f,0xbf]
-; GFX9-NEXT: s_setprio 3 ; encoding: [0x03,0x00,0x8f,0xbf]
-; GFX9-NEXT: s_setprio 10 ; encoding: [0x0a,0x00,0x8f,0xbf]
-; GFX9-NEXT: s_setprio -1 ; encoding: [0xff,0xff,0x8f,0xbf]
-; GFX9-NEXT: s_setprio 0 ; encoding: [0x00,0x00,0x8f,0xbf]
-; GFX9-NEXT: s_setprio 1 ; encoding: [0x01,0x00,0x8f,0xbf]
-; GFX9-NEXT: s_setprio -1 ; encoding: [0xff,0xff,0x8f,0xbf]
-; GFX9-NEXT: s_setpc_b64 s[30:31] ; encoding: [0x1e,0x1d,0x80,0xbe]
-;
-; SI-LABEL: test_llvm_amdgcn_s_setprio.relaxed:
-; SI: ; %bb.0:
-; SI-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0) ; encoding: [0x00,0x00,0x8c,0xbf]
-; SI-NEXT: s_setprio 0 ; encoding: [0x00,0x00,0x8f,0xbf]
-; SI-NEXT: s_setprio 1 ; encoding: [0x01,0x00,0x8f,0xbf]
-; SI-NEXT: s_setprio 2 ; encoding: [0x02,0x00,0x8f,0xbf]
-; SI-NEXT: s_setprio 3 ; encoding: [0x03,0x00,0x8f,0xbf]
-; SI-NEXT: s_setprio 10 ; encoding: [0x0a,0x00,0x8f,0xbf]
-; SI-NEXT: s_setprio -1 ; encoding: [0xff,0xff,0x8f,0xbf]
-; SI-NEXT: s_setprio 0 ; encoding: [0x00,0x00,0x8f,0xbf]
-; SI-NEXT: s_setprio 1 ; encoding: [0x01,0x00,0x8f,0xbf]
-; SI-NEXT: s_setprio -1 ; encoding: [0xff,0xff,0x8f,0xbf]
-; SI-NEXT: s_setpc_b64 s[30:31] ; encoding: [0x1e,0x20,0x80,0xbe]
- call void @llvm.amdgcn.s.setprio.relaxed(i16 0)
- call void @llvm.amdgcn.s.setprio.relaxed(i16 1)
- call void @llvm.amdgcn.s.setprio.relaxed(i16 2)
- call void @llvm.amdgcn.s.setprio.relaxed(i16 3)
- call void @llvm.amdgcn.s.setprio.relaxed(i16 10)
- call void @llvm.amdgcn.s.setprio.relaxed(i16 65535)
- call void @llvm.amdgcn.s.setprio.relaxed(i16 65536)
- call void @llvm.amdgcn.s.setprio.relaxed(i16 65537)
- call void @llvm.amdgcn.s.setprio.relaxed(i16 -1)
- ret void
-}
-
-attributes #0 = { nounwind }
diff --git a/llvm/test/CodeGen/AMDGPU/required-export-priority.ll b/llvm/test/CodeGen/AMDGPU/required-export-priority.ll
index 5d464a4fd314d..b6e2e75c9d28c 100644
--- a/llvm/test/CodeGen/AMDGPU/required-export-priority.ll
+++ b/llvm/test/CodeGen/AMDGPU/required-export-priority.ll
@@ -445,9 +445,7 @@ define amdgpu_ps void @test_export_in_callee_prio(float %v) #0 {
; GFX11: ; %bb.0:
; GFX11-NEXT: s_mov_b32 s32, 0
; GFX11-NEXT: v_add_f32_e32 v0, 1.0, v0
-; GFX11-NEXT: ; sched_barrier mask(0x00000000)
; GFX11-NEXT: s_setprio 0
-; GFX11-NEXT: ; sched_barrier mask(0x00000000)
; GFX11-NEXT: s_getpc_b64 s[0:1]
; GFX11-NEXT: s_add_u32 s0, s0, test_export_gfx at gotpcrel32@lo+4
; GFX11-NEXT: s_addc_u32 s1, s1, test_export_gfx at gotpcrel32@hi+12
@@ -461,9 +459,7 @@ define amdgpu_ps void @test_export_in_callee_prio(float %v) #0 {
; GFX1150-NEXT: s_setprio 2
; GFX1150-NEXT: s_mov_b32 s32, 0
; GFX1150-NEXT: v_add_f32_e32 v0, 1.0, v0
-; GFX1150-NEXT: ; sched_barrier mask(0x00000000)
; GFX1150-NEXT: s_setprio 2
-; GFX1150-NEXT: ; sched_barrier mask(0x00000000)
; GFX1150-NEXT: s_getpc_b64 s[0:1]
; GFX1150-NEXT: s_add_u32 s0, s0, test_export_gfx at gotpcrel32@lo+4
; GFX1150-NEXT: s_addc_u32 s1, s1, test_export_gfx at gotpcrel32@hi+12
diff --git a/llvm/test/CodeGen/AMDGPU/sched-setprio-relaxed.ll b/llvm/test/CodeGen/AMDGPU/sched-setprio-relaxed.ll
deleted file mode 100644
index 63a51aa4420c8..0000000000000
--- a/llvm/test/CodeGen/AMDGPU/sched-setprio-relaxed.ll
+++ /dev/null
@@ -1,134 +0,0 @@
-; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
-; RUN: llc -mtriple=amdgpu9.08-amd-amdhsa < %s | FileCheck --check-prefix=GCN %s
-
-; Test that llvm.amdgcn.s.setprio.relaxed allows instructions to be scheduled across it,
-; unlike llvm.amdgcn.s.setprio which is as a scheduling boundary.
-
-declare void @llvm.amdgcn.s.setprio.relaxed(i16)
-declare float @llvm.amdgcn.rcp.f32(float)
-
-define amdgpu_kernel void @test_relaxed_allows_salu(ptr addrspace(1) %out, i32 inreg %x, i32 inreg %y) {
-; GCN-LABEL: test_relaxed_allows_salu:
-; GCN: ; %bb.0:
-; GCN-NEXT: s_load_dwordx4 s[0:3], s[8:9], 0x0
-; GCN-NEXT: v_mov_b32_e32 v0, 0
-; GCN-NEXT: s_setprio 1
-; GCN-NEXT: s_waitcnt lgkmcnt(0)
-; GCN-NEXT: s_add_i32 s2, s2, s3
-; GCN-NEXT: s_add_i32 s3, s3, s3
-; GCN-NEXT: s_add_i32 s2, s2, s3
-; GCN-NEXT: v_mov_b32_e32 v1, s2
-; GCN-NEXT: global_store_dword v0, v1, s[0:1]
-; GCN-NEXT: s_endpgm
- %add1 = add i32 %x, %y
- call void @llvm.amdgcn.s.setprio.relaxed(i16 1)
- %add2 = add i32 %y, %y
- %sum = add i32 %add1, %add2
- store i32 %sum, ptr addrspace(1) %out
- ret void
-}
-
-define amdgpu_kernel void @test_relaxed_allows_valu(ptr addrspace(1) %out, float %x, float %y) {
-; GCN-LABEL: test_relaxed_allows_valu:
-; GCN: ; %bb.0:
-; GCN-NEXT: s_load_dwordx4 s[0:3], s[8:9], 0x0
-; GCN-NEXT: v_mov_b32_e32 v0, 0
-; GCN-NEXT: s_setprio 1
-; GCN-NEXT: s_waitcnt lgkmcnt(0)
-; GCN-NEXT: v_add_f32_e64 v1, s2, 1.0
-; GCN-NEXT: v_add_f32_e64 v2, s3, 2.0
-; GCN-NEXT: v_add_f32_e32 v1, v1, v2
-; GCN-NEXT: global_store_dword v0, v1, s[0:1]
-; GCN-NEXT: s_endpgm
- %add1 = fadd float %x, 1.0
- call void @llvm.amdgcn.s.setprio.relaxed(i16 1)
- %add2 = fadd float %y, 2.0
- %sum = fadd float %add1, %add2
- store float %sum, ptr addrspace(1) %out
- ret void
-}
-
-define amdgpu_kernel void @test_relaxed_allows_trans(ptr addrspace(1) %out, float %x, float %y) {
-; GCN-LABEL: test_relaxed_allows_trans:
-; GCN: ; %bb.0:
-; GCN-NEXT: s_load_dwordx4 s[0:3], s[8:9], 0x0
-; GCN-NEXT: v_mov_b32_e32 v2, 0
-; GCN-NEXT: s_setprio 1
-; GCN-NEXT: s_waitcnt lgkmcnt(0)
-; GCN-NEXT: v_rcp_f32_e32 v0, s2
-; GCN-NEXT: v_rcp_f32_e32 v1, s3
-; GCN-NEXT: v_add_f32_e32 v0, v0, v1
-; GCN-NEXT: global_store_dword v2, v0, s[0:1]
-; GCN-NEXT: s_endpgm
- %rcp1 = call float @llvm.amdgcn.rcp.f32(float %x)
- call void @llvm.amdgcn.s.setprio.relaxed(i16 1)
- %rcp2 = call float @llvm.amdgcn.rcp.f32(float %y)
- %sum = fadd float %rcp1, %rcp2
- store float %sum, ptr addrspace(1) %out
- ret void
-}
-
-define amdgpu_kernel void @test_relaxed_allows_mfma(ptr addrspace(1) %out, <4 x float> %in) {
-; GCN-LABEL: test_relaxed_allows_mfma:
-; GCN: ; %bb.0:
-; GCN-NEXT: s_load_dwordx2 s[4:5], s[8:9], 0x0
-; GCN-NEXT: s_load_dwordx4 s[0:3], s[8:9], 0x10
-; GCN-NEXT: v_mov_b32_e32 v0, 1.0
-; GCN-NEXT: v_mov_b32_e32 v2, 0x40400000
-; GCN-NEXT: v_mov_b32_e32 v1, 2.0
-; GCN-NEXT: v_mov_b32_e32 v4, 0
-; GCN-NEXT: s_waitcnt lgkmcnt(0)
-; GCN-NEXT: v_mov_b32_e32 v5, s0
-; GCN-NEXT: v_mov_b32_e32 v3, s1
-; GCN-NEXT: v_mov_b32_e32 v6, s2
-; GCN-NEXT: v_accvgpr_write_b32 a0, v5
-; GCN-NEXT: v_mov_b32_e32 v5, s3
-; GCN-NEXT: v_accvgpr_write_b32 a1, v3
-; GCN-NEXT: v_accvgpr_write_b32 a2, v6
-; GCN-NEXT: v_accvgpr_write_b32 a3, v5
-; GCN-NEXT: s_setprio 1
-; GCN-NEXT: v_mfma_f32_4x4x1f32 a[0:3], v0, v1, a[0:3]
-; GCN-NEXT: v_mov_b32_e32 v0, 4.0
-; GCN-NEXT: s_nop 1
-; GCN-NEXT: v_mfma_f32_4x4x1f32 a[0:3], v2, v0, a[0:3]
-; GCN-NEXT: s_nop 3
-; GCN-NEXT: v_accvgpr_read_b32 v0, a0
-; GCN-NEXT: v_accvgpr_read_b32 v1, a1
-; GCN-NEXT: v_accvgpr_read_b32 v2, a2
-; GCN-NEXT: v_accvgpr_read_b32 v3, a3
-; GCN-NEXT: s_nop 1
-; GCN-NEXT: global_store_dwordx4 v4, v[0:3], s[4:5]
-; GCN-NEXT: s_endpgm
- %mfma1 = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float 1.0, float 2.0, <4 x float> %in, i32 0, i32 0, i32 0)
- call void @llvm.amdgcn.s.setprio.relaxed(i16 1)
- %mfma2 = call <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float 3.0, float 4.0, <4 x float> %mfma1, i32 0, i32 0, i32 0)
- store <4 x float> %mfma2, ptr addrspace(1) %out
- ret void
-}
-
-declare <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float, float, <4 x float>, i32 immarg, i32 immarg, i32 immarg)
-
-; Memory operations cannot move across s_setprio due to its side effects.
-define amdgpu_kernel void @test_relaxed_blocks_memory(ptr addrspace(1) %out1, ptr addrspace(1) %out2, ptr addrspace(1) %out3, float %x) {
-; GCN-LABEL: test_relaxed_blocks_memory:
-; GCN: ; %bb.0:
-; GCN-NEXT: s_load_dwordx4 s[0:3], s[8:9], 0x0
-; GCN-NEXT: s_load_dword s6, s[8:9], 0x18
-; GCN-NEXT: v_mov_b32_e32 v0, 0
-; GCN-NEXT: v_mov_b32_e32 v1, 1.0
-; GCN-NEXT: s_load_dwordx2 s[4:5], s[8:9], 0x10
-; GCN-NEXT: s_waitcnt lgkmcnt(0)
-; GCN-NEXT: global_store_dword v0, v1, s[0:1]
-; GCN-NEXT: v_rcp_f32_e32 v1, s6
-; GCN-NEXT: v_mov_b32_e32 v2, 2.0
-; GCN-NEXT: s_setprio 1
-; GCN-NEXT: global_store_dword v0, v2, s[2:3]
-; GCN-NEXT: global_store_dword v0, v1, s[4:5]
-; GCN-NEXT: s_endpgm
- store float 1.0, ptr addrspace(1) %out1
- %rcp = call float @llvm.amdgcn.rcp.f32(float %x)
- call void @llvm.amdgcn.s.setprio.relaxed(i16 1)
- store float 2.0, ptr addrspace(1) %out2
- store float %rcp, ptr addrspace(1) %out3
- ret void
-}
diff --git a/llvm/test/CodeGen/AMDGPU/sched-setprio.ll b/llvm/test/CodeGen/AMDGPU/sched-setprio.ll
index 3cb868734dc5b..c6b83fd715782 100644
--- a/llvm/test/CodeGen/AMDGPU/sched-setprio.ll
+++ b/llvm/test/CodeGen/AMDGPU/sched-setprio.ll
@@ -4,9 +4,7 @@ declare void @llvm.amdgcn.s.setprio(i16)
declare <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float, float, <4 x float>, i32, i32, i32)
; GCN-LABEL: {{^}}test_mfma_f32_4x4x1f32:
-; GCN: sched_barrier mask(0x00000000)
-; GCN-NEXT: s_setprio 1
-; GCN-NEXT: sched_barrier mask(0x00000000)
+; GCN: s_setprio 1
; GCN: v_mfma
; GCN: v_mfma
; GCN: s_setprio 0
>From e717be03462eca1c669a0c524c63a7b030a8e0bc Mon Sep 17 00:00:00 2001
From: Frederik Harwath <fharwath at amd.com>
Date: Thu, 27 Aug 2026 10:48:27 -0400
Subject: [PATCH 10/10] [AMDGPU] Weaken S_SETPRIO scheduling barrier behavior
The S_SETPRIO instruction was made a scheduling boundary to prevent
the movement of vector instructions (VALU, MFMA, TRANS) across
it. This movement would contradict the intention of the wave priority
change. However, the full scheduling boundary prevents other types of
instructions from moving across the S_SETPRIO which were not
necessarily meant to be affected, in particular SALU instructions.
Remove the scheduling boundary property from S_SETPRIO and change
AMDGPUIGroupLP to treat the instruction like a SCHED_BARRIER whose
mask allows only SALU instructions to cross the barrier. Memory
instructions could also be admitted, but this is not possible with the
current side effects of S_SETPRIO. Improving this is left as follow-up
work.
---
llvm/lib/Target/AMDGPU/AMDGPUIGroupLP.cpp | 38 ++++--
llvm/lib/Target/AMDGPU/SIInstrInfo.cpp | 1 -
.../AMDGPU/required-export-priority.ll | 12 +-
llvm/test/CodeGen/AMDGPU/sched-setprio.ll | 126 +++++++++++++++++-
4 files changed, 157 insertions(+), 20 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUIGroupLP.cpp b/llvm/lib/Target/AMDGPU/AMDGPUIGroupLP.cpp
index df7e80d62e065..a452e99722a1c 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUIGroupLP.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUIGroupLP.cpp
@@ -2485,7 +2485,10 @@ class IGroupLPDAGMutation : public ScheduleDAGMutation {
DenseMap<int, SUnitsToCandidateSGsMap> SyncedInstrs;
// Add DAG edges that enforce SCHED_BARRIER ordering.
- void addSchedBarrierEdges(SUnit &SU);
+ // AllowedMask specifies which instruction types are allowed to
+ // cross the barrier and is inverted internally to get a blocked
+ // mask.
+ void addSchedBarrierEdges(SUnit &Barrier, SchedGroupMask AllowedMask);
// Use a SCHED_BARRIER's mask to identify instruction SchedGroups that should
// not be reordered accross the SCHED_BARRIER. This is used for the base
@@ -2747,11 +2750,20 @@ void IGroupLPDAGMutation::apply(ScheduleDAGInstrs *DAGInstrs) {
unsigned Opc = R->getInstr()->getOpcode();
// SCHED_[GROUP_]BARRIER and IGLP are mutually exclusive.
if (Opc == AMDGPU::SCHED_BARRIER) {
- addSchedBarrierEdges(*R);
+ auto AllowedMask = (SchedGroupMask)R->getInstr()->getOperand(0).getImm();
+ addSchedBarrierEdges(*R, AllowedMask);
FoundSB = true;
} else if (Opc == AMDGPU::SCHED_GROUP_BARRIER) {
initSchedGroupBarrierPipelineStage(R);
FoundSB = true;
+ } else if (Opc == AMDGPU::S_SETPRIO) {
+ // Moving VALU, MFMA, TRANS instructions across
+ // S_SETPRIO would contradict the intention of
+ // the instruction.
+ // FIXME Memory instructions could be allowed here,
+ // but are blocked by the side effects of S_SETPRIO.
+ addSchedBarrierEdges(*R, SchedGroupMask::SALU);
+ FoundSB = true;
} else if (Opc == AMDGPU::IGLP_OPT) {
if (!FoundSB && !FoundIGLP) {
FoundIGLP = true;
@@ -2769,22 +2781,24 @@ void IGroupLPDAGMutation::apply(ScheduleDAGInstrs *DAGInstrs) {
}
}
-void IGroupLPDAGMutation::addSchedBarrierEdges(SUnit &SchedBarrier) {
- MachineInstr &MI = *SchedBarrier.getInstr();
- assert(MI.getOpcode() == AMDGPU::SCHED_BARRIER);
- LLVM_DEBUG(dbgs() << "Building SchedGroup for SchedBarrier with Mask: "
- << MI.getOperand(0).getImm() << "\n");
- auto InvertedMask =
- invertSchedBarrierMask((SchedGroupMask)MI.getOperand(0).getImm());
- SchedGroup SG(InvertedMask, std::nullopt, DAG, TII);
+void IGroupLPDAGMutation::addSchedBarrierEdges(SUnit &Barrier,
+ SchedGroupMask AllowedMask) {
+ MachineInstr &MI = *Barrier.getInstr();
+
+ auto BlockedMask = invertSchedBarrierMask(AllowedMask);
+ LLVM_DEBUG(dbgs() << "Building SchedGroup for "
+ << TII->getName(MI.getOpcode()) << " with blocked mask: "
+ << format_hex((int)BlockedMask, 10, true) << "\n");
+
+ SchedGroup SG(BlockedMask, std::nullopt, DAG, TII);
for (SUnit &SU : DAG->SUnits)
if (SG.canAddSU(SU))
SG.add(SU);
- // Preserve original instruction ordering relative to the SCHED_BARRIER.
+ // Preserve original instruction ordering relative to the barrier.
SG.link(
- SchedBarrier,
+ Barrier,
(function_ref<bool(const SUnit *A, const SUnit *B)>)[](
const SUnit *A, const SUnit *B) { return A->NodeNum > B->NodeNum; });
}
diff --git a/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp b/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp
index b2c3772cb0041..793c30d8369cb 100644
--- a/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp
@@ -4489,7 +4489,6 @@ bool SIInstrInfo::isSchedulingBoundary(const MachineInstr &MI,
return MI.modifiesRegister(AMDGPU::EXEC, &RI) ||
MI.getOpcode() == AMDGPU::S_SETREG_IMM32_B32 ||
MI.getOpcode() == AMDGPU::S_SETREG_B32 ||
- MI.getOpcode() == AMDGPU::S_SETPRIO ||
MI.getOpcode() == AMDGPU::S_SETPRIO_INC_WG ||
changesVGPRIndexingMode(MI);
}
diff --git a/llvm/test/CodeGen/AMDGPU/required-export-priority.ll b/llvm/test/CodeGen/AMDGPU/required-export-priority.ll
index b6e2e75c9d28c..48002356be00a 100644
--- a/llvm/test/CodeGen/AMDGPU/required-export-priority.ll
+++ b/llvm/test/CodeGen/AMDGPU/required-export-priority.ll
@@ -443,13 +443,13 @@ define amdgpu_ps void @test_export_in_callee(float %v) #0 {
define amdgpu_ps void @test_export_in_callee_prio(float %v) #0 {
; GFX11-LABEL: test_export_in_callee_prio:
; GFX11: ; %bb.0:
-; GFX11-NEXT: s_mov_b32 s32, 0
-; GFX11-NEXT: v_add_f32_e32 v0, 1.0, v0
-; GFX11-NEXT: s_setprio 0
; GFX11-NEXT: s_getpc_b64 s[0:1]
; GFX11-NEXT: s_add_u32 s0, s0, test_export_gfx at gotpcrel32@lo+4
; GFX11-NEXT: s_addc_u32 s1, s1, test_export_gfx at gotpcrel32@hi+12
+; GFX11-NEXT: v_add_f32_e32 v0, 1.0, v0
; GFX11-NEXT: s_load_b64 s[0:1], s[0:1], 0x0
+; GFX11-NEXT: s_mov_b32 s32, 0
+; GFX11-NEXT: s_setprio 0
; GFX11-NEXT: s_waitcnt lgkmcnt(0)
; GFX11-NEXT: s_swappc_b64 s[30:31], s[0:1]
; GFX11-NEXT: s_endpgm
@@ -457,13 +457,13 @@ define amdgpu_ps void @test_export_in_callee_prio(float %v) #0 {
; GFX1150-LABEL: test_export_in_callee_prio:
; GFX1150: ; %bb.0:
; GFX1150-NEXT: s_setprio 2
-; GFX1150-NEXT: s_mov_b32 s32, 0
-; GFX1150-NEXT: v_add_f32_e32 v0, 1.0, v0
-; GFX1150-NEXT: s_setprio 2
; GFX1150-NEXT: s_getpc_b64 s[0:1]
; GFX1150-NEXT: s_add_u32 s0, s0, test_export_gfx at gotpcrel32@lo+4
; GFX1150-NEXT: s_addc_u32 s1, s1, test_export_gfx at gotpcrel32@hi+12
+; GFX1150-NEXT: v_add_f32_e32 v0, 1.0, v0
; GFX1150-NEXT: s_load_b64 s[0:1], s[0:1], 0x0
+; GFX1150-NEXT: s_mov_b32 s32, 0
+; GFX1150-NEXT: s_setprio 2
; GFX1150-NEXT: s_waitcnt lgkmcnt(0)
; GFX1150-NEXT: s_swappc_b64 s[30:31], s[0:1]
; GFX1150-NEXT: s_endpgm
diff --git a/llvm/test/CodeGen/AMDGPU/sched-setprio.ll b/llvm/test/CodeGen/AMDGPU/sched-setprio.ll
index c6b83fd715782..051feae73f2bb 100644
--- a/llvm/test/CodeGen/AMDGPU/sched-setprio.ll
+++ b/llvm/test/CodeGen/AMDGPU/sched-setprio.ll
@@ -1,7 +1,9 @@
-; RUN: llc -mtriple=amdgpu9.08 < %s | FileCheck --check-prefix=GCN %s
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
+; RUN: llc -mtriple=amdgpu9.08-amd-amdhsa < %s | FileCheck --check-prefix=GCN %s
declare void @llvm.amdgcn.s.setprio(i16)
declare <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float, float, <4 x float>, i32, i32, i32)
+declare float @llvm.amdgcn.rcp.f32(float)
; GCN-LABEL: {{^}}test_mfma_f32_4x4x1f32:
; GCN: s_setprio 1
@@ -9,6 +11,39 @@ declare <4 x float> @llvm.amdgcn.mfma.f32.4x4x1f32(float, float, <4 x float>, i3
; GCN: v_mfma
; GCN: s_setprio 0
define amdgpu_kernel void @test_mfma_f32_4x4x1f32(ptr addrspace(1) %arg) #0 {
+; GCN-LABEL: test_mfma_f32_4x4x1f32:
+; GCN: ; %bb.0: ; %bb
+; GCN-NEXT: s_load_dwordx2 s[4:5], s[8:9], 0x0
+; GCN-NEXT: v_mov_b32_e32 v4, 0
+; GCN-NEXT: s_waitcnt lgkmcnt(0)
+; GCN-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x0
+; GCN-NEXT: s_setprio 1
+; GCN-NEXT: v_mov_b32_e32 v0, 1.0
+; GCN-NEXT: v_mov_b32_e32 v1, 0x40400000
+; GCN-NEXT: v_mov_b32_e32 v2, 2.0
+; GCN-NEXT: s_waitcnt lgkmcnt(0)
+; GCN-NEXT: v_mov_b32_e32 v5, s0
+; GCN-NEXT: v_mov_b32_e32 v3, s1
+; GCN-NEXT: v_mov_b32_e32 v6, s2
+; GCN-NEXT: v_accvgpr_write_b32 a0, v5
+; GCN-NEXT: v_mov_b32_e32 v5, s3
+; GCN-NEXT: v_accvgpr_write_b32 a1, v3
+; GCN-NEXT: v_accvgpr_write_b32 a2, v6
+; GCN-NEXT: v_accvgpr_write_b32 a3, v5
+; GCN-NEXT: s_nop 0
+; GCN-NEXT: v_mfma_f32_4x4x1f32 a[0:3], v0, v2, a[0:3]
+; GCN-NEXT: v_mov_b32_e32 v0, 4.0
+; GCN-NEXT: s_nop 1
+; GCN-NEXT: v_mfma_f32_4x4x1f32 a[0:3], v1, v0, a[0:3]
+; GCN-NEXT: s_setprio 0
+; GCN-NEXT: s_nop 2
+; GCN-NEXT: v_accvgpr_read_b32 v0, a0
+; GCN-NEXT: v_accvgpr_read_b32 v1, a1
+; GCN-NEXT: v_accvgpr_read_b32 v2, a2
+; GCN-NEXT: v_accvgpr_read_b32 v3, a3
+; GCN-NEXT: s_nop 1
+; GCN-NEXT: global_store_dwordx4 v4, v[0:3], s[4:5]
+; GCN-NEXT: s_endpgm
bb:
%in.1 = load <4 x float>, ptr addrspace(1) %arg
call void @llvm.amdgcn.s.setprio(i16 1)
@@ -18,3 +53,92 @@ bb:
store <4 x float> %mai.2, ptr addrspace(1) %arg
ret void
}
+
+; Test that SALU can be scheduled across S_SETPRIO
+define amdgpu_kernel void @test_allows_salu(ptr addrspace(1) %out, i32 inreg %x, i32 inreg %y) {
+; GCN-LABEL: test_allows_salu:
+; GCN: ; %bb.0:
+; GCN-NEXT: s_load_dwordx4 s[0:3], s[8:9], 0x0
+; GCN-NEXT: v_mov_b32_e32 v0, 0
+; GCN-NEXT: s_setprio 1
+; GCN-NEXT: s_waitcnt lgkmcnt(0)
+; GCN-NEXT: s_add_i32 s2, s2, s3
+; GCN-NEXT: s_add_i32 s3, s3, s3
+; GCN-NEXT: s_add_i32 s2, s2, s3
+; GCN-NEXT: v_mov_b32_e32 v1, s2
+; GCN-NEXT: global_store_dword v0, v1, s[0:1]
+; GCN-NEXT: s_endpgm
+ %add1 = add i32 %x, %y
+ call void @llvm.amdgcn.s.setprio(i16 1)
+ %add2 = add i32 %y, %y
+ %sum = add i32 %add1, %add2
+ store i32 %sum, ptr addrspace(1) %out
+ ret void
+}
+
+; Test that VALU is blocked by S_SETPRIO
+define amdgpu_kernel void @test_blocks_valu(ptr addrspace(1) %out, float %x, float %y) {
+; GCN-LABEL: test_blocks_valu:
+; GCN: ; %bb.0:
+; GCN-NEXT: s_load_dwordx4 s[0:3], s[8:9], 0x0
+; GCN-NEXT: v_mov_b32_e32 v0, 0
+; GCN-NEXT: s_waitcnt lgkmcnt(0)
+; GCN-NEXT: v_add_f32_e64 v1, s2, 1.0
+; GCN-NEXT: s_setprio 1
+; GCN-NEXT: v_add_f32_e64 v2, s3, 2.0
+; GCN-NEXT: v_add_f32_e32 v1, v1, v2
+; GCN-NEXT: global_store_dword v0, v1, s[0:1]
+; GCN-NEXT: s_endpgm
+ %add1 = fadd float %x, 1.0
+ call void @llvm.amdgcn.s.setprio(i16 1)
+ %add2 = fadd float %y, 2.0
+ %sum = fadd float %add1, %add2
+ store float %sum, ptr addrspace(1) %out
+ ret void
+}
+
+; Test that TRANS is blocked by S_SETPRIO
+define amdgpu_kernel void @test_blocks_trans(ptr addrspace(1) %out, float %x, float %y) {
+; GCN-LABEL: test_blocks_trans:
+; GCN: ; %bb.0:
+; GCN-NEXT: s_load_dwordx4 s[0:3], s[8:9], 0x0
+; GCN-NEXT: v_mov_b32_e32 v0, 0
+; GCN-NEXT: s_waitcnt lgkmcnt(0)
+; GCN-NEXT: v_rcp_f32_e32 v1, s2
+; GCN-NEXT: s_setprio 1
+; GCN-NEXT: v_rcp_f32_e32 v2, s3
+; GCN-NEXT: v_add_f32_e32 v1, v1, v2
+; GCN-NEXT: global_store_dword v0, v1, s[0:1]
+; GCN-NEXT: s_endpgm
+ %rcp1 = call float @llvm.amdgcn.rcp.f32(float %x)
+ call void @llvm.amdgcn.s.setprio(i16 1)
+ %rcp2 = call float @llvm.amdgcn.rcp.f32(float %y)
+ %sum = fadd float %rcp1, %rcp2
+ store float %sum, ptr addrspace(1) %out
+ ret void
+}
+
+; Test that memory operations are blocked by S_SETPRIO
+define amdgpu_kernel void @test_memory_ops(ptr addrspace(1) %out1, ptr addrspace(1) %out2, ptr addrspace(1) %out3, float %x) {
+; GCN-LABEL: test_memory_ops:
+; GCN: ; %bb.0:
+; GCN-NEXT: s_load_dwordx4 s[0:3], s[8:9], 0x0
+; GCN-NEXT: s_load_dword s6, s[8:9], 0x18
+; GCN-NEXT: v_mov_b32_e32 v0, 0
+; GCN-NEXT: v_mov_b32_e32 v1, 1.0
+; GCN-NEXT: s_load_dwordx2 s[4:5], s[8:9], 0x10
+; GCN-NEXT: s_waitcnt lgkmcnt(0)
+; GCN-NEXT: global_store_dword v0, v1, s[0:1]
+; GCN-NEXT: v_rcp_f32_e32 v1, s6
+; GCN-NEXT: s_setprio 1
+; GCN-NEXT: v_mov_b32_e32 v2, 2.0
+; GCN-NEXT: global_store_dword v0, v2, s[2:3]
+; GCN-NEXT: global_store_dword v0, v1, s[4:5]
+; GCN-NEXT: s_endpgm
+ store float 1.0, ptr addrspace(1) %out1
+ %rcp = call float @llvm.amdgcn.rcp.f32(float %x)
+ call void @llvm.amdgcn.s.setprio(i16 1)
+ store float 2.0, ptr addrspace(1) %out2
+ store float %rcp, ptr addrspace(1) %out3
+ ret void
+}
More information about the cfe-commits
mailing list