From 2cc4c0834188e2f0d6c6e917f95c2364ce9b8fba Mon Sep 17 00:00:00 2001 From: Tony Davis Date: Fri, 2 Oct 2026 22:36:14 -0400 Subject: [PATCH 01/11] fix(rocisa): AIHPBLAS-5022 split 64-bit add immediates into low and high halves On targets without s_add_u64 or v_add_nc_u64 (gfx90a, gfx942, gfx950), SAddU64 and VAddNCU64 lower to a 32-bit carry chain. splitSrcs copied an immediate into both halves, which is right for packed operations but makes a 64-bit add of imm also add imm * 2^32, moving a pointer by hundreds of GiB (AMD-ROCm-Internal PR 2557). An immediate above 2^31 also reached the assembly as a double. splitSrcs64 gives the high half the sign extension of a 32-bit immediate, splits a larger one as a 64-bit integer, and rejects a symbol it cannot split. Nothing on develop lowers an immediate this way today; the immediate callers are gfx1250 TDM code with a native 64-bit add. test_lowering64.py checks the emitted pair for register and immediate operands (positive, negative, and above 2^31) on gfx90a, gfx942 and gfx950. Tensile.Utilities.address_carry_lint reports a 64-bit address whose low dword is updated without a carry into the high dword, the ROCM-32046 shape. test_address_carry_lint.py checks it on synthetic cases and on kernels generated from designed configs for gfx90a, gfx942, gfx950 and gfx1250. Co-authored-by: Cursor --- .../Tests/unit/test_address_carry_lint.py | 132 ++++++++ .../Tensile/Utilities/address_carry_lint.py | 298 ++++++++++++++++++ .../rocisa/include/instruction/common.hpp | 4 +- .../include/instruction/instruction.hpp | 50 +++ .../rocisa/test/test_lowering64.py | 114 +++++++ 5 files changed, 596 insertions(+), 2 deletions(-) create mode 100644 projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_address_carry_lint.py create mode 100644 projects/hipblaslt/tensilelite/Tensile/Utilities/address_carry_lint.py create mode 100644 projects/hipblaslt/tensilelite/rocisa/test/test_lowering64.py diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_address_carry_lint.py b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_address_carry_lint.py new file mode 100644 index 000000000000..eefe6f7c1de7 --- /dev/null +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_address_carry_lint.py @@ -0,0 +1,132 @@ +# Copyright Advanced Micro Devices, Inc., or its affiliates. +# SPDX-License-Identifier: MIT + +"""The address-carry lint must report a 64-bit address update that drops the carry, and stay +quiet on correct carry chains and on registers that are only reused for other arithmetic.""" + +import os + +import pytest + +from Tensile.Utilities.address_carry_lint import lint + +pytestmark = pytest.mark.unit + +_DESIGNED = os.path.join( + os.path.dirname(__file__), "characterization", "_codegen", "data", "test_data", "_designed" +) + +# Generated kernels from designed configs that cover the address arithmetic most at risk: batch +# and stride address setup, Stream-K and GSU workspace addressing, fp8 and MX scale loads, and +# TDM descriptors. gfx90a, gfx942 and gfx950 have no native 64-bit add, so every address update +# is a 32-bit carry chain there. +_GENERATED = [ + ("gfx90a", "rich_gemm.yaml"), + ("gfx942", "asmaddr_initstrides.yaml"), + ("gfx942", "s00_loadbatchedaddress_non_stridedba.yaml"), + ("gfx942", "streamk.yaml"), + ("gfx942", "streamk_fixup_tree.yaml"), + ("gfx942", "gsu.yaml"), + ("gfx942", "fp8_gr_conv.yaml"), + ("gfx950", "mx_fp8_scale_swizzle.yaml"), + ("gfx950", "mx_bias_act_gsu.yaml"), + ("gfx950", "subtile.yaml"), + ("gfx1250", "streamk.yaml"), + ("gfx1250", "streamk_tdm_prefetchgl2.yaml"), + ("gfx1250", "mx_f4_nn_umlds0.yaml"), +] + + +@pytest.mark.parametrize("arch,config", _GENERATED, ids=lambda x: str(x).replace(".yaml", "")) +def test_generated_kernels_carry_every_address_update(arch, config): + from config_harness import emit_kernels_from_config + + results = emit_kernels_from_config( + os.path.join(_DESIGNED, arch, config), limit=4, arch=arch, canonical=False + ) + assert results, f"{arch}/{config} emitted no kernels" + for base, src, err in results: + assert err == 0, f"{base} failed to emit" + findings = lint(src) + assert not findings, f"{base}:\n" + "\n".join(str(f) for f in findings) + + +def _reasons(asm): + return [f.reason for f in lint(asm)] + + +def test_scalar_carry_chain_on_a_descriptor_base_is_clean(): + asm = """ + s_add_u32 s[sgprSrdA+0], s[sgprSrdA+0], s[sgprGlobalReadIncsA] + s_addc_u32 s[sgprSrdA+1], s[sgprSrdA+1], 0 + buffer_load_dwordx4 v[0:3], v4, s[sgprSrdA:sgprSrdA+3], 0 offen + """ + assert _reasons(asm) == [] + + +def test_scalar_carry_that_never_reaches_the_high_dword_is_reported(): + # The ROCM-32046 shape: the low dword is advanced, the carry is dropped. + asm = """ + s_add_u32 s[sgprSrdA+0], s[sgprSrdA+0], s[sgprGlobalReadIncsA] + s_mov_b32 s[sgprTmp], 0 + buffer_load_dwordx4 v[0:3], v4, s[sgprSrdA:sgprSrdA+3], 0 offen + """ + reasons = _reasons(asm) + assert len(reasons) == 1 + assert "s[sgprSrdA+0]" in reasons[0] and "s[sgprSrdA+1]" in reasons[0] + + +def test_add_without_carry_on_an_address_low_dword_is_reported(): + asm = """ + s_add_i32 s8, s8, 64 + s_load_dwordx2 s[10:11], s[8:9], 0x0 + """ + reasons = _reasons(asm) + assert len(reasons) == 1 and "s_add_i32" in reasons[0] and "s9" in reasons[0] + + +def test_vector_global_address_needs_a_carry_chain(): + good = """ + v_add_co_u32 v4, vcc, v4, v6 + v_addc_co_u32 v5, vcc, v5, 0, vcc + global_load_dwordx2 v[0:1], v[4:5], off + """ + bad = """ + v_add_u32 v4, v4, v6 + global_load_dwordx2 v[0:1], v[4:5], off + """ + assert _reasons(good) == [] + reasons = _reasons(bad) + assert len(reasons) == 1 and "v4" in reasons[0] + + +def test_a_register_reused_for_other_arithmetic_is_not_reported(): + # s20 is an address elsewhere, but this sum is overwritten before any address use. + asm = """ + s_load_dwordx2 s[30:31], s[20:21], 0x0 + s_sub_u32 s20, s2, s20 + s_add_u32 s20, s20, s3 + s_mov_b64 s[20:21], s[0:1] + s_load_dwordx2 s[32:33], s[20:21], 0x8 + """ + assert _reasons(asm) == [] + + +def test_data_registers_of_a_buffer_load_are_not_addresses(): + asm = """ + buffer_load_dwordx4 v[34:37], v5, s[40:43], 0 offen + v_sub_u32 v34, s12, v34 + buffer_load_dwordx4 v[34:37], v5, s[40:43], 0 offen + """ + assert _reasons(asm) == [] + + +def test_registers_are_judged_within_their_own_kernel(): + asm = """ + .amdgpu_hsa_kernel first + s_add_i32 s4, s4, 1 + s_endpgm + .amdgpu_hsa_kernel second + s_load_dwordx2 s[10:11], s[4:5], 0x0 + """ + assert _reasons(asm) == [] diff --git a/projects/hipblaslt/tensilelite/Tensile/Utilities/address_carry_lint.py b/projects/hipblaslt/tensilelite/Tensile/Utilities/address_carry_lint.py new file mode 100644 index 000000000000..c41a25c752ba --- /dev/null +++ b/projects/hipblaslt/tensilelite/Tensile/Utilities/address_carry_lint.py @@ -0,0 +1,298 @@ +# Copyright Advanced Micro Devices, Inc., or its affiliates. +# SPDX-License-Identifier: MIT + +"""Lint AMDGPU assembly for 64-bit address updates that drop the carry. + +A 64-bit address held in a register pair must be advanced with a carry chain: a low add that +produces a carry (s_add_u32, v_add_co_u32) followed by a high add that consumes it (s_addc_u32, +v_addc_co_u32). An add to the low dword with no carry into the high dword is correct only until +the low dword wraps, which happens when the buffer crosses a 4 GiB boundary: the access then +lands 4 GiB away (ROCM-32046, class C on AIHPBLAS-4988). + +The lint finds the registers used as the low dword of an address: the base of a buffer resource +descriptor (s[N:N+3] in a buffer instruction), the address pair of a scalar load (s[A:A+1]), and +the address pair or saddr of a global or flat access. For each write to such a register by an add +or subtract, it requires either a carry-producing instruction followed, within a short window, +by a carry-consuming write to the next register, or a full redefinition of the pair. Anything +else is reported, if the updated value is next used as an address before it is overwritten. The +scan is linear and ignores control flow, so a finding is a lead to read, not a proof. + +It is meant for Tensile-generated and hand-written kernels, which keep a 64-bit value in an +adjacent register pair. Compiler-generated code may keep the two halves of a sum in unrelated +registers, so the lint reports false findings there and is not meant for it. + + python -m Tensile.Utilities.address_carry_lint kernel.s + python -m Tensile.Utilities.address_carry_lint --disassemble TensileLibrary_gfx942.co +""" + +from __future__ import annotations + +import argparse +import re +import shutil +import subprocess +import sys +from dataclasses import dataclass +from pathlib import Path +from typing import Iterable, Optional + +# Low adds that set a carry, and the high adds that must consume it. +CARRY_OUT = {"s_add_u32", "s_sub_u32", "v_add_co_u32", "v_sub_co_u32", "v_subrev_co_u32"} +CARRY_IN = { + "s_addc_u32", + "s_subb_u32", + "v_addc_co_u32", + "v_subb_co_u32", + "v_subbrev_co_u32", + "v_add_co_ci_u32", + "v_sub_co_ci_u32", +} +# Adds that update one dword and drop any carry. +NO_CARRY = { + "s_add_i32", + "s_sub_i32", + "s_addk_i32", + "v_add_u32", + "v_sub_u32", + "v_add_nc_u32", + "v_sub_nc_u32", + "v_add_i32", +} +# Instructions that rewrite a whole pair, which ends any obligation on its low dword. +PAIR_DEFS = {"s_mov_b64", "v_mov_b64", "v_lshlrev_b64", "s_lshl_b64", "s_add_u64", "v_add_nc_u64"} + +WINDOW = 16 +# How far a written low dword is followed to its next use. +FLOW_WINDOW = 400 + +_REG = re.compile(r"^(?P[sv])(?:(?P\d+)|\[(?P[^\]:]+)(?::(?P[^\]]+))?\])") + + +@dataclass(frozen=True) +class Reg: + kind: str # "s" or "v" + base: str # number as text, or a symbol such as sgprSrdA + offset: int + + def plus(self, n: int) -> "Reg": + return Reg(self.kind, self.base, self.offset + n) + + def __str__(self) -> str: + if self.base.isdigit(): + return f"{self.kind}{int(self.base) + self.offset}" + return f"{self.kind}[{self.base}+{self.offset}]" + + +def _single(text: str, kind: str) -> Reg: + text = text.strip() + if text.isdigit(): + return Reg(kind, "0", int(text)) + m = re.match(r"^([A-Za-z_]\w*)(?:\s*\+\s*(\d+))?$", text) + if m: + return Reg(kind, m.group(1), int(m.group(2) or 0)) + return Reg(kind, text, 0) + + +def parse_regs(operand: str) -> list[Reg]: + """The registers an operand names, low to high; empty for a non-register operand.""" + m = _REG.match(operand.strip()) + if not m: + return [] + kind = m.group("kind") + if m.group("num") is not None: + return [Reg(kind, "0", int(m.group("num")))] + lo = _single(m.group("lo"), kind) + if m.group("hi") is None: + return [lo] + hi = _single(m.group("hi"), kind) + if hi.base != lo.base or hi.offset < lo.offset: + return [lo] + return [lo.plus(i) for i in range(hi.offset - lo.offset + 1)] + + +@dataclass +class Instruction: + line: int + text: str + mnemonic: str + operands: list[str] + + +def parse(asm: str) -> list[Instruction]: + out = [] + for n, raw in enumerate(asm.splitlines(), 1): + text = re.split(r"//|;", raw)[0].strip() + if not text or text.startswith(".") or text.endswith(":"): + continue + mnemonic, _, rest = text.partition(" ") + if not re.match(r"^[sv]_|^buffer_|^global_|^flat_|^scratch_", mnemonic): + continue + mnemonic = re.sub(r"_e(32|64)$", "", mnemonic) + operands = [o.strip() for o in rest.split(",")] if rest.strip() else [] + out.append(Instruction(n, text, mnemonic, operands)) + return out + + +def address_operands(inst: Instruction) -> list[Reg]: + """The registers this instruction uses as the low dword of an address or descriptor base.""" + regs = [parse_regs(o) for o in inst.operands] + m = inst.mnemonic + if m.startswith(("buffer_", "s_buffer_")): + # The descriptor is the SGPR quadruple; a VGPR quadruple is data. + return [r[0] for r in regs if len(r) == 4 and r[0].kind == "s"] + if m.startswith(("s_load_", "s_store_", "s_scratch_")): + return [regs[1][0]] if len(regs) > 1 and len(regs[1]) == 2 else [] + if m.startswith(("global_", "flat_")): + # Leave out the data and destination operands: a load writes operand 0, a store reads + # its data from operand 1, and a returning atomic does both. + if "_atomic" in m: + data = {0, 2} if (m.endswith("_rtn") or "glc" in inst.text) else {1} + elif "_store" in m: + data = {1} + else: + data = {0} + return [r[0] for k, r in enumerate(regs) if k not in data and len(r) == 2] + return [] + + +def written(inst: Instruction) -> list[Reg]: + """The registers this instruction writes: its first operand, unless it has no destination.""" + if not inst.operands or inst.mnemonic.startswith( + ("s_cmp", "s_bitcmp", "s_cbranch", "s_branch", "s_setpc", "s_waitcnt", "s_nop") + ): + return [] + if re.search(r"_store|_atomic", inst.mnemonic) and not ( + inst.mnemonic.endswith("_rtn") or "glc" in inst.text + ): + return [] + return parse_regs(inst.operands[0]) + + +def address_lows(insts: Iterable[Instruction]) -> set[Reg]: + """Registers used as the low dword of a 64-bit address or a descriptor base.""" + lows: set[Reg] = set() + for inst in insts: + lows.update(address_operands(inst)) + return lows + + +@dataclass +class Finding: + line: int + text: str + reason: str + + def __str__(self) -> str: + return f"line {self.line}: {self.text}\n {self.reason}" + + +# A kernel starts at Tensile's label_ASM_Start, at an assembler kernel directive, or, in +# disassembly, at a function symbol header such as :. +_KERNEL_START = re.compile( + r"label_ASM_Start>?:|^\s*\.amdgpu_hsa_kernel\b|^\s*\.globl\b|^<(?!label_)[^>\s]+>:\s*$", + re.M, +) + + +def kernels(asm: str) -> list[str]: + """Splits assembly into one chunk per kernel, so registers are judged within their kernel.""" + starts = [m.start() for m in _KERNEL_START.finditer(asm)] + if not starts: + return [asm] + # Keep line numbers: each chunk is padded with the newlines that precede it. + bounds = [0] + starts[1:] + [len(asm)] + return ["\n" * asm.count("\n", 0, a) + asm[a:b] for a, b in zip(bounds, bounds[1:])] + + +def lint(asm: str) -> list[Finding]: + findings = [] + for chunk in kernels(asm): + findings += _lint_kernel(chunk) + return findings + + +def _lint_kernel(asm: str) -> list[Finding]: + insts = parse(asm) + uses = [address_operands(inst) for inst in insts] + writes = [written(inst) for inst in insts] + lows = {r for u in uses for r in u} + + def flows_to_address(low: Reg, start: int) -> bool: + """Whether the value written at start is next used as an address, not overwritten.""" + for j in range(start + 1, min(len(insts), start + 1 + FLOW_WINDOW)): + if low in uses[j]: + return True + if low in writes[j]: + return False + return False + + findings = [] + for i, inst in enumerate(insts): + dst = writes[i] + if len(dst) != 1 or dst[0] not in lows: + continue + low, high = dst[0], dst[0].plus(1) + if inst.mnemonic in NO_CARRY: + if flows_to_address(low, i): + findings.append( + Finding( + inst.line, + inst.text, + f"{inst.mnemonic} updates {low}, the low dword of an address, with no " + f"carry into {high}", + ) + ) + elif inst.mnemonic in CARRY_OUT: + carried = any( + insts[j].mnemonic in CARRY_IN + and writes[j][:1] == [high] + or insts[j].mnemonic in PAIR_DEFS + and writes[j][:1] == [low] + for j in range(i + 1, min(len(insts), i + 1 + WINDOW)) + ) + if not carried and flows_to_address(low, i): + findings.append( + Finding( + inst.line, + inst.text, + f"{inst.mnemonic} sets a carry out of {low}, the low dword of an " + f"address, but nothing within {WINDOW} instructions adds it into {high}", + ) + ) + return findings + + +def disassemble(code_object: Path) -> str: + objdump = shutil.which("llvm-objdump") or "/opt/rocm/llvm/bin/llvm-objdump" + return subprocess.run( + [objdump, "-d", "--no-show-raw-insn", "--no-leading-addr", str(code_object)], + capture_output=True, + text=True, + check=True, + ).stdout + + +def main(argv: Optional[list[str]] = None) -> int: + parser = argparse.ArgumentParser( + description=__doc__, formatter_class=argparse.RawDescriptionHelpFormatter + ) + parser.add_argument("paths", nargs="+", type=Path, help="assembly files or code objects") + parser.add_argument( + "--disassemble", + action="store_true", + help="treat the paths as code objects and disassemble them with llvm-objdump", + ) + args = parser.parse_args(argv) + total = 0 + for path in args.paths: + asm = disassemble(path) if args.disassemble else path.read_text(errors="replace") + findings = lint(asm) + total += len(findings) + for f in findings: + print(f"{path}: {f}") + print(f"{total} finding(s) in {len(args.paths)} file(s)") + return 1 if total else 0 + + +if __name__ == "__main__": + sys.exit(main()) diff --git a/projects/hipblaslt/tensilelite/rocisa/rocisa/include/instruction/common.hpp b/projects/hipblaslt/tensilelite/rocisa/rocisa/include/instruction/common.hpp index e5dfbbd44f7e..3d309e042485 100644 --- a/projects/hipblaslt/tensilelite/rocisa/rocisa/include/instruction/common.hpp +++ b/projects/hipblaslt/tensilelite/rocisa/rocisa/include/instruction/common.hpp @@ -303,7 +303,7 @@ namespace rocisa = std::dynamic_pointer_cast(dst)->splitRegContainer(); std::vector srcs1; std::vector srcs2; - splitSrcs(srcs, srcs1, srcs2); + splitSrcs64(srcs, srcs1, srcs2); // s_add_u32 sets SCC, s_addc_u32 consumes it (carry is implicit). instructions = {std::make_shared(dst1, srcs1[0], srcs1[1], comment), @@ -3564,7 +3564,7 @@ namespace rocisa = std::dynamic_pointer_cast(dst)->splitRegContainer(); std::vector srcs1; std::vector srcs2; - splitSrcs(srcs, srcs1, srcs2); + splitSrcs64(srcs, srcs1, srcs2); auto vcc = std::make_shared(); instructions = {std::make_shared(dst1, vcc, srcs1[0], srcs1[1], comment), diff --git a/projects/hipblaslt/tensilelite/rocisa/rocisa/include/instruction/instruction.hpp b/projects/hipblaslt/tensilelite/rocisa/rocisa/include/instruction/instruction.hpp index f05c789cdbda..cc398e6e5d6b 100644 --- a/projects/hipblaslt/tensilelite/rocisa/rocisa/include/instruction/instruction.hpp +++ b/projects/hipblaslt/tensilelite/rocisa/rocisa/include/instruction/instruction.hpp @@ -27,9 +27,12 @@ #include "format.hpp" #include "helper.hpp" +#include +#include #include #include #include +#include #include #include #include @@ -116,6 +119,53 @@ namespace rocisa } } + // Splits the sources of a 64-bit integer add into its low and high 32-bit halves. A register + // pair splits into its two registers. An immediate splits into its low 32 bits and its high + // 32 bits, the sign extension for a 32-bit value; copying the immediate into both halves, as + // splitSrcs does for packed operations, would also add it times 2^32. An immediate above 2^31 + // reaches here as a double holding an integer. A symbol cannot be split and is rejected. + inline void splitSrcs64(const std::vector& srcs, + std::vector& srcs1, + std::vector& srcs2) + { + for(const auto& s : srcs) + { + std::visit( + [&srcs1, &srcs2](auto&& arg) -> void { + using T = std::decay_t; + if constexpr(std::is_same_v>) + { + auto regContainer = std::dynamic_pointer_cast(arg); + + auto [r1, r2] = regContainer->splitRegContainer(); + srcs1.push_back(r1); + srcs2.push_back(r2); + } + else if constexpr(std::is_same_v) + { + srcs1.push_back(arg); + srcs2.push_back(arg < 0 ? -1 : 0); + } + else if constexpr(std::is_same_v) + { + if(arg != std::trunc(arg) || std::fabs(arg) >= 0x1p63) + throw std::invalid_argument( + "64-bit add: cannot split non-integer immediate " + + std::to_string(arg) + " into 32-bit halves"); + const uint64_t v = static_cast(static_cast(arg)); + srcs1.push_back(static_cast(static_cast(v))); + srcs2.push_back(static_cast(static_cast(v >> 32))); + } + else + { + throw std::invalid_argument("64-bit add: cannot split immediate '" + arg + + "' into 32-bit halves"); + } + }, + s); + } + } + struct Instruction : public Item { InstType instType; diff --git a/projects/hipblaslt/tensilelite/rocisa/test/test_lowering64.py b/projects/hipblaslt/tensilelite/rocisa/test/test_lowering64.py new file mode 100644 index 000000000000..b2aa6edce328 --- /dev/null +++ b/projects/hipblaslt/tensilelite/rocisa/test/test_lowering64.py @@ -0,0 +1,114 @@ +# Copyright Advanced Micro Devices, Inc., or its affiliates. +# SPDX-License-Identifier: MIT + +"""Lowering of 64-bit integer adds on targets without a native 64-bit add. + +SAddU64 and VAddNCU64 lower to a pair of 32-bit adds on targets that lack s_add_u64 and +v_add_nc_u64 (gfx942, gfx950, ...): the low add sets the carry and the high add consumes it. +A register pair splits into its two halves. An immediate must split into its low 32 bits and +the sign extension of those into the high half, never the immediate again: with it, ptr += imm +would also add imm * 2^32 and move the pointer hundreds of GiB (AMD-ROCm-Internal PR 2557, +class C2 on AIHPBLAS-4988). +""" + +import os +import shutil + +import pytest + +import rocisa +from rocisa.container import sgpr, vgpr +from rocisa.instruction import SAddU64, VAddNCU64 + +# Targets without a native 64-bit add, so both instructions take the lowered path. +_ISAS = [(9, 0, 10), (9, 4, 2), (9, 5, 0)] + + +def _init(isa): + rocm_path = os.environ.get("ROCM_PATH", "/opt/rocm") + search_path = os.pathsep.join( + [os.path.join(rocm_path, "bin"), os.path.join(rocm_path, "lib", "llvm", "bin")] + ) + assembler = shutil.which("amdclang++", path=search_path) or "amdclang++" + ri = rocisa.rocIsa.getInstance() + ri.init(isa, assembler, False) + ri.setKernel(isa, 64) + return ri + + +@pytest.fixture(params=_ISAS, ids=lambda isa: rocisa.isaToGfx(isa)) +def lowered_isa(request): + ri = _init(request.param) + caps = ri.getAsmCaps() + if caps["s_add_u64"] or caps["v_add_nc_u64"]: + pytest.skip(f"{rocisa.isaToGfx(request.param)} has a native 64-bit add") + yield request.param + _init((9, 4, 2)) + + +def _lines(inst): + """The emitted instructions, without comments, as (mnemonic, operands) pairs.""" + out = [] + for line in str(inst).splitlines(): + line = line.split("//")[0].strip() + if line: + mnemonic, _, rest = line.partition(" ") + out.append((mnemonic, [o.strip() for o in rest.split(",")])) + return out + + +def _as_int(text): + return int(text, 0) + + +# (immediate, expected low 32 bits as a signed value, expected high 32 bits) +_IMMEDIATES = [ + (5, 5, 0), + (0, 0, 0), + (-8, -8, -1), + (0x7FFFFFFF, 0x7FFFFFFF, 0), + (-(2**31), -(2**31), -1), + # Above 2^31: Python passes these through as doubles, which must split as 64-bit integers. + (2**31 + 5, -(2**31) + 5, 0), + (2**32 + 3, 3, 1), + (3 * 2**32 + 7, 7, 3), +] + + +def test_saddu64_registers_lower_to_a_carry_chain(lowered_isa): + lines = _lines(SAddU64(dst=sgpr(4, 2), src0=sgpr(6, 2), src1=sgpr(8, 2))) + assert lines == [("s_add_u32", ["s4", "s6", "s8"]), ("s_addc_u32", ["s5", "s7", "s9"])] + + +@pytest.mark.parametrize("imm,low,high", _IMMEDIATES) +def test_saddu64_immediate_high_half_is_the_sign_extension(lowered_isa, imm, low, high): + lines = _lines(SAddU64(dst=sgpr(4, 2), src0=sgpr(4, 2), src1=imm)) + assert [m for m, _ in lines] == ["s_add_u32", "s_addc_u32"] + (_, lo_ops), (_, hi_ops) = lines + assert lo_ops[:2] == ["s4", "s4"] and hi_ops[:2] == ["s5", "s5"] + assert _as_int(lo_ops[2]) == low + assert _as_int(hi_ops[2]) == high, f"high half of {imm:#x} must be {high}, got {hi_ops[2]}" + + +def test_vaddncu64_registers_lower_to_a_carry_chain(lowered_isa): + lines = _lines(VAddNCU64(dst=vgpr(4, 2), src0=vgpr(6, 2), src1=vgpr(8, 2))) + assert [m for m, _ in lines] == ["v_add_co_u32", "v_addc_co_u32"] + (_, lo_ops), (_, hi_ops) = lines + assert lo_ops[0] == "v4" and lo_ops[2:] == ["v6", "v8"] + assert hi_ops[0] == "v5" and hi_ops[2:4] == ["v7", "v9"] + # The low add writes the carry and the high add reads the same carry register. + assert lo_ops[1] == hi_ops[1] == hi_ops[4] + + +@pytest.mark.parametrize("imm,low,high", _IMMEDIATES) +def test_vaddncu64_immediate_high_half_is_the_sign_extension(lowered_isa, imm, low, high): + lines = _lines(VAddNCU64(dst=vgpr(4, 2), src0=vgpr(4, 2), src1=imm)) + assert [m for m, _ in lines] == ["v_add_co_u32", "v_addc_co_u32"] + (_, lo_ops), (_, hi_ops) = lines + assert _as_int(lo_ops[3]) == low + assert _as_int(hi_ops[3]) == high, f"high half of {imm:#x} must be {high}, got {hi_ops[3]}" + + +def test_symbolic_immediate_is_rejected(lowered_isa): + with pytest.raises(Exception, match="cannot split"): + str(SAddU64(dst=sgpr(4, 2), src0=sgpr(4, 2), src1="SomeSymbol")) From 17d084d57db60f051613614c6cd6a3b680e0c2db Mon Sep 17 00:00:00 2001 From: Tony Davis Date: Fri, 2 Oct 2026 23:56:16 -0400 Subject: [PATCH 02/11] fix(tensilelite): AIHPBLAS-5022 stop the address-carry lint at jumps and constant adds On gfx950 the lint reported every kernel with a long branch: s_getpc_b64, s_add_i32 sN, , 4, a carried add into the PC and s_setpc_b64. The linear scan walked past the jump and met an unrelated address use of the same register in the next block. Following a written value now stops at an unconditional jump (s_branch, s_setpc_b64, s_endpgm), and an add of two constants is not an address update. Two synthetic tests cover both, and an address update followed by a jump is still reported. Co-authored-by: Cursor --- .../Tests/unit/test_address_carry_lint.py | 25 +++++++++++++++++++ .../Tensile/Utilities/address_carry_lint.py | 11 +++++--- 2 files changed, 33 insertions(+), 3 deletions(-) diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_address_carry_lint.py b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_address_carry_lint.py index eefe6f7c1de7..8f655b2f6284 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_address_carry_lint.py +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_address_carry_lint.py @@ -121,6 +121,31 @@ def test_data_registers_of_a_buffer_load_are_not_addresses(): assert _reasons(asm) == [] +def test_a_long_branch_offset_is_not_an_address(): + # The gfx950 long-branch sequence: s64 holds a constant branch offset, and the next use of + # s[64:65] as an address is in another block, past the jump. + asm = """ + s_getpc_b64 s[62:63] + s_add_i32 s64, 0x424f4, 4 + s_add_u32 s62, s62, s64 + s_addc_u32 s63, s63, 0 + s_setpc_b64 s[62:63] + s_cmp_eq_u64 s[64:65], 0 + s_load_dword s8, s[64:65], 0x0 + """ + assert _reasons(asm) == [] + + +def test_an_address_update_before_a_jump_is_still_reported(): + asm = """ + s_add_i32 s8, s8, 64 + s_load_dwordx2 s[10:11], s[8:9], 0x0 + s_branch label_next + """ + reasons = _reasons(asm) + assert len(reasons) == 1 and "s8" in reasons[0] + + def test_registers_are_judged_within_their_own_kernel(): asm = """ .amdgpu_hsa_kernel first diff --git a/projects/hipblaslt/tensilelite/Tensile/Utilities/address_carry_lint.py b/projects/hipblaslt/tensilelite/Tensile/Utilities/address_carry_lint.py index c41a25c752ba..b5dc16372869 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Utilities/address_carry_lint.py +++ b/projects/hipblaslt/tensilelite/Tensile/Utilities/address_carry_lint.py @@ -14,8 +14,9 @@ the address pair or saddr of a global or flat access. For each write to such a register by an add or subtract, it requires either a carry-producing instruction followed, within a short window, by a carry-consuming write to the next register, or a full redefinition of the pair. Anything -else is reported, if the updated value is next used as an address before it is overwritten. The -scan is linear and ignores control flow, so a finding is a lead to read, not a proof. +else is reported, if the updated value is next used as an address before it is overwritten or an +unconditional jump. An add of two constants sets the register rather than advancing an address, +so it is not reported. The scan is otherwise linear, so a finding is a lead to read, not a proof. It is meant for Tensile-generated and hand-written kernels, which keep a 64-bit value in an adjacent register pair. Compiler-generated code may keep the two halves of a sum in unrelated @@ -60,6 +61,8 @@ } # Instructions that rewrite a whole pair, which ends any obligation on its low dword. PAIR_DEFS = {"s_mov_b64", "v_mov_b64", "v_lshlrev_b64", "s_lshl_b64", "s_add_u64", "v_add_nc_u64"} +# Unconditional transfers: the next instruction in the listing is not the next one executed. +JUMPS = {"s_branch", "s_setpc_b64", "s_endpgm"} WINDOW = 16 # How far a written low dword is followed to its next use. @@ -222,7 +225,7 @@ def flows_to_address(low: Reg, start: int) -> bool: for j in range(start + 1, min(len(insts), start + 1 + FLOW_WINDOW)): if low in uses[j]: return True - if low in writes[j]: + if low in writes[j] or insts[j].mnemonic in JUMPS: return False return False @@ -233,6 +236,8 @@ def flows_to_address(low: Reg, start: int) -> bool: continue low, high = dst[0], dst[0].plus(1) if inst.mnemonic in NO_CARRY: + if not any(parse_regs(o) for o in inst.operands[1:]): + continue if flows_to_address(low, i): findings.append( Finding( From d1739751c212fab6da26d63b6eabc65c5891effa Mon Sep 17 00:00:00 2001 From: Tony Davis Date: Sat, 3 Oct 2026 10:48:37 -0400 Subject: [PATCH 03/11] fix(tensilelite): AIHPBLAS-5022 drop the gfx1250 MX config from the lint's generated kernels Develop now rejects every solution in gfx1250/mx_f4_nn_umlds0.yaml ("M-major MX-scale local read for A requires VectorWidthA >= MatrixInstK // MXBlockA"), so the generated-kernel lint case emitted no kernels and failed the TensileLite test, coverage and Jenkins codecov checks. The gfx950 MX configs still cover MX scale addressing. Co-authored-by: Cursor --- .../tensilelite/Tensile/Tests/unit/test_address_carry_lint.py | 1 - 1 file changed, 1 deletion(-) diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_address_carry_lint.py b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_address_carry_lint.py index 8f655b2f6284..5fac57c61219 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_address_carry_lint.py +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_address_carry_lint.py @@ -33,7 +33,6 @@ ("gfx950", "subtile.yaml"), ("gfx1250", "streamk.yaml"), ("gfx1250", "streamk_tdm_prefetchgl2.yaml"), - ("gfx1250", "mx_f4_nn_umlds0.yaml"), ] From a2198886e13c1ea034ab1ebae72bf5f8d473a3e2 Mon Sep 17 00:00:00 2001 From: Tony Davis Date: Sat, 3 Oct 2026 10:59:01 -0400 Subject: [PATCH 04/11] fix(tensilelite): AIHPBLAS-5022 follow scalar carries until SCC changes, and split kernels in linear time The gfx950 audit found 105 false findings in fp8 kernels where the scheduler placed s_addc_u32 about 50 instructions after its s_add_u32, past the lint's 16-instruction window, with nothing in between writing SCC. A scalar carry is now followed until the matching carry-in, a jump, or the first instruction that writes SCC. Vector carries keep the short window. The audit also timed out on 33 large code objects: kernels() kept line numbers by padding each kernel's text with every newline before it, which is quadratic in the file size. Kernels now carry their first line number instead. New synthetic tests cover a far carry-in, a carry lost to an SCC write, and line numbers in a later kernel. Co-authored-by: Cursor --- .../Tests/unit/test_address_carry_lint.py | 39 +++++++++ .../Tensile/Utilities/address_carry_lint.py | 81 +++++++++++++------ 2 files changed, 96 insertions(+), 24 deletions(-) diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_address_carry_lint.py b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_address_carry_lint.py index 5fac57c61219..4800e846a598 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_address_carry_lint.py +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_address_carry_lint.py @@ -145,6 +145,45 @@ def test_an_address_update_before_a_jump_is_still_reported(): assert len(reasons) == 1 and "s8" in reasons[0] +def test_a_scalar_carry_moved_far_by_the_scheduler_is_clean(): + # The gfx950 fp8 kernels: the carry-in comes 40 vector instructions after the carry-out, and + # nothing in between writes SCC. + filler = "\n".join(" v_perm_b32 v27, v55, v54, s78" for _ in range(40)) + asm = f""" + s_add_u32 s68, s68, s76 +{filler} + s_addc_u32 s69, s69, 0 + buffer_load_dwordx4 v8, s[68:71], 0 offen lds + """ + assert _reasons(asm) == [] + + +def test_a_scalar_carry_lost_to_an_scc_write_is_reported(): + asm = """ + s_add_u32 s68, s68, s76 + s_cmp_eq_u32 s73, 0 + s_addc_u32 s69, s69, 0 + buffer_load_dwordx4 v8, s[68:71], 0 offen lds + """ + reasons = _reasons(asm) + assert len(reasons) == 1 and "before SCC changes" in reasons[0] + + +def test_findings_keep_their_line_numbers_across_kernels(): + asm = "\n".join( + [ + ".amdgpu_hsa_kernel first", + " s_load_dwordx2 s[10:11], s[4:5], 0x0", + ".amdgpu_hsa_kernel second", + " s_nop 0", + " s_add_i32 s8, s8, 64", + " s_load_dwordx2 s[10:11], s[8:9], 0x0", + ] + ) + findings = lint(asm) + assert [f.line for f in findings] == [5] + + def test_registers_are_judged_within_their_own_kernel(): asm = """ .amdgpu_hsa_kernel first diff --git a/projects/hipblaslt/tensilelite/Tensile/Utilities/address_carry_lint.py b/projects/hipblaslt/tensilelite/Tensile/Utilities/address_carry_lint.py index b5dc16372869..46e52df99e21 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Utilities/address_carry_lint.py +++ b/projects/hipblaslt/tensilelite/Tensile/Utilities/address_carry_lint.py @@ -12,11 +12,13 @@ The lint finds the registers used as the low dword of an address: the base of a buffer resource descriptor (s[N:N+3] in a buffer instruction), the address pair of a scalar load (s[A:A+1]), and the address pair or saddr of a global or flat access. For each write to such a register by an add -or subtract, it requires either a carry-producing instruction followed, within a short window, -by a carry-consuming write to the next register, or a full redefinition of the pair. Anything -else is reported, if the updated value is next used as an address before it is overwritten or an -unconditional jump. An add of two constants sets the register rather than advancing an address, -so it is not reported. The scan is otherwise linear, so a finding is a lead to read, not a proof. +or subtract, it requires either a carry-producing instruction followed by a carry-consuming write +to the next register, or a full redefinition of the pair. A scalar carry is followed until SCC +changes, since the scheduler can move the carry-in far from the carry-out; a vector carry, for a +short window. Anything else is reported, if the updated value is next used as an address before it +is overwritten or an unconditional jump. An add of two constants sets the register rather than +advancing an address, so it is not reported. The scan is otherwise linear, so a finding is a lead +to read, not a proof. It is meant for Tensile-generated and hand-written kernels, which keep a 64-bit value in an adjacent register pair. Compiler-generated code may keep the two halves of a sum in unrelated @@ -63,6 +65,11 @@ PAIR_DEFS = {"s_mov_b64", "v_mov_b64", "v_lshlrev_b64", "s_lshl_b64", "s_add_u64", "v_add_nc_u64"} # Unconditional transfers: the next instruction in the listing is not the next one executed. JUMPS = {"s_branch", "s_setpc_b64", "s_endpgm"} +# Scalar instructions that write SCC, which ends a scalar carry chain. +_SCC_WRITERS = re.compile( + r"^s_(add|sub|addc|subb|addk|cmp|cmpk|bitcmp|and|or|xor|andn[12]|orn[12]|nand|nor|xnor|not|" + r"lshl|lshr|ashr|bfe|abs|absdiff|min|max|bcnt|quadmask|wqm)" +) WINDOW = 16 # How far a written low dword is followed to its next use. @@ -121,9 +128,9 @@ class Instruction: operands: list[str] -def parse(asm: str) -> list[Instruction]: +def parse(asm: str, first_line: int = 1) -> list[Instruction]: out = [] - for n, raw in enumerate(asm.splitlines(), 1): + for n, raw in enumerate(asm.splitlines(), first_line): text = re.split(r"//|;", raw)[0].strip() if not text or text.startswith(".") or text.endswith(":"): continue @@ -197,25 +204,34 @@ def __str__(self) -> str: ) -def kernels(asm: str) -> list[str]: - """Splits assembly into one chunk per kernel, so registers are judged within their kernel.""" +def _kernel_chunks(asm: str) -> list[tuple[int, str]]: + """One (first line number, text) chunk per kernel.""" starts = [m.start() for m in _KERNEL_START.finditer(asm)] if not starts: - return [asm] - # Keep line numbers: each chunk is padded with the newlines that precede it. + return [(1, asm)] bounds = [0] + starts[1:] + [len(asm)] - return ["\n" * asm.count("\n", 0, a) + asm[a:b] for a, b in zip(bounds, bounds[1:])] + chunks, line, pos = [], 1, 0 + for a, b in zip(bounds, bounds[1:]): + line += asm.count("\n", pos, a) + pos = a + chunks.append((line, asm[a:b])) + return chunks + + +def kernels(asm: str) -> list[str]: + """Splits assembly into one chunk per kernel, so registers are judged within their kernel.""" + return [text for _, text in _kernel_chunks(asm)] def lint(asm: str) -> list[Finding]: findings = [] - for chunk in kernels(asm): - findings += _lint_kernel(chunk) + for first_line, chunk in _kernel_chunks(asm): + findings += _lint_kernel(chunk, first_line) return findings -def _lint_kernel(asm: str) -> list[Finding]: - insts = parse(asm) +def _lint_kernel(asm: str, first_line: int = 1) -> list[Finding]: + insts = parse(asm, first_line) uses = [address_operands(inst) for inst in insts] writes = [written(inst) for inst in insts] lows = {r for u in uses for r in u} @@ -229,6 +245,23 @@ def flows_to_address(low: Reg, start: int) -> bool: return False return False + def consumes(j: int, low: Reg, high: Reg) -> bool: + """Whether instruction j adds a carry into high, or redefines the whole pair.""" + return (insts[j].mnemonic in CARRY_IN and writes[j][:1] == [high]) or ( + insts[j].mnemonic in PAIR_DEFS and writes[j][:1] == [low] + ) + + def scalar_carried(start: int, low: Reg, high: Reg) -> bool: + """Whether a scalar carry reaches high. The scheduler may move the carry-in well away + from the carry-out, so it is followed until SCC changes rather than for a fixed count. + """ + for j in range(start + 1, min(len(insts), start + 1 + FLOW_WINDOW)): + if consumes(j, low, high): + return True + if insts[j].mnemonic in JUMPS or _SCC_WRITERS.match(insts[j].mnemonic): + return False + return False + findings = [] for i, inst in enumerate(insts): dst = writes[i] @@ -248,20 +281,20 @@ def flows_to_address(low: Reg, start: int) -> bool: ) ) elif inst.mnemonic in CARRY_OUT: - carried = any( - insts[j].mnemonic in CARRY_IN - and writes[j][:1] == [high] - or insts[j].mnemonic in PAIR_DEFS - and writes[j][:1] == [low] - for j in range(i + 1, min(len(insts), i + 1 + WINDOW)) - ) + if inst.mnemonic.startswith("s_"): + carried, where = scalar_carried(i, low, high), "before SCC changes" + else: + carried = any( + consumes(j, low, high) for j in range(i + 1, min(len(insts), i + 1 + WINDOW)) + ) + where = f"within {WINDOW} instructions" if not carried and flows_to_address(low, i): findings.append( Finding( inst.line, inst.text, f"{inst.mnemonic} sets a carry out of {low}, the low dword of an " - f"address, but nothing within {WINDOW} instructions adds it into {high}", + f"address, but nothing {where} adds it into {high}", ) ) return findings From 83a84d1404becd91e62ee5fa6e613c3324de449a Mon Sep 17 00:00:00 2001 From: Tony Davis Date: Sat, 3 Oct 2026 15:48:00 -0400 Subject: [PATCH 05/11] fix(rocisa): AIHPBLAS-5022 accept -2^63 as a 64-bit add immediate splitSrcs64 rejected any immediate whose magnitude reached 2^63, which also rejected -2^63, a valid signed 64-bit value. The bound is now [-2^63, 2^63). The lowering test splits -2^63 and checks that 2^63 is still rejected. Co-authored-by: Cursor --- .../rocisa/rocisa/include/instruction/instruction.hpp | 7 ++++--- .../hipblaslt/tensilelite/rocisa/test/test_lowering64.py | 6 ++++++ 2 files changed, 10 insertions(+), 3 deletions(-) diff --git a/projects/hipblaslt/tensilelite/rocisa/rocisa/include/instruction/instruction.hpp b/projects/hipblaslt/tensilelite/rocisa/rocisa/include/instruction/instruction.hpp index cc398e6e5d6b..cc504d8c507e 100644 --- a/projects/hipblaslt/tensilelite/rocisa/rocisa/include/instruction/instruction.hpp +++ b/projects/hipblaslt/tensilelite/rocisa/rocisa/include/instruction/instruction.hpp @@ -148,10 +148,11 @@ namespace rocisa } else if constexpr(std::is_same_v) { - if(arg != std::trunc(arg) || std::fabs(arg) >= 0x1p63) + // The signed 64-bit range is [-2^63, 2^63). + if(arg != std::trunc(arg) || arg < -0x1p63 || arg >= 0x1p63) throw std::invalid_argument( - "64-bit add: cannot split non-integer immediate " - + std::to_string(arg) + " into 32-bit halves"); + "64-bit add: cannot split immediate " + std::to_string(arg) + + ", which is not a 64-bit integer, into 32-bit halves"); const uint64_t v = static_cast(static_cast(arg)); srcs1.push_back(static_cast(static_cast(v))); srcs2.push_back(static_cast(static_cast(v >> 32))); diff --git a/projects/hipblaslt/tensilelite/rocisa/test/test_lowering64.py b/projects/hipblaslt/tensilelite/rocisa/test/test_lowering64.py index b2aa6edce328..c6ff8ad3a35f 100644 --- a/projects/hipblaslt/tensilelite/rocisa/test/test_lowering64.py +++ b/projects/hipblaslt/tensilelite/rocisa/test/test_lowering64.py @@ -72,6 +72,7 @@ def _as_int(text): (2**31 + 5, -(2**31) + 5, 0), (2**32 + 3, 3, 1), (3 * 2**32 + 7, 7, 3), + (-(2**63), 0, -(2**31)), ] @@ -112,3 +113,8 @@ def test_vaddncu64_immediate_high_half_is_the_sign_extension(lowered_isa, imm, l def test_symbolic_immediate_is_rejected(lowered_isa): with pytest.raises(Exception, match="cannot split"): str(SAddU64(dst=sgpr(4, 2), src0=sgpr(4, 2), src1="SomeSymbol")) + + +def test_immediate_outside_the_64_bit_range_is_rejected(lowered_isa): + with pytest.raises(Exception, match="cannot split"): + str(SAddU64(dst=sgpr(4, 2), src0=sgpr(4, 2), src1=2**63)) From 2daea089044a81c476f79a6db95d114c0ff2f971 Mon Sep 17 00:00:00 2001 From: Tony Davis Date: Sat, 3 Oct 2026 15:48:00 -0400 Subject: [PATCH 06/11] fix(tensilelite): AIHPBLAS-5022 parse tab-separated lines, track the vector carry, skip unsupported targets - The parser split the mnemonic from its operands on a space only, so tab-separated hand-written kernels were skipped entirely and reported clean. Tabs are now treated as spaces. - A vector carry chain was accepted if any carry-in to the high half followed within the window, even after another instruction overwrote the carry register. The carry-in must now read the register the low add wrote, before anything else writes it. - The generated-kernel cases now skip a target the toolchain does not support (the SupportedISA capability the other gfx1250 codegen tests use) instead of failing because no kernels were emitted. Co-authored-by: Cursor --- .../Tests/unit/test_address_carry_lint.py | 38 +++++++++++++++++++ .../Tensile/Utilities/address_carry_lint.py | 36 +++++++++++++----- 2 files changed, 65 insertions(+), 9 deletions(-) diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_address_carry_lint.py b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_address_carry_lint.py index 4800e846a598..672b1ba4227e 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_address_carry_lint.py +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_address_carry_lint.py @@ -4,6 +4,7 @@ """The address-carry lint must report a 64-bit address update that drops the carry, and stay quiet on correct carry chains and on registers that are only reused for other arithmetic.""" +import functools import os import pytest @@ -36,10 +37,23 @@ ] +@functools.lru_cache(maxsize=None) +def _assembler_supports(arch): + from Tensile.Common.Architectures import gfxToIsa + from Tensile.Common.Capabilities import makeIsaInfoMap + from Tensile.Toolchain.Validators import validateToolchain + + isa = gfxToIsa(arch) + return bool(makeIsaInfoMap([isa], validateToolchain("amdclang++"))[isa].asmCaps["SupportedISA"]) + + @pytest.mark.parametrize("arch,config", _GENERATED, ids=lambda x: str(x).replace(".yaml", "")) def test_generated_kernels_carry_every_address_update(arch, config): from config_harness import emit_kernels_from_config + if not _assembler_supports(arch): + pytest.skip(f"amdclang++ in this environment does not support {arch}") + results = emit_kernels_from_config( os.path.join(_DESIGNED, arch, config), limit=4, arch=arch, canonical=False ) @@ -184,6 +198,30 @@ def test_findings_keep_their_line_numbers_across_kernels(): assert [f.line for f in findings] == [5] +def test_tab_separated_instructions_are_parsed(): + asm = "\ts_add_i32\ts8, s8, 64\n\ts_load_dwordx2\ts[10:11], s[8:9], 0x0\n" + reasons = _reasons(asm) + assert len(reasons) == 1 and "s8" in reasons[0] + + +def test_a_vector_carry_overwritten_before_the_carry_in_is_reported(): + good = """ + v_add_co_u32 v4, vcc, v4, v6 + v_add_co_u32 v10, s[20:21], v10, v12 + v_addc_co_u32 v5, vcc, v5, 0, vcc + global_load_dwordx2 v[0:1], v[4:5], off + """ + bad = """ + v_add_co_u32 v4, vcc, v4, v6 + v_add_co_u32 v10, vcc, v10, v12 + v_addc_co_u32 v5, vcc, v5, 0, vcc + global_load_dwordx2 v[0:1], v[4:5], off + """ + assert _reasons(good) == [] + reasons = _reasons(bad) + assert len(reasons) == 1 and "v4" in reasons[0] + + def test_registers_are_judged_within_their_own_kernel(): asm = """ .amdgpu_hsa_kernel first diff --git a/projects/hipblaslt/tensilelite/Tensile/Utilities/address_carry_lint.py b/projects/hipblaslt/tensilelite/Tensile/Utilities/address_carry_lint.py index 46e52df99e21..3d730b5a908e 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Utilities/address_carry_lint.py +++ b/projects/hipblaslt/tensilelite/Tensile/Utilities/address_carry_lint.py @@ -15,10 +15,10 @@ or subtract, it requires either a carry-producing instruction followed by a carry-consuming write to the next register, or a full redefinition of the pair. A scalar carry is followed until SCC changes, since the scheduler can move the carry-in far from the carry-out; a vector carry, for a -short window. Anything else is reported, if the updated value is next used as an address before it -is overwritten or an unconditional jump. An add of two constants sets the register rather than -advancing an address, so it is not reported. The scan is otherwise linear, so a finding is a lead -to read, not a proof. +short window and only until its carry register is written again. Anything else is reported, if the +updated value is next used as an address before it is overwritten or an unconditional jump. An add +of two constants sets the register rather than advancing an address, so it is not reported. The +scan is otherwise linear, so a finding is a lead to read, not a proof. It is meant for Tensile-generated and hand-written kernels, which keep a 64-bit value in an adjacent register pair. Compiler-generated code may keep the two halves of a sum in unrelated @@ -134,7 +134,8 @@ def parse(asm: str, first_line: int = 1) -> list[Instruction]: text = re.split(r"//|;", raw)[0].strip() if not text or text.startswith(".") or text.endswith(":"): continue - mnemonic, _, rest = text.partition(" ") + # Hand-written kernels separate the mnemonic from its operands with a tab. + mnemonic, _, rest = text.replace("\t", " ").partition(" ") if not re.match(r"^[sv]_|^buffer_|^global_|^flat_|^scratch_", mnemonic): continue mnemonic = re.sub(r"_e(32|64)$", "", mnemonic) @@ -251,6 +252,25 @@ def consumes(j: int, low: Reg, high: Reg) -> bool: insts[j].mnemonic in PAIR_DEFS and writes[j][:1] == [low] ) + def vector_carried(start: int, low: Reg, high: Reg) -> bool: + """Whether a vector carry reaches high: a carry-in to high that reads the carry register + the low add wrote, before anything else writes that register.""" + carry = insts[start].operands[1] if len(insts[start].operands) > 1 else "" + for j in range(start + 1, min(len(insts), start + 1 + WINDOW)): + ops = insts[j].operands + if insts[j].mnemonic in PAIR_DEFS and writes[j][:1] == [low]: + return True + if insts[j].mnemonic in CARRY_IN and writes[j][:1] == [high]: + return bool(ops) and ops[-1] == carry + # A carry register is written as the destination of a compare or move, or as the + # carry-out of another add. + if ops and ( + ops[0] == carry + or (insts[j].mnemonic in CARRY_OUT | CARRY_IN and len(ops) > 1 and ops[1] == carry) + ): + return False + return False + def scalar_carried(start: int, low: Reg, high: Reg) -> bool: """Whether a scalar carry reaches high. The scheduler may move the carry-in well away from the carry-out, so it is followed until SCC changes rather than for a fixed count. @@ -284,10 +304,8 @@ def scalar_carried(start: int, low: Reg, high: Reg) -> bool: if inst.mnemonic.startswith("s_"): carried, where = scalar_carried(i, low, high), "before SCC changes" else: - carried = any( - consumes(j, low, high) for j in range(i + 1, min(len(insts), i + 1 + WINDOW)) - ) - where = f"within {WINDOW} instructions" + carried = vector_carried(i, low, high) + where = f"within {WINDOW} instructions, from the carry it wrote," if not carried and flows_to_address(low, i): findings.append( Finding( From ba84c8cab3135e47fb3452c375b9ce32b3f4b804 Mon Sep 17 00:00:00 2001 From: Tony Davis Date: Sat, 3 Oct 2026 18:22:17 -0400 Subject: [PATCH 07/11] fix(rocisa): AIHPBLAS-5022 split only immediates a double holds exactly A Python integer above 2^31 reaches the 64-bit add lowering as a double, which rounds integers at or beyond 2^53 in magnitude: 2^63 - 1 arrives as 2^63 and 2^53 + 1 as 2^53. Splitting such a value could add a different offset than the one asked for, so the lowering now rejects any immediate of magnitude 2^53 or more. The lowering test splits +/-(2^53 - 1) and rejects 2^53, 2^53 + 1, -2^53, 2^63 - 1, -2^63 and 2^63. Co-authored-by: Cursor --- .../rocisa/rocisa/include/instruction/instruction.hpp | 8 +++++--- .../tensilelite/rocisa/test/test_lowering64.py | 10 +++++++--- 2 files changed, 12 insertions(+), 6 deletions(-) diff --git a/projects/hipblaslt/tensilelite/rocisa/rocisa/include/instruction/instruction.hpp b/projects/hipblaslt/tensilelite/rocisa/rocisa/include/instruction/instruction.hpp index cc504d8c507e..676bc78958fb 100644 --- a/projects/hipblaslt/tensilelite/rocisa/rocisa/include/instruction/instruction.hpp +++ b/projects/hipblaslt/tensilelite/rocisa/rocisa/include/instruction/instruction.hpp @@ -124,6 +124,8 @@ namespace rocisa // 32 bits, the sign extension for a 32-bit value; copying the immediate into both halves, as // splitSrcs does for packed operations, would also add it times 2^32. An immediate above 2^31 // reaches here as a double holding an integer. A symbol cannot be split and is rejected. + // A double is exact only below 2^53 in magnitude; at or beyond that, a Python integer may + // have been rounded on its way in, so it is rejected rather than split as another value. inline void splitSrcs64(const std::vector& srcs, std::vector& srcs1, std::vector& srcs2) @@ -148,11 +150,11 @@ namespace rocisa } else if constexpr(std::is_same_v) { - // The signed 64-bit range is [-2^63, 2^63). - if(arg != std::trunc(arg) || arg < -0x1p63 || arg >= 0x1p63) + if(arg != std::trunc(arg) || std::fabs(arg) >= 0x1p53) throw std::invalid_argument( "64-bit add: cannot split immediate " + std::to_string(arg) - + ", which is not a 64-bit integer, into 32-bit halves"); + + ", which is not an integer below 2^53 in magnitude, into" + " 32-bit halves"); const uint64_t v = static_cast(static_cast(arg)); srcs1.push_back(static_cast(static_cast(v))); srcs2.push_back(static_cast(static_cast(v >> 32))); diff --git a/projects/hipblaslt/tensilelite/rocisa/test/test_lowering64.py b/projects/hipblaslt/tensilelite/rocisa/test/test_lowering64.py index c6ff8ad3a35f..45d72e616544 100644 --- a/projects/hipblaslt/tensilelite/rocisa/test/test_lowering64.py +++ b/projects/hipblaslt/tensilelite/rocisa/test/test_lowering64.py @@ -72,7 +72,9 @@ def _as_int(text): (2**31 + 5, -(2**31) + 5, 0), (2**32 + 3, 3, 1), (3 * 2**32 + 7, 7, 3), - (-(2**63), 0, -(2**31)), + # The largest magnitudes a double carries exactly. + (2**53 - 1, -1, 2**21 - 1), + (-(2**53 - 1), 1, -(2**21)), ] @@ -115,6 +117,8 @@ def test_symbolic_immediate_is_rejected(lowered_isa): str(SAddU64(dst=sgpr(4, 2), src0=sgpr(4, 2), src1="SomeSymbol")) -def test_immediate_outside_the_64_bit_range_is_rejected(lowered_isa): +# A double cannot tell these from a neighbor: 2^53 + 1 and 2^63 - 1 round to 2^53 and 2^63. +@pytest.mark.parametrize("imm", [2**53, 2**53 + 1, -(2**53), 2**63 - 1, -(2**63), 2**63]) +def test_immediate_a_double_cannot_hold_exactly_is_rejected(lowered_isa, imm): with pytest.raises(Exception, match="cannot split"): - str(SAddU64(dst=sgpr(4, 2), src0=sgpr(4, 2), src1=2**63)) + str(SAddU64(dst=sgpr(4, 2), src0=sgpr(4, 2), src1=imm)) From 685081e0c2e7d729e7c7f8f5a4f626d2a8c46cea Mon Sep 17 00:00:00 2001 From: Tony Davis Date: Sun, 4 Oct 2026 11:51:13 -0400 Subject: [PATCH 08/11] fix(tensilelite): AIHPBLAS-5022 lint subtracts without carry and every returning-atomic form The lint missed v_sub_i32 and v_subrev_u32 (and the gfx10+ v_subrev_nc_u32, v_add_nc_i32 and v_sub_nc_i32), so a low-dword decrement with any of them went unchecked. It also recognized a returning global or flat atomic only by a _rtn suffix or glc; gfx94x and gfx950 mark it with sc0 and gfx12 with th:TH_ATOMIC_RETURN. For those it took the address (operand 1) for data and could take the data pair for an address. Returning atomics are now recognized by any of the four, and stores never count as writing a register. Tests cover each opcode and each return form. Co-authored-by: Cursor --- .../Tests/unit/test_address_carry_lint.py | 45 +++++++++++++++++++ .../Tensile/Utilities/address_carry_lint.py | 12 ++++- 2 files changed, 55 insertions(+), 2 deletions(-) diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_address_carry_lint.py b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_address_carry_lint.py index 672b1ba4227e..173a7826a23b 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_address_carry_lint.py +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_address_carry_lint.py @@ -125,6 +125,51 @@ def test_a_register_reused_for_other_arithmetic_is_not_reported(): assert _reasons(asm) == [] +@pytest.mark.parametrize( + "op", + ["v_sub_u32", "v_subrev_u32", "v_sub_i32", "v_subrev_nc_u32", "v_add_nc_i32", "v_sub_nc_i32"], +) +def test_vector_subtract_without_carry_on_an_address_is_reported(op): + asm = f""" + {op} v4, v4, v6 + global_load_dwordx2 v[0:1], v[4:5], off + """ + reasons = _reasons(asm) + assert len(reasons) == 1 and op in reasons[0] and "v4" in reasons[0] + + +# A returning atomic is dst, vaddr, data[, saddr]; one that does not return is vaddr, data[, saddr]. +@pytest.mark.parametrize("marker", ["glc", "sc0", "th:TH_ATOMIC_RETURN"]) +def test_a_returning_atomic_address_is_its_second_operand(marker): + bad = f""" + v_add_u32 v2, v2, v6 + global_atomic_add_u64 v[0:1], v[2:3], v[4:5], off {marker} + """ + reasons = _reasons(bad) + assert len(reasons) == 1 and "v2" in reasons[0] + data_only = f""" + v_add_u32 v4, v4, v6 + v_add_co_u32 v2, vcc, v2, v6 + v_addc_co_u32 v3, vcc, v3, 0, vcc + global_atomic_add_u64 v[0:1], v[2:3], v[4:5], off {marker} + """ + assert _reasons(data_only) == [] + + +def test_a_non_returning_atomic_address_is_its_first_operand(): + asm = """ + v_add_u32 v4, v4, v6 + global_atomic_add_u64 v[2:3], v[4:5], off + """ + assert _reasons(asm) == [] + asm = """ + v_add_u32 v2, v2, v6 + global_atomic_add_u64 v[2:3], v[4:5], off + """ + reasons = _reasons(asm) + assert len(reasons) == 1 and "v2" in reasons[0] + + def test_data_registers_of_a_buffer_load_are_not_addresses(): asm = """ buffer_load_dwordx4 v[34:37], v5, s[40:43], 0 offen diff --git a/projects/hipblaslt/tensilelite/Tensile/Utilities/address_carry_lint.py b/projects/hipblaslt/tensilelite/Tensile/Utilities/address_carry_lint.py index 3d730b5a908e..7feebb2b94cb 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Utilities/address_carry_lint.py +++ b/projects/hipblaslt/tensilelite/Tensile/Utilities/address_carry_lint.py @@ -57,10 +57,18 @@ "s_addk_i32", "v_add_u32", "v_sub_u32", + "v_subrev_u32", "v_add_nc_u32", "v_sub_nc_u32", + "v_subrev_nc_u32", "v_add_i32", + "v_sub_i32", + "v_add_nc_i32", + "v_sub_nc_i32", } +# How a global or flat atomic asks for the old value: a _rtn mnemonic, glc (gfx90a), sc0 +# (gfx94x and gfx950) or th:TH_ATOMIC_RETURN (gfx12 and later). +_ATOMIC_RETURN = re.compile(r"_rtn\b|\bglc\b|\bsc0\b|\bTH_ATOMIC_RETURN\b") # Instructions that rewrite a whole pair, which ends any obligation on its low dword. PAIR_DEFS = {"s_mov_b64", "v_mov_b64", "v_lshlrev_b64", "s_lshl_b64", "s_add_u64", "v_add_nc_u64"} # Unconditional transfers: the next instruction in the listing is not the next one executed. @@ -157,7 +165,7 @@ def address_operands(inst: Instruction) -> list[Reg]: # Leave out the data and destination operands: a load writes operand 0, a store reads # its data from operand 1, and a returning atomic does both. if "_atomic" in m: - data = {0, 2} if (m.endswith("_rtn") or "glc" in inst.text) else {1} + data = {0, 2} if _ATOMIC_RETURN.search(inst.text) else {1} elif "_store" in m: data = {1} else: @@ -173,7 +181,7 @@ def written(inst: Instruction) -> list[Reg]: ): return [] if re.search(r"_store|_atomic", inst.mnemonic) and not ( - inst.mnemonic.endswith("_rtn") or "glc" in inst.text + "_atomic" in inst.mnemonic and _ATOMIC_RETURN.search(inst.text) ): return [] return parse_regs(inst.operands[0]) From ab76f47591288eeec4074320534a6abd5d310d0d Mon Sep 17 00:00:00 2001 From: Tony Davis Date: Sun, 4 Oct 2026 15:21:39 -0400 Subject: [PATCH 09/11] fix(tensilelite): AIHPBLAS-5022 follow carries out of carry-in adds and partial carry writes A carry-in add (v_add_co_ci_u32, v_sub_co_ci_u32, v_addc_co_u32 and the rest) also sets a carry out, but the lint only followed carries from the plain carry-out adds, so one updating an address low dword without carrying into the high dword was accepted. It also compared the carry register as text, so a write to part of it, such as s_mov_b32 s20 while the carry sits in s[20:21], or vcc_lo while it sits in vcc, did not count as changing it. Carry registers are now compared by the registers (or vcc halves) they occupy. Tests cover three carry-in forms and three partial writes. Co-authored-by: Cursor --- .../Tests/unit/test_address_carry_lint.py | 42 ++++++++++++++++++ .../Tensile/Utilities/address_carry_lint.py | 43 ++++++++++++++----- 2 files changed, 75 insertions(+), 10 deletions(-) diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_address_carry_lint.py b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_address_carry_lint.py index 173a7826a23b..eb94c9eb095e 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_address_carry_lint.py +++ b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_address_carry_lint.py @@ -267,6 +267,48 @@ def test_a_vector_carry_overwritten_before_the_carry_in_is_reported(): assert len(reasons) == 1 and "v4" in reasons[0] +# A carry-in add also sets a carry out, so on a low dword its carry must reach the high dword. +@pytest.mark.parametrize( + "op,carry", + [("v_add_co_ci_u32", "vcc_lo"), ("v_sub_co_ci_u32", "vcc_lo"), ("v_addc_co_u32", "vcc")], +) +def test_a_carry_in_add_on_a_low_dword_must_carry_on(op, carry): + good = f""" + {op} v4, {carry}, v4, v6, {carry} + v_add_co_ci_u32 v5, {carry}, v5, 0, {carry} + global_load_dwordx2 v[0:1], v[4:5], off + """ + bad = f""" + {op} v4, {carry}, v4, v6, {carry} + global_load_dwordx2 v[0:1], v[4:5], off + """ + assert _reasons(good) == [] + reasons = _reasons(bad) + assert len(reasons) == 1 and op in reasons[0] and "v4" in reasons[0] + + +@pytest.mark.parametrize( + "carry,partial", + [("s[20:21]", "s_mov_b32 s20, 0"), ("s[20:21]", "s_mov_b32 s21, 0"), ("vcc", "s_mov_b32 vcc_lo, 0")], +) +def test_a_partial_write_to_the_carry_register_is_reported(carry, partial): + good = f""" + v_add_co_u32 v4, {carry}, v4, v6 + s_mov_b32 s30, 0 + v_addc_co_u32 v5, {carry}, v5, 0, {carry} + global_load_dwordx2 v[0:1], v[4:5], off + """ + bad = f""" + v_add_co_u32 v4, {carry}, v4, v6 + {partial} + v_addc_co_u32 v5, {carry}, v5, 0, {carry} + global_load_dwordx2 v[0:1], v[4:5], off + """ + assert _reasons(good) == [] + reasons = _reasons(bad) + assert len(reasons) == 1 and "v4" in reasons[0] + + def test_registers_are_judged_within_their_own_kernel(): asm = """ .amdgpu_hsa_kernel first diff --git a/projects/hipblaslt/tensilelite/Tensile/Utilities/address_carry_lint.py b/projects/hipblaslt/tensilelite/Tensile/Utilities/address_carry_lint.py index 7feebb2b94cb..f6f539853483 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Utilities/address_carry_lint.py +++ b/projects/hipblaslt/tensilelite/Tensile/Utilities/address_carry_lint.py @@ -49,7 +49,10 @@ "v_subbrev_co_u32", "v_add_co_ci_u32", "v_sub_co_ci_u32", + "v_subrev_co_ci_u32", } +# A carry-in add sets a carry out as well, so on a low dword it must be carried on too. +CARRY_SETTERS = CARRY_OUT | CARRY_IN # Adds that update one dword and drop any carry. NO_CARRY = { "s_add_i32", @@ -128,6 +131,19 @@ def parse_regs(operand: str) -> list[Reg]: return [lo.plus(i) for i in range(hi.offset - lo.offset + 1)] +_VCC_PARTS = {"vcc": {"vcc_lo", "vcc_hi"}, "vcc_lo": {"vcc_lo"}, "vcc_hi": {"vcc_hi"}} + + +def _carry_parts(operand: str) -> set: + """What a carry operand occupies: its registers, or the halves of vcc, so that a write to + any part of a carry register is seen as changing it.""" + operand = operand.strip() + if operand in _VCC_PARTS: + return set(_VCC_PARTS[operand]) + regs = parse_regs(operand) + return set(regs) if regs else ({operand} if operand else set()) + + @dataclass class Instruction: line: int @@ -260,22 +276,29 @@ def consumes(j: int, low: Reg, high: Reg) -> bool: insts[j].mnemonic in PAIR_DEFS and writes[j][:1] == [low] ) + def carry_writes(j: int) -> set: + """The carry-register parts instruction j writes: its destination, including vcc, and + the carry-out of another add.""" + ops = insts[j].operands + out = set(writes[j]) + if ops and ops[0] in _VCC_PARTS: + out |= _carry_parts(ops[0]) + if insts[j].mnemonic in CARRY_SETTERS and len(ops) > 1: + out |= _carry_parts(ops[1]) + return out + def vector_carried(start: int, low: Reg, high: Reg) -> bool: """Whether a vector carry reaches high: a carry-in to high that reads the carry register - the low add wrote, before anything else writes that register.""" - carry = insts[start].operands[1] if len(insts[start].operands) > 1 else "" + the low add wrote, before anything writes any part of that register.""" + ops = insts[start].operands + carry = _carry_parts(ops[1]) if len(ops) > 1 else set() for j in range(start + 1, min(len(insts), start + 1 + WINDOW)): ops = insts[j].operands if insts[j].mnemonic in PAIR_DEFS and writes[j][:1] == [low]: return True if insts[j].mnemonic in CARRY_IN and writes[j][:1] == [high]: - return bool(ops) and ops[-1] == carry - # A carry register is written as the destination of a compare or move, or as the - # carry-out of another add. - if ops and ( - ops[0] == carry - or (insts[j].mnemonic in CARRY_OUT | CARRY_IN and len(ops) > 1 and ops[1] == carry) - ): + return bool(ops) and bool(carry) and _carry_parts(ops[-1]) == carry + if carry_writes(j) & carry: return False return False @@ -308,7 +331,7 @@ def scalar_carried(start: int, low: Reg, high: Reg) -> bool: f"carry into {high}", ) ) - elif inst.mnemonic in CARRY_OUT: + elif inst.mnemonic in CARRY_SETTERS: if inst.mnemonic.startswith("s_"): carried, where = scalar_carried(i, low, high), "before SCC changes" else: From a3b9f86e92919312918ac7fbc2e8099bd52d616b Mon Sep 17 00:00:00 2001 From: Tony Davis Date: Sun, 4 Oct 2026 17:48:41 -0400 Subject: [PATCH 10/11] fix(rocisa): AIHPBLAS-5022 put a 64-bit vector add's immediate where gfx9 can encode it The lowered VAddNCU64 put an immediate in src1 of v_add_co_u32, where VOP2 takes only a VGPR and VOP3 on gfx9 takes no literal, so any immediate outside the inline-constant range failed to assemble (as it did before the split fix). The immediate now goes in src0, which takes a literal. v_addc_co_u32 reads vcc as well, which leaves no room for a literal at all, so an immediate whose high half is not an inline constant (-16 to 64) is rejected; every int immediate has a high half of 0 or -1. The lowering test now also assembles every emitted pair for gfx90a, gfx942 and gfx950. With the immediate left in src1, 33 of its cases fail. Co-authored-by: Cursor --- .../rocisa/include/instruction/common.hpp | 16 +++++++ .../rocisa/test/test_lowering64.py | 46 +++++++++++++++++-- 2 files changed, 59 insertions(+), 3 deletions(-) diff --git a/projects/hipblaslt/tensilelite/rocisa/rocisa/include/instruction/common.hpp b/projects/hipblaslt/tensilelite/rocisa/rocisa/include/instruction/common.hpp index 3d309e042485..2ff156bbee1a 100644 --- a/projects/hipblaslt/tensilelite/rocisa/rocisa/include/instruction/common.hpp +++ b/projects/hipblaslt/tensilelite/rocisa/rocisa/include/instruction/common.hpp @@ -3565,6 +3565,22 @@ namespace rocisa std::vector srcs1; std::vector srcs2; splitSrcs64(srcs, srcs1, srcs2); + // VOP2 takes a literal only in src0, so an immediate goes first (the add is + // commutative). The carry-in add reads vcc as well, which leaves no room for a + // literal at all: its half of an immediate must be an inline constant. + if(std::holds_alternative(srcs1[1]) && !std::holds_alternative(srcs1[0])) + { + std::swap(srcs1[0], srcs1[1]); + std::swap(srcs2[0], srcs2[1]); + } + for(const auto& s : srcs2) + if(std::holds_alternative(s) + && (std::get(s) < -16 || std::get(s) > 64)) + throw std::invalid_argument( + "64-bit vector add: cannot split an immediate whose high half " + + std::to_string(std::get(s)) + + " is not an inline constant (-16 to 64); v_addc_co_u32 cannot " + "encode it next to its carry"); auto vcc = std::make_shared(); instructions = {std::make_shared(dst1, vcc, srcs1[0], srcs1[1], comment), diff --git a/projects/hipblaslt/tensilelite/rocisa/test/test_lowering64.py b/projects/hipblaslt/tensilelite/rocisa/test/test_lowering64.py index 45d72e616544..af74ce34d9ad 100644 --- a/projects/hipblaslt/tensilelite/rocisa/test/test_lowering64.py +++ b/projects/hipblaslt/tensilelite/rocisa/test/test_lowering64.py @@ -13,6 +13,7 @@ class C2 on AIHPBLAS-4988). import os import shutil +import subprocess import pytest @@ -103,13 +104,52 @@ def test_vaddncu64_registers_lower_to_a_carry_chain(lowered_isa): assert lo_ops[1] == hi_ops[1] == hi_ops[4] +def _inline(v): + return -16 <= v <= 64 + + @pytest.mark.parametrize("imm,low,high", _IMMEDIATES) def test_vaddncu64_immediate_high_half_is_the_sign_extension(lowered_isa, imm, low, high): - lines = _lines(VAddNCU64(dst=vgpr(4, 2), src0=vgpr(4, 2), src1=imm)) + inst = VAddNCU64(dst=vgpr(4, 2), src0=vgpr(4, 2), src1=imm) + if not _inline(high): + # v_addc_co_u32 reads vcc, which leaves no room for a literal next to it. + with pytest.raises(Exception, match="not an inline constant"): + str(inst) + return + lines = _lines(inst) assert [m for m, _ in lines] == ["v_add_co_u32", "v_addc_co_u32"] (_, lo_ops), (_, hi_ops) = lines - assert _as_int(lo_ops[3]) == low - assert _as_int(hi_ops[3]) == high, f"high half of {imm:#x} must be {high}, got {hi_ops[3]}" + # VOP2 takes a literal only in src0, so the immediate comes first. + assert _as_int(lo_ops[2]) == low and lo_ops[3] == "v4" + assert _as_int(hi_ops[2]) == high and hi_ops[3] == "v5" + + +def _assemble(isa, text, tmp_path): + rocm_path = os.environ.get("ROCM_PATH", "/opt/rocm") + search_path = os.pathsep.join( + [os.path.join(rocm_path, "bin"), os.path.join(rocm_path, "lib", "llvm", "bin")] + ) + clang = shutil.which("amdclang++", path=search_path) or shutil.which("amdclang++") + if not clang: + pytest.skip("amdclang++ not found") + src = tmp_path / "lowered.s" + src.write_text(text + "\n") + return subprocess.run( + [clang, "-x", "assembler", "-target", "amdgcn-amd-amdhsa", + f"-mcpu={rocisa.isaToGfx(isa)}", "-c", str(src), "-o", str(tmp_path / "lowered.o")], + capture_output=True, text=True, + ) # fmt: skip + + +# String checks miss an operand the hardware cannot encode; the assembler does not. +@pytest.mark.parametrize("imm,low,high", _IMMEDIATES) +def test_lowered_immediate_adds_assemble(lowered_isa, imm, low, high, tmp_path): + insts = [SAddU64(dst=sgpr(4, 2), src0=sgpr(4, 2), src1=imm)] + if _inline(high): + insts.append(VAddNCU64(dst=vgpr(4, 2), src0=vgpr(4, 2), src1=imm)) + text = "\n".join(str(i) for i in insts) + result = _assemble(lowered_isa, text, tmp_path) + assert result.returncode == 0, f"{text}\n{result.stderr}" def test_symbolic_immediate_is_rejected(lowered_isa): From 2ac9fca28440087070e0e8629b4df818b63620ac Mon Sep 17 00:00:00 2001 From: Tony Davis Date: Sun, 4 Oct 2026 17:53:26 -0400 Subject: [PATCH 11/11] refactor(tensilelite): AIHPBLAS-5022 move the address-carry lint to its own PR The lint and its tests move to #13060 for their own review, so this PR is only the rocisa lowering fix and its test. Co-authored-by: Cursor --- .../Tests/unit/test_address_carry_lint.py | 320 --------------- .../Tensile/Utilities/address_carry_lint.py | 385 ------------------ 2 files changed, 705 deletions(-) delete mode 100644 projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_address_carry_lint.py delete mode 100644 projects/hipblaslt/tensilelite/Tensile/Utilities/address_carry_lint.py diff --git a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_address_carry_lint.py b/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_address_carry_lint.py deleted file mode 100644 index eb94c9eb095e..000000000000 --- a/projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_address_carry_lint.py +++ /dev/null @@ -1,320 +0,0 @@ -# Copyright Advanced Micro Devices, Inc., or its affiliates. -# SPDX-License-Identifier: MIT - -"""The address-carry lint must report a 64-bit address update that drops the carry, and stay -quiet on correct carry chains and on registers that are only reused for other arithmetic.""" - -import functools -import os - -import pytest - -from Tensile.Utilities.address_carry_lint import lint - -pytestmark = pytest.mark.unit - -_DESIGNED = os.path.join( - os.path.dirname(__file__), "characterization", "_codegen", "data", "test_data", "_designed" -) - -# Generated kernels from designed configs that cover the address arithmetic most at risk: batch -# and stride address setup, Stream-K and GSU workspace addressing, fp8 and MX scale loads, and -# TDM descriptors. gfx90a, gfx942 and gfx950 have no native 64-bit add, so every address update -# is a 32-bit carry chain there. -_GENERATED = [ - ("gfx90a", "rich_gemm.yaml"), - ("gfx942", "asmaddr_initstrides.yaml"), - ("gfx942", "s00_loadbatchedaddress_non_stridedba.yaml"), - ("gfx942", "streamk.yaml"), - ("gfx942", "streamk_fixup_tree.yaml"), - ("gfx942", "gsu.yaml"), - ("gfx942", "fp8_gr_conv.yaml"), - ("gfx950", "mx_fp8_scale_swizzle.yaml"), - ("gfx950", "mx_bias_act_gsu.yaml"), - ("gfx950", "subtile.yaml"), - ("gfx1250", "streamk.yaml"), - ("gfx1250", "streamk_tdm_prefetchgl2.yaml"), -] - - -@functools.lru_cache(maxsize=None) -def _assembler_supports(arch): - from Tensile.Common.Architectures import gfxToIsa - from Tensile.Common.Capabilities import makeIsaInfoMap - from Tensile.Toolchain.Validators import validateToolchain - - isa = gfxToIsa(arch) - return bool(makeIsaInfoMap([isa], validateToolchain("amdclang++"))[isa].asmCaps["SupportedISA"]) - - -@pytest.mark.parametrize("arch,config", _GENERATED, ids=lambda x: str(x).replace(".yaml", "")) -def test_generated_kernels_carry_every_address_update(arch, config): - from config_harness import emit_kernels_from_config - - if not _assembler_supports(arch): - pytest.skip(f"amdclang++ in this environment does not support {arch}") - - results = emit_kernels_from_config( - os.path.join(_DESIGNED, arch, config), limit=4, arch=arch, canonical=False - ) - assert results, f"{arch}/{config} emitted no kernels" - for base, src, err in results: - assert err == 0, f"{base} failed to emit" - findings = lint(src) - assert not findings, f"{base}:\n" + "\n".join(str(f) for f in findings) - - -def _reasons(asm): - return [f.reason for f in lint(asm)] - - -def test_scalar_carry_chain_on_a_descriptor_base_is_clean(): - asm = """ - s_add_u32 s[sgprSrdA+0], s[sgprSrdA+0], s[sgprGlobalReadIncsA] - s_addc_u32 s[sgprSrdA+1], s[sgprSrdA+1], 0 - buffer_load_dwordx4 v[0:3], v4, s[sgprSrdA:sgprSrdA+3], 0 offen - """ - assert _reasons(asm) == [] - - -def test_scalar_carry_that_never_reaches_the_high_dword_is_reported(): - # The ROCM-32046 shape: the low dword is advanced, the carry is dropped. - asm = """ - s_add_u32 s[sgprSrdA+0], s[sgprSrdA+0], s[sgprGlobalReadIncsA] - s_mov_b32 s[sgprTmp], 0 - buffer_load_dwordx4 v[0:3], v4, s[sgprSrdA:sgprSrdA+3], 0 offen - """ - reasons = _reasons(asm) - assert len(reasons) == 1 - assert "s[sgprSrdA+0]" in reasons[0] and "s[sgprSrdA+1]" in reasons[0] - - -def test_add_without_carry_on_an_address_low_dword_is_reported(): - asm = """ - s_add_i32 s8, s8, 64 - s_load_dwordx2 s[10:11], s[8:9], 0x0 - """ - reasons = _reasons(asm) - assert len(reasons) == 1 and "s_add_i32" in reasons[0] and "s9" in reasons[0] - - -def test_vector_global_address_needs_a_carry_chain(): - good = """ - v_add_co_u32 v4, vcc, v4, v6 - v_addc_co_u32 v5, vcc, v5, 0, vcc - global_load_dwordx2 v[0:1], v[4:5], off - """ - bad = """ - v_add_u32 v4, v4, v6 - global_load_dwordx2 v[0:1], v[4:5], off - """ - assert _reasons(good) == [] - reasons = _reasons(bad) - assert len(reasons) == 1 and "v4" in reasons[0] - - -def test_a_register_reused_for_other_arithmetic_is_not_reported(): - # s20 is an address elsewhere, but this sum is overwritten before any address use. - asm = """ - s_load_dwordx2 s[30:31], s[20:21], 0x0 - s_sub_u32 s20, s2, s20 - s_add_u32 s20, s20, s3 - s_mov_b64 s[20:21], s[0:1] - s_load_dwordx2 s[32:33], s[20:21], 0x8 - """ - assert _reasons(asm) == [] - - -@pytest.mark.parametrize( - "op", - ["v_sub_u32", "v_subrev_u32", "v_sub_i32", "v_subrev_nc_u32", "v_add_nc_i32", "v_sub_nc_i32"], -) -def test_vector_subtract_without_carry_on_an_address_is_reported(op): - asm = f""" - {op} v4, v4, v6 - global_load_dwordx2 v[0:1], v[4:5], off - """ - reasons = _reasons(asm) - assert len(reasons) == 1 and op in reasons[0] and "v4" in reasons[0] - - -# A returning atomic is dst, vaddr, data[, saddr]; one that does not return is vaddr, data[, saddr]. -@pytest.mark.parametrize("marker", ["glc", "sc0", "th:TH_ATOMIC_RETURN"]) -def test_a_returning_atomic_address_is_its_second_operand(marker): - bad = f""" - v_add_u32 v2, v2, v6 - global_atomic_add_u64 v[0:1], v[2:3], v[4:5], off {marker} - """ - reasons = _reasons(bad) - assert len(reasons) == 1 and "v2" in reasons[0] - data_only = f""" - v_add_u32 v4, v4, v6 - v_add_co_u32 v2, vcc, v2, v6 - v_addc_co_u32 v3, vcc, v3, 0, vcc - global_atomic_add_u64 v[0:1], v[2:3], v[4:5], off {marker} - """ - assert _reasons(data_only) == [] - - -def test_a_non_returning_atomic_address_is_its_first_operand(): - asm = """ - v_add_u32 v4, v4, v6 - global_atomic_add_u64 v[2:3], v[4:5], off - """ - assert _reasons(asm) == [] - asm = """ - v_add_u32 v2, v2, v6 - global_atomic_add_u64 v[2:3], v[4:5], off - """ - reasons = _reasons(asm) - assert len(reasons) == 1 and "v2" in reasons[0] - - -def test_data_registers_of_a_buffer_load_are_not_addresses(): - asm = """ - buffer_load_dwordx4 v[34:37], v5, s[40:43], 0 offen - v_sub_u32 v34, s12, v34 - buffer_load_dwordx4 v[34:37], v5, s[40:43], 0 offen - """ - assert _reasons(asm) == [] - - -def test_a_long_branch_offset_is_not_an_address(): - # The gfx950 long-branch sequence: s64 holds a constant branch offset, and the next use of - # s[64:65] as an address is in another block, past the jump. - asm = """ - s_getpc_b64 s[62:63] - s_add_i32 s64, 0x424f4, 4 - s_add_u32 s62, s62, s64 - s_addc_u32 s63, s63, 0 - s_setpc_b64 s[62:63] - s_cmp_eq_u64 s[64:65], 0 - s_load_dword s8, s[64:65], 0x0 - """ - assert _reasons(asm) == [] - - -def test_an_address_update_before_a_jump_is_still_reported(): - asm = """ - s_add_i32 s8, s8, 64 - s_load_dwordx2 s[10:11], s[8:9], 0x0 - s_branch label_next - """ - reasons = _reasons(asm) - assert len(reasons) == 1 and "s8" in reasons[0] - - -def test_a_scalar_carry_moved_far_by_the_scheduler_is_clean(): - # The gfx950 fp8 kernels: the carry-in comes 40 vector instructions after the carry-out, and - # nothing in between writes SCC. - filler = "\n".join(" v_perm_b32 v27, v55, v54, s78" for _ in range(40)) - asm = f""" - s_add_u32 s68, s68, s76 -{filler} - s_addc_u32 s69, s69, 0 - buffer_load_dwordx4 v8, s[68:71], 0 offen lds - """ - assert _reasons(asm) == [] - - -def test_a_scalar_carry_lost_to_an_scc_write_is_reported(): - asm = """ - s_add_u32 s68, s68, s76 - s_cmp_eq_u32 s73, 0 - s_addc_u32 s69, s69, 0 - buffer_load_dwordx4 v8, s[68:71], 0 offen lds - """ - reasons = _reasons(asm) - assert len(reasons) == 1 and "before SCC changes" in reasons[0] - - -def test_findings_keep_their_line_numbers_across_kernels(): - asm = "\n".join( - [ - ".amdgpu_hsa_kernel first", - " s_load_dwordx2 s[10:11], s[4:5], 0x0", - ".amdgpu_hsa_kernel second", - " s_nop 0", - " s_add_i32 s8, s8, 64", - " s_load_dwordx2 s[10:11], s[8:9], 0x0", - ] - ) - findings = lint(asm) - assert [f.line for f in findings] == [5] - - -def test_tab_separated_instructions_are_parsed(): - asm = "\ts_add_i32\ts8, s8, 64\n\ts_load_dwordx2\ts[10:11], s[8:9], 0x0\n" - reasons = _reasons(asm) - assert len(reasons) == 1 and "s8" in reasons[0] - - -def test_a_vector_carry_overwritten_before_the_carry_in_is_reported(): - good = """ - v_add_co_u32 v4, vcc, v4, v6 - v_add_co_u32 v10, s[20:21], v10, v12 - v_addc_co_u32 v5, vcc, v5, 0, vcc - global_load_dwordx2 v[0:1], v[4:5], off - """ - bad = """ - v_add_co_u32 v4, vcc, v4, v6 - v_add_co_u32 v10, vcc, v10, v12 - v_addc_co_u32 v5, vcc, v5, 0, vcc - global_load_dwordx2 v[0:1], v[4:5], off - """ - assert _reasons(good) == [] - reasons = _reasons(bad) - assert len(reasons) == 1 and "v4" in reasons[0] - - -# A carry-in add also sets a carry out, so on a low dword its carry must reach the high dword. -@pytest.mark.parametrize( - "op,carry", - [("v_add_co_ci_u32", "vcc_lo"), ("v_sub_co_ci_u32", "vcc_lo"), ("v_addc_co_u32", "vcc")], -) -def test_a_carry_in_add_on_a_low_dword_must_carry_on(op, carry): - good = f""" - {op} v4, {carry}, v4, v6, {carry} - v_add_co_ci_u32 v5, {carry}, v5, 0, {carry} - global_load_dwordx2 v[0:1], v[4:5], off - """ - bad = f""" - {op} v4, {carry}, v4, v6, {carry} - global_load_dwordx2 v[0:1], v[4:5], off - """ - assert _reasons(good) == [] - reasons = _reasons(bad) - assert len(reasons) == 1 and op in reasons[0] and "v4" in reasons[0] - - -@pytest.mark.parametrize( - "carry,partial", - [("s[20:21]", "s_mov_b32 s20, 0"), ("s[20:21]", "s_mov_b32 s21, 0"), ("vcc", "s_mov_b32 vcc_lo, 0")], -) -def test_a_partial_write_to_the_carry_register_is_reported(carry, partial): - good = f""" - v_add_co_u32 v4, {carry}, v4, v6 - s_mov_b32 s30, 0 - v_addc_co_u32 v5, {carry}, v5, 0, {carry} - global_load_dwordx2 v[0:1], v[4:5], off - """ - bad = f""" - v_add_co_u32 v4, {carry}, v4, v6 - {partial} - v_addc_co_u32 v5, {carry}, v5, 0, {carry} - global_load_dwordx2 v[0:1], v[4:5], off - """ - assert _reasons(good) == [] - reasons = _reasons(bad) - assert len(reasons) == 1 and "v4" in reasons[0] - - -def test_registers_are_judged_within_their_own_kernel(): - asm = """ - .amdgpu_hsa_kernel first - s_add_i32 s4, s4, 1 - s_endpgm - .amdgpu_hsa_kernel second - s_load_dwordx2 s[10:11], s[4:5], 0x0 - """ - assert _reasons(asm) == [] diff --git a/projects/hipblaslt/tensilelite/Tensile/Utilities/address_carry_lint.py b/projects/hipblaslt/tensilelite/Tensile/Utilities/address_carry_lint.py deleted file mode 100644 index f6f539853483..000000000000 --- a/projects/hipblaslt/tensilelite/Tensile/Utilities/address_carry_lint.py +++ /dev/null @@ -1,385 +0,0 @@ -# Copyright Advanced Micro Devices, Inc., or its affiliates. -# SPDX-License-Identifier: MIT - -"""Lint AMDGPU assembly for 64-bit address updates that drop the carry. - -A 64-bit address held in a register pair must be advanced with a carry chain: a low add that -produces a carry (s_add_u32, v_add_co_u32) followed by a high add that consumes it (s_addc_u32, -v_addc_co_u32). An add to the low dword with no carry into the high dword is correct only until -the low dword wraps, which happens when the buffer crosses a 4 GiB boundary: the access then -lands 4 GiB away (ROCM-32046, class C on AIHPBLAS-4988). - -The lint finds the registers used as the low dword of an address: the base of a buffer resource -descriptor (s[N:N+3] in a buffer instruction), the address pair of a scalar load (s[A:A+1]), and -the address pair or saddr of a global or flat access. For each write to such a register by an add -or subtract, it requires either a carry-producing instruction followed by a carry-consuming write -to the next register, or a full redefinition of the pair. A scalar carry is followed until SCC -changes, since the scheduler can move the carry-in far from the carry-out; a vector carry, for a -short window and only until its carry register is written again. Anything else is reported, if the -updated value is next used as an address before it is overwritten or an unconditional jump. An add -of two constants sets the register rather than advancing an address, so it is not reported. The -scan is otherwise linear, so a finding is a lead to read, not a proof. - -It is meant for Tensile-generated and hand-written kernels, which keep a 64-bit value in an -adjacent register pair. Compiler-generated code may keep the two halves of a sum in unrelated -registers, so the lint reports false findings there and is not meant for it. - - python -m Tensile.Utilities.address_carry_lint kernel.s - python -m Tensile.Utilities.address_carry_lint --disassemble TensileLibrary_gfx942.co -""" - -from __future__ import annotations - -import argparse -import re -import shutil -import subprocess -import sys -from dataclasses import dataclass -from pathlib import Path -from typing import Iterable, Optional - -# Low adds that set a carry, and the high adds that must consume it. -CARRY_OUT = {"s_add_u32", "s_sub_u32", "v_add_co_u32", "v_sub_co_u32", "v_subrev_co_u32"} -CARRY_IN = { - "s_addc_u32", - "s_subb_u32", - "v_addc_co_u32", - "v_subb_co_u32", - "v_subbrev_co_u32", - "v_add_co_ci_u32", - "v_sub_co_ci_u32", - "v_subrev_co_ci_u32", -} -# A carry-in add sets a carry out as well, so on a low dword it must be carried on too. -CARRY_SETTERS = CARRY_OUT | CARRY_IN -# Adds that update one dword and drop any carry. -NO_CARRY = { - "s_add_i32", - "s_sub_i32", - "s_addk_i32", - "v_add_u32", - "v_sub_u32", - "v_subrev_u32", - "v_add_nc_u32", - "v_sub_nc_u32", - "v_subrev_nc_u32", - "v_add_i32", - "v_sub_i32", - "v_add_nc_i32", - "v_sub_nc_i32", -} -# How a global or flat atomic asks for the old value: a _rtn mnemonic, glc (gfx90a), sc0 -# (gfx94x and gfx950) or th:TH_ATOMIC_RETURN (gfx12 and later). -_ATOMIC_RETURN = re.compile(r"_rtn\b|\bglc\b|\bsc0\b|\bTH_ATOMIC_RETURN\b") -# Instructions that rewrite a whole pair, which ends any obligation on its low dword. -PAIR_DEFS = {"s_mov_b64", "v_mov_b64", "v_lshlrev_b64", "s_lshl_b64", "s_add_u64", "v_add_nc_u64"} -# Unconditional transfers: the next instruction in the listing is not the next one executed. -JUMPS = {"s_branch", "s_setpc_b64", "s_endpgm"} -# Scalar instructions that write SCC, which ends a scalar carry chain. -_SCC_WRITERS = re.compile( - r"^s_(add|sub|addc|subb|addk|cmp|cmpk|bitcmp|and|or|xor|andn[12]|orn[12]|nand|nor|xnor|not|" - r"lshl|lshr|ashr|bfe|abs|absdiff|min|max|bcnt|quadmask|wqm)" -) - -WINDOW = 16 -# How far a written low dword is followed to its next use. -FLOW_WINDOW = 400 - -_REG = re.compile(r"^(?P[sv])(?:(?P\d+)|\[(?P[^\]:]+)(?::(?P[^\]]+))?\])") - - -@dataclass(frozen=True) -class Reg: - kind: str # "s" or "v" - base: str # number as text, or a symbol such as sgprSrdA - offset: int - - def plus(self, n: int) -> "Reg": - return Reg(self.kind, self.base, self.offset + n) - - def __str__(self) -> str: - if self.base.isdigit(): - return f"{self.kind}{int(self.base) + self.offset}" - return f"{self.kind}[{self.base}+{self.offset}]" - - -def _single(text: str, kind: str) -> Reg: - text = text.strip() - if text.isdigit(): - return Reg(kind, "0", int(text)) - m = re.match(r"^([A-Za-z_]\w*)(?:\s*\+\s*(\d+))?$", text) - if m: - return Reg(kind, m.group(1), int(m.group(2) or 0)) - return Reg(kind, text, 0) - - -def parse_regs(operand: str) -> list[Reg]: - """The registers an operand names, low to high; empty for a non-register operand.""" - m = _REG.match(operand.strip()) - if not m: - return [] - kind = m.group("kind") - if m.group("num") is not None: - return [Reg(kind, "0", int(m.group("num")))] - lo = _single(m.group("lo"), kind) - if m.group("hi") is None: - return [lo] - hi = _single(m.group("hi"), kind) - if hi.base != lo.base or hi.offset < lo.offset: - return [lo] - return [lo.plus(i) for i in range(hi.offset - lo.offset + 1)] - - -_VCC_PARTS = {"vcc": {"vcc_lo", "vcc_hi"}, "vcc_lo": {"vcc_lo"}, "vcc_hi": {"vcc_hi"}} - - -def _carry_parts(operand: str) -> set: - """What a carry operand occupies: its registers, or the halves of vcc, so that a write to - any part of a carry register is seen as changing it.""" - operand = operand.strip() - if operand in _VCC_PARTS: - return set(_VCC_PARTS[operand]) - regs = parse_regs(operand) - return set(regs) if regs else ({operand} if operand else set()) - - -@dataclass -class Instruction: - line: int - text: str - mnemonic: str - operands: list[str] - - -def parse(asm: str, first_line: int = 1) -> list[Instruction]: - out = [] - for n, raw in enumerate(asm.splitlines(), first_line): - text = re.split(r"//|;", raw)[0].strip() - if not text or text.startswith(".") or text.endswith(":"): - continue - # Hand-written kernels separate the mnemonic from its operands with a tab. - mnemonic, _, rest = text.replace("\t", " ").partition(" ") - if not re.match(r"^[sv]_|^buffer_|^global_|^flat_|^scratch_", mnemonic): - continue - mnemonic = re.sub(r"_e(32|64)$", "", mnemonic) - operands = [o.strip() for o in rest.split(",")] if rest.strip() else [] - out.append(Instruction(n, text, mnemonic, operands)) - return out - - -def address_operands(inst: Instruction) -> list[Reg]: - """The registers this instruction uses as the low dword of an address or descriptor base.""" - regs = [parse_regs(o) for o in inst.operands] - m = inst.mnemonic - if m.startswith(("buffer_", "s_buffer_")): - # The descriptor is the SGPR quadruple; a VGPR quadruple is data. - return [r[0] for r in regs if len(r) == 4 and r[0].kind == "s"] - if m.startswith(("s_load_", "s_store_", "s_scratch_")): - return [regs[1][0]] if len(regs) > 1 and len(regs[1]) == 2 else [] - if m.startswith(("global_", "flat_")): - # Leave out the data and destination operands: a load writes operand 0, a store reads - # its data from operand 1, and a returning atomic does both. - if "_atomic" in m: - data = {0, 2} if _ATOMIC_RETURN.search(inst.text) else {1} - elif "_store" in m: - data = {1} - else: - data = {0} - return [r[0] for k, r in enumerate(regs) if k not in data and len(r) == 2] - return [] - - -def written(inst: Instruction) -> list[Reg]: - """The registers this instruction writes: its first operand, unless it has no destination.""" - if not inst.operands or inst.mnemonic.startswith( - ("s_cmp", "s_bitcmp", "s_cbranch", "s_branch", "s_setpc", "s_waitcnt", "s_nop") - ): - return [] - if re.search(r"_store|_atomic", inst.mnemonic) and not ( - "_atomic" in inst.mnemonic and _ATOMIC_RETURN.search(inst.text) - ): - return [] - return parse_regs(inst.operands[0]) - - -def address_lows(insts: Iterable[Instruction]) -> set[Reg]: - """Registers used as the low dword of a 64-bit address or a descriptor base.""" - lows: set[Reg] = set() - for inst in insts: - lows.update(address_operands(inst)) - return lows - - -@dataclass -class Finding: - line: int - text: str - reason: str - - def __str__(self) -> str: - return f"line {self.line}: {self.text}\n {self.reason}" - - -# A kernel starts at Tensile's label_ASM_Start, at an assembler kernel directive, or, in -# disassembly, at a function symbol header such as :. -_KERNEL_START = re.compile( - r"label_ASM_Start>?:|^\s*\.amdgpu_hsa_kernel\b|^\s*\.globl\b|^<(?!label_)[^>\s]+>:\s*$", - re.M, -) - - -def _kernel_chunks(asm: str) -> list[tuple[int, str]]: - """One (first line number, text) chunk per kernel.""" - starts = [m.start() for m in _KERNEL_START.finditer(asm)] - if not starts: - return [(1, asm)] - bounds = [0] + starts[1:] + [len(asm)] - chunks, line, pos = [], 1, 0 - for a, b in zip(bounds, bounds[1:]): - line += asm.count("\n", pos, a) - pos = a - chunks.append((line, asm[a:b])) - return chunks - - -def kernels(asm: str) -> list[str]: - """Splits assembly into one chunk per kernel, so registers are judged within their kernel.""" - return [text for _, text in _kernel_chunks(asm)] - - -def lint(asm: str) -> list[Finding]: - findings = [] - for first_line, chunk in _kernel_chunks(asm): - findings += _lint_kernel(chunk, first_line) - return findings - - -def _lint_kernel(asm: str, first_line: int = 1) -> list[Finding]: - insts = parse(asm, first_line) - uses = [address_operands(inst) for inst in insts] - writes = [written(inst) for inst in insts] - lows = {r for u in uses for r in u} - - def flows_to_address(low: Reg, start: int) -> bool: - """Whether the value written at start is next used as an address, not overwritten.""" - for j in range(start + 1, min(len(insts), start + 1 + FLOW_WINDOW)): - if low in uses[j]: - return True - if low in writes[j] or insts[j].mnemonic in JUMPS: - return False - return False - - def consumes(j: int, low: Reg, high: Reg) -> bool: - """Whether instruction j adds a carry into high, or redefines the whole pair.""" - return (insts[j].mnemonic in CARRY_IN and writes[j][:1] == [high]) or ( - insts[j].mnemonic in PAIR_DEFS and writes[j][:1] == [low] - ) - - def carry_writes(j: int) -> set: - """The carry-register parts instruction j writes: its destination, including vcc, and - the carry-out of another add.""" - ops = insts[j].operands - out = set(writes[j]) - if ops and ops[0] in _VCC_PARTS: - out |= _carry_parts(ops[0]) - if insts[j].mnemonic in CARRY_SETTERS and len(ops) > 1: - out |= _carry_parts(ops[1]) - return out - - def vector_carried(start: int, low: Reg, high: Reg) -> bool: - """Whether a vector carry reaches high: a carry-in to high that reads the carry register - the low add wrote, before anything writes any part of that register.""" - ops = insts[start].operands - carry = _carry_parts(ops[1]) if len(ops) > 1 else set() - for j in range(start + 1, min(len(insts), start + 1 + WINDOW)): - ops = insts[j].operands - if insts[j].mnemonic in PAIR_DEFS and writes[j][:1] == [low]: - return True - if insts[j].mnemonic in CARRY_IN and writes[j][:1] == [high]: - return bool(ops) and bool(carry) and _carry_parts(ops[-1]) == carry - if carry_writes(j) & carry: - return False - return False - - def scalar_carried(start: int, low: Reg, high: Reg) -> bool: - """Whether a scalar carry reaches high. The scheduler may move the carry-in well away - from the carry-out, so it is followed until SCC changes rather than for a fixed count. - """ - for j in range(start + 1, min(len(insts), start + 1 + FLOW_WINDOW)): - if consumes(j, low, high): - return True - if insts[j].mnemonic in JUMPS or _SCC_WRITERS.match(insts[j].mnemonic): - return False - return False - - findings = [] - for i, inst in enumerate(insts): - dst = writes[i] - if len(dst) != 1 or dst[0] not in lows: - continue - low, high = dst[0], dst[0].plus(1) - if inst.mnemonic in NO_CARRY: - if not any(parse_regs(o) for o in inst.operands[1:]): - continue - if flows_to_address(low, i): - findings.append( - Finding( - inst.line, - inst.text, - f"{inst.mnemonic} updates {low}, the low dword of an address, with no " - f"carry into {high}", - ) - ) - elif inst.mnemonic in CARRY_SETTERS: - if inst.mnemonic.startswith("s_"): - carried, where = scalar_carried(i, low, high), "before SCC changes" - else: - carried = vector_carried(i, low, high) - where = f"within {WINDOW} instructions, from the carry it wrote," - if not carried and flows_to_address(low, i): - findings.append( - Finding( - inst.line, - inst.text, - f"{inst.mnemonic} sets a carry out of {low}, the low dword of an " - f"address, but nothing {where} adds it into {high}", - ) - ) - return findings - - -def disassemble(code_object: Path) -> str: - objdump = shutil.which("llvm-objdump") or "/opt/rocm/llvm/bin/llvm-objdump" - return subprocess.run( - [objdump, "-d", "--no-show-raw-insn", "--no-leading-addr", str(code_object)], - capture_output=True, - text=True, - check=True, - ).stdout - - -def main(argv: Optional[list[str]] = None) -> int: - parser = argparse.ArgumentParser( - description=__doc__, formatter_class=argparse.RawDescriptionHelpFormatter - ) - parser.add_argument("paths", nargs="+", type=Path, help="assembly files or code objects") - parser.add_argument( - "--disassemble", - action="store_true", - help="treat the paths as code objects and disassemble them with llvm-objdump", - ) - args = parser.parse_args(argv) - total = 0 - for path in args.paths: - asm = disassemble(path) if args.disassemble else path.read_text(errors="replace") - findings = lint(asm) - total += len(findings) - for f in findings: - print(f"{path}: {f}") - print(f"{total} finding(s) in {len(args.paths)} file(s)") - return 1 if total else 0 - - -if __name__ == "__main__": - sys.exit(main())