Menú
DSL-specific tuning contract for CuTe/CUTLASS C++ template kernels. Loaded by croq-tune when dsl=cute-cpp.
Instalar con Codex o Claude Copia este prompt, pégalo en Codex, Claude u otro asistente, y deja que revise la página de la skill y la instale por ti.
Basado en la clasificación ocupacional SOC
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.