Triton Language and Compiler

Gluon provides two primary layers of APIs for kernel development:

1. Runtime API: Handles the execution lifecycle, including JIT decorators, result handling, and host-side tensor descriptors required to launch Gluon kernels.
2. Language API: Defines the core programming model, including memory operations, math operations, layouts, and compile-time helpers.

ConSan is a tool that instruments Triton IR with runtime checks to detect illegal concurrent access to shared memory and tensor memory. It tracks read and write frontiers over buffer state lanes and models synchronization for various operations, including mbarrier synchronization and commit-count synchronization for asynchronous operations like cp.async, WGMMA, TMA stores, and AMD TDM copies.

ConSan is target-hook based:

• cuda:* uses NVIDIA hooks.
• hip:* uses AMD hooks.

It uses BufferRegion analysis to collect exact shared-memory and tensor-memory address sets and barrier allocations, creating auxiliary state in distributed tensors and shared-cluster global scratch memory.

Triton provides several Python modules for kernel development, testing, and semantics:

• triton: The core entry point.
• triton.language: The primary module for writing Triton kernels.
• triton.testing: Utilities for testing Triton code.
• triton.language.extra.cuda: CUDA-specific extensions for the language module.
• Triton Semantics: Documentation regarding the formal behavior and semantics of the language.

f2reduce is a lightweight library for converting a binary matrix to row reduced echelon form (RREF) using Gaussian elimination over GF(2). It is optimized for autovectorization (GCC/LLVM) and high instruction-level parallelism, and it can utilize AVX512's vpternlogq instruction if available.

To use the library, you can call the primary RREF function or use the utility function to determine an optimal stride for your matrix width.

// Primary function to perform in-place RREF
void inplace_rref_strided(uint64_t *matrix, uint64_t rows, uint64_t cols, uint64_t stride);

// Utility to get a recommended stride based on column count
uint64_t get_recommended_stride(uint64_t cols);

Triton-shared is a project aimed at lowering Triton IR to core MLIR dialects (such as linalg and memref) to facilitate running Triton on CPUs and improving compiler modularity.

Current Capabilities:

• Modular Compiler Passes: Decoupled data extraction from lowering via paths like Triton-to-structured $\rightarrow$ triton-arith-to-linalg $\rightarrow$ Structured-to-memref.
• Improved Pointer Analysis: Supports nested loops and non-contiguous memory access.
• Unstructured Access: Supports lowering unstructured access using a single base pointer.
• Op Support: Lowers Triton operations like split, join, and cat to MLIR/linalg.

Roadmap Goals:

• Complete support for non-contiguous pointers.
• Detection of complex memory access patterns (e.g., row-gather/scatter sequences).
• Extension to control flow operations.

Triton is designed for high-performance GPU computing with the following support:

Supported Platforms

• Linux

Supported Hardware

• NVIDIA GPUs: Compute Capability 8.0+
• AMD GPUs: ROCm 6.2+
• CPUs: Under development

The triton Python package provides several key modules for writing and optimizing GPU kernels:

• triton.jit: The primary decorator used to compile Python functions into Triton kernels.
• triton.autotune: A module used to automatically find the best configuration (e.g., block sizes) for a kernel by running it with multiple parameter sets.
• triton.heuristics: Provides mechanisms for making compile-time decisions based on input parameters.
• triton.Config: A class used to define specific configurations (like BLOCK_SIZE) for use with the autotuner.

The triton.testing module provides utilities for benchmarking Triton kernels, comparing numerical results, and generating performance reports. It is primarily used to validate the correctness of kernels via assert_close and to measure execution performance using various benchmarking functions like do_bench and do_bench_proton.

import triton.testing as tt

# Example usage pattern for benchmarking
# tt.do_bench(lambda: my_kernel(...))

# Example usage pattern for numerical verification
# tt.assert_close(actual, expected)

Gluon is Triton's lower-level GPU programming model. It is designed for advanced kernel development where developers need to trade the convenience of high-level Triton abstractions for direct control over hardware-specific features.

Use Gluon when you need to explicitly manage:

• Layouts: Controlling how data is organized in memory.
• Shared Memory: Manual management of on-chip memory.
• Warp Specialization: Assigning specific tasks to different warps within a thread block.
• Target-specific features: Leveraging unique hardware capabilities of specific GPU architectures.