NVIDIA Stack — Reverse-Engineered GPU Compute Stack

⚠️ NOT OPEN SOURCE — Sovereign corporate product. Commercial use requires a Sovereign Node Key.

Architecture

flowchart TB
    subgraph LOGICAL["Logical Specification (Datalog)"]
        DL["paged_attention.dl<br/>Souffle Datalog"]
        RT["root_table<br/>seq_id -> block_table_ptr"]
        BTE["block_table_entry<br/>table_id, block_idx, base, refcount"]
        VT["virtual_token<br/>seq_id, token_pos, block_idx, offset"]
        SB["swapped_block<br/>CPU fallback path"]
        RKV["resolved_kv_address<br/>final physical address"]
    end

    subgraph PHYSICAL["Physical Implementation (HIP/CUDA)"]
        BA["BlockAllocator<br/>Lock-free LIFO free list"]
        PAM["PagedAttentionManager<br/>Block table CRUD + swap"]
        RV["resolve_kv_address<br/>Fused device function"]
        PK["paged_attention_kernel<br/>Attention with paged KV"]
        FB["Fragmentation Benchmark<br/>ShareGPT workload"]
    end

    subgraph HARDWARE["gfx942 Hardware"]
        LDS["LDS<br/>Bank conflict avoidance"]
        MFMA["MFMA<br/>v_mfma_f32_16x16x16f16"]
        MEM["Global Memory<br/>Paged KV cache blocks"]
    end

    RT --> BA
    BTE --> PAM
    VT --> RV
    SB --> PAM
    RKV --> RV
    BA --> PAM
    PAM --> PK
    RV --> PK
    FB --> PAM
    PK --> LDS
    LDS --> MFMA
    MFMA --> MEM
    MEM --> BTE

What This Is

A complete reverse-engineered GPU compute stack covering the full chain from high-level tensor operations down to hardware cycles:

PyTorch/CuTe Layouts → PTX/SASS ISA → Tensor Core/MFMA Microarchitecture → Hardware Signals

Coverage

Layer	NVIDIA	AMD	x86-64	Quantum
Tensor Layout	CuTe layouts (Rust)	A/B row-major / column-major (Python)	—	—
Instruction Set	SASS HMMA/LDG/STG (Rust)	AMDGPU MFMA ISA (asm)	AVX2 FMA (NASM)	QIR intrinsics
Microarchitecture	Tensor Core MAC simulation (Rust)	Matrix Core wave simulation	OoO core scheduling model	Linear type verifier
Memory	Global/L1/L2 cache model	LDS bank conflict avoidance + XOR swizzle	Cache-blocked GEMV	—
KV Cache	—	PagedAttention block table manager (HIP/CUDA)	—	—
Logical Spec	—	Datalog/Souffle PagedAttention schema	—	#q dialect (MLIR TableGen)
SSM Backbone	Mamba-2 SSD selective scan (CUDA)	Mamba-2 SSD selective scan (HIP)	—	—
Waveform Synthesis	—	—	LW-LGM latent-to-waveform (Rust/NASM)	—
FSL Dialect	Mamba-2 SSM state transition (C++)	Selective SSM with SiLU gating (C++)	—	FSM + continuous hybrid semantics
Quantum Circuits	—	—	—	Rust-Q + QIR lowering (Rust)
MFMA Core	OCaml→C→HLS pipeline	HIP gfx942 kernel	CUDA SM_86 WMMA	—
High-Level API	—	HIP/rocwmma GEMM (fragment loads, mfma_sync)	—	Circuit builder
Validation	—	Fragment map validator + structural checks	Linearity + energy tests	No-cloning + angle domain
Layout Search	—	Padding + XOR swizzle optimizer	—	Clifford+T rewrite patterns
Assembly	—	gfx942 MFMA GEMM kernels	x86-64 AVX2 GEMV kernel	—

Repository Structure

