docs/nr-reverse-engineering.md

NR architecture and PyTorch inspection workbench

We recovered the deployed model's weight container, named block structure and GPU intermediate code from the verified NVIDIA runtime. The merged PyTorch export contains 851 tensors covering all 153 source records, with unresolved fields preserved as raw bytes. Repacking it reproduces every original payload byte. The first FFN operator matches NVIDIA's original kernel in 24 native comparisons, and a complete latent ViT block supports autograd through a continuous mathematical reference. Whole-model numerical equivalence and fine-tuning remain incomplete. See numerical validation, temporal reconstruction, and the texture/noise contract.

What is established

Finding	Evidence
Approximately 154 million parameters; transformer; predominantly FP8 with selected FP16 GEMMs	NVIDIA technical report, §2.6
71 numbered blocks; hierarchical Swin and ViT encoder–decoder with skip inputs	Native descriptor constructor RVA 0x39780, reached from the shipping descriptor at 0x3d670
153 named packed records, totaling 147,683,778 payload bytes	Exact parsing to EOF; record inventory
Every exported byte round-trips through PyTorch	inspection manifest
Eight FFN-expansion matrices, totaling 33,554,432 E4M3 values, recovered in output/input order	matrix extraction manifest and PTX derivation
Remaining ViT matrices, residual scales and query scales	remaining ViT extraction
Split-Swin512 dense/grouped matrices, score biases, scales and final head	Swin512 extraction
Small-channel fused Swin, pre/post, down/up and decoder weight layouts	fused-stage extraction
851 tensors rebuild all 153 original payloads, including unresolved bytes	merged export manifest
24 native FFN comparisons / 4,751,360 output bytes all exact	native numerical validation
Trainable ViT block with recovered custom attention and residual ordering	continuous reference, math evidence
Differentiable temporal pre/post components with explicit sampler/noise inputs	temporal reference
D3D12 texture descriptors and exact integer/uniform noise reconstruction	texture/noise evidence
One embedded shipping model in this runtime	Registry bounds and preset lookup; static analysis

The technical report link was previously resolved incorrectly. Its correct URL uses the research site's HTML base directory, /labs/adlr/files/. Its SHA-256 is b97e0744375a0c69a808ae22a8847a53e82dfd430518e8b8552439f334646d7c. The published parameter count describes NVIDIA's model; it is not independently derived from this binary's payload length.

Recovered structure

These are native descriptor channel arguments and kernel families. This is a structural reconstruction, not a verified reproduction of the forward math.

flowchart LR
  A["0–4: Swin, C32"] --> B["5–8: Swin, C64"]
  B --> C["9–14: Swin, C128"]
  C --> D["15–22: Swin, C256"]
  D --> E["23–30: split Swin, C512"]
  E --> F["31–38: ViT 1D, C1024"]
  F --> G["39–47: upsample + Swin, C512"]
  G --> H["48–55: Swin, C256"]
  H --> I["56–61: Swin, C128"]
  I --> J["62–65: Swin, C64"]
  J --> K["66–69: Swin, C32"]
  K --> L["70: fused post / RGB"]
  E -. "block 30" .-> G
  D -. "block 22" .-> H
  C -. "block 14" .-> I
  B -. "block 8" .-> J
  A -. "block 4" .-> K
  A -. "block 0" .-> L

The native families identify 1h, 2h, 4h, 8h, and 16h Swin variants. Descriptor loops, call sites, builder RVAs, and skip input/output selectors are recorded in the static report, with disassembly evidence. Names alone are not used to claim a particular attention-window size, tensor permutation, or complete mathematical definition.

The shipping builder rewrites descriptors to use _fp8 kernels. It selects one embedded resource, WEIGHTS_HT, for the model identifier:

CC_Control_History_Blend_Quantize_With_Teacher_honest_tench_2026_07_04_22_30_weights

The local one-model registry does not contradict NVIDIA supplying other models in other SDK packages. Style controls are a scalar conditioning input in the recovered preblock; they should not be described as three embedded checkpoints.

Recovered GPU code

The runtime embeds seven CUDA fatbinaries. Each contains a Zstandard-compressed PTX entry and a compiled CUDA ELF entry. We extracted 223 visible kernel entries across seven PTX modules, targeting sm_120, PTX version 9.4. These include alternative scheduling and precision variants; 223 is not a layer count or a claim that all kernels run on every frame.

The PTX/ELF manifest records exact source offsets and hashes for every extracted artifact. Extraction uses the existing system libzstd and executes no DLL or GPU code:

python3 scripts/extract_nr_ptx.py

