[llvm] [AMDGPU] Add scheduling mask operand to S_SETPRIO (PR #204850)
Frederik Harwath via llvm-commits
llvm-commits at lists.llvm.org
Wed Jul 29 03:36:38 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 1/5] [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 2/5] [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 3/5] [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 4/5] 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 8b018669c034c1cbb86fe539d509b879a673bb4d Mon Sep 17 00:00:00 2001
From: Frederik Harwath <fharwath at amd.com>
Date: Wed, 29 Jul 2026 05:10:28 -0400
Subject: [PATCH 5/5] Reorganize tests
- Move setprio verifier testing into separate file
- Restore baseline sched-setprio.ll
- Move PR test cases into new sched-setprio-mask.ll file
---
.../test/CodeGen/AMDGPU/sched-setprio-mask.ll | 281 ++++++++++++++++++
llvm/test/CodeGen/AMDGPU/sched-setprio.ll | 275 +----------------
.../AMDGPU/verify-setprio-mask.mir | 37 +++
3 files changed, 332 insertions(+), 261 deletions(-)
create mode 100644 llvm/test/CodeGen/AMDGPU/sched-setprio-mask.ll
create mode 100644 llvm/test/MachineVerifier/AMDGPU/verify-setprio-mask.mir
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_endpgm
+ %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
+}
+
+define amdgpu_cs void @test_mask4_allows_salu(ptr addrspace(1) %out, i32 inreg %x, i32 inreg %y) {
+; GCN-LABEL: test_mask4_allows_salu:
+; GCN: ; %bb.0:
+; GCN-NEXT: s_add_i32 s0, s0, s1
+; GCN-NEXT: s_add_i32 s1, s1, s1
+; GCN-NEXT: s_add_i32 s0, s0, s1
+; GCN-NEXT: s_setprio 1
+; 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 4)
+ %add2 = add i32 %y, %y
+ %sum = add i32 %add1, %add2
+ store i32 %sum, ptr addrspace(1) %out
+ ret void
+}
+
+define amdgpu_cs void @test_mask4_blocks_valu(ptr addrspace(1) %out, float %x, float %y) {
+; GCN-LABEL: test_mask4_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 4)
+ %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_mask2_allows_valu(ptr addrspace(1) %out, float %x, float %y) {
+; GCN-LABEL: test_mask2_allows_valu:
+; GCN: ; %bb.0:
+; GCN-NEXT: v_add_f32_e32 v2, 1.0, v2
+; GCN-NEXT: v_add_f32_e32 v3, 2.0, v3
+; GCN-NEXT: v_add_f32_e32 v2, v2, v3
+; GCN-NEXT: s_setprio 1
+; 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 2)
+ %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_mask2_blocks_salu(ptr addrspace(1) %out, i32 inreg %x, i32 inreg %y) {
+; GCN-LABEL: test_mask2_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 2)
+ %add2 = add i32 %y, %y
+ %sum = add i32 %add1, %add2
+ store i32 %sum, ptr addrspace(1) %out
+ ret void
+}
+
+define amdgpu_cs void @test_mask1_allows_salu(ptr addrspace(1) %out, i32 inreg %x, i32 inreg %y) {
+; GCN-LABEL: test_mask1_allows_salu:
+; GCN: ; %bb.0:
+; GCN-NEXT: s_add_i32 s0, s0, s1
+; 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: s_setprio 1
+; 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 1)
+ %add2 = add i32 %y, %y
+ %sum = add i32 %add1, %add2
+ store i32 %sum, ptr addrspace(1) %out
+ ret void
+}
+
+define amdgpu_cs void @test_mask1_allows_valu(ptr addrspace(1) %out, float %x, float %y) {
+; GCN-LABEL: test_mask1_allows_valu:
+; GCN: ; %bb.0:
+; GCN-NEXT: v_add_f32_e32 v2, 1.0, v2
+; GCN-NEXT: v_add_f32_e32 v3, 2.0, v3
+; GCN-NEXT: v_add_f32_e32 v2, v2, v3
+; GCN-NEXT: s_setprio 1
+; 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 1)
+ %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_mask1024_allows_trans(ptr addrspace(1) %out, float %x, float %y) {
+; GCN-LABEL: test_mask1024_allows_trans:
+; GCN: ; %bb.0:
+; GCN-NEXT: v_rcp_f32_e32 v2, v2
+; GCN-NEXT: v_rcp_f32_e32 v3, v3
+; GCN-NEXT: s_setprio 1
+; GCN-NEXT: v_add_f32_e32 v2, v2, v3
+; GCN-NEXT: global_store_dword v[0:1], v2, off
+; GCN-NEXT: s_endpgm
+ %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
+}
+
+define amdgpu_cs void @test_mask1024_blocks_salu(ptr addrspace(1) %out, i32 inreg %x, i32 inreg %y) {
+; GCN-LABEL: test_mask1024_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 1024)
+ %add2 = add i32 %y, %y
+ %sum = add i32 %add1, %add2
+ store i32 %sum, ptr addrspace(1) %out
+ ret void
+}
+
+define amdgpu_cs void @test_mask6_allows_salu_and_valu(ptr addrspace(1) %out, i32 inreg %x, i32 inreg %y) {
+; GCN-LABEL: test_mask6_allows_salu_and_valu:
+; GCN: ; %bb.0:
+; GCN-NEXT: s_add_i32 s0, s0, s1
+; 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: s_setprio 1
+; 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 6)
+ %add2 = add i32 %y, %y
+ %sum = add i32 %add1, %add2
+ store i32 %sum, ptr addrspace(1) %out
+ ret void
+}
+
+define amdgpu_cs void @test_mask8_allows_mfma(ptr addrspace(1) %out, <4 x float> %in) {
+; GCN-LABEL: test_mask8_allows_mfma:
+; GCN: ; %bb.0:
+; GCN-NEXT: v_mov_b32_e32 v7, 1.0
+; GCN-NEXT: v_mov_b32_e32 v6, 2.0
+; GCN-NEXT: v_accvgpr_write_b32 a0, v2
+; GCN-NEXT: v_accvgpr_write_b32 a1, v3
+; GCN-NEXT: s_setprio 1
+; GCN-NEXT: v_mov_b32_e32 v2, 0x40400000
+; GCN-NEXT: v_accvgpr_write_b32 a2, v4
+; GCN-NEXT: v_accvgpr_write_b32 a3, v5
+; GCN-NEXT: v_mov_b32_e32 v3, 4.0
+; GCN-NEXT: v_mfma_f32_4x4x1f32 a[0:3], v7, v6, a[0:3]
+; GCN-NEXT: s_nop 0
+; GCN-NEXT: v_mfma_f32_4x4x1f32 a[0:3], v2, v3, a[0:3]
+; GCN-NEXT: s_nop 3
+; GCN-NEXT: v_accvgpr_read_b32 v5, a3
+; GCN-NEXT: v_accvgpr_read_b32 v4, a2
+; GCN-NEXT: v_accvgpr_read_b32 v3, a1
+; GCN-NEXT: v_accvgpr_read_b32 v2, a0
+; GCN-NEXT: s_nop 1
+; GCN-NEXT: global_store_dwordx4 v[0:1], v[2:5], off
+; 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.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
+}
+
+define amdgpu_cs void @test_mask0_blocks_mfma(ptr addrspace(1) %out, <4 x float> %in) {
+; GCN-LABEL: test_mask0_blocks_mfma:
+; GCN: ; %bb.0:
+; GCN-NEXT: v_mov_b32_e32 v7, 1.0
+; GCN-NEXT: v_accvgpr_write_b32 a0, v2
+; GCN-NEXT: v_accvgpr_write_b32 a1, v3
+; GCN-NEXT: v_accvgpr_write_b32 a2, v4
+; GCN-NEXT: v_accvgpr_write_b32 a3, v5
+; GCN-NEXT: v_mov_b32_e32 v6, 2.0
+; GCN-NEXT: s_nop 1
+; GCN-NEXT: v_mfma_f32_4x4x1f32 a[0:3], v7, v6, a[0:3]
+; GCN-NEXT: s_setprio 1
+; GCN-NEXT: v_mov_b32_e32 v2, 0x40400000
+; GCN-NEXT: v_mov_b32_e32 v3, 4.0
+; GCN-NEXT: s_nop 1
+; GCN-NEXT: v_mfma_f32_4x4x1f32 a[0:3], v2, v3, a[0:3]
+; GCN-NEXT: s_nop 3
+; GCN-NEXT: v_accvgpr_read_b32 v5, a3
+; GCN-NEXT: v_accvgpr_read_b32 v4, a2
+; GCN-NEXT: v_accvgpr_read_b32 v3, a1
+; GCN-NEXT: v_accvgpr_read_b32 v2, a0
+; GCN-NEXT: s_nop 1
+; GCN-NEXT: global_store_dwordx4 v[0:1], v[2:5], off
+; 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.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
+}
diff --git a/llvm/test/CodeGen/AMDGPU/sched-setprio.ll b/llvm/test/CodeGen/AMDGPU/sched-setprio.ll
index c0db14447e11b..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 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.
-
-;--- 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
}
diff --git a/llvm/test/MachineVerifier/AMDGPU/verify-setprio-mask.mir b/llvm/test/MachineVerifier/AMDGPU/verify-setprio-mask.mir
new file mode 100644
index 0000000000000..75620af223e48
--- /dev/null
+++ b/llvm/test/MachineVerifier/AMDGPU/verify-setprio-mask.mir
@@ -0,0 +1,37 @@
+# RUN: not --crash llc -mtriple=amdgpu9.08 -run-pass=none -filetype=null %s 2>&1 | FileCheck %s
+
+# CHECK: *** Bad machine code: S_SETPRIO mask contains invalid memory operation bits ***
+# CHECK: - instruction: S_SETPRIO 1, 16
+
+# CHECK: *** Bad machine code: S_SETPRIO mask contains invalid memory operation bits ***
+# CHECK: - instruction: S_SETPRIO 1, 32
+
+# CHECK: *** Bad machine code: S_SETPRIO mask contains invalid memory operation bits ***
+# CHECK: - instruction: S_SETPRIO 1, 64
+
+# CHECK: *** Bad machine code: S_SETPRIO mask contains invalid memory operation bits ***
+# CHECK: - instruction: S_SETPRIO 1, 128
+
+# CHECK: *** Bad machine code: S_SETPRIO mask contains invalid memory operation bits ***
+# CHECK: - instruction: S_SETPRIO 1, 256
+
+# CHECK: *** Bad machine code: S_SETPRIO mask contains invalid memory operation bits ***
+# CHECK: - instruction: S_SETPRIO 1, 512
+
+# CHECK: *** Bad machine code: S_SETPRIO mask contains invalid memory operation bits ***
+# CHECK: - instruction: S_SETPRIO 1, 2048
+
+---
+name: test_invalid_masks
+body: |
+ bb.0:
+ S_SETPRIO 1, 16 ; vmem
+ S_SETPRIO 1, 32 ; vmem-read
+ S_SETPRIO 1, 64 ; vmem-write
+ S_SETPRIO 1, 128 ; ds
+ S_SETPRIO 1, 256 ; ds-read
+ S_SETPRIO 1, 512 ; ds-write
+ S_SETPRIO 1, 2048 ; ldsdma
+
+ S_ENDPGM 0
+...
More information about the llvm-commits
mailing list