nvidia-stack/
├── src/
│   └── main.rs                          Rust NVIDIA stack simulator
│       ├── CuTe Layouts                 Tensor-to-memory coordinate mapping
│       ├── SASS ISA                     HMMA/LDG/STG instruction model
│       ├── Tensor Core Hardware          MAC units, pipeline, clock simulation
│       └── Stack Orchestrator            Full chain execution + timing
├── asm/
│   ├── mfma_f16_16x16x16.s             AMDGPU MFMA basic tile (gfx90a)
│   ├── mfma_lds_staging.s              gfx942 MFMA with LDS ping-pong staging
│   └── mfma_lds_xor_swizzle.s          gfx942 MFMA with XOR swizzle bank conflict avoidance
├── datalog/
│   └── paged_attention.dl              Souffle Datalog: PagedAttention KV cache logical spec
│       ├── Schema Declarations          root_table, block_table_entry, virtual_token
│       ├── Integrity Constraints        Alignment, bounds, refcount checks
│       ├── Core Rules                   resolved_kv_address (GPU + CPU swap paths)
│       └── Test Dataset                 Multi-sequence block sharing, swap demo
├── hip/
│   ├── gemm_kernel.cpp                 HIP/rocwmma GEMM (16x16 MFMA, multi-wave, shared memory)
│   └── paged_attention.cu              PagedAttention block manager + fused attention kernel
│       ├── BlockAllocator              Lock-free free list (LIFO, atomic ops)
│       ├── PagedAttentionManager       Block table CRUD, prefix caching, swap logic
│       ├── resolve_kv_address          Fused device function (matches Datalog rules)
│       ├── paged_attention_kernel      Attention with paged KV cache reads
│       └── Fragmentation Benchmark     ShareGPT workload validation
├── kernels/
│   ├── mamba2_torch.py                  PyTorch Mamba-2 SSD module (pure-PyTorch + CUDA dispatch)
│   ├── mamba2.cu                        Mamba-2 SSD CUDA kernel (sm_86/sm_89+, fp8 quantisation)
│   └── build_mamba2.py                 Build libmamba2.so (nvcc compile + link)
├── waveforms/
│   ├── Cargo.toml                      lw-lgm package (ndarray + rand)
│   ├── src/
│   │   ├── lib.rs                      build_dictionary + latent_to_waveform (Rust)
│   │   └── main.rs                     CLI demo
│   ├── latent_to_waveform_nasm.asm     x86-64 AVX2 GEMV kernel (NASM)
│   └── lw_lgm.py                       Python reference implementation + validation
├── fsl/
│   ├── include/
│   │   ├── FSLTypes.td                 MLIR TableGen: statevector, tokenvector, ssmmatrices types
│   │   └── FSLOps.td                   MLIR TableGen: mamba_step, selective_mamba_step, output_projection ops
│   └── kernels/
│       ├── fsl_mamba_step.cpp          Basic SSM state transition kernel (C)
│       ├── fsl_selective_mamba_step.cpp Selective Mamba-2 SSM kernel with SiLU gating (C)
│       └── fsl_mamba_test.cpp          Unit tests for FSL kernels
├── quantum/
│   ├── include/
│   │   ├── QuantumTypes.td             MLIR TableGen: qubit, qureg, pauli types
│   │   └── QuantumOps.td              MLIR TableGen: alloc, unitary, entangle, measure ops
│   ├── lib/
│   │   ├── QuantumVerifier.cpp         Linear-type verifier (no-cloning, bounds, angles)
│   │   └── QuantumRewritePatterns.cpp  Algebraic rewrites (H²=I, T³=S², Rz merge)
│   └── rustq/
│       ├── Cargo.toml                  rustq crate (zero dependencies)
│       └── src/
│           └── lib.rs                  Circuit builder + QIR lowering (Rust)
├── mfma-core/
│   ├── src/
│   │   ├── mfma_core.ml               OCaml algorithm specification
│   │   ├── mfma_hls_wrapper.c         HLS-compatible C wrapper
│   │   ├── mfma_core.h                Public C interface
│   │   ├── mfma_core_hip.cpp          AMD gfx942 HIP kernel
│   │   └── mfma_core.cu               NVIDIA RTX 3080 CUDA kernel
│   ├── rtl/
│   │   └── fpga_mfma_accelerator.sv   SystemVerilog FPGA implementation
│   ├── analog/
│   │   └── mfma_power_supply_droop.vams  Verilog-A power/droop model
│   ├── formal/
│   │   └── mfma_nan.why               Why3 NaN propagation proof
│   ├── fpga/scripts/                   Vivado flow scripts
│   ├── asic/scripts/                   Synopsys DC + PrimeTime + KLayout
│   ├── Makefile                        Master build pipeline
│   └── README.md                       MFMA Core documentation
├── python/
│   ├── fragment_map.py                  Opcode-accurate fragment map + layout search
│   ├── structural_validator.py          Bijectivity, per-lane, VGPR, C/D checks
│   └── lds_padding.py                   ds_read_b128 padding calculator
├── LICENSE                              Business Source License 1.1
├── LICENSE-AGPL                         GNU AGPL v3.0
└── README.md                            This file

