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 <[email protected]> Date: Thu, 18 Jun 2026 09:58:18 -0400 Subject: [PATCH 1/8] [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 <[email protected]> Date: Mon, 22 Jun 2026 04:10:41 -0400 Subject: [PATCH 2/8] [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 <[email protected]> Date: Mon, 22 Jun 2026 05:47:25 -0400 Subject: [PATCH 3/8] [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 <[email protected]> Date: Tue, 28 Jul 2026 12:16:35 -0400 Subject: [PATCH 4/8] 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 <[email protected]> Date: Wed, 26 Aug 2026 03:15:02 -0400 Subject: [PATCH 5/8] 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 <[email protected]> Date: Wed, 26 Aug 2026 03:16:43 -0400 Subject: [PATCH 6/8] 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 <[email protected]> Date: Wed, 26 Aug 2026 03:17:33 -0400 Subject: [PATCH 7/8] 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 60b1b643b95f42f18a90cc7b4d5f279865c2bab1 Mon Sep 17 00:00:00 2001 From: Frederik Harwath <[email protected]> Date: Wed, 26 Aug 2026 04:49:17 -0400 Subject: [PATCH 8/8] [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..6a661de4b4ffe 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_PSEUDO: { + 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..e1a4acc03ad00 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_PSEUDO : 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@gotpcrel32@lo+4 ; GFX11-NEXT: s_addc_u32 s1, s1, test_export_gfx@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@gotpcrel32@lo+4 ; GFX1150-NEXT: s_addc_u32 s1, s1, test_export_gfx@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 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
