feat(tools): add a hardware oracle for RDNA instruction semantics

Co-Authored-By: Claude <noreply@anthropic.com>

docs/dev/HW_ORACLE.md

@@ -0,0 +1,64 @@
+# Hardware oracle
+
+`tools/hw-oracle/` runs a few lines of RDNA assembly on a local AMD GPU, one input row per lane, and returns what the
+hardware computed. Use it to settle instruction semantics that the ISA documentation, LLVM or Mesa leave open, then pin the
+measured rows as expected values in `tests/execution/`.
+
+The PS5 GPU is gfx1013. Other RDNA2 GPUs (gfx103x) share the VALU, SALU, DS, MUBUF and MIMG encodings, apart from ops that
+exist only on gfx10.1 (`v_mad_legacy_f32`, `v_mac_legacy_f32`, `v_mul_lo_i32`, `s_dcache_discard`, ...). RDNA3 also runs
+the tool, but differs from gfx1013 in more places. Say which GPU you measured on in the PR.
+
+## Requirements
+
+- Linux with the `amdgpu` driver and read/write access to `/dev/kfd` (usually the `render` group). Root is not needed.
+- The ROCm HSA runtime with its headers: `libhsa-runtime-dev` (Debian, Ubuntu), `rocm-runtime-devel` (Fedora),
+  `hsa-rocr` (Arch), or AMD's ROCm packages. A runtime under `ROCM_PATH` or `/opt/rocm` is used first. The rest of ROCm
+  is not needed.
+- A C compiler, and `clang` + `ld.lld` with the AMDGPU target (`llc --version` lists `amdgcn`). If they aren't in
+  `PATH`, they are taken from `$ROCM_PATH/llvm/bin`.
+
+The first run builds `oracle.c` into `~/.cache/anyps5-hw-oracle/` (`HW_ORACLE_CACHE` overrides it). The target is read
+from the GPU (`HW_ORACLE_TARGET` overrides it).
+
+## Kernel
+
+`template.s` wraps the body under test. Per lane:
+
+- `v4`-`v7` hold the row's 4 input dwords. `v10`-`v25` start at 0 and are stored as the result (16 dwords).
+- Leave `v0` (lane id), `v1` (input offset), `v2` (output offset) and `s[4:7]` (input and output addresses) unchanged.
+- 4 KiB of LDS. Rows run in workgroups of up to 1024 lanes, padded to whole waves with zero rows.
+- Bytes passed as `extra` follow the rows of each dispatch, at `s[4:5] + 16 * rows`: buffer contents, texels, etc.
+  Build buffer and image descriptors in SGPRs from that address.
+
+The FLOAT_MODE of the dispatch is set with `ieee` (0/1), `denorm32` and `denorm16` (0-3, 3 = denormals kept). Float
+atomics need `coarse`: on fine-grained system memory they do nothing.
+
+## Use
+
+```sh
+cat > body.s <<'EOF'
+  v_add_f16 v10, v4, v5
+EOF
+printf '0x3c00 0x3c00 0 0\n0x7e01 0x3c00 0 0\n' > rows.txt
+python3 tools/hw-oracle/hw_oracle.py body.s rows.txt --outs 1 --ieee 0
+```
+
+prints one line per row with the output dwords in hex. From Python:
+
+```python
+import sys
+sys.path.insert(0, "tools/hw-oracle")
+from hw_oracle import run
+
+rows = [(a, b, 0, 0) for a in values for b in values]
+results = run(body, rows, wave64=True, denorm32=0)
+```
+
+Usual workflow: run the instruction on rows that separate the candidate models (edge values, NaN payloads, rounding
+ties), write a model that matches every row, implement it, and put measured rows in the execution test.
+
+## Limits
+
+Only compute kernels. Pixel-shader ops that read their inputs from LDS also work: `v_interp_*` reads the attribute
+parameters from LDS at `M0`, so writing them there by hand is enough. Behaviour that depends on the graphics pipeline
+(parameter cache, exports, rasterization, `SPI_*` state) can't be measured with it.

tools/hw-oracle/hw_oracle.py

