### Install RotorQuant and Triton

Source: https://github.com/scrya-com/rotorquant/blob/main/README.md

Install the RotorQuant package and the Triton library using pip. This is required for using the Python/Triton examples.

```bash
pip install -e . && pip install triton
```

--------------------------------

### RotorQuant Perplexity Benchmark Script

Source: https://github.com/scrya-com/rotorquant/blob/main/grok_prompt.md

This script sets up and runs the RotorQuant perplexity benchmark. It parses command-line arguments for model configuration, loads the dataset and model, and then computes perplexity for FP16 and various RotorQuant bit settings. Ensure necessary libraries like transformers, datasets, and torch are installed.

```python
parser.add_argument("--max-length", type=int, default=2048)
    parser.add_argument("--stride", type=int, default=512)
    parser.add_argument("--max-tokens", type=int, default=0,
                        help="Limit dataset tokens (0=full test set)")
    args = parser.parse_args()

    os.environ['TRANSFORMERS_NO_ADVISORY_WARNINGS'] = '1'
    import logging; logging.disable(logging.WARNING)

    print()
    print("=" * 70)
    print("  RotorQuant Perplexity Benchmark (wikitext-2)")
    print(f"  Model: {args.model}")
    print(f"  Bits: {args.bits}")
    print(f"  Window: {args.max_length}, stride: {args.stride}")
    print(f"  GPU: {torch.cuda.get_device_name()}")
    print("=" * 70)

    # Load dataset
    print("\nLoading wikitext-2...", flush=True)
    from datasets import load_dataset
    dataset = load_dataset("wikitext", "wikitext-2-raw-v1", split="test")
    text = "\n\n".join(dataset["text"])

    # Load model
    print("Loading model...", flush=True)
    from transformers import AutoModelForCausalLM, AutoTokenizer, BitsAndBytesConfig

    tokenizer = AutoTokenizer.from_pretrained(args.model)
    model = AutoModelForCausalLM.from_pretrained(
        args.model,
        quantization_config=BitsAndBytesConfig(
            load_in_4bit=True,
            bnb_4bit_compute_dtype=torch.float16,
            bnb_4bit_quant_type="nf4",
        ),
        device_map="auto",
        dtype=torch.float16,
    )
    model.eval()

    if args.max_tokens > 0:
        tokens = tokenizer(text, return_tensors="pt")
        text = tokenizer.decode(tokens.input_ids[0][:args.max_tokens])

    total_tokens = len(tokenizer(text).input_ids)
    print(f"Dataset: {total_tokens:,} tokens")
    print()

    # FP16 baseline
    print("Computing FP16 baseline perplexity...", flush=True)
    t0 = time.perf_counter()
    ppl_fp16, n_tok = compute_perplexity(
        model, tokenizer, text,
        max_length=args.max_length, stride=args.stride,
    )
    t_fp16 = time.perf_counter() - t0
    print(f"  FP16:     PPL = {ppl_fp16:.2f}  ({n_tok:,} tokens, {t_fp16:.1f}s)")
    print()

    # RotorQuant: roundtrip during forward pass (worst case — compounding)
    print("  Roundtrip quantization (keys quantized during forward pass):")
    print(f"  {'Method':>15s}  {'PPL':>8s}  {'Delta':>8s}  {'%change':>8s}  {'Time':>8s}")
    print(f"  {'─'*15}  {'─'*8}  {'─'*8}  {'─'*8}  {'─'*8}")
    print(f"  {'FP16':>15s}  {ppl_fp16:>8.2f}  {'—':>8s}  {'—':>8s}  {t_fp16:>6.1f}s")

    for bits in args.bits:
        torch.cuda.empty_cache()
        gc.collect()

        t0 = time.perf_counter()
        ppl_rq, n_tok = compute_perplexity_with_rq(
            model, tokenizer, text, bits=bits,
            max_length=args.max_length, stride=args.stride,
        )
        t_rq = time.perf_counter() - t0

        delta = ppl_rq - ppl_fp16
        pct = (ppl_rq - ppl_fp16) / ppl_fp16 * 100

        print(f"  {'RQ ' + str(bits) + '-bit':>15s}  {ppl_rq:>8.2f}  {delta:>+8.2f}  {pct:>+7.1f}%  {t_rq:>6.1f}s")

    print()
    print("  NOTE: Roundtrip quantization degrades perplexity for ALL methods")
    print("  (TurboQuant roundtrip is even worse: PPL ~12,000).")
    print("  Google's 'zero accuracy loss' requires the fused attention kernel")
    print("  which computes Q@K^T directly from compressed keys without roundtrip.")
    print("  The post-prefill strategy (used for generation) avoids this by running")
    print("  prefill at full FP16 precision.")

    print()
    print("=" * 70)
    print("DONE")
    print("=" * 70)


if __name__ == "__main__":
    main()
```

--------------------------------

### Install Fused Attention

Source: https://github.com/scrya-com/rotorquant/blob/main/grok_prompt.md

Patches all attention layers in a Gemma3 model to use fused TurboQuant. Returns a CompressedKVCache for use with generate().

