Customizing Kernels with hipBLASLt TensileLite GEMM Tuning - Advanced User Guide

2026-04-06 · Source: AMD ROCm Blogs · Field: Technology & Digital — Artificial Intelligence & Machine Learning, Software Development & Engineering, Hardware Acceleration · Depth: Advanced, extended

Summary

This guide details TensileLite Tuning, an advanced optimization method within AMD's hipBLASLt library for General Matrix Multiply (GEMM) operations on AMD Instinct MI300X GPUs. Unlike Offline and Online tuning, which select from existing kernel pools, TensileLite generates entirely new, custom GEMM kernels tailored to specific problem sizes (M, N, K). The process involves capturing GEMM shapes, generating a configuration YAML, customizing parameters like thread organization, loop control, parallelism strategies (Split-K), memory access, and instruction configurations, then building and executing the tuning workflow. Benchmarks on an AMD Instinct MI300X with ROCm 6.4.3 and hipBLASLt 1.1.0 show TensileLite Tuning delivers an average 119% speedup over Offline Tuning and 225% over the Baseline, significantly improving efficiency for small, medium, and large workloads.

Key takeaway

For AI Engineers optimizing LLM inference on AMD Instinct MI300X GPUs, TensileLite Tuning is crucial for achieving peak GEMM performance. If your workloads have unique or critical (M, N, K) shapes, you should implement this advanced workflow to generate custom kernels, as it can yield up to 3.2x speedups over baseline performance and significantly surpass pre-compiled binaries, ensuring your applications run at maximum efficiency.

Key insights

TensileLite Tuning generates custom GEMM kernels for AMD GPUs, significantly outperforming existing solutions.

Principles

• Custom kernel generation maximizes hardware utilization.
• Parameter tuning balances GPU resource usage.
• Optimized kernels reduce latency and boost throughput.

Method

The TensileLite Tuning workflow involves capturing GEMM shapes, generating and customizing a configuration YAML, executing the tuning script to benchmark kernel candidates, and merging the best-performing kernel into the hipBLASLt library.

In practice

• Use `hipblaslt-bench` to capture target GEMM shapes.
• Adjust `WorkGroup`, `ThreadTile`, `LoopUnroll`, and `SplitU` parameters.
• Rebuild hipBLASLt after merging new kernel logic.

Topics

• hipBLASLt
• TensileLite Tuning
• GEMM Optimization
• AMD Instinct MI300X
• Kernel Generation

Code references

• ROCm/rocm-libraries

Best for: Machine Learning Engineer, AI Engineer, AI Hardware Engineer

Related on AIssential

• Building and Deploying Custom hipBLASLt Libraries on AMD Instinct GPUs — Custom "hipBLASLt" builds offer precise GPU optimization and flexible deployment for AMD Instinct GPUs, avoiding system-wide conflicts.
• hipBLASLt Online GEMM Tuning — hipBLASLt Online GEMM Tuning dynamically optimizes matrix multiplication at runtime for improved LLM performance.
• Fast transformer inference with Metal Performance Shaders — Metal Performance Shaders in Thinc PyTorch layers accelerate spaCy transformer inference on Apple Silicon Macs by up to 4.7x.
• AscendKernelGen: A Systematic Study of LLM-Based Kernel Generation for Neural Processing Units — Domain-specific reasoning and rigorous evaluation are crucial for LLM-based NPU kernel generation.
• AutoMegaKernel: A Statically-Checked Agent Harness for Self-Retargeting Megakernel Synthesis — AutoMegaKernel synthesizes safe, efficient, and retargetable megakernels for Llama models, outperforming cuBLAS on inference-class GPUs.
• From Tokens to Regions: CUDA-Sensitive Instruction Tuning for GPU Kernel Generation — CUDA sensitivity in LLM-generated kernels manifests at token and region levels, informing a low-cost instruction tuning method.

Open in AIssential →

Editorial summary, takeaway, and curation by AIssential. Original article published by AMD ROCm Blogs.