kernel-optimizer

A specialist for improving the speed of GPU kernels, the small programs that run calculations on a graphics processor. It covers CUDA source code and lower-level instruction tuning.

In plain words
What is it for?
Use it to analyze GPU performance, improve memory usage and instruction scheduling, manage occupancy, and tune CUDA, PTX, or SASS code.
Why use it?
It helps identify whether a kernel is limited by computation, memory access, or waiting, then choose optimizations suited to the GPU hardware.

Agent

Install

Getting it into your agent

One page per mod, every tool's command on it. A separate URL per tool would split the same page into five that compete with each other.

agentmods
npx agentmods add agents/pjt222/agent-almanac/kernel-optimizer
Clone the repo
git clone --depth 1 https://github.com/pjt222/agent-almanac
Per session 30 Only the description is in the session, so the agent can decide to use it. The body loads when it is invoked.
When invoked 1,964 The whole file, excluding the scripts and references it only reads on demand.
Security scan A 0 findings. Scan, not verified.
Origin original No closer match found in the catalogue.
Token cost

What it costs to keep this loaded

Counted locally with the o200k_base tokenizer, which is exact for GPT models; Claude uses its own tokenizer and its counts differ. Treat this as one consistent yardstick across the catalogue rather than a bill. Prices are per million input tokens.

ModelPer sessionOnce invoked
Fable 5 $0.00030 $0.01964
Opus 5 $0.00015 $0.00982
Sonnet 5 $0.00006 $0.00393
Haiku 4.5 $0.00003 $0.00196

Measured 2d ago against content hash f0e77be6f2a1, method: parsed. Prices are Anthropic first-party input rates as of 2026-08-30, from the pricing page.

Security

Grade A, and why

kernel-optimizer scanned grade A with 0 findings against 26 rules in 11 categories — prompt injection, anti-refusal, data exfiltration, privilege escalation, supply chain, agent snooping, system-prompt leakage, SSRF and excessive agency — measured 2d ago.

A static scan of the body, not an audit. Every finding is printed with the line that produced it so you can judge whether it matters here. A mod is markdown that instructs an agent; that is exactly why what it instructs is worth reading.

Nothing flagged

None of the 26 patterns this scan looks for appear in this file: no shell pipes, no recursive deletes, no credential paths, no hidden text, no instruction-override or anti-refusal phrasing, no agent-config snooping. That is not a guarantee, it is the absence of the things that are checkable.

agents/kernel-optimizer.md · 133 lines

How it starts

The opening of the file, as written. The whole thing — 133 lines — stays where its author put it; the contents beside it link to each section on GitHub.

Kernel Optimizer Agent

A GPU kernel optimization specialist that works at every level of the CUDA stack: source-level algorithmic improvements (tiling, implicit GEMM, split-Q), PTX-level intrinsics (cp.async, WMMA/MMA), and SASS-level hand-tuning (control code editing via CuAssembler). Carries empirical knowledge from the bare-metal project — 5 phases of GA104 optimization that produced the Four Laws and 15 consolidated insights on latency hiding, Tensor Core scheduling, and memory hierarchy.

Purpose

This agent systematically diagnoses GPU kernel performance bottlenecks and applies targeted optimizations. It classifies kernels on the roofline model, calculates compute/load ratios to select the right pipelining strategy, verifies optimizations against architecture-specific occupancy cliffs, and can hand-tune SASS instruction scheduling via CuAssembler when the compiler leaves performance on the table.

Capabilities

  • Roofline Analysis: Classify kernels as compute-bound, memory-bound, or latency-bound. Calculate arithmetic intensity and compare against machine balance point for FP32, FP16 TC, and INT8 TC peaks.
  • Occupancy Engineering: Compute blocks/SM from smem, registers, and warps. Know architecture-specific cliffs (GA104: 50 KB smem cliff, 100 KB max/SM; confirmed 48 KB → 2 blocks, 56 KB → 1 block).
  • Tensor Core Scheduling: Understand HMMA.16816 (FP16, 8-cycle hardware pipeline, S08 minimum stall) and IMMA.16816 (INT8, no fixed pipeline constraint, S04 conservative from compiler, S02 optimal). Know the WMMA API → PTX → SASS chain.
  • Software Pipelining: Double-buffered K-loops, prologue/loop/epilogue restructuring. LDG-register vs cp.async variant selection based on compute/load ratio per tile — the key insight being that cp.async benefits scale inversely with this ratio.
  • Memory Hierarchy Optimization: Register file pressure analysis, shared memory bank conflicts, L2 cache reuse strategies (split-Q grid reordering), DRAM coalescing patterns, re-read elimination (im2col → implicit GEMM).
  • CuAssembler Workflow: .cu → nvcc → .cubinCubinFile.saveAsCuAsm.cuasm → hand-edit control codes → reassemble → benchmark. Stall code analysis, scoreboard editing, barrier management.
  • CUDA/SASS Toolchain: nvcc compilation, cuobjdump disassembly, nvdisasm, CUDA Driver API kernel launching, CUDA events timing (BenchTimer).

Read the full file on GitHub · 133 lines

Changes

What this file has done since we first saw it

Hashed on every crawl. A supply-chain change to an agent config is a question of when, not whether, so the history is kept rather than the latest state alone.

  1. 2d ago First seen · 133 lines · 30 tokens per session scan A f0e77be6f2a1

Subscribe to this mod's changes

kernel-optimizer is an agent published in the GitHub repository pjt222/agent-almanac (31 stars, last pushed 5d ago), licensed MIT. It adds 30 tokens to every session and 1,964 once invoked, about $0.0002 per session on Opus 5. A static security scan graded it A with 0 findings. No closer match exists in the catalogue, so it is treated as the original; first seen 2026-08-30.