dlss5-nr-reverse-engineering / docs /nr-reverse-engineering.md

Update files, manifest and model card (36 files)

2f3db25 verified about 1 month ago

15.5 kB

# 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](nr-numerical-validation.md),
[temporal reconstruction](nr-temporal-reconstruction.md), and the
[texture/noise contract](nr-texture-contract.md).

## What is established

| Finding | Evidence |
| --- | --- |
| Approximately 154 million parameters; transformer; predominantly FP8 with selected FP16 GEMMs | [NVIDIA technical report](https://research.nvidia.com/labs/adlr/files/DLSS5_Report.pdf), §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](nr-numerical-validation.md) |
| Trainable ViT block with recovered custom attention and residual ordering | [continuous reference](../tools/nr_vit_reference.py), math evidence |
| Differentiable temporal pre/post components with explicit sampler/noise inputs | [temporal reference](../tools/nr_temporal_reference.py) |
| D3D12 texture descriptors and exact integer/uniform noise reconstruction | [texture/noise evidence](nr-texture-contract.md) |
| 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.

```mermaid
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:

```text
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:

```bash
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:

```text
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](../tools/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.

```python
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:

```bash
.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:

```bash
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:

```python
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:

```bash
.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](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-fragment-mma-16832)
gives this packed byte address for matrix element `W[N,K]`:

```text
(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.

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

```python
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](nr-numerical-validation.md) now validate the tested FFN
operation. Other operators require their own comparisons.

The [notebook](../tools/nr_workbench.ipynb) 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`](../tools/nr_ffn_reference.py).

## Trainable latent ViT reference

[`NrVitBlockReference`](../tools/nr_vit_reference.py) 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.Parameter`s.

```python
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.

	# 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](nr-numerical-validation.md),
	[temporal reconstruction](nr-temporal-reconstruction.md), and the
	[texture/noise contract](nr-texture-contract.md).

	## What is established

	\| Finding \| Evidence \|
	\| --- \| --- \|
	\| Approximately 154 million parameters; transformer; predominantly FP8 with selected FP16 GEMMs \| [NVIDIA technical report](https://research.nvidia.com/labs/adlr/files/DLSS5_Report.pdf), §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](nr-numerical-validation.md) \|
	\| Trainable ViT block with recovered custom attention and residual ordering \| [continuous reference](../tools/nr_vit_reference.py), math evidence \|
	\| Differentiable temporal pre/post components with explicit sampler/noise inputs \| [temporal reference](../tools/nr_temporal_reference.py) \|
	\| D3D12 texture descriptors and exact integer/uniform noise reconstruction \| [texture/noise evidence](nr-texture-contract.md) \|
	\| 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.

	```mermaid
	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:

	```text
	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:

	```bash
	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:

	```text
	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](../tools/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.

	```python
	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:

	```bash
	.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:

	```bash
	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:

	```python
	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:

	```bash
	.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](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-fragment-mma-16832)
	gives this packed byte address for matrix element `W[N,K]`:

	```text
	(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.

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

	```python
	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](nr-numerical-validation.md) now validate the tested FFN
	operation. Other operators require their own comparisons.

	The [notebook](../tools/nr_workbench.ipynb) 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`](../tools/nr_ffn_reference.py).

	## Trainable latent ViT reference

	[`NrVitBlockReference`](../tools/nr_vit_reference.py) 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.Parameter`s.

	```python
	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.