name: croq-dsl-cuda description: DSL-specific tuning contract for CUDA kernels. Loaded by croq-tune when dsl=cuda.

Croq-DSL: CUDA

Source extension: .cu | Compiler: nvcc -> static binary | Group: compiled-binary

Environment Validation

No additional checks beyond the standard ncu/nvcc/GPU validation.

BUILD / RUN Templates

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>/cuda/bin/<shape_key>/<model>/iter<NNN>_<tag> \
     tuning/<gpu>/cuda/srcs/<shape_key>/<model>/iter<NNN>_<tag>.cu \
     2>&1 | tee tuning/<gpu>/cuda/perf/<shape_key>/<model>/build_iter<NNN>.txt

run_iter<NNN>.sh:

#!/usr/bin/env bash
tuning/<gpu>/cuda/bin/<shape_key>/<model>/iter<NNN>_<tag> \
    2>&1 | tee tuning/<gpu>/cuda/perf/<shape_key>/<model>/timing_iter<NNN>.txt

Binary must print: TFLOPS: <value> time_ms: <value>

PROFILE — ncu Command

ncu --set full \
    --export tuning/<gpu>/cuda/perf/<shape_key>/<model>/ncu_iter<NNN>.ncu-rep \
    --force-overwrite \
    tuning/<gpu>/cuda/bin/<shape_key>/<model>/iter<NNN>_<tag> [args]

ncu --import tuning/<gpu>/cuda/perf/<shape_key>/<model>/ncu_iter<NNN>.ncu-rep \
    --csv --page raw \
    > tuning/<gpu>/cuda/perf/<shape_key>/<model>/ncu_iter<NNN>.csv

Pure Implementation Rule

Allowed: Raw CUDA kernels, CUDA intrinsics (__shfl_sync, __ldg, etc.), PTX inline asm (asm volatile), cooperative groups, CuTe/CUTLASS atoms only (copy atoms, MMA atoms, layout algebra).

Forbidden: cublasSgemm, cublasGemmEx, cutlass::gemm::device::Gemm, cuSPARSELt, cuDNN compute calls, PyTorch/TensorFlow ops.

IDEA Menu

Key levers: BM/BN/BK (multiples of 16, typically 64-256), warp count (blockDim), pipeline stages (1-5), smem padding, register pressure (#pragma unroll N, --maxrregcount), vectorized loads, warp primitives, PTX inline (ldmatrix, mma.sync), TMA.

VERIFY / MEASURE

The binary must accept --verify (or run verify by default) and:

1. Compute with custom kernel
2. Compute reference (cuBLAS in harness, or prior verified iter)
3. Print VERIFY: PASS or VERIFY: FAIL max_abs_err=<v>

Verification Tolerance:

NEVER print "Test Passed" without actual numerical comparison. Keep reference tensors alive in host code. Timing via CUDA events (10 warmup + 50 timed).

BASELINE (iter000)

cuBLAS harness (.cu) — calls cublasGemmEx. Library calls allowed only in iter000.