diff --git a/docs/windows-performance.md b/docs/windows-performance.md index acdf8ad..df23a5f 100644 --- a/docs/windows-performance.md +++ b/docs/windows-performance.md @@ -342,8 +342,53 @@ sensor/window and its discrepancy is unresolved. Do not infer a physical CPU limit, isolated shader occupancy or a FSR4 speed comparison. Query overhead is median 0.094 ms per 500 ms poll, with a roughly 12 ms first query. Results: `test-results/silent-hill2-k-telemetry-unprofiled-4k`. The game stability -summary is still pending its existing ten-minute bound. Remaining K work now -prioritizes transformer/output-kernel cost; frame scheduling remains measured. +summary completes 37187 K frames with zero backend failures, CPU image copies, +previous-frame outputs or recorded crashes, stopping at its ten-minute bound. +Remaining K work now prioritizes transformer/output-kernel cost; frame +scheduling remains measured. + +### Rejected packed FP16 accumulator experiment + +An isolated `D4R_K_F16_WMMA` enc1 build uses gfx12's documented packed FP16 +C/D intrinsic, retaining the current operand/accumulator lane mapping. It +removes most conversion instructions (1066 static F16/F32 conversions become +two) and reduces main-kernel VGPRs from 218 to 201; LDS remains 24576 bytes, +with zero spills. These are compiler/ISA counts, not GPU time or occupancy. +The candidate is rejected before moving to another layer: its real replay +changes 1182727 half codes in the output allocation, max absolute difference +0.0390625 from the validated baseline. All values are finite. Against the +current twelve-window NumPy reference, PSNR drops from baseline 78.7702 dB +to 69.2377 dB; max absolute error rises from 0.009765625 to 0.0390625. +Two independent PTX windows give 70.4022/70.08 dB. These pass the existing +coarse reference thresholds but do not satisfy this optimization's baseline +preservation gate. No game installation or performance claim follows. + +A public 64-matrix primitive reproducer isolates the difference to direct +instruction arithmetic. The packed intrinsic and the layout adapter agree +bit-for-bit. Direct FP16 output differs from FP32 WMMA rounded once to FP16 +even after one K16 step (1674 of 4096 elements, max absolute 0.00390625). +After eight steps, 2874 elements differ, max absolute 0.0625. This rules out +the adapter packing as the source of this fixture's difference; the hardware +internal accumulation/rounding cause is not established. Production K keeps +its existing FP32-instruction/FP16-rounding baseline and the experimental flag +is off in every normal build. Results: `test-results/k-f16-wmma-enc1` and +`test-results/wmma-f16-arithmetic-compare-fixed`. +The finished `test-wmma-f16.ps1` repeats the same 64-case result. Rebuilding +every ordinary K/M object leaves all sixteen byte-identical; all seventeen +CTest gates pass and enc1 still compiles for gfx1101 (compile-only). Logs: +`test-results/wmma-f16-script-check` and `test-results/k-f16-baseline-ctest.log`. + +```powershell +.\scripts\windows\test-wmma-f16.ps1 +.\scripts\windows\build-native-k-f16-wmma.ps1 -Layer enc1 +``` + +The first command compiles and runs the public arithmetic diagnostic without +NVIDIA inputs. Its success means the reproducer executed correctly, not that +the two instruction paths are equivalent; JSON retains `baselineEquivalent` +and `gameOptimizationAccepted`. The second compiles the rejected experiment +in a separate directory and generates no override manifest. Intrinsic API: +[Clang gfx12 FP16 WMMA](https://clang.llvm.org/docs/AMDGPUBuiltinReference.html#builtin-amdgcn-wmma-f16-16x16x16-f16-w32-gfx12). Further game coverage (HIP kernel profiling disabled, GPU output checks enabled): M at 1920x1080 completes 4833 frames, all finite. K at 3840x2160 completes diff --git a/docs/windows-rdna4-port.md b/docs/windows-rdna4-port.md index 3b5f653..789ec35 100644 --- a/docs/windows-rdna4-port.md +++ b/docs/windows-rdna4-port.md @@ -140,6 +140,21 @@ metrics as absent and preserving the driver's raw timestamp. See The combined capture smoke test passes, followed by the user-confirmed 1935-present/60-poll measurement recorded at the top of this document. Logs and runtime/DLL hashes: `test-results/silent-hill2-k-telemetry-unprofiled-4k`. +Its ten-minute game summary completes **37187 K/4K frames**, zero backend +errors/CPU image copies/previous-frame outputs or recorded crashes. + +The next isolated K enc1 experiment replaces FP32 WMMA + FP16 rounding with +gfx12's packed FP16 accumulator instruction. ISA conversions and VGPRs fall, +but the real output changes and reference PSNR falls. A 64-matrix public GPU +reproducer shows the direct instruction differs from FP32 + FP16 rounding +already at a single K16 step while the adapter is exact against the packed +instruction. The experiment is rejected at enc1; no subsequent layer or game +uses it. Baseline preservation takes priority over the instruction reduction. +`D4R_K_F16_WMMA` remains disabled in every normal build. Reproducer/results and +exact metrics are in `windows-performance.md`. +All sixteen ordinary K/M objects remain byte-identical after rebuilding with +the disabled experimental code. All seventeen CTest gates pass; gfx1101 enc1 +compile-only compatibility and the ready-made public WMMA script also pass. The public prerelease `windows-rdna4-dev-20261002-a346d76` is published to `xdfnx-dev/d4r`, containing source commit `a346d76` and ZLUDA `a1c506f`. diff --git a/kernels/common/wmma_layout.h b/kernels/common/wmma_layout.h index e417b0c..6335874 100644 --- a/kernels/common/wmma_layout.h +++ b/kernels/common/wmma_layout.h @@ -229,6 +229,20 @@ WM_FN wm_f8v wm_mma(const wm_op& a, const wm_op& b, wm_f8v c) { return __builtin_amdgcn_wmma_f32_16x16x16_f16_w32_gfx12(__builtin_bit_cast(wm_h8, a), __builtin_bit_cast(wm_h8, b), c); } +#ifdef D4R_K_F16_WMMA +// Rejected K experiment: direct packed F16 changes the arithmetic on gfx1201. +// The public WMMA reproducer and enc1 replay are in tools/windows and docs. +// Experimental gfx12 packed FP16 C/D. K's strict baseline enters each K16 +// step with an already FP16-rounded accumulator. Keep the existing logical +// eight-row contract; only this instruction/register representation changes. +WM_FN wm_f8v wm_mma_f16_step(const wm_op& a, const wm_op& b, wm_f8v c) +{ + const wm_h8 hc = __builtin_convertvector(c, wm_h8); + const wm_h8 hd = __builtin_amdgcn_wmma_f16_16x16x16_f16_w32_gfx12( + __builtin_bit_cast(wm_h8, a), __builtin_bit_cast(wm_h8, b), hc); + return __builtin_convertvector(hd, wm_f8v); +} +#endif #else WM_FN wm_f8v wm_mma11(const wm_u8v& a, const wm_u8v& b, wm_f8v c) { diff --git a/kernels/k/pwin_common.h b/kernels/k/pwin_common.h index 3867fa1..78cae77 100644 --- a/kernels/k/pwin_common.h +++ b/kernels/k/pwin_common.h @@ -80,6 +80,12 @@ __device__ __forceinline__ half_t l2_sum(const half_t* row) // PWIN_F32ACC: keep the accumulator in f32 through the chain (rounded to f16 where the values are used) __device__ __forceinline__ f8v mma16(const op_t& a, const op_t& b, f8v c) { +#ifdef D4R_K_F16_WMMA +#if !defined(__GFX12__) || D4R_WMMA_LAYOUT != 12 || !defined(D4R_K_FP16_BASELINE) +#error "experimental packed F16 WMMA requires gfx12 native layout and strict K baseline" +#endif + return wm_mma_f16_step(a, b, c); +#else f8v d = wm_mma(a, b, c); #if !defined(PWIN_F32ACC) || defined(D4R_K_FP16_BASELINE) #pragma unroll @@ -87,6 +93,7 @@ __device__ __forceinline__ f8v mma16(const op_t& a, const op_t& b, f8v c) d[i] = (float)(half_t)d[i]; #endif return d; +#endif } __device__ __forceinline__ f8v splat8(float v) diff --git a/scripts/windows/build-native-k-f16-wmma.ps1 b/scripts/windows/build-native-k-f16-wmma.ps1 new file mode 100644 index 0000000..56b02bf --- /dev/null +++ b/scripts/windows/build-native-k-f16-wmma.ps1 @@ -0,0 +1,40 @@ +[CmdletBinding()] +param([ValidateSet('enc0','enc1','enc2','enc3','enc4','dec5','dec4','dec3','dec2','dec1','dec0')] + [string]$Layer='enc1', [string]$HipRoot, [string]$OutputDirectory) +$ErrorActionPreference='Stop' +$repo=[IO.Path]::GetFullPath((Join-Path $PSScriptRoot '../..')) +if (!$HipRoot) { $HipRoot=Join-Path $repo '.tools/therock-10.2.0a20260929/_rocm_sdk_core' } +if (!$OutputDirectory) { $OutputDirectory=Join-Path $repo 'build/native-k-f16-wmma-experimental' } +$OutputDirectory=[IO.Path]::GetFullPath($OutputDirectory) +if ($OutputDirectory -eq [IO.Path]::GetFullPath((Join-Path $repo 'build/native-k-gfx1201'))) { + throw 'Keep the validated baseline separate from this experimental build.' +} +$compiler=Join-Path $HipRoot 'lib/llvm/bin/clang++.exe' +if (!(Test-Path -LiteralPath $compiler)) { $compiler=Join-Path $HipRoot 'bin/clang++.exe' } +if (!(Test-Path -LiteralPath $compiler)) { throw "HIP compiler missing under $HipRoot" } +$name="dltss_pwin_${Layer}_layer" +$source=Join-Path $repo "kernels/k/$name.hip" +$flags=@() +if ((Get-Content -LiteralPath $source -Raw) -match '// d4r-build-flags: ([^\r\n]+)') { $flags=$Matches[1].Trim().Split(' ') } +New-Item -ItemType Directory -Force $OutputDirectory | Out-Null +$object=Join-Path $OutputDirectory "$name.hsaco" +$ErrorActionPreference='Continue' +& $compiler -x hip -std=c++17 --offload-arch=gfx1201 --offload-device-only --no-gpu-bundle-output ` + -mno-wavefrontsize64 -nogpuinc -nogpulib -O3 -DD4R_DEVICE_ONLY_MINIMAL ` + -DD4R_K_FP16_BASELINE -DD4R_K_F16_WMMA @flags $source -o $object ` + > (Join-Path $OutputDirectory "$name.build.stdout.log") 2> (Join-Path $OutputDirectory "$name.build.stderr.log") +$ErrorActionPreference='Stop' +if ($LASTEXITCODE) { throw "Experimental K FP16 WMMA compile failed: $Layer" } +$isa=Join-Path $OutputDirectory "$name.isa.txt" +& (Join-Path (Split-Path $compiler) 'llvm-objdump.exe') --disassemble --mcpu=gfx1201 $object > $isa +if ($LASTEXITCODE -or !(Select-String -LiteralPath $isa -Pattern 'v_wmma_f16_16x16x16_f16')) { + throw 'Packed FP16 WMMA is absent from ISA; no alternative instruction is accepted.' +} +& (Join-Path (Split-Path $compiler) 'llvm-readobj.exe') --notes $object ` + > (Join-Path $OutputDirectory "$name.metadata.txt") +if ($LASTEXITCODE) { throw 'Native object metadata extraction failed.' } +@{layer=$Layer; architecture='gfx1201'; experimental=$true; numericalValidation='pending'; + compiler=$compiler; sha256=(Get-FileHash -LiteralPath $object -Algorithm SHA256).Hash} | + ConvertTo-Json | Set-Content -LiteralPath (Join-Path $OutputDirectory "$name.build-info.json") -Encoding UTF8 +Write-Host "Compiled experimental gfx1201 packed FP16 WMMA: $Layer. Validate before selecting another layer." +# This script does not generate an override manifest or install into the game. diff --git a/scripts/windows/test-wmma-f16.ps1 b/scripts/windows/test-wmma-f16.ps1 new file mode 100644 index 0000000..0a6c2da --- /dev/null +++ b/scripts/windows/test-wmma-f16.ps1 @@ -0,0 +1,24 @@ +[CmdletBinding()] +param([string]$HipRoot, [string]$PackageRoot, [string]$OutputDirectory) +$ErrorActionPreference='Stop' +$repo=[IO.Path]::GetFullPath((Join-Path $PSScriptRoot '../..')) +if (!$HipRoot) { $HipRoot=Join-Path $repo '.tools/therock-10.2.0a20260929/_rocm_sdk_core' } +if (!$PackageRoot) { $PackageRoot=Join-Path $repo 'dist/windows-rdna4-command-list' } +if (!$OutputDirectory) { $OutputDirectory=Join-Path $repo ('test-results/wmma-f16-' + (Get-Date -Format 'yyyyMMdd-HHmmss-fff')) } +$compiler=Join-Path $HipRoot 'lib/llvm/bin/clang++.exe' +if (!(Test-Path -LiteralPath $compiler)) { $compiler=Join-Path $HipRoot 'bin/clang++.exe' } +New-Item -ItemType Directory -Force $OutputDirectory | Out-Null +$module=Join-Path $OutputDirectory 'wmma_f16_compare.hsaco' +& $compiler -x hip -std=c++17 --offload-arch=gfx1201 --offload-device-only --no-gpu-bundle-output ` + -mno-wavefrontsize64 -nogpuinc -nogpulib -O3 (Join-Path $repo 'tools/windows/wmma_f16_compare.hip') -o $module +if ($LASTEXITCODE) { throw 'WMMA arithmetic probe compilation failed.' } +$previousPython=$env:PYTHONPATH +try { + $env:PYTHONPATH="$(Join-Path $repo '.tools/python/vendor');$previousPython" + & python.exe (Join-Path $repo 'tools/windows/wmma_f16_compare.py') ` + --probe (Join-Path $PackageRoot 'bin/d4r_native_replay_probe.exe') --module $module ` + --hip-root $HipRoot --output-dir $OutputDirectory + if ($LASTEXITCODE) { throw "WMMA arithmetic diagnostic failed; logs: $OutputDirectory" } +} finally { $env:PYTHONPATH=$previousPython } +# A successful diagnostic is not acceptance of native FP16 accumulators for K. +# Read baselineEquivalent/gameOptimizationAccepted in validation.json. diff --git a/tools/windows/wmma_f16_compare.hip b/tools/windows/wmma_f16_compare.hip new file mode 100644 index 0000000..79a357f --- /dev/null +++ b/tools/windows/wmma_f16_compare.hip @@ -0,0 +1,40 @@ +// Public synthetic reproducer: compare packed F16 WMMA with F32 WMMA followed +// by one FP16 rounding after each K16 step. Never part of game packaging. +#define D4R_K_F16_WMMA +#include "../../kernels/common/wmma_layout.h" + +// The replay utility passes a single by-value parameter block, like K/M. +struct CompareParams { const _Float16* a; const _Float16* b; const float* c; float* output; }; +extern "C" __attribute__((global)) void d4r_wmma_f16_compare(CompareParams params) +{ + const _Float16* a = params.a; + const _Float16* b = params.b; + const float* c = params.c; + float* output = params.output; + const unsigned test = __builtin_amdgcn_workgroup_id_x(); + const unsigned token = wm_lane() & 15; + const unsigned base = test * 256; + const wm_op aa = wm_op_load(a + base + token * 16); + const wm_op bb = wm_op_load(b + base + token * 16); + wm_f8v baseline, adapted; + wm_h8 packed; + for (unsigned i = 0; i < 8; ++i) { + baseline[i] = c[base + wm_acc_row(i) * 16 + token]; + packed[i] = (_Float16)baseline[i]; + } + adapted = baseline; + const unsigned steps = 1u << (test / 16); + for (unsigned step = 0; step < steps; ++step) { + baseline = wm_mma(aa, bb, baseline); + for (unsigned i = 0; i < 8; ++i) baseline[i] = (float)(_Float16)baseline[i]; + packed = __builtin_amdgcn_wmma_f16_16x16x16_f16_w32_gfx12( + __builtin_bit_cast(wm_h8, aa), __builtin_bit_cast(wm_h8, bb), packed); + adapted = wm_mma_f16_step(aa, bb, adapted); + } + for (unsigned i = 0; i < 8; ++i) { + const unsigned index = wm_acc_row(i) * 16 + token; + output[test * 768 + index] = baseline[i]; + output[test * 768 + 256 + index] = (float)packed[i]; + output[test * 768 + 512 + index] = adapted[i]; + } +} diff --git a/tools/windows/wmma_f16_compare.py b/tools/windows/wmma_f16_compare.py new file mode 100644 index 0000000..9f2ddef --- /dev/null +++ b/tools/windows/wmma_f16_compare.py @@ -0,0 +1,67 @@ +"""Public gfx1201 WMMA arithmetic comparison; does not load DLSS or weights.""" +import argparse +import json +import os +import pathlib +import struct +import subprocess + +import numpy as np + + +def main(): + p = argparse.ArgumentParser(description=__doc__) + p.add_argument('--probe', type=pathlib.Path, required=True) + p.add_argument('--module', type=pathlib.Path, required=True) + p.add_argument('--hip-root', type=pathlib.Path, required=True) + p.add_argument('--output-dir', type=pathlib.Path, required=True) + args = p.parse_args() + root = args.output_dir.resolve() + root.mkdir(parents=True, exist_ok=True) + fixture = root / 'fixture' + fixture.mkdir(exist_ok=True) + rng = np.random.default_rng(90701201) + a = rng.uniform(-1, 1, (64, 16, 16)).astype(np.float16) + b = rng.uniform(-1, 1, a.shape).astype(np.float16) + c = rng.uniform(-.25, .25, a.shape).astype(np.float16).astype(np.float32) + initial_output = np.full((64, 3, 16, 16), np.nan, np.float32) + arrays = (a, b, c, initial_output) + bases = [0x10000000 * (i + 1) for i in range(4)] + (fixture / 'args.bin').write_bytes(struct.pack('<4Q', *bases)) + lines = ['kernel d4r_wmma_f16_compare', 'launch 64 1 1 32 1 1 0', 'args 32'] + for i, data in enumerate(arrays): + data.tofile(fixture / f'alloc-{i}.bin') + lines += [f'alloc {i} {bases[i]:#x} {data.nbytes}', f'pointer {i * 8} {i} 0'] + (fixture / 'manifest.txt').write_text('\n'.join(lines) + '\n', encoding='utf-8') + env = os.environ.copy() + env['D4R_DIAG_DIR'] = str(root) + with (root / 'gpu.stdout.log').open('wb') as out, (root / 'gpu.stderr.log').open('wb') as err: + status = subprocess.run([str(args.probe.resolve()), '--hip-root', str(args.hip_root.resolve()), + '--module', str(args.module.resolve()), '--fixture-dir', str(fixture), '--output-dir', + str(root / 'output'), '--iterations', '1'], env=env, stdout=out, stderr=err, timeout=120) + if status.returncode: + raise RuntimeError(f'GPU probe exit {status.returncode}; see {root}') + output = np.fromfile(root / 'output/alloc-3.bin', np.float32).reshape(64, 3, 16, 16) + if not np.isfinite(output).all(): + raise RuntimeError('Nonfinite/unwritten WMMA output') + rows = [] + for group, steps in enumerate((1, 2, 4, 8)): + values = output[group * 16:(group + 1) * 16] + difference = np.abs(values[:, 0].astype(np.float64) - values[:, 1].astype(np.float64)) + row = dict(steps=steps, elements=difference.size, differing=int(np.count_nonzero(difference)), + maxAbsoluteError=float(difference.max()), rms=float(np.sqrt(np.mean(difference ** 2)))) + rows.append(row) + print('WMMA_F16_COMPARE ' + json.dumps(row), flush=True) + adapter_exact = np.array_equal(output[:, 1].view(np.uint32), output[:, 2].view(np.uint32)) + if not adapter_exact: + raise RuntimeError('Direct packed instruction and layout adapter differ') + result = dict(passed=True, architecture='gfx1201', publicSynthetic=True, adapterExact=adapter_exact, + baselineEquivalent=all(row['differing'] == 0 for row in rows), cases=rows, + gameOptimizationAccepted=False) + (root / 'validation.json').write_text(json.dumps(result, indent=2) + '\n', encoding='utf-8') + print('PASS WMMA_F16_ARITHMETIC_DIAGNOSTIC adapter_exact=1 baseline_equivalent=' + + str(int(result['baselineEquivalent'])), flush=True) + + +if __name__ == '__main__': + main()