RDNA2 has native kernels for K's layers only; L's and M's test frames still
compiled their transformer layers with emulated matrix products for most of
an hour. The builder now builds e, f, k, k-ldr on gfx103x (detected like the
native kernels: D4R_GPU_ARCH, else the KFD node with the most SIMDs);
D4R_ENGINE_ALL_MODELS=1 builds L and M anyway. install.sh takes the model
list from build-engine-models.sh --list-sr-models and keys the shared cache
on it.
Models build in alphabetical order (E, F, K, L, M). The first frame renders
alone to create the Proton prefix; the rest render at the same time, each
model built as soon as its frame is done (ZLUDA compiles in one thread per
process). Cold install.sh --yes on a Steam Deck: 185 s.
DLSS loads every module of a model at feature creation and ZLUDA compiled
each one, even when the bridge had already linked a native kernel for every
visible entry and the compiled code never ran. On RDNA2 the transformer
layers emulate MMA and each took minutes in LLVM's register allocator, so
install.sh's K test frame looked hung for about half an hour on a Steam Deck.
Such a module (no globals, 64-bit) now loads the first native code object as
a stand-in and keeps its PTX; it is compiled only if a function without a
native replacement is requested. The K frame takes ~40 s on the Deck.
The models list and its completeness check (have) only knew rr-* and
k/m/l: an up-to-date stamp skipped the builder, so engine\e and engine\f
(the DLSS 3 CNN, the RDNA2 route) were never built on installs that
already had the other models, and copy_models skipped them from the
shared cache. The prompt now names the RDNA2 route.
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.
rozaxlia