@@ -0,0 +1,106 @@
+import argparse
+import os
+import shutil
+import struct
+import subprocess
+import sys
+import tempfile
+from functools import cache
+from pathlib import Path
+
+HERE = Path(__file__).resolve().parent
+CACHE = Path(os.environ.get("HW_ORACLE_CACHE") or Path(os.environ.get("XDG_CACHE_HOME") or Path.home() / ".cache") / "anyps5-hw-oracle")
+ROWS_PER_DISPATCH = 1024
+
+
+def rocm_root():
+    for root in (os.environ.get("ROCM_PATH"), "/opt/rocm"):
+        if root and (Path(root) / "include" / "hsa" / "hsa.h").exists():
+            return Path(root)
+    return None
+
+
+def tool(name):
+    found = shutil.which(name)
+    if found:
+        return found
+    root = rocm_root()
+    if root and (root / "llvm" / "bin" / name).exists():
+        return str(root / "llvm" / "bin" / name)
+    sys.exit(f"{name} not found: install LLVM with the AMDGPU target or set ROCM_PATH")
+
+
+def oracle():
+    binary = CACHE / "oracle"
+    source = HERE / "oracle.c"
+    if binary.exists() and binary.stat().st_mtime >= source.stat().st_mtime:
+        return binary
+    CACHE.mkdir(parents=True, exist_ok=True)
+    command = [os.environ.get("CC", "cc"), "-O1", str(source), "-o", str(binary)]
+    root = rocm_root()
+    if root:
+        lib = root / "lib"
+        command += [f"-I{root / 'include'}", f"-L{lib}", f"-Wl,-rpath,{lib}"]
+    subprocess.run(command + ["-lhsa-runtime64"], check=True)
+    return binary
+
+
+@cache
+def target():
+    return os.environ.get("HW_ORACLE_TARGET") or subprocess.run([oracle(), "--target"], capture_output=True, text=True, check=True).stdout.strip()
+
+
+def assemble(body, work, wave64, ieee, denorm32, denorm16):
+    text = (HERE / "template.s").read_text()
+    for key, value in (("@TARGET@", target()), ("@WAVE32@", "0" if wave64 else "1"), ("@WAVESIZE@", "64" if wave64 else "32"),
+                       ("@DENORM@", str(denorm32)), ("@DENORM16@", str(denorm16)), ("@IEEE@", str(ieee)), ("@BODY@", body)):
+        text = text.replace(key, value)
+    (work / "k.s").write_text(text)
+    subprocess.run([tool("clang"), "-x", "assembler", "-target", "amdgcn-amd-amdhsa", f"-mcpu={target()}", "-c", str(work / "k.s"), "-o", str(work / "k.o")], check=True)
+    subprocess.run([tool("ld.lld"), "-shared", str(work / "k.o"), "-o", str(work / "k.co")], check=True)
+    return work / "k.co"
+
+
+def run(body, rows, extra=b"", wave64=False, ieee=1, denorm32=3, denorm16=3, coarse=False):
+    rows = [tuple(r) for r in rows]
+    wave = 64 if wave64 else 32
+    padded = rows + [(0, 0, 0, 0)] * (-len(rows) % wave)
+    env = dict(os.environ)
+    env.pop("COARSE", None)
+    if coarse:
+        env["COARSE"] = "1"
+    out = []
+    with tempfile.TemporaryDirectory() as tmp:
+        work = Path(tmp)
+        code = assemble(body, work, wave64, ieee, denorm32, denorm16)
+        for start in range(0, len(padded), ROWS_PER_DISPATCH):
+            chunk = padded[start:start + ROWS_PER_DISPATCH]
+            (work / "in.bin").write_bytes(b"".join(struct.pack("<4I", *r) for r in chunk) + extra)
+            subprocess.run([oracle(), str(code), str(len(chunk)), str(work / "in.bin"), str(work / "out.bin"), "16"], check=True, env=env)
+            data = (work / "out.bin").read_bytes()
+            out += [struct.unpack_from("<16I", data, lane * 64) for lane in range(len(chunk))]
+    return out[:len(rows)]
+
+
+def main():
+    parser = argparse.ArgumentParser(description="Run an RDNA kernel body on the local AMD GPU, one input row per lane.")
+    parser.add_argument("body", type=Path, help="assembly inserted into template.s")
+    parser.add_argument("rows", type=Path, help="one row per line: 4 u32 values (decimal or 0x hex)")
+    parser.add_argument("--wave64", action="store_true")
+    parser.add_argument("--ieee", type=int, choices=(0, 1), default=1)
+    parser.add_argument("--denorm32", type=int, choices=range(4), default=3)
+    parser.add_argument("--denorm16", type=int, choices=range(4), default=3)
+    parser.add_argument("--coarse", action="store_true", help="input and output in coarse-grained GPU memory")
+    parser.add_argument("--extra", type=Path, help="bytes appended after the rows")
+    parser.add_argument("--outs", type=int, choices=range(1, 17), default=16, metavar="N", help="print v10..v(10+N-1)")
+    args = parser.parse_args()
+    rows = [tuple(int(v, 0) for v in line.split()) for line in args.rows.read_text().splitlines() if line.strip()]
+    if any(len(r) != 4 for r in rows):
+        sys.exit("every row needs 4 values")
+    extra = args.extra.read_bytes() if args.extra else b""
+    for result in run(args.body.read_text(), rows, extra, args.wave64, args.ieee, args.denorm32, args.denorm16, args.coarse):
+        print(" ".join(f"{v:08x}" for v in result[:args.outs]))
+
+
+if __name__ == "__main__":
+    main()