```python
def install_fused_attention(model, bits: int = 4) -> CompressedKVCache:
    """Patch all attention layers in a Gemma3 model to use fused TurboQuant.

    Returns a CompressedKVCache to pass as past_key_values to generate().
    """
    # Detect head_dim
    config = model.config
    if hasattr(config, 'text_config'):
        text_config = config.text_config
    else:
        text_config = config
    head_dim = getattr(text_config, 'head_dim', 256)

    # Create quantizer
    tq = TurboQuantMSE(d=head_dim, bits=bits, device="cuda")

    # Create compressed cache
    cache = CompressedKVCache(tq)

    # Find and patch text attention layers (not vision encoder)
    patched = 0
    layer_idx = 0
    for name, module in model.named_modules():
        if all(hasattr(module, attr) for attr in ['q_proj', 'k_proj', 'v_proj', 'out_proj', 'num_heads']):
            module.forward = make_fused_attention_forward(module, cache, tq, layer_idx)
            patched += 1
            layer_idx += 1

    print(f"  Installed fused TurboQuant ({bits}-bit) on {patched} attention layers")
    return cache
```

--------------------------------

### Clifford Algebra Cl(3,0) Imports

Source: https://github.com/scrya-com/rotorquant/blob/main/grok_prompt.md

Initial setup and imports for the Clifford algebra module used in RotorQuant.

```python
import torch
import torch.nn as nn
import math
from typing import Tuple
```

--------------------------------

### Fused TurboQuant Runner Usage Example

Source: https://github.com/scrya-com/rotorquant/blob/main/grok_prompt.md

Demonstrates how to instantiate and use the `FusedTurboQuantRunner` for text generation with a Gemma 3 model. Requires the model, processor, and desired bit precision.

```python
from turboquant_fused import FusedTurboQuantRunner
runner = FusedTurboQuantRunner(model, processor, bits=4)
text = runner.generate("What is 2+2?", max_new_tokens=30)
```

--------------------------------

### Install Fused Rotor Attention

Source: https://github.com/scrya-com/rotorquant/blob/main/grok_prompt.md

Patches all attention layers in a model to use the fused RotorQuant attention mechanism. Returns a RotorQuantCompressedCache to manage past key-value states.

```python
def install_fused_rotor_attention(model, bits: int = 3) -> RotorQuantCompressedCache:
    """Patch all attention layers to use fused RotorQuant.

    Returns a RotorQuantCompressedCache to pass as past_key_values.
    """
    config = model.config
    text_config = getattr(config, 'text_config', config)
    head_dim = getattr(text_config, 'head_dim',
                       text_config.hidden_size // text_config.num_attention_heads)

    rq = RotorQuantMSE(head_dim, bits, device="cuda")
    cache = RotorQuantCompressedCache(rq, device="cuda")

    patched = 0
    layer_idx = 0
    for name, module in model.named_modules():
        has_projs = all(hasattr(module, a) for a in ['q_proj', 'k_proj', 'v_proj'])
        has_out = hasattr(module, 'o_proj') or hasattr(module, 'out_proj')
        if has_projs and has_out:
            # Inject head counts if not on module (Qwen2 stores in config)
            if not hasattr(module, 'num_heads'):
                module.num_heads = text_config.num_attention_heads
            if not hasattr(module, 'num_key_value_heads'):
                module.num_key_value_heads = getattr(
                    text_config, 'num_key_value_heads', text_config.num_attention_heads)
            module.forward = make_fused_rotor_attention_forward(
                module, cache, layer_idx)
            patched += 1
            layer_idx += 1

    print(f"  Installed fused RotorQuant ({bits}-bit) on {patched} attention layers")
    return cache
```

--------------------------------

### Install Fused Rotor Attention

Source: https://context7.com/scrya-com/rotorquant/llms.txt

Patches transformer attention layers to use RotorQuant compression with the QJL two-term unbiased estimator.

```python
import torch
from transformers import AutoModelForCausalLM, AutoTokenizer
from turboquant.fused_attention import (
    install_fused_rotor_attention,
    RotorQuantCompressedCache,
)

# Load model
model = AutoModelForCausalLM.from_pretrained(
    "Qwen/Qwen2.5-3B",
    torch_dtype=torch.float16,
    device_map="cuda"
)
tokenizer = AutoTokenizer.from_pretrained("Qwen/Qwen2.5-3B")

# Install fused RotorQuant + QJL attention (patches all layers)
cache = install_fused_rotor_attention(model, bits=3)

# Generate with compressed KV cache
inputs = tokenizer("The future of AI is", return_tensors="pt").to("cuda")
outputs = model.generate(
    **inputs,
    max_new_tokens=50,
    past_key_values=cache,
    use_cache=True,
)
print(tokenizer.decode(outputs[0]))

# Inspect compressed cache
for layer_idx in range(len(cache._compressed_keys)):
    compressed = cache.get_compressed_key(layer_idx)
    if compressed:
        print(f"Layer {layer_idx}:")
        print(f"  MSE indices shape: {compressed['mse_indices'].shape}")
        print(f"  QJL signs shape: {compressed['qjl_signs'].shape}")
```

--------------------------------

### Run llama-server with K-only PlanarQuant Quantization

Source: https://github.com/scrya-com/rotorquant/blob/main/README.md

Launch the llama-server using K-only PlanarQuant for the KV cache. This offers zero PPL loss and significant compression.

```bash
./build/bin/llama-server -m model.gguf --jinja -ngl 99 \
    --cache-type-k planar3 --cache-type-v f16 --host 0.0.0.0 --port 8080
```

--------------------------------

### Execute Kernel Tests and Benchmarks

Source: https://github.com/scrya-com/rotorquant/blob/main/grok_prompt.md

Main entry point to run the accuracy tests and performance benchmarks.

```python
if __name__ == "__main__":
    print("=" * 60)
    print("Triton fused attention kernel tests")
    print("=" * 60)

    ok = test_fused_kernel()
    if not ok:
        print("Kernel test FAILED — skipping benchmark")
        exit(1)

    print("Benchmarking fused vs standard attention scores:")
    benchmark_fused_vs_standard()
```

