From f617b40a458f817b98ade39f0294dc6b633f4042 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Adri=C3=A0=20Franch?= Date: Tue, 6 Oct 2026 13:54:28 +0200 Subject: [PATCH] feat(tools): add a hardware oracle for RDNA instruction semantics Co-Authored-By: Claude --- docs/dev/HW_ORACLE.md | 64 ++++++++++++++ tools/hw-oracle/hw_oracle.py | 106 ++++++++++++++++++++++ tools/hw-oracle/oracle.c | 165 +++++++++++++++++++++++++++++++++++ tools/hw-oracle/template.s | 70 +++++++++++++++ 4 files changed, 405 insertions(+) create mode 100644 docs/dev/HW_ORACLE.md create mode 100644 tools/hw-oracle/hw_oracle.py create mode 100644 tools/hw-oracle/oracle.c create mode 100644 tools/hw-oracle/template.s diff --git a/docs/dev/HW_ORACLE.md b/docs/dev/HW_ORACLE.md new file mode 100644 index 000000000..917efdc65 --- /dev/null +++ b/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. diff --git a/tools/hw-oracle/hw_oracle.py b/tools/hw-oracle/hw_oracle.py new file mode 100644 index 000000000..eee171f50 --- /dev/null +++ b/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() diff --git a/tools/hw-oracle/oracle.c b/tools/hw-oracle/oracle.c new file mode 100644 index 000000000..cb240017e --- /dev/null +++ b/tools/hw-oracle/oracle.c @@ -0,0 +1,165 @@ +#include +#include +#include +#include +#include + +#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; +} diff --git a/tools/hw-oracle/template.s b/tools/hw-oracle/template.s new file mode 100644 index 000000000..ad5426247 --- /dev/null +++ b/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