skills.rest
xchang1121
xchang1121COMMUNITY
Explore all skills by this author →
GitHub

tilelang-cuda-gemm

Generates and optimizes TileLang CUDA GEMM kernels using swizzling, pipelining, and autotuning techniques.

6|1|Updated Apr 19, 2026
Search Tags:#gpu#cuda#tilelang#gemm#kernel optimization#autotuning#tensor core
One-click install
npx skills add https://github.com/xchang1121/op-autoresearch --skill tilelang-cuda-gemm-xchang1121
Or copy as Structured Prompt for Agent▼
Please help me install this Agent Skill.
Skill: tilelang-cuda-gemm
Source: https://github.com/xchang1121/op-autoresearch/tree/main/skills/tilelang-cuda/guides/tilelang-cuda-gemm
Command: npx skills add https://github.com/xchang1121/op-autoresearch --skill tilelang-cuda-gemm-xchang1121

SYSTEM DOCUMENTATION & REQUIREMENTS

💡 This Skill requires tilelang, torch.

What problem does it solve? Writing high-performance matrix multiplication kernels in TileLang requires deep knowledge of memory hierarchies, Tensor Core layouts, and GPU-specific tuning parameters. This Skill provides a complete, structured reference for generating and optimizing GEMM kernels on CUDA, from basic templates to advanced techniques like Split-K and fine-grained MMA. ## Core Features & Use Cases - Layered Optimization Guidance: Covers five optimization levels, from basic tiled GEMM templates through swizzling, L2 rasterization, auto-pipelining, persistent kernels, and Split-K/Stream-K decomposition. - Autotuning Workflows: Includes AutoTuner search-space construction, Roller (BitBLAS) hint-based configs, and SM-version heuristics for A100/H100 targets. - Specialized GEMM Variants: Provides patterns for FP8 GEMM, transpose-B layouts, dequant GEMM, and fine-grained MMA via TensorCoreIntrinEmitter, plus profiling with do_bench and a common-errors checklist. - Use Case: When asked to write a fast FP16 matmul kernel for an H100, follow the guide to start from the basic template, add swizzle layouts and rasterization, then autotune block sizes and pipeline stages to reach target TFLOPs. ## Quick Start Use the tilelang-cuda-gemm skill to write and optimize a TileLang GEMM kernel for my matrix multiplication workload.

Frequently Asked Questions about tilelang-cuda-gemm

High-intent search queries and answers about installing and using this skill.

FAQPage Schema
How do I write a GEMM kernel in TileLang for CUDA?▼

Start with the basic template: allocate shared memory with T.alloc_shared, fragments with T.alloc_fragment, then loop over K with T.Pipelined using T.copy for data movement and T.gemm for compute. Use block sizes of 128x128x32, 128 threads, and num_stages=3 as defaults.

How to autotune TileLang GEMM block sizes and pipeline stages?▼

Use AutoTuner.from_kernel with a config list spanning block_M, block_N, block_K, num_stages, and thread counts, then call run with warmup and rep parameters. Alternatively, use Roller (BitBLAS) MatmulTemplate to generate recommended hints, or apply SM-version heuristics for A100 and H100.

Does TileLang support FP8 GEMM on CUDA?▼

Yes, FP8 GEMM is supported using dtypes from determine_fp8_type() such as e4m3 or e5m2. It requires transpose_B=True with B shaped (N, K), float32 accumulation, and validation via calc_diff rather than torch.testing.assert_close.

When should I use persistent kernels vs Split-K in TileLang?▼

Use persistent kernels for large matrices (M, N > 4096) to reduce launch overhead by limiting the grid to the SM count. Use Split-K when K is much larger than M and N, splitting K across an extra grid dimension with atomic accumulation of partial results.

Why does my TileLang GEMM produce wrong results or bank conflicts?▼

Common causes include forgetting T.clear(C_local) before accumulation, missing make_mma_swizzle_layout annotations causing shared memory bank conflicts, or omitting transpose_B=True for FP8. Also check that num_stages does not exceed shared memory limits.

Related Skills

Explore More Software Engineering Skills→
vipshopvipshop
1.3k

cuda-cpp-kernel

Implements, debugs, and optimizes CUDA C++ and PTX kernels across NVIDIA GPU architectures.

Official
Advanced
mattmirelesmattmireles
57

coreml-validate

Validate CoreML model numerical accuracy against PyTorch reference models.

Community
Advanced
NVIDIANVIDIA
2.8k

perf-nsight-compute-analysis

Analyze Nsight Compute .ncu-rep reports to classify CUDA kernel bottlenecks.

Official
Advanced
FutureMLS-LabFutureMLS-Lab
556

add-sgl-kernel

Guide adding AOT CUDA or C++ kernels with Torch registration, CMake, tests, and benchmarks.

Official
Advanced