Quick Start

Rust (NVIDIA Stack Simulator)

cd nvidia-stack
cargo run

Output:

--- Starting Stack Execution ---
[Stack] Layouts Generated: A([16, 16], [16, 1]), B([16, 16], [16, 1])
[HW] Memory Load (L1/L2 Cache Hit)
[HW] Memory Load (L1/L2 Cache Hit)
[HW] Executing HMMA 16x16x16 | Cycles: 1.00 | Latency: 6.19ns
[HW] Memory Store
--- Stack Execution Complete ---
Total Wall-Clock Time (Simulated): 36.1905 ns

Python (Fragment Map + Layout Optimizer)

cd python
python fragment_map.py

Output:

Fragment map validation passed.

=== Operand A (row-major) ===
Layout: padded
 Padding: 0 FP16 elements
 Row stride: 16 FP16 elements
 = 32 bytes

=== Operand B (column-major) ===
Layout: padded
 Padding: 0 FP16 elements
 Column stride: 16 FP16 elements
 = 32 bytes

=== Layout Certificate ===
{
  "target": "gfx942",
  "opcode": "v_mfma_f32_16x16x16f16",
  "wavefront_size": 64,
  "mfma_tile": {"M": 16, "N": 16, "K": 16},
  "operand_A": {
    "load": "ds_read_b64",
    "conflicts": []
  },
  "operand_B": {
    "load": "ds_read_b64",
    "conflicts": []
  }
}

Structural Validator

cd python
python structural_validator.py

Validates:

Element count (256 A, 256 B, 256 C, 256 D)
Coordinate bijectivity (no duplicates, no missing)
Per-lane occupancy (4 FP16 A, 4 FP16 B, 4 FP32 C per lane)
Packed FP16 register pairs (one low, one high per VGPR)
C/D accumulator correspondence

AMDGPU Assembly

# Assemble for gfx942
llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx942 -filetype=obj asm/mfma_lds_xor_swizzle.s -o mfma.o

# Assemble for gfx90a
llvm-mc -triple=amdgcn-amd-amdhsa -mcpu=gfx90a -filetype=obj asm/mfma_f16_16x16x16.s -o mfma_basic.o

HIP/rocwmma GEMM

# Compile for gfx942
hipcc -std=c++17 -offload-arch=gfx942 hip/gemm_kernel.cpp -o gemm -lrocwmma

# Run
./gemm

Features:

16x16x16 MFMA tiles via rocwmma fragments
Multi-wave execution (4 waves per block, 256 threads)
Shared memory staging for A/B tiles
Bounds-safe zero-padding for non-multiple dimensions
FP16 inputs, FP32 accumulation
NaN propagation per IEEE-754 FMA rules

