OpenDLSS-NR: Een Vulkan-reimplementatie van het DLSS 5 Neural Rendering-netwerk van NVIDIA

Het project implementeert hetzelfde netwerk van 71 Swin / ViT-blokken als DLSS-NR build 310.8.0 en maakt gebruik van FP8 op de tensor cores. Niet alleen het uiteindelijke beeld komt overeen, maar ook de tussenstappen: alle 75 blokgrenzen zijn byte voor byte identiek.

Daarnaast bevat de map ports/browser-webgpu/ een tweede, onafhankelijke implementatie. Deze draait in een browser zonder gebruik van tensor cores of FP8. De gebruiker dient zelf de gewichten aan te leveren in een modeldirectory volgens de beschreven structuur.

Het Netwerk

Het netwerk is een U-net van shifted-window transformer blocks met een globale ViT aan de basis. De specificaties zijn als volgt:

• Structuur: 71 blokken verdeeld over zes pooling-niveaus.
• Precisie: FP8 (E4M3) activaties met FP16 accumulatie.
• Gewichten: 141 MiB.

Het betreft een generatief neural rendering-netwerk (naar de terminologie van NVIDIA). In plaats van een upscaler, renderen deze netwerken het frame dat de engine al heeft getekend opnieuw. Hierbij wordt detail gegenereerd uit geïnjecteerde ruis en worden toon, structuur en huid aangepast op basis van een stijlinstelling. De input en output hebben dezelfde resolutie.

Input en Output

Het netwerk neemt één gerenderd frame als input (een low dynamic range proxy daarvan, drie banen van Gaussian-ruis, de output van het vorige frame na herprojectie en vijf conditionele scalairen). De output bestaat uit vier f32-kanalen per pixel: een RGB-residu en één temporal-blend logit.

Voor een volledige weergave van de graaf kan worden verwezen naar docs/network.md. NVIDIA beschrijft het model in het rapport DLSS 5: Generative Neural Rendering.

Bouwen en Uitvoeren

Installatie en Build-stappen

Gebruik de volgende PowerShell-commando's voor de installatie en het bouwen van het project:

# Eenmalig: installeert tools (glslang, Vulkan-Headers, volk, CMake, Ninja)
powershell -File scripts\fetch_tools.ps1 [-Npm]

# Bouwt shaders, PTX en build\dlss5vk.exe
powershell -File scripts\build.ps1

# Eenmalig voor de demo: installeert third_party\filament inclusief patch
powershell -File scripts\fetch_filament.ps1

# Eenmalig voor de demo: bouwt third_party\filament-install
powershell -File scripts\build_filament.ps1

# Bouwt build\demo\dlss5-demo.exe
powershell -File scripts\build_demo.ps1

Gebruik van de tool dlss5vk.exe

• Benchmark: build\dlss5vk.exe bench --model <dir> --width 768 --height 768
• Profilering (per-dispatch timings): build\dlss5vk.exe profile --model <dir> --width 768 --height 768
• Pariteitstest (bit-exactness tegen een fixture): build\dlss5vk.exe parity --model <dir> --fixture <dir>
• Verificatie (block-0 kernel-by-kernel bisect): build\dlss5vk.exe verify --model <dir> --fixture <dir>
• PTX-test: python scripts\ptx\testfastdivmod.py (test de PTX-deler voor elke n < 2^24 met numpy).

De demo kan via een dubbelklik worden gestart. Deze laadt scènes uit build\scenes via een dropdown-menu of via een glTF-bestand via de command line. De modeldirectory wordt bepaald door --model <dir>, anders via de omgevingsvariabele DLSS5VK_MODEL, of standaard in models\nr. Zie demo/README.md voor details over de renderer, toetsen, scènes en view.json.

Prestaties

Gemeten op een RTX 4070 SUPER, volledig netwerk per frame, minimum over 40 frames (241 dispatches bij elke resolutie).

Resolutie	Tijd
768x768	2.8 ms
1920x1080	7.8 ms
2560x1440	12.6 ms
3840x2160	29.3 ms

Opmerking: Onder constante belasting wisselt de GPU tussen twee kloksnelheden, waardoor medianen enkele procenten hoger liggen. Vergelijk daarom de minima.

Projectcomponenten

Onderdeel	Bestanden	Notities
Host	src/ (C++20)	Vulkan-context, model laden en weight re-layout, kernel wrappers, netwerkgraaf, CPU-referentie van de rekenkunde, dlss5vk tool.
GLSL kernels	shaders/	De referentieroute: cooperative-matrix FP8 GEMMs, fused 32-channel block, fused QKV + window attention, expert MLP, global attention, elementwise ops.
PTX kernels	scripts/ptx/	Python-generatoren die PTX genereren voor de snelle route: mma.sync E4M3 met f16 accumulatie, cp.async rings, barrier-free chaining, split-K GEMMs, streamed global attention.
Demo	demo/, third_party/filament.patch	Het netwerk geïntegreerd in een Filament (Apache-2.0) frame. Filament is gepatcht voor per-object motion vectors en een Vulkan interop hook. Gebruikt gltfio voor scènes en ImGui voor controls.
WebGPU port	ports/browser-webgpu/	Dezelfde netwerkstructuur in een browser. Bit-exact tegen dezelfde captures, maar zonder tensor cores, FP8, fusion of chaining. (72 ms bij 512x512 versus 2.7 ms op GPU).

