OpenDLSS-NR: the network is a sequence of roundings

mdjsonmcpcopy↗ai

2026-09-26 · 24 min · explainer · gpu · kernels · quantization · performance · reproducibility · realtime · open-source

On 21 September 2026, @maanalaolaqy posted this:

I reverse engineered and reimplemented DLSS 5's neural rendering in Vulkan, from the architecture down to the last rounding. same input, same output as Nvidia's DLSS5, 7.8ms per 1080P on RTX 4070 super.

Each of its three claims names something the repository contains: a network it documents, a test it runs, a benchmark it reports. So I read OpenDLSS-NR against its own docs.

It re-renders the frame; it does not upscale it

DLSS has meant super resolution and frame generation for years, so "DLSS 5 neural rendering" reads like a better upscaler. The README says otherwise: "Input and output are the same resolution; it is not an upscaler." The network "re-renders the frame the engine already drew, generating detail from injected noise and adjusting tone, structure and skin under a style setting." DLSS super resolution is "a different network" and is not implemented.

The input is 16 f32 lanes per pixel (docs/network.md):

A DLSS-SR reader would look for depth, jitter and motion vectors, and none is a lane. Motion vectors do their work outside the network, where the demo's preprocess samples the previous output at the reprojected position with a five-tap Catmull-Rom filter, plus a flag saying whether a history exists at all. The output (the "head") is four f32 channels, composed like this:

neural  = clamp(proxy + rgb / 4, 0, 1)              # RGB residual, in sRGB code space
weight  = clamp(sigmoid(logit) * blendScale, 0, 1)  # blendScale: a learned f16, 0.7397 per docs/frame.md
display = lerp(neural, reprojected_history, weight)

So the network predicts a residual on the frame it was given, plus its own per-pixel opinion of how much history to keep. NVIDIA's project page describes DLSS 5 as "a one-step pixel-space diffusion model" conditioned on "the current rendered frame, engine motion vectors, carried temporal state, and artistic-direction values", running on GeForce RTX 50 Series GPUs. Reasoned: the lane list reads naturally as that design: the Gaussian lanes are the noise, one pass is the one step, and the reprojected output is the carried state. The repository does not make that mapping itself.

expand

expand

A U-net of windows, with a ViT at the bottom

The graph is recorded into one command buffer by src/nr_graph.cpp. It is a U-net of shifted-window transformer blocks over six pooling levels. Blocks 0 and 70 run at full resolution with 32 channels. The encoder widens 32, 64, 128, 256, 512 over blocks 1-30 while it halves the image five times; blocks 31-38 are a global ViT with 1024 channels at the bottom; the decoder mirrors it back up to block 70, with a skip from each encoder stage. Block 39 is the ViT's exit, a 1024-to-512 projection merged onto the skip, and the code runs it as a merge rather than an attention block. That is the "71-block" network. Its 153 weight records hold 141 MiB: FP8 (E4M3) matrices, with the f16 scale vectors and attention priors laid out beside them.

The graph does not run on your frame. It runs on a padded field, because six exact halvings and a decoder that upsamples whole 8-pixel windows must land back on the encoder's sizes. Each level is alignUp(ceil(size / 2), 4); the field is aligned to two to the power of the number of halvings that actually shrink the axis (plus one when level 0 would not be whole 8-pixel windows), floored at 320, and then widened once more under a condition the docs say has no visible reason: "The last line has no reason that is visible anywhere, and it still has to be reproduced: the field size decides the window grid, and therefore the result inside the valid rectangle too." The padding is a mirrored copy of the image, while the noise still hashes the padded coordinate.

The widget below is that rule, ported line for line from Geometry::fromValid. It reproduces every example the docs give: 512x512 runs on 576x512, 768x768 on 832x768, 644x768 on 768x768, 1920x1080 on 1920x1152 and 3840x2160 on 3840x2176. At the five sizes the repository times, it also shows the reported network time against a frame budget.

For a 1080p frame that means 6.7% more pixels than the frame, a level-5 grid of 32x20, and a ViT attending over 640 tokens. Widths from 1 to 16 and from 25 to 32 are refused: native crops there in the last block, the port does not implement the crop, and it throws rather than differ silently.