PagedAttention KV Cache Manager

# Compile for gfx942
hipcc -std=c++17 -offload-arch=gfx942 -O3 hip/paged_attention.cu -o paged_attention

# Run (runs built-in fragmentation benchmark)
./paged_attention

Features:

Lock-free block allocator (LIFO free list, atomic ops)
Atomic 16-bit reference counting (prefix caching / beam search)
Fused resolve_kv_address device function (no indirection overhead)
Swap logic for GPU memory pressure (CPU fallback path)
Fragmentation benchmark: ShareGPT workload (50% short / 30% medium / 20% long)
Matches Datalog schema: root_table, block_table_entry, virtual_token

Datalog PagedAttention Schema

# Run with Souffle
cd datalog
souffle paged_attention.dl -F . -D .

# Output: resolved_kv_address.csv
cat resolved_kv_address.csv

Logical specification:

root_table(seq_id, block_table_ptr) -- sequence -> block table pointer
block_table_entry(table_id, block_idx, base_addr, refcount) -- physical block mapping
virtual_token(seq_id, token_pos, block_idx, offset) -- position decomposition
swapped_block(table_id, block_idx, cpu_addr) -- CPU-resident fallback
resolved_kv_address(seq_id, token_pos, phys_addr) -- final KV cache address

Constraints enforced:

256-byte alignment (Base mod 256 == 0)
Offset bounds (0 <= Offset < 256)
Non-negative refcount

Mamba-2 SSD Selective Scan

# Pure PyTorch (no nvcc required, runs on RTX 3080)
cd kernels
python mamba2_torch.py

# Build CUDA extension (requires nvcc on bbqbaddie)
python build_mamba2.py --arch sm_86   # RTX 3080
python build_mamba2.py --arch sm_89   # RTX 5000 Ada

Three execution modes (auto-selected):

CUDA .so — fastest; requires compiled libmamba2.so
torch.ops — JIT compile via torch.utils.cpp_extension.load()
Pure PyTorch — reference implementation; numerically identical to CUDA kernel

from kernels.mamba2_torch import Mamba2Layer, Mamba2Block, Mamba2Model

# Single layer
layer = Mamba2Layer(d_model=512, d_state=16, d_conv=4)
x = torch.randn(2, 128, 512)          # [B, L, D]
y, h = layer(x)                        # y: [B, L, D], h: [B, D, N] state

# Autoregressive step
x_step = torch.randn(2, 1, 512)
y_step, h = layer(x_step, recurrent_state=h)

# Full model (stack of Mamba-2 blocks)
model = Mamba2Model(d_model=512, n_layers=4, vocab_size=512)
tokens = torch.randint(0, 512, (2, 128))
out, states = model(tokens)             # out: [2, 128, 512]

Features:

Mamba-2 SSD (Structured State-Space Duality) selective scan
Causal depthwise conv with cache for autoregressive inference
Recurrent state carry: (ssm_h, conv_cache) per layer
FP8 quantisation in CUDA kernel (simulated on sm_86, native on sm_89+)
Chunk-parallel SSD kernel for long sequences
Haskell FFI: mamba2_step_fp8() / mamba2_forward_fp8()

LW-LGM Latent-to-Waveform Synthesis

# Rust (recommended)
cd waveforms
cargo run

# Python reference
cd waveforms
python lw_lgm.py

# NASM assembly kernel
nasm -f elf64 -o latent_to_waveform_nasm.o latent_to_waveform_nasm.asm

Mathematical construction:

Mother waveform: φ(t) = Gaussian(σ₀)
Dictionary atoms: ψ_i(t) = (1/√|a_i|) φ((t - b_i)/a_i)
Affine grid: Logarithmic dilation + uniform translation
Mapping: x(t) = z^T W^T Ψ(t) (linear expansion in fixed dictionary)

use lw_lgm::{build_dictionary, latent_to_waveform};

