[llvm] [AMDGPU] Account for inline asm size in inst_pref_size calculation (PR #192133)
Adel Ejjeh via llvm-commits
llvm-commits at lists.llvm.org
Tue Apr 14 13:59:07 PDT 2026
https://github.com/adelejjeh created https://github.com/llvm/llvm-project/pull/192133
SIProgramInfo::getFunctionCodeSize() with IsLowerBound=true was completely skipping inline assembly instructions, treating them as zero bytes. This caused amdhsa_inst_pref_size to be severely underestimated for kernels containing inline asm.
Override getInlineAsmLength() in SIInstrInfo to provide a lower-bound estimate using MinInstAlignment (4 bytes) per instruction line instead of MaxInstLength. Handle .space, .fill, and data-emitting directives (.byte/.short/.long/.quad/.zero) with exact sizes. Skip comments, labels, and non-emitting directives.
Update getFunctionCodeSize() to call the new getInlineAsmLength() with IsLowerBound forwarded, instead of skipping inline asm entirely.
Update the existing inst-prefetch-hint.ll lit test and add a new inst-prefetch-inline-asm.ll lit test.
>From 280fe34f888925b84ff5dcd97ed898a0ae7a89c9 Mon Sep 17 00:00:00 2001
From: Adel Ejjeh <adel.ejjeh at amd.com>
Date: Tue, 14 Apr 2026 15:38:38 -0500
Subject: [PATCH] [AMDGPU] Account for inline asm size in inst_pref_size
calculation
SIProgramInfo::getFunctionCodeSize() with IsLowerBound=true was
completely skipping inline assembly instructions, treating them as
zero bytes. This caused amdhsa_inst_pref_size to be severely
underestimated for kernels containing inline asm, defeating
instruction prefetch on gfx11+.
Override getInlineAsmLength() in SIInstrInfo to provide a lower-bound
estimate using MinInstAlignment (4 bytes) per instruction line instead
of MaxInstLength. Handle .space, .fill, and data-emitting directives
(.byte/.short/.long/.quad/.zero) with exact sizes. Skip comments,
labels, and non-emitting directives.
Update getFunctionCodeSize() to call the new getInlineAsmLength()
with IsLowerBound forwarded, instead of skipping inline asm entirely.
---
llvm/lib/Target/AMDGPU/SIInstrInfo.cpp | 130 ++++++++++++++++++
llvm/lib/Target/AMDGPU/SIInstrInfo.h | 11 ++
llvm/lib/Target/AMDGPU/SIProgramInfo.cpp | 19 ++-
.../test/CodeGen/AMDGPU/inst-prefetch-hint.ll | 6 +-
.../AMDGPU/inst-prefetch-inline-asm.ll | 120 ++++++++++++++++
5 files changed, 280 insertions(+), 6 deletions(-)
create mode 100644 llvm/test/CodeGen/AMDGPU/inst-prefetch-inline-asm.ll
diff --git a/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp b/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp
index bc3052b139d18..055f2e08433a3 100644
--- a/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/SIInstrInfo.cpp
@@ -32,6 +32,7 @@
#include "llvm/CodeGen/ScheduleDAG.h"
#include "llvm/IR/DiagnosticInfo.h"
#include "llvm/IR/IntrinsicsAMDGPU.h"
+#include "llvm/MC/MCAsmInfo.h"
#include "llvm/MC/MCContext.h"
#include "llvm/Support/CommandLine.h"
#include "llvm/Target/TargetMachine.h"
@@ -9911,6 +9912,135 @@ unsigned SIInstrInfo::getInstSizeInBytes(const MachineInstr &MI) const {
}
}
+unsigned SIInstrInfo::getInlineAsmLength(const char *Str, const MCAsmInfo &MAI,
+ const TargetSubtargetInfo *STI) const {
+ return getInlineAsmLength(Str, MAI, STI, /*IsLowerBound=*/false);
+}
+
+unsigned SIInstrInfo::getInlineAsmLength(const char *Str, const MCAsmInfo &MAI,
+ const TargetSubtargetInfo *STI,
+ bool IsLowerBound) const {
+ if (!IsLowerBound)
+ return TargetInstrInfo::getInlineAsmLength(Str, MAI, STI);
+
+ // For a lower-bound estimate, count each non-comment, non-directive
+ // instruction line as MinInstAlignment bytes (4 for AMDGPU), and parse
+ // data-emitting directives for their exact sizes.
+ //
+ // TODO: For more accurate inline asm size estimation, consider assembling
+ // the asm string through the MC layer (MCAsmParser -> MCCodeEmitter) to get
+ // exact encoded sizes. This would automatically handle new instructions via
+ // TableGen and correctly size instructions with literal operands. See
+ // mlir::ROCDL::assembleIsa() for an AMDGPU-specific example of this
+ // approach. The main challenges are: (1) operand placeholders ($0, $1) need
+ // substitution with dummy registers before parsing, and (2) MC pipeline
+ // creation overhead per inline asm block (mitigable via caching).
+ const unsigned MinInstSize = MAI.getMinInstAlignment();
+ unsigned Length = 0;
+ bool AtInsnStart = true;
+
+ for (; *Str; ++Str) {
+ if (*Str == '\n' || strncmp(Str, MAI.getSeparatorString(),
+ strlen(MAI.getSeparatorString())) == 0) {
+ AtInsnStart = true;
+ continue;
+ }
+
+ if (strncmp(Str, MAI.getCommentString().data(),
+ MAI.getCommentString().size()) == 0) {
+ // Comment — skip rest of line.
+ AtInsnStart = false;
+ continue;
+ }
+
+ if (AtInsnStart && !isSpace(static_cast<unsigned char>(*Str))) {
+ AtInsnStart = false;
+
+ // Handle .space directive.
+ if (strncmp(Str, ".space", 6) == 0) {
+ char *End;
+ int SpaceSize = strtol(Str + 6, &End, 10);
+ SpaceSize = SpaceSize < 0 ? 0 : SpaceSize;
+ while (*End != '\n' && *End != '\0' &&
+ isSpace(static_cast<unsigned char>(*End)))
+ ++End;
+ if (*End == '\0' || *End == '\n' ||
+ strncmp(End, MAI.getCommentString().data(),
+ MAI.getCommentString().size()) == 0)
+ Length += SpaceSize;
+ else
+ Length += MinInstSize;
+ continue;
+ }
+
+ // Handle .fill directive: .fill count [, size [, value]]
+ if (strncmp(Str, ".fill", 5) == 0 &&
+ isSpace(static_cast<unsigned char>(Str[5]))) {
+ char *End;
+ long Count = strtol(Str + 5, &End, 10);
+ Count = Count < 0 ? 0 : Count;
+ long Size = 1; // default size for .fill
+ while (*End != '\0' && isSpace(static_cast<unsigned char>(*End)))
+ ++End;
+ if (*End == ',') {
+ ++End;
+ Size = strtol(End, &End, 10);
+ Size = Size < 0 ? 0 : Size;
+ }
+ Length += Count * Size;
+ continue;
+ }
+
+ // Handle data directives.
+ if (strncmp(Str, ".byte", 5) == 0) {
+ Length += 1;
+ continue;
+ }
+ if (strncmp(Str, ".short", 6) == 0 || strncmp(Str, ".hword", 6) == 0) {
+ Length += 2;
+ continue;
+ }
+ if (strncmp(Str, ".long", 5) == 0 || strncmp(Str, ".word", 5) == 0) {
+ Length += 4;
+ continue;
+ }
+ if (strncmp(Str, ".quad", 5) == 0) {
+ Length += 8;
+ continue;
+ }
+ if (strncmp(Str, ".zero", 5) == 0) {
+ char *End;
+ int ZeroSize = strtol(Str + 5, &End, 10);
+ ZeroSize = ZeroSize < 0 ? 0 : ZeroSize;
+ Length += ZeroSize;
+ continue;
+ }
+
+ // Skip known non-emitting directives.
+ if (*Str == '.') {
+ // Labels, .set, .globl, .local, .type, .size, .section, .align,
+ // .file, .loc, etc. — these don't emit code bytes.
+ continue;
+ }
+
+ // Skip labels (e.g. "my_label:").
+ {
+ const char *P = Str;
+ while (*P && *P != '\n' && *P != ':' &&
+ !isSpace(static_cast<unsigned char>(*P)))
+ ++P;
+ if (*P == ':')
+ continue;
+ }
+
+ // Default: assume it's an instruction — count minimum size.
+ Length += MinInstSize;
+ }
+ }
+
+ return Length;
+}
+
bool SIInstrInfo::mayAccessFlatAddressSpace(const MachineInstr &MI) const {
if (!isFLAT(MI))
return false;
diff --git a/llvm/lib/Target/AMDGPU/SIInstrInfo.h b/llvm/lib/Target/AMDGPU/SIInstrInfo.h
index 3c1232ac098a0..1fe0958c225bf 100644
--- a/llvm/lib/Target/AMDGPU/SIInstrInfo.h
+++ b/llvm/lib/Target/AMDGPU/SIInstrInfo.h
@@ -1582,6 +1582,17 @@ class SIInstrInfo final : public AMDGPUGenInstrInfo {
unsigned getInstBundleSize(const MachineInstr &MI) const;
unsigned getInstSizeInBytes(const MachineInstr &MI) const override;
+ unsigned getInlineAsmLength(const char *Str, const MCAsmInfo &MAI,
+ const TargetSubtargetInfo *STI) const override;
+
+ /// Compute the size of inline assembly with an option for lower-bound
+ /// estimation. When \p IsLowerBound is true, uses the minimum instruction
+ /// size (4 bytes for AMDGPU) per instruction line, giving a conservative
+ /// lower-bound estimate suitable for instruction prefetch sizing.
+ unsigned getInlineAsmLength(const char *Str, const MCAsmInfo &MAI,
+ const TargetSubtargetInfo *STI,
+ bool IsLowerBound) const;
+
bool mayAccessFlatAddressSpace(const MachineInstr &MI) const;
std::pair<unsigned, unsigned>
diff --git a/llvm/lib/Target/AMDGPU/SIProgramInfo.cpp b/llvm/lib/Target/AMDGPU/SIProgramInfo.cpp
index a3f261b87e80b..20471ceb1b5ee 100644
--- a/llvm/lib/Target/AMDGPU/SIProgramInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/SIProgramInfo.cpp
@@ -18,7 +18,9 @@
#include "GCNSubtarget.h"
#include "SIDefines.h"
#include "Utils/AMDGPUBaseInfo.h"
+#include "llvm/MC/MCAsmInfo.h"
#include "llvm/MC/MCExpr.h"
+#include "llvm/Target/TargetMachine.h"
using namespace llvm;
@@ -213,6 +215,16 @@ uint64_t SIProgramInfo::getFunctionCodeSize(const MachineFunction &MF,
uint64_t CodeSize = 0;
+ // TODO: For more accurate inline asm size estimation, consider assembling
+ // the asm string through the MC layer (MCAsmParser -> MCCodeEmitter) to get
+ // exact encoded sizes. This would automatically handle new instructions via
+ // TableGen and correctly size instructions with literal operands. See
+ // mlir::ROCDL::assembleIsa() for an AMDGPU-specific example of this
+ // approach. The main challenges are: (1) operand placeholders ($0, $1) need
+ // substitution with dummy registers before parsing, and (2) MC pipeline
+ // creation overhead per inline asm block (mitigable via caching).
+ const MCAsmInfo *MAI = MF.getTarget().getMCAsmInfo();
+
for (const MachineBasicBlock &MBB : MF) {
// The amount of padding to align code can be both underestimated and
// overestimated. In case of inline asm used getInstSizeInBytes() will
@@ -227,10 +239,11 @@ uint64_t SIProgramInfo::getFunctionCodeSize(const MachineFunction &MF,
if (MI.isMetaInstruction())
continue;
- // We cannot properly estimate inline asm size. It can be as small as zero
- // if that is just a comment.
- if (IsLowerBound && MI.isInlineAsm())
+ if (MI.isInlineAsm()) {
+ const char *AsmStr = MI.getOperand(0).getSymbolName();
+ CodeSize += TII->getInlineAsmLength(AsmStr, *MAI, &STM, IsLowerBound);
continue;
+ }
CodeSize += TII->getInstSizeInBytes(MI);
}
diff --git a/llvm/test/CodeGen/AMDGPU/inst-prefetch-hint.ll b/llvm/test/CodeGen/AMDGPU/inst-prefetch-hint.ll
index 580167076e1f0..b44fad5523a4c 100644
--- a/llvm/test/CodeGen/AMDGPU/inst-prefetch-hint.ll
+++ b/llvm/test/CodeGen/AMDGPU/inst-prefetch-hint.ll
@@ -20,11 +20,11 @@ bb:
ret void
}
-; Ignore inline asm in size calculation
+; Account for inline asm in size calculation
; GCN-LABEL: .amdhsa_kernel inline_asm
-; GCN: .amdhsa_inst_pref_size 1
-; GCN: codeLenInByte = {{[0-9]$}}
+; GCN: .amdhsa_inst_pref_size 9
+; GCN: codeLenInByte = 1028
define amdgpu_kernel void @inline_asm() {
bb:
call void asm sideeffect ".fill 256, 4, 0", ""()
diff --git a/llvm/test/CodeGen/AMDGPU/inst-prefetch-inline-asm.ll b/llvm/test/CodeGen/AMDGPU/inst-prefetch-inline-asm.ll
new file mode 100644
index 0000000000000..2d8b03c6362e0
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/inst-prefetch-inline-asm.ll
@@ -0,0 +1,120 @@
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1100 < %s | FileCheck %s
+
+;; Test that inline assembly is properly accounted for in
+;; amdhsa_inst_pref_size calculation. Each kernel exercises a specific
+;; directive or pattern handled by SIInstrInfo::getInlineAsmLength()
+;; in lower-bound mode.
+
+; --- .fill directive: .fill count, size, value => count * size bytes ---
+
+; CHECK-LABEL: .amdhsa_kernel test_fill
+; CHECK: .amdhsa_inst_pref_size 9
+; CHECK: codeLenInByte = 1028
+define amdgpu_kernel void @test_fill() {
+ call void asm sideeffect ".fill 256, 4, 0", ""()
+ ret void
+}
+
+; --- .fill with size=1: .fill 1024, 1, 0 => 1024 bytes ---
+
+; CHECK-LABEL: .amdhsa_kernel test_fill_size1
+; CHECK: .amdhsa_inst_pref_size 9
+; CHECK: codeLenInByte = 1028
+define amdgpu_kernel void @test_fill_size1() {
+ call void asm sideeffect ".fill 1024, 1, 0", ""()
+ ret void
+}
+
+; --- .space directive: .space N => N bytes ---
+
+; CHECK-LABEL: .amdhsa_kernel test_space
+; CHECK: .amdhsa_inst_pref_size 9
+; CHECK: codeLenInByte = 1028
+define amdgpu_kernel void @test_space() {
+ call void asm sideeffect ".space 1024", ""()
+ ret void
+}
+
+; --- .zero directive: .zero N => N zero bytes ---
+
+; CHECK-LABEL: .amdhsa_kernel test_zero
+; CHECK: .amdhsa_inst_pref_size 9
+; CHECK: codeLenInByte = 1028
+define amdgpu_kernel void @test_zero() {
+ call void asm sideeffect ".zero 1024", ""()
+ ret void
+}
+
+; --- Data directives: .byte (1) + .short (2) + .long (4) + .quad (8) = 15 bytes ---
+;; Lower bound estimate: 15 + 4 (s_endpgm) = 19
+
+; CHECK-LABEL: .amdhsa_kernel test_data_directives
+; CHECK: .amdhsa_inst_pref_size 1
+; CHECK: codeLenInByte = 19
+define amdgpu_kernel void @test_data_directives() {
+ call void asm sideeffect ".byte 0\0A.short 0\0A.long 0\0A.quad 0", ""()
+ ret void
+}
+
+; --- .hword and .word aliases: .hword (2) + .word (4) = 6 bytes ---
+
+; CHECK-LABEL: .amdhsa_kernel test_data_aliases
+; CHECK: .amdhsa_inst_pref_size 1
+; CHECK: codeLenInByte = 10
+define amdgpu_kernel void @test_data_aliases() {
+ call void asm sideeffect ".hword 0\0A.word 0", ""()
+ ret void
+}
+
+; --- Instructions: s_nop lines counted at MinInstAlignment (4 bytes) each ---
+
+; CHECK-LABEL: .amdhsa_kernel test_instructions
+; CHECK: .amdhsa_inst_pref_size 8
+; CHECK: codeLenInByte = 1020
+define amdgpu_kernel void @test_instructions() {
+ call void asm sideeffect "s_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0\0As_nop 0", ""()
+ ret void
+}
+
+; --- Comments should not contribute to size ---
+;; Only s_endpgm (4 bytes) — same as an empty kernel.
+
+; CHECK-LABEL: .amdhsa_kernel test_comments
+; CHECK: .amdhsa_inst_pref_size 1
+; CHECK: codeLenInByte = 4
+define amdgpu_kernel void @test_comments() {
+ call void asm sideeffect "; comment 1\0A; comment 2\0A; comment 3", ""()
+ ret void
+}
+
+; --- Labels should not contribute to size ---
+
+; CHECK-LABEL: .amdhsa_kernel test_labels
+; CHECK: .amdhsa_inst_pref_size 1
+; CHECK: codeLenInByte = 8
+define amdgpu_kernel void @test_labels() {
+ call void asm sideeffect "my_label:\0As_nop 0", ""()
+ ret void
+}
+
+; --- Non-emitting directives (.set) should not contribute to size ---
+;; Only s_endpgm (4 bytes) — same as an empty kernel.
+
+; CHECK-LABEL: .amdhsa_kernel test_non_emitting
+; CHECK: .amdhsa_inst_pref_size 1
+; CHECK: codeLenInByte = 4
+define amdgpu_kernel void @test_non_emitting() {
+ call void asm sideeffect ".set my_const, 42", ""()
+ ret void
+}
+
+; --- Mixed: .fill (256) + s_nop (4) + comment (0) + .space (128) + s_nop (4) ---
+;; Lower bound estimate: 256 + 4 + 0 + 128 + 4 + 4 (s_endpgm) = 396
+
+; CHECK-LABEL: .amdhsa_kernel test_mixed
+; CHECK: .amdhsa_inst_pref_size 4
+; CHECK: codeLenInByte = 396
+define amdgpu_kernel void @test_mixed() {
+ call void asm sideeffect ".fill 64, 4, 0\0As_nop 0\0A; comment\0A.space 128\0As_nop 0", ""()
+ ret void
+}
More information about the llvm-commits
mailing list