One block is FFN, then QKV, then window attention, then a projection, with two learned per-channel skip scales:

y   = FFN(x) + x * ffnScale                  # f16; E4M3 for the QKV input
out = Proj(WindowAttention(QKV(y))) + y * attnScale

Each skip is seeded into the accumulator as the C operand of the first matrix multiply, not added afterwards, "because the rounding differs". The FFN depends on width. At 32 channels it is dense, 32 to 128 to 32. At 64, 128 and 256 channels it is C/32 parallel paths of C -> 128 -> 32 concatenated and contracted; the weights are named experts, "but there is no router: every path runs on every token". The 512 stage splits into eight 64-wide branches, and the ViT is dense, 1024 to 4096 to 1024.

The attention is scaled cosine attention on 8x8 windows. Q and K are normalized to unit length in f16, Q is multiplied by a learned per-head temperature, and a learned 64x64 prior per head is the accumulator's starting value. The exponential is not a transcendental. It is a bit trick on a half: x = clamp(f16(0.044921875 s + 1.30078125), 1.03125, 1.5693359375), then (bits(x) << 5) ^ 0x8000 reinterpreted as a half, which is 2^(1.4375 s - 5.375) with a linear mantissa. There is no max subtraction in the softmax; the clamp alone keeps the exponentials in range, which is what lets the denominator live in f16. The window grid cycles through four origins per level, (0,0), (-4,-4), (-4,0), (0,-4), rather than the usual two, and a decoder stage continues the phase count its encoder stage left. Windows that hang off the field are not clipped: out-of-field tokens are zero vectors whose exp(prior) still enters the softmax denominator, and "Masking them out instead changes the result."

The docs are candid about how much of this is understood. A section titled "What has no derivation" lists constants, orderings and layouts that are simply what native does, the SiLU cubic's coefficients (-0.055908203125, 0.447265625, 0.89453125) among them: "the only account of them I can give is that they are what they are." A table beside it files eight 2-byte ViT tensors that the graph never reads under "inherited, no visible reason".

The contract is the roundings

This is the sentence the whole repository hangs from (docs/README.md):

Matching NVIDIA's output therefore means matching an ordering of roundings, not an error bound.

Every value that crosses a kernel boundary is either an E4M3 byte (4 exponent bits, 3 mantissa bits, maximum 448, no infinities) or an IEEE half. The publication rules are fixed. Round to half first, then from the half to E4M3, "never f32 straight to E4M3". Round to nearest even at both steps, and saturate at ±448. NaN publishes as +0, not as the E4M3 NaN code, because a zero row's cosine normalization is 0 * inf for every out-of-field token and the NaN code would poison the multiply that reads it. And the sign of a zero survives, because the captures are compared as bytes and -0 is 0x80.

The first rule sounds pedantic until you count what it changes. Near 1.0, E4M3 codes are 1/8 apart, so the boundary between two of them sits 1/16 above the lower one. That boundary is itself an exact half value. A float within half a half-ulp of it rounds onto it in f16, and then the E4M3 step sees an exact tie and sends it to the even code, which for half of those floats is the other side of the boundary from where a single rounding would put it.

Measured, by me, over every float32 in [1, 2): 32,768 of the 8,388,608 values, one in 256, publish a different byte depending on which order the roundings run in. I get the same count in every binade from 1/64 up to 256. Most values in this repository reach a publication already on the half grid, straight out of an f16 accumulator, and for those the order cannot matter. It matters wherever arithmetic produces an f32 that is not a half. The ViT's attention output is one: an f16 value accumulator times an f16 reciprocal is an exact f32 product, and global_attend.comp passes it through roundF16 before e4m3CodeFromF32. A port that kept activations in f32 and quantized them to FP8 in one step would take the shortcut at every layer. Reasoned: it would differ wherever a value lands in one of those strips; how many do across a frame depends on activations that are in fixtures I do not have.

