---
license: other
license_name: sovereign-source-license-v2
library_name: custom
tags:
- code
- sovereign-compute
---
# NVIDIA Stack — Reverse-Engineered GPU Compute Stack
[](https://github.com/SNAPKITTYWEST/nvidia-stack/blob/main/LICENSE)
[](https://github.com/SNAPKITTYWEST/nvidia-stack/blob/main/LICENSE-AGPL)
[](https://www.rust-lang.org/)
[](https://www.python.org/)
[](https://developer.nvidia.com/cuda-toolkit)
[](https://rocm.docs.amd.com/)
[](https://github.com/SNAPKITTYWEST)
**⚠️ NOT OPEN SOURCE** — Sovereign corporate product. Commercial use requires a Sovereign Node Key.
---
## Architecture
```mermaid
flowchart TB
subgraph LOGICAL["Logical Specification (Datalog)"]
DL["paged_attention.dl
Souffle Datalog"]
RT["root_table
seq_id -> block_table_ptr"]
BTE["block_table_entry
table_id, block_idx, base, refcount"]
VT["virtual_token
seq_id, token_pos, block_idx, offset"]
SB["swapped_block
CPU fallback path"]
RKV["resolved_kv_address
final physical address"]
end
subgraph PHYSICAL["Physical Implementation (HIP/CUDA)"]
BA["BlockAllocator
Lock-free LIFO free list"]
PAM["PagedAttentionManager
Block table CRUD + swap"]
RV["resolve_kv_address
Fused device function"]
PK["paged_attention_kernel
Attention with paged KV"]
FB["Fragmentation Benchmark
ShareGPT workload"]
end
subgraph HARDWARE["gfx942 Hardware"]
LDS["LDS
Bank conflict avoidance"]
MFMA["MFMA
v_mfma_f32_16x16x16f16"]
MEM["Global Memory
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)
```bash
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)
```bash
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
```bash
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
```bash
# 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
```bash
# 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
```bash
# 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
```bash
# 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
```bash
# 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):
1. **CUDA .so** — fastest; requires compiled `libmamba2.so`
2. **torch.ops** — JIT compile via `torch.utils.cpp_extension.load()`
3. **Pure PyTorch** — reference implementation; numerically identical to CUDA kernel
```python
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
```bash
# 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)
```rust
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::::eye(64);
let z = ndarray::Array1::::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
```bash
# 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:
```cpp
#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
```bash
# 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:
```rust
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:
```tablegen
// 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)
```bash
# 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
(* 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
1. REVERSE-ENGINEERED NVIDIA TENSOR CORE STACK
Complete CuTe → SASS → Hardware chain simulation with MAC unit
counting, pipeline depth modeling, and cycle-accurate timing.
2. 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.
3. LDS BANK CONFLICT PADDING OPTIMIZER
Automated search over row-major padding and XOR swizzle
parameters to eliminate ds_read_b128 bank conflicts.
4. CROSS-VENDOR GPU COMPUTE MODEL
Unified abstraction covering NVIDIA HMMA and AMD MFMA with
hardware-specific lane-to-fragment mappings.
5. 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.
6. 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).
7. 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.
8. 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.
9. 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.
10. 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.
11. 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
```bibtex
@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