let psi = build_dictionary(1.0, 0.5, 2.0, -5.0, 5.0, 64, -10.0, 10.0, 0.01);
let W = ndarray::Array2::<f64>::eye(64);
let z = ndarray::Array1::<f64>::random(64, rand::distributions::Uniform::new(-1.0, 1.0));
let x = latent_to_waveform(&z, &W, &psi);  // x ∈ ℝ^N

Features:

Linearity: L(αz₁ + βz₂) = αL(z₁) + βL(z₂)
Frame expansion in L^2(ℝ) with affine dictionary
Energy preservation via tight frame design
AVX2 FMA kernel with cache-blocking for large matrices
Python reference with linearity + energy validation tests

FSL Dialect — Mamba Step Kernels

# Compile and run FSL kernel tests
cd fsl/kernels
g++ -O2 -o fsl_test fsl_mamba_step.cpp fsl_selective_mamba_step.cpp fsl_mamba_test.cpp
./fsl_test

Hybrid continuous-discrete semantics for Mamba-2 SSM:

#include "fsl_mamba_step.cpp"

// Basic Mamba step: s_{t+1} = A * s_t + B * u_t
float state[16], input[512], A[16*16], B[16*512], next_state[16], output[512];
fsl_mamba_step(state, input, A, B, next_state, output, 16, 512);

// Selective Mamba-2 step with SiLU gating
float A_log[16], W_conv[512*4];
fsl_selective_mamba_step(state, input, A_log, B, W_conv,
                         next_state, output, 16, 512, 4);

// FSM transition (discrete state)
int new_state = fsl_fsm_transition(0, 1, condition_flag);

// Scan complete check
int done = fsl_scan_complete(next_state, 16, 1e-6f);

Features:

Basic SSM: s_{t+1} = A * s_t + B * u_t (fixed A, B)
Selective SSM: depthwise conv + SiLU gating + SSM update
FSM semantics: discrete state transitions gated by conditions
YAML-configured parameters (d_state=16, d_model=512, d_conv=4)
MLIR TableGen ops: fsl.mamba_step, fsl.selective_mamba_step
Hybrid continuous-discrete: SSM state evolves continuously, FSM gates actions

