[llvm] [AMDGPU] Add scheduling mask operand to S_SETPRIO (PR #204850)
via llvm-commits
llvm-commits at lists.llvm.org
Wed Jul 29 06:11:33 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-backend-amdgpu
Author: Frederik Harwath (frederik-h)
<details>
<summary>Changes</summary>
The S_SETPRIO instruction is a scheduling boundary. The intention of this is to prevent moving VALU instructions across S_SETPRIO (see commit [68901fd "[AMDGPU] Consider S_SETPRIO a scheduling boundary"](https://github.com/llvm/llvm-project/commit/68901fdbebb73617067cbab9888364ba9f8a2843)). However, making the instruction a scheduling boundary also prevents movement of other instruction types, e.g. SALU instructions.
Add a mask operand to S_SETPRIO that controls which instruction types can be scheduled across it, using the same semantics as the mask of the SCHED_BARRIER instruction. The default mask 0 creates a full scheduling boundary, preserving existing behavior. The machine verifier checks that the mask is compatible with inherent dependencies of S_SETPRIO, i.e. memory operations cannot cross it.
The existing llvm.amdgcn.s.setprio intrinsic is changed to create S_SETPRIO instructions with mask 0 and a new llvm.amdgcn.s.setprio.mask intrinsic with an explicit mask parameter is added.
---
Patch is 29.79 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/204850.diff
13 Files Affected:
- (modified) llvm/docs/AMDGPUUsage.rst (+10)
- (modified) llvm/include/llvm/IR/IntrinsicsAMDGPU.td (+6)
- (modified) llvm/lib/Target/AMDGPU/AMDGPUIGroupLP.cpp (+12-6)
- (modified) llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp (+1)
- (modified) llvm/lib/Target/AMDGPU/AMDGPUSetWavePriority.cpp (+2-1)
- (modified) llvm/lib/Target/AMDGPU/GCNHazardRecognizer.cpp (+6-3)
- (modified) llvm/lib/Target/AMDGPU/SIInstrInfo.cpp (+29-2)
- (modified) llvm/lib/Target/AMDGPU/SIInstrInfo.h (+3)
- (modified) llvm/lib/Target/AMDGPU/SOPInstructions.td (+16-2)
- (modified) llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.setprio.ll (+33)
- (modified) llvm/test/CodeGen/AMDGPU/required-export-priority.mir (+25-25)
- (added) llvm/test/CodeGen/AMDGPU/sched-setprio-mask.ll (+281)
- (added) llvm/test/MachineVerifier/AMDGPU/verify-setprio-mask.mir (+37)
``````````diff
diff --git a/llvm/docs/AMDGPUUsage.rst b/llvm/docs/AMDGPUUsage.rst
index 443dd7f9772d7..e2af2257e8402 100644
--- a/llvm/docs/AMDGPUUsage.rst
+++ b/llvm/docs/AMDGPUUsage.rst
@@ -1915,6 +1915,16 @@ 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. 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.
+
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 565637b36131c..d6d5b69b28168 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]>;
+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 df7e80d62e065..04bf0525eeb52 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 (Opc == AMDGPU::SCHED_BARRIER) {
+ if (TII->isSchedBarrierLike(*R->getInstr())) {
addSchedBarrierEdges(*R);
FoundSB = true;
} else if (Opc == AMDGPU::SCHED_GROUP_BARRIER) {
@@ -2771,11 +2771,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 881afaf920037..3eab1bc2c4c4a 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp
@@ -1821,6 +1821,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 e3448be9f44f5..df3fe8c9562f3 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 af326d60aacb4..3970f701fe164 100644
--- a/llvm/lib/Target/AMDGPU/GCNHazardRecognizer.cpp
+++ b/llvm/lib/Target/AMDGPU/GCNHazardRecognizer.cpp
@@ -3699,7 +3699,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;
}
@@ -3777,7 +3778,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.
@@ -3792,7 +3794,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 b2c3772cb0041..42a8d52f0a5b5 100644
--- a/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp
@@ -4462,6 +4462,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 {
@@ -4480,7 +4496,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
@@ -4489,7 +4505,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);
}
@@ -5958,6 +5973,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/lib/Target/AMDGPU/SIInstrInfo.h b/llvm/lib/Target/AMDGPU/SIInstrInfo.h
index 4c8641a6091d7..75c4eab1be937 100644
--- a/llvm/lib/Target/AMDGPU/SIInstrInfo.h
+++ b/llvm/lib/Target/AMDGPU/SIInstrInfo.h
@@ -1149,6 +1149,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 8fa54de570e67..b6b8ba1720994 100644
--- a/llvm/lib/Target/AMDGPU/SOPInstructions.td
+++ b/llvm/lib/Target/AMDGPU/SOPInstructions.td
@@ -1754,10 +1754,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 87df1e152bef5..98dea66d3d99e 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 -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:
@@ -46,4 +47,36 @@ 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:
+; 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: ; %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_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_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]
+; 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 cf12eff2c0527..34924e1e62381 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
@@ -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
@@ -362,28 +362,28 @@ 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
+ 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-mask.ll b/llvm/test/CodeGen/AMDGPU/sched-setprio-mask.ll
new file mode 100644
index 0000000000000..08517898e2b1a
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/sched-setprio-mask.ll
@@ -0,0 +1,281 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
+; RUN: llc -mtriple=amdgpu9.08 < %s | FileCheck --check-prefix=GCN %s
+
+; Demonstrate that the mask bits that are compatible with s_set_prio
+; allow corresponding instruction types to be scheduled across the s_setprio
+; instruction.
+
+declare void @llvm.amdgcn.s.setprio(i16)
+declare void @llvm.amdgcn.s.setprio.mask(i16, i32)
+declare float @llvm.amdgcn.rcp.f32(float)
+
+define amdgpu_cs void @test_mask0_blocks_salu(ptr addrspace(1) %out, i32 inreg %x, i32 inreg %y) {
+; GCN-LABEL: test_mask0_blocks_salu:
+; GCN: ; %bb.0:
+; GCN-NEXT: s_add_i32 s0, s0, s1
+; GCN-NEXT: s_setprio 1
+; GCN-NEXT: s_add_i32 s1, s1, s1
+; GCN-NEXT: s_add_i32 s0, s0, s1
+; GCN-NEXT: v_mov_b32_e32 v2, s0
+; GCN-NEXT: global_store_dword v[0:1], v2, off
+; GCN-NEXT: s_endpgm
+ %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
+}
+
+define amdgpu_cs void @test_mask0_blocks_valu(ptr addrspace(1) %out, float %x, float %y) {
+; GCN-LABEL: test_mask0_blocks_valu:
+; GCN: ; %bb.0:
+; GCN-NEXT: v_add_f32_e32 v2, 1.0, v2
+; GCN-NEXT: s_setprio 1
+; GCN-NEXT: v_add_f32_e32 v3, 2.0, v3
+; GCN-NEXT: v_add_f32_e32 v2, v2, v3
+; GCN-NEXT: global_store_dword v[0:1], v2, off
+; GCN-NEXT: s_endpgm
+ %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
+}
+
+define amdgpu_cs void @test_mask0_blocks_trans(ptr addrspace(1) %out, float %x, float %y) {
+; GCN-LABEL: test_mask0_blocks_trans:
+; GCN: ; %bb.0:
+; GCN-NEXT: v_rcp_f32_e32 v2, v2
+; GCN-NEXT: s_setprio 1
+; GCN-NEXT: v_rcp_f32_e32 v3, v3
+; GCN-NEXT: v_add_f32_e32 v2, v2, v3
+; GCN-NEXT: global_store_dword v[0:1], v2, off
+; GCN-NEXT: s_...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/204850
More information about the llvm-commits
mailing list