Omar SobhandClaude Opus 4.7
7b23f61632
Hand-authored skill catalog anchored to real 2026-07 versions:
- Rust 1.97.1 (stable), edition 2024
- React 19.2.7, Server Components + Actions
- TailwindCSS 4.3.3 (CSS-first config, Oxide engine)
- three.js r185 (WebGPURenderer stable, BatchedMesh matured)
- React Native 0.86 / Expo SDK 54+ (New Architecture default)
- cargo-nextest 0.9.140, gitleaks 8.20+, cargo-audit 0.21+
- Postgres 17 (18 in beta, don't rely on)
- CUDA Blackwell, Metal Apple7+, ROCm CDNA3
Ships 15 skills across the categories:
foundation/ workspace-repo-commit-protocol
small-focused-commits
tdd-red-green-refactor
code-review-checklist
int-xx-marker-protocol
decompose-int-items
rust/ write-rust-current-edition
rust-error-handling
cargo-test-driven-development
rust-async-tokio-idioms
backend/ postgres-migrations-forward-only
postgres-index-selection
api-pagination-day-1
frontend/ react-19-server-components
tailwind-v4-idioms
component-4-state-model
mobile/ expo-managed-vs-bare
rn-flashlist-perf
gpu/ gpu-coalescing-and-occupancy
roofline-model
threejs/ threejs-perf-and-teardown
security/ cargo-audit-workflow
secret-scanning-gitleaks
skills_loader.rs walks skills/**/*.md, parses YAML frontmatter
(name, description, when_to_use, tags), upserts via
skills_catalog::upsert_builtin. Idempotent per boot — bumps version
+ appends skill_versions row ONLY when body changes. Deterministic
sha256-derived ids so builtins are stable across boots.
Dockerfile copies skills/ to /etc/clawmates/skills. Server boot
task spawns loader alongside team_template_loader.
Follow-ups (Slice 3.5c continuation, future PRs):
- 20-30 more skills (duckdb, shadcn composition, a11y, WebGPU
migration, metal frame capture, rocprof, deep gitea forge
integration, semgrep rulepacks)
- Bind skills to team template roles (add [role.skills] refs to
templates/teams/*.toml + wire template_role_skills population
in team_template_loader)
Co-Authored-By: Claude Opus 4.7 <[email protected]>
3.0 KiB
3.0 KiB
name, description, when_to_use, tags
| name | description | when_to_use | tags | |||||
|---|---|---|---|---|---|---|---|---|
| gpu-coalescing-and-occupancy | The two knobs that dominate GPU perf — memory coalescing + occupancy targets across CUDA (Blackwell), Metal (Apple 7+), ROCm (RDNA3/CDNA3). | You're a kernel_author or bench_engineer profiling a GPU kernel. GPU team pins. |
|
Coalescing + occupancy
Two dominant knobs on ANY modern GPU. Get them right first; then worry about smaller wins.
Coalescing
Rule: consecutive threads in a warp/wavefront should read consecutive 32/64/128-bit words.
- CUDA (Blackwell, SM_100+) — warp = 32 threads. Threads 0..31 accessing
arr[tid]= 1 memory transaction.arr[tid * stride]for stride > 1 = multiple transactions. Violation cost: 4×–32× slowdown depending on stride. - Metal (Apple 7+, i.e., A17 / M3+) — SIMD-group = 32 threads, same principle.
metal::simd_shufflefor cross-lane comms without shared memory. - ROCm (RDNA3, CDNA3) — wavefront = 32 (RDNA) or 64 (CDNA). Same rule; note the 64-wide wavefronts on CDNA change your indexing math.
Prove coalescing in a comment on the kernel:
// Coalesced: thread tid reads global[tid], stride-1.
// Warp-level transaction: 1× 128-byte load per warp.
float x = global_in[gid];
Occupancy
Rough targets:
- Memory-bound kernel: aim for ≥ 50% occupancy to hide DRAM latency.
- Compute-bound kernel with high ILP: 25%–50% is often fine; more threads compete for registers + shared memory.
Occupancy is bounded by whichever hits first:
- Registers per thread × threads per block ≤ register file per SM.
- Shared / threadgroup memory per block ≤ per-SM budget.
- Threads per block ≤ max (typically 1024).
Tools:
- Nsight Compute (CUDA):
--section=Occupancyprints the bottleneck. - Xcode GPU Frame Capture (Metal): occupancy chart in the performance report.
- rocprof (ROCm):
--statsfor kernel occupancy.
Memory hierarchy strategy
Order-of-magnitude comparison — always prefer the fastest tier your data footprint allows:
| Tier | CUDA name | Latency | Notes |
|---|---|---|---|
| Register | reg | ~1 cycle | Per-thread; too many spills to L1 |
| Shared | shared / threadgroup / LDS | ~30 cycles | Per-block. CUDA: 48–100 KB. Metal: 32 KB (Apple7+). ROCm: 64 KB LDS |
| L1 / texture | L1 tex | ~200 cycles | Cache; some HW has separate paths |
| L2 | L2 | ~500 cycles | Shared across SMs |
| Global | DRAM / HBM | 400–800 cycles | Coalescing matters MOST here |
Divergence
- Warp-uniform control flow keeps all lanes active.
if (tid % 2)splits the warp — half the lanes idle per branch.- On CUDA Volta+ (7.0+), Independent Thread Scheduling makes divergence CORRECT under all conditions but doesn't make it FAST.
The one bottleneck rule
After profiling, report ONE bottleneck to attack next. Reporting five doesn't help — you'll fix one and re-profile anyway.
Format:
KERNEL foo achieves 340 GB/s of 3.35 TB/s HBM = 10% BW-bound. Occupancy 62%. Bottleneck: stride-2 loads from arr — refactor tile shape to coalesce.