macrogrok-harness

Rust harness for MACROGROK fixed-point macro-models and Transformer-X CUDA (Tensor Core WMMA, attention memory, Tavily RAG).

Architecture

• Rust host (tokio, cudarc 0.11, reqwest, half) + CUDA cudarc / custom kernels
• Attention — Tensor Core WMMA 16x16x16 (mma.sync path available), explicit AttentionMemory KV-cache (contiguous/paged), FlashAttention online-softmax tiling
• MACROGROK — Q1.14/Q3.12 fixed-point INFER4 harness (src/macrogrok/)
• Tavily — async HTTP RAG (POST https://api.tavily.com/search, search_depth=advanced)
• X algorithm — sparse/hybrid attention hook inside attention_tc.cu (insert sparsity mask / linear state / IndexPool before/after WMMA)

Layout

Cargo.toml
src/
  main.rs              — tokio main, Transformer load, AttentionMemory, Tavily loop
  memory.rs            — AttentionMemory {k_cache, v_cache} [layers, heads, max_seq, head_dim]
  attention.rs         — TensorCoreAttention {forward()} (WMMA)
  transformer.rs       — TransformerConfig / Block / forward
  tavily.rs            — TavilyClient::search()
  macrogrok/mod.rs     — Q-format INFER4 Rust port (mirrors MACROGROK/src/sim.py)
cuda/
  attention_tc.cu      — WMMA kernel (QK^T + softmax + PV, head_dim tiling, smem)
  tma_mbarrier.ptx     — Hopper TMA + mbarrier bulk tensor (2D shared::cluster)
src/macrogrok/tests    — fixed-point vectors

Build

# Cargo.toml excerpt
[dependencies]
cudarc = { version = "0.11", features = ["cuda-12000"] }
anyhow = "1"
tokio = { version = "1", features = ["full"] }
reqwest = { version = "0.12", features = ["json"] }
serde = { version = "1", features = ["derive"] }
half = "2"
bytemuck = "1"

cargo build --release
cargo test
# CUDA kernels: nvcc -arch=sm_80 -ptx cuda/attention_tc.cu -o attention_tc.ptx
# Run
TAVILY_API_KEY=... cargo run --release -- --prompt "hello"
python -m pytest # if Python harness tests

MACROGROK Harness

use macrogrok_harness::macrogrok::{infer4, Q1_14};
let input = [Q1_14(16384), Q1_14(0), Q1_14(0), Q1_14(0)];
let (output, score, flags) = infer4(input); // Q1.14, Q3.12, flags

Matches MACROGROK/examples/infer4.asm semantics: Q1.14*Q1.14>>4 accum Q3.24 → >>12 → +BIAS_Q3_12 → SAT → THRESHOLD → UPDATE_STATE_3_4 (3*STATE+TARGET)/4.

Transformer-X Loop

tokenize → Transformer::forward (QKV proj GEMM → TensorCoreAttention::forward → MLP)
         → if tool call JSON → TavilyClient::search → inject → continue
         → AttentionMemory grows (reset on overflow)

CUDA Notes

• half inputs, float accumulator, scale=1/sqrt(head_dim), causal mask, __shfl_xor max/sum reductions, __float2half store via smem_out (wmma::store_matrix_sync → shared → half).
• Head-dim tiling required for head_dim>16 (64/128): outer d_tile loops for QK^T reduction and AV.
• Hopper TMA: cp.async.bulk.tensor.2d.shared::cluster.global.mbarrier::complete_tx::bytes + mbarrier.init/arrive.expect_tx/try_wait (see cuda/tma_mbarrier.ptx). On Hopper cuda::memcpy_async lowers to TMA when 16-byte aligned + barrier in SMEM.

License

BSL-1.1 / AGPL-3.0 / MPL-2.0 — Patent Pending, BEL ESPRIT D ACCORD TRUST HOLDINGS INC.