All skills
nvidia avatar

/tilegym-converting-cutile-triton-to-cutile-rs

@48a26cf
by NVIDIA Corporationnvidia/tilegym821 stars
89

Use this skill to convert, port, or translate Triton-TileIR or cuTile-Python GPU kernels to cutile-rs (Rust). The orchestrator runs scripts/preflight.sh, then drives a bounded Agent A -> B -> D -> E pipeline (Agent C is diagnostic, Agent F optional), delegating all kernel/host/correctness/perf work to sub-agents and routing by each stage's single-line VERDICT.

  • 44 files
  • 452.9 KB
  • CC-BY-4
  • Updated 2 months ago
  • GitHub

Use this Skill: https://skilld.dev/gh/nvidia/tilegym/tilegym-converting-cutile-triton-to-cutile-rs

This session only. Nothing lands on disk.

referencesenv-block.md

≈838 tokens on demand. Your agent reads this file only when SKILL.md points to it.

Environment Block for Agent Prompts

Required env vars

Everything now lives in the tilegym checkout — there is no separate cutile-rs checkout and no CUTILE_RS_ROOT. The kernel Rust and the aggregated cutile_kernels crate live under $TILEGYM_PATH/src/tilegym/ops/cutile_rs/, build against crates.io pins (cutile="=0.2.0", …; Rule 33), and produce a single libcutile_kernels.so. Set these once per shell session; scripts/preflight.sh verifies them.

# ─── Git repo ─────────────────────────────────────────────────────────────
export TILEGYM_PATH=<your-tilegym-checkout>     # tests + Python wrappers + cutile_rs Rust

# ─── Toolchain (existence-checked) ────────────────────────────────────────
export TILEIRAS_BIN=<path-to-tileiras>          # cuTile compiler binary (IR->cubin at launch)
export CUDA_TILE_OPT_BIN=<path-to-cuda-tile-opt> # MLIR canonicalizer (Agent B/C)
export TRITON_TILEIR_PYTHONPATH=<triton-tileir/python>  # Triton-TileIR backend bindings
export CUDA_TOOLKIT_PATH=/usr/local/cuda        # CUDA root (cuda-bindings build.rs); default /usr/local/cuda

Building the cutile_kernels crate. The crate uses PINNED crates.io dependencies (no path deps), so there is nothing to sed-replace. The Python loader autobuilds it on first use — CUTILE_RS_AUTOBUILD is ON by default; set it to 0 to use a prebuilt library. Override the crate directory (default $TILEGYM_PATH/src/tilegym/ops/cutile_rs/cutile_kernels) with CUTILE_RS_KERNELS_DIR. The cuda-bindings build step reads CUDA_TOOLKIT_PATH (default /usr/local/cuda); tileiras lowers IR to cubin at launch time.

CRITICAL: tileiras alignment

cutile-python and cutile-rs MUST use the SAME tileiras. The binary pointed at by $TILEIRAS_BIN must come BEFORE /usr/local/cuda/bin/ on PATH so cuTile-py JIT picks it up.

Full manual environment

# Required env vars (incl. CUDA_TOOLKIT_PATH) — see top of file.

# tileiras must precede system CUDA on PATH
export PATH=$(dirname "$TILEIRAS_BIN"):$PATH

# Triton-TileIR bindings + tilegym source on PYTHONPATH
export PYTHONPATH=$TRITON_TILEIR_PYTHONPATH:$TILEGYM_PATH/src:$PYTHONPATH

export ENABLE_TILE=1
export TILEIR_ENABLE_FTZ=1
export TILEIR_ENABLE_APPROX=1

# Benchmark defaults
export CUPTI=1
export WARMUP=100
export REP=50
export MIN_REP=2

# CUDA_TOOLKIT_PATH is one of the required vars set at the top of this file.

Tool paths

# tileiras (for cutile-rs JIT → cubin, and cutile-python compilation)
which tileiras   # MUST be: $TILEIRAS_BIN
                 # NOT: /usr/local/cuda/bin/tileiras

# cuda-tile-opt (for canonicalize/CSE on dumped IR — used by Agent B / Agent C)
echo "$CUDA_TILE_OPT_BIN"

Verification

which tileiras
# Expected: $TILEIRAS_BIN

python -c "import triton; print(triton.__file__)"
# Expected: contains $TRITON_TILEIR_PYTHONPATH

python -c "import tilegym; print('OK')"
# Expected: OK

cd "$TILEGYM_PATH/src/tilegym/ops/cutile_rs/cutile_kernels" && cargo build --release 2>&1 | tail -1
# Expected: Finished `release` profile ...   (produces libcutile_kernels.so)

Source: SKILL.md on GitHub

No third-party reports yet.

Signed by skilld at 48a26cf. This ties the file your Agent reads to that commit on GitHub. It does not review the instructions.

Last checked against GitHub 4 days ago.

Activeupdated 2 months ago
metadata
{
  "author": "TileGym Team <TileGym@nvidia.com>"
}

README badge

README badge for nvidia/tilegym/tilegym-converting-cutile-triton-to-cutile-rs