rtx/SOV_RTX_GAMEPLAN.md

SOV-RTX: Complete the 40% — Agent Research & Build Prompt

Context

Repository: github.com/SNAPKITTYWEST/sov-kernel-monster
Target: rtx/ directory — sovereign GPU runtime, zero-libc, zero-CRT
Language: C (zero-libc), PTX assembly (sm_89), C-- (GHC backend), Fortran 2018

What Already Exists (60%)

rtx/src/cuda/flash_attention.ptx   — PagedAttention + online softmax + tensor core WMMA
                                     rmsnorm_fused, silu_fused — all real PTX sm_89
rtx/windows_rtx/cuda_driver_loader.c — PEB walk, 25 CUDA functions, zero headers
rtx/windows_rtx/main.c             — zero-CRT entry, PEB→kernel32, boots full stack
rtx/windows_rtx/power_handler.c    — suspend/resume/battery power events
rtx/src/c--/scheduler.cmm          — continuous batching FSM: IDLE/PREFILL/GENERATE/SWAP/CHECKPOINT/RESUME
rtx/src/fortran/transformer_kernel.f90 — rope_fused, gqa_attention_paged, kv stubs
rtx/src/loader/gguf.c              — GGUF v3 parser, zero libc, VirtualAlloc/mmap
rtx/include/sov_rtx.h              — SOV_MAX_SEQS=256, SOV_MAX_HEADS=128, KV_BLOCK_SIZE=16

The Missing 40% — Exactly Four Components

---

MISSING 1: gemm.ptx — Matrix Multiply (sm_89 tensor cores)

Why it's the hardest missing piece:
Every linear layer in a transformer is a GEMM: Q/K/V projection, output projection, FFN gate/up/down. Without GEMM, you cannot do any forward pass. The scheduler.cmm already calls sov_cuda_flash_attention — but before that you need: x_proj = Linear(x) — that's GEMM.

What to build:
PTX kernel using mma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32 (tensor core WMMA). This is the same instruction family already in flash_attention.ptx.

Shortcut:
Use the CUTLASS PTX templates as reference — they're Apache 2.0. Specifically CUTLASS include/cutlass/arch/mma_sm89.h shows the exact PTX idiom. No need to copy CUTLASS — just use it as a reference for the instruction pattern.

Signature to implement:

// gemm.ptx: C = A × B + C (M×K @ K×N → M×N)
// grid = (M/16, N/8, 1)   block = (32, 1, 1)  — one warp per output tile
// Uses mma.sync.aligned.m16n8k16 (Ampere/Ada tensor core)
// Inputs: f16 (quantized weights from GGUF), f32 accumulator

Files to create:

• rtx/src/cuda/gemm.ptx — the kernel
• rtx/src/cuda/gemm_dispatch.c — C dispatch: picks tile size, launches grid

---

MISSING 2: kv_allocator.c — Block Allocator

Why it's needed:
sov_rtx.h declares SOV_MAX_SEQS=256, KV_BLOCK_SIZE=16 but there's no implementation of the free-list that manages which KV blocks are allocated to which sequences. kv_allocate_blocks and kv_append_tokens in transformer_kernel.f90 are stubs. The scheduler calls them but they do nothing.

What to build:
A zero-malloc block allocator using a static free-list array. VirtualAlloc one large region upfront (GPU side: cuMemAlloc via driver API).

Data structure:

// Static: SOV_MAX_SEQS * max_blocks_per_seq block table
// Free list: stack of available block IDs (int32 array, top pointer)
// GPU-side: copy block_table to device before every flash_attention launch

Shortcut:
vLLM's PagedAttention paper (Kwon et al. 2023) describes the block table layout exactly. The block_table in flash_attention.ptx already matches it: block_table[seq_id * max_blocks + b]. Build the CPU-side allocator, GPU side just reads block_table.

Files to create:

• rtx/src/kv_allocator.c — free-list + block_table management
• rtx/src/kv_allocator.h — types + API