For example, cc_vit_1d_ffn_expand_fp8 in module_5.ptx contains mma.sync.aligned.m16n8k32.row.col.f16.e4m3.e4m3.f16 and cvt.rn.satfinite.e4m3x2.f16x2 instructions. This establishes E4M3 arithmetic on that kernel path, including FP16 accumulation instructions, rather than inferring precision from the weight blob's byte histogram. It does not by itself recover the full matrix packing or activation scales.

Container format

All integers below are little-endian. The top-level map is:

u64 total_file_bytes
repeat until exact EOF:
    u64 name_byte_count
    u8  name[name_byte_count]
    u64 record_byte_count
    record:
        u64 record_byte_count_including_header
        u64 payload_byte_count
        u32 device_code
        u8  payload[payload_byte_count]
        u32 field0
        u32 field1
        u64 dimension_count
        u32 dimensions[dimension_count]

The parser follows the native decoders at RVAs 0x44ee0 and 0x44cf0. device_code is not dtype: the runtime rejects device 2 with a CPU/None error. All shipping records have device 1, fields 0/0 and rank 1. field1 participates in storage-width validation, but does not establish the logical types of all values inside a fused payload.

Each flat dimension is half the payload byte count. It would be incorrect to interpret this as a count of FP16 learned weights: fused data contains packed components, and the active GPU kernels use FP8. The loader copies payload bytes directly; no encryption or compression is used on this weight-loading path. GPU code is stored separately, including compressed PTX.

Use it in PyTorch

The local environment uses official PyTorch 2.14.0+cpu, installed from https://download.pytorch.org/whl/cpu. It inherits existing Python dependencies and requires no GPU allocation. The original runtime and resource remain pinned by SHA-256; these tools do not write to the installed game.

An interactive entry point is research/nr_workbench.ipynb. Select the project CPU interpreter in your notebook editor. The tools also run as ordinary Python scripts; no notebook server was installed or launched.

The most complete tensor export is recovered_tensors.pt. It contains 143,831,040 E4M3 values, 1,923,744 FP16 values, 714 FP32 values and 2,394 unresolved/raw bytes. These are extraction counts, including auxiliary values; they are not an independently verified trainable parameter count. The raw fields include reserved gaps and tails, eight unused attention records, and a two-byte blend-scale record whose exact host binding remains untraced. The decoded values are the deployed quantized weights, not NVIDIA's original pre-quantization training checkpoint.

import torch

recovered = torch.load(
    "sdk-private/nr-model-inspection/pytorch/recovered/recovered_tensors.pt",
    map_location="cpu", weights_only=True,
)
print(len(recovered))  # 851

To reproduce extraction and merge from the pinned local source:

.venv-nr-inspect/bin/python scripts/nr_unpack_vit.py
.venv-nr-inspect/bin/python scripts/nr_unpack_vit_remaining.py
.venv-nr-inspect/bin/python scripts/nr_unpack_swin512.py
.venv-nr-inspect/bin/python scripts/nr_unpack_fused_swin.py
.venv-nr-inspect/bin/python scripts/nr_export_recovered.py

The merger validates family hashes, unique source-record coverage, tensor bits, and inverse reconstruction of every original payload from the merged checkpoint. The older raw inspection interface remains useful for tracing packed bytes:

From the workspace root:

python3 scripts/nr_weight_format.py \
  sdk-private/nr-model-inspection/weights_ht.bin \
  --output reports/nr-weight-inventory.json

.venv-nr-inspect/bin/python scripts/nr_torch_inspect.py --export

The existing export is packed_state.pt. Its keys preserve native names such as block31.layer0.layer; all values are CPU torch.uint8 tensors. The module listing groups them by recovered block family. These are registered buffers, with zero trainable parameters; calling forward() raises an explicit error.

Load the tensor dictionary that we produced:

import torch

packed = torch.load(
    "sdk-private/nr-model-inspection/pytorch/packed_state.pt",
    map_location="cpu",
    weights_only=True,
)
print(packed["block31.layer0.layer"].shape, packed["block31.layer0.layer"].dtype)
# torch.Size([4194320]) torch.uint8

For an explicit unscaled candidate interpretation, with byte boundaries chosen by the investigator:

.venv-nr-inspect/bin/python scripts/nr_torch_inspect.py \
  --candidate block31.layer0.layer --interpret-as fp8_e4m3fn \
  --start-byte 0 --byte-count 4194304 \
  --output-dir sdk-private/nr-model-inspection/candidates/ffn31

That command does not establish matrix ordering or scaling. Finite/plausible numbers are not sufficient evidence that a tensor interpretation is correct. Do not pass the vendor blob directly to torch.load; it is not a PyTorch file.

