[llvm] [AMDGPU] Account for inline asm size in inst_pref_size calculation (PR #192133)

[Robert Imschweiler via llvm-commits]

================
@@ -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' &&
----------------
ro-i wrote:

maybe some of that could be simplified a lot by using `llvm::StringReg`?

https://github.com/llvm/llvm-project/pull/192133