port-cuda-kernel-to-triton

Skill: Port a CUDA Kernel to Triton

Purpose

Guide the agent through systematically porting an existing CUDA kernel to Triton, mapping the CUDA execution model to Triton's tile-based program model, preserving numerical correctness, and identifying the patterns that do not translate directly.

Use this when

• An existing CUDA kernel must be made available as a Python-callable Triton kernel without rewriting the entire algorithm.
• The CUDA kernel implements a well-defined tiled operation (GEMM, softmax, layernorm, elementwise, reduction) where the tile structure is already clear.
• Prototyping speed matters and maintaining the Triton version is preferable to maintaining CUDA C++ for the target team.
• The kernel needs to run on hardware with good Triton support (NVIDIA A-series, H-series; AMD MI-series via ROCm Triton).

Do not use this when

• The CUDA kernel relies on warp shuffle instructions (__shfl_sync, __shfl_xor_sync) for its critical computation path. Triton has no direct warp shuffle API; the logic must be restructured to use tl.sum/tl.max or eliminated, which is a non-trivial redesign.
• The CUDA kernel uses complex intra-block data exchange patterns (e.g., warp-level matrix multiply with explicit register fragments via WMMA) that have no natural Triton equivalent. Porting will require restructuring the algorithm, not just translating syntax.
• The kernel depends on __threadfence, __threadfence_block, or other fine-grained memory fence semantics not present in Triton.
• The CUDA kernel uses dynamic shared memory in a way that depends on runtime-determined offsets or aliased smem regions. Triton manages smem implicitly and cannot be directed at this level.
• The kernel is already performance-critical and well-tuned in CUDA; a Triton port may not match its throughput without significant autotuning. Evaluate this tradeoff first.

Inputs the agent should gather first

• The complete CUDA kernel source, including all device functions it calls.
• The kernel's inputs and outputs: tensor shapes, dtypes, memory layouts (row-major, column-major, strided).
• The block dimensions (blockDim.x/y/z) and grid dimensions (gridDim.x/y/z) used in the launch.
• Which shared memory loads correspond to which input tensors, and which smem regions are reused across iterations.
• Whether the kernel contains warp shuffles, atomics, or texture reads.
• The target hardware architecture and Triton version (Triton API changes between versions for some ops).
• Whether autotuning of BLOCK_M, BLOCK_N, BLOCK_K, num_warps, and num_stages is planned or if a fixed configuration is required.