woutermans
b9254f7f96
Presets E (5) and F (6) run on d4r's own Vulkan engine without cooperative matrices: the input and reconstruction kernels translated from the user's DLL (engine/compile_sr_e.py, rr/ptx2rr.py), the network engine/sr/e_net.comp with packed v_pk_fma_f16 accumulation (engine/cnn_sr/mma_f16.glsl, the default head; --head-accumulation f32 keeps the f32 head). Runtime runtime_e.inc (EModel/EEngine), game adapter game_e.cpp, replay d4r-e; the shim serves them from engine\e and engine\f (D4R_ENGINE_MODEL_E_DIR / _F_DIR), LDR or HDR, any render size, not with AutoExposure. build-engine-models.sh builds both. Translator fixes that apply to every translated kernel: .relu on f16/f16x2 fma/add/mul was ignored; add.rz was always taken for the roundf idiom (lost the log1pf correction term); formatted image stores expand at the call site (SPIR-V image type fix); f16 mma.sync m16n8k8/k16, .const loads, shared 16-bit stores, aligned 32-bit arena words (D4R_WORD_ARENA), unified arena. Steam Deck gfx1033: 2.55 ms at 640x360 -> 1280x720, 6.20 ms at 1080p output; Rounded Quality sizes take mid_3x3 (reconstruction 2.44 -> 1.23 ms). Ray Reconstruction CNN presets A-C (compile_rr3.py, the packed RR flavor) stay out of this commit; engine/README.md already documents their tuning.
25 lines
1.6 KiB
Bash
Executable File
25 lines
1.6 KiB
Bash
Executable File
#!/usr/bin/env bash
|
|
set -euo pipefail
|
|
ROOT="$(cd "$(dirname "${BASH_SOURCE[0]}")/.." && pwd)"
|
|
OUT="${1:-$ROOT/build/native-engine}"
|
|
CXX="${CXX:-g++}"
|
|
mkdir -p "$OUT"
|
|
# The native HIP network (hip_net.h) is shared by both runners when ROCm is installed; without it they fall back
|
|
# to the engine's own Vulkan network. The library itself never links HIP.
|
|
ROCM="${D4R_ROCM_DIR:-/opt/rocm}"
|
|
HIP=()
|
|
if [[ -f "$ROCM/include/hip/hip_runtime_api.h" ]]; then
|
|
HIP=(-DD4R_ENGINE_HIP "$ROOT/engine/hip_net.cpp" -I"$ROCM/include" -L"$ROCM/lib" -lamdhip64 "-Wl,-rpath,$ROCM/lib")
|
|
fi
|
|
"$CXX" -std=c++17 -O3 -Wall -Wextra -Wno-missing-field-initializers -fPIC -shared \
|
|
"$ROOT/engine/runtime.cpp" "$ROOT/engine/game.cpp" "$ROOT/engine/game_m.cpp" "$ROOT/engine/game_rr.cpp" "$ROOT/engine/game_e.cpp" -lvulkan -o "$OUT/libd4r_engine.so"
|
|
"$CXX" -std=c++17 -O3 -Wall -Wextra -Wno-missing-field-initializers \
|
|
"$ROOT/engine/replay.cpp" "${HIP[@]}" -L"$OUT" -ld4r_engine -lvulkan '-Wl,-rpath,$ORIGIN' -o "$OUT/d4r-k"
|
|
"$CXX" -std=c++17 -O3 -Wall -Wextra -Wno-missing-field-initializers \
|
|
"$ROOT/engine/replay_m.cpp" "${HIP[@]}" -L"$OUT" -ld4r_engine -lvulkan '-Wl,-rpath,$ORIGIN' -o "$OUT/d4r-m"
|
|
"$CXX" -std=c++17 -O3 -Wall -Wextra -Wno-missing-field-initializers \
|
|
"$ROOT/engine/replay_rr.cpp" -L"$OUT" -ld4r_engine -lvulkan '-Wl,-rpath,$ORIGIN' -o "$OUT/d4r-rr"
|
|
"$CXX" -std=c++17 -O3 -Wall -Wextra -Wno-missing-field-initializers \
|
|
"$ROOT/engine/replay_e.cpp" -L"$OUT" -ld4r_engine -lvulkan '-Wl,-rpath,$ORIGIN' -o "$OUT/d4r-e"
|
|
printf 'Built %s, %s, %s and %s\n' "$OUT/d4r-k" "$OUT/d4r-m" "$OUT/d4r-rr" "$OUT/d4r-e"
|