mcpbeat Sign in

Tilegym Adding Cutile Kernel Agent Skill

by NVIDIA

Add a new cuTile GPU kernel operator to TileGym. Covers dispatch registration in ops.py, cuTile backend implementation, __init__.py exports, test creation, and benchmark in tests/benchmark. Use when adding, creating, or implementing a new cuTile operator/kernel in TileGym, or when asking how to register a new cuTile op.

7k tokens
context cost
the whole folder, loaded on every use
5
files
instructions only
0
copies elsewhere
how many repositories repackaged it
789
stars on the repo
on the repository, not the skill itself

Install

one command, takes just this skill from the repository 
npx skills add https://github.com/NVIDIA/TileGym --skill tilegym-adding-cutile-kernel

What comes with it

18 164 bytes besides the instruction
BENCHMARK.md
evals/evals.json
skill-card.md
skill.oms.sig

The instruction itself

9 sections, as written by the author

Adding a cuTile Kernel to TileGym

End-to-end workflow for adding a new operator (e.g., my_op) with cuTile backend.

Execution Rules

MUST follow these rules strictly:

  • Use TodoWrite to create the checklist below BEFORE writing any code
  • Execute steps in order — do NOT skip ahead or combine steps
  • Mark each todo as completed after finishing, in_progress when starting
  • If a step is not applicable (e.g., no cuTile impl), mark it completed with a note, do NOT silently skip
  • Each step MUST result in a file write or explicit skip decision — no silent omissions

Instructions

MUST copy this checklist to TodoWrite at the start:

- [ ] Step 1: Register dispatch interface in ops.py
- [ ] Step 2: Implement cuTile backend
- [ ] Step 3: Register in __init__.py (cutile)
- [ ] Step 4: Add tests
- [ ] Step 5: Add benchmark to tests/benchmark
- [ ] Step 6: Verify (run pytest + lint)

Step 1: Register dispatch interface

File: src/tilegym/ops/ops.py

Add a @dispatch function — this is the single entry point for all backends.

@dispatch(
    "my_op",
)
def my_op(
    input: torch.Tensor,
    out: Optional[torch.Tensor] = None,
    **kwargs: Any,
):
    """
    Description of my_op.

    Args:
        input: Input tensor
        out: Optional preallocated output tensor
        **kwargs: Additional arguments for backend-specific configurations

    Returns:
        torch.Tensor
    """
    raise NotImplementedError(f"my_op is not implemented for {get_current_backend()}")

Key rules:

  • Function body only raises NotImplementedError
  • Include **kwargs for backend-specific parameters

Reference: See existing ops in src/tilegym/ops/ops.py (e.g., silu_and_mul, softmax)

Step 2: Implement cuTile backend

File: src/tilegym/ops/cutile/my_op.py

The file structure follows this template:

import torch
import cuda.tile as ct

from tilegym.backend import register_impl


@ct.kernel
def my_op_kernel_ct(x, output, n_elements: ct.Constant[int], BLOCK_SIZE: ct.Constant[int]):
    bid = ct.bid(0)
    indices = bid * BLOCK_SIZE + ct.arange(0, BLOCK_SIZE)
    x_val = ct.gather(x, indices)
    # ... compute ...
    ct.scatter(output, indices, result)