First recovered semantic matrices

Blocks 31–38 each have an FFN-expansion matrix with 4096 output channels and 1024 input channels. Tracing the weight loads in cc_vit_1d_ffn_expand_fp8 against NVIDIA's MMA B-fragment mapping gives this packed byte address for matrix element W[N,K]:

(K//32)*131072 + (N//16)*512 + (N%8)*64 + ((K%16)//4)*16
    + ((N%16)//8)*8 + ((K%32)//16)*4 + (K%4)

We checked the full permutation for one-to-one coverage, independently compared the documented first-tile MMA coordinates and an alternative reshape/transpose implementation, and repacked all eight matrices byte-exactly. The trailing 16 bytes of each payload are preserved in the manifest; their role is unresolved. The E4M3 format is bit-compatible with torch.float8_e4m3fn for this inspection.

.venv-nr-inspect/bin/python scripts/nr_unpack_vit.py

import torch

matrices = torch.load(
    "sdk-private/nr-model-inspection/pytorch/vit_ffn/ffn_expand_e4m3_unscaled.pt",
    map_location="cpu", weights_only=True,
)
weight = matrices["block31.ffn_expand.weight_e4m3_unscaled"]
print(weight.shape, weight.dtype)  # (4096, 1024), torch.float8_e4m3fn
weight_fp32 = weight.float()      # exact decoded values; no extra scaling applied

The extraction alone did not establish numerical equivalence. Subsequent reconstruction recovered the cubic activation and half/FP8 rounding path, and native comparisons now validate the tested FFN operation. Other operators require their own comparisons.

The notebook loads the merged tensor export, reproduces a saved native capture on CPU, and demonstrates finite, nonzero gradients through the smooth FFN using actual decoded weights. The numerical reference and the differentiable relaxation are separate APIs in nr_ffn_reference.py.

Trainable latent ViT reference

NrVitBlockReference loads the actual recovered weights for blocks 31–38. It implements expansion/activation, contraction with scaled residual, normalized Q/K/V, global custom attention, and projection with a second scaled residual. All nine decoded learned tensors are nn.Parameters.

import sys
import torch
sys.path.insert(0, "scripts")
from nr_vit_reference import NrVitBlockReference

torch.set_num_threads(2)
block = NrVitBlockReference.from_recovered(31)
x = torch.randn(64, 1024) * 0.01
y = block(x)
y.square().mean().backward()

Input and output are latent [tokens,1024] tensors, not images. The token count must be a positive multiple of 64; unvalidated native tail behavior is rejected. The notebook checks finite, nonzero gradients for all nine parameter tensors. Query chunking limits temporary forward tensors, but autograd retains attention graphs and training memory still grows quadratically with token count.

The native attention uses a clamped half-bit transformation to produce positive weights, rather than ordinary softmax. The reference derives a continuous interpolation of that transformation. It deliberately relaxes FP8/FP16 rounding and approximate reciprocal/normalization instructions. Native attention also quantizes numerator weights differently from denominator weights. Those differences need separate numerical reproduction; FFN parity does not establish full-block parity. See the equations and evidence.

The workbench also exercises temporal preprocessing/postprocessing, separation of recurrent output from display intensity, and coordinate/counter-based noise. The integer hash and uniform conversion are exact reconstructions of the PTX; CPU Gaussian transcendental math and texture sampling are not yet native-matched.

What remains before fine-tuning

1. Implement and assemble the recovered fused Swin/pre/post operations and cross-stage spatial transforms. All source records now have tensor/raw coverage; reserved fields remain explicitly unresolved.
2. Extend native comparisons beyond FFN expansion to projection, attention, grouped operations, and cross-stage spatial transforms. Recovered attention includes a custom approximation; a stock softmax is not an established substitute.
3. Join the components with native-matched sampling, noise, controls and temporal state, then validate full frames and moving sequences.
4. For whole-model training, choose a deliberate quantization/backpropagation policy and training data/losses after validating the forward. Local smooth operator gradients already work, but do not establish whole-model fidelity. Repacking or replacing the live runtime remains a separate engineering task.

There is no verified full NR forward pass or whole-model training checkpoint in this workbench yet. Recovered matrices can be decoded to floating-point tensors; that does not recover NVIDIA's original pre-quantization training weights.

Validation is recorded in the research report: CPU parser, layout, math, gradient and upload-ordering tests pass. All twelve current notebook code cells execute on CPU, and independent checks match the exported payloads and recovered matrices against the original resource. Original source hashes remain pinned; the native harness executes separate extracted code.