Chronos PoC: PTX transform selection via RLVR (DA-GRPO)

README.md

@@ -0,0 +1,156 @@
+# Chronos PoC: PTX Transform Selection via RLVR
+
+Proof-of-concept RL agent that selects PTX assembly transforms to optimize GPU kernel performance on NVIDIA L4 (sm_89).
+
+## What This Is
+
+An MLP policy trained with DA-GRPO (Demonstration-Anchored Group Relative Policy Optimization) to select sequences of PTX-level transforms that reduce GPU kernel execution cycles. Trained on 64 gemm_tile kernel variants, validated on 33 diverse Triton kernels.
+
+## Results
+
+| Metric | Value |
+|--------|-------|
+| Mean cycle reduction (gemm_tile) | **-29.2%** |
+| Best single kernel | **-53.8%** (gemm_tile 4,6,8: 1839 -> 849 cycles) |
+| Generalization to Triton kernels | 17/33 kernels improve |
+| Best Triton improvement | **-47.0%** (attention_d64_kv64 with maxnreg_255) |
+| Training time | ~6 hours on single NVIDIA L4 |
+| Model parameters | ~20K |
+
+## Model Architecture
+
+```
+Input:  25 kernel features + 21 action mask + 21 action history = 67 dims
+Hidden: Linear(67, 128) -> ReLU -> Dropout(0.1) -> Linear(128, 128) -> ReLU -> Dropout(0.1)
+Output: Linear(128, 21) -> mask -> softmax
+```
+
+**25 features**: instruction counts (loads, stores, FMA, branches), vectorization ratios, cache hint coverage, register state, instruction mix ratios.
+
+**21 actions**: 20 PTX transforms + stop. Transforms include vectorize loads/stores, cache hints (cs/cg/ca/cv), store cache hints (cs/wt/wb), register budget limits (32/64/128/255), instruction reorder (critical_path/interleave/loads_first/stores_last), prefetch, and split vector loads.
+
+## Files
+
+```
+checkpoint_best.pt       # Best checkpoint (epoch 250, -29.2% mean)
+checkpoint_latest.pt     # Final checkpoint (epoch 500)
+inference.py             # Self-contained inference script
+training_result.json     # Per-kernel results (64 gemm_tile kernels)
+bc_stats.json            # Behavior cloning warm-start statistics
+measure_triton_results.json  # Triton kernel measurement results (33 kernels)
+```
+
+## Inference
+
+### Requirements
+
+```
+pip install torch
+```
+
+No GPU required for inference (model runs on CPU). GPU only needed for actually applying transforms and measuring cycles.
+
+### Quick Start
+
+```python
+import torch
+from inference import load_model, predict_transforms
+
+# Load model
+model = load_model("checkpoint_best.pt")
+
+# Predict transforms for a PTX kernel
+with open("your_kernel.ptx") as f:
+    ptx = f.read()
+actions = predict_transforms(model, ptx)
+# -> ['maxnreg_128', 'vec_ld', 'vec_st']
+```
+
+### Command Line
+
+```bash
+# Demo with synthetic features
+python inference.py --checkpoint checkpoint_best.pt
+
+# Run on a PTX file
+python inference.py --checkpoint checkpoint_best.pt --ptx path/to/kernel.ptx
+```
+
+### Programmatic Usage
+
+```python
+import torch
+from inference import TransformPolicy, extract_features_from_ptx, get_action_mask, get_action_history, ACTION_NAMES
+
+# Load
+model = TransformPolicy(hidden=128)
+ckpt = torch.load("checkpoint_best.pt", map_location="cpu", weights_only=False)
+model.load_state_dict(ckpt["policy"])
+model.eval()
+
+# Extract features from PTX
+ptx_source = open("kernel.ptx").read()
+features = extract_features_from_ptx(ptx_source)
+
+# Predict step by step
+applied = set()
+for step in range(6):
+    feat_t = torch.tensor(features, dtype=torch.float32)
+    mask_t = torch.tensor(get_action_mask(applied), dtype=torch.float32)
+    hist_t = torch.tensor(get_action_history(applied), dtype=torch.float32)
+
+    action_id = model.get_greedy_action(feat_t, mask_t, hist_t)
+    action = ACTION_NAMES[action_id]
+    if action == "stop":
+        break
+    print(f"Step {step+1}: apply {action}")
+    applied.add(action)
+```
+
+## Training Details
+
+### Algorithm: DA-GRPO
+
+1. **BC warm-start** (50 epochs): Clone greedy search trajectories. Best accuracy: 64.5%.
+2. **GRPO training** (450 epochs): Hardware-in-the-loop RL with SM clock() cycle measurement.
+   - Group size: 8 rollouts per kernel (1 anchor from reference policy + 7 with forced diverse first actions)
+   - Advantage: MC-GRPO (median baseline per kernel, global z-normalization)
+   - Reward: log(cycles_before / cycles_after) — outcome-only, terminal
+   - KL penalty: beta=0.01 against BC reference policy
+   - Clipped surrogate: epsilon=0.2
+
+### Action Space
+
+20 PTX transforms organized into 5 conflict groups (only one per group):
+- **Cache hints** (load): cs, cg, ca, cv
+- **Store cache hints**: cs, wt, wb
+- **Register budget**: maxnreg 32, 64, 128, 255
+- **Instruction reorder**: critical_path, interleave, loads_first, stores_last
+- **Prefetch**: L1, L2
+- **Vectorize**: loads, stores (independent)
+- **Split**: vector loads (independent)
+
+### Hardware
+
+- NVIDIA L4 GPU (sm_89, Ada Lovelace)
+- SM clock() cycle counter (1-cycle std dev, 200 samples per measurement)
+- pip-installed CUDA 12.9 ptxas
+
+## Limitations
+
+- Trained on gemm_tile kernels only (64 variants, m,n,k in {2,4,6,8})
+- Mode collapse: 52/64 kernels get the same sequence (vec_st + vec_ld + maxnreg_128)
+- MLP can't read PTX code — relies on 25 scalar features
+- 3% measurement error rate from cudaErrorMisalignedAddress on gemm_tile(4,6,4)
+- Reorder transforms deadlock on kernels with bar.sync barriers
+
+## References
+
+- [CuAsmRL (CGO 2025)](https://arxiv.org/abs/2501.08071): PPO on SASS scheduling
+- [Dr. Kernel (2026)](https://arxiv.org/abs/2602.05885): REINFORCE for Triton kernels
+- [DeepSeek-R1 (2025)](https://arxiv.org/abs/2501.12948): GRPO algorithm
+- [MC-GRPO (2025)](https://arxiv.org/abs/2601.22582): Median-centered baseline
+
+## License
+
+Research prototype. Contact for usage terms.

bc_stats.json

@@ -0,0 +1,107 @@
+{
+  "bc_loss": [
+    2.0958820213222134,
+    1.7892995409983925,
+    1.7074809244700841,
+    1.5575190510989156,
+    1.441180571626052,
+    1.413911653761698,
+    1.4011170574136682,
+    1.364859262028256,
+    1.342358095765574,
+    1.338929759949791,
+    1.3145962235550162,
+    1.2938860436203856,
+    1.2686388336093268,
+    1.2415368708864603,
+    1.2314557232911982,
+    1.1912865268217552,
+    1.2434098955286976,
+    1.223619041755853,
+    1.192927201742371,
+    1.2041826229758243,
+    1.209502972230948,
+    1.178917215137408,
+    1.1955577655188365,
+    1.178613586315317,
+    1.1684499366863355,
+    1.1616653519247488,
+    1.1482060889940005,
+    1.1699149070099053,
+    1.1195348450576017,
+    1.1293887990782159,
+    1.1804475940792718,
+    1.112368142052507,
+    1.0938795485091486,
+    1.1037644244529106,
+    1.0885746230489959,
+    1.0920597482832242,
+    1.0800647707979651,
+    1.049601698474074,
+    1.0656030951319515,
+    1.1000592331168275,
+    1.0485037210825328,
+    1.0458013574128906,
+    1.0802548773960718,
+    1.0371340227863504,
+    1.0327755164455723,
+    1.0279540154637057,
+    1.0051751578636612,
+    1.0235914099515635,
+    1.0102946024143558,
+    0.9888515271045066
+  ],
+  "bc_accuracy": [
+    0.3938223938223938,
+    0.47104247104247104,
+    0.5752895752895753,
+    0.5752895752895753,
+    0.5444015444015444,
+    0.5598455598455598,
+    0.5752895752895753,
+    0.5752895752895753,
+    0.5598455598455598,
+    0.5598455598455598,
+    0.5714285714285714,
+    0.5791505791505791,
+    0.5714285714285714,
+    0.5752895752895753,
+    0.5598455598455598,
+    0.5945945945945946,
+    0.5637065637065637,
+    0.5868725868725869,
+    0.6023166023166023,
+    0.583011583011583,
+    0.6061776061776062,
+    0.6061776061776062,
+    0.5675675675675675,
+    0.5868725868725869,
+    0.5945945945945946,
+    0.6138996138996139,
+    0.5984555984555985,
+    0.5984555984555985,
+    0.6254826254826255,
+    0.5907335907335908,
+    0.5868725868725869,
+    0.6023166023166023,
+    0.6138996138996139,
+    0.5945945945945946,
+    0.6332046332046332,
+    0.6216216216216216,
+    0.6138996138996139,
+    0.6177606177606177,
+    0.6447876447876448,
+    0.6254826254826255,
+    0.6216216216216216,
+    0.6061776061776062,
+    0.5907335907335908,
+    0.6061776061776062,
+    0.6293436293436293,
+    0.6216216216216216,
+    0.6447876447876448,
+    0.6254826254826255,
+    0.6254826254826255,
+    0.6447876447876448
+  ],
+  "best_accuracy": 0.6447876447876448
+}

checkpoint_best.pt

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:4601dcda039965c4a57461e7fa32d71259302f334be34f3fe8ac99cc08f6f937
+size 469964

checkpoint_latest.pt

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:bf261b22976e88acde2df07bc5f886bae7c218c03270e1ded086c31fad4bcfc8
+size 476816

inference.py

@@ -0,0 +1,421 @@
+"""Chronos PoC: PTX transform selection inference.
+
+Loads a trained TransformPolicy checkpoint and predicts the optimal
+sequence of PTX transforms for a given kernel.
+
+Usage:
+    python inference.py --checkpoint checkpoint_best.pt --kernel gemm_tile --m 4 --n 6 --k 8
+    python inference.py --checkpoint checkpoint_best.pt --ptx path/to/kernel.ptx
+"""
+
+import argparse
+import sys
+import os
+import json
+
+import torch
+import torch.nn as nn
+import torch.nn.functional as F
+from torch.distributions import Categorical
+
+
+# ---------------------------------------------------------------------------
+# Model definition (self-contained, no external dependencies for inference)
+# ---------------------------------------------------------------------------
+
+N_FEATURES = 25  # Model was trained with 25 scalar features
+N_ACTIONS = 21
+
+ACTION_NAMES = [
+    "vec_ld", "vec_st",
+    "cache_cs", "cache_cg", "cache_ca", "cache_cv",
+    "st_cache_cs", "st_cache_wt", "st_cache_wb",
+    "maxnreg_32", "maxnreg_64", "maxnreg_128", "maxnreg_255",
+    "reorder_cp", "reorder_il", "reorder_lf", "reorder_sl",
+    "prefetch_L1", "prefetch_L2",
+    "split_ld",
+    "stop",
+]
+
+FEATURE_NAMES = [
+    "n_instructions", "n_ld_global", "n_st_global", "n_fma",
+    "n_ld_param", "n_prefetch", "n_branch",
+    "n_ld_global_vec", "n_st_global_vec", "vec_ld_ratio", "vec_st_ratio",
+    "n_cache_hint_ld", "n_cache_hint_st", "hint_ld_ratio", "hint_st_ratio",
+    "load_ratio", "store_ratio", "fma_ratio", "compute_ratio",
+    "mem_ratio", "compute_to_mem",
+    "total_regs", "n_f32_regs", "n_b64_regs", "maxnreg",
+]
+
+CONFLICT_GROUPS = {
+    "cache_hints": {"cache_cs", "cache_cg", "cache_ca", "cache_cv"},
+    "store_cache_hints": {"st_cache_cs", "st_cache_wt", "st_cache_wb"},
+    "register_budget": {"maxnreg_32", "maxnreg_64", "maxnreg_128", "maxnreg_255"},
+    "prefetch": {"prefetch_L1", "prefetch_L2"},
+    "reorder": {"reorder_cp", "reorder_il", "reorder_lf", "reorder_sl"},
+}
+
+
+class TransformPolicy(nn.Module):
+    """MLP policy for PTX transform selection.
+
+    Input:  25 features + 21 action mask + 21 action history = 67 dims
+    Output: 21 logits (masked before softmax)
+    """
+
+    def __init__(self, hidden=128):
+        super().__init__()
+        input_dim = N_FEATURES + N_ACTIONS + N_ACTIONS  # 67
+        self.net = nn.Sequential(
+            nn.Linear(input_dim, hidden),
+            nn.ReLU(),
+            nn.Dropout(0.1),
+            nn.Linear(hidden, hidden),
+            nn.ReLU(),
+            nn.Dropout(0.1),
+            nn.Linear(hidden, N_ACTIONS),
+        )
+
+    def forward(self, features, action_mask, action_history):
+        x = torch.cat([features, action_mask, action_history], dim=-1)
+        logits = self.net(x)
+        logits = logits.masked_fill(action_mask == 0, float('-inf'))
+        return logits
+
+    @torch.no_grad()
+    def get_greedy_action(self, features, action_mask, action_history):
+        logits = self.forward(
+            features.unsqueeze(0), action_mask.unsqueeze(0),
+            action_history.unsqueeze(0),
+        )
+        return logits.argmax(dim=-1).item()
+
+    @torch.no_grad()
+    def get_action_probs(self, features, action_mask, action_history):
+        logits = self.forward(
+            features.unsqueeze(0), action_mask.unsqueeze(0),
+            action_history.unsqueeze(0),
+        )
+        probs = F.softmax(logits, dim=-1)
+        return probs.squeeze(0)
+
+
+# ---------------------------------------------------------------------------
+# Feature extraction (self-contained, regex-based)
+# ---------------------------------------------------------------------------
+
+import re
+
+_LD_GLOBAL = re.compile(r'ld\.global')
+_LD_GLOBAL_VEC = re.compile(r'ld\.global(?:\.\w+)*\.v[24]')
+_ST_GLOBAL = re.compile(r'st\.global')
+_ST_GLOBAL_VEC = re.compile(r'st\.global(?:\.\w+)*\.v[24]')
+_FMA = re.compile(r'\bfma\.')
+_MUL = re.compile(r'\bmul\.')
+_ADD = re.compile(r'\badd\.')
+_LD_PARAM = re.compile(r'ld\.param')
+_PREFETCH = re.compile(r'prefetch\.global')
+_CACHE_HINT_LD = re.compile(r'ld\.global\.(?:cs|cg|ca|cv)')
+_CACHE_HINT_ST = re.compile(r'st\.global\.(?:wb|wt|cs)')
+_MAXNREG = re.compile(r'\.maxnreg\s+(\d+)')
+
+
+def extract_features_from_ptx(ptx_source):
+    """Extract 25 scalar features from PTX source text."""
+    n_instr = 0
+    n_ld_global = 0
+    n_ld_global_vec = 0
+    n_st_global = 0
+    n_st_global_vec = 0
+    n_fma = 0
+    n_mul = 0
+    n_add = 0
+    n_ld_param = 0
+    n_prefetch = 0
+    n_cache_hint_ld = 0
+    n_cache_hint_st = 0
+    n_branch = 0
+
+    # Parse register declarations
+    reg_decls = {}
+    for line in ptx_source.split('\n'):
+        m = re.search(r'\.reg\s+(\.\w+)\s+%\w+<(\d+)>\s*;', line)
+        if m:
+            reg_decls[m.group(1)] = int(m.group(2))
+
+    # Count instructions (lines between { and })
+    in_body = False
+    for line in ptx_source.split('\n'):
+        stripped = line.strip()
+
+        if stripped == '{':
+            in_body = True
+            continue
+        if stripped == '}':
+            in_body = False
+            continue
+        if not in_body:
+            continue
+
+        # Skip non-instructions
+        if not stripped or stripped.startswith('//') or stripped.startswith('.'):
+            continue
+        if stripped.endswith(':'):  # label
+            continue
+        if stripped in ('ret;', 'exit;', ')', ','):
+            continue
+
+        # Check for branch
+        if 'bra ' in stripped or 'bra\t' in stripped:
+            n_branch += 1
+            continue
+
+        n_instr += 1
+
+        if _LD_GLOBAL.search(line):
+            n_ld_global += 1
+            if _LD_GLOBAL_VEC.search(line):
+                n_ld_global_vec += 1
+            if _CACHE_HINT_LD.search(line):
+                n_cache_hint_ld += 1
+        if _ST_GLOBAL.search(line):
+            n_st_global += 1
+            if _ST_GLOBAL_VEC.search(line):
+                n_st_global_vec += 1
+            if _CACHE_HINT_ST.search(line):
+                n_cache_hint_st += 1
+        if _FMA.search(line):
+            n_fma += 1
+        if _MUL.search(line):
+            n_mul += 1
+        if _ADD.search(line):
+            n_add += 1
+        if _LD_PARAM.search(line):
+            n_ld_param += 1
+        if _PREFETCH.search(line):
+            n_prefetch += 1
+
+    maxnreg = 0
+    m = _MAXNREG.search(ptx_source)
+    if m:
+        maxnreg = int(m.group(1))
+
+    total_regs = sum(reg_decls.values())
+    n_f32_regs = reg_decls.get('.f32', 0)
+    n_b64_regs = reg_decls.get('.b64', 0)
+    n_total = max(n_instr, 1)
+    n_compute = n_fma + n_mul + n_add
+    n_mem = n_ld_global + n_st_global
+
+    return [
+        n_instr,
+        n_ld_global,
+        n_st_global,
+        n_fma,
+        n_ld_param,
+        n_prefetch,
+        n_branch,
+        n_ld_global_vec,
+        n_st_global_vec,
+        round(n_ld_global_vec / max(n_ld_global, 1), 4),  # vec_ld_ratio
+        round(n_st_global_vec / max(n_st_global, 1), 4),  # vec_st_ratio
+        n_cache_hint_ld,
+        n_cache_hint_st,
+        round(n_cache_hint_ld / max(n_ld_global, 1), 4),  # hint_ld_ratio
+        round(n_cache_hint_st / max(n_st_global, 1), 4),  # hint_st_ratio
+        round(n_ld_global / n_total, 4),  # load_ratio
+        round(n_st_global / n_total, 4),  # store_ratio
+        round(n_fma / n_total, 4),  # fma_ratio
+        round(n_compute / n_total, 4),  # compute_ratio
+        round(n_mem / n_total, 4),  # mem_ratio
+        round(n_compute / max(n_mem, 1), 4),  # compute_to_mem
+        total_regs,
+        n_f32_regs,
+        n_b64_regs,
+        maxnreg,
+    ]
+
+
+# ---------------------------------------------------------------------------
+# Action mask and history
+# ---------------------------------------------------------------------------
+
+def get_action_mask(applied_set):
+    mask = []
+    for label in ACTION_NAMES:
+        if label == "stop":
+            mask.append(1)
+            continue
+        if label in applied_set:
+            mask.append(0)
+            continue
+        conflict = False
+        for group_labels in CONFLICT_GROUPS.values():
+            if label in group_labels and applied_set & group_labels:
+                conflict = True
+                break
+        mask.append(0 if conflict else 1)
+    return mask
+
+
+def get_action_history(applied_set):
+    return [1 if name in applied_set else 0 for name in ACTION_NAMES]
+
+
+# ---------------------------------------------------------------------------
+# Inference
+# ---------------------------------------------------------------------------
+
+def load_model(checkpoint_path, device="cpu"):
+    """Load trained TransformPolicy from checkpoint."""
+    ckpt = torch.load(checkpoint_path, map_location=device, weights_only=False)
+    model = TransformPolicy(hidden=128)
+    model.load_state_dict(ckpt["policy"])
+    model.eval()
+    model.to(device)
+    epoch = ckpt.get("epoch", "unknown")
+    print(f"Loaded checkpoint from epoch {epoch}")
+    if "eval_result" in ckpt:
+        mean_imp = ckpt["eval_result"].get("mean_improvement", 0)
+        print(f"  Eval mean improvement: {mean_imp*100:.1f}%")
+    return model
+
+
+def predict_transforms(model, ptx_source, max_steps=6, verbose=True):
+    """Predict optimal transform sequence for a PTX kernel.
+
+    Returns list of transform labels (excluding 'stop').
+    """
+    features = extract_features_from_ptx(ptx_source)
+    applied = set()
+    actions = []
+
+    if verbose:
+        print(f"\nKernel: {features[0]} instructions, "
+              f"{features[1]} global loads, {features[2]} global stores, "
+              f"{features[3]} FMA, {features[21]} total regs")
+
+    for step in range(max_steps):
+        feat_t = torch.tensor(features, dtype=torch.float32)
+        mask = get_action_mask(applied)
+        mask_t = torch.tensor(mask, dtype=torch.float32)
+        hist = get_action_history(applied)
+        hist_t = torch.tensor(hist, dtype=torch.float32)
+
+        action_id = model.get_greedy_action(feat_t, mask_t, hist_t)
+        action_label = ACTION_NAMES[action_id]
+
+        if verbose:
+            probs = model.get_action_probs(feat_t, mask_t, hist_t)
+            top5 = torch.topk(probs, min(5, probs.size(0)))
+            top5_str = ", ".join(
+                f"{ACTION_NAMES[i]}={p:.2f}"
+                for p, i in zip(top5.values.tolist(), top5.indices.tolist())
+            )
+            print(f"  Step {step+1}: {action_label} (top5: {top5_str})")
+
+        if action_label == "stop":
+            break
+
+        actions.append(action_label)
+        applied.add(action_label)
+
+    if verbose:
+        print(f"\nPredicted sequence: {' -> '.join(actions) if actions else '(no transforms)'}")
+
+    return actions
+
+
+# ---------------------------------------------------------------------------
+# Main
+# ---------------------------------------------------------------------------
+
+def main():
+    parser = argparse.ArgumentParser(description="Chronos PoC: PTX transform inference")
+    parser.add_argument("--checkpoint", required=True, help="Path to .pt checkpoint")
+    parser.add_argument("--ptx", help="Path to PTX file")
+    parser.add_argument("--kernel", default="gemm_tile",
+                        help="Kernel type (for generating PTX if --ptx not provided)")
+    parser.add_argument("--m", type=int, default=4)
+    parser.add_argument("--n", type=int, default=6)
+    parser.add_argument("--k", type=int, default=8)
+    args = parser.parse_args()
+
+    model = load_model(args.checkpoint)
+
+    if args.ptx:
+        with open(args.ptx) as f:
+            ptx_source = f.read()
+        print(f"\nLoaded PTX from: {args.ptx}")
+    else:
+        print(f"\nTo run on a specific kernel, use: --ptx path/to/kernel.ptx")
+        print("Showing demo with a sample feature vector...")
+
+        # Demo: create a synthetic feature vector matching gemm_tile(4,6,8)
+        # (the best kernel from training: -53.8% improvement)
+        demo_features = [
+            170,  # n_instructions
+            16,   # n_ld_global
+            8,    # n_st_global
+            48,   # n_fma
+            12,   # n_ld_param
+            0,    # n_prefetch
+            2,    # n_branch
+            0,    # n_ld_global_vec
+            0,    # n_st_global_vec
+            0.0,  # vec_ld_ratio
+            0.0,  # vec_st_ratio
+            0,    # n_cache_hint_ld
+            0,    # n_cache_hint_st
+            0.0,  # hint_ld_ratio
+            0.0,  # hint_st_ratio
+            0.094,  # load_ratio
+            0.047,  # store_ratio
+            0.282,  # fma_ratio
+            0.388,  # compute_ratio
+            0.141,  # mem_ratio
+            2.75,   # compute_to_mem
+            95,   # total_regs
+            48,   # n_f32_regs
+            16,   # n_b64_regs
+            0,    # maxnreg
+        ]
+
+        applied = set()
+        actions = []
+        print(f"\nDemo: gemm_tile({args.m},{args.n},{args.k})-like features")
+        print(f"Features: {len(demo_features)} dims")
+
+        for step in range(6):
+            feat_t = torch.tensor(demo_features, dtype=torch.float32)
+            mask = get_action_mask(applied)
+            mask_t = torch.tensor(mask, dtype=torch.float32)
+            hist = get_action_history(applied)
+            hist_t = torch.tensor(hist, dtype=torch.float32)
+
+            action_id = model.get_greedy_action(feat_t, mask_t, hist_t)
+            action_label = ACTION_NAMES[action_id]
+
+            probs = model.get_action_probs(feat_t, mask_t, hist_t)
+            top3 = torch.topk(probs, min(3, probs.size(0)))
+            top3_str = ", ".join(
+                f"{ACTION_NAMES[i]}={p:.2f}"
+                for p, i in zip(top3.values.tolist(), top3.indices.tolist())
+            )
+            print(f"  Step {step+1}: {action_label} (probs: {top3_str})")
+
+            if action_label == "stop":
+                break
+            actions.append(action_label)
+            applied.add(action_label)
+
+        print(f"\nPredicted: {' -> '.join(actions)}")
+        print(f"Expected for gemm_tile(4,6,8): maxnreg_128 -> vec_ld -> vec_st -> stop")
+        return
+
+    actions = predict_transforms(model, ptx_source)
+    print(f"\nTo apply these transforms, use the Chronos transform pipeline.")
+
+
+if __name__ == "__main__":
+    main()

measure_triton_results.json

@@ -0,0 +1,2733 @@
+{
+  "n_kernels": 33,
+  "n_improved": 17,
+  "elapsed_s": 83.0,
+  "transform_stats": {
+    "cache_cs": {
+      "n": 29,
+      "mean_delta": -1.16,
+      "median_delta": -0.12,
+      "best": -13.13,
+      "worst": 1.09,
+      "improved_count": 6,
+      "degraded_count": 1,
+      "errors": 0
+    },
+    "cache_cg": {
+      "n": 29,
+      "mean_delta": 9.63,
+      "median_delta": 0.26,
+      "best": -3.28,
+      "worst": 96.43,
+      "improved_count": 4,
+      "degraded_count": 11,
+      "errors": 0
+    },
+    "cache_ca": {
+      "n": 29,
+      "mean_delta": -0.92,
+      "median_delta": -0.12,
+      "best": -10.58,
+      "worst": 1.15,
+      "improved_count": 5,
+      "degraded_count": 1,
+      "errors": 0
+    },
+    "cache_cv": {
+      "n": 29,
+      "mean_delta": 9.67,
+      "median_delta": 0.26,
+      "best": -3.37,
+      "worst": 96.17,
+      "improved_count": 4,
+      "degraded_count": 10,
+      "errors": 0
+    },
+    "st_cache_cs": {
+      "n": 33,
+      "mean_delta": -0.75,
+      "median_delta": -0.12,
+      "best": -10.99,
+      "worst": 1.2,
+      "improved_count": 5,
+      "degraded_count": 1,
+      "errors": 0
+    },
+    "st_cache_wt": {
+      "n": 33,
+      "mean_delta": 0.08,
+      "median_delta": 0.0,
+      "best": -2.98,
+      "worst": 5.72,
+      "improved_count": 1,
+      "degraded_count": 1,
+      "errors": 0
+    },
+    "st_cache_wb": {
+      "n": 33,
+      "mean_delta": -0.81,
+      "median_delta": -0.03,
+      "best": -11.5,
+      "worst": 1.15,
+      "improved_count": 5,
+      "degraded_count": 1,
+      "errors": 0
+    },
+    "maxnreg_32": {
+      "n": 33,
+      "mean_delta": 36.02,
+      "median_delta": 0.38,
+      "best": -4.26,
+      "worst": 366.65,
+      "improved_count": 5,
+      "degraded_count": 12,
+      "errors": 0
+    },
+    "maxnreg_64": {
+      "n": 33,
+      "mean_delta": 12.76,
+      "median_delta": 0.0,
+      "best": -7.28,
+      "worst": 191.22,
+      "improved_count": 5,
+      "degraded_count": 8,
+      "errors": 0
+    },
+    "maxnreg_128": {
+      "n": 33,
+      "mean_delta": 3.72,
+      "median_delta": 0.04,
+      "best": -3.88,
+      "worst": 89.88,
+      "improved_count": 2,
+      "degraded_count": 5,
+      "errors": 0
+    },
+    "maxnreg_255": {
+      "n": 33,
+      "mean_delta": -2.26,
+      "median_delta": -0.21,
+      "best": -46.95,
+      "worst": 1.07,
+      "improved_count": 7,
+      "degraded_count": 1,
+      "errors": 0
+    },
+    "reorder_cp": {
+      "n": 16,
+      "mean_delta": -0.06,
+      "median_delta": -0.04,
+      "best": -7.65,
+      "worst": 5.3,
+      "improved_count": 2,
+      "degraded_count": 3,
+      "errors": 14
+    },
+    "reorder_il": {
+      "n": 16,
+      "mean_delta": -2.31,
+      "median_delta": -0.15,
+      "best": -19.02,
+      "worst": 2.27,
+      "improved_count": 5,
+      "degraded_count": 1,
+      "errors": 14
+    },
+    "reorder_lf": {
+      "n": 16,
+      "mean_delta": -0.37,
+      "median_delta": 0.0,
+      "best": -9.98,
+      "worst": 3.42,
+      "improved_count": 2,
+      "degraded_count": 3,
+      "errors": 14
+    },
+    "reorder_sl": {
+      "n": 16,
+      "mean_delta": -0.89,
+      "median_delta": -0.23,
+      "best": -9.72,
+      "worst": 9.94,
+      "improved_count": 4,
+      "degraded_count": 2,
+      "errors": 14
+    }
+  },
+  "kernel_results": {
+    "triton_vector_add_256": {
+      "source": "triton_kernels",
+      "baseline": 944,
+      "baseline_std": 29.49187981462016,
+      "baseline_error": null,
+      "transforms": {
+        "cache_cs": {
+          "cycles": 941,
+          "std": 27.570924902875493,
+          "delta_pct": -0.32
+        },
+        "cache_cg": {
+          "cycles": 934,
+          "std": 27.46245436955699,
+          "delta_pct": -1.06
+        },
+        "cache_ca": {
+          "cycles": 937,
+          "std": 26.23790197405273,
+          "delta_pct": -0.74
+        },
+        "cache_cv": {
+          "cycles": 933,
+          "std": 28.072007409517404,
+          "delta_pct": -1.17
+        },
+        "st_cache_cs": {
+          "cycles": 939,
+          "std": 25.358580401907357,
+          "delta_pct": -0.53
+        },
+        "st_cache_wt": {
+          "cycles": 941,
+          "std": 27.90693417414389,
+          "delta_pct": -0.32
+        },
+        "st_cache_wb": {
+          "cycles": 937,
+          "std": 26.86858946800148,
+          "delta_pct": -0.74
+        },
+        "maxnreg_32": {
+          "cycles": 939,
+          "std": 28.326063969425757,
+          "delta_pct": -0.53
+        },
+        "maxnreg_64": {
+          "cycles": 940,
+          "std": 27.253641499806957,
+          "delta_pct": -0.42
+        },
+        "maxnreg_128": {
+          "cycles": 941,
+          "std": 23.985612354075933,
+          "delta_pct": -0.32
+        },
+        "maxnreg_255": {
+          "cycles": 938,
+          "std": 28.52321160037908,
+          "delta_pct": -0.64
+        },
+        "reorder_cp": {
+          "cycles": 941,
+          "std": 28.676645549994163,
+          "delta_pct": -0.32
+        },
+        "reorder_il": {
+          "cycles": 938,
+          "std": 25.202558897857973,
+          "delta_pct": -0.64
+        },
+        "reorder_lf": {
+          "cycles": 935,
+          "std": 27.466561033372926,
+          "delta_pct": -0.95
+        },
+        "reorder_sl": {
+          "cycles": 937,
+          "std": 30.11105735440056,
+          "delta_pct": -0.74
+        }
+      }
+    },
+    "triton_vector_add_512": {
+      "source": "triton_kernels",
+      "baseline": 1068,
+      "baseline_std": 60.74178442390378,
+      "baseline_error": null,
+      "transforms": {
+        "cache_cs": {
+          "cycles": 1074,
+          "std": 63.366669274942964,
+          "delta_pct": 0.56
+        },
+        "cache_cg": {
+          "cycles": 1033,
+          "std": 43.91783208447339,
+          "delta_pct": -3.28
+        },
+        "cache_ca": {
+          "cycles": 1062,
+          "std": 58.1985736509066,
+          "delta_pct": -0.56
+        },
+        "cache_cv": {
+          "cycles": 1032,
+          "std": 46.53583135606368,
+          "delta_pct": -3.37
+        },
+        "st_cache_cs": {
+          "cycles": 1076,
+          "std": 58.655327976237594,
+          "delta_pct": 0.75
+        },
+        "st_cache_wt": {
+          "cycles": 1070,
+          "std": 55.24308983936362,
+          "delta_pct": 0.19
+        },
+        "st_cache_wb": {
+          "cycles": 1072,
+          "std": 61.67726465238224,
+          "delta_pct": 0.37
+        },
+        "maxnreg_32": {
+          "cycles": 1065,
+          "std": 58.49837412270532,
+          "delta_pct": -0.28
+        },
+        "maxnreg_64": {
+          "cycles": 1071,
+          "std": 58.60743894762848,
+          "delta_pct": 0.28
+        },
+        "maxnreg_128": {
+          "cycles": 1075,
+          "std": 59.72745578877439,
+          "delta_pct": 0.66
+        },
+        "maxnreg_255": {
+          "cycles": 1066,
+          "std": 56.25154908978063,
+          "delta_pct": -0.19
+        },
+        "reorder_cp": {
+          "cycles": 1069,
+          "std": 61.26006019422442,
+          "delta_pct": 0.09
+        },
+        "reorder_il": {
+          "cycles": 1069,
+          "std": 61.40023697022675,
+          "delta_pct": 0.09
+        },
+        "reorder_lf": {
+          "cycles": 1070,
+          "std": 61.17932473475006,
+          "delta_pct": 0.19
+        },
+        "reorder_sl": {
+          "cycles": 1063,
+          "std": 59.396992979442985,
+          "delta_pct": -0.47
+        }
+      }
+    },
+    "triton_vector_add_1024": {
+      "source": "triton_kernels",
+      "baseline": 1219,
+      "baseline_std": 42.60126171840454,
+      "baseline_error": null,
+      "transforms": {
+        "cache_cs": {
+          "cycles": 1217,
+          "std": 55.49619356316251,
+          "delta_pct": -0.16
+        },
+        "cache_cg": {
+          "cycles": 1233,
+          "std": 38.63293898993448,
+          "delta_pct": 1.15
+        },
+        "cache_ca": {
+          "cycles": 1217,
+          "std": 37.32690048745007,
+          "delta_pct": -0.16
+        },
+        "cache_cv": {
+          "cycles": 1231,
+          "std": 55.58792584725571,
+          "delta_pct": 0.98
+        },
+        "st_cache_cs": {
+          "cycles": 1220,
+          "std": 29.24901323121859,
+          "delta_pct": 0.08
+        },
+        "st_cache_wt": {
+          "cycles": 1219,
+          "std": 51.917569039776886,
+          "delta_pct": 0.0
+        },
+        "st_cache_wb": {
+          "cycles": 1218,
+          "std": 55.40694879705974,
+          "delta_pct": -0.08
+        },
+        "maxnreg_32": {
+          "cycles": 1219,
+          "std": 41.19195036654613,
+          "delta_pct": 0.0
+        },
+        "maxnreg_64": {
+          "cycles": 1219,
+          "std": 56.34494121036954,
+          "delta_pct": 0.0
+        },
+        "maxnreg_128": {
+          "cycles": 1215,
+          "std": 37.623588345611054,
+          "delta_pct": -0.33
+        },
+        "maxnreg_255": {
+          "cycles": 1218,
+          "std": 42.37277309782781,
+          "delta_pct": -0.08
+        },
+        "reorder_cp": {
+          "cycles": 1217,
+          "std": 47.77099302924318,
+          "delta_pct": -0.16
+        },
+        "reorder_il": {
+          "cycles": 1217,
+          "std": 35.28370693393766,
+          "delta_pct": -0.16
+        },
+        "reorder_lf": {
+          "cycles": 1220,
+          "std": 56.07920091263784,
+          "delta_pct": 0.08
+        },
+        "reorder_sl": {
+          "cycles": 1220,
+          "std": 46.88487922560961,
+          "delta_pct": 0.08
+        }
+      }
+    },
+    "triton_softmax_1024": {
+      "source": "triton_kernels",
+      "baseline": 3838,
+      "baseline_std": 84.44705974751281,
+      "baseline_error": null,
+      "transforms": {
+        "cache_cs": {
+          "cycles": 3880,
+          "std": 76.23343361019494,
+          "delta_pct": 1.09
+        },
+        "cache_cg": {
+          "cycles": 3843,
+          "std": 62.285206710743125,
+          "delta_pct": 0.13
+        },
+        "cache_ca": {
+          "cycles": 3882,
+          "std": 79.20490499331466,
+          "delta_pct": 1.15
+        },
+        "cache_cv": {
+          "cycles": 3850,
+          "std": 58.920025246091,
+          "delta_pct": 0.31
+        },
+        "st_cache_cs": {
+          "cycles": 3884,
+          "std": 72.98933466610036,
+          "delta_pct": 1.2
+        },
+        "st_cache_wt": {
+          "cycles": 3842,
+          "std": 85.59785496728291,
+          "delta_pct": 0.1
+        },
+        "st_cache_wb": {
+          "cycles": 3882,
+          "std": 74.53467045610385,
+          "delta_pct": 1.15
+        },
+        "maxnreg_32": {
+          "cycles": 3802,
+          "std": 83.56720035396663,
+          "delta_pct": -0.94
+        },
+        "maxnreg_64": {
+          "cycles": 3889,
+          "std": 72.41214262815319,
+          "delta_pct": 1.33
+        },
+        "maxnreg_128": {
+          "cycles": 3848,
+          "std": 86.11081000664201,
+          "delta_pct": 0.26
+        },
+        "maxnreg_255": {
+          "cycles": 3879,
+          "std": 72.32222600971295,
+          "delta_pct": 1.07
+        },
+        "reorder_cp": {
+          "error": "skipped_barrier"
+        },
+        "reorder_il": {
+          "error": "skipped_barrier"
+        },
+        "reorder_lf": {
+          "error": "skipped_barrier"
+        },
+        "reorder_sl": {
+          "error": "skipped_barrier"
+        }
+      }
+    },
+    "triton_layernorm_1024": {
+      "source": "triton_kernels",
+      "baseline": 4556,
+      "baseline_std": 112.65114469014507,
+      "baseline_error": null,
+      "transforms": {
+        "cache_cs": {
+          "cycles": 4523,
+          "std": 100.84496008725473,
+          "delta_pct": -0.72
+        },
+        "cache_cg": {
+          "cycles": 6911,
+          "std": 457.00405643713924,
+          "delta_pct": 51.69
+        },
+        "cache_ca": {
+          "cycles": 4529,
+          "std": 121.65385772757064,
+          "delta_pct": -0.59
+        },
+        "cache_cv": {
+          "cycles": 6876,
+          "std": 449.7087875503435,
+          "delta_pct": 50.92
+        },
+        "st_cache_cs": {
+          "cycles": 4507,
+          "std": 91.85497863480236,
+          "delta_pct": -1.08
+        },
+        "st_cache_wt": {
+          "cycles": 4562,
+          "std": 106.59615893642697,
+          "delta_pct": 0.13
+        },
+        "st_cache_wb": {
+          "cycles": 4534,
+          "std": 115.7763961263262,
+          "delta_pct": -0.48
+        },
+        "maxnreg_32": {
+          "cycles": 4567,
+          "std": 110.52709520746485,
+          "delta_pct": 0.24
+        },
+        "maxnreg_64": {
+          "cycles": 4523,
+          "std": 111.35946019535116,
+          "delta_pct": -0.72
+        },
+        "maxnreg_128": {
+          "cycles": 4565,
+          "std": 125.39940580002762,
+          "delta_pct": 0.2
+        },
+        "maxnreg_255": {
+          "cycles": 4527,
+          "std": 115.90962761996951,
+          "delta_pct": -0.64
+        },
+        "reorder_cp": {
+          "error": "skipped_barrier"
+        },
+        "reorder_il": {
+          "error": "skipped_barrier"
+        },
+        "reorder_lf": {
+          "error": "skipped_barrier"
+        },
+        "reorder_sl": {
+          "error": "skipped_barrier"
+        }
+      }
+    },
+    "triton_matmul_64x64x32_w4": {
+      "source": "triton_kernels",
+      "baseline": 171410,
+      "baseline_std": 501.0587783633772,
+      "baseline_error": null,
+      "transforms": {
+        "st_cache_cs": {
+          "cycles": 171382,
+          "std": 432.3299073335084,
+          "delta_pct": -0.02
+        },
+        "st_cache_wt": {
+          "cycles": 171423,
+          "std": 470.07234483109085,
+          "delta_pct": 0.01
+        },
+        "st_cache_wb": {
+          "cycles": 171367,
+          "std": 460.0975798404073,
+          "delta_pct": -0.03
+        },
+        "maxnreg_32": {
+          "cycles": 425115,
+          "std": 4083.747927991516,
+          "delta_pct": 148.01
+        },
+        "maxnreg_64": {
+          "cycles": 245521,
+          "std": 4866.736257041262,
+          "delta_pct": 43.24
+        },
+        "maxnreg_128": {
+          "cycles": 194120,
+          "std": 4007.7541684059165,
+          "delta_pct": 13.25
+        },
+        "maxnreg_255": {
+          "cycles": 171367,
+          "std": 462.19168531573564,
+          "delta_pct": -0.03
+        },
+        "reorder_cp": {
+          "error": "skipped_barrier"
+        },
+        "reorder_il": {
+          "error": "skipped_barrier"
+        },
+        "reorder_lf": {
+          "error": "skipped_barrier"
+        },
+        "reorder_sl": {
+          "error": "skipped_barrier"
+        }
+      }
+    },
+    "triton_matmul_64x64x32_w8": {
+      "source": "triton_kernels",
+      "baseline": 186277,
+      "baseline_std": 555.2674566368895,
+      "baseline_error": null,
+      "transforms": {
+        "st_cache_cs": {
+          "cycles": 186202,
+          "std": 563.2344040450654,
+          "delta_pct": -0.04
+        },
+        "st_cache_wt": {
+          "cycles": 186236,
+          "std": 542.4974524364146,
+          "delta_pct": -0.02
+        },
+        "st_cache_wb": {
+          "cycles": 186132,
+          "std": 537.6838005510301,
+          "delta_pct": -0.08
+        },
+        "maxnreg_32": {
+          "cycles": 319430,
+          "std": 4796.146295076392,
+          "delta_pct": 71.48
+        },
+        "maxnreg_64": {
+          "cycles": 208291,
+          "std": 4566.849550376605,
+          "delta_pct": 11.82
+        },
+        "maxnreg_128": {
+          "cycles": 186231,
+          "std": 542.7691460464569,
+          "delta_pct": -0.02
+        },
+        "maxnreg_255": {
+          "cycles": 185762,
+          "std": 473.6743501605296,
+          "delta_pct": -0.28
+        },
+        "reorder_cp": {
+          "error": "skipped_barrier"
+        },
+        "reorder_il": {
+          "error": "skipped_barrier"
+        },
+        "reorder_lf": {
+          "error": "skipped_barrier"
+        },
+        "reorder_sl": {
+          "error": "skipped_barrier"
+        }
+      }
+    },
+    "triton_matmul_128x128x32_w4": {
+      "source": "triton_kernels",
+      "baseline": 240682,
+      "baseline_std": 342.63614286878726,
+      "baseline_error": null,
+      "transforms": {
+        "st_cache_cs": {
+          "cycles": 240657,
+          "std": 354.66443206924487,
+          "delta_pct": -0.01
+        },
+        "st_cache_wt": {
+          "cycles": 240660,
+          "std": 354.1097540311479,
+          "delta_pct": -0.01
+        },
+        "st_cache_wb": {
+          "cycles": 240703,
+          "std": 353.70178395789867,
+          "delta_pct": 0.01
+        },
+        "maxnreg_32": {
+          "cycles": 1123144,
+          "std": 1845.049125064154,
+          "delta_pct": 366.65
+        },
+        "maxnreg_64": {
+          "cycles": 700917,
+          "std": 1328.5777667867244,
+          "delta_pct": 191.22
+        },
+        "maxnreg_128": {
+          "cycles": 457011,
+          "std": 1192.7186457731766,
+          "delta_pct": 89.88
+        },
+        "maxnreg_255": {
+          "cycles": 238426,
+          "std": 471.0170941430894,
+          "delta_pct": -0.94
+        }
+      }
+    },
+    "triton_matmul_128x128x32_w8": {
+      "source": "triton_kernels",
+      "baseline": 211438,
+      "baseline_std": 210.94903596603612,
+      "baseline_error": null,
+      "transforms": {
+        "st_cache_cs": {
+          "cycles": 211484,
+          "std": 207.6386823306293,
+          "delta_pct": 0.02
+        },
+        "st_cache_wt": {
+          "cycles": 211417,
+          "std": 192.11210763509936,
+          "delta_pct": -0.01
+        },
+        "st_cache_wb": {
+          "cycles": 211456,
+          "std": 205.09111999304113,
+          "delta_pct": 0.01
+        },
+        "maxnreg_32": {
+          "cycles": 710161,
+          "std": 1178.3205980440978,
+          "delta_pct": 235.87
+        },
+        "maxnreg_64": {
+          "cycles": 425841,
+          "std": 724.849326946642,
+          "delta_pct": 101.4
+        },
+        "maxnreg_128": {
+          "cycles": 244024,
+          "std": 483.0836308290729,
+          "delta_pct": 15.41
+        },
+        "maxnreg_255": {
+          "cycles": 211445,
+          "std": 217.40122119022237,
+          "delta_pct": 0.0
+        },
+        "reorder_cp": {
+          "error": "skipped_barrier"
+        },
+        "reorder_il": {
+          "error": "skipped_barrier"
+        },
+        "reorder_lf": {
+          "error": "skipped_barrier"
+        },
+        "reorder_sl": {
+          "error": "skipped_barrier"
+        }
+      }
+    },
+    "triton_fused_add_mul_256": {
+      "source": "triton_kernels",
+      "baseline": 1332,
+      "baseline_std": 21.35080267811962,
+      "baseline_error": null,
+      "transforms": {
+        "cache_cs": {
+          "cycles": 1331,
+          "std": 34.96618938346013,
+          "delta_pct": -0.08
+        },
+        "cache_cg": {
+          "cycles": 1332,
+          "std": 23.14809711401782,
+          "delta_pct": 0.0
+        },
+        "cache_ca": {
+          "cycles": 1328,
+          "std": 33.876983336773066,
+          "delta_pct": -0.3
+        },
+        "cache_cv": {
+          "cycles": 1329,
+          "std": 23.705819960507586,
+          "delta_pct": -0.23
+        },
+        "st_cache_cs": {
+          "cycles": 1326,
+          "std": 32.04271524075324,
+          "delta_pct": -0.45
+        },
+        "st_cache_wt": {
+          "cycles": 1332,
+          "std": 23.455223298873108,
+          "delta_pct": 0.0
+        },
+        "st_cache_wb": {
+          "cycles": 1332,
+          "std": 37.35344662812255,
+          "delta_pct": 0.0
+        },
+        "maxnreg_32": {
+          "cycles": 1337,
+          "std": 22.76026746327907,
+          "delta_pct": 0.38
+        },
+        "maxnreg_64": {
+          "cycles": 1326,
+          "std": 34.47232186842076,
+          "delta_pct": -0.45
+        },
+        "maxnreg_128": {
+          "cycles": 1334,
+          "std": 24.19508162829793,
+          "delta_pct": 0.15
+        },
+        "maxnreg_255": {
+          "cycles": 1330,
+          "std": 35.5125879090781,
+          "delta_pct": -0.15
+        },
+        "reorder_cp": {
+          "cycles": 1336,
+          "std": 20.016243403795826,
+          "delta_pct": 0.3
+        },
+        "reorder_il": {
+          "cycles": 1330,
+          "std": 35.21678435064735,
+          "delta_pct": -0.15
+        },
+        "reorder_lf": {
+          "cycles": 1335,
+          "std": 20.668061713668266,
+          "delta_pct": 0.23
+        },
+        "reorder_sl": {
+          "cycles": 1330,
+          "std": 33.829409912086845,
+          "delta_pct": -0.15
+        }
+      }
+    },
+    "triton_fused_add_mul_512": {
+      "source": "triton_kernels",
+      "baseline": 1379,
+      "baseline_std": 59.30477552440444,
+      "baseline_error": null,
+      "transforms": {
+        "cache_cs": {
+          "cycles": 1363,
+          "std": 50.025993243512914,
+          "delta_pct": -1.16
+        },
+        "cache_cg": {
+          "cycles": 1356,
+          "std": 58.66862449384679,
+          "delta_pct": -1.67
+        },
+        "cache_ca": {
+          "cycles": 1356,
+          "std": 51.74146886202595,
+          "delta_pct": -1.67
+        },
+        "cache_cv": {
+          "cycles": 1355,
+          "std": 57.26626493844347,
+          "delta_pct": -1.74
+        },
+        "st_cache_cs": {
+          "cycles": 1354,
+          "std": 50.8162766739162,
+          "delta_pct": -1.81
+        },
+        "st_cache_wt": {
+          "cycles": 1380,
+          "std": 60.92277386166851,
+          "delta_pct": 0.07
+        },
+        "st_cache_wb": {
+          "cycles": 1359,
+          "std": 48.36108326950504,
+          "delta_pct": -1.45
+        },
+        "maxnreg_32": {
+          "cycles": 1376,
+          "std": 59.94648425887877,
+          "delta_pct": -0.22
+        },
+        "maxnreg_64": {
+          "cycles": 1355,
+          "std": 48.759710571331325,
+          "delta_pct": -1.74
+        },
+        "maxnreg_128": {
+          "cycles": 1375,
+          "std": 62.22318780004766,
+          "delta_pct": -0.29
+        },
+        "maxnreg_255": {
+          "cycles": 1361,
+          "std": 51.13456365316908,
+          "delta_pct": -1.31
+        },
+        "reorder_cp": {
+          "cycles": 1373,
+          "std": 60.520011359879966,
+          "delta_pct": -0.44
+        },
+        "reorder_il": {
+          "cycles": 1356,
+          "std": 48.52316946573049,
+          "delta_pct": -1.67
+        },
+        "reorder_lf": {
+          "cycles": 1374,
+          "std": 61.745087051521764,
+          "delta_pct": -0.36
+        },
+        "reorder_sl": {
+          "cycles": 1357,
+          "std": 51.92411674742287,
+          "delta_pct": -1.6
+        }
+      }
+    },
+    "triton_fused_add_mul_1024": {
+      "source": "triton_kernels",
+      "baseline": 1647,
+      "baseline_std": 126.03052715513016,
+      "baseline_error": null,
+      "transforms": {
+        "cache_cs": {
+          "cycles": 1620,
+          "std": 163.05165676864493,
+          "delta_pct": -1.64
+        },
+        "cache_cg": {
+          "cycles": 1652,
+          "std": 136.04815728263281,
+          "delta_pct": 0.3
+        },
+        "cache_ca": {
+          "cycles": 1636,
+          "std": 157.5926254619803,
+          "delta_pct": -0.67
+        },
+        "cache_cv": {
+          "cycles": 1651,
+          "std": 137.32488157650093,
+          "delta_pct": 0.24
+        },
+        "st_cache_cs": {
+          "cycles": 1638,
+          "std": 164.70073398439973,
+          "delta_pct": -0.55
+        },
+        "st_cache_wt": {
+          "cycles": 1645,
+          "std": 132.36343112808765,
+          "delta_pct": -0.12
+        },
+        "st_cache_wb": {
+          "cycles": 1636,
+          "std": 158.1481551425751,
+          "delta_pct": -0.67
+        },
+        "maxnreg_32": {
+          "cycles": 1645,
+          "std": 131.74864363628188,
+          "delta_pct": -0.12
+        },
+        "maxnreg_64": {
+          "cycles": 1633,
+          "std": 164.37551247068404,
+          "delta_pct": -0.85
+        },
+        "maxnreg_128": {
+          "cycles": 1644,
+          "std": 136.379041553312,
+          "delta_pct": -0.18
+        },
+        "maxnreg_255": {
+          "cycles": 1637,
+          "std": 164.06369182424245,
+          "delta_pct": -0.61
+        },
+        "reorder_cp": {
+          "cycles": 1645,
+          "std": 142.6154956517699,
+          "delta_pct": -0.12
+        },
+        "reorder_il": {
+          "cycles": 1638,
+          "std": 166.29074658561132,
+          "delta_pct": -0.55
+        },
+        "reorder_lf": {
+          "cycles": 1644,
+          "std": 132.60286233713055,
+          "delta_pct": -0.18
+        },
+        "reorder_sl": {
+          "cycles": 1634,
+          "std": 162.15513983528243,
+          "delta_pct": -0.79
+        }
+      }
+    },
+    "reduction_sum_1024": {
+      "source": "diverse",
+      "baseline": 2283,
+      "baseline_std": 28.91735810892828,
+      "baseline_error": null,
+      "transforms": {
+        "cache_cs": {
+          "cycles": 2279,
+          "std": 29.059206802664107,
+          "delta_pct": -0.18
+        },
+        "cache_cg": {
+          "cycles": 2261,
+          "std": 34.178703530122384,
+          "delta_pct": -0.96
+        },
+        "cache_ca": {
+          "cycles": 2280,
+          "std": 30.294712409923946,
+          "delta_pct": -0.13
+        },
+        "cache_cv": {
+          "cycles": 2257,
+          "std": 36.78278265710739,
+          "delta_pct": -1.14
+        },
+        "st_cache_cs": {
+          "cycles": 2280,
+          "std": 29.946478924908686,
+          "delta_pct": -0.13
+        },
+        "st_cache_wt": {
+          "cycles": 2287,
+          "std": 28.96892775026373,
+          "delta_pct": 0.18
+        },
+        "st_cache_wb": {
+          "cycles": 2287,
+          "std": 31.985745262538437,
+          "delta_pct": 0.18
+        },
+        "maxnreg_32": {
+          "cycles": 2518,
+          "std": 45.40396898950575,
+          "delta_pct": 10.29
+        },
+        "maxnreg_64": {
+          "cycles": 2283,
+          "std": 29.030967948037834,
+          "delta_pct": 0.0
+        },
+        "maxnreg_128": {
+          "cycles": 2284,
+          "std": 27.902578733873323,
+          "delta_pct": 0.04
+        },
+        "maxnreg_255": {
+          "cycles": 2286,
+          "std": 30.912740092072067,
+          "delta_pct": 0.13
+        },
+        "reorder_cp": {
+          "error": "skipped_barrier"
+        },
+        "reorder_il": {
+          "error": "skipped_barrier"
+        },
+        "reorder_lf": {
+          "error": "skipped_barrier"
+        },
+        "reorder_sl": {
+          "error": "skipped_barrier"
+        }
+      }
+    },
+    "reduction_sum_512": {
+      "source": "diverse",
+      "baseline": 1442,
+      "baseline_std": 19.446850644770223,
+      "baseline_error": null,
+      "transforms": {
+        "cache_cs": {
+          "cycles": 1436,
+          "std": 21.160753767292885,
+          "delta_pct": -0.42
+        },
+        "cache_cg": {
+          "cycles": 1440,
+          "std": 17.872425688753054,
+          "delta_pct": -0.14
+        },
+        "cache_ca": {
+          "cycles": 1437,
+          "std": 18.555635262636525,
+          "delta_pct": -0.35
+        },
+        "cache_cv": {
+          "cycles": 1440,
+          "std": 22.142944135773813,
+          "delta_pct": -0.14
+        },
+        "st_cache_cs": {
+          "cycles": 1438,
+          "std": 20.739546161861885,
+          "delta_pct": -0.28
+        },
+        "st_cache_wt": {
+          "cycles": 1441,
+          "std": 19.672681566070246,
+          "delta_pct": -0.07
+        },
+        "st_cache_wb": {
+          "cycles": 1440,
+          "std": 20.03422072355199,
+          "delta_pct": -0.14
+        },
+        "maxnreg_32": {
+          "cycles": 1513,
+          "std": 17.36521738994361,
+          "delta_pct": 4.92
+        },
+        "maxnreg_64": {
+          "cycles": 1438,
+          "std": 18.31485667429587,
+          "delta_pct": -0.28
+        },
+        "maxnreg_128": {
+          "cycles": 1442,
+          "std": 23.397187865211496,
+          "delta_pct": 0.0
+        },
+        "maxnreg_255": {
+          "cycles": 1439,
+          "std": 19.986604889275217,
+          "delta_pct": -0.21
+        },
+        "reorder_cp": {
+          "error": "skipped_barrier"
+        },
+        "reorder_il": {
+          "error": "skipped_barrier"
+        },
+        "reorder_lf": {
+          "error": "skipped_barrier"
+        },
+        "reorder_sl": {
+          "error": "skipped_barrier"
+        }
+      }
+    },
+    "reduction_max_1024": {
+      "source": "diverse",
+      "baseline": 2287,
+      "baseline_std": 29.404587652269498,
+      "baseline_error": null,
+      "transforms": {
+        "cache_cs": {
+          "cycles": 2288,
+          "std": 29.20560049031692,
+          "delta_pct": 0.04
+        },
+        "cache_cg": {
+          "cycles": 2261,
+          "std": 38.18233099222728,
+          "delta_pct": -1.14
+        },
+        "cache_ca": {
+          "cycles": 2290,
+          "std": 31.04470325192367,
+          "delta_pct": 0.13
+        },
+        "cache_cv": {
+          "cycles": 2265,
+          "std": 36.98072335690582,
+          "delta_pct": -0.96
+        },
+        "st_cache_cs": {
+          "cycles": 2283,
+          "std": 33.27393837525098,
+          "delta_pct": -0.17
+        },
+        "st_cache_wt": {
+          "cycles": 2285,
+          "std": 30.970121084684184,
+          "delta_pct": -0.09
+        },
+        "st_cache_wb": {
+          "cycles": 2294,
+          "std": 31.33937459490856,
+          "delta_pct": 0.31
+        },
+        "maxnreg_32": {
+          "cycles": 2525,
+          "std": 47.04580321346421,
+          "delta_pct": 10.41
+        },
+        "maxnreg_64": {
+          "cycles": 2285,
+          "std": 28.196409345872393,
+          "delta_pct": -0.09
+        },
+        "maxnreg_128": {
+          "cycles": 2289,
+          "std": 29.040254819818646,
+          "delta_pct": 0.09
+        },
+        "maxnreg_255": {
+          "cycles": 2285,
+          "std": 31.75385960792798,
+          "delta_pct": -0.09
+        },
+        "reorder_cp": {
+          "error": "skipped_barrier"
+        },
+        "reorder_il": {
+          "error": "skipped_barrier"
+        },
+        "reorder_lf": {
+          "error": "skipped_barrier"
+        },
+        "reorder_sl": {
+          "error": "skipped_barrier"
+        }
+      }
+    },
+    "reduction_max_512": {
+      "source": "diverse",
+      "baseline": 1443,
+      "baseline_std": 21.411211432331427,
+      "baseline_error": null,
+      "transforms": {
+        "cache_cs": {
+          "cycles": 1440,
+          "std": 19.682082587978336,
+          "delta_pct": -0.21
+        },
+        "cache_cg": {
+          "cycles": 1444,
+          "std": 20.720013875477978,
+          "delta_pct": 0.07
+        },
+        "cache_ca": {
+          "cycles": 1442,
+          "std": 19.579233897167683,
+          "delta_pct": -0.07
+        },
+        "cache_cv": {
+          "cycles": 1441,
+          "std": 19.73320488415402,
+          "delta_pct": -0.14
+        },
+        "st_cache_cs": {
+          "cycles": 1440,
+          "std": 18.96362570818144,
+          "delta_pct": -0.21
+        },
+        "st_cache_wt": {
+          "cycles": 1444,
+          "std": 23.070044971781048,
+          "delta_pct": 0.07
+        },
+        "st_cache_wb": {
+          "cycles": 1443,
+          "std": 19.77969666097031,
+          "delta_pct": 0.0
+        },
+        "maxnreg_32": {
+          "cycles": 1514,
+          "std": 18.103860776088617,
+          "delta_pct": 4.92
+        },
+        "maxnreg_64": {
+          "cycles": 1442,
+          "std": 19.005459741874176,
+          "delta_pct": -0.07
+        },
+        "maxnreg_128": {
+          "cycles": 1442,
+          "std": 21.203138918565806,
+          "delta_pct": -0.07
+        },
+        "maxnreg_255": {
+          "cycles": 1440,
+          "std": 19.25943340288078,
+          "delta_pct": -0.21
+        },
+        "reorder_cp": {
+          "error": "skipped_barrier"
+        },
+        "reorder_il": {
+          "error": "skipped_barrier"
+        },
+        "reorder_lf": {
+          "error": "skipped_barrier"
+        },
+        "reorder_sl": {
+          "error": "skipped_barrier"
+        }
+      }
+    },
+    "prefix_scan_1024": {
+      "source": "diverse",
+      "baseline": 1914,
+      "baseline_std": 18.85006896539108,
+      "baseline_error": null,
+      "transforms": {
+        "cache_cs": {
+          "cycles": 1913,
+          "std": 19.857842279563002,
+          "delta_pct": -0.05
+        },
+        "cache_cg": {
+          "cycles": 1951,
+          "std": 26.21342404189121,
+          "delta_pct": 1.93
+        },
+        "cache_ca": {
+          "cycles": 1909,
+          "std": 19.845399970774082,
+          "delta_pct": -0.26
+        },
+        "cache_cv": {
+          "cycles": 1950,
+          "std": 25.577528711742268,
+          "delta_pct": 1.88
+        },
+        "st_cache_cs": {
+          "cycles": 1910,
+          "std": 18.901454309126585,
+          "delta_pct": -0.21
+        },
+        "st_cache_wt": {
+          "cycles": 1913,
+          "std": 18.244280610646175,
+          "delta_pct": -0.05
+        },
+        "st_cache_wb": {
+          "cycles": 1909,
+          "std": 20.177750122350115,
+          "delta_pct": -0.26
+        },
+        "maxnreg_32": {
+          "cycles": 1919,
+          "std": 21.419439768584052,
+          "delta_pct": 0.26
+        },
+        "maxnreg_64": {
+          "cycles": 1914,
+          "std": 24.30736462473873,
+          "delta_pct": 0.0
+        },
+        "maxnreg_128": {
+          "cycles": 1918,
+          "std": 22.96361469803916,
+          "delta_pct": 0.21
+        },
+        "maxnreg_255": {
+          "cycles": 1916,
+          "std": 23.2384331442548,
+          "delta_pct": 0.1
+        },
+        "reorder_cp": {
+          "error": "skipped_barrier"
+        },
+        "reorder_il": {
+          "error": "skipped_barrier"
+        },
+        "reorder_lf": {
+          "error": "skipped_barrier"
+        },
+        "reorder_sl": {
+          "error": "skipped_barrier"
+        }
+      }
+    },
+    "prefix_scan_512": {
+      "source": "diverse",
+      "baseline": 1470,
+      "baseline_std": 24.63637097869733,
+      "baseline_error": null,
+      "transforms": {
+        "cache_cs": {
+          "cycles": 1467,
+          "std": 29.877576876313114,
+          "delta_pct": -0.2
+        },
+        "cache_cg": {
+          "cycles": 1473,
+          "std": 36.70606353179267,
+          "delta_pct": 0.2
+        },
+        "cache_ca": {
+          "cycles": 1474,
+          "std": 33.85005022152848,
+          "delta_pct": 0.27
+        },
+        "cache_cv": {
+          "cycles": 1467,
+          "std": 29.68908722072809,
+          "delta_pct": -0.2
+        },
+        "st_cache_cs": {
+          "cycles": 1468,
+          "std": 29.363446238478208,
+          "delta_pct": -0.14
+        },
+        "st_cache_wt": {
+          "cycles": 1470,
+          "std": 21.9029472902621,
+          "delta_pct": 0.0
+        },
+        "st_cache_wb": {
+          "cycles": 1474,
+          "std": 24.60645240582234,
+          "delta_pct": 0.27
+        },
+        "maxnreg_32": {
+          "cycles": 1478,
+          "std": 20.00399960007998,
+          "delta_pct": 0.54
+        },
+        "maxnreg_64": {
+          "cycles": 1470,
+          "std": 26.488219551340173,
+          "delta_pct": 0.0
+        },
+        "maxnreg_128": {
+          "cycles": 1474,
+          "std": 30.10612852892248,
+          "delta_pct": 0.27
+        },
+        "maxnreg_255": {
+          "cycles": 1474,
+          "std": 30.13604320411026,
+          "delta_pct": 0.27
+        },
+        "reorder_cp": {
+          "error": "skipped_barrier"
+        },
+        "reorder_il": {
+          "error": "skipped_barrier"
+        },
+        "reorder_lf": {
+          "error": "skipped_barrier"
+        },
+        "reorder_sl": {
+          "error": "skipped_barrier"
+        }
+      }
+    },
+    "attention_d64_kv64": {
+      "source": "diverse",
+      "baseline": 53610,
+      "baseline_std": 93.32397320624536,
+      "baseline_error": null,
+      "transforms": {
+        "cache_cs": {
+          "cycles": 53613,
+          "std": 102.94325123581439,
+          "delta_pct": 0.01
+        },
+        "cache_cg": {
+          "cycles": 58146,
+          "std": 188.2566184759516,
+          "delta_pct": 8.46
+        },
+        "cache_ca": {
+          "cycles": 53613,
+          "std": 101.8963491740504,
+          "delta_pct": 0.01
+        },
+        "cache_cv": {
+          "cycles": 58174,
+          "std": 181.6218515900551,
+          "delta_pct": 8.51
+        },
+        "st_cache_cs": {
+          "cycles": 53615,
+          "std": 101.64258150499721,
+          "delta_pct": 0.01
+        },
+        "st_cache_wt": {
+          "cycles": 53611,
+          "std": 101.52303925710656,
+          "delta_pct": 0.0
+        },
+        "st_cache_wb": {
+          "cycles": 53589,
+          "std": 97.34385239962512,
+          "delta_pct": -0.04
+        },
+        "maxnreg_32": {
+          "cycles": 116641,
+          "std": 835.3467224901286,
+          "delta_pct": 117.57
+        },
+        "maxnreg_64": {
+          "cycles": 61921,
+          "std": 213.7253255349024,
+          "delta_pct": 15.5
+        },
+        "maxnreg_128": {
+          "cycles": 53249,
+          "std": 162.7886285954888,
+          "delta_pct": -0.67
+        },
+        "maxnreg_255": {
+          "cycles": 28440,
+          "std": 22.207642265670618,
+          "delta_pct": -46.95
+        }
+      }
+    },
+    "attention_d64_kv32": {
+      "source": "diverse",
+      "baseline": 56910,
+      "baseline_std": 182.19544992946447,
+      "baseline_error": null,
+      "transforms": {
+        "cache_cs": {
+          "cycles": 56888,
+          "std": 181.49333321089236,
+          "delta_pct": -0.04
+        },
+        "cache_cg": {
+          "cycles": 59204,
+          "std": 185.004154277681,
+          "delta_pct": 4.03
+        },
+        "cache_ca": {
+          "cycles": 56881,
+          "std": 194.19240452448184,
+          "delta_pct": -0.05
+        },
+        "cache_cv": {
+          "cycles": 59155,
+          "std": 197.05218845524146,
+          "delta_pct": 3.94
+        },
+        "st_cache_cs": {
+          "cycles": 56877,
+          "std": 198.4987042778869,
+          "delta_pct": -0.06
+        },
+        "st_cache_wt": {
+          "cycles": 56894,
+          "std": 179.84295864725982,
+          "delta_pct": -0.03
+        },
+        "st_cache_wb": {
+          "cycles": 56919,
+          "std": 187.93011227581385,
+          "delta_pct": 0.02
+        },
+        "maxnreg_32": {
+          "cycles": 72623,
+          "std": 203.6142676729703,
+          "delta_pct": 27.61
+        },
+        "maxnreg_64": {
+          "cycles": 55899,
+          "std": 111.39286512160463,
+          "delta_pct": -1.78
+        },
+        "maxnreg_128": {
+          "cycles": 56444,
+          "std": 115.63091109214697,
+          "delta_pct": -0.82
+        },
+        "maxnreg_255": {
+          "cycles": 56451,
+          "std": 98.92903504532934,
+          "delta_pct": -0.81
+        },
+        "reorder_cp": {
+          "error": "skipped_barrier"
+        },
+        "reorder_il": {
+          "error": "skipped_barrier"
+        },
+        "reorder_lf": {
+          "error": "skipped_barrier"
+        },
+        "reorder_sl": {
+          "error": "skipped_barrier"
+        }
+      }
+    },
+    "attention_d128_kv64": {
+      "source": "diverse",
+      "baseline": 90703,
+      "baseline_std": 331.9657663073107,
+      "baseline_error": null,
+      "transforms": {
+        "cache_cs": {
+          "cycles": 90984,
+          "std": 289.66283153867016,
+          "delta_pct": 0.31
+        },
+        "cache_cg": {
+          "cycles": 96840,
+          "std": 153.33229592946165,
+          "delta_pct": 6.77
+        },
+        "cache_ca": {
+          "cycles": 90723,
+          "std": 357.56033952327545,
+          "delta_pct": 0.02
+        },
+        "cache_cv": {
+          "cycles": 96850,
+          "std": 155.450086442562,
+          "delta_pct": 6.78
+        },
+        "st_cache_cs": {
+          "cycles": 90717,
+          "std": 396.2587030905441,
+          "delta_pct": 0.02
+        },
+        "st_cache_wt": {
+          "cycles": 90690,
+          "std": 369.9314664907542,
+          "delta_pct": -0.01
+        },
+        "st_cache_wb": {
+          "cycles": 90692,
+          "std": 395.4016678960775,
+          "delta_pct": -0.01
+        },
+        "maxnreg_32": {
+          "cycles": 274793,
+          "std": 5722.283842660638,
+          "delta_pct": 202.96
+        },
+        "maxnreg_64": {
+          "cycles": 146437,
+          "std": 401.55479713234655,
+          "delta_pct": 61.45
+        },
+        "maxnreg_128": {
+          "cycles": 94028,
+          "std": 815.9024439079711,
+          "delta_pct": 3.67
+        },
+        "maxnreg_255": {
+          "cycles": 90439,
+          "std": 214.42241714662205,
+          "delta_pct": -0.29
+        }
+      }
+    },
+    "relu_1024": {
+      "source": "diverse",
+      "baseline": 853,
+      "baseline_std": 10.05211296195979,
+      "baseline_error": null,
+      "transforms": {
+        "cache_cs": {
+          "cycles": 852,
+          "std": 8.13938572620809,
+          "delta_pct": -0.12
+        },
+        "cache_cg": {
+          "cycles": 859,
+          "std": 37.65465409481276,
+          "delta_pct": 0.7
+        },
+        "cache_ca": {
+          "cycles": 852,
+          "std": 8.371676952678,
+          "delta_pct": -0.12
+        },
+        "cache_cv": {
+          "cycles": 860,
+          "std": 37.299769101162006,
+          "delta_pct": 0.82
+        },
+        "st_cache_cs": {
+          "cycles": 852,
+          "std": 9.529657653871938,
+          "delta_pct": -0.12
+        },
+        "st_cache_wt": {
+          "cycles": 852,
+          "std": 7.353903725233286,
+          "delta_pct": -0.12
+        },
+        "st_cache_wb": {
+          "cycles": 853,
+          "std": 7.97044383958635,
+          "delta_pct": 0.0
+        },
+        "maxnreg_32": {
+          "cycles": 860,
+          "std": 9.282078161704954,
+          "delta_pct": 0.82
+        },
+        "maxnreg_64": {
+          "cycles": 852,
+          "std": 7.753444395879808,
+          "delta_pct": -0.12
+        },
+        "maxnreg_128": {
+          "cycles": 853,
+          "std": 8.314323484204833,
+          "delta_pct": 0.0
+        },
+        "maxnreg_255": {
+          "cycles": 853,
+          "std": 8.720344029910747,
+          "delta_pct": 0.0
+        },
+        "reorder_cp": {
+          "cycles": 852,
+          "std": 8.963810573634406,
+          "delta_pct": -0.12
+        },
+        "reorder_il": {
+          "cycles": 853,
+          "std": 8.406293773120233,
+          "delta_pct": 0.0
+        },
+        "reorder_lf": {
+          "cycles": 852,
+          "std": 9.209710907514959,
+          "delta_pct": -0.12
+        },
+        "reorder_sl": {
+          "cycles": 853,
+          "std": 6.763798858629668,
+          "delta_pct": 0.0
+        }
+      }
+    },
+    "relu_512": {
+      "source": "diverse",
+      "baseline": 764,
+      "baseline_std": 16.43397699888861,
+      "baseline_error": null,
+      "transforms": {
+        "cache_cs": {
+          "cycles": 764,
+          "std": 17.24999710144903,
+          "delta_pct": 0.0
+        },
+        "cache_cg": {
+          "cycles": 766,
+          "std": 16.810089232362806,
+          "delta_pct": 0.26
+        },
+        "cache_ca": {
+          "cycles": 764,
+          "std": 16.147148355050188,
+          "delta_pct": 0.0
+        },
+        "cache_cv": {
+          "cycles": 766,
+          "std": 16.470385544971315,
+          "delta_pct": 0.26
+        },
+        "st_cache_cs": {
+          "cycles": 764,
+          "std": 17.5565877948991,
+          "delta_pct": 0.0
+        },
+        "st_cache_wt": {
+          "cycles": 764,
+          "std": 17.564737401965335,
+          "delta_pct": 0.0
+        },
+        "st_cache_wb": {
+          "cycles": 766,
+          "std": 15.71966841253339,
+          "delta_pct": 0.26
+        },
+        "maxnreg_32": {
+          "cycles": 764,
+          "std": 17.01982300142983,
+          "delta_pct": 0.0
+        },
+        "maxnreg_64": {
+          "cycles": 766,
+          "std": 15.90412210717712,
+          "delta_pct": 0.26
+        },
+        "maxnreg_128": {
+          "cycles": 766,
+          "std": 15.540193048993954,
+          "delta_pct": 0.26
+        },
+        "maxnreg_255": {
+          "cycles": 765,
+          "std": 15.783535725559087,
+          "delta_pct": 0.13
+        },
+        "reorder_cp": {
+          "cycles": 764,
+          "std": 16.173768113831727,
+          "delta_pct": 0.0
+        },
+        "reorder_il": {
+          "cycles": 767,
+          "std": 16.0684909061181,
+          "delta_pct": 0.39
+        },
+        "reorder_lf": {
+          "cycles": 764,
+          "std": 16.771750057760816,
+          "delta_pct": 0.0
+        },
+        "reorder_sl": {
+          "cycles": 764,
+          "std": 17.16017482428428,
+          "delta_pct": 0.0
+        }
+      }
+    },
+    "gelu_1024": {
+      "source": "diverse",
+      "baseline": 1371,
+      "baseline_std": 11.377819430804832,
+      "baseline_error": null,
+      "transforms": {
+        "cache_cs": {
+          "cycles": 1372,
+          "std": 12.677648638450272,
+          "delta_pct": 0.07
+        },
+        "cache_cg": {
+          "cycles": 1363,
+          "std": 7.351979325324575,
+          "delta_pct": -0.58
+        },
+        "cache_ca": {
+          "cycles": 1371,
+          "std": 11.97738598359425,
+          "delta_pct": 0.0
+        },
+        "cache_cv": {
+          "cycles": 1364,
+          "std": 11.346672419700852,
+          "delta_pct": -0.51
+        },
+        "st_cache_cs": {
+          "cycles": 1371,
+          "std": 11.756273856966756,
+          "delta_pct": 0.0
+        },
+        "st_cache_wt": {
+          "cycles": 1371,
+          "std": 11.591272578970782,
+          "delta_pct": 0.0
+        },
+        "st_cache_wb": {
+          "cycles": 1370,
+          "std": 8.280940465913277,
+          "delta_pct": -0.07
+        },
+        "maxnreg_32": {
+          "cycles": 1326,
+          "std": 9.086781608468424,
+          "delta_pct": -3.28
+        },
+        "maxnreg_64": {
+          "cycles": 1371,
+          "std": 8.118768379501907,
+          "delta_pct": 0.0
+        },
+        "maxnreg_128": {
+          "cycles": 1371,
+          "std": 12.319496742968035,
+          "delta_pct": 0.0
+        },
+        "maxnreg_255": {
+          "cycles": 1370,
+          "std": 11.341242215912681,
+          "delta_pct": -0.07
+        },
+        "reorder_cp": {
+          "cycles": 1370,
+          "std": 9.817228478547294,
+          "delta_pct": -0.07
+        },
+        "reorder_il": {
+          "cycles": 1370,
+          "std": 12.119792696246913,
+          "delta_pct": -0.07
+        },
+        "reorder_lf": {
+          "cycles": 1371,
+          "std": 12.60591527815414,
+          "delta_pct": 0.0
+        },
+        "reorder_sl": {
+          "cycles": 1371,
+          "std": 11.214205277236545,
+          "delta_pct": 0.0
+        }
+      }
+    },
+    "gelu_512": {
+      "source": "diverse",
+      "baseline": 958,
+      "baseline_std": 17.80862431520189,
+      "baseline_error": null,
+      "transforms": {
+        "cache_cs": {
+          "cycles": 955,
+          "std": 20.503441174593107,
+          "delta_pct": -0.31
+        },
+        "cache_cg": {
+          "cycles": 962,
+          "std": 26.21575432826605,
+          "delta_pct": 0.42
+        },
+        "cache_ca": {
+          "cycles": 957,
+          "std": 19.55654366190508,
+          "delta_pct": -0.1
+        },
+        "cache_cv": {
+          "cycles": 963,
+          "std": 26.990672462908364,
+          "delta_pct": 0.52
+        },
+        "st_cache_cs": {
+          "cycles": 955,
+          "std": 23.808225469362473,
+          "delta_pct": -0.31
+        },
+        "st_cache_wt": {
+          "cycles": 957,
+          "std": 20.593517790800096,
+          "delta_pct": -0.1
+        },
+        "st_cache_wb": {
+          "cycles": 957,
+          "std": 17.687107592820258,
+          "delta_pct": -0.1
+        },
+        "maxnreg_32": {
+          "cycles": 956,
+          "std": 18.6074118296984,
+          "delta_pct": -0.21
+        },
+        "maxnreg_64": {
+          "cycles": 957,
+          "std": 25.47200374921455,
+          "delta_pct": -0.1
+        },
+        "maxnreg_128": {
+          "cycles": 957,
+          "std": 20.030021842224738,
+          "delta_pct": -0.1
+        },
+        "maxnreg_255": {
+          "cycles": 955,
+          "std": 21.280082236683203,
+          "delta_pct": -0.31
+        },
+        "reorder_cp": {
+          "cycles": 959,
+          "std": 19.8063903829042,
+          "delta_pct": 0.1
+        },
+        "reorder_il": {
+          "cycles": 957,
+          "std": 22.347339886438384,
+          "delta_pct": -0.1
+        },
+        "reorder_lf": {
+          "cycles": 957,
+          "std": 22.097415572867337,
+          "delta_pct": -0.1
+        },
+        "reorder_sl": {
+          "cycles": 955,
+          "std": 21.986995588301735,
+          "delta_pct": -0.31
+        }
+      }
+    },
+    "dropout_1024": {
+      "source": "diverse",
+      "baseline": 2137,
+      "baseline_std": 0.0,
+      "baseline_error": null,
+      "transforms": {
+        "cache_cs": {
+          "cycles": 2137,
+          "std": 0.09949874371066199,
+          "delta_pct": 0.0
+        },
+        "cache_cg": {
+          "cycles": 2137,
+          "std": 0.0,
+          "delta_pct": 0.0
+        },
+        "cache_ca": {
+          "cycles": 2137,
+          "std": 0.0,
+          "delta_pct": 0.0
+        },
+        "cache_cv": {
+          "cycles": 2137,
+          "std": 0.0,
+          "delta_pct": 0.0
+        },
+        "st_cache_cs": {
+          "cycles": 2137,
+          "std": 0.07053367989832941,
+          "delta_pct": 0.0
+        },
+        "st_cache_wt": {
+          "cycles": 2137,
+          "std": 0.07053367989832941,
+          "delta_pct": 0.0
+        },
+        "st_cache_wb": {
+          "cycles": 2137,
+          "std": 0.0,
+          "delta_pct": 0.0
+        },
+        "maxnreg_32": {
+          "cycles": 2046,
+          "std": 0.0,
+          "delta_pct": -4.26
+        },
+        "maxnreg_64": {
+          "cycles": 2120,
+          "std": 0.0,
+          "delta_pct": -0.8
+        },
+        "maxnreg_128": {
+          "cycles": 2149,
+          "std": 0.0,
+          "delta_pct": 0.56
+        },
+        "maxnreg_255": {
+          "cycles": 2149,
+          "std": 0.21160103969498834,
+          "delta_pct": 0.56
+        },
+        "reorder_cp": {
+          "cycles": 2210,
+          "std": 5.169020700287435,
+          "delta_pct": 3.42
+        },
+        "reorder_il": {
+          "cycles": 2014,
+          "std": 1.7173744495595593,
+          "delta_pct": -5.76
+        },
+        "reorder_lf": {
+          "cycles": 2210,
+          "std": 5.540794166904236,
+          "delta_pct": 3.42
+        },
+        "reorder_sl": {
+          "cycles": 2210,
+          "std": 8.832144416844644,
+          "delta_pct": 3.42
+        }
+      }
+    },
+    "dropout_512": {
+      "source": "diverse",
+      "baseline": 1625,
+      "baseline_std": 3.360621222333752,
+      "baseline_error": null,
+      "transforms": {
+        "cache_cs": {
+          "cycles": 1625,
+          "std": 3.2819354046050324,
+          "delta_pct": 0.0
+        },
+        "cache_cg": {
+          "cycles": 1623,
+          "std": 3.996845631244719,
+          "delta_pct": -0.12
+        },
+        "cache_ca": {
+          "cycles": 1625,
+          "std": 3.205854020382089,
+          "delta_pct": 0.0
+        },
+        "cache_cv": {
+          "cycles": 1623,
+          "std": 3.873744958047703,
+          "delta_pct": -0.12
+        },
+        "st_cache_cs": {
+          "cycles": 1625,
+          "std": 3.0862234202986665,
+          "delta_pct": 0.0
+        },
+        "st_cache_wt": {
+          "cycles": 1625,
+          "std": 3.241774205585578,
+          "delta_pct": 0.0
+        },
+        "st_cache_wb": {
+          "cycles": 1625,
+          "std": 3.2674760901956117,
+          "delta_pct": 0.0
+        },
+        "maxnreg_32": {
+          "cycles": 1601,
+          "std": 3.7460779490021294,
+          "delta_pct": -1.48
+        },
+        "maxnreg_64": {
+          "cycles": 1562,
+          "std": 4.086682639990534,
+          "delta_pct": -3.88
+        },
+        "maxnreg_128": {
+          "cycles": 1562,
+          "std": 3.956008088970497,
+          "delta_pct": -3.88
+        },
+        "maxnreg_255": {
+          "cycles": 1561,
+          "std": 4.047144054762568,
+          "delta_pct": -3.94
+        },
+        "reorder_cp": {
+          "cycles": 1628,
+          "std": 2.705087798944796,
+          "delta_pct": 0.18
+        },
+        "reorder_il": {
+          "cycles": 1627,
+          "std": 2.4992798962901293,
+          "delta_pct": 0.12
+        },
+        "reorder_lf": {
+          "cycles": 1627,
+          "std": 2.6220173531080984,
+          "delta_pct": 0.12
+        },
+        "reorder_sl": {
+          "cycles": 1627,
+          "std": 2.5333574560255014,
+          "delta_pct": 0.12
+        }
+      }
+    },
+    "cross_entropy_1024": {
+      "source": "diverse",
+      "baseline": 4663,
+      "baseline_std": 74.4845136588808,
+      "baseline_error": null,
+      "transforms": {
+        "cache_cs": {
+          "cycles": 4683,
+          "std": 71.8598217016992,
+          "delta_pct": 0.43
+        },
+        "cache_cg": {
+          "cycles": 4620,
+          "std": 68.62513952626982,
+          "delta_pct": -0.92
+        },
+        "cache_ca": {
+          "cycles": 4683,
+          "std": 71.5853749029786,
+          "delta_pct": 0.43
+        },
+        "cache_cv": {
+          "cycles": 4630,
+          "std": 69.17326777737192,
+          "delta_pct": -0.71
+        },
+        "st_cache_cs": {
+          "cycles": 4672,
+          "std": 64.96739162841618,
+          "delta_pct": 0.19
+        },
+        "st_cache_wt": {
+          "cycles": 4656,
+          "std": 67.25985411075466,
+          "delta_pct": -0.15
+        },
+        "st_cache_wb": {
+          "cycles": 4675,
+          "std": 77.72003843925967,
+          "delta_pct": 0.26
+        },
+        "maxnreg_32": {
+          "cycles": 4686,
+          "std": 53.9437818474011,
+          "delta_pct": 0.49
+        },
+        "maxnreg_64": {
+          "cycles": 4666,
+          "std": 63.665802241077586,
+          "delta_pct": 0.06
+        },
+        "maxnreg_128": {
+          "cycles": 4671,
+          "std": 70.02021136786149,
+          "delta_pct": 0.17
+        },
+        "maxnreg_255": {
+          "cycles": 4679,
+          "std": 68.15793405759891,
+          "delta_pct": 0.34
+        },
+        "reorder_cp": {
+          "error": "skipped_barrier"
+        },
+        "reorder_il": {
+          "error": "skipped_barrier"
+        },
+        "reorder_lf": {
+          "error": "skipped_barrier"
+        },
+        "reorder_sl": {
+          "error": "skipped_barrier"
+        }
+      }
+    },
+    "cross_entropy_512": {
+      "source": "diverse",
+      "baseline": 2768,
+      "baseline_std": 27.59124815951609,
+      "baseline_error": null,
+      "transforms": {
+        "cache_cs": {
+          "cycles": 2767,
+          "std": 27.64353405409663,
+          "delta_pct": -0.04
+        },
+        "cache_cg": {
+          "cycles": 2850,
+          "std": 24.742711553101852,
+          "delta_pct": 2.96
+        },
+        "cache_ca": {
+          "cycles": 2765,
+          "std": 24.116241415278623,
+          "delta_pct": -0.11
+        },
+        "cache_cv": {
+          "cycles": 2849,
+          "std": 24.128014837528596,
+          "delta_pct": 2.93
+        },
+        "st_cache_cs": {
+          "cycles": 2771,
+          "std": 28.70365133567505,
+          "delta_pct": 0.11
+        },
+        "st_cache_wt": {
+          "cycles": 2762,
+          "std": 27.35715217269517,
+          "delta_pct": -0.22
+        },
+        "st_cache_wb": {
+          "cycles": 2764,
+          "std": 31.47713416116531,
+          "delta_pct": -0.14
+        },
+        "maxnreg_32": {
+          "cycles": 2694,
+          "std": 29.4571616928719,
+          "delta_pct": -2.67
+        },
+        "maxnreg_64": {
+          "cycles": 2764,
+          "std": 28.905667264396442,
+          "delta_pct": -0.14
+        },
+        "maxnreg_128": {
+          "cycles": 2763,
+          "std": 28.315727078780796,
+          "delta_pct": -0.18
+        },
+        "maxnreg_255": {
+          "cycles": 2762,
+          "std": 28.16327218204589,
+          "delta_pct": -0.22
+        },
+        "reorder_cp": {
+          "error": "skipped_barrier"
+        },
+        "reorder_il": {
+          "error": "skipped_barrier"
+        },
+        "reorder_lf": {
+          "error": "skipped_barrier"
+        },
+        "reorder_sl": {
+          "error": "skipped_barrier"
+        }
+      }
+    },
+    "batch_norm_1024": {
+      "source": "diverse",
+      "baseline": 3818,
+      "baseline_std": 305.4150782132408,
+      "baseline_error": null,
+      "transforms": {
+        "cache_cs": {
+          "cycles": 3493,
+          "std": 265.58175346774107,
+          "delta_pct": -8.51
+        },
+        "cache_cg": {
+          "cycles": 6958,
+          "std": 271.8136611725025,
+          "delta_pct": 82.24
+        },
+        "cache_ca": {
+          "cycles": 3567,
+          "std": 268.9106727149371,
+          "delta_pct": -6.57
+        },
+        "cache_cv": {
+          "cycles": 6788,
+          "std": 266.913621982843,
+          "delta_pct": 77.79
+        },
+        "st_cache_cs": {
+          "cycles": 3532,
+          "std": 293.2227704323114,
+          "delta_pct": -7.49
+        },
+        "st_cache_wt": {
+          "cycles": 3843,
+          "std": 244.83664023180845,
+          "delta_pct": 0.65
+        },
+        "st_cache_wb": {
+          "cycles": 3549,
+          "std": 286.46786276997983,
+          "delta_pct": -7.05
+        },
+        "maxnreg_32": {
+          "cycles": 3678,
+          "std": 269.71796264802236,
+          "delta_pct": -3.67
+        },
+        "maxnreg_64": {
+          "cycles": 3540,
+          "std": 278.07100527563097,
+          "delta_pct": -7.28
+        },
+        "maxnreg_128": {
+          "cycles": 3818,
+          "std": 268.53938775345415,
+          "delta_pct": 0.0
+        },
+        "maxnreg_255": {
+          "cycles": 3557,
+          "std": 297.45760097028955,
+          "delta_pct": -6.84
+        },
+        "reorder_cp": {
+          "cycles": 3526,
+          "std": 234.53006012023278,
+          "delta_pct": -7.65
+        },
+        "reorder_il": {
+          "cycles": 3092,
+          "std": 262.68993314362086,
+          "delta_pct": -19.02
+        },
+        "reorder_lf": {
+          "cycles": 3437,
+          "std": 303.6824212479214,
+          "delta_pct": -9.98
+        },
+        "reorder_sl": {
+          "cycles": 3447,
+          "std": 274.7649700653269,
+          "delta_pct": -9.72
+        }
+      }
+    },
+    "batch_norm_512": {
+      "source": "diverse",
+      "baseline": 1960,
+      "baseline_std": 24.425166836687115,
+      "baseline_error": null,
+      "transforms": {
+        "cache_cs": {
+          "cycles": 1912,
+          "std": 47.514229184529555,
+          "delta_pct": -2.45
+        },
+        "cache_cg": {
+          "cycles": 3850,
+          "std": 65.87921599412064,
+          "delta_pct": 96.43
+        },
+        "cache_ca": {
+          "cycles": 1914,
+          "std": 51.70893902411845,
+          "delta_pct": -2.35
+        },
+        "cache_cv": {
+          "cycles": 3845,
+          "std": 78.12604799297095,
+          "delta_pct": 96.17
+        },
+        "st_cache_cs": {
+          "cycles": 1917,
+          "std": 48.36344073781351,
+          "delta_pct": -2.19
+        },
+        "st_cache_wt": {
+          "cycles": 1959,
+          "std": 27.981314747523925,
+          "delta_pct": -0.05
+        },
+        "st_cache_wb": {
+          "cycles": 1916,
+          "std": 45.44963806236525,
+          "delta_pct": -2.24
+        },
+        "maxnreg_32": {
+          "cycles": 2007,
+          "std": 37.455332063672856,
+          "delta_pct": 2.4
+        },
+        "maxnreg_64": {
+          "cycles": 1928,
+          "std": 44.08847808668383,
+          "delta_pct": -1.63
+        },
+        "maxnreg_128": {
+          "cycles": 1963,
+          "std": 24.53813155071103,
+          "delta_pct": 0.15
+        },
+        "maxnreg_255": {
+          "cycles": 1920,
+          "std": 48.55923805003534,
+          "delta_pct": -2.04
+        },
+        "reorder_cp": {
+          "cycles": 1851,
+          "std": 35.36270775831511,
+          "delta_pct": -5.56
+        },
+        "reorder_il": {
+          "cycles": 1892,
+          "std": 55.189219961872986,
+          "delta_pct": -3.47
+        },
+        "reorder_lf": {
+          "cycles": 1892,
+          "std": 45.385106587954596,
+          "delta_pct": -3.47
+        },
+        "reorder_sl": {
+          "cycles": 1847,
+          "std": 44.671076772336704,
+          "delta_pct": -5.77
+        }
+      }
+    },
+    "embedding_lookup_256": {
+      "source": "diverse",
+      "baseline": 2583,
+      "baseline_std": 431.5732507639926,
+      "baseline_error": null,
+      "transforms": {
+        "cache_cs": {
+          "cycles": 2415,
+          "std": 374.6212551030707,
+          "delta_pct": -6.5
+        },
+        "cache_cg": {
+          "cycles": 2663,
+          "std": 439.39535952829544,
+          "delta_pct": 3.1
+        },
+        "cache_ca": {
+          "cycles": 2498,
+          "std": 385.7563738941976,
+          "delta_pct": -3.29
+        },
+        "cache_cv": {
+          "cycles": 3023,
+          "std": 315.9150834876359,
+          "delta_pct": 17.03
+        },
+        "st_cache_cs": {
+          "cycles": 2299,
+          "std": 398.473873045649,
+          "delta_pct": -10.99
+        },
+        "st_cache_wt": {
+          "cycles": 2506,
+          "std": 330.7936619329337,
+          "delta_pct": -2.98
+        },
+        "st_cache_wb": {
+          "cycles": 2286,
+          "std": 390.9568873609979,
+          "delta_pct": -11.5
+        },
+        "maxnreg_32": {
+          "cycles": 2602,
+          "std": 392.489609543998,
+          "delta_pct": 0.74
+        },
+        "maxnreg_64": {
+          "cycles": 2592,
+          "std": 382.27569619320553,
+          "delta_pct": 0.35
+        },
+        "maxnreg_128": {
+          "cycles": 2556,
+          "std": 583.0537709122547,
+          "delta_pct": -1.05
+        },
+        "maxnreg_255": {
+          "cycles": 2401,
+          "std": 406.8704358883796,
+          "delta_pct": -7.05
+        },
+        "reorder_cp": {
+          "cycles": 2688,
+          "std": 396.48086851700674,
+          "delta_pct": 4.07
+        },
+        "reorder_il": {
+          "cycles": 2372,
+          "std": 383.4321869314051,
+          "delta_pct": -8.17
+        },
+        "reorder_lf": {
+          "cycles": 2630,
+          "std": 488.5738096746488,
+          "delta_pct": 1.82
+        },
+        "reorder_sl": {
+          "cycles": 2368,
+          "std": 343.44249791631785,
+          "delta_pct": -8.32
+        }
+      }
+    },
+    "embedding_lookup_512": {
+      "source": "diverse",
+      "baseline": 3130,
+      "baseline_std": 522.3538258450875,
+      "baseline_error": null,
+      "transforms": {
+        "cache_cs": {
+          "cycles": 2719,
+          "std": 470.5250943095384,
+          "delta_pct": -13.13
+        },
+        "cache_cg": {
+          "cycles": 4012,
+          "std": 490.27805896246264,
+          "delta_pct": 28.18
+        },
+        "cache_ca": {
+          "cycles": 2799,
+          "std": 481.8130207611662,
+          "delta_pct": -10.58
+        },
+        "cache_cv": {
+          "cycles": 3815,
+          "std": 484.65131342027746,
+          "delta_pct": 21.88
+        },
+        "st_cache_cs": {
+          "cycles": 3121,
+          "std": 460.9627250657042,
+          "delta_pct": -0.29
+        },
+        "st_cache_wt": {
+          "cycles": 3309,
+          "std": 564.0860440571101,
+          "delta_pct": 5.72
+        },
+        "st_cache_wb": {
+          "cycles": 2993,
+          "std": 544.558885130892,
+          "delta_pct": -4.38
+        },
+        "maxnreg_32": {
+          "cycles": 3118,
+          "std": 572.5250244312471,
+          "delta_pct": -0.38
+        },
+        "maxnreg_64": {
+          "cycles": 3586,
+          "std": 367.3977695903991,
+          "delta_pct": 14.57
+        },
+        "maxnreg_128": {
+          "cycles": 3302,
+          "std": 497.0922078447821,
+          "delta_pct": 5.5
+        },
+        "maxnreg_255": {
+          "cycles": 3032,
+          "std": 532.9646478895199,
+          "delta_pct": -3.13
+        },
+        "reorder_cp": {
+          "cycles": 3296,
+          "std": 559.147051297778,
+          "delta_pct": 5.3
+        },
+        "reorder_il": {
+          "cycles": 3201,
+          "std": 469.6084406183517,
+          "delta_pct": 2.27
+        },
+        "reorder_lf": {
+          "cycles": 3234,
+          "std": 527.5769300301142,
+          "delta_pct": 3.32
+        },
+        "reorder_sl": {
+          "cycles": 3441,
+          "std": 550.5740747392671,
+          "delta_pct": 9.94
+        }
+      }
+    }
+  },
+  "best_improvements": [
+    {
+      "kernel": "attention_d64_kv32",
+      "transform": "maxnreg_64",
+      "delta_pct": -1.78
+    },
+    {
+      "kernel": "attention_d64_kv64",
+      "transform": "maxnreg_255",
+      "delta_pct": -46.95
+    },
+    {
+      "kernel": "batch_norm_1024",
+      "transform": "reorder_il",
+      "delta_pct": -19.02
+    },
+    {
+      "kernel": "batch_norm_512",
+      "transform": "reorder_sl",
+      "delta_pct": -5.77
+    },
+    {
+      "kernel": "cross_entropy_512",
+      "transform": "maxnreg_32",
+      "delta_pct": -2.67
+    },
+    {
+      "kernel": "dropout_1024",
+      "transform": "reorder_il",
+      "delta_pct": -5.76
+    },
+    {
+      "kernel": "dropout_512",
+      "transform": "maxnreg_255",
+      "delta_pct": -3.94
+    },
+    {
+      "kernel": "embedding_lookup_256",
+      "transform": "st_cache_wb",
+      "delta_pct": -11.5
+    },
+    {
+      "kernel": "embedding_lookup_512",
+      "transform": "cache_cs",
+      "delta_pct": -13.13
+    },
+    {
+      "kernel": "gelu_1024",
+      "transform": "maxnreg_32",
+      "delta_pct": -3.28
+    },
+    {
+      "kernel": "reduction_max_1024",
+      "transform": "cache_cg",
+      "delta_pct": -1.14
+    },
+    {
+      "kernel": "reduction_sum_1024",
+      "transform": "cache_cv",
+      "delta_pct": -1.14
+    },
+    {
+      "kernel": "triton_fused_add_mul_1024",
+      "transform": "cache_cs",
+      "delta_pct": -1.64
+    },
+    {
+      "kernel": "triton_fused_add_mul_512",
+      "transform": "st_cache_cs",
+      "delta_pct": -1.81
+    },
+    {
+      "kernel": "triton_layernorm_1024",
+      "transform": "st_cache_cs",
+      "delta_pct": -1.08
+    },
+    {
+      "kernel": "triton_vector_add_256",
+      "transform": "cache_cv",
+      "delta_pct": -1.17
+    },
+    {
+      "kernel": "triton_vector_add_512",
+      "transform": "cache_cv",
+      "delta_pct": -3.37
+    }
+  ]
+}

training_result.json

@@ -0,0 +1,725 @@
+{
+  "mean_improvement": -0.29151748216617807,
+  "n_kernels": 64,
+  "per_kernel": {
+    "gemm_tile(2,2,2)": {
+      "improvement": -0.0191,
+      "actions": [
+        "vec_st",
+        "vec_ld",
+        "stop"
+      ],
+      "baseline_cycles": 419,
+      "final_cycles": 411
+    },
+    "gemm_tile(2,2,4)": {
+      "improvement": -0.0363,
+      "actions": [
+        "vec_ld",
+        "vec_st",
+        "stop"
+      ],
+      "baseline_cycles": 441,
+      "final_cycles": 425
+    },
+    "gemm_tile(2,2,6)": {
+      "improvement": -0.1475,
+      "actions": [
+        "vec_ld",
+        "vec_st",
+        "maxnreg_128",
+        "stop"
+      ],
+      "baseline_cycles": 522,
+      "final_cycles": 445
+    },
+    "gemm_tile(2,2,8)": {
+      "improvement": -0.1373,
+      "actions": [
+        "cache_cs",
+        "vec_ld",
+        "vec_st",
+        "maxnreg_128",
+        "stop"
+      ],
+      "baseline_cycles": 539,
+      "final_cycles": 465
+    },
+    "gemm_tile(2,4,2)": {
+      "improvement": -0.1237,
+      "actions": [
+        "vec_st",
+        "vec_ld",
+        "stop"
+      ],
+      "baseline_cycles": 477,
+      "final_cycles": 418
+    },
+    "gemm_tile(2,4,4)": {
+      "improvement": -0.1654,
+      "actions": [
+        "vec_ld",
+        "vec_st",
+        "maxnreg_128",
+        "stop"
+      ],
+      "baseline_cycles": 544,
+      "final_cycles": 454
+    },
+    "gemm_tile(2,4,6)": {
+      "improvement": -0.1932,
+      "actions": [
+        "cache_cs",
+        "vec_ld",
+        "maxnreg_128",
+        "vec_st",
+        "stop"
+      ],
+      "baseline_cycles": 647,
+      "final_cycles": 522
+    },
+    "gemm_tile(2,4,8)": {
+      "improvement": -0.361,
+      "actions": [
+        "maxnreg_128",
+        "vec_ld",
+        "vec_st",
+        "stop"
+      ],
+      "baseline_cycles": 856,
+      "final_cycles": 547
+    },
+    "gemm_tile(2,6,2)": {
+      "improvement": -0.1899,
+      "actions": [
+        "vec_st",
+        "vec_ld",
+        "stop"
+      ],
+      "baseline_cycles": 537,
+      "final_cycles": 435
+    },
+    "gemm_tile(2,6,4)": {
+      "improvement": -0.1651,
+      "actions": [
+        "cache_cs",
+        "vec_ld",
+        "maxnreg_128",
+        "vec_st",
+        "stop"
+      ],
+      "baseline_cycles": 630,
+      "final_cycles": 526
+    },
+    "gemm_tile(2,6,6)": {
+      "improvement": -0.3322,
+      "actions": [
+        "maxnreg_128",
+        "vec_ld",
+        "vec_st",
+        "stop"
+      ],
+      "baseline_cycles": 891,
+      "final_cycles": 595
+    },
+    "gemm_tile(2,6,8)": {
+      "improvement": -0.493,
+      "actions": [
+        "maxnreg_128",
+        "vec_ld",
+        "vec_st",
+        "stop"
+      ],
+      "baseline_cycles": 1278,
+      "final_cycles": 648
+    },
+    "gemm_tile(2,8,2)": {
+      "improvement": -0.1929,
+      "actions": [
+        "vec_st",
+        "vec_ld",
+        "maxnreg_128",
+        "stop"
+      ],
+      "baseline_cycles": 565,
+      "final_cycles": 456
+    },
+    "gemm_tile(2,8,4)": {
+      "improvement": -0.4123,
+      "actions": [
+        "prefetch_L1",
+        "st_cache_cs",
+        "maxnreg_128",
+        "vec_ld",
+        "vec_st",
+        "stop"
+      ],
+      "baseline_cycles": 912,
+      "final_cycles": 536
+    },
+    "gemm_tile(2,8,6)": {
+      "improvement": -0.2716,
+      "actions": [
+        "maxnreg_128",
+        "vec_ld",
+        "vec_st",
+        "stop"
+      ],
+      "baseline_cycles": 902,
+      "final_cycles": 657
+    },
+    "gemm_tile(2,8,8)": {
+      "improvement": -0.2637,
+      "actions": [
+        "maxnreg_128",
+        "vec_ld",
+        "vec_st",
+        "stop"
+      ],
+      "baseline_cycles": 1005,
+      "final_cycles": 740
+    },
+    "gemm_tile(4,2,2)": {
+      "improvement": -0.0695,
+      "actions": [
+        "vec_st",
+        "vec_ld",
+        "stop"
+      ],
+      "baseline_cycles": 475,
+      "final_cycles": 442
+    },
+    "gemm_tile(4,2,4)": {
+      "improvement": -0.1343,
+      "actions": [
+        "vec_ld",
+        "vec_st",
+        "maxnreg_128",
+        "stop"
+      ],
+      "baseline_cycles": 536,
+      "final_cycles": 464
+    },
+    "gemm_tile(4,2,6)": {
+      "improvement": -0.1936,
+      "actions": [
+        "cache_cs",
+        "vec_ld",
+        "maxnreg_128",
+        "vec_st",
+        "stop"
+      ],
+      "baseline_cycles": 625,
+      "final_cycles": 504
+    },
+    "gemm_tile(4,2,8)": {
+      "improvement": -0.2722,
+      "actions": [
+        "maxnreg_128",
+        "vec_ld",
+        "vec_st",
+        "stop"
+      ],
+      "baseline_cycles": 753,
+      "final_cycles": 548
+    },
+    "gemm_tile(4,4,2)": {
+      "improvement": -0.1865,
+      "actions": [
+        "vec_st",
+        "vec_ld",
+        "maxnreg_128",
+        "stop"
+      ],
+      "baseline_cycles": 547,
+      "final_cycles": 445
+    },
+    "gemm_tile(4,4,4)": {
+      "improvement": -0.1693,
+      "actions": [
+        "prefetch_L1",
+        "vec_st",
+        "maxnreg_128",
+        "vec_ld",
+        "stop"
+      ],
+      "baseline_cycles": 691,
+      "final_cycles": 574
+    },
+    "gemm_tile(4,4,6)": {
+      "improvement": -0.3213,
+      "actions": [
+        "maxnreg_128",
+        "vec_ld",
+        "vec_st",
+        "stop"
+      ],
+      "baseline_cycles": 940,
+      "final_cycles": 638
+    },
+    "gemm_tile(4,4,8)": {
+      "improvement": -0.4308,
+      "actions": [
+        "maxnreg_128",
+        "vec_ld",
+        "vec_st",
+        "stop"
+      ],
+      "baseline_cycles": 1258,
+      "final_cycles": 716
+    },
+    "gemm_tile(4,6,2)": {
+      "improvement": -0.2639,
+      "actions": [
+        "vec_st",
+        "vec_ld",
+        "maxnreg_128",
+        "stop"
+      ],
+      "baseline_cycles": 648,
+      "final_cycles": 477
+    },
+    "gemm_tile(4,6,4)": {
+      "improvement": -0.238,
+      "actions": [
+        "maxnreg_128",
+        "vec_ld",
+        "vec_st",
+        "stop"
+      ],
+      "baseline_cycles": 832,
+      "final_cycles": 634
+    },
+    "gemm_tile(4,6,6)": {
+      "improvement": -0.502,
+      "actions": [
+        "maxnreg_128",
+        "vec_ld",
+        "vec_st",
+        "stop"
+      ],
+      "baseline_cycles": 1502,
+      "final_cycles": 748
+    },
+    "gemm_tile(4,6,8)": {
+      "improvement": -0.5383,
+      "actions": [
+        "maxnreg_128",
+        "vec_ld",
+        "vec_st",
+        "stop"
+      ],
+      "baseline_cycles": 1839,
+      "final_cycles": 849
+    },
+    "gemm_tile(4,8,2)": {
+      "improvement": -0.2437,
+      "actions": [
+        "vec_st",
+        "maxnreg_128",
+        "vec_ld",
+        "stop"
+      ],
+      "baseline_cycles": 714,
+      "final_cycles": 540
+    },
+    "gemm_tile(4,8,4)": {
+      "improvement": -0.2336,
+      "actions": [
+        "maxnreg_128",
+        "vec_ld",
+        "vec_st",
+        "stop"
+      ],
+      "baseline_cycles": 959,
+      "final_cycles": 735
+    },
+    "gemm_tile(4,8,6)": {
+      "improvement": -0.2232,
+      "actions": [
+        "maxnreg_128",
+        "vec_ld",
+        "vec_st",
+        "stop"
+      ],
+      "baseline_cycles": 1102,
+      "final_cycles": 856
+    },
+    "gemm_tile(4,8,8)": {
+      "improvement": -0.3196,
+      "actions": [
+        "maxnreg_128",
+        "vec_ld",
+        "vec_st",
+        "stop"
+      ],
+      "baseline_cycles": 1430,
+      "final_cycles": 973
+    },
+    "gemm_tile(6,2,2)": {
+      "improvement": -0.1468,
+      "actions": [
+        "vec_st",
+        "vec_ld",
+        "stop"
+      ],
+      "baseline_cycles": 545,
+      "final_cycles": 465
+    },
+    "gemm_tile(6,2,4)": {
+      "improvement": -0.1749,
+      "actions": [
+        "cache_cs",
+        "vec_ld",
+        "maxnreg_128",
+        "vec_st",
+        "stop"
+      ],
+      "baseline_cycles": 606,
+      "final_cycles": 500
+    },
+    "gemm_tile(6,2,6)": {
+      "improvement": -0.336,
+      "actions": [
+        "maxnreg_128",
+        "vec_ld",
+        "vec_st",
+        "stop"
+      ],
+      "baseline_cycles": 866,
+      "final_cycles": 575
+    },
+    "gemm_tile(6,2,8)": {
+      "improvement": -0.36,
+      "actions": [
+        "maxnreg_128",
+        "vec_ld",
+        "vec_st",
+        "stop"
+      ],
+      "baseline_cycles": 989,
+      "final_cycles": 633
+    },
+    "gemm_tile(6,4,2)": {
+      "improvement": -0.2897,
+      "actions": [
+        "vec_st",
+        "vec_ld",
+        "maxnreg_128",
+        "stop"
+      ],
+      "baseline_cycles": 673,
+      "final_cycles": 478
+    },
+    "gemm_tile(6,4,4)": {
+      "improvement": -0.2296,
+      "actions": [
+        "maxnreg_128",
+        "vec_ld",
+        "vec_st",
+        "stop"
+      ],
+      "baseline_cycles": 832,
+      "final_cycles": 641
+    },
+    "gemm_tile(6,4,6)": {
+      "improvement": -0.4646,
+      "actions": [
+        "maxnreg_128",
+        "vec_ld",
+        "vec_st",
+        "stop"
+      ],
+      "baseline_cycles": 1386,
+      "final_cycles": 742
+    },
+    "gemm_tile(6,4,8)": {
+      "improvement": -0.4832,
+      "actions": [
+        "maxnreg_128",
+        "vec_ld",
+        "vec_st",
+        "stop"
+      ],
+      "baseline_cycles": 1546,
+      "final_cycles": 799
+    },
+    "gemm_tile(6,6,2)": {
+      "improvement": -0.2968,
+      "actions": [
+        "vec_st",
+        "maxnreg_128",
+        "vec_ld",
+        "stop"
+      ],
+      "baseline_cycles": 775,
+      "final_cycles": 545
+    },
+    "gemm_tile(6,6,4)": {
+      "improvement": -0.3052,
+      "actions": [
+        "maxnreg_128",
+        "vec_ld",
+        "vec_st",
+        "stop"
+      ],
+      "baseline_cycles": 983,
+      "final_cycles": 683
+    },
+    "gemm_tile(6,6,6)": {
+      "improvement": -0.4544,
+      "actions": [
+        "maxnreg_128",
+        "vec_ld",
+        "vec_st",
+        "stop"
+      ],
+      "baseline_cycles": 1635,
+      "final_cycles": 892
+    },
+    "gemm_tile(6,6,8)": {
+      "improvement": -0.4978,
+      "actions": [
+        "maxnreg_255",
+        "vec_ld",
+        "vec_st",
+        "stop"
+      ],
+      "baseline_cycles": 2059,
+      "final_cycles": 1034
+    },
+    "gemm_tile(6,8,2)": {
+      "improvement": -0.2758,
+      "actions": [
+        "vec_st",
+        "maxnreg_128",
+        "vec_ld",
+        "stop"
+      ],
+      "baseline_cycles": 881,
+      "final_cycles": 638
+    },
+    "gemm_tile(6,8,4)": {
+      "improvement": -0.264,
+      "actions": [
+        "vec_st",
+        "maxnreg_128",
+        "vec_ld",
+        "stop"
+      ],
+      "baseline_cycles": 1197,
+      "final_cycles": 881
+    },
+    "gemm_tile(6,8,6)": {
+      "improvement": -0.3264,
+      "actions": [
+        "maxnreg_128",
+        "vec_ld",
+        "vec_st",
+        "stop"
+      ],
+      "baseline_cycles": 1578,
+      "final_cycles": 1063
+    },
+    "gemm_tile(6,8,8)": {
+      "improvement": -0.3011,
+      "actions": [
+        "maxnreg_255",
+        "vec_ld",
+        "vec_st",
+        "stop"
+      ],
+      "baseline_cycles": 1737,
+      "final_cycles": 1214
+    },
+    "gemm_tile(8,2,2)": {
+      "improvement": -0.1706,
+      "actions": [
+        "vec_st",
+        "vec_ld",
+        "maxnreg_128",
+        "stop"
+      ],
+      "baseline_cycles": 592,
+      "final_cycles": 491
+    },
+    "gemm_tile(8,2,4)": {
+      "improvement": -0.2521,
+      "actions": [
+        "prefetch_L1",
+        "st_cache_cs",
+        "maxnreg_128",
+        "vec_ld",
+        "vec_st",
+        "stop"
+      ],
+      "baseline_cycles": 722,
+      "final_cycles": 540
+    },
+    "gemm_tile(8,2,6)": {
+      "improvement": -0.4012,
+      "actions": [
+        "maxnreg_128",
+        "vec_ld",
+        "vec_st",
+        "stop"
+      ],
+      "baseline_cycles": 1042,
+      "final_cycles": 624
+    },
+    "gemm_tile(8,2,8)": {
+      "improvement": -0.3817,
+      "actions": [
+        "maxnreg_128",
+        "vec_ld",
+        "vec_st",
+        "stop"
+      ],
+      "baseline_cycles": 1158,
+      "final_cycles": 716
+    },
+    "gemm_tile(8,4,2)": {
+      "improvement": -0.297,
+      "actions": [
+        "vec_st",
+        "maxnreg_128",
+        "vec_ld",
+        "stop"
+      ],
+      "baseline_cycles": 744,
+      "final_cycles": 523
+    },
+    "gemm_tile(8,4,4)": {
+      "improvement": -0.3152,
+      "actions": [
+        "maxnreg_128",
+        "vec_ld",
+        "vec_st",
+        "stop"
+      ],
+      "baseline_cycles": 1050,
+      "final_cycles": 719
+    },
+    "gemm_tile(8,4,6)": {
+      "improvement": -0.4074,
+      "actions": [
+        "maxnreg_128",
+        "vec_ld",
+        "vec_st",
+        "stop"
+      ],
+      "baseline_cycles": 1436,
+      "final_cycles": 851
+    },
+    "gemm_tile(8,4,8)": {
+      "improvement": -0.4238,
+      "actions": [
+        "maxnreg_128",
+        "vec_ld",
+        "vec_st",
+        "stop"
+      ],
+      "baseline_cycles": 1647,
+      "final_cycles": 949
+    },
+    "gemm_tile(8,6,2)": {
+      "improvement": -0.2839,
+      "actions": [
+        "vec_st",
+        "maxnreg_128",
+        "vec_ld",
+        "stop"
+      ],
+      "baseline_cycles": 863,
+      "final_cycles": 618
+    },
+    "gemm_tile(8,6,4)": {
+      "improvement": -0.4151,
+      "actions": [
+        "vec_st",
+        "maxnreg_128",
+        "vec_ld",
+        "stop"
+      ],
+      "baseline_cycles": 1349,
+      "final_cycles": 789
+    },
+    "gemm_tile(8,6,6)": {
+      "improvement": -0.4326,
+      "actions": [
+        "maxnreg_128",
+        "vec_ld",
+        "vec_st",
+        "stop"
+      ],
+      "baseline_cycles": 1789,
+      "final_cycles": 1015
+    },
+    "gemm_tile(8,6,8)": {
+      "improvement": -0.4535,
+      "actions": [
+        "maxnreg_255",
+        "vec_ld",
+        "vec_st",
+        "stop"
+      ],
+      "baseline_cycles": 2192,
+      "final_cycles": 1198
+    },
+    "gemm_tile(8,8,2)": {
+      "improvement": -0.2404,
+      "actions": [
+        "vec_st",
+        "maxnreg_128",
+        "vec_ld",
+        "stop"
+      ],
+      "baseline_cycles": 965,
+      "final_cycles": 733
+    },
+    "gemm_tile(8,8,4)": {
+      "improvement": -0.3331,
+      "actions": [
+        "vec_st",
+        "maxnreg_128",
+        "vec_ld",
+        "stop"
+      ],
+      "baseline_cycles": 1540,
+      "final_cycles": 1027
+    },
+    "gemm_tile(8,8,6)": {
+      "improvement": -0.3909,
+      "actions": [
+        "vec_st",
+        "maxnreg_255",
+        "vec_ld",
+        "stop"
+      ],
+      "baseline_cycles": 2062,
+      "final_cycles": 1256
+    },
+    "gemm_tile(8,8,8)": {
+      "improvement": -0.4082,
+      "actions": [
+        "maxnreg_255",
+        "vec_ld",
+        "vec_st",
+        "stop"
+      ],
+      "baseline_cycles": 2452,
+      "final_cycles": 1451
+    }
+  },
+  "action_distribution": {
+    "vec_st": 0.246,
+    "vec_ld": 0.246,
+    "stop": 0.246,
+    "maxnreg_128": 0.204,
+    "cache_cs": 0.019,
+    "maxnreg_255": 0.019,
+    "prefetch_L1": 0.012,
+    "st_cache_cs": 0.008
+  },
+  "unique_sequences": 12
+}