Niet geïmplementeerd: DLSS-SR (dit is een ander netwerk). Het temporele pad is wel geïmplementeerd in de demo: de history-input lanes en de per-pixel blend logit sturen een herprojecteerde feedback-loop aan (docs/frame.md). De dlss5vk tool draait enkelvoudige frames zonder historie.

Vereisten

Hardware en Drivers

• Windows.
• NVIDIA Ada (of nieuwer) GPU.
• Driver met ondersteuning voor: VKKHRcooperativematrix, VKNVcooperativematrix2, VKEXTshaderfloat8 en VKNVcudakernel_launch.

Software

• Visual Studio 2022 of later met C++ x64 toolset.
• git.
• Python 3 (voor PTX-generatoren).
• Node.js + npm en Pillow (voor de scene converter).

De portable toolchain in tools/ (git-ignored) wordt gedownload via scripts\fetch_tools.ps1 en bevat glslang 16.6.0, Vulkan-Headers v1.4.363, volk, CMake 3.31 en Ninja 1.13. Een Vulkan SDK installatie is niet vereist.

Voor de demo is Filament v1.77.0 nodig. Dit wordt beheerd via scripts\fetchfilament.ps1 en scripts\buildfilament.ps1. De build-tree (ca. 6 GB) wordt geplaatst in %LOCALAPPDATA%\dlss5-vulkan of DLSS5FILAMENTBUILD_DIR.

Modeldirectory

De functie nr::Model leest een manifest.json bestand dat bestaat uit:

1. Een stages array: (id, bestand relatief aan de directory, packedByteLength, sha256).
2. Een tensors array: (naam, blok, layer, parameter, stage, stageOffset, byteLength).

De stage-bestanden bevatten de E4M3-gewichten als packed bytes; de host herstructureert deze naar de matrixvormen die de kernels gebruiken (src/nr_model.cpp). Deze repository produceert zelf geen modeldirectories. Alleen modellen met exact 71 blokken (conform build 310.8.0) worden geaccepteerd.

Fixtures

De parity-functie vergelijkt resultaten met opgenomen captures van het origineel. Een fixture is een directory met een manifest.json bevattende:

• Sleutels: sourceDimensions en fullDimensions.
• Input: proxy of inputFeatures (RGBA f32 image of f32 features).
• Checks: Wat er gevalideerd moet worden (boundaries, head, output).
• Referenties:
• boundaries: E4M3 referenties voor blokken en transities.
• referenceHead: De f32 RGBA head.
• nativeOutput: Het gecomponeerde beeld (dtype "f32" of "u8").

Alle gegevens worden gevalideerd voordat de GPU start. Oordelen zijn bit-exact (pass), gelijk tot aan het teken van nul (failure), binnen één code (alleen voor 8-bit captures) of een mismatch. Zie docs/numerics.md voor details.

Tuning Switches

Alle switches staan standaard op de snelle, exacte route. Elke switch behoudt de byte-identieke output.

• DLSS5VK_UNFUSED=1: Gebruikt de GLSL referentieroute (tot 2560x1440), waarbij elk intermediair wordt gematerialiseerd.
• DLSS5VKPTXDIR: Pad naar de PTX directory (standaard build/ptx).
• DLSS5VK_CHAIN=0: Plaatst barriers tussen elke launch in plaats van counter chaining.
• DLSS5VKPTXGEMM, GEMMT, GEMMV, BLOCK32, FFN, QKV, ATTN: Indien op 0 gezet, schakelt dit de betreffende kernel-familie terug naar de GLSL-versie.
• DLSS5VKATTNSTREAM: Forceert streamed global attention uit (0) of aan (1).
• DLSS5VKSPLITK=0, DLSS5VKVITCHAIN=0, DLSS5VKNOPTXMLP=1: Schakelt respectievelijk split-K, de ViT chain en de PTX MLP uit.
• DLSS5VKNOFUSE_PRE, POOL, UPRES, POST: Indien op 1 gezet, wordt de betreffende fusion verwijderd.
• DLSS5VKCHAINMASK: Bitmasker voor chaining (1: expert stages, 2: c32 blocks, 4: split GEMMs). Standaard 3.
• DLSS5VKDEFERMAX: De breedste stage waarvan de projection GEMM is gefused. Standaard 128.
• DLSS5VK_VALIDATION=1: Draait onder de Khronos validation layer.
• DLSS5VK_DEBUG=1: Print driver-berichten (inclusief die van de PTX-compiler).
• DLSS5VKLISTEXTENSIONS=1: Lijst beschikbare extensies op.

Documentatie

De index is te vinden in docs/README.md. Specifieke documenten zijn:

• network.md: De graaf.
• numerics.md: Het contract voor exactheid.
• weights.md: Layouts van gewichten.
• execution.md: Scheduling.
• frame.md: Het frame van de demo.

Disclaimer en Licentie

Geen affiliatie met NVIDIA Dit project is niet geaffilieerd met, onderschreven door of ondersteund door NVIDIA. Het bevat geen NVIDIA-software, gewichten, headers of instructies om deze te verkrijgen. Er worden geen rechten verleend op intellectueel eigendom van NVIDIA. De gebruiker is zelf verantwoordelijk voor de licenties van de modeldata die zij gebruiken.

Licentie Alles in deze repository is gelicentieerd onder de MIT-licentie. Third-party componenten staan vermeld in het NOTICE bestand.