Repository navigation
fix(shader): round the v_mad_legacy_f32 product before the add - #1281
Merged
adriadam10 merged 1 commit intoOct 10, 2026
Merged
Conversation
9 tasks done
Contributor
|
I had Hermes go over the technical debt and found a divergence from my card and this PR: Measured VOP3 0x140 on RDNA2 desktop hardware (RX 6700 XT, gfx1031) with the HSA oracle from tools/hw-oracle/template.s, in all four IEEE/denorm MODE combinations: it is fused there. Your fix is still correct for the console family this emulator targets — filing the divergence so nobody "fixes" it back based on a desktop measurement.
Plus six CPU-fmaf-vs-mul-add discriminator rows: on gfx1031 the legacy result equals v_fma_f32 bit-exactly in every mode combination, and differs from mul-then-add by 1 ULP. The same kernel also confirms your zero rule and denormal flush carry over to RDNA2: +tiny x 1 + (-0) and -tiny x 1 + (-0) both give +0 for the legacy instruction while v_fma_f32 keeps the tiny — only the fusion differs between generations. Corroborating LLVM evidence: for gfx1031, clang assembles this encoding only as v_fma_legacy_f32 and rejects v_mad_legacy_f32 with "instruction not supported on this GPU", and llvm-objdump disassembles D540 000A 041A 0B04 as v_fma_legacy_f32. The rename looks like a real behavior change at RDNA2, so the TechnicalDebt entry could note that RDNA2 desktop fuses (measured on RX 6700 XT) — otherwise a future contributor with RDNA2 hardware measures fused and flips the translation back. |
Contributor
Author
|
Thanks for measuring it on gfx1031, that's useful. It matches the LLVM split ( I added a commit that records it in the TechnicalDebt entry: gfx1013 is decoded as |
9 tasks done
adriadam10
requested changes
Oct 9, 2026
adriadam10
left a comment
adriadam10
left a comment
Collaborator
There was a problem hiding this comment.
Thanks, the gfx1010 measurement and @pandather's gfx1031 cross-check make the case for a rounded product on gfx10.1 (it also matches LLVM decoding 0x140 as v_mad_legacy_f32 up to gfx10.1 and only selecting mad when f32 denormals are flushed). It now conflicts with main in a way that isn't a mechanical rebase, though:
FloatInstructions.cppvFmaLegacyF32: since 1910d14mainpicks the NaN result on the bits (nanResultF32({lhsFactor, rhsFactor, addend}, result, &invalidProductF32(lhsFactor, rhsFactor))). The rebased version should keep that wrapper around the new mul + add. Otherwise the NaN payload is left to the driver, and rows 49/50 ofFmaLegacy.cpp(0xffc00000/0x7fc00000) will fail on NVIDIA, which returns0x7fffffff.- Denormals: since #1850,
mainflushes f32 denormals according to the shader'sFLOAT_MODE(flushF32Denormal,OperandAccess.cpp). Your measurement says the mad form flushes in every mode, which is fine, but please keep it as a deliberate exception to that model rather than a local lambda that silently overrides it. Then say so in the TechnicalDebt entry. docs/dev/TechnicalDebt.md: besides the 0x140 entry, two lines onmainbecome false. The NaN line listsfma_legacy_f32among the fused forms and the "separately rounded mad/mac forms" as left to the driver; the f32-denormal line lists the mad forms as not covered. Please update both.- The new entry says "measured with the hardware oracle", but the description says the kernel was launched through
amd_comgr/amdhip64on Windows. Please describe it that way, and attach the probe and raw output (30 rows × 4 modes) to the PR. We can't reproduce gfx1010 here (our oracle runs on gfx1035, which fuses 0x140), so your measurement is the only evidence for these rows.
#1290 is stacked on this one and will need the same rebase.
SirDarky
force-pushed
the
fix/shader-mad-legacy-rounded-product
branch
from
October 10, 2026 04:03
8f34e9f to
8566772
Compare
Contributor
Author
|
Thanks for the review. Rebased onto
Raw output: 32 rows × 5 encodings, wave32 and wave64, IEEE 0 and 1, f32 denormal modes 0 to 3, on an RX 5700 XT (gfx1010), Adrenalin driver on Windows 11. All 16 combinations give the same results, except rows 30 and 31 (signaling NaNs), which are quieted only with IEEE=1. The legacy-mad-gfx1010-raw.txtprobe_legacy_mad.py (the probe)import json
import sys
from pathlib import Path
sys.path.insert(0, str(Path(__file__).resolve().parent))
import hw_oracle_win as oracle
FORMS = {
"vop3_0x140": (
"v_mad_legacy_f32 (VOP3 0x140): plain, -src0, |src0| -src2, src0=0, -src0 -|src1| (the FmaLegacy.cpp words)",
[0xd540000a, 0x041a0b04, 0xd540000b, 0x241a0b04, 0xd540010c, 0x841a0b04, 0xd540000d, 0x041a0a80, 0xd540020e, 0x641a0b04],
5,
),
"vop2_0x06": (
"v_mac_legacy_f32 (VOP2 0x06 / VOP3 0x106): VOP2, VOP3 -src0, VOP3 |src0| -|src1|, VOP2 src0=0 (the MacLegacy.cpp words)",
[0x7e140306, 0x0c140b04, 0x7e160306, 0xd506000b, 0x20020b04, 0x7e180306, 0xd506030c, 0x40020b04, 0x7e1a0306, 0x0c1a0a80],
4,
),
}
ROWS = [
(0x00000000, 0x7f800000, 0x3f800000), (0x80000000, 0x7f800000, 0xc0200000), (0x7f800000, 0x00000000, 0x3f400000),
(0xff800000, 0x80000000, 0x40400000), (0x00000000, 0x7fc00000, 0x3fc00000), (0x7fc00000, 0x80000000, 0xc0800000),
(0x00000000, 0x40000000, 0x7f800000), (0x80000000, 0xc0400000, 0xff800000), (0x00000000, 0x3f800000, 0x7fc00000),
(0x3fc00000, 0x40100000, 0xbf400000), (0xc0c00000, 0x3f000000, 0x41200000), (0x7f800000, 0x40000000, 0x3f800000),
(0x7f800000, 0xbf800000, 0x7f800000), (0x7fc00000, 0x3f800000, 0x3f800000), (0x40400000, 0x40800000, 0xc1400000),
(0x7f7fffff, 0x40000000, 0x3f800000), (0x3ec00000, 0xc1480000, 0x42c88000), (0xc4800000, 0xbd800000, 0xc2800000),
(0x40e00000, 0x41100000, 0x3f000000), (0x80000000, 0x80000000, 0x40a00000),
(0x3f800800, 0x3f800800, 0xbf800000), (0x3f800001, 0x3f7fffff, 0xbf800000), (0x3faaaaab, 0x40400000, 0xc0800000),
(0x7f7fffff, 0x40000000, 0xff7fffff), (0x00800000, 0x3f000000, 0x00000000), (0x00000001, 0x4b000000, 0x00000000),
(0x00000001, 0x3f800000, 0x80000000), (0x80000001, 0x3f800000, 0x00000000), (0x00800000, 0x3f800000, 0x80400000),
(0x3f800000, 0x00400000, 0x00000000), (0x7f812345, 0x3f800000, 0x3f800000), (0x3f800000, 0x3f800000, 0x7f812345),
]
MODES = [(ieee, denorm32) for ieee in (0, 1) for denorm32 in (0, 1, 2, 3)]
def main():
target = oracle.target()
report = {"target": target, "forms": {}}
lines = [f"target {target}; tools/hw-oracle/template.s assembled by amd_comgr_2.dll and launched by amdhip64_6.dll; denorm16=3"]
for key, (label, words, outputs) in FORMS.items():
body = "\n".join(f" .long 0x{w:08x}" for w in words)
results = {}
for wave64 in (False, True):
for ieee, denorm32 in MODES:
rows = oracle.run(body, [(a, b, c, 0) for a, b, c in ROWS], wave64=wave64, ieee=ieee, denorm32=denorm32, denorm16=3)
results[f"wave{64 if wave64 else 32} ieee={ieee} denorm32={denorm32}"] = [list(r[:outputs]) for r in rows]
reference = results["wave32 ieee=1 denorm32=3"]
lines.append("")
lines.append(f"== {label}")
for mode, rows in results.items():
differing = [i for i, (x, y) in enumerate(zip(rows, reference)) if x != y]
lines.append(f"{mode}: rows differing from wave32 ieee=1 denorm32=3: {differing}")
lines.append("src0 src1 src2 -> results (wave32 ieee=0 denorm32=0, the mode of a shader without a known float mode)")
for (a, b, c), r in zip(ROWS, results["wave32 ieee=0 denorm32=0"]):
lines.append(f"{a:08x} {b:08x} {c:08x} -> " + " ".join(f"{v:08x}" for v in r))
report["forms"][key] = {"label": label, "words": [f"0x{w:08x}" for w in words], "rows": [[f"0x{v:08x}" for v in row] for row in ROWS], "results": {m: [[f"0x{v:08x}" for v in r] for r in rows] for m, rows in results.items()}}
Path(__file__).with_name("legacy-mad-gfx1010-raw.txt").write_text("\n".join(lines) + "\n")
Path(__file__).with_name("legacy-mad-gfx1010-raw.json").write_text(json.dumps(report, indent=1))
print("\n".join(lines))
if __name__ == "__main__":
main()hw_oracle_win.py (the runner: template.s through amd_comgr + HIP)import argparse
import ctypes
import os
import struct
import sys
from functools import cache
from pathlib import Path
ROWS_PER_DISPATCH = 1024
TEMPLATE = Path(os.environ.get("HW_ORACLE_TEMPLATE", Path(__file__).resolve().parent / "template.s"))
DATA_SOURCE = 0x1
DATA_LOG = 0x5
DATA_RELOCATABLE = 0x7
DATA_EXECUTABLE = 0x8
ASSEMBLE_ACTIONS = (0x8, 0xA)
LINK_ACTIONS = (0x7, 0x9)
HIP_MEMCPY_HOST_TO_DEVICE = 1
HIP_MEMCPY_DEVICE_TO_HOST = 2
class Handle(ctypes.Structure):
_fields_ = [("handle", ctypes.c_uint64)]
def check(status, what):
if status != 0:
raise RuntimeError(f"{what} failed with status {status}")
@cache
def hip():
library = ctypes.WinDLL(os.environ.get("HW_ORACLE_HIP", "amdhip64_6.dll"))
check(library.hipInit(0), "hipInit")
return library
@cache
def comgr():
library = ctypes.WinDLL(os.environ.get("HW_ORACLE_COMGR", "amd_comgr_2.dll"))
library.amd_comgr_create_data.argtypes = [ctypes.c_int, ctypes.POINTER(Handle)]
library.amd_comgr_create_data_set.argtypes = [ctypes.POINTER(Handle)]
library.amd_comgr_create_action_info.argtypes = [ctypes.POINTER(Handle)]
library.amd_comgr_set_data.argtypes = [Handle, ctypes.c_size_t, ctypes.c_char_p]
library.amd_comgr_set_data_name.argtypes = [Handle, ctypes.c_char_p]
library.amd_comgr_data_set_add.argtypes = [Handle, Handle]
library.amd_comgr_action_info_set_isa_name.argtypes = [Handle, ctypes.c_char_p]
library.amd_comgr_action_info_set_logging.argtypes = [Handle, ctypes.c_bool]
library.amd_comgr_do_action.argtypes = [ctypes.c_int, Handle, Handle, Handle]
library.amd_comgr_action_data_count.argtypes = [Handle, ctypes.c_int, ctypes.POINTER(ctypes.c_size_t)]
library.amd_comgr_action_data_get_data.argtypes = [Handle, ctypes.c_int, ctypes.c_size_t, ctypes.POINTER(Handle)]
library.amd_comgr_get_data.argtypes = [Handle, ctypes.POINTER(ctypes.c_size_t), ctypes.c_char_p]
library.amd_comgr_release_data.argtypes = [Handle]
library.amd_comgr_destroy_data_set.argtypes = [Handle]
library.amd_comgr_destroy_action_info.argtypes = [Handle]
return library
@cache
def target():
if os.environ.get("HW_ORACLE_TARGET"):
return os.environ["HW_ORACLE_TARGET"]
props = ctypes.create_string_buffer(8192)
check(hip().hipGetDevicePropertiesR0600(props, 0), "hipGetDeviceProperties")
start = props.raw.find(b"gfx")
return props.raw[start:props.raw.index(b"\0", start)].decode().split(":")[0]
def data_bytes(data_set, kind):
library = comgr()
count = ctypes.c_size_t()
check(library.amd_comgr_action_data_count(data_set, kind, ctypes.byref(count)), "amd_comgr_action_data_count")
chunks = []
for index in range(count.value):
data = Handle()
check(library.amd_comgr_action_data_get_data(data_set, kind, index, ctypes.byref(data)), "amd_comgr_action_data_get_data")
size = ctypes.c_size_t()
check(library.amd_comgr_get_data(data, ctypes.byref(size), None), "amd_comgr_get_data")
buffer = ctypes.create_string_buffer(size.value)
check(library.amd_comgr_get_data(data, ctypes.byref(size), buffer), "amd_comgr_get_data")
chunks.append(buffer.raw[:size.value])
library.amd_comgr_release_data(data)
return chunks
def single_input(kind, name, payload):
library = comgr()
data = Handle()
check(library.amd_comgr_create_data(kind, ctypes.byref(data)), "amd_comgr_create_data")
check(library.amd_comgr_set_data(data, len(payload), payload), "amd_comgr_set_data")
check(library.amd_comgr_set_data_name(data, name), "amd_comgr_set_data_name")
data_set = Handle()
check(library.amd_comgr_create_data_set(ctypes.byref(data_set)), "amd_comgr_create_data_set")
check(library.amd_comgr_data_set_add(data_set, data), "amd_comgr_data_set_add")
library.amd_comgr_release_data(data)
return data_set
def comgr_step(candidates, inputs, kind):
library = comgr()
logs = []
for action in candidates:
info = Handle()
check(library.amd_comgr_create_action_info(ctypes.byref(info)), "amd_comgr_create_action_info")
check(library.amd_comgr_action_info_set_isa_name(info, f"amdgcn-amd-amdhsa--{target()}".encode()), "amd_comgr_action_info_set_isa_name")
check(library.amd_comgr_action_info_set_logging(info, True), "amd_comgr_action_info_set_logging")
outputs = Handle()
check(library.amd_comgr_create_data_set(ctypes.byref(outputs)), "amd_comgr_create_data_set")
status = library.amd_comgr_do_action(action, info, inputs, outputs)
library.amd_comgr_destroy_action_info(info)
result = data_bytes(outputs, kind)
logs.append(f"action 0x{action:x} status {status}: " + b"".join(data_bytes(outputs, DATA_LOG)).decode(errors="replace").strip())
library.amd_comgr_destroy_data_set(outputs)
if status == 0 and len(result) == 1:
return result[0]
raise RuntimeError("comgr produced no output\n" + "\n".join(logs))
def assemble(body, wave64, ieee, denorm32, denorm16):
text = TEMPLATE.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)
source = single_input(DATA_SOURCE, b"k.s", text.encode())
relocatable = comgr_step(ASSEMBLE_ACTIONS, source, DATA_RELOCATABLE)
comgr().amd_comgr_destroy_data_set(source)
objects = single_input(DATA_RELOCATABLE, b"k.o", relocatable)
executable = comgr_step(LINK_ACTIONS, objects, DATA_EXECUTABLE)
comgr().amd_comgr_destroy_data_set(objects)
return executable
def dispatch(image, payload, threads, coarse):
library = hip()
module = ctypes.c_void_p()
check(library.hipModuleLoadData(ctypes.byref(module), ctypes.c_char_p(image)), "hipModuleLoadData")
function = ctypes.c_void_p()
check(library.hipModuleGetFunction(ctypes.byref(function), module, b"probe"), "hipModuleGetFunction")
out_size = threads * 64
allocate, release = (library.hipMalloc, library.hipFree) if coarse else (library.hipHostMalloc, library.hipHostFree)
device_in, device_out = ctypes.c_void_p(), ctypes.c_void_p()
if coarse:
check(allocate(ctypes.byref(device_in), ctypes.c_size_t(len(payload))), "hipMalloc")
check(allocate(ctypes.byref(device_out), ctypes.c_size_t(out_size)), "hipMalloc")
else:
check(allocate(ctypes.byref(device_in), ctypes.c_size_t(len(payload)), 0), "hipHostMalloc")
check(allocate(ctypes.byref(device_out), ctypes.c_size_t(out_size), 0), "hipHostMalloc")
check(library.hipMemcpy(device_in, ctypes.c_char_p(payload), ctypes.c_size_t(len(payload)), HIP_MEMCPY_HOST_TO_DEVICE), "hipMemcpy")
check(library.hipMemset(device_out, 0xee, ctypes.c_size_t(out_size)), "hipMemset")
arguments = (ctypes.c_void_p * 2)(ctypes.addressof(device_in), ctypes.addressof(device_out))
check(library.hipModuleLaunchKernel(function, 1, 1, 1, ctypes.c_uint(threads), 1, 1, 0, None, arguments, None), "hipModuleLaunchKernel")
check(library.hipDeviceSynchronize(), "hipDeviceSynchronize")
output = ctypes.create_string_buffer(out_size)
check(library.hipMemcpy(output, device_out, ctypes.c_size_t(out_size), HIP_MEMCPY_DEVICE_TO_HOST), "hipMemcpy")
check(release(device_in), "free")
check(release(device_out), "free")
check(library.hipModuleUnload(module), "hipModuleUnload")
return output.raw
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)
image = assemble(body, wave64, ieee, denorm32, denorm16)
out = []
for start in range(0, len(padded), ROWS_PER_DISPATCH):
chunk = padded[start:start + ROWS_PER_DISPATCH]
payload = b"".join(struct.pack("<4I", *r) for r in chunk) + extra
data = dispatch(image, payload, len(chunk), coarse)
out += [struct.unpack_from("<16I", data, lane * 64) for lane in range(len(chunk))]
return out[:len(rows)]
def main():
parser = argparse.ArgumentParser(description="Windows (HIP + comgr from the AMD driver) runner for tools/hw-oracle/template.s.")
parser.add_argument("body", type=Path)
parser.add_argument("rows", type=Path)
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")
parser.add_argument("--extra", type=Path)
parser.add_argument("--outs", type=int, choices=range(1, 17), default=16, metavar="N")
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()#1290 will get the same rebase once this one is settled. |
This was referenced Oct 10, 2026
VOP3 0x140 is v_mad_legacy_f32 on gfx10.1. Measured on gfx1010 (RX 5700 XT) in wave32 and wave64, IEEE 0 and 1 and f32 denormal modes 0 to 3, it rounds the product before the add, flushes f32 denormal sources, product and result in every mode, and gives a +0 product for a zero factor. The translation multiplies, applies the zero rule and flushTinyProduct, adds, and keeps the nanResultF32 wrapper. f32DenormalFlushFor returns 3 for the instruction before reading FLOAT_MODE. A sum of two zeros takes its sign on the bits, as lavapipe returned -0 for +0 + -0. The execution test pins the measured results and compares NaN bits exactly. RDNA2 (gfx1031), where the encoding is v_fma_legacy_f32 and fused, is noted in TechnicalDebt. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
adriadam10
force-pushed
the
fix/shader-mad-legacy-rounded-product
branch
from
October 10, 2026 06:26
8566772 to
5377db4
Compare
adriadam10
approved these changes
Oct 10, 2026
adriadam10
left a comment
adriadam10
left a comment
Collaborator
There was a problem hiding this comment.
main moved after the review (#2279 changed nanResultF32 for three sources), so the PR conflicted in FloatInstructions.cpp and TechnicalDebt.md. I rebased it and pushed the result:
vFmaLegacyF32: dropped theinvalidProductF32argument, which no longer exists; with the zero rule both factors are 0 whenever one is, so the product is never invalid and the NaN choice is the same.TechnicalDebt.md: thev_fma_legacy_f32entry onmainis replaced by thev_mad_legacy_f32one, and the two sentences are re-applied to the current NaN and f32 denormal entries.
Build + scoped ctest green, agc_driver_fma_legacy, vector_mad_legacy and vector_fma_nan_result pass (the first and the last also on CI's lavapipe). Merging once CI is green; #1290 can rebase on top.
SleglJan
pushed a commit
to SleglJan/AnyPS5
that referenced
this pull request
Oct 10, 2026
…y-rounded-product fix(shader): round the v_mad_legacy_f32 product before the add
9 tasks done
Merged
11 tasks done
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
3 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
What
VOP3 0x140 is
v_mad_legacy_f32on gfx10.1, the family of the PS5's gfx1013 (LLVM decodes it that way up to gfx10.1). It was translated as a fused multiply-add, an assumption recorded in TechnicalDebt as unmeasured. Measured on gfx1010, it is not fused:(1+2^-12)^2 - 1gives0x3a000000(fused:0x3a000400),max * 2 - maxgives+inf(fused:max);0 * inf + 1 = 1,-0 * 5 + -0 = +0);mainuses for the f32 arithmetic: a NaN source quieted only with IEEE=1,inf - infgiving0xffc00000, the zero rule hiding a NaN factor, and a NaN src2 winning.The translation now multiplies, applies the zero rule and
flushTinyProduct, adds, and keeps thenanResultF32wrapper withinvalidProductF32. A sum of two zeros takes its sign on the bits (+0 unless both are -0), because lavapipe returned -0 for +0 + -0. The flushing is an explicit exception inf32DenormalFlushFor: it returns 3 forVMadLegacyF32before readingFLOAT_MODE, soreadOperand,flushTinyProductandwriteOperandflush in every mode.FmaLegacy.cpppins the measured results of its five encodings (plain,-src0,|src0| -src2,src0=0,-src0 -|src1|). It covers the original 20 rows, 10 rows that separate fused from rounded and keep from flush, and 2 signaling NaN rows. NaN bits are compared exactly. TechnicalDebt: the 0x140 entry is replaced by the measurement, and the NaN and f32-denormal lines are updated. The entry also records @pandather's gfx1031 measurement, where the encoding isv_fma_legacy_f32and fused, so an RDNA2 measurement is not taken as the console's behaviour.Source of the semantics: measured on hardware, an AMD Radeon RX 5700 XT (gfx1010), with
tools/hw-oracle/template.sunchanged. It was assembled by the AMD Windows driver'samd_comgr_2.dlland launched through itsamdhip64_6.dllinstead oforacle.cand ROCr. 32 rows × 5 encodings, wave32 and wave64, IEEE 0 and 1, f32 denormal modes 0 to 3: all 16 combinations agree, except the signaling NaN rows, which are quieted only with IEEE=1. The probe, the runner and the raw output are attached in a comment.Tested
Windows 11, MinGW-w64 GCC 15.2.0 (WinLibs r7), AMD Radeon RX 5700 XT, CI configure flags, full
ctest --output-on-failure --timeout 120onmainc963432 plus this change: 570 of 573 pass.agc_driver_pixel_interlock(wave64 BDA fault: pixel 0 counted 1 of 0 overlapping read-modify-writes),agc_driver_sdwa_float_selectorsandagc_driver_sdwa_ldexp_f16fail with the same messages on unmodifiedmainon this GPU and are unrelated.agc_driver_fma_legacyandvector_mad_legacypass. Withmain's fused translation,agc_driver_fma_legacyfails (lane 20 (0x3f800800, 0x3f800800, 0xbf800000) v_mad_legacy_f32 is 0x3a000400, expected 0x3a000000). Without thef32DenormalFlushForexception, it fails at lane 26 (-src0 is 0x80000000, expected 0x00000000).Checklist
main; no other open PR implements the same functionsNotImplemented_nid_no_patch); silent stubs are listed in TechnicalDebt.mdfiles or images added to the repository (attach them to this PR)🤖 Generated with Claude Code