wouter/d4r
1
0 Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Improve GPU compatibility, RDNA2 kernels and startup diagnostics

Browse Source 
- Match render/compute adapters and gate RR models by GPU capabilities
- Add RDNA2 native texture kernels and fix FP16 rounding
- Harden model builds, cache validation and packaging
- Add startup tracing, timeouts, documentation and regression tests
This commit is contained in:
ayo committed 2026-10-19 11:36:40 -04:00
1 parent ca93d0fbf7
commit 2558fcbd39
32 files changed
+1131 -81
Download Patch File Download Diff File Expand all files Collapse all files

No files matched your search

kernels/build.sh
+25 -8
View File
@@ -5,8 +5,10 @@
# usage: kernels/build.sh [all|k|l|m|tex|rr|rrswin] [OUT_DIR] (default: all, kernels/out/native)
#
# D4R_ROCM_DIR ROCm installation with clang and the HIP device libraries (default /opt/rocm)
# D4R_GPU_ARCH target GPU (default gfx1101): an RDNA3 (gfx110x) or RDNA4 (gfx120x) target; the kernels
# need wave32 WMMA, and kernels/common/wmma_layout.h picks the register layout per target
# D4R_GPU_ARCH target GPU (default gfx1101): an RDNA3 (gfx110x) or RDNA4 (gfx120x) target, where
# kernels/common/wmma_layout.h picks the WMMA register layout per target, or an RDNA2 (gfx103x)
# target: there only K's layers (software matrix products) and the texture kernels that
# need no WMMA (tex: K's natively translated hiluma output and input kernels) are built
# D4R_NATIVE_FP8 1 (RDNA4 only): the variant for ZLUDA's native FP8 WMMA (d4r.ini [Kernels] NativeFp8):
# M's layers on FP8 WMMA with NVIDIA's e4m3 activations, texture tails matching ZLUDA's
# D4R_ZLUDA_WMMA_FP8_NATIVE lowering. K is the same in both variants.
@@ -45,6 +47,12 @@ CLANG="$ROCM/lib/llvm/bin/clang++"
[[ -x "$CLANG" ]] || { echo "clang++ not found in $ROCM/lib/llvm/bin (set D4R_ROCM_DIR)" >&2; exit 2; }
mkdir -p "$OUT"
case "$WHAT" in all|k|l|m|tex|rr|rrswin) ;; *) echo "unknown kernel family: $WHAT" >&2; exit 2 ;; esac
# RDNA2 has no WMMA instruction: M/L's Swin layers, their texture tails and Ray Reconstruction call it directly
RDNA2=0
if [[ "$ARCH" == gfx103* ]]; then
RDNA2=1
case "$WHAT" in m|l|rr|rrswin) echo "$WHAT kernels need WMMA, which RDNA2 ($ARCH) does not have" >&2; exit 2 ;; esac
fi
# Never certify a directory containing older fast binaries as an accuracy set.
if [[ "$ACCURACY" == 1 ]] && compgen -G "$OUT/*.hsaco" >/dev/null &&
{ [[ ! -f "$OUT/d4r-accuracy.txt" ]] || [[ "$(cat "$OUT/d4r-accuracy.txt")" != 1 ]]; }; then
@@ -82,14 +90,15 @@ if [[ "$WHAT" == all || "$WHAT" == k ]]; then
echo "== DLSS 4 (preset K) transformer layers"
for src in "$HERE"/k/dltss_pwin_*.hip; do build_hip "$src"; done
fi
if [[ "$WHAT" == all || "$WHAT" == m || "$WHAT" == l ]]; then
if [[ "$RDNA2" == 0 && ( "$WHAT" == all || "$WHAT" == m || "$WHAT" == l ) ]]; then
echo "== DLSS 4.5 (presets L/M) shared Swin layers"
for src in "$HERE"/m/rrlite_*.hip; do build_hip "$src" "$([[ "$FP8" == 1 ]] && echo -DD4R_FP8_WMMA)"; done
fi
if [[ "$WHAT" == all || "$WHAT" == tex || "$WHAT" == l ]]; then
echo "== texture kernels (NVIDIA PTX with native parts)"
: "${D4R_DLSS_DLL:?set D4R_DLSS_DLL to nvngx_dlss.dll (310.7)}"
: "${D4R_ZLUDA_EMIT:?set D4R_ZLUDA_EMIT to d4r_emit from a ZLUDA build with patches/zluda}"
# RDNA2 builds only the PTX-to-HIP translations below, which ZLUDA does not compile
[[ "$RDNA2" == 1 ]] || : "${D4R_ZLUDA_EMIT:?set D4R_ZLUDA_EMIT to d4r_emit from a ZLUDA build with patches/zluda}"
PTX_DIR="${D4R_DLSS_PTX_DIR:-$HERE/extracted/ptx}"
if [[ -z "${D4R_DLSS_PTX_DIR:-}" ]]; then
python3 "$HERE/tools/extract_dlss_ptx.py" "$D4R_DLSS_DLL" "$PTX_DIR"
@@ -152,7 +161,10 @@ if [[ "$WHAT" == all || "$WHAT" == tex || "$WHAT" == l ]]; then
for range in hdr ldr; do native_ins+=("hiluma_engine_input_${depth}_${mv}_${range}_v2_rel"); done
done
done
for spec in "${specs[@]}"; do
# RDNA2: none of the ZLUDA-compiled texture builds above (their tails call WMMA, and the sust_only builds
# are not validated there); only the natively translated hiluma kernels below
[[ "$RDNA2" == 1 ]] && specs=()
for spec in "${specs[@]+"${specs[@]}"}"; do
IFS=: read -r kernel src mode <<< "$spec"
[[ "$ACCURACY" == 0 && " ${native_outs[*]} " == *" $kernel "* ]] && continue
# accuracy sets are wave32, except M's post and downsample kernels on RDNA4: their wave64 builds keep the
@@ -180,8 +192,13 @@ if [[ "$WHAT" == all || "$WHAT" == tex || "$WHAT" == l ]]; then
python3 "$NAT/ptx2hip.py" --ctile --as1 "$ptx" "$native_out" > "$NT/output_translated.hip"
python3 "$NAT/rewrite_output.py" "$NT/output_translated.hip" > "$NT/$native_out.hip"
depth_tile=(); [[ "$native_out" == *_mvhi_* ]] && depth_tile=(-DD4R_K_NO_DEPTH_TILE)
D4R_ROCM_DIR="$ROCM" D4R_GPU_ARCH="$ARCH" "$NAT/build_native.sh" "$NT/$native_out.hip" "$OUT/$native_out.hsaco" \
-DD4R_NATIVE_F2I "${depth_tile[@]}" -DD4R_AS1_STORES
# RDNA3/4: wave64 (FP32 dual issue; 0.49 -> 0.44 ms on an RX 7700 XT, same bytes). RDNA2 issues wave64
# in two passes, so it gets wave32 (no stack frame, 96 VGPRs), and WGP mode: a 256-thread workgroup
# then takes 2 of each SIMD's 10 wave slots instead of 4, so 10 waves fit per SIMD instead of 8
# (measured with RDNA2's occupancy on the RX 7700 XT: 0.596 -> 0.568 ms, same bytes).
wave=(); [[ "$RDNA2" == 1 ]] && wave=(WAVE=32 CUMODE=0)
env "${wave[@]}" D4R_ROCM_DIR="$ROCM" D4R_GPU_ARCH="$ARCH" "$NAT/build_native.sh" "$NT/$native_out.hip" \
"$OUT/$native_out.hsaco" -DD4R_NATIVE_F2I "${depth_tile[@]}" -DD4R_AS1_STORES
done
for native_in in "${native_ins[@]}"; do
ptx="$(ptx_of "$native_in")"; [[ -n "$ptx" ]] || { echo "no PTX defines $native_in" >&2; exit 2; }
@@ -192,7 +209,7 @@ if [[ "$WHAT" == all || "$WHAT" == tex || "$WHAT" == l ]]; then
done
fi
fi
if [[ "$WHAT" == rr || ( "$WHAT" == all && -n "${D4R_DLSSD_DLL:-}" ) ]]; then
if [[ "$WHAT" == rr || ( "$WHAT" == all && "$RDNA2" == 0 && -n "${D4R_DLSSD_DLL:-}" ) ]]; then
echo "== Ray Reconstruction texture kernels"
: "${D4R_DLSSD_DLL:?set D4R_DLSSD_DLL to the CUDA-capable nvngx_dlssd.dll}"
: "${D4R_ZLUDA_EMIT:?set D4R_ZLUDA_EMIT to d4r_emit from a patched ZLUDA build}"
kernels/k/pwin_common.h
+11
View File
@@ -61,6 +61,17 @@ __device__ __forceinline__ uint32_t lane_id()
#undef PWIN_NATIVE_F16ACC
#undef PWIN_F16ACC_NATIVE
#endif
// RDNA2: the weight images use plain global loads. The raw-buffer path below builds its resource with gfx11's
// dword3 (0x31004000: unified format 4, 8_SSCALED), which gfx11 ignores for untyped loads; gfx10 descriptors want
// a 32-bit format there (AMD's own libraries use 0x31014000 on gfx103x), and no gfx10 encoding has been validated
// on RDNA2 hardware, so RDNA2 takes no descriptor at all. Keyed on D4R_WMMA_SOFT (always set on gfx10) so that the
// RDNA3 validation build of the RDNA2 kernels runs the same loads.
#ifdef D4R_WMMA_SOFT
#ifndef PWIN_NO_BUFFER_IMG
#define PWIN_NO_BUFFER_IMG
#endif
#undef PWIN_GELU_LUT_BUFFER
#endif
typedef wm_op op_t; // WMMA operand: u8v (gfx11) or u4v (gfx12)
// one k16 step with the f16 accumulator of NVIDIA's f16 wmma (rounded after the step)
kernels/native/build_native.sh
+3 -1
View File
@@ -1,5 +1,6 @@
#!/usr/bin/env bash
# build_native.sh SRC.hip OUT.hsaco [extra clang flags]: natively compiled DLSS texture kernel (gfx1101, wave64 default)
# WAVE=32|64 (default 64); CUMODE=1|0 (default 1): CU mode, or WGP mode (a workgroup's waves on all four SIMDs)
set -euo pipefail
ND=$(cd "$(dirname "$0")" && pwd)
SRC=$(realpath "$1") OUT=$(realpath -m "$2"); shift 2
@@ -8,9 +9,10 @@ R=${D4R_ROCM_DIR:-/opt/rocm}
ARCH=${D4R_GPU_ARCH:-gfx1101}
W=${WAVE:-64}
WF=$([[ $W == 64 ]] && echo -mwavefrontsize64 || echo -mno-wavefrontsize64)
CU=$([[ ${CUMODE:-1} == 1 ]] && echo -mcumode || echo -mno-cumode)
D=$(cd "$(dirname "$SRC")" && pwd)
T=$(mktemp -d); trap 'rm -rf "$T"' EXIT
(cd "$T" && "$R/lib/llvm/bin/clang++" -x hip -std=c++20 --offload-arch=$ARCH --offload-device-only -O3 $WF -mcumode \
(cd "$T" && "$R/lib/llvm/bin/clang++" -x hip -std=c++20 --offload-arch=$ARCH --offload-device-only -O3 $WF $CU \
-ffp-contract=off -fgpu-flush-denormals-to-zero -fno-fast-math -I"$D" -I"$ND" -I"$R/include" --rocm-path="$R" \
--rocm-device-lib-path="$R/amdgcn/bitcode" "$@" -o "$OUT" "$SRC" -save-temps=cwd --no-gpu-bundle-output)
S=$(ls "$T"/*.s | head -1)
kernels/native/output_tile.h
+3 -2
View File
@@ -24,8 +24,9 @@ RT void output_color_pack(f4 color, float exposure, uint32_t* a, uint32_t* b)
const float q = div_a(ln, 9.0f) + asf(0xBC4FAE29u);
weight = rcp_a(y) * q;
}
*a = pack16(ashu((h1)(y * weight)), ashu((h1)((exposure * co) * weight)));
*b = pack16(ashu((h1)((exposure * cg) * weight)), ashu((h1)color[3]));
// f2h_rn: rounded to f32 and then to f16, as ZLUDA's compile does (no v_fma_mixlo of the product)
*a = pack16(ashu(f2h_rn(y * weight)), ashu(f2h_rn((exposure * co) * weight)));
*b = pack16(ashu(f2h_rn((exposure * cg) * weight)), ashu(f2h_rn(color[3])));
}
RT int output_tile_coord(int relative, int lo, int hi)
kernels/native/ptx2hip.py
+2 -2
View File
@@ -593,10 +593,10 @@ class Translator:
if dt == 'f16' and stp == 'f32':
if 'rn' not in mods:
raise Err(op)
A(a[0], f'(h1){V(a[1], "f32")}', 'f16'); return
A(a[0], f'f2h_rn({V(a[1], "f32")})', 'f16'); return
if dt == 'f16x2' and stp == 'f32':
# d = {cvt(a) in the high half, cvt(b) in the low half}
A(a[0], f'pack16(ashu((h1){V(a[2], "f32")}), ashu((h1){V(a[1], "f32")}))', 'u32'); return
A(a[0], f'pack16(ashu(f2h_rn({V(a[2], "f32")})), ashu(f2h_rn({V(a[1], "f32")})))', 'u32'); return
if dt == 'f16' and stp == 'f16':
x = V(a[1], 'f16')
if 'rni' in mods:
kernels/native/ptx_rt.h
+10
View File
@@ -77,6 +77,16 @@ RT uint32_t f2u_rz(float x)
}
// ---------------------------------------------------------------- f16 (denormals kept)
// cvt.rn.f16.f32: the f32 value rounded to f16. The empty asm keeps LLVM from folding the conversion into the
// f32 operation that produced the value (fptrunc(fma(fpext ...)) -> v_fma_mixlo/mixhi_f16), which rounds the exact
// result to f16 once instead of to f32 and then to f16: on an RX 7700 XT that differed from the two-step result for
// 1 in ~29,000 random f16 operand triples, and in the hiluma output kernel it left single 1-ulp seeds that the
// temporal history then spread (ZLUDA keeps both roundings).
RT h1 f2h_rn(float x)
{
__asm__ volatile("" : "+v"(x));
return (h1)x;
}
RT h1 hfma(h1 a, h1 b, h1 c) { return __builtin_fmaf16(a, b, c); }
RT h2 h2fma(h2 a, h2 b, h2 c) { return __builtin_elementwise_fma(a, b, c); }
RT h1 hmaxn(h1 a, h1 b) { return __builtin_fmaximum_numf16(a, b); }