Quantum Dialect (#q) + Rust-Q

# Rust-Q circuit builder + QIR lowering
cd quantum/rustq
cargo test

# MLIR dialect (requires LLVM/MLIR build)
cd quantum
mlir-tblgen --gen-op-decls include/QuantumOps.td -I include/
mlir-tblgen --gen-op-defs include/QuantumOps.td -I include/

Linear-type quantum IR with no-cloning enforcement:

use rustq::{Circuit, QirLowering, ControlOperand};

let mut c = Circuit::new();
let q0 = c.alloca_qubit();   // !quantum.qubit (linear resource)
let q1 = c.alloca_qubit();

c.h(q0);                      // H gate (no controls)
c.cx(q0, q1);                 // CNOT (controlled-X)

// Controlled gate with register as control
let reg = c.alloca_veq(3);
c.controlled("h", vec![ControlOperand::Veq(reg)], vec![q1], vec![], false);

let r0 = c.mz(q0);           // Measurement → i1
let r1 = c.mz(q1);

let qir = QirLowering::lower(&c);  // → __quantum__qis__* calls

MLIR TableGen definitions:

// Linear qubit type (no cloning)
!quantum.qubit

// Unitary with exact algebraic angles
quantum.unitary %q [0.5] axis "Y" : (!quantum.qubit) -> !quantum.qubit

// Controlled operation
quantum.entangle [%c0, %c1] %t : (!quantum.qubit, !quantum.qubit) -> ...

// Measurement
quantum.measure %q -> "c" : (!quantum.qubit) -> (i1, !quantum.qubit)

Features:

Linear-type enforcement: every qubit has exactly one use
Exact algebraic angles (rational, not floating-point)
Controlled gates: single Veq, multi-qubit, multi-target
QIR lowering: __quantum__qis__* / __quantum__rt__* symbols
Algebraic rewrites: H²=I, T³=S², Rz(a)+Rz(b)=Rz(a+b)
No-cloning verifier + bounds checking + angle domain validation

MFMA Core (OCaml → C → HLS → RTL → FPGA/ASIC)

# Build HLS library (OCaml → C → .so)
cd mfma-core
make all

# Build HIP kernel (AMD gfx942)
make hip

# Build CUDA kernel (NVIDIA RTX 3080)
make cuda

# FPGA synthesis (AMD Vivado)
make fpga

# ASIC synthesis (Synopsys DC + PrimeTime)
make asic

Complete hardware design flow for 16x16x16 FP16 → FP32 MFMA tile:

(* OCaml algorithm specification *)
let mfma_tile a_tile b_tile c_tile =
  Array.init 16 (fun m ->
    Array.init 16 (fun n ->
      let acc = ref (Array.get c_tile m n) in
      for k = 0 to 15 do
        let va = half_to_float a_tile.(m * 16 + k) in
        let vb = half_to_float b_tile.(k * 16 + n) in
        acc := !acc +. (va *. vb)
      done;
      !acc
    )
  )

Features:

OCaml → C: ocamlopt -output-obj with zero runtime in HLS region
HLS Pragmas: PIPELINE II=1, UNROLL, m_axi interface binding
NaN Propagation: IEEE-754 compliant, verified in Why3 (zero sorries)
HIP kernel: Maps to v_mfma_f32_16x16x16f16 on gfx942
CUDA kernel: Uses wmma::mma_sync on SM_86 Tensor Cores
FPGA: SystemVerilog RTL, Vivado flow for Alveo U55C/U250
ASIC: Synopsys DC + PrimeTime STA, GDSII tape-out ready
Formal: Why3 proof of NaN safety (mfma_nan.why)

Fragment Map (v_mfma_f32_16x16x16f16)

The canonical lane-to-fragment mapping for gfx942:

A Operand (M×K = 16×16 FP16)

m = lane >> 2 (row, 0..15)
k0 = (lane & 0x3) << 2 (column start, step 4)
4 FP16 elements per lane → 2 packed VGPRs (v4, v5)

B Operand (K×N = 16×16 FP16)

k0 = (lane >> 4) << 2 (row start, step 4)
n = lane & 0xF (column, 0..15)
4 FP16 elements per lane → 2 packed VGPRs (v8, v9)

C/D Operand (M×N = 16×16 FP32)

n = lane & 0xF (column, 0..15)
m0 = lane >> 4 (row start, step 4)
4 FP32 elements per lane → 4 accumulator VGPRs (v0, v1, v2, v3)

LDS Bank Conflict Avoidance

ds_read_b128 Lane Groups (gfx942)

G0: lanes 0-3 + 20-23    G4: lanes 32-35 + 52-55
G1: lanes 4-7 + 16-19    G5: lanes 36-39 + 48-51
G2: lanes 8-11 + 28-31   G6: lanes 40-43 + 60-63
G3: lanes 12-15 + 24-27  G7: lanes 44-47 + 56-59

XOR Swizzle Formula

physical_col_word = logical_col_word XOR (row >> row_shift) << xor_shift

Eliminates bank conflicts without increasing LDS consumption.

Protected Inventions

REVERSE-ENGINEERED NVIDIA TENSOR CORE STACK Complete CuTe → SASS → Hardware chain simulation with MAC unit counting, pipeline depth modeling, and cycle-accurate timing.
AMD MFMA FRAGMENT MAP VALIDATOR Structural validation proving bijection, per-lane occupancy, packed FP16 register pairs, and C/D accumulator correspondence for v_mfma_f32_16x16x16f16.
LDS BANK CONFLICT PADDING OPTIMIZER Automated search over row-major padding and XOR swizzle parameters to eliminate ds_read_b128 bank conflicts.
CROSS-VENDOR GPU COMPUTE MODEL Unified abstraction covering NVIDIA HMMA and AMD MFMA with hardware-specific lane-to-fragment mappings.
PAGEDATTENTION LOGICAL SPECIFICATION (DATALOG) Formal Datalog schema for PagedAttention KV cache address translation with integrity constraints, block sharing, and CPU swap fallback paths. Proves zero fragmentation via fixed-size block indirection.
LOCK-FREE PAGED BLOCK MANAGER (HIP/CUDA) Production-ready block allocator with atomic reference counting for prefix caching, fused address translation in attention kernels, and ShareGPT-validated fragmentation benchmarks (<5% vs 40-60% contiguous).
MAMBA-2 SSD SELECTIVE SCAN (CUDA/PYTORCH) Sovereign Mamba-2 implementation with fp8 quantisation, chunk-parallel SSD kernel, recurrent state carry for autoregressive inference, and Haskell FFI for BOB Architecture integration. Numerically equivalent CUDA and pure-PyTorch paths.
LW-LGM LATENT-TO-WAVEFORM LINEAR GEOMETRIC MAP Explicit construction of analog waveforms from latent vectors via affine group action on a mother Gaussian, with frame-theoretic energy bounds, AVX2 FMA assembly kernel, and cache-blocked GEMV for large dictionary matrices.
LINEAR-TYPE QUANTUM DIALECT (#q) + RUST-Q Strict linear-type refinement of CUDA-Q Quake with no-cloning enforcement at the type level, exact algebraic angles (rational, not floating-point), and explicit QIR lowering to quantum__qis* / quantum__rt* symbols. Includes algebraic rewrite patterns (H²=I, T³=S², Rz merge) and multi-target controlled-gate support.
FSL DIALECT — HYBRID CONTINUOUS-DISCRETE MAMBA-2 Hand-rolled C kernels implementing the Mamba-2 selective SSM with FSM hybrid semantics. Basic and selective variants with depthwise convolution, SiLU gating, and discrete state transitions. MLIR TableGen ops for compiler integration.
MFMA CORE — OCAML-TO-SILICON HARDWARE DESIGN FLOW Complete OCaml → C → HLS → RTL → FPGA/ASIC pipeline for 16x16x16 FP16 → FP32 MFMA tile computation. Includes HIP (gfx942), CUDA (SM_86), SystemVerilog FPGA, Verilog-A analog model, Why3 NaN propagation proof, and GDSII tape-out scripts for TSMC N6.

License

⚠️ THIS IS NOT OPEN SOURCE

This project is a sovereign corporate product licensed under Business Source License 1.1 (BSL-1.1) with GNU AGPL v3.0 copyleft for network services.

Component	License	File	Scope
Core Stack & Simulators	BSL-1.1	`LICENSE`	Rust simulator, Python validators
API/Network	GNU AGPL v3.0	`LICENSE-AGPL`	Any network service exposure

Citation

@misc{nvidiastack2026,
  title={NVIDIA Stack: Reverse-Engineered GPU Compute Stack},
  author={Ahmad Ali Parr and Jessica Westerhoff},
  year={2026},
  note={CuTe/SASS/MFMA simulator, PagedAttention, Mamba-2 SSD, LW-LGM, FSL dialect, #q quantum dialect, MFMA Core},
  publisher={SNAPKITTYWEST},
  howpublished={\url{https://github.com/SNAPKITTYWEST/nvidia-stack}},
  license={BSL-1.1}
}

Contact

Ahmad Ali Parr - ahmedparr93@gmail.com Jessica Westerhoff - jessicalw34@gmail.com

Bel Esprit d'Accord Trust — 50/50 equal sovereigns

Downloads last month: -; Downloads are not tracked for this model. How to track

Inference Providers NEW

This model isn't deployed by any Inference Provider. 🙋 Ask for provider support