--------------------------------

### Deploy with llama.cpp

Source: https://context7.com/scrya-com/rotorquant/llms.txt

Builds and runs the RotorQuant-supported llama.cpp fork for production deployment and benchmarking.

```bash
# Clone and build llama.cpp with RotorQuant support
git clone https://github.com/johndpope/llama-cpp-turboquant.git
cd llama-cpp-turboquant && git checkout feature/planarquant-kv-cache

# Build with CUDA
cmake -B build -DGGML_CUDA=ON -DCMAKE_BUILD_TYPE=Release
cmake --build build -j

# Run server with symmetric 3-bit KV compression (best quality per bit)
./build/bin/llama-server -m model.gguf --jinja -ngl 99 \
    --cache-type-k iso3 --cache-type-v iso3 --host 0.0.0.0 --port 8080

# K-only compression (zero PPL loss, 5x compression)
./build/bin/llama-server -m model.gguf --jinja -ngl 99 \
    --cache-type-k planar3 --cache-type-v f16 --host 0.0.0.0 --port 8080

# Benchmark decode/prefill speed
./build/bin/llama-bench -m model.gguf -ngl 99 -ctk planar3 -ctv planar3 -p 512 -n 128

# Measure perplexity on WikiText-2
./build/bin/llama-perplexity -m model.gguf -f /tmp/wiki.txt -ngl 99 -c 2048 \
    --cache-type-k iso3 --cache-type-v iso3
```

--------------------------------

### Run llama-server with Symmetric 3-bit Quantization

Source: https://github.com/scrya-com/rotorquant/blob/main/README.md

Launch the llama-server using symmetric 3-bit IsoQuant for the KV cache. This provides the best quality per bit.

```bash
./build/bin/llama-server -m model.gguf --jinja -ngl 99 \
    --cache-type-k iso3 --cache-type-v iso3 --host 0.0.0.0 --port 8080
```

--------------------------------

### RotorQuantProd Initialization

Source: https://github.com/scrya-com/rotorquant/blob/main/grok_prompt.md

Initializes the RotorQuantProd module, setting up the MSE quantizer and the QJL projection matrix.

```python
def __init__(self, d: int, bits: int, qjl_dim: Optional[int] = None,
                 seed: int = 42, device: str = "cpu"):
        super().__init__()
        self.d = d
        self.bits = bits
        self.mse_bits = max(bits - 1, 1)
        self.qjl_dim = qjl_dim or d
        self.device = device

        # Stage 1: Rotor MSE quantizer
        self.mse = RotorQuantMSE(d, self.mse_bits, seed=seed, device=device)

        # Stage 2: QJL projection matrix (same as TurboQuant)
        gen = torch.Generator(device='cpu')
        gen.manual_seed(seed + 1)
        S = torch.randn(self.qjl_dim, d, generator=gen)
        self.register_buffer("S", S.to(device))
```

--------------------------------

### Initialize and use PlanarQuantMSE

Source: https://context7.com/scrya-com/rotorquant/llms.txt

PlanarQuantMSE utilizes 2D Givens rotations for efficient quantization. It is the fastest method and provides optimal 3-bit perplexity.

```python
import torch
from turboquant import PlanarQuantMSE

# Initialize 3-bit PlanarQuant
pq = PlanarQuantMSE(d=128, bits=3, seed=42, device='cuda')

# Quantize attention keys
keys = torch.randn(64, 128, device='cuda')  # batch of 64 tokens
x_hat, indices = pq(keys)

# Manual quantize/dequantize workflow
v_q, indices_dict = pq.quantize(keys)
reconstructed = pq.dequantize(indices_dict)

# Access internal rotation parameters
print(f"Number of 2D groups: {pq.n_groups}")       # 64 groups for d=128
print(f"Rotation params shape: {pq.rot2.shape}")   # (64, 2) as [cos θ, sin θ]
print(f"Centroids: {pq.centroids}")                # Lloyd-Max optimal centroids

# Compare reconstruction quality
cosine_sim = torch.nn.functional.cosine_similarity(keys, reconstructed, dim=-1)
print(f"Mean cosine similarity: {cosine_sim.mean():.4f}")
```

--------------------------------

### Benchmark llama.cpp with PlanarQuant KV Cache

Source: https://github.com/scrya-com/rotorquant/blob/main/README.md

Run llama-bench to benchmark performance using PlanarQuant for the KV cache. Adjust parameters like -p and -n as needed.

```bash
./build/bin/llama-bench -m model.gguf -ngl 99 -ctk planar3 -ctv planar3 -p 512 -n 128
```

--------------------------------

### Initialize and Use IsoQuantMSE in Python

Source: https://github.com/scrya-com/rotorquant/blob/main/README.md

Initialize an IsoQuantMSE object for 4-bit quantization and apply it to input data `x`. Requires CUDA device.

```python
from turboquant import IsoQuantMSE, PlanarQuantMSE

# IsoQuant: best 4-bit quality (PPL 9.03)
iq = IsoQuantMSE(d=128, bits=4, mode='fast', device='cuda')
x_hat, indices = iq(x)
```

--------------------------------

### Initialize and Use PlanarQuantMSE in Python

Source: https://github.com/scrya-com/rotorquant/blob/main/README.md

Initialize a PlanarQuantMSE object for 3-bit quantization and apply it to input data `x`. Requires CUDA device.