---

MISSING 3: sampler.c — Token Sampling

Why it's needed:
After the forward pass produces logits (vocab_size floats), you need to:

1. Apply temperature scaling
2. Run top-p (nucleus) or top-k filtering
3. Sample from the resulting distribution
4. Return the next token ID

The scheduler.cmm increments tokens_generated and checks for EOS — but there's no implementation of what produces the next token from logits.

What to build:
A CPU-side sampler that receives logits from GPU memory. No GPU kernel needed — sampling is memory-bandwidth bound not compute bound at batch_size=1, and pulling 32K floats (Llama vocab) from GPU is fast enough.

Shortcut:
llama.cpp's sampling.cpp is MIT licensed. The core top-p algorithm is 40 lines. Use it as direct reference. Zero-libc version: replace stdlib qsort with a simple partial sort (only need top-k candidates, not full sort).

Files to create:

• rtx/src/sampler.c — temperature, top-p, top-k, greedy (zero-libc)
• Declare sov_sample_token in sov_rtx.h

---

MISSING 4: Wiring — cuda_driver_loader.c → actual kernel launches

Why it's needed:
cuda_driver_loader.c resolves 25 CUDA driver API functions from nvcuda.dll via PEB walk. But the actual launch code — loading the PTX module, getting function handles, setting up launch parameters for flash_attention_paged and the new gemm kernel — is missing. There's no cuModuleLoadData, cuModuleGetFunction, cuLaunchKernel wiring.

What to build:
A cuda_kernels.c that:

1. Calls cuModuleLoadData(ptx_blob) to load flash_attention.ptx and gemm.ptx
2. Calls cuModuleGetFunction to get handles for all 5 kernels
3. Implements sov_cuda_flash_attention() and sov_cuda_gemm() as thin wrappers that set up CUlaunchKernel calls with correct grid/block dims

Shortcut:
The PTX blobs can be embedded as C char arrays (xxd -i flash_attention.ptx). No file I/O needed — zero-libc compatible.

Files to create:

• rtx/src/cuda_kernels.c — module load + function dispatch

---

Build Order (strict)

1. kv_allocator.c           ← no dependencies, unblocks everything
2. gemm.ptx + gemm_dispatch  ← depends on kv_allocator for buffer management
3. cuda_kernels.c            ← depends on both PTX files existing
4. sampler.c                 ← independent, but final step in the inference loop

---

Shortcuts and Reference Materials

Gap	Best Reference	License	Effort
GEMM PTX (tensor cores)	CUTLASS sm89 mma patterns	Apache 2.0	~200 lines PTX
KV block allocator	vLLM PagedAttention paper + llama.cpp kv_cache.cpp	Apache 2.0	~150 lines C
Sampler	llama.cpp sampling.cpp	MIT	~100 lines C
CUDA driver wiring	CUDA driver API docs + cuModuleLoadData examples	Public	~200 lines C

---

Invariants to Preserve

The following invariants already exist in the codebase — do NOT break them:

1. Zero-libc: No #include <stdlib.h>, <stdio.h>, <string.h>. Use VirtualAlloc/mmap for memory. Manual memcpy loops or PTX ld/st.

2. Zero-CRT: main.c has no C runtime. Entry point is sov_main(). No malloc. No global constructors.

3. WORM every 64 tokens: scheduler.cmm calls sov_worm_checkpoint(kv_ptr) every 64 tokens. This must remain. New code must not bypass this.

4. Power checkpoint semantics: Any GPU kernel must check power_state global before heavy compute. flash_attention.ptx already does this. gemm.ptx must too.

5. Janet config array: scheduler_janet_array[32] at fixed address. Slots 0-7 are scheduler config. Slots 8-31 are WORM receipt. New code uses janet_get/janet_set from scheduler.cmm — never raw array access.