The tensor cores are modelled just as literally. Every FP8 GEMM is a chain of 16x16x32 E4M3 multiplies that accumulate in f16, the same mma.sync.aligned.m16n8k32.row.col.f16.e4m3.e4m3.f16 instruction the docs say native uses. The docs, and ref::adaFp8Fdpa16 in src/reference.cpp, treat each k32 step as two groups of 16 products summed in fixed point:

E      = max(exp(accumulator), max over the 16 pairs of exp(a) + exp(b))
sum    = trunc(acc * 2^(13 - E)) + sum_i trunc(a_i * b_i * 2^(13 - E))   # 13 fractional bits, truncated
result = roundF16(sum * 2^E)

Three consequences follow, and the code depends on all of them. The K order is part of the result. The residual has to be seeded as the accumulator, not added at the end. And split-K is not a tuning knob: each partition accumulates from zero and the partial sums are added in f16, so the split points are specification. "The native ViT launches carry them as the z dimension of the grid": 4096/1024 for the FFN contraction, 1024/512 for qkv, 1024/256 for the projection. The softmax reduction tree is fixed the same way, down to which pairs of keys are added first.

My favourite detail is where this bit the author first. shaders/common.glsl rounds to half by manipulating the bit pattern rather than writing float(float16_t(x)), and the comment says why:

// Don't write float(float16_t(x)): the NVIDIA compiler drops it as a no-op and the rounding
// goes with it. That was the first parity bug, inside the SiLU.

A compiler that removes a round trip through f16 is doing something every numerical programmer would call an optimization. Here it is a different network.

Two routes through Vulkan

The repository computes the same bytes twice, on purpose. The GLSL route is 13 compute shaders on VK_KHR_cooperative_matrix and VK_EXT_shader_float8: a fused FP8 GEMM with a residual prologue and a SiLU-or-quantize epilogue, a whole 32-channel block in one dispatch, fused QKV plus window attention, the expert FFN, the global attention and the elementwise ops. The docs call it "the specification", and DLSS5VK_UNFUSED=1 runs its least-fused form.

The PTX route is what runs by default, and the most unusual engineering in the repository. VK_NV_cuda_kernel_launch lets a Vulkan command buffer launch a CUDA module JIT-compiled from PTX text. The PTX is emitted by Python: ptxgen.py (133 lines) is a tiny assembler with virtual registers and loops unrolled at generation time, and swin.py holds the network's arithmetic once, the MMA wrappers, the SiLU, the normalization, the bit-trick exponential and the softmax. Measured from scripts/build_shaders.ps1: the build emits 78 PTX variants, one per shape the graph asks for, for sm_89. The docs give the reason for the detour: GLSL cannot express "cp.async rings, mma.sync with explicit fragment placement, .maxnreg control, nanosleep backoff, and, most of all, release/acquire traffic between launches." It is the same lesson as FlashAttention-3: once the arithmetic is fixed, the speed is in the schedule. CUDA Rust is another answer to wanting to write the kernel yourself rather than only launch it; here the answer is a 133-line Python assembler.

The schedule has three layers. Fusion takes a 512x512 frame from about 538 dispatches, unfused, to 317 with fused attention and MLPs, to the default 241 with whole 32-channel blocks in one kernel. Chaining then removes the barriers between most of those 241: a producer does a release-add on a device counter, and the consumer spins on it with an acquire load and a nanosleep backoff instead of waiting behind a pipeline barrier. That is only safe if every wait finishes, and the docs make the argument explicitly. The waits must form a forward DAG, checked by Kernels::checkChainOrder at every recording. And the GPU must issue every workgroup of a launch before any of a later launch, which "is what NVIDIA hardware does, and it is not a Vulkan guarantee". Because nothing documents that order, it is made to fail safe: any wait that lasts one second sets bit 31 on the counters it is stuck on, releases every other waiter, and reports a failed frame instead of hanging the GPU.

The weights need work before any kernel sees them. The model stores E4M3 weights as MMA fragments, and the host re-lays them into plain [K/batchK][batchK/32][N][32] matrices at load. It also folds a permutation into the weight rows, the inverse of a native rotation of bits 1-3 of the channel index. That is legal only because the permutation stays inside each group of 16 products, which is exactly the group the fixed-point sum runs over; move it across a group and the arithmetic changes.

