Menu
DSL-specific tuning contract for CuTe/CUTLASS C++ template kernels. Loaded by croq-tune when dsl=cute-cpp.
Install with Codex or Claude Copy this prompt, paste it into Codex, Claude, or another assistant, and let it review the skill page and install it for you.
Based on SOC occupation classification
DSL-specific tuning contract for CuTe DSL (Python JIT) kernels. Loaded by croq-tune when dsl=cute-dsl.
Launch an infinite AI-driven kernel optimization loop for GPU kernels. Use when the user asks to "tune", "ai-tune", "optimize", "auto-tune", or "perf-tune" a kernel. Profiles with ncu, iterates optimizations indefinitely until interrupted.
Cursor-only kernel tuning entrypoint for GPU kernels. Use when the user asks to /croq-tune, continue tuning, resume tuning, ai-tune, auto-tune, optimize, or perf-tune a kernel in this repository. Profiles with ncu, iterates optimizations indefinitely until interrupted.
Reference collection of Choreo GPU kernel implementations. Use when studying kernel patterns, tiling strategies, warp specialization, TMA/DMA forms, or pipeline staging before editing .co files.
Launch an infinite AI-driven kernel optimization loop for GPU kernels. Use when the user asks to "tune", "ai-tune", "optimize", or "auto-tune" a kernel. Profiles with ncu, iterates optimizations indefinitely until interrupted.
Durable wrapper for GPU kernel tuning. Use when the user wants to combine durable-request with croq-tune, asks for a checkpointed tuning workflow, or wants human confirmation between autonomous tuning invocation leases while keeping the base /croq-tune protocol unchanged.
| name | croq-dsl-cute-cpp |
| description | DSL-specific tuning contract for CuTe/CUTLASS C++ template kernels. Loaded by croq-tune when dsl=cute-cpp. |
Source extension: .cu or .cpp | Compiler: nvcc (CUTLASS templates) | Group: compiled-binary
Note: This is the C++ CuTe/CUTLASS template interface. For Python JIT via CuTe DSL, use
cute-dslinstead.
python3 -c "import cutlass" # optional, for Python template instantiation
build_iter<NNN>.sh:
#!/usr/bin/env bash
set -e
nvcc -O3 -arch=sm_90 -std=c++17 -I/usr/local/cuda/include \
-o tuning/<gpu>/cute-cpp/bin/<shape_key>/<model>/iter<NNN>_<tag> \
tuning/<gpu>/cute-cpp/srcs/<shape_key>/<model>/iter<NNN>_<tag>.cu \
2>&1 | tee tuning/<gpu>/cute-cpp/perf/<shape_key>/<model>/build_iter<NNN>.txt
run_iter<NNN>.sh:
#!/usr/bin/env bash
tuning/<gpu>/cute-cpp/bin/<shape_key>/<model>/iter<NNN>_<tag> \
2>&1 | tee tuning/<gpu>/cute-cpp/perf/<shape_key>/<model>/timing_iter<NNN>.txt
Binary must print: TFLOPS: <value> time_ms: <value>
ncu --set full \
--export tuning/<gpu>/cute-cpp/perf/<shape_key>/<model>/ncu_iter<NNN>.ncu-rep \
--force-overwrite \
tuning/<gpu>/cute-cpp/bin/<shape_key>/<model>/iter<NNN>_<tag> [args]
ncu --import tuning/<gpu>/cute-cpp/perf/<shape_key>/<model>/ncu_iter<NNN>.ncu-rep \
--csv --page raw \
> tuning/<gpu>/cute-cpp/perf/<shape_key>/<model>/ncu_iter<NNN>.csv
Allowed: CUTLASS C++ tile primitives, CuTe atoms, raw CUDA intrinsics, PTX.
Forbidden: cutlass::gemm::device::Gemm (top-level library GEMM), cuBLAS.
| Bottleneck | Ideas |
|---|---|
| Memory-bound | Increase ThreadblockShape; add pipeline stages; TMA (Hopper) |
| Compute-bound | Change WarpShape; increase InstructionShape |
| Latency-bound | Increase kStages; use persistent kernel |
| L2 locality | Apply threadblock swizzle epilogue |
The binary must accept --verify (or run verify by default) and:
VERIFY: PASS or VERIFY: FAIL max_abs_err=<v>Verification Tolerance:
| Precision | base_tol | rel_tol |
|---|---|---|
| FP16 input, FP32 accum | 1.0 | 0.01 |
| FP16 input, FP16 accum | 16.0 | 0.05 |
| FP8 E4M3 input, FP16 accum | 0.5 | 0.01 |
NEVER print "Test Passed" without actual numerical comparison. Timing via CUDA events. Default: 10 warmup + 50 timed.
cuBLAS harness (.cu) — calls cublasGemmEx. Library calls allowed only in iter000.