README.md

---
license: other
license_name: snapkitty-tri-license
license_link: https://huggingface.co/Snapkitty/macrogrok-harness/blob/main/LICENSE.tri
library_name: custom
tags:
- code
- sovereign-compute
- snapkitty
- compilers
- python
---

> Source: [github.com/SNAPKITTYWEST/macrogrok-harness](https://github.com/SNAPKITTYWEST/macrogrok-harness)

# 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

```toml
# 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"
```

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

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

### 💼 Commercial License

Snapkitty code is free and open under **AGPL-3.0** for open-source use. Building a commercial product or service? A **proprietary commercial license** from Snapkitty Collective LLC lets you ship this code without the AGPL's source-sharing and network-use obligations.

**[→ Get a commercial license](mailto:A.parr@belespritdaccord.uk?subject=Commercial%20license:%20macrogrok-harness)** · A.parr@belespritdaccord.uk