What "same output" is measured to mean

dlss5vk parity --model <dir> --fixture <dir> is the gate. A fixture is a directory of recorded captures of native's run: an input (a proxy image or the features themselves), the checks it gates (boundaries, head, output), and reference bytes for each. Every verdict is named:

The fixture contract is strict in the way test suites rarely are. A boundary fixture must account for all 75 comparable boundaries (blocks 0-69 plus the five encoder transitions) with a reference or a stated reason, so "A fixture cut short therefore fails; it cannot pass as a smaller suite." The head is compared on the production schedule, resubmitted three times by default, and must also equal the same graph with a barrier after every launch. The boundary captures come from a separate instrumented run whose head must equal production's.

The results are Reported. At 512x512 the boundary fixture is 75 boundaries and 57,704,448 E4M3 bytes, all bit-exact, signed zeros included. The 8-bit capture of native's final image matches within one code: 786,180 channels exact and 252 one code off. Eleven further fixtures from 644x768 to 3840x2160 are bit-exact on the composed RGB halves, and four of them, at 1024x768, 1920x1080, 2560x1440 and 3840x2160, on all 75 boundaries as well. Every route switch gives the same bytes.

That is a precise claim, and it has edges the docs state themselves:

• Single frames, no history. The dlss5vk tool runs frames without history, "which is what the reference captures were made with". The temporal loop, the sigmoid blend and the stored history, is implemented in the demo and not gated against native.
• Two places matched by effect. Native computes the cosine norm's reciprocal square root and the softmax's reciprocal with rsqrt.approx.ftz.f32 and rcp.approx.ftz.f32; the port uses other spellings (inversesqrt and 1.0 / x in GLSL, rsqrt.approx.f32 and the correctly rounded rcp.rn.f32 in PTX) and argues the difference is absorbed by the E4M3 grid: it is "the one arithmetic step I match by outcome rather than by instruction." The other is the NaN code, which native emits for a zero row and the port replaces with +0, in an intermediate byte no fixture captures.
• A partial CPU reference. dlss5vk verify bisects block 0 kernel by kernel against src/reference.cpp, but "The expert FFN, the 512 split block and the ViT have no CPU reference". They are checked only by two GPU routes agreeing with native's captures.

The larger edge is that none of it can be reproduced from the repository alone. The fixtures are "recorded captures of the original, which are not part of this repository", and the weights are not either. What the repository does ship is internal consistency: a WebGPU self-test that checks a JavaScript oracle and the WGSL against web/fixtures/numerics.bin, generated from the Vulkan tree's own CPU reference, "exhaustive over all 65 536 half bit patterns wherever the domain allows it", and a numpy test of the PTX divider over every n < 2^24. Those prove the implementations agree with each other. Agreement with NVIDIA rests on captures only the author has.

The WebGPU port is the strongest evidence that the specification, not the hardware, carries the exactness. With no tensor core, no FP8 type and no inline assembly, it writes the k32 step out as fixed point, places every half rounding by hand on f32 bit patterns, and builds powers of two from the exponent field because WGSL gives exp2 an accuracy allowance. Its README reports it passing both fixtures it runs with the same verdicts as the Vulkan tool: 75 bit-exact boundaries and the 8-bit capture within one code on one, the head and composed image bit-exact on the other. It costs about 73 ms a frame at 512x512 across 451 dispatches (the top-level README says 72 ms), against 2.7 ms for the Vulkan route. It also records a bug worth any verification engineer's time: the port followed a variable named adapter and skipped from the wrong tensor, and every block boundary stayed exact because the mistake sat downstream of the last one. Only the fixture that checks the head caught it. The same discipline shows up in auto-gpu-kernel, and its absence in Lexing on the GPU: correctness has to be gated somewhere a mistake can actually reach.

What 7.8 ms measures

The number in the post is in the README, and docs/execution.md has the full table (Reported, RTX 4070 SUPER, minimum over 40 frames):