tools/hw-oracle/oracle.c

@@ -0,0 +1,165 @@
+#include <hsa/hsa.h>
+#include <hsa/hsa_ext_amd.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
+
+#define CHECK(x) do { hsa_status_t s_ = (x); if (s_ != HSA_STATUS_SUCCESS && s_ != HSA_STATUS_INFO_BREAK) { const char* m_; hsa_status_string(s_, &m_); fprintf(stderr, "%s: %s\n", #x, m_); exit(1); } } while (0)
+
+static hsa_agent_t gpu;
+static int haveGpu;
+static hsa_amd_memory_pool_t kernargPool, finePool;
+static int haveKernarg, haveFine;
+static hsa_amd_memory_pool_t coarsePool;
+static int haveCoarse;
+
+static hsa_status_t findCoarse(hsa_amd_memory_pool_t pool, void* data) {
+    hsa_amd_segment_t segment;
+    hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_SEGMENT, &segment);
+    if (segment != HSA_AMD_SEGMENT_GLOBAL) return HSA_STATUS_SUCCESS;
+    uint32_t flags;
+    hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS, &flags);
+    if ((flags & HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_COARSE_GRAINED) && !haveCoarse) { coarsePool = pool; haveCoarse = 1; }
+    return HSA_STATUS_SUCCESS;
+}
+
+static hsa_status_t findGpu(hsa_agent_t agent, void* data) {
+    hsa_device_type_t type;
+    hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &type);
+    if (type == HSA_DEVICE_TYPE_GPU) { gpu = agent; haveGpu = 1; return HSA_STATUS_INFO_BREAK; }
+    return HSA_STATUS_SUCCESS;
+}
+
+static hsa_status_t findPools(hsa_amd_memory_pool_t pool, void* data) {
+    hsa_amd_segment_t segment;
+    hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_SEGMENT, &segment);
+    if (segment != HSA_AMD_SEGMENT_GLOBAL) return HSA_STATUS_SUCCESS;
+    uint32_t flags;
+    hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS, &flags);
+    if ((flags & HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_KERNARG_INIT) && !haveKernarg) { kernargPool = pool; haveKernarg = 1; }
+    if ((flags & HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_FINE_GRAINED) && !haveFine) { finePool = pool; haveFine = 1; }
+    return HSA_STATUS_SUCCESS;
+}
+
+static hsa_status_t findCpu(hsa_agent_t agent, void* data) {
+    hsa_device_type_t type;
+    hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &type);
+    if (type == HSA_DEVICE_TYPE_CPU) { hsa_amd_agent_iterate_memory_pools(agent, findPools, NULL); }
+    return HSA_STATUS_SUCCESS;
+}
+
+static void* readFile(const char* path, size_t* size) {
+    FILE* f = fopen(path, "rb");
+    if (!f) { perror(path); exit(1); }
+    fseek(f, 0, SEEK_END); *size = (size_t)ftell(f); fseek(f, 0, SEEK_SET);
+    void* data = malloc(*size);
+    if (fread(data, 1, *size, f) != *size) { perror(path); exit(1); }
+    fclose(f);
+    return data;
+}
+
+int main(int argc, char** argv) {
+    if (argc == 2 && strcmp(argv[1], "--target") == 0) {
+        CHECK(hsa_init());
+        CHECK(hsa_iterate_agents(findGpu, NULL));
+        if (!haveGpu) { fprintf(stderr, "no GPU agent\n"); return 1; }
+        char name[64] = {0};
+        CHECK(hsa_agent_get_info(gpu, HSA_AGENT_INFO_NAME, name));
+        printf("%s\n", name);
+        return 0;
+    }
+    if (argc != 6) { fprintf(stderr, "usage: oracle --target | oracle code.o threads in.bin out.bin out_dwords\n"); return 2; }
+    const uint32_t threads = (uint32_t)atoi(argv[2]);
+    const uint32_t outDwords = (uint32_t)atoi(argv[5]);
+    CHECK(hsa_init());
+    CHECK(hsa_iterate_agents(findGpu, NULL));
+    if (!haveGpu) { fprintf(stderr, "no GPU agent\n"); return 1; }
+    CHECK(hsa_iterate_agents(findCpu, NULL));
+    if (!haveKernarg || !haveFine) { fprintf(stderr, "no system memory pools\n"); return 1; }
+
+    size_t codeSize; void* code = readFile(argv[1], &codeSize);
+    hsa_code_object_reader_t reader;
+    CHECK(hsa_code_object_reader_create_from_memory(code, codeSize, &reader));
+    hsa_executable_t exe;
+    CHECK(hsa_executable_create_alt(HSA_PROFILE_FULL, HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT, NULL, &exe));
+    CHECK(hsa_executable_load_agent_code_object(exe, gpu, reader, NULL, NULL));
+    CHECK(hsa_executable_freeze(exe, NULL));
+    hsa_executable_symbol_t sym;
+    CHECK(hsa_executable_get_symbol_by_name(exe, "probe.kd", &gpu, &sym));
+    uint64_t kernelObject; uint32_t kernargSize, groupSize, privateSize;
+    hsa_executable_symbol_get_info(sym, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT, &kernelObject);
+    hsa_executable_symbol_get_info(sym, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_SIZE, &kernargSize);
+    hsa_executable_symbol_get_info(sym, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_GROUP_SEGMENT_SIZE, &groupSize);
+    hsa_executable_symbol_get_info(sym, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_PRIVATE_SEGMENT_SIZE, &privateSize);
+
+    size_t inSize; void* inHost = readFile(argv[3], &inSize);
+    void *in, *out, *args;
+    const size_t outSize = (size_t)threads * outDwords * 4u;
+    const int coarse = getenv("COARSE") != NULL;
+    if (coarse) {
+        CHECK(hsa_amd_agent_iterate_memory_pools(gpu, findCoarse, NULL));
+        if (!haveCoarse) { fprintf(stderr, "no coarse-grained GPU pool\n"); return 1; }
+    }
+    void *inStage = NULL, *outStage = NULL;
+    CHECK(hsa_amd_memory_pool_allocate(coarse ? coarsePool : finePool, inSize, 0, &in));
+    CHECK(hsa_amd_memory_pool_allocate(coarse ? coarsePool : finePool, outSize, 0, &out));
+    if (coarse) {
+        CHECK(hsa_amd_memory_pool_allocate(finePool, inSize, 0, &inStage));
+        CHECK(hsa_amd_memory_pool_allocate(finePool, outSize, 0, &outStage));
+        CHECK(hsa_amd_agents_allow_access(1, &gpu, NULL, inStage));
+        CHECK(hsa_amd_agents_allow_access(1, &gpu, NULL, outStage));
+    }
+    CHECK(hsa_amd_memory_pool_allocate(kernargPool, 16, 0, &args));
+    if (!coarse) {
+        CHECK(hsa_amd_agents_allow_access(1, &gpu, NULL, in));
+        CHECK(hsa_amd_agents_allow_access(1, &gpu, NULL, out));
+    }
+    CHECK(hsa_amd_agents_allow_access(1, &gpu, NULL, args));
+    if (coarse) {
+        memcpy(inStage, inHost, inSize);
+        memset(outStage, 0xee, outSize);
+        CHECK(hsa_memory_copy(in, inStage, inSize));
+        CHECK(hsa_memory_copy(out, outStage, outSize));
+    } else {
+        memcpy(in, inHost, inSize);
+        memset(out, 0xee, outSize);
+    }
+    ((void**)args)[0] = in;
+    ((void**)args)[1] = out;
+
+    hsa_queue_t* queue;
+    CHECK(hsa_queue_create(gpu, 64, HSA_QUEUE_TYPE_SINGLE, NULL, NULL, UINT32_MAX, UINT32_MAX, &queue));
+    hsa_signal_t done;
+    CHECK(hsa_signal_create(1, 0, NULL, &done));
+    const uint64_t index = hsa_queue_add_write_index_relaxed(queue, 1);
+    hsa_kernel_dispatch_packet_t* packet = (hsa_kernel_dispatch_packet_t*)queue->base_address + (index & (queue->size - 1));
+    memset((uint8_t*)packet + 4, 0, sizeof(*packet) - 4);
+    packet->setup = 1u << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS;
+    packet->workgroup_size_x = (uint16_t)(threads < 1024 ? threads : 1024);
+    packet->workgroup_size_y = 1;
+    packet->workgroup_size_z = 1;
+    packet->grid_size_x = threads;
+    packet->grid_size_y = 1;
+    packet->grid_size_z = 1;
+    packet->kernel_object = kernelObject;
+    packet->kernarg_address = args;
+    packet->private_segment_size = privateSize;
+    packet->group_segment_size = groupSize;
+    packet->completion_signal = done;
+    const uint16_t header = (HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE) |
+        (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE) | (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE);
+    __atomic_store_n((uint32_t*)packet, header | ((uint32_t)packet->setup << 16), __ATOMIC_RELEASE);
+    hsa_signal_store_screlease(queue->doorbell_signal, (hsa_signal_value_t)index);
+    if (hsa_signal_wait_scacquire(done, HSA_SIGNAL_CONDITION_LT, 1, 5000000000ull, HSA_WAIT_STATE_BLOCKED) != 0) {
+        fprintf(stderr, "kernel timed out\n");
+        return 1;
+    }
+    if (coarse) {
+        CHECK(hsa_memory_copy(outStage, out, outSize));
+        out = outStage;
+    }
+    FILE* f = fopen(argv[4], "wb");
+    if (!f || fwrite(out, 1, outSize, f) != outSize) { perror(argv[4]); return 1; }
+    fclose(f);
+    return 0;
+}