```python
# PlanarQuant: best 3-bit quality (PPL 10.12)
pq = PlanarQuantMSE(d=128, bits=3, device='cuda')
x_hat, indices = pq(x)
```

--------------------------------

### Initialize RotorQuantCompressedCache

Source: https://github.com/scrya-com/rotorquant/blob/main/grok_prompt.md

Initializes the cache with RotorQuant parameters and prepares packed rotors for Triton kernels.

```python
class RotorQuantCompressedCache(DynamicCache):
    """KV cache that stores compressed keys (uint8 indices + fp16 norms).

    Keys are quantized on insertion via rotor embedding + grade-aware quantization.
    During attention, the fused Triton kernel reads compressed keys directly.
    Values are stored in fp16 (standard).
    """

    def __init__(self, rq: RotorQuantMSE, device: str = "cuda"):
        super().__init__()
        self.rq = rq
        self.device = device
        self.packed_rotors = pack_rotors_for_triton(rq.rotors).to(device)
        self.centroids_vector = getattr(rq, 'centroids_vector').to(device)
        self.centroids_trivector = getattr(rq, 'centroids_trivector').to(device)

        # Build a single flat centroid array for the gather-dot kernel.
        # Vector and trivector use different codebooks, but we can use
        # vector centroids for all 4 components (3 vector + 1 trivector)
        # since they have similar scale after the codebook fix.
        # For maximum accuracy, use vector centroids for grade-1 and
        # trivector centroids for grade-3.
        self.n_groups = rq.n_groups

        # Per-layer compressed key storage
        self._compressed_keys: list[Optional[dict]] = []
```

--------------------------------

### Initialize and Use RotorQuantKVCache

Source: https://context7.com/scrya-com/rotorquant/llms.txt

Initialize a Clifford rotor-based KV cache and perform batch append, attention scoring, and value retrieval.

```python
import torch
from turboquant import RotorQuantKVCache

# Initialize with Clifford rotor compression
cache = RotorQuantKVCache(
    d_key=128,
    d_value=128,
    bits=3,
    device='cuda'
)

# Batch append multiple KV pairs
batch_keys = torch.randn(8, 128, device='cuda')
batch_values = torch.randn(8, 128, device='cuda')
cache.append(batch_keys, batch_values)

# Query attention
queries = torch.randn(8, 128, device='cuda')
attention_scores = cache.attention_scores(queries)

# Retrieve values
cached_values = cache.get_values()
print(f"Cached values shape: {cached_values.shape}")
```

--------------------------------

### Clone and Build llama-cpp-turboquant

Source: https://github.com/scrya-com/rotorquant/blob/main/README.md

Clone the llama-cpp-turboquant repository and build it with CUDA or Metal support. This is the recommended way to use RotorQuant for the fastest performance.

```bash
git clone https://github.com/johndpope/llama-cpp-turboquant.git
cd llama-cpp-turboquant && git checkout feature/planarquant-kv-cache

# CUDA
cmake -B build -DGGML_CUDA=ON -DCMAKE_BUILD_TYPE=Release && cmake --build build -j

# Metal (Apple Silicon)
cmake -B build -DGGML_METAL=ON -DGGML_METAL_EMBED_LIBRARY=ON -DCMAKE_BUILD_TYPE=Release && cmake --build build -j
```

--------------------------------

### Initialize and use IsoQuantMSE

Source: https://context7.com/scrya-com/rotorquant/llms.txt

IsoQuantMSE uses quaternion sandwich products for 4D block rotation. Use 'full' mode for higher precision or 'fast' mode for increased speed.

```python
import torch
from turboquant import IsoQuantMSE

# Initialize quantizer for 128-dimensional vectors with 4-bit precision
iq = IsoQuantMSE(d=128, bits=4, mode='full', device='cuda')

# Create sample KV cache vectors (batch of 32 tokens)
keys = torch.randn(32, 128, device='cuda')

# Full quantize-dequantize cycle
x_hat, indices = iq(keys)

# Manual quantize/dequantize for KV cache storage
v_q, indices_dict = iq.quantize(keys)
reconstructed = iq.dequantize(indices_dict)

# Access stored norms and indices
norms = indices_dict['_norms']      # (32,) original vector norms
idx = indices_dict['indices']       # (32, 128) centroid indices

# Check reconstruction error
mse = ((keys - reconstructed) ** 2).mean()
print(f"Reconstruction MSE: {mse.item():.6f}")

# Use 'fast' mode for ~2x faster rotation (3 DOF vs 6 DOF)
iq_fast = IsoQuantMSE(d=128, bits=3, mode='fast', device='cuda')
x_hat_fast, _ = iq_fast(keys)
```

--------------------------------

### Run RotorQuant Benchmark Scripts

Source: https://github.com/scrya-com/rotorquant/blob/main/README.md

Execute various benchmark scripts provided by RotorQuant to evaluate perplexity and Triton kernel performance. Use --bits to specify quantization bits.

```bash
python -m turboquant.benchmark_google_parity          # PPL (post-prefill)
python -m turboquant.benchmark_perplexity --bits 3 4   # PPL (roundtrip)
python -m turboquant.benchmark_triton                  # Triton kernel speed
python -m turboquant.poc_high_context --backend planar  # High-context generation
```

--------------------------------

### RotorQuantMSE Initialization

Source: https://github.com/scrya-com/rotorquant/blob/main/grok_prompt.md

Initializes the RotorQuantMSE module, setting up parameters for Clifford algebra-based quantization. It configures grade-aware bit allocation and pre-computes random rotors for decorrelation.