Reading runBench in src/main.cpp says exactly what that is. The input features are synthetic, roundF16(sin(i * 0.0017) * 0.125), uploaded once. One warm-up frame compiles everything; then each frame writes a top-of-pipe timestamp, records the graph, writes a bottom-of-pipe timestamp and submits. So the number is GPU time for the network's 241 dispatches and nothing else: not the game's render, not the feature preprocess, not the composite. The tool defaults to 10 frames; the tables say 40. The minimum is the published figure because the card "alternates between two clock states under sustained load", with medians 3-5% higher.

The FLOP column explains the shape. From 512x512 to 768x768 the arithmetic grows 2.1x and the time 1.04x: nothing is saturated, and the 241 launches plus stages too small to fill 56 SMs set the floor. From 1080p to 4K the arithmetic grows 3.83x and the time 3.77x, compute-bound at 133-137 TFLOP/s, which the docs put at "roughly half the part's dense FP8 tensor rate". The ViT is the outlier: 13% of the FLOPs and about 17% of the time, "because 192 tokens cannot fill the machine" at 768x768.

Reasoned from the table: 7.77 ms is 47% of a 60 Hz frame, spent after the game has rendered a native-resolution frame for the network to re-render. At 1440p the network takes 12.6 ms, three quarters of that frame. At 4K the network alone, at 29.3 ms, takes most of a 30 Hz frame on this card.

What NVIDIA's own implementation costs is not in the repository, and I could not find it either. The project page says DLSS 5 enables "real-time rendering at up to 4K resolution" on GeForce RTX 50 Series GPUs and gives no millisecond figure; the technical report it links returned HTTP 403 to my fetch. So there is no like-for-like comparison here: the port's timing is on an RTX 40-series card NVIDIA does not list, and NVIDIA's is not published where I could read it.

Provenance and licence

The licence is MIT, "for everything in this repository". The README states what the repository is not:

This project is not affiliated with, endorsed by, or supported by NVIDIA. It contains no NVIDIA software, weights, headers, or instructions for obtaining them.

The NOTICE adds that "DLSS" is an NVIDIA trademark "used here only to describe what the network implemented by this code is compatible with". The weights are the user's to supply as a model directory: a manifest.json of eleven stage files, each with a SHA-256, and 153 tensor records. The README says "Nothing in this repository produces such a directory", and I found nothing in the tree that does. The graph is pinned to one model, "the 71-block network of 310.8.0", and refuses a model with a different block count.

The post says the network was reverse engineered. The repository does not describe how the architecture was recovered or how the captures were recorded, and I am not going to guess. What I can say is what is checkable: the code, the docs and the licence are all in the tree, and the claim of exactness rests on data that is not.

What it adds up to

The post compresses three careful statements into a tweet. "Neural rendering" is a same-resolution generative re-render, not an upscaler. "Same input, same output" is byte equality at every block boundary, on single frames without history, against captures only the author holds. "7.8ms per 1080P" is the network alone on a padded field, as a minimum over 40 frames.

The docs are more precise than the post at every turn, and the repository's best idea is the one its first design note states: a quantized network's identity is the order of its roundings, so a port is right when it reproduces that order and wrong otherwise, however good the frame looks. Quantized inference code usually treats rounding as an error to be bounded. This treats it as the specification, writes it down, and builds two independent implementations that agree to the bit.

Everything above was read from maanHimself/OpenDLSS-NR at commit 9d08f41, MIT, cloned and not executed: README.md, NOTICE and docs/ for the claims; src/nr_graph.cpp for the graph and the padded field; src/reference.cpp and shaders/common.glsl for the arithmetic; scripts/ptx/swin.py, scripts/ptx/ptxgen.py and scripts/build_shaders.ps1 for the PTX route; src/main.cpp for parity and bench; demo/ and docs/frame.md for the frame around the network; and ports/browser-webgpu/README.md for the WebGPU port. Line counts are wc -l over git ls-files. NVIDIA's description of DLSS 5 is quoted from its project page, fetched on 26 September 2026. The two figures are the repository's own renders, served locally with its NOTICE and licence in public/articles/opendlss-nr/.