tools/hw-oracle/template.s

@@ -0,0 +1,70 @@
+.amdgcn_target "amdgcn-amd-amdhsa--@TARGET@"
+.text
+.globl probe
+.p2align 8
+.type probe,@function
+probe:
+  s_load_dwordx4 s[4:7], s[0:1], 0x0
+  v_lshlrev_b32 v1, 4, v0
+  v_lshlrev_b32 v2, 6, v0
+  s_waitcnt lgkmcnt(0)
+  global_load_dwordx4 v[4:7], v1, s[4:5]
+  s_waitcnt vmcnt(0)
+  v_mov_b32 v10, 0
+  v_mov_b32 v11, 0
+  v_mov_b32 v12, 0
+  v_mov_b32 v13, 0
+  v_mov_b32 v14, 0
+  v_mov_b32 v15, 0
+  v_mov_b32 v16, 0
+  v_mov_b32 v17, 0
+  v_mov_b32 v18, 0
+  v_mov_b32 v19, 0
+  v_mov_b32 v20, 0
+  v_mov_b32 v21, 0
+  v_mov_b32 v22, 0
+  v_mov_b32 v23, 0
+  v_mov_b32 v24, 0
+  v_mov_b32 v25, 0
+@BODY@
+  global_store_dwordx4 v2, v[10:13], s[6:7]
+  global_store_dwordx4 v2, v[14:17], s[6:7] offset:16
+  global_store_dwordx4 v2, v[18:21], s[6:7] offset:32
+  global_store_dwordx4 v2, v[22:25], s[6:7] offset:48
+  s_endpgm
+.Lend:
+.size probe, .Lend-probe
+
+.rodata
+.p2align 6
+.amdhsa_kernel probe
+  .amdhsa_user_sgpr_kernarg_segment_ptr 1
+  .amdhsa_kernarg_size 16
+  .amdhsa_group_segment_fixed_size 4096
+  .amdhsa_next_free_vgpr 128
+  .amdhsa_next_free_sgpr 48
+  .amdhsa_wavefront_size32 @WAVE32@
+  .amdhsa_float_denorm_mode_32 @DENORM@
+  .amdhsa_float_denorm_mode_16_64 @DENORM16@
+  .amdhsa_ieee_mode @IEEE@
+.end_amdhsa_kernel
+
+.amdgpu_metadata
+---
+amdhsa.version: [ 1, 2 ]
+amdhsa.kernels:
+  - .name: probe
+    .symbol: probe.kd
+    .kernarg_segment_size: 16
+    .group_segment_fixed_size: 4096
+    .private_segment_fixed_size: 0
+    .kernarg_segment_align: 8
+    .wavefront_size: @WAVESIZE@
+    .sgpr_count: 48
+    .vgpr_count: 64
+    .max_flat_workgroup_size: 1024
+    .args:
+      - { .size: 8, .offset: 0, .value_kind: global_buffer, .address_space: global }
+      - { .size: 8, .offset: 8, .value_kind: global_buffer, .address_space: global }
+...
+.end_amdgpu_metadata