6. GGUF v3 layout: gguf.c reads weights. The GEMM kernel receives weight pointers from GGUF tensor data directly. No copy/convert step.

7. sm_89 target: All PTX must be .target sm_89. No sm_80, no sm_90.

---

Exact PTX Instruction for GEMM

The tensor core instruction already used in flash_attention.ptx for the warp-reduce is shfl.sync.bfly. The GEMM needs:

// Load f16 tiles from A and B into registers
// A tile: 8x16 f16 = 4 x .b32 registers (each .b32 holds 2 f16)
// B tile: 16x8 f16 = 4 x .b32 registers
// C/D accumulator: 4 x .f32 registers

mma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32
    {%f0, %f1, %f2, %f3},    // D (f32 accumulator out)
    {%a0, %a1, %a2, %a3},    // A matrix (f16, 4 registers = 8 f16 values)
    {%b0, %b1},              // B matrix (f16, 2 registers = 4 f16 values)
    {%f0, %f1, %f2, %f3};   // C (f32 accumulator in)

This computes a 16×8×16 output tile in one instruction. Tile the full GEMM by iterating K dimension in chunks of 16.

---

Agent Research Tasks

For each agent you spawn, give it this context and one specific task:

Agent A — GEMM PTX researcher: "Research CUTLASS sm89 PTX mma.sync patterns for f16 matrix multiply. Find the exact register layout for m16n8k16, ldmatrix instruction for loading f16 tiles from shared memory, and the loop structure for tiling a full GEMM. Output: annotated PTX pseudocode with exact register names."

Agent B — KV allocator researcher: "Research vLLM PagedAttention block table format and llama.cpp kv_cache allocation. Specifically: how does the block_table[seq_id * max_blocks + b] map to physical GPU memory addresses? What is the free-list data structure? Output: C struct definitions and allocate/free pseudocode, zero-malloc."

Agent C — CUDA driver loader researcher: "Research cuModuleLoadData, cuModuleGetFunction, cuLaunchKernel from the CUDA driver API (not CUDA runtime). How do you load a PTX string at runtime? How do you pass pointer parameters to a kernel via cuKernelParamsSetAttribute? Output: complete working example in C, no CUDA headers (use typedefs only)."

Agent D — Zero-libc sampler researcher: "Research top-p nucleus sampling algorithm. Implement it in C with no stdlib: no qsort, no malloc. Given float* logits of size vocab_size, apply temperature, partial sort top-k candidates, compute cumulative probabilities, return sampled token index. Output: complete C function, zero external dependencies."

---

Final Integration Point

When all four components exist, the full inference loop in scheduler.cmm becomes:

IDLE → PREFILL:
  kv_allocator_alloc(seq_id, num_tokens)     ← kv_allocator.c
  gemm(embedding_weights, token_ids, x)      ← gemm.ptx via cuda_kernels.c

GENERATE (per token):
  gemm(q_proj, x, q)                         ← gemm.ptx
  gemm(k_proj, x, k)                         ← gemm.ptx
  gemm(v_proj, x, v)                         ← gemm.ptx
  kv_append_tokens(seq_id, k, v)             ← kv_allocator.c
  flash_attention_paged(q, block_table, out) ← flash_attention.ptx (EXISTS)
  rmsnorm_fused(out, w, out)                 ← EXISTS
  gemm(o_proj, out, x)                       ← gemm.ptx
  silu_fused + gemm(gate_proj + up_proj)     ← gemm.ptx
  gemm(down_proj, x, x)                      ← gemm.ptx
  rmsnorm_fused(x, w, x)                     ← EXISTS
  gemm(lm_head, x, logits)                   ← gemm.ptx
  token = sampler(logits, temperature, top_p) ← sampler.c
  if token == EOS or len >= max: → IDLE
  else: repeat GENERATE
  every 64 tokens: worm_checkpoint()         ← EXISTS

That loop is the complete sovereign LLM inference engine. Everything except the four missing pieces already exists in the repository.