@register_impl("my_op", backend="cutile")
def my_op(input: torch.Tensor, out: torch.Tensor = None, **kwargs) -> torch.Tensor:
    n = input.numel()
    if out is None:
        out = torch.empty_like(input)
    grid = ((n + 1023) // 1024,)
    ct.launch(stream, grid, kernel, (some args, ...))
    return out

Reference: src/tilegym/ops/cutile/silu_and_mul.py

Step 3: Register in __init__.py (CRITICAL)

Missing this step means the cuTile backend implementation never gets loaded.

File: src/tilegym/ops/cutile/__init__.py

Add inside if is_backend_available("cutile"): block (alphabetically):

from . import my_op

And in the function import section:

from .my_op import my_op

And add "my_op" to __all__.

Step 4: Add tests

File: tests/ops/test_my_op.py

CRITICAL: Always import from tilegym.ops, NEVER from tilegym.ops.cutile.my_op.

import pytest
import torch

from tilegym.backend import is_backend_available, set_backend
from .. import common

_backends = ["cutile"]


class Test_MY_OP(common.PyTestCase):
    @staticmethod
    def reference(input):
        """Reference implementation using PyTorch."""
        return torch.some_reference(input)

    @pytest.mark.parametrize("shape, dtype", [
        ((1024,), torch.float16),
        ((1024, 512), torch.float32),
        ((64, 64, 64), torch.bfloat16),
    ])
    @pytest.mark.parametrize("backend", _backends)
    def test_op(self, shape, dtype, backend, arch):
        if backend == "cutile" and not is_backend_available("cutile"):
            pytest.skip("Cutile backend not available")
        try:
            set_backend(backend)
        except Exception as e:
            pytest.skip(f"Backend is not supported: {e}")

        self.setUp()

        from tilegym.ops import my_op

        A = torch.randn(*shape, dtype=dtype, device="cuda")
        self.assertCorrectness(
            my_op, self.reference, {"input": A},
            atol=1e-3, rtol=1e-3,
        )

Key patterns:

  • _backends = ["cutile"]
  • test_op: use set_backend(backend) with try-except, call self.setUp()

Reference: tests/ops/test_silu_and_mul.py

Below is the common errors.

1. Missing _backends list (inside class)
2. test_op / test_op_xxx — missing @pytest.mark.parametrize("backend", _backends), backend parameter, and tilegym.is_backend_available / tilegym.set_backend pattern

Step 5: Add benchmark to tests/benchmark

File: tests/benchmark/bench_my_op.py

Key rules from benchmark_rules.md:

  • Call the op via tilegym.ops.my_op(a, b, ..., backend=backend) — do not use set_backend.
  • Define ALL_BACKENDS (include at least cutile and torch), filter with get_supported_backends().
  • Implement reference_my_op(...) and register it: register_impl("my_op", "torch")(reference_my_op).
  • Use create_benchmark_config() to build triton.testing.Benchmark configs (e.g. by shape/dtype).
  • Use @triton.testing.perf_report([...]) on bench_my_op(...); inside the bench function: correctness check with torch.testing.assert_close(fn(), ref(), ...), then ms = triton.testing.do_bench(fn) (or do_bench_cudagraph), compute GB/s or TFLOPS, and return the metric.
  • Entry point: if __name__ == "__main__": bench_my_op.run(print_data=True).

Template structure:

import torch
import triton
import triton.testing

import tilegym
from tilegym.backend import is_backend_available, register_impl

ALL_BACKENDS = [
    ("cutile", "cuTile", ("orange", "-")) if is_backend_available("cutile") else None,
    ("torch", "PyTorch", ("green", "-")),
]

def get_supported_backends():
    return [p for p in ALL_BACKENDS if p is not None]

def reference_my_op(input: torch.Tensor, out: torch.Tensor = None, **kwargs):
    """Reference implementation using PyTorch."""
    ...

register_impl("my_op", "torch")(reference_my_op)

def create_benchmark_config(datatype, ...):
    available_backends = get_supported_backends()
    if not available_backends:
        return None
    backends, names, styles = zip(*available_backends)
    return triton.testing.Benchmark(
        x_names=["M"],  # or other dimension names
        x_vals=[...],
        line_arg="backend",
        line_vals=list(backends),
        line_names=list(names),
        styles=list(styles),
        ylabel="GB/s",  # or TFLOPS
        plot_name="my-op-...",
        args={"datatype": datatype, ...},
    )

@triton.testing.perf_report([
    create_benchmark_config(datatype, ...)
    for datatype in [torch.float16, torch.float32]
    for ... in [...]
])
def bench_my_op(M, backend, datatype, ..., device="cuda"):
    x = torch.randn(..., dtype=datatype, device=device)

    fn = lambda: tilegym.ops.my_op(x, backend=backend)
    ref = lambda: reference_my_op(x)
    torch.testing.assert_close(fn(), ref(), rtol=1e-2, atol=1e-2)

    ms = triton.testing.do_bench(fn)  # or do_bench_cudagraph(fn)
    # Compute metric (e.g. GB/s or TFLOPS) from ms and problem size
    return metric

if __name__ == "__main__":
    bench_my_op.run(print_data=True)

Benchmark Plot Names: Must include -TFLOPS or -GBps suffix

  • Example: plot_name=f"persistent-layer-norm-M{num_rows}-{dtype_name}-GBps"

Step 6: Verify

# Run tests
pytest tests/ops/test_my_op.py -v

# Run benchmark (optional)
python tests/benchmark/bench_my_op.py

# Lint
pre-commit run -a

Other skills for the same job

different authors, same section of the catalogue
Scikit Learn
by christophacham
×3

Machine learning in Python with scikit-learn. Use when working with supervised learning (classification, regression), unsupervised learning (clustering, dimensionality reduction), model evaluation, hyperparameter tuning, preprocessing, or building ML pipelines. Provides comprehensive reference documentation for algorithms, preprocessing techniques, pipelines, and best practices.

30k tokens scripts
Scikit Learn
by ComeOnOliver
×3

Machine learning in Python with scikit-learn. Use when working with supervised learning (classification, regression), unsupervised learning (clustering, dimensionality reduction), model evaluation, hyperparameter tuning, preprocessing, or building ML pipelines. Provides comprehensive reference documentation for algorithms, preprocessing techniques, pipelines, and best practices.

32k tokens scripts
LLM Evaluation
by ComeOnOliver
×2

Implement comprehensive evaluation strategies for LLM applications using automated metrics, human feedback, and benchmarking. Use when testing LLM performance, measuring AI application quality, or establishing evaluation frameworks.

6k tokens
LLM Evaluation
by ComeOnOliver
×2

Implement comprehensive evaluation strategies for LLM applications using automated metrics, human feedback, and benchmarking. Use when testing LLM performance, measuring AI application quality, or establishing evaluation frameworks.

6k tokens
Agent Evaluation
by lingxling
×1

Testing and benchmarking LLM agents including behavioral testing, capability assessment, reliability metrics, and production monitoring—where even top agents achieve less than 50% on real-world benchmarks

9k tokens
Celltypist Cell Annotation
by BioTender-max
×1

Automated scRNA-seq cell type annotation via pre-trained logistic regression. 45+ models: immune, gut, lung, brain, fetal, cancer microenvironments. Input normalized AnnData; outputs per-cell labels, majority-vote cluster labels, confidence scores. Use for fast, reference-backed annotation without manual marker inspection.

5k tokens
Scikit Learn Machine Learning
by BioTender-max
×1

Classical ML in Python: classification, regression, clustering, dim reduction, evaluation, tuning, preprocessing pipelines. Linear models, tree ensembles, SVMs, K-Means, PCA, t-SNE. Use PyTorch/TF for deep learning; XGBoost/LightGBM for scale.

4k tokens
Statsmodels Statistical Modeling
by BioTender-max
×1

Python statistical modeling: regression (OLS, WLS, GLM), discrete (Logit, Poisson, NegBin), time series (ARIMA, SARIMAX, VAR), with rigorous inference, diagnostics, and hypothesis tests. Use scikit-learn for ML; statistical-analysis for test choice.

4k tokens

How to use it

Copy the folder

Take nvidia/tilegym-adding-cutile-kernel from the repository into ~/.claude/skills for personal use, or into .claude/skills inside a project.

Check the name does not clash

The agent identifies a skill by the name field in its header. Two skills with the same name cannot sit side by side — one of them will be ignored.