tilelang-cuda-skills

Use torch.profiler as a lightweight, always-available alternative to Nsight Compute/Systems for profiling, debugging, and optimizing TileLang GPU kernels. Trigger this skill whenever the user wants to profile a TileLang kernel without installing ncu/nsys, get an operator/kernel breakdown of a forward pass that includes TileLang custom ops, hunt launch-overhead in a multi-kernel pipeline, classify whether a TileLang kernel is IO-bound / CUDA-core bound / Tensor-core bound, generate a Chrome trace, profile GPU memory allocations, or compare a TileLang kernel against a torch/cuBLAS reference inside one trace. Also trigger on phrases like "profile this kernel", "why is this slow", "how do I see the kernel timeline", "lightweight profiling", "torch profiler", "chrome trace", or "perfetto", even if the user does not mention TileLang explicitly but TileLang code is on the table.

How to optimize TileLang GPU kernels for better performance. Use this skill whenever the user has a working TileLang kernel that is too slow, wants to tune tile sizes or pipeline stages, needs to improve TFLOPS or bandwidth utilization, wants to use the AutoTuner, or asks questions like "how do I make my kernel faster", "what tile sizes should I use", "how do I autotune", or "why is my kernel slower than cuBLAS". Also trigger when the user mentions block sizes, num_stages, threads per block, shared memory pressure, occupancy, swizzle, vectorization, split-K, or any performance tuning in TileLang context. Even for vague requests like "improve performance" or "speed up this kernel" when TileLang is involved.

How to write and test TileLang kernels that need both forward and backward passes. Use this skill whenever the user is implementing custom operators with gradients, writing attention forward+backward, linear attention fwd+bwd, any op used inside torch.autograd.Function, or debugging gradient mismatches. Also trigger when the user mentions backward pass, gradients, dQ/dK/dV, autograd, ctx.save_for_backward, atomic_add for gradient accumulation, or gradient testing in TileLang context. Use even for questions like "how do I test my TileLang backward kernel" or "my gradients don't match the reference".

How to diagnose and fix errors in TileLang programs. Use this skill whenever a TileLang kernel fails to compile, crashes at runtime, produces incorrect results, or shows numerical mismatches. Also use when the user encounters TVM/TIR errors, CUDA codegen failures, shape mismatches, assertion failures, NaN/inf outputs, tilelang.compile errors, or needs to inspect generated CUDA code, use T.print for debugging, run AutoDD for minimal reproduction, or compare against a PyTorch reference. Trigger even for vague complaints like "my kernel doesn't work" or "wrong output" when TileLang is involved. Also trigger for TMA descriptor errors, illegal memory access, or kernel launch failures.

How to benchmark and profile TileLang GPU kernels for performance measurement. Use this skill whenever the user wants to measure kernel latency, compare TileLang performance against PyTorch/cuBLAS/Triton, use do_bench or get_profiler, calculate TFLOPS or bandwidth, run Nsight Compute or Nsight Systems on a TileLang kernel, understand occupancy or shared memory usage, or figure out why a kernel is slow. Also trigger when the user mentions benchmarking, profiling, timing, throughput, performance numbers, or latency in the context of TileLang or GPU kernels. Even for simple questions like "how fast is my kernel" or "how do I time this".

How to write TileLang GPU kernels from scratch or by adapting existing patterns. Use this skill whenever the user wants to create a new TileLang kernel, implement a GPU operator in TileLang, write a prim_func, set up tiling/blocking, define shared memory and fragment allocations, use T.gemm or T.copy, create a JIT-compiled kernel, or port an algorithm (GEMM, elementwise, reduction, attention, normalization, softmax, convolution) to TileLang. Also use when the user asks about tilelang.jit, T.prim_func, T.Kernel, T.Tensor, T.alloc_shared, T.alloc_fragment, T.Pipelined, T.Parallel, or any TileLang DSL construct for kernel authoring. Even for questions like "how do I start with TileLang" or "show me a basic TileLang example", this skill provides the right templates and workflow.