Tilus Documentation

Tilus is a research domain-specific language (DSL) designed for tile-level GPU kernel programming. It is intended for high-performance computing, particularly for low-precision workloads.

Key Features:

• Thread-block-level granularity: Kernels are defined at the thread-block level using tensors as the primary data type.
• Explicit memory control: Provides explicit control over shared memory and register tensors (unlike Triton).
• Low-precision support: Supports low-precision types with arbitrary bit-widths ranging from 1 to 8 bits.
• Developer productivity: Includes automatic tuning, caching, and a Pythonic interface.

The Blackwell Matmul tutorial series provides a progressive path from a minimal single-warp kernel to a high-performance kernel capable of reaching vendor-library-level performance on NVIDIA Blackwell GPUs. The series follows this evolution:

• V0: Minimal single-warp kernel.
• V1: Addition of TMA (Tensor Memory Accelerator) loads.
• V2: Implementation of software pipelining.
• V3: Implementation of warp specialization.
• V4: Addition of tile rasterization and a pipeline abstraction.
• V5: Implementation of CLC (Cluster-Level Control) persistent scheduling with a pipelined epilogue.
• V6: Implementation of 2-CTA clusters with distributed MMA (Matrix Multiply-Accumulate).

Tilus is a domain-specific language (DSL) for GPU programming that operates at thread-block-level granularity. It uses tensors as its core data type and provides developers with explicit control over shared memory and tensor layouts, distinguishing it from other DSLs like Triton.

Key capabilities include:

• Tensor-Centric Design: Tensors are the primary data abstraction.
• Explicit Memory Control: Direct management of shared memory and tensor layouts.
• Low-Precision Support: Support for low-precision types with arbitrary bit-widths.
• Optimization Features: Includes automatic tuning, caching, and a Pythonic interface for ease of use.

The tilus.ir module is the core Intermediate Representation (IR) package for Tilus. It provides the fundamental building blocks for defining GPU kernels at a tile level. The module is organized into three primary conceptual areas:

1. Scalar and Pointer Types: Defines basic data types and memory addressing models.
2. Tensors: Provides abstractions for different memory hierarchies, including registers, shared memory, T-memory, and global memory.
3. Layouts: Defines how data is organized in memory (e.g., register, shared, or global layouts) to facilitate efficient tile-level access.

In standard GPU kernels, each output tile of a result matrix (C) is typically computed by a single thread block iterating over the entire K dimension. This is efficient for large M and N dimensions.

Split-K is an optimization for workloads where M and N are small (not enough output tiles to saturate all SMs) but K is large. It partitions the K dimension into split_k_factor segments, assigning each segment to a separate thread block. These blocks compute partial results which are then aggregated in-place using semaphore-based synchronization.

A GlobalTensor is a tensor stored in the GPU's global memory. It is characterized by three main attributes:

• dtype: The scalar data type of the elements.
• shape: A tuple of integers representing dimension sizes. These can be constants or grid-invariant expressions (e.g., kernel parameters).
• layout: The mapping that defines how multi-dimensional indices correspond to linear global memory addresses.

A RegisterLayout (class tilus.ir.RegisterLayout) defines how elements of a register tensor are distributed across the local registers of all threads within a thread block.

Unlike global or shared memory layouts which are distributed layouts (defining how to find an element's position in memory), a register layout is a mapping that specifies:

1. Which thread(s) are storing the element (thread_id).
2. The position of the element within the local register memory of those threads (local_id).

Formally, it maps a (thread_id, local_id) pair to a logical tensor index, or vice versa.

An mbarrier (memory barrier) is a 64-bit synchronization object residing in shared memory. It is used on Hopper and Blackwell GPUs to track the completion of both thread arrivals and asynchronous hardware operations (like TMA data transfers).

A phase in an mbarrier completes only when two conditions are met:

1. pending arrivals = 0: All participating threads have arrived.
2. pending tx-count = 0: All tracked asynchronous transactions (in bytes) have completed.

When both conditions are met, the hardware automatically flips the phase bit, resets the pending arrivals to the expected arrival count, resets the pending tx-count to 0, and wakes any waiting threads.

A RegisterTensor is a tensor stored directly in the GPU registers. Unlike global or shared memory tensors, a register tensor is distributed among the threads within a thread block.

Key properties of a RegisterTensor include:

• dtype: The scalar data type of the elements (e.g., float32).
• shape: A tuple of integers representing the dimensions of the tensor.
• layout (optional): A RegisterLayout object that defines how elements are distributed across threads. The layout significantly impacts which operations are available and the overall kernel performance.

Global layout determines how Tilus calculates the address of an element within a global tensor. When creating a global tensor via tilus.Script.global_view (from a pointer) or tilus.Script.global_tensor (allocation), you can control the memory organization using the layout or strides arguments.

There are three ways to define the layout:

1. Default: If neither layout nor strides is provided, Tilus uses a row-major compact layout.
2. Strides: Providing a strides argument tells Tilus to construct a layout based on those specific strides.
3. Explicit Layout: Providing a layout argument tells Tilus to use that specific layout object directly.

Regardless of the method used, Tilus represents the resulting configuration as a tilus.ir.GlobalLayout object.