```python
def __init__(self, d: int, bits: int, seed: int = 42,
                 grade_bits: Optional[Dict[str, int]] = None,
                 device: str = "cpu"):
        super().__init__()
        self.d = d
        self.bits = bits
        self.device = device

        self.n_groups = (d + 2) // 3
        self.mv_dim = self.n_groups * MV_DIM

        if grade_bits is None:
            grade_bits = {
                'vector': bits,
                'trivector': max(bits - 1, 1),
            }
        self.grade_bits = grade_bits

        d_eff_vector = d
        d_eff_trivector = max(d // 2, 8)
        self.codebooks = nn.ModuleDict()
        for grade_name, gb in grade_bits.items():
            if grade_name == 'trivector':
                cb = LloydMaxCodebook(d_eff_trivector, gb)
            else:
                cb = LloydMaxCodebook(d_eff_vector, gb)
            self.register_buffer(f'centroids_{grade_name}',
                                 cb.centroids.to(device))

        self.grade_map = {
            'vector':   [1, 2, 3],
            'trivector': [7],
        }

        rotors = []
        for i in range(self.n_groups):
            r = make_random_rotor((), device=device, seed=seed + i)
            rotors.append(r)
        self.register_buffer('rotors', torch.stack(rotors))
```

--------------------------------

### Calculate Perplexity with llama.cpp

Source: https://github.com/scrya-com/rotorquant/blob/main/README.md

Calculate the perplexity of a model using a text file. Ensure the dataset is downloaded and the text file is created.

```bash
pip install datasets
python3 -c "from datasets import load_dataset; open('/tmp/wiki.txt','w').write('
'.join(load_dataset('wikitext','wikitext-2-raw-v1',split='test')['text']))"
./build/bin/llama-perplexity -m model.gguf -f /tmp/wiki.txt -ngl 99 -c 2048 \
    --cache-type-k iso3 --cache-type-v iso3
```

--------------------------------

### TurboQuantKVCache Wrapper

Source: https://context7.com/scrya-com/rotorquant/llms.txt

Provides a drop-in wrapper for transformer KV caches to manage memory usage via TurboQuant compression.

```python
import torch
from turboquant import TurboQuantKVCache

# Initialize KV cache for typical transformer dimensions
cache = TurboQuantKVCache(
    d_key=128,     # Key dimension
    d_value=128,   # Value dimension
    bits=3,        # 3-bit quantization
    device='cuda'
)

# Simulate autoregressive generation
for step in range(100):
    # New KV pairs from attention layer
    new_keys = torch.randn(1, 1, 128, device='cuda')    # (batch, seq, dim)
    new_values = torch.randn(1, 1, 128, device='cuda')

    cache.append(new_keys, new_values)

# Compute attention scores with current query
query = torch.randn(128, device='cuda')
scores = cache.attention_scores(query)  # Scores for all 100 cached positions

# Get reconstructed values for attention output
values = cache.get_values()  # (100, 128)
```

--------------------------------

### RotorQuantKVCache Implementation

Source: https://github.com/scrya-com/rotorquant/blob/main/grok_prompt.md

A drop-in replacement for TurboQuantKVCache that uses RotorQuant compression for key and value caching.

```python
class RotorQuantKVCache:
    """
    KV cache using RotorQuant compression.
    Drop-in replacement for TurboQuantKVCache.
    """

    def __init__(self, d_key: int, d_value: int, bits: int = 3,
                 seed: int = 42, device: str = "cpu"):
        self.d_key = d_key
        self.d_value = d_value
        self.bits = bits
        self.device = device

        self.key_quantizer = RotorQuantProd(d_key, bits, seed=seed, device=device)
        self.value_quantizer = RotorQuantMSE(d_value, bits, seed=seed + 100, device=device)

        self.key_cache = []
        self.value_cache = []

    def append(self, keys: torch.Tensor, values: torch.Tensor):
        flat_keys = keys.reshape(-1, self.d_key)
        flat_values = values.reshape(-1, self.d_value)

        compressed_keys = self.key_quantizer.quantize(flat_keys)
        _, value_indices = self.value_quantizer(flat_values)

        self.key_cache.append(compressed_keys)
        self.value_cache.append(value_indices)

    def attention_scores(self, queries: torch.Tensor) -> torch.Tensor:
        scores = []
        for cached in self.key_cache:
            s = self.key_quantizer.inner_product(queries, cached)
            scores.append(s)
        return torch.cat(scores, dim=-1) if scores else torch.tensor([])

    def get_values(self) -> torch.Tensor:
        values = []
        for indices in self.value_cache:
            v = self.value_quantizer.dequantize(indices)
            values.append(v)
        return torch.cat(values, dim=0) if values else torch.tensor([])

    def __len__(self):
        return sum(
            c['qjl_signs'].shape[0] for c in self.key_cache
        ) if self.key_cache else 0
```

--------------------------------

### Initialize and use RotorQuantMSE

Source: https://context7.com/scrya-com/rotorquant/llms.txt

RotorQuantMSE employs Clifford algebra Cl(3,0) rotors for grade-aware quantization. It allows for distinct bit budgets for vector and trivector components.

```python
import torch
from turboquant import RotorQuantMSE

# Initialize with grade-aware bit allocation
grade_bits = {'vector': 3, 'trivector': 2}
rq = RotorQuantMSE(d=128, bits=3, grade_bits=grade_bits, device='cuda')

# Quantize vectors
keys = torch.randn(16, 128, device='cuda')
x_hat, indices = rq(keys)

# Inspect multivector structure
print(f"Number of Cl(3,0) groups: {rq.n_groups}")  # ceil(128/3) = 43 groups
print(f"MV dimension per group: 8")                 # [1, e1, e2, e3, e12, e13, e23, e123]

# Access grade-specific indices
vector_indices = indices['vector']       # Indices for e1, e2, e3 components
trivector_indices = indices['trivector'] # Indices for e123 component
stored_norms = indices['_norms']         # Original vector norms

# Manual dequantization
reconstructed = rq.dequantize(indices)
```

--------------------------------

### Python Wrapper for Triton Inverse Rotor Sandwich Kernel

Source: https://github.com/scrya-com/rotorquant/blob/main/grok_prompt.md

Provides a Python interface to the Triton kernel for the inverse rotor sandwich operation. It handles tensor preparation, grid configuration, and kernel launching.

```python
def triton_rotor_inverse_sandwich(
    input_mv: torch.Tensor,
    rotors: torch.Tensor,
    emb_dim: int,
) -> torch.Tensor:
    """Inverse rotor sandwich R̃ x R using Triton."""
    batch_size, n_groups, _ = input_mv.shape

    input_f32 = input_mv.float().contiguous()
    rotors_f32 = rotors.float().contiguous()

    output = torch.empty(batch_size, emb_dim,
                         device=input_mv.device, dtype=torch.float32)

    BLOCK_G = min(triton.next_power_of_2(n_groups), 128)
    grid = (batch_size, triton.cdiv(n_groups, BLOCK_G))

    _rotor_inverse_sandwich_kernel[grid](
        input_f32, rotors_f32, output,
        batch_size, emb_dim, n_groups,
        input_f32.stride(0), input_f32.stride(1), input_f32.stride(2),
        output.stride(0), output.stride(1),
        BLOCK_G=BLOCK_G,
    )

    return output.to(input_mv.dtype)
```

--------------------------------

### PyTorch C++ Template for QJL Score

Source: https://github.com/scrya-com/rotorquant/blob/main/grok_prompt.md

Template function to prepare tensors and launch the CUDA kernel for score calculation.

```cpp
template <typename T, typename Tproj>
torch::Tensor QJLScoreCudaTemplate(
    torch::Tensor key_quant,
    torch::Tensor key_outlier_quant,
    torch::Tensor key_norm,
    torch::Tensor key_outlier_norm,
    torch::Tensor outlier_indices,
    torch::Tensor query_sketch,
    torch::Tensor query_states,
    torch::Tensor rand_prj) {

    auto options = torch::TensorOptions().device(torch::kCUDA, 0).dtype(torch::kFloat);

    int batch = key_quant.size(0);
    int head = key_quant.size(1);
    int n = key_quant.size(2);
    int group_size = key_quant.size(3);
    int emb_dim = query_states.size(3);
    int sketch_dim = rand_prj.size(1);
    int outlier_sketch_dim = 8*key_outlier_quant.size(4);
    int outlier_counts = outlier_indices.size(3);

    auto scores = torch::zeros({batch, head, n * group_size, 1}, options).contiguous();
    
    auto query_states_ptr = query_states.data_ptr<T>();
    auto key_norm_ptr = key_norm.data_ptr<T>();
    auto key_outlier_norm_ptr = key_outlier_norm.data_ptr<T>();
    auto rand_prj_ptr = rand_prj.data_ptr<Tproj>();

    int blocksPerGroup = (group_size + WARP_SIZE - 1) / WARP_SIZE;
    dim3 numBlocks(batch * head, n, blocksPerGroup);
    dim3 threadsPerBlockDim(WARP_SIZE, WARPS_PER_BLOCK, 1);

    calc_score_kernel<<<numBlocks, threadsPerBlockDim>>>(
        query_states_ptr,
        key_quant.data_ptr<uint8_t>(),
        key_outlier_quant.data_ptr<uint8_t>(),
        key_norm_ptr,
        key_outlier_norm_ptr,
        outlier_indices.data_ptr<uint8_t>(),
        query_sketch.data_ptr<float>(),
        rand_prj_ptr,
        scores.data_ptr<float>(),
        batch, head, n, group_size, sketch_dim, outlier_sketch_dim, emb_dim, outlier_counts);

    return scores;
}
```

--------------------------------

### Manage Compressed Key Storage

Source: https://github.com/scrya-com/rotorquant/blob/main/grok_prompt.md

Methods to store and retrieve compressed key states from the cache.

```python
    def store_compressed_key(self, key_states: torch.Tensor, layer_idx: int):
        """Quantize and store key states."""
        compressed = self._quantize_keys(key_states)

        while len(self._compressed_keys) <= layer_idx:
            self._compressed_keys.append(None)

        if self._compressed_keys[layer_idx] is None:
            self._compressed_keys[layer_idx] = compressed
        else:
            prev = self._compressed_keys[layer_idx]
            self._compressed_keys[layer_idx] = {
                'indices': torch.cat([prev['indices'], compressed['indices']], dim=2),
                'norms': torch.cat([prev['norms'], compressed['norms']], dim=2),
            }

    def get_compressed_key(self, layer_idx: int) -> Optional[dict]:
        if layer_idx < len(self._compressed_keys):
            return self._compressed_keys[layer_idx]
        return None
```

--------------------------------

### Query Pre-processing: Project Through S

Source: https://github.com/scrya-com/rotorquant/blob/main/grok_prompt.md

This Python code snippet shows how to project the query vectors through the random projection matrix S to obtain the query sketch. This sketch is used in the QJL correction term.

```python
query_sketch = Q @ S.T  # [batch, n_heads, q_len, m]
```

--------------------------------

### Triton Rotor Sandwich Kernel

Source: https://github.com/scrya-com/rotorquant/blob/main/grok_prompt.md

A JIT-compiled Triton kernel that performs the rotor sandwich operation by loading packed rotors and input vectors, applying geometric product transformations, and storing the resulting multivectors.

```python
@triton.jit
def _rotor_sandwich_kernel(
    input_ptr, rotors_ptr, output_ptr,
    batch_size, emb_dim, n_groups: tl.constexpr,
    stride_in_b, stride_in_d,
    stride_out_b, stride_out_g, stride_out_c,
    BLOCK_G: tl.constexpr,
):
    pid_b = tl.program_id(0)
    pid_g = tl.program_id(1)

    g_offs = pid_g * BLOCK_G + tl.arange(0, BLOCK_G)
    g_mask = g_offs < n_groups

    # Load rotor: R = [s, p12, p13, p23]
    r_s   = tl.load(rotors_ptr + g_offs * 4 + 0, mask=g_mask, other=1.0)
    r_p12 = tl.load(rotors_ptr + g_offs * 4 + 1, mask=g_mask, other=0.0)
    r_p13 = tl.load(rotors_ptr + g_offs * 4 + 2, mask=g_mask, other=0.0)
    r_p23 = tl.load(rotors_ptr + g_offs * 4 + 3, mask=g_mask, other=0.0)

    # Embed: load 3 vector components per group
    d0 = g_offs * 3
    v1 = tl.load(input_ptr + pid_b * stride_in_b + d0 * stride_in_d,
                  mask=g_mask & (d0 < emb_dim), other=0.0)
    v2 = tl.load(input_ptr + pid_b * stride_in_b + (d0 + 1) * stride_in_d,
                  mask=g_mask & ((d0 + 1) < emb_dim), other=0.0)
    v3 = tl.load(input_ptr + pid_b * stride_in_b + (d0 + 2) * stride_in_d,
                  mask=g_mask & ((d0 + 2) < emb_dim), other=0.0)

    z = tl.zeros_like(v1)

    # Step 1: temp = R * x (rotor on LEFT)
    t0, t1, t2, t3, t4, t5, t6, t7 = _gp_rotor_mv(
        r_s, r_p12, r_p13, r_p23,
        z, v1, v2, v3, z, z, z, z,
    )

    # Step 2: result = temp * R̃ (rotor on RIGHT, reverse negates bivectors)
    o0, o1, o2, o3, o4, o5, o6, o7 = _gp_mv_rotor(
        t0, t1, t2, t3, t4, t5, t6, t7,
        r_s, -r_p12, -r_p13, -r_p23,
    )

    # Store output multivectors [batch, n_groups, 8]
    out_base = pid_b * stride_out_b + g_offs * stride_out_g
    tl.store(output_ptr + out_base + 0 * stride_out_c, o0, mask=g_mask)
    tl.store(output_ptr + out_base + 1 * stride_out_c, o1, mask=g_mask)
    tl.store(output_ptr + out_base + 2 * stride_out_c, o2, mask=g_mask)
    tl.store(output_ptr + out_base + 3 * stride_out_c, o3, mask=g_mask)
    tl.store(output_ptr + out_base + 4 * stride_out_c, o4, mask=g_mask)
    tl.store(output_ptr + out_base + 5 * stride_out_c, o5, mask=g_mask)
    tl.store(output_ptr + out_base + 6 * stride_out_c, o6, mask=g_mask)
    tl.store(output_ptr + out_base + 7 * stride_out_c, o7, mask=g_mask)
```

--------------------------------

### Execute Triton PlanarQuant Kernels

Source: https://context7.com/scrya-com/rotorquant/llms.txt

Perform fused GPU-accelerated quantization and dequantization using Triton kernels for rotation-based operations.

```python
import torch
from turboquant import PlanarQuantMSE
from turboquant.triton_planarquant import (
    triton_planar2_fused,
    triton_planar2_quantize,
    triton_planar2_dequantize,
)

# Setup quantizer to get rotation params and centroids
pq = PlanarQuantMSE(d=128, bits=3, device='cuda')

# Fused quantize-dequantize roundtrip (single kernel)
keys = torch.randn(1024, 128, device='cuda')
reconstructed = triton_planar2_fused(keys, pq.rot2, pq.centroids)

# Quantize only - returns indices for KV cache storage
indices, norms = triton_planar2_quantize(keys, pq.rot2, pq.centroids)
print(f"Indices dtype: {indices.dtype}")  # int8
print(f"Norms shape: {norms.shape}")      # (1024,)

# Dequantize from stored indices
reconstructed_from_indices = triton_planar2_dequantize(
    indices, norms, pq.rot2, pq.centroids
)

# Verify consistency
diff = (reconstructed - reconstructed_from_indices).abs().max()
print(f"Max difference: {diff.item()}")  # Should be ~0
```

--------------------------------

### RotorQuant Pipeline Wrapper

Source: https://github.com/scrya-com/rotorquant/blob/main/grok_prompt.md

Python wrapper function to prepare inputs, calculate norms, and launch the fused Triton kernel. It handles normalization and rescaling to maintain numerical stability.

```python
def triton_rotor_full_fused(
    input: torch.Tensor,
    rotors: torch.Tensor,
    c_scalar: torch.Tensor,
    c_vector: torch.Tensor,
    c_bivector: torch.Tensor,
    c_trivector: torch.Tensor,
) -> torch.Tensor:
    """Fused RotorQuant pipeline: normalize→embed→rotor→quantize→unrotor→extract→rescale.

    Single kernel launch for the full quantize-dequantize roundtrip.
    Only quantizes non-zero grades (vector + trivector); scalar and bivector
    are always zero after sandwich of grade-1 input and are skipped.
    """
    batch_size, emb_dim = input.shape
    n_groups = rotors.shape[0]

    # Norm separation: quantize unit vectors, store norms
    input_f32 = input.float()
    norms = input_f32.norm(dim=-1, keepdim=True).clamp(min=1e-8)
    input_f32 = (input_f32 / norms).contiguous()
    rotors_f32 = rotors.float().contiguous()
    # Scalar/bivector centroids not used (zero grades), but kernel signature needs them
    c_s = c_scalar.float().contiguous() if c_scalar is not None else c_vector.float().contiguous()
    c_v = c_vector.float().contiguous()
    c_b = c_bivector.float().contiguous() if c_bivector is not None else c_vector.float().contiguous()
    c_t = c_trivector.float().contiguous()

    output = torch.empty_like(input_f32)

    BLOCK_G = min(triton.next_power_of_2(n_groups), 128)
    grid = (batch_size, triton.cdiv(n_groups, BLOCK_G))

    _rotor_full_fused_kernel[grid](
        input_f32, output, rotors_f32,
        c_s, c_v, c_b, c_t,
        batch_size, emb_dim, n_groups,
        len(c_s), len(c_v), len(c_b), len(c_t),
        input_f32.stride(0), input_f32.stride(1),
        output.stride(0), output.stride(1),
        BLOCK_G=BLOCK_G,
    )

    # Rescale by original norms
    output = output * norms
    return output.to(input.dtype)
```

--------------------------------

### Manage LiteratiQuantKVCache

Source: https://context7.com/scrya-com/rotorquant/llms.txt

Insert KV vectors into a 1-bit compressed cache, retrieve decompressed data, and monitor memory usage.

```python
import torch
from turboquant import LiteratiQuantKVCache

# Initialize 1-bit KV cache
cache = LiteratiQuantKVCache(d=128, group_size=128, device='cuda')

# Insert KV vectors (batch, n_heads, seq_len, head_dim)
keys = torch.randn(2, 8, 32, 128, device='cuda')  # 2 batches, 8 heads, 32 tokens
cache.insert(keys)

# Continue inserting during generation
new_keys = torch.randn(2, 8, 1, 128, device='cuda')
cache.insert(new_keys)

print(f"Cache sequence length: {cache.seq_len}")  # 33

# Retrieve decompressed cache
all_keys = cache.get_all()  # (2, 8, 33, 128)

# Memory usage
bytes_used = cache.memory_bytes()
fp16_equiv = 2 * 8 * 33 * 128 * 2  # FP16 baseline
print(f"Memory: {bytes_used} bytes")
print(f"FP16 equivalent: {fp16_equiv} bytes")
print(f"Compression: {fp16_equiv / bytes_used:.1f}x")

# Clear cache for new sequence
cache.clear()
```

--------------------------------

### Compressed Key-Value Cache Implementation

Source: https://github.com/scrya-com/rotorquant/blob/main/grok_prompt.md

The `CompressedKVCache` class extends `DynamicCache` to store keys as quantized uint8 indices and norms, optimizing attention by avoiding fp16 dequantization. Values are stored in standard fp16.

```python
import torch
import math
from transformers import DynamicCache
from turboquant_core import TurboQuantMSE
from triton_attention import fused_qk_scores


class CompressedKVCache(DynamicCache):
    """KV cache that stores compressed keys (uint8 indices + norms).

    Keys are quantized on insertion. During attention, the fused Triton
    kernel reads compressed keys directly — no fp16 dequantization needed.

    Values are stored in fp16 (standard) since the softmax@V matmul
    benefits less from compression.
    """

    def __init__(self, quantizer: TurboQuantMSE):
        super().__init__()
        self.tq = quantizer
        # Per-layer compressed key storage
        self._compressed_keys: list[dict | None] = []

    def store_compressed_key(self, key_states: torch.Tensor, layer_idx: int):
        """Quantize and store key states. Called from patched attention."""
        while len(self._compressed_keys) <= layer_idx:
            self._compressed_keys.append(None)

        q = self.tq.quantize(key_states.float())

        if self._compressed_keys[layer_idx] is None:
            self._compressed_keys[layer_idx] = q
        else:
            prev = self._compressed_keys[layer_idx]
            self._compressed_keys[layer_idx] = {
                "idx": torch.cat([prev["idx"], q["idx"]], dim=2),
                "norms": torch.cat([prev["norms"], q["norms"]], dim=2),
            }

    def get_compressed_key(self, layer_idx: int) -> dict | None:
        if layer_idx < len(self._compressed_keys):
            return self._compressed_keys[layer_idx]
        return None

    def get_kv_seq_length(self, layer_idx: int = 0) -> int:
        if layer_idx < len(self._compressed_keys) and self._compressed_keys[layer_idx] is not None:
            return self._compressed_keys[layer_idx]["idx"].shape[2]
        return 0
```

--------------------------------

### Key Compression: Compute QJL After MSE Quantization

Source: https://github.com/scrya-com/rotorquant/blob/main/grok_prompt.md

This Python code demonstrates the steps for computing QJL correction components during key compression. It involves calculating the residual, its norm, and the sign projections using a random projection matrix S.

```python
k_mse = dequantize(indices, key_norms)
residual = k_original - k_mse
residual_norm = ||residual||
qjl_signs = sign(S @ residual.T)  # 1-bit per dim
```