[llvm] [AMDGPU] Don't assume a hazard for empty inline asm (PR #223526)
Lixun Zhang via llvm-commits
llvm-commits at lists.llvm.org
Mon Sep 14 14:48:03 PDT 2026
https://github.com/zhanglx13 updated https://github.com/llvm/llvm-project/pull/223526
>From 2605f09a659607cffeac13d5d43f0e49b164aaee Mon Sep 17 00:00:00 2001
From: Lixun Zhang <lixun.zhang at amd.com>
Date: Mon, 14 Sep 2026 21:01:13 +0000
Subject: [PATCH] [AMDGPU] Don't assume a hazard for empty inline asm
An inline asm whose asm string is empty or only whitespace emits no
instruction. Frontends use one as a pure register-class constraint on a
tied operand (e.g. a `"=a,0"` constraint that keeps an MFMA accumulator
in AGPRs), so it cannot produce a hazard of its own. Any hazard carried
by the value is still found at the instruction that produced it, since
wait-state counting already looks past inline asm.
Don't assume a dst-sel forwarding hazard for such inline asm in
`checkVALUHazards` and `checkInlineAsmHazards`, and don't let it take a
slot in the scheduler's lookahead window, where it could push that
producer out of view.
For a Triton GEMM kernel with its MFMA accumulators pinned this way,
this removes about 90 `s_nop` from the loop.
---
.../lib/Target/AMDGPU/GCNHazardRecognizer.cpp | 21 +++-
.../inlineasm-empty-no-hazard-lookahead.ll | 45 +++++++
.../AMDGPU/inlineasm-empty-no-hazard-mfma.ll | 63 ++++++++++
.../AMDGPU/inlineasm-empty-no-hazard.ll | 119 ++++++++++++++++++
4 files changed, 245 insertions(+), 3 deletions(-)
create mode 100644 llvm/test/CodeGen/AMDGPU/inlineasm-empty-no-hazard-lookahead.ll
create mode 100644 llvm/test/CodeGen/AMDGPU/inlineasm-empty-no-hazard-mfma.ll
create mode 100644 llvm/test/CodeGen/AMDGPU/inlineasm-empty-no-hazard.ll
diff --git a/llvm/lib/Target/AMDGPU/GCNHazardRecognizer.cpp b/llvm/lib/Target/AMDGPU/GCNHazardRecognizer.cpp
index 33c367bbcd748..2958ffa99e809 100644
--- a/llvm/lib/Target/AMDGPU/GCNHazardRecognizer.cpp
+++ b/llvm/lib/Target/AMDGPU/GCNHazardRecognizer.cpp
@@ -74,6 +74,15 @@ static cl::opt<bool> EnableWMMAVnopHoisting(
static bool shouldRunLdsBranchVmemWARHazardFixup(const MachineFunction &MF,
const GCNSubtarget &ST);
+/// An empty inline asm emits nothing, so it has no hazard of its own. A hazard
+/// on its tied operand is still found at the instruction that produced it.
+static bool isEmptyInlineAsm(const MachineInstr &MI) {
+ if (!MI.isInlineAsm())
+ return false;
+ const char *Asm = MI.getOperand(InlineAsm::MIOp_AsmString).getSymbolName();
+ return !Asm || StringRef(Asm).trim().empty();
+}
+
GCNHazardRecognizer::GCNHazardRecognizer(
const MachineFunction &MF, GCNHazardRecognizer::OperatingMode Mode,
MachineLoopInfo *MLI)
@@ -833,7 +842,9 @@ void GCNHazardRecognizer::AdvanceCycle() {
}
unsigned NumWaitStates = TII.getNumWaitStates(*CurrCycleInstr);
- if (!NumWaitStates) {
+ // An empty inline asm emits nothing, so it must not take a slot in the
+ // lookahead window either.
+ if (!NumWaitStates || isEmptyInlineAsm(*CurrCycleInstr)) {
CurrCycleInstr = nullptr;
return;
}
@@ -1564,7 +1575,7 @@ int GCNHazardRecognizer::checkVALUHazards(MachineInstr *VALU) const {
return consumesDstSelForwardingOperand(VALU, ForwardedDst, TRI);
}
- if (ProducerMI.isInlineAsm()) {
+ if (ProducerMI.isInlineAsm() && !isEmptyInlineAsm(ProducerMI)) {
// Assume inline asm has dst forwarding hazard
for (auto &Def : ProducerMI.all_defs()) {
if (consumesDstSelForwardingOperand(VALU, &Def, TRI))
@@ -1668,6 +1679,10 @@ int GCNHazardRecognizer::checkInlineAsmHazards(MachineInstr *IA) const {
!ST.hasCvtScaleForwardingHazard())
return 0;
+ // An empty inline asm executes nothing, so it cannot consume a hazard.
+ if (isEmptyInlineAsm(*IA))
+ return 0;
+
const MachineRegisterInfo &MRI = MF.getRegInfo();
int WaitStatesNeeded = 0;
@@ -1694,7 +1709,7 @@ int GCNHazardRecognizer::checkInlineAsmHazards(MachineInstr *IA) const {
return IA->modifiesRegister(Dst->getReg(), &TRI) ||
IA->readsRegister(Dst->getReg(), &TRI);
- if (ProducerMI.isInlineAsm()) {
+ if (ProducerMI.isInlineAsm() && !isEmptyInlineAsm(ProducerMI)) {
// If MI is inline asm, assume it has dst forwarding hazard
for (auto &Def : ProducerMI.all_defs()) {
if (IA->modifiesRegister(Def.getReg(), &TRI) ||
diff --git a/llvm/test/CodeGen/AMDGPU/inlineasm-empty-no-hazard-lookahead.ll b/llvm/test/CodeGen/AMDGPU/inlineasm-empty-no-hazard-lookahead.ll
new file mode 100644
index 0000000000000..a0f6701a6a862
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/inlineasm-empty-no-hazard-lookahead.ll
@@ -0,0 +1,45 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
+; RUN: llc -mtriple=amdgpu9.42 -O3 --enable-misched=0 -o - %s | FileCheck %s
+; RUN: llc -mtriple=amdgpu9.50 -O3 --enable-misched=0 -o - %s | FileCheck %s
+
+; Empty inline asm takes no slot in the scheduler's lookahead window, so the
+; dst-sel producer behind five of them is still seen, and the scheduler fills
+; the wait with the independent add instead of an s_nop.
+
+declare i32 @llvm.amdgcn.cvt.scalef32.sr.fp8.f32(i32 %old, float %src, i32 %seed, float %scale, i32 %dst_sel)
+
+define amdgpu_ps void @hazard_behind_five_pins(ptr addrspace(1) %out, ptr addrspace(1) %out2, ptr addrspace(3) %lds, float %src, i32 %seed, float %scale, i32 %x, i32 %y) {
+; CHECK-LABEL: hazard_behind_five_pins:
+; CHECK: ; %bb.0:
+; CHECK-NEXT: global_load_dword v0, v[0:1], off
+; CHECK-NEXT: s_waitcnt vmcnt(0)
+; CHECK-NEXT: v_mov_b32_e32 v1, v0
+; CHECK-NEXT: v_cvt_scalef32_sr_fp8_f32 v1, v5, v6, v7 op_sel:[0,0,0,1]
+; CHECK-NEXT: ;;#ASMSTART
+; CHECK-NEXT: ;;#ASMEND
+; CHECK-NEXT: ;;#ASMSTART
+; CHECK-NEXT: ;;#ASMEND
+; CHECK-NEXT: ;;#ASMSTART
+; CHECK-NEXT: ;;#ASMEND
+; CHECK-NEXT: ;;#ASMSTART
+; CHECK-NEXT: ;;#ASMEND
+; CHECK-NEXT: ;;#ASMSTART
+; CHECK-NEXT: ;;#ASMEND
+; CHECK-NEXT: v_add_u32_e32 v0, v0, v9
+; CHECK-NEXT: v_add_u32_e32 v1, v1, v8
+; CHECK-NEXT: ds_write_b32 v4, v1
+; CHECK-NEXT: global_store_dword v[2:3], v0, off
+; CHECK-NEXT: s_endpgm
+ %old = load i32, ptr addrspace(1) %out, align 4
+ %cvt = call i32 @llvm.amdgcn.cvt.scalef32.sr.fp8.f32(i32 %old, float %src, i32 %seed, float %scale, i32 2)
+ %p0 = call i32 asm "", "=v,0"(i32 %cvt)
+ %p1 = call i32 asm "", "=v,0"(i32 %p0)
+ %p2 = call i32 asm "", "=v,0"(i32 %p1)
+ %p3 = call i32 asm "", "=v,0"(i32 %p2)
+ %p4 = call i32 asm "", "=v,0"(i32 %p3)
+ %sum = add i32 %p4, %x
+ store i32 %sum, ptr addrspace(3) %lds, align 4
+ %other = add i32 %old, %y
+ store i32 %other, ptr addrspace(1) %out2, align 4
+ ret void
+}
diff --git a/llvm/test/CodeGen/AMDGPU/inlineasm-empty-no-hazard-mfma.ll b/llvm/test/CodeGen/AMDGPU/inlineasm-empty-no-hazard-mfma.ll
new file mode 100644
index 0000000000000..7b3160722517a
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/inlineasm-empty-no-hazard-mfma.ll
@@ -0,0 +1,63 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
+; RUN: llc -mtriple=amdgpu9.50 -O3 -o - %s | FileCheck %s
+
+; No s_nop between an empty "=a,0" inline asm and the MFMA that reads it.
+
+declare <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half>, <8 x half>, <4 x float>, i32, i32, i32)
+
+define amdgpu_kernel void @empty_pin(ptr addrspace(3) %a, ptr addrspace(3) %b, ptr addrspace(3) %out) {
+; CHECK-LABEL: empty_pin:
+; CHECK: ; %bb.0:
+; CHECK-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x24
+; CHECK-NEXT: s_waitcnt lgkmcnt(0)
+; CHECK-NEXT: v_mov_b32_e32 v0, s0
+; CHECK-NEXT: ds_read_b128 v[4:7], v0
+; CHECK-NEXT: v_mov_b32_e32 v0, s1
+; CHECK-NEXT: ds_read_b128 v[8:11], v0
+; CHECK-NEXT: s_waitcnt lgkmcnt(0)
+; CHECK-NEXT: v_mfma_f32_16x16x32_f16 a[0:3], v[4:7], v[8:11], 0
+; CHECK-NEXT: ;;#ASMSTART
+; CHECK-NEXT: ;;#ASMEND
+; CHECK-NEXT: v_mfma_f32_16x16x32_f16 a[0:3], v[4:7], v[8:11], a[0:3]
+; CHECK-NEXT: v_mov_b32_e32 v4, s2
+; CHECK-NEXT: s_nop 6
+; CHECK-NEXT: ds_write_b128 v4, a[0:3]
+; CHECK-NEXT: s_endpgm
+ %va = load <8 x half>, ptr addrspace(3) %a, align 16
+ %vb = load <8 x half>, ptr addrspace(3) %b, align 16
+ %m0 = tail call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half> %va, <8 x half> %vb, <4 x float> zeroinitializer, i32 0, i32 0, i32 0)
+ %pin = tail call <4 x float> asm "", "=a,0"(<4 x float> %m0)
+ %m1 = tail call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half> %va, <8 x half> %vb, <4 x float> %pin, i32 0, i32 0, i32 0)
+ store <4 x float> %m1, ptr addrspace(3) %out, align 16
+ ret void
+}
+
+; A comment-only inline asm keeps its s_nop.
+define amdgpu_kernel void @comment_only_asm(ptr addrspace(3) %a, ptr addrspace(3) %b, ptr addrspace(3) %out) {
+; CHECK-LABEL: comment_only_asm:
+; CHECK: ; %bb.0:
+; CHECK-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x24
+; CHECK-NEXT: s_waitcnt lgkmcnt(0)
+; CHECK-NEXT: v_mov_b32_e32 v0, s0
+; CHECK-NEXT: ds_read_b128 v[4:7], v0
+; CHECK-NEXT: v_mov_b32_e32 v0, s1
+; CHECK-NEXT: ds_read_b128 v[8:11], v0
+; CHECK-NEXT: s_waitcnt lgkmcnt(0)
+; CHECK-NEXT: v_mfma_f32_16x16x32_f16 a[0:3], v[4:7], v[8:11], 0
+; CHECK-NEXT: ;;#ASMSTART
+; CHECK-NEXT: ; def a[0:3]
+; CHECK-NEXT: ;;#ASMEND
+; CHECK-NEXT: s_nop 0
+; CHECK-NEXT: v_mfma_f32_16x16x32_f16 a[0:3], v[4:7], v[8:11], a[0:3]
+; CHECK-NEXT: v_mov_b32_e32 v4, s2
+; CHECK-NEXT: s_nop 6
+; CHECK-NEXT: ds_write_b128 v4, a[0:3]
+; CHECK-NEXT: s_endpgm
+ %va = load <8 x half>, ptr addrspace(3) %a, align 16
+ %vb = load <8 x half>, ptr addrspace(3) %b, align 16
+ %m0 = tail call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half> %va, <8 x half> %vb, <4 x float> zeroinitializer, i32 0, i32 0, i32 0)
+ %pin = tail call <4 x float> asm "; def $0", "=a,0"(<4 x float> %m0)
+ %m1 = tail call <4 x float> @llvm.amdgcn.mfma.f32.16x16x32.f16(<8 x half> %va, <8 x half> %vb, <4 x float> %pin, i32 0, i32 0, i32 0)
+ store <4 x float> %m1, ptr addrspace(3) %out, align 16
+ ret void
+}
diff --git a/llvm/test/CodeGen/AMDGPU/inlineasm-empty-no-hazard.ll b/llvm/test/CodeGen/AMDGPU/inlineasm-empty-no-hazard.ll
new file mode 100644
index 0000000000000..d38ef7755292b
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/inlineasm-empty-no-hazard.ll
@@ -0,0 +1,119 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
+; RUN: llc -mtriple=amdgpu9.42 -o - %s | FileCheck %s
+; RUN: llc -mtriple=amdgpu9.50 -o - %s | FileCheck %s
+
+; An empty inline asm has no hazard of its own; any other inline asm keeps one.
+
+declare i32 @llvm.amdgcn.cvt.scalef32.sr.fp8.f32(i32 %old, float %src, i32 %seed, float %scale, i32 %dst_sel)
+
+; No s_nop after an empty inline asm.
+define amdgpu_ps void @empty_pin(ptr addrspace(1) %out, i32 %x, i32 %y, i32 %z) {
+; CHECK-LABEL: empty_pin:
+; CHECK: ; %bb.0:
+; CHECK-NEXT: v_add_u32_e32 v2, v2, v3
+; CHECK-NEXT: ;;#ASMSTART
+; CHECK-NEXT: ;;#ASMEND
+; CHECK-NEXT: v_add_u32_e32 v2, v2, v4
+; CHECK-NEXT: global_store_dword v[0:1], v2, off
+; CHECK-NEXT: s_endpgm
+ %a = add i32 %x, %y
+ %pin = call i32 asm "", "=v,0"(i32 %a)
+ %b = add i32 %pin, %z
+ store i32 %b, ptr addrspace(1) %out, align 4
+ ret void
+}
+
+; A real dst-sel hazard across an empty inline asm is still padded, once.
+define amdgpu_ps void @real_hazard_through_empty_pin(ptr addrspace(1) %out, float %src, i32 %seed, float %scale, i32 %x) {
+; CHECK-LABEL: real_hazard_through_empty_pin:
+; CHECK: ; %bb.0:
+; CHECK-NEXT: global_load_dword v6, v[0:1], off
+; CHECK-NEXT: s_waitcnt vmcnt(0)
+; CHECK-NEXT: v_cvt_scalef32_sr_fp8_f32 v6, v2, v3, v4 op_sel:[0,0,0,1]
+; CHECK-NEXT: ;;#ASMSTART
+; CHECK-NEXT: ;;#ASMEND
+; CHECK-NEXT: s_nop 0
+; CHECK-NEXT: v_add_u32_e32 v2, v6, v5
+; CHECK-NEXT: global_store_dword v[0:1], v2, off
+; CHECK-NEXT: s_endpgm
+ %old = load i32, ptr addrspace(1) %out, align 4
+ %cvt = call i32 @llvm.amdgcn.cvt.scalef32.sr.fp8.f32(i32 %old, float %src, i32 %seed, float %scale, i32 2)
+ %pin = call i32 asm "", "=v,0"(i32 %cvt)
+ %sum = add i32 %pin, %x
+ store i32 %sum, ptr addrspace(1) %out, align 4
+ ret void
+}
+
+; Non-empty inline asm keeps its s_nop.
+define amdgpu_ps void @nonempty_asm(ptr addrspace(1) %out, i32 %x, i32 %y, i32 %z) {
+; CHECK-LABEL: nonempty_asm:
+; CHECK: ; %bb.0:
+; CHECK-NEXT: v_add_u32_e32 v2, v2, v3
+; CHECK-NEXT: ;;#ASMSTART
+; CHECK-NEXT: v_mov_b32_e32 v2, v2
+; CHECK-NEXT: ;;#ASMEND
+; CHECK-NEXT: s_nop 0
+; CHECK-NEXT: v_add_u32_e32 v2, v2, v4
+; CHECK-NEXT: global_store_dword v[0:1], v2, off
+; CHECK-NEXT: s_endpgm
+ %a = add i32 %x, %y
+ %pin = call i32 asm "v_mov_b32_e32 $0, $1", "=v,0"(i32 %a)
+ %b = add i32 %pin, %z
+ store i32 %b, ptr addrspace(1) %out, align 4
+ ret void
+}
+
+; A comment-only inline asm is not empty.
+define amdgpu_ps void @comment_only_asm(ptr addrspace(1) %out, i32 %x, i32 %y, i32 %z) {
+; CHECK-LABEL: comment_only_asm:
+; CHECK: ; %bb.0:
+; CHECK-NEXT: v_add_u32_e32 v2, v2, v3
+; CHECK-NEXT: ;;#ASMSTART
+; CHECK-NEXT: ; def v2
+; CHECK-NEXT: ;;#ASMEND
+; CHECK-NEXT: s_nop 0
+; CHECK-NEXT: v_add_u32_e32 v2, v2, v4
+; CHECK-NEXT: global_store_dword v[0:1], v2, off
+; CHECK-NEXT: s_endpgm
+ %a = add i32 %x, %y
+ %pin = call i32 asm "; def $0", "=v,0"(i32 %a)
+ %b = add i32 %pin, %z
+ store i32 %b, ptr addrspace(1) %out, align 4
+ ret void
+}
+
+; No s_nop between an empty inline asm and a non-empty one that reads it.
+define amdgpu_ps void @asm_after_empty_pin(ptr addrspace(1) %out, i32 %x, i32 %y) {
+; CHECK-LABEL: asm_after_empty_pin:
+; CHECK: ; %bb.0:
+; CHECK-NEXT: v_add_u32_e32 v2, v2, v3
+; CHECK-NEXT: ;;#ASMSTART
+; CHECK-NEXT: ;;#ASMEND
+; CHECK-NEXT: ;;#ASMSTART
+; CHECK-NEXT: v_mov_b32_e32 v2, v2
+; CHECK-NEXT: ;;#ASMEND
+; CHECK-NEXT: global_store_dword v[0:1], v2, off
+; CHECK-NEXT: s_endpgm
+ %a = add i32 %x, %y
+ %pin = call i32 asm "", "=v,0"(i32 %a)
+ %mov = call i32 asm "v_mov_b32_e32 $0, $1", "=v,0"(i32 %pin)
+ store i32 %mov, ptr addrspace(1) %out, align 4
+ ret void
+}
+
+; UTC_ARGS: --disable
+; A whitespace-only inline asm is empty too.
+define amdgpu_ps void @whitespace_pin(ptr addrspace(1) %out, i32 %x, i32 %y, i32 %z) {
+; CHECK-LABEL: whitespace_pin:
+; CHECK: v_add_u32_e32 v2, v2, v3
+; CHECK-NEXT: ;;#ASMSTART
+; CHECK-NEXT: {{^[[:blank:]]+$}}
+; CHECK-NEXT: ;;#ASMEND
+; CHECK-NEXT: v_add_u32_e32 v2, v2, v4
+ %a = add i32 %x, %y
+ %pin = call i32 asm " ", "=v,0"(i32 %a)
+ %b = add i32 %pin, %z
+ store i32 %b, ptr addrspace(1) %out, align 4
+ ret void
+}
+; UTC_ARGS: --enable
More information about the llvm-commits
mailing list