tile-ai/tilelang

Features

• Python GPU Kernels - Provides a Python-embedded DSL for writing high-performance GPU kernels with automatic thread mapping and shared memory management.
• Autotuning GPU Kernel Compilers - Automatically searches over tile sizes and thread counts to optimize GPU kernel performance.
• GPU Kernel DSL Compilers - Provides the core Pythonic DSL for structuring high-performance kernels with automatic memory management and thread mapping.
• Attention Kernel Optimizers - Implements efficient attention mechanisms for transformers with custom tiling, pipelining, and memory access patterns.
• Memory-Compute Overlaps - Overlaps memory copies with computation in GPU kernels to hide memory latency.
• Convolution Kernel Configurators - Defines matrix-matrix convolution computations with configurable dimensions, data types, and optional bias.
• GEMV Shared Memory Tilers - Writes GEMV kernels that tile the vector and matrix into shared memory for efficient data reuse.
• Python Tile Kernel Interfaces - Provides a Pythonic DSL to express tile-based GPU kernels with automatic thread mapping and memory management.
• Tile-Based GPU Kernel Programming - Provides a tile-level programming model that abstracts individual thread management.
• GPU Hardware Architecture Modelings - Models GPU hardware specifications like register capacity and warp size to guide code generation.
• Hardware-Specific - Provides hardware-specific kernel configuration retrieval for automatic tuning across GPU architectures.
• Attention Compute Graph Builders - Sets up the compute graph for attention kernels including matrix multiplication, bias, and type casting.
• Multi-SM Attention Parallelizers - Divides attention computation across multiple streaming multiprocessors for parallel execution and merges results.
• Tile Data Loaders - Copies tiles between global, shared, and register memory with bounds checking.
• Automatic Tiling Configurators - Automatically generates and ranks tiling configurations for matrix and reduction kernels.
• Heuristic Tiling Optimizations - Ships a heuristic policy that automatically selects optimal tile configurations for GPU kernels.
• Multi-Level Tiling Optimizations - Implements multi-level tiling to maximize bandwidth and reduce latency across memory hierarchy.
• GPU Warp Specializations - Specializes GPU warps into producer-consumer configurations for automatic synchronization.
• Tensor Core Programming Frameworks - Provides a Pythonic framework for writing tensor core matrix multiply-accumulate operations.
• Tensor Core Matrix Multiplications - Performs matrix multiply-accumulate operations using tensor core instructions and shared memory.
• Thread-Level All-Reducers - Replaces global atomic adds with a thread-level all-reduce to reduce synchronization overhead in GPU kernels.
• GPU Buffer Allocators - TVM's feature to allocate a multi-dimensional buffer in the shared, local, fragment, or global memory space for use inside a GPU kernel.
• Buffer Initialization Operations - TVM's feature to fill every element of a buffer with a specified constant value, including a dedicated operation to zero it out.
• Single-Element Buffer Allocations - TVM's feature to allocate a single-element buffer in local or variable memory, optionally with an initial value.
• GPU Availability Detections - Checks GPU acceleration availability and queries target-specific hardware features for kernel compilation.
• Split-K Reduction Distributors - Distributes the reduction along the K dimension across multiple threads and combines results with atomic adds.
• Multi-Stage Memory Pipelines - Configures multi-stage pipelines to overlap memory access with computation in GPU kernels.
• Annotated Kernel Definitions - Defines GPU or CPU kernels by annotating functions with tensor shapes and data types.
• GPU Kernel Configuration Generators - Generates top-ranked GPU kernel configurations tailored to the target architecture for automatic tuning.
• Architecture Detection - Checks target device architecture to guide GPU kernel compilation and optimization.
• GPU Architecture Detections - Detects GPU architecture and compute capability to correctly target kernel compilation at runtime.
• Atomic Memory Operations - Provides atomic read-modify-write operations on shared and global GPU memory for safe concurrency.
• GPU Device Property Queries - Models GPU resources such as register capacity and shared memory limits for compiler tuning.
• Buffer Reshaping and Reinterpretation - TVM's feature to reinterpret a buffer with a new shape and optionally a new data type without copying the underlying data.
• GPU Architecture Identifications - Identifies GPU generation to guide architecture-specific kernel optimizations during compilation.
• JIT Kernel Compilers - Compiles GPU kernels on the fly and executes them immediately or returns a reusable object.
• CUDA Kernel Compilers - Compiles GPU kernels to CUDA device binaries with configurable architecture options.
• GPU Backend Selections - Selects GPU or CPU backend for kernel compilation with optional architecture specification for tuning.
• Runtime GPU Kernel Compilation Libraries - Compiles and caches GPU kernels at runtime for dynamic specialization without a separate build step.
• Compiled Kernel Executions - Executes compiled GPU kernels on the active stream with automatic pointer binding.
• TVM JIT Compilers - Compiles high-level kernel descriptions into device binaries at runtime using TVM's IR and LLVM backends.
• Dense and Sparse MMA Executions - Executes both dense and sparse matrix multiply-accumulate operations on Tensor Cores.
• GPU Kernel Code Generators - Generates executable GPU code from intermediate representations for target platforms.
• Grid-Wide Synchronizations - Provides grid-wide thread block synchronization for cooperative kernel launches.
• Kernel Autotuning Frameworks - Automatically searches over tile sizes, thread counts, and scheduling policies to find the fastest kernel configuration.
• Kernel Fusion Operations - Combines successive operations into fused kernels to minimize global memory traffic.
• Scoped GPU Memory Allocations - TVM's feature to allocate on-chip shared memory, per-thread fragments, or scalar variables under a block-scoped zone.
• Software Pipelining Optimizations - Implements automatic multistage software pipelining to overlap memory loads with computation in GPU kernels.
• TMA Data Loads - TVM's feature to use special TMA hardware to efficiently load multidimensional tensors from global into shared memory.
• TMA Data Stores - TVM's feature to use special TMA hardware to write data from shared back to global memory, optionally performing atomic reduction.
• TMA Data Transfer Initiations - TVM's feature to initiate an asynchronous data load from global memory to shared memory using a TMA descriptor, supporting standard and 2SM configurations.
• Configurable Matrix Multiplication Kernels - Configures and generates high-performance matrix multiplication kernels with adjustable transposition, data types, and bias.
• Tile-Based Matrix Multiplications - Performs tile-sized matrix multiplication using shared memory and tensor cores.
• Kernel Scheduling Hints - Specifies tiling and tensor core hints to guide optimized GPU kernel compilation.
• Threadblock Swizzling Configurations - TVM's feature to remap the scheduling order of threadblocks using configurable patterns to improve L2 cache hit rates.
• Memory Access Pattern Optimizers - Optimizes memory access patterns using layout annotations, swizzling, and pipelining for GPU kernels.
• GPU Compute Capability Inspections - Inspects GPU compute capability and instruction sets to specialize kernels for different hardware generations.
• GPU Capability Queries - Detects GPU compute capability, tensor core availability, and data type support at compile time.
• Hardware-Specific Attention Configurators - Fetches hardware-specific configuration hints for attention kernels on the target GPU.
• Hardware-Aware Reduction Configurators - Generates hardware-aware configurations for reduction kernels tailored to the target architecture.
• Hardware-Aware Convolution Configurators - Retrieves optimized convolution configurations tailored to the target GPU architecture.
• Dynamic Tensor Shapes - Accepts symbolic dimension values so a single compiled kernel adapts to varying input sizes without recompilation.
• Tensor Instruction Shape Queries - Queries supported tensor instruction shapes from the target architecture for instruction selection.
• Elementwise Tile Math Operations - Applies elementwise math functions like exp, log, and sigmoid on tile fragments.
• Elementary Math Functions - TVM's feature to compute exponential, logarithmic, trigonometric, and other elementary math operations within GPU/CPU kernel code.
• FP8 Data Type Resolutions - Resolves the correct FP8 data type representation for the current GPU platform architecture.
• Low-Precision Decoding - TVM's feature to decode low-precision integer or floating-point data into half-precision formats using inline PTX operations.
• Kernel Caching Systems - Persists GPU binaries and launcher code to disk for reuse across sessions.
• Cached Kernel Loading - Loads compiled kernels from cache to avoid recompilation across sessions.
• Sparse Metadata Loading Implementations - Loads sparsity metadata into shared memory for use with sparse tensor core instructions.
• PyTorch Kernel Integrations - Wraps native function calls to execute in the current PyTorch GPU stream and device context.
• Cumulative Sum Calculators - Calculates cumulative sum or maximum along a sequence in one or two dimensions, optionally in reverse.
• Warp-Level Reductions - Provides warp and block-level reduction operations for aggregating values across GPU threads.
• Tile Reductions - Computes sum, min, max, and cumulative reductions on tile fragments.
• Tile Iteration Pattern Configurations - Provides configurable tile traversal patterns to optimize memory access patterns.
• Buffer Layout Inferences - Automatically infers buffer layouts from operator specifications to generate efficient GPU kernel code.
• Sparse Tensor Core Instruction Generation - Generates low-level sparse tensor core instructions for matrix multiply-accumulate.
• Tensor Core Instruction Generation - Produces the final tensor core machine instructions from shared memory fragments.
• Tensor Core Operand Loadings - Loads matrix operands from shared memory into registers for tensor core matrix multiply-accumulate.
• Tensor Core Result Storage Implementations - Implements the final write-back step for tensor core matrix results into shared memory.
• DSL Compilers - Converts DSL source code into a compiled object executable on GPU or CPU hardware.
• Tensor Core Layout Translations - Translates between shared memory tile layouts and MMA register file layouts for tensor cores.
• DLPack Protocols - Wraps tensors from DLPack-compatible frameworks for direct kernel execution.
• Vectorized Memory Access - Provides vectorized memory reads to maximize memory bandwidth in GPU kernels.
• GPU Kernel Print Debugging - Provides assertion and print functions that operate on the GPU device for kernel debugging.
• Asynchronous Operation Synchronizations - Synchronizes asynchronous data movement with computation using barrier primitives.
• Synchronization Barrier Allocations - Allocates synchronization barriers for coordinating thread arrivals in GPU kernels.
• Math Library Accelerators - Provides accelerated implementations of common math functions on GPU and CPU.
• GPU Occupancy Controllers - Controls GPU occupancy by hinting minimum thread blocks per SM to optimize register usage.
• Schedule Block Retrievals - Retrieves specific schedule blocks from the compiler for targeted optimization and analysis.
• Kernel-to-Template Bindings - Binds kernel functions to hardware-aware templates for targeted configuration generation.
• GPU Kernel Assertions - Implements GPU kernel assertions that halt execution and optionally print a message when a condition fails.
• Persistent Thread Block Patterns - Implements persistent thread-block patterns for dynamic work distribution across GPU thread blocks.
• Shared Memory Swizzling - Provides shared memory swizzling to reduce bank conflicts in GPU kernel execution.
• Sparse MMA Layout Conversions - Implements shared memory to MMA sparse layout conversion for tensor core instructions.
• Warp-Level Matrix Multiply-Accumulates - Generates warp-level tensor core instructions for matrix multiply-accumulate operations.
• C++ Kernel Launcher Generators - Generates C++ launcher code for GPU kernels with automatic memory management.
• Cluster Thread Block Synchronizations - Coordinates thread blocks within GPU clusters using barrier primitives.
• Async Proxy Operation Fencings - Guarantees correct ordering between generic and async memory accesses in GPU kernels.
• AMD GPU Kernel Compilers - Compiles GPU kernels for AMD architectures with automatic device library and linker location.
• HIP Kernel Compilers - Compiles GPU kernels into HIP device binaries for AMD architectures.
• Kernel Artifact Caches - Caches compiled kernel binaries and source code to disk for reuse across sessions.
• NVRTC Kernel Compilers - Uses NVRTC to compile GPU kernels at runtime for immediate execution.
• Batch Kernel Compilations - Batch-compiles kernel configurations in parallel to accelerate autotuning.
• Parallel Kernel Compilations - Compiles kernel parameterizations in parallel using multiple workers.
• Kernel Performance Benchmarks - Benchmarks compiled GPU kernel execution time with test tensors to measure latency.
• GPU Asynchronous Data Transfers - TVM's feature to move data from global to shared memory without blocking, enabling computation to overlap with data movement.
• Inter-Level GPU Memory Copies - TVM's feature to copy data between specified memory buffers or address spaces, such as between global and shared memory in GPU kernels.
• Kernel Launch Context Setups - Configures grid and thread dimensions for GPU kernel launches.
• Kernel Library Compilations - Compiles kernel source code into shared libraries for target devices and manages their lifecycle.
• Kernel Shared Library Generators - Generates compiled shared libraries from kernel specifications for integration with Python.
• Matrix Fragment Transfers - TVM's feature to transfer matrix fragments by loading or storing 8×8 fragments between shared memory and registers using PTX matrix instructions.
• Memory Hierarchy Data Movements - TVM's feature to move data between global, shared, and fragment memories and overlay computation with communication via software pipelining.
• MMA Shape Policies - Configures tensor core shape policies for generating efficient MMA instructions.
• Multi-Pass Compiler Pipelines - Defines a sequence of compiler passes to lower IR modules for specific hardware targets.
• Multi-Platform Code Generators - Generates GPU kernel code for multiple hardware backends through a unified interface.
• Thread and Block Index Queries - Provides built-in primitives to query thread and block indices within GPU kernels.
• Tile-Based GPU Memory Movements - TVM's feature to copy tiles between global, shared, and register‑file memory scopes, with an asynchronous variant for manual prefetch pipelining.
• TMA Descriptor Creation - Provides TMA descriptor creation for asynchronous data movement in GPU kernels.
• Parameterized Numeric Types - Defines parameterized numeric types with configurable precision and FP4 detection for GPU kernels.
• Warp-Group Matrix Multiply-Accumulates - Performs warp-group level matrix multiply-accumulate for high-performance GPU kernels.
• Sparse Matrix Tile Loaders - Loads and transposes matrix tiles from shared memory for sparse tensor core operations.
• Structured - Implements 2:4 structured sparse matrix multiplication using specialized tensor core instructions.
• Multi-Dimensional Buffer Accesses - Supports multi-dimensional indexing and slicing of buffers within GPU kernels.
• Arithmetic Operations - Translates Python arithmetic operators into device-side computations in GPU kernels.
• Sparse Matrix Tile Store Operations - Stores matrix tiles in swizzled shared memory layout for sparse tensor core use.
• Vector Dot Product Kernels - Performs packed 4-element dot product accumulation for efficient integer matrix multiply.
• Hardware-Specific Loop Structuring - Structures loop iteration with serial, unrolled, parallel, and software-pipelined constructs for GPU hardware.
• Architecture-Specific - Analyzes kernel functions to suggest top-k optimization hints for the target architecture.
• Buffer Performance Hints - Annotates buffers with layout and cache hints to guide compiler optimization of GPU kernels.
• Tuning Result Persistence - Persists tuning results to disk so the best configuration can be reused across sessions.
• Block Index Swizzlings - TVM's feature to map block indices to row-major or column-major swizzled rasterization coordinates to improve memory coalescing.
• GPU Kernel Profilers - Measures kernel latency using a built-in profiler to evaluate GPU kernel performance.
• GPU Kernel Code Block Definitions - Marks code blocks with decorators for compile-time inlining as reusable device functions.

Star history