### One-time Setup for Development Source: https://github.com/openai/triton/blob/main/docs/getting-started/installation.rst Perform a one-time setup for development environment, including installing necessary tools and dependencies. ```bash make dev-install ``` -------------------------------- ### Build LLVM from Source (Example) Source: https://github.com/openai/triton/blob/main/README.md Example commands for building LLVM from source. This process can take a significant amount of time. ```shell $ cd $HOME/llvm-project # your clone of LLVM. $ mkdir build $ cd build $ cmake -G Ninja -DCMAKE_BUILD_TYPE=Release -DLLVM_ENABLE_ASSERTIONS=ON ../llvm -DLLVM_ENABLE_PROJECTS="mlir;llvm;lld;clang" -DLLVM_TARGETS_TO_BUILD="host;NVPTX;AMDGPU" $ ninja ``` -------------------------------- ### Install Triton from Source Source: https://github.com/openai/triton/blob/main/README.md Clones the Triton repository and installs it from source, including build-time dependencies. ```shell git clone https://github.com/triton-lang/triton.git cd triton pip install -r python/requirements.txt # build-time dependencies pip install -e . ``` -------------------------------- ### Install Proton Source: https://github.com/openai/triton/blob/main/third_party/proton/README.md Clone the Triton repository and install Proton from the python directory. ```bash git clone https://github.com/triton-lang/triton cd triton/python pip install . ``` -------------------------------- ### Profiler Setup and Callback Handling Source: https://github.com/openai/triton/blob/main/third_party/proton/README.md Implementation of profiler setup and callback handling, utilizing an externally linked profiling API. ```cpp void MyProfilerImplementation::MyProfilerImplementationPimpl::doStart() { // Any setup required for profiling, make use of MyProfilingApi which was linked to // the libproton build using the cmake add_proton_backend_external_lib function earlier. MyProfilingApi::registerCallback(callback_function) } void MyProfilerImplementation::MyProfilerImplementationPimpl::callbackFunction(...) { // Do correlation and record callback/activity in a similar manor to Cupti/Roctracer // assuming the profiling API has a similar design. If not introduce your own abstractions // to handle correlation and callbacks. const auto kernelMetric = ... entry.upsertMetric(std::move(kernelMetric)) } ``` -------------------------------- ### Install Triton Tutorial Dependencies Source: https://github.com/openai/triton/blob/main/python/tutorials/README.rst Installs the required dependencies for running Triton tutorials. Navigate to the triton directory and execute this command. ```bash cd triton pip install -r python/tutorials/requirements.txt ``` -------------------------------- ### Build and Install Triton with Custom LLVM Source: https://github.com/openai/triton/blob/main/README.md Builds LLVM and installs Triton using the custom LLVM build. This is a convenience command. ```shell make dev-install-llvm ``` -------------------------------- ### Install Triton with Custom LLVM Build Source: https://github.com/openai/triton/blob/main/README.md Installs Triton from source, pointing to a pre-built custom LLVM. Ensure LLVM_BUILD_DIR is set correctly. ```shell # Modify as appropriate to point to your LLVM build. $ export LLVM_BUILD_DIR=$HOME/llvm-project/build $ cd $ LLVM_INCLUDE_DIRS=$LLVM_BUILD_DIR/include \ LLVM_LIBRARY_DIR=$LLVM_BUILD_DIR/lib \ LLVM_SYSPATH=$LLVM_BUILD_DIR \ pip install -e . ``` -------------------------------- ### Install Triton from Source with Virtualenv Source: https://github.com/openai/triton/blob/main/README.md Clones the Triton repository and installs it from source using a virtual environment, including build-time dependencies. ```shell git clone https://github.com/triton-lang/triton.git cd triton python -m venv .venv --prompt triton source .venv/bin/activate pip install -r python/requirements.txt # build-time dependencies pip install -e . ``` -------------------------------- ### Install Triton Library Source: https://github.com/openai/triton/blob/main/CMakeLists.txt Installs the Triton library target to the appropriate destination directory based on CMAKE_INSTALL_LIBDIR. ```cmake install(TARGETS triton COMPONENT libraries LIBRARY DESTINATION ${CMAKE_INSTALL_LIBDIR} ARCHIVE DESTINATION ${CMAKE_INSTALL_LIBDIR}) ``` -------------------------------- ### Install and Use Hatchet for Profile Visualization Source: https://github.com/openai/triton/blob/main/third_party/proton/README.md Instructions for installing the Hatchet library and using the `proton-viewer` command to visualize profile data on the terminal. This is used for aggregated profile data. ```bash pip install llnl-hatchet proton-viewer -m time/s ``` -------------------------------- ### Registering a Third-Party Backend with Proton Source: https://github.com/openai/triton/blob/main/third_party/proton/README.md Example CMakeLists.txt to register a custom backend, device type, and external library with Proton. ```cmake # Register my backend with proton so calls to it's hook can be generated. # This also adds the MyBackend target to the libproton build. add_proton_backend(MyBackend MyBackendNamespace MyBackendSourceFile.cpp ) # Register an extension to the DeviceType enum if your backend introduces a new device. add_proton_device_type(MY_DEVICE) # Link an external library into the build, for example a custom profiling api. add_proton_backend_external_lib(MyProfilingApi) # Add any includes/sources/dependencies to MyBackend target using standard CMake practices. ... ``` -------------------------------- ### Install Triton Python Package from Source Source: https://github.com/openai/triton/blob/main/docs/getting-started/installation.rst Install the Triton Python package in editable mode from the cloned source code. This command builds and installs the package. ```bash pip install -e . ``` -------------------------------- ### Start Instrumentation Profiling (Default Mode) Source: https://github.com/openai/triton/blob/main/third_party/proton/README.md Initiates instrumentation profiling with a specified name and backend. If no mode is specified, it defaults to profiling kernel cycles. ```python import triton.profiler as proton proton.start( name="profile_name", backend="instrumentation", mode="=:=:..." ) ``` -------------------------------- ### Running with Plugin Loaded (No Pass Insertion) Source: https://github.com/openai/triton/blob/main/examples/plugins/README.md Demonstrates running the toy kernel example after loading a plugin library. The output shows that loading the plugin alone does not alter the compiler pipeline's behavior. ```bash TRITON_PLUGIN_PATHS=/home/triton/python/triton/plugins/libTritonPluginsTestLib.so python test.py ``` -------------------------------- ### Install Triton using Pip Source: https://github.com/openai/triton/blob/main/docs/getting-started/installation.rst Install the latest stable release of Triton using pip. Binary wheels are available for specific Python versions. ```bash pip install triton ``` -------------------------------- ### Start Proton Profiler with Trace Data Source: https://github.com/openai/triton/blob/main/third_party/proton/README.md This Python snippet demonstrates how to start the Proton profiler with the `data` option set to `trace`. This ensures the entire trace is dumped, not just aggregated data, for detailed analysis. ```python import triton.profiler as proton proton.start(name="profile_name", data="trace") ``` -------------------------------- ### Toy Kernel Example Source: https://github.com/openai/triton/blob/main/examples/plugins/README.md A basic Triton kernel example used to demonstrate compiler pipeline modifications. It includes necessary imports and a simple kernel definition. ```python import torch import os import triton import triton.language as tl from triton._C.libtriton import ir, passes from triton import knobs DEVICE = triton.runtime.driver.active.get_active_torch_device() @triton.jit def kernel(BLOCK_SIZE: tl.constexpr): return if __name__ == '__main__': size = 98432 x = torch.rand(size, device=DEVICE) output = torch.empty_like(x) n_elements = output.numel() grid = lambda meta: (triton.cdiv(n_elements, meta['BLOCK_SIZE']), ) h = kernel[grid](BLOCK_SIZE=1024) print(h.asm["ttgir"]) ``` -------------------------------- ### Start Profiling with Launch Hook Source: https://github.com/openai/triton/blob/main/third_party/proton/README.md Enables Proton to invoke a launch_metadata function before launching the GPU kernel. The metadata_fn provides kernel launch information. ```python import triton.profiler as proton from typing import NamedTuple # hook: When hook="triton", it enables proton to invoke launch_metadata function before launching the GPU kernel proton.start("profile_name", hook="triton") def metadata_fn( grid: tuple, metadata: NamedTuple, args: dict ): return {"name": "", "flops8": 1.0} @triton.jit(launch_metadata=metadata_fn) def foo(x, y): tl.store(y, tl.load(x)) ``` -------------------------------- ### Start Profiling Session Source: https://github.com/openai/triton/blob/main/third_party/proton/README.md Starts a profiling session with a given name. This is an experimental API for retrieving profile data in memory without dumping to files. ```python import triton.profiler as proton session_id = proton.start(name="profile_name") ``` -------------------------------- ### Start Instrumentation Profiling (Default Mode Object) Source: https://github.com/openai/triton/blob/main/third_party/proton/README.md Initiates instrumentation profiling using a mode object for collecting metrics from every warp. ```python import triton.profiler.mode as pmode proton.start( name="profile_name", backend="instrumentation", mode=pmode.Default() # collect metrics from every warp ) ``` -------------------------------- ### Plugin Interface and Custom Dialects Source: https://github.com/openai/triton/blob/main/docs/meetups/01-06-2026/notes.md Details on the plugin interface, which uses PyBind names, and how to register custom dialects for new operations. Examples are available in the demo. ```text Uses PyBind names. See demo for how to register plugin with API See demo for creating a transformation pass (e.g. like invoking pass after a loop unrolled.) New operations! Registered similarly to passes. ``` -------------------------------- ### Slack Reminder Message Example Source: https://github.com/openai/triton/blob/main/docs/meetups/for_moderators/README.md Example of a reminder message to be posted in Slack for the community meetup. Includes a placeholder for the agenda. ```text Reminder, this month's community meetup is tomorrow at 10am PST. Agenda: Topic #1 Topic #2 ``` -------------------------------- ### Kernel Example with Plugin and Hook Integration Source: https://github.com/openai/triton/blob/main/examples/plugins/README.md An extended Triton kernel example that loads a plugin and integrates a custom hook to modify the compiler pipeline. This version includes the necessary functions for plugin interaction and hook registration. ```python import torch import os import triton import triton.language as tl from triton._C.libtriton import ir, passes from triton import knobs import pathlib import hashlib DEVICE = triton.runtime.driver.active.get_active_torch_device() @triton.jit def kernel(BLOCK_SIZE: tl.constexpr): return #These two methods must be implemented by the plugin def get_key(): return pathlib.Path(__file__).read_text() def get_hash(): return hashlib.sha256(get_key().encode('utf-8')).hexdigest() def inspect_stages_hook(self=None, stages=None, options=None, language=None, capability=None): # If the hook is called with no arguments we assume were just after the key and hash and don't want to # actually execute the pipeline yet. # This no argument early return must be implemented. if all(arg is None for arg in (stages, options, language, capability)): return get_key(), get_hash() def make_ttir_wrapper(mod, metadata, opt, capability): mod = self.make_ttir(mod, metadata, opt, capability) pm = ir.pass_manager(mod.context) pm.enable_debug() passes.plugin.add_plugin(pm) pm.run(mod, 'make_ttir_plugin') return mod stages["ttir"] = lambda src, metadata: make_ttir_wrapper(src, metadata, options, capability) return get_key(), get_hash() if __name__ == '__main__': size = 98432 x = torch.rand(size, device=DEVICE) output = torch.empty_like(x) n_elements = output.numel() grid = lambda meta: (triton.cdiv(n_elements, meta['BLOCK_SIZE']), ) h = kernel[grid](BLOCK_SIZE=1024) print(h.asm["ttgir"]) if "TRITON_PLUGIN_PATHS" in os.environ: knobs.runtime.add_stages_inspection_hook = inspect_stages_hook h = kernel[grid](BLOCK_SIZE=1024) print(h.asm["ttgir"]) # Unset the hook to go back to the standard pipeline knobs.runtime.add_stages_inspection_hook = None h = kernel[grid](BLOCK_SIZE=1024) print(h.asm["ttgir"]) ``` -------------------------------- ### Kernel-Side Instrumentation with Gluon Source: https://github.com/openai/triton/blob/main/third_party/proton/README.md Example of instrumenting a kernel written in Triton DSL using Gluon. Requires enabling semantic support for Triton. ```python from triton.experimental import gluon from triton.experimental.gluon import language as gl import triton.profiler.language as pl @gluon.jit def kernel(...): pl.enter_scope("scope0") for i in range(iters): gl.load(...) pl.exit_scope("scope0") with pl.scope("scope1"): for i in range(iters): gl.load(...) ``` -------------------------------- ### Develop and Run a Custom Pass with triton-opt Source: https://github.com/openai/triton/blob/main/examples/plugins/README.md This example demonstrates how to enable Triton plugins and run a custom pass using triton-opt. It requires setting the TRITON_EXT_ENABLED environment variable and specifying the path to the plugin library. ```bash export TRITON_EXT_ENABLED=1; make dev-install-llvm TRITON_PLUGIN_PATHS=/home/triton/python/triton/plugins/libTritonPluginsTestLib.so triton-opt -tritongpu-plugin test/Plugins/test-plugin.mlir ``` -------------------------------- ### Install Python Package Dependencies Source: https://github.com/openai/triton/blob/main/docs/getting-started/installation.rst Install build-time dependencies for the Python package from source. This ensures all necessary libraries are available before building. ```bash pip install -r python/requirements.txt ``` -------------------------------- ### Clone Triton Repository Source: https://github.com/openai/triton/blob/main/docs/getting-started/installation.rst Clone the Triton repository from GitHub to install from source. This is the first step in building Triton locally. ```bash git clone https://github.com/triton-lang/triton.git cd triton ``` -------------------------------- ### Basic State Annotation Example Source: https://github.com/openai/triton/blob/main/third_party/proton/README.md Demonstrates the basic usage of `proton.scope` and `proton.state` to customize the call path of GPU operations. The inner-most state overwrites outer states. ```python with proton.scope("test"): with proton.state("state0"): with proton.scope("test0"): foo0[1,](x, y) with proton.scope("test1"): foo1[1,](x, y) ``` -------------------------------- ### Start CUDA Graph Profiling Source: https://github.com/openai/triton/blob/main/third_party/proton/README.md Initializes Proton for profiling CUDA graph launched kernels. It captures and concatenates call paths and supports flexible metric aggregation. ```python import triton.profiler as proton proton.start(name="profile_name", context="shadow") # Capture the CUDA graph graph = torch.cuda.CUDAGraph() with torch.cuda.graph(graph): with proton.scope("graph"): ... proton.deactivate() ``` -------------------------------- ### Speed up pip installs with no-build-isolation Source: https://github.com/openai/triton/blob/main/README.md Pass '--no-build-isolation' to 'pip install' to accelerate no-builds by preventing repeated cmake symlink usage and ninja rebuilds. ```shell pip install -e . --no-build-isolation ``` -------------------------------- ### Find and Configure MLIR and LLD Source: https://github.com/openai/triton/blob/main/CMakeLists.txt Locates MLIR and LLD installations using CMake's find_package command. It also appends necessary directories to CMAKE_MODULE_PATH and includes MLIR-specific CMake modules. ```cmake if(NOT MLIR_DIR) set(MLIR_DIR ${LLVM_LIBRARY_DIR}/cmake/mlir) endif() if(NOT LLD_DIR) set(LLD_DIR ${LLVM_LIBRARY_DIR}/cmake/lld) endif() # MLIR find_package(MLIR REQUIRED CONFIG PATHS ${MLIR_DIR}) list(APPEND CMAKE_MODULE_PATH "${MLIR_CMAKE_DIR}") list(APPEND CMAKE_MODULE_PATH "${LLVM_CMAKE_DIR}") include(TableGen) # required by AddMLIR include(AddLLVM) include(AddMLIR) ``` -------------------------------- ### Command-line Profiling with Proton Source: https://github.com/openai/triton/blob/main/third_party/proton/README.md Examples of using Proton as a command-line tool to profile Python scripts and Pytest tests. The `proton` command or `python -m triton.profiler.proton` can be used with various options. ```bash proton [options] script.py [script_args] [script_options] proton [options] pytest [pytest_args] [script_options] python -m triton.profiler.proton [options] script.py [script_args] [script_options] proton --instrument=[instrumentation pass] script.py ``` -------------------------------- ### Get Recommended Stride for Matrix Source: https://github.com/openai/triton/blob/main/third_party/f2reduce/README.md Returns a recommended stride value for a binary matrix based on its number of columns. This helps optimize memory access patterns for improved performance. ```c++ uint64_t get_recommended_stride(uint64_t cols); ``` -------------------------------- ### Basic C Loop for Polyhedral Analysis Source: https://github.com/openai/triton/blob/main/docs/programming-guide/chapter-2/related-work.rst A simple C nested loop structure used to illustrate the concept of iteration domains in polyhedral compilation. This example helps define the bounds and structure for static analysis. ```C for(int i = 0; i < 3; i++) for(int j = i; j < 5; j++) A[i][j] = 0; ``` -------------------------------- ### Install Triton without Proton Source: https://github.com/openai/triton/blob/main/third_party/proton/README.md Set the TRITON_BUILD_PROTON environment variable to OFF to prevent Proton from building during installation. ```bash TRITON_BUILD_PROTON=OFF pip install . ``` -------------------------------- ### Visualize Sorted Profile Data with Proton-Viewer Source: https://github.com/openai/triton/blob/main/third_party/proton/README.md Command to visualize sorted profile data using `proton-viewer`, specifying metrics like time in nanoseconds and percentage. This helps in identifying performance bottlenecks. ```bash proton-viewer -m time/ns,time/% --print-sorted ``` -------------------------------- ### View Proton-Viewer Help Source: https://github.com/openai/triton/blob/main/third_party/proton/README.md Command to display the help message for `proton-viewer`, listing all available options and commands for profile data visualization. ```bash proton-viewer -h ``` -------------------------------- ### C Loop Example for Matrix Multiplication Source: https://github.com/openai/triton/blob/main/docs/programming-guide/chapter-2/related-work.rst This C code snippet demonstrates a nested loop structure for matrix multiplication, where the inner loop's bound depends on an external array K. This is an example of a computation with irregular iteration spaces. ```C for(int i = 0; i < 4; i++) for(int j = 0; j < 4; j++) float acc = 0; for(int k = 0; k < K[i]; k++) acc += A[i][col[i, k]] * B[k][j] C[i][j] = acc; ``` -------------------------------- ### Configure FileCheck Utility Source: https://github.com/openai/triton/blob/main/CMakeLists.txt Copies the FileCheck utility from the LLVM syspath to the Triton wheel directory, ensuring it's available for wheel building. ```cmake configure_file( "${LLVM_SYSPATH}/bin/FileCheck" "${TRITON_WHEEL_DIR}/FileCheck" COPYONLY) ``` -------------------------------- ### Visualize CPU and GPU Time with Proton-Viewer Source: https://github.com/openai/triton/blob/main/third_party/proton/README.md Command to visualize both CPU time (`cpu_time`) and accelerator time (`time`) using `proton-viewer`. This allows for a comprehensive analysis of performance across different execution environments. ```bash proton-viewer -m time/ns,cpu_time/ns ``` -------------------------------- ### Get CMake Build Directory Source: https://github.com/openai/triton/blob/main/AGENTS.md This command determines the CMake build directory for Triton. It's used to set up the build environment. ```bash BUILD_DIR := $(shell PYTHONPATH="./python" python3 -c 'from build_helpers import get_cmake_dir; print(get_cmake_dir())') ``` -------------------------------- ### Instruction Sampling with CUPTI Backend Source: https://github.com/openai/triton/blob/main/third_party/proton/README.md Enable instruction sampling on NVIDIA GPUs using the 'cupti' backend and 'pcsampling' mode. This incurs significant CPU-side overhead. ```python import triton.profiler as proton proton.start(name="profile_name", context="shadow", backend="cupti", mode="pcsampling") ``` -------------------------------- ### Configure LIT Site Configuration File Source: https://github.com/openai/triton/blob/main/test/CMakeLists.txt Generates the `lit.site.cfg.py` file, which is essential for configuring the LIT testing framework. It takes an input template and an output path, along with specific configuration options. ```cmake configure_lit_site_cfg( ${CMAKE_CURRENT_SOURCE_DIR}/lit.site.cfg.py.in ${CMAKE_CURRENT_BINARY_DIR}/lit.site.cfg.py MAIN_CONFIG ${CMAKE_CURRENT_BINARY_DIR}/lit.site.cfg.py ) ``` -------------------------------- ### MLIR Output with Pass Modification Source: https://github.com/openai/triton/blob/main/examples/plugins/README.md The MLIR output after the custom pass has been applied via the hook. This demonstrates that the pass successfully modified the kernel, for example, by changing its name. ```MLIR module attributes {"ttg.num-ctas" = 1 : i32, "ttg.num-warps" = 4 : i32, ttg.target = "cuda:90", "ttg.threads-per-warp" = 32 : i32} { tt.func public @kernel() attributes {noinline = false} { tt.return loc(#loc1) } loc(#loc) } loc(#loc) #loc = loc("/home/triton/test.py":13:0) #loc1 = loc("/home/triton/test.py":14:4) module attributes {"ttg.num-ctas" = 1 : i32, "ttg.num-warps" = 4 : i32, ttg.target = "cuda:90", "ttg.threads-per-warp" = 32 : i32} { tt.func public @foo() attributes {noinline = false} { tt.return loc(#loc1) } loc(#loc) } loc(#loc) #loc = loc("/home/triton/test.py":13:0) #loc1 = loc("/home/triton/test.py":14:4) module attributes {"ttg.num-ctas" = 1 : i32, "ttg.num-warps" = 4 : i32, ttg.target = "cuda:90", "ttg.threads-per-warp" = 32 : i32} { tt.func public @kernel() attributes {noinline = false} { tt.return loc(#loc1) } loc(#loc) ``` -------------------------------- ### Supported Metadata Keys for Launch Hook Source: https://github.com/openai/triton/blob/main/third_party/proton/README.md Lists the keys that can be returned by the metadata_fn function to provide metadata for the GPU kernel. These include kernel name and various floating-point operation counts. ```python name: str # The name of the kernel flops8: float # The number of 8-bit floating-point operations flops16: float # The number of 16-bit floating-point operations flops32: float # The number of 32-bit floating-point operations flops64: float # The number of 64-bit floating-point operations bytes: int # The number of bytes expected to be transferred ``` -------------------------------- ### Configuring GSan Include Flags Source: https://github.com/openai/triton/blob/main/third_party/nvidia/CMakeLists.txt Sets up include paths for the GSan build process by iterating through CUDA runtime paths. Empty paths are ignored. ```cmake set(TRITON_GSAN_INCLUDE_FLAGS "") foreach(include_path IN ITEMS "${TRITON_CUDACRT_PATH}" "${TRITON_CUDART_PATH}") if(NOT "${include_path}" STREQUAL "") list(APPEND TRITON_GSAN_INCLUDE_FLAGS "-I${include_path}") endif() endforeach() ``` -------------------------------- ### Run Lit Tests Source: https://github.com/openai/triton/blob/main/AGENTS.md This command executes lit tests from the build directory. It's used for testing MLIR components. ```bash cd BUILD_DIR; ninja triton-opt; lit -v test/.mlir ``` -------------------------------- ### Create Testing Backend Library Source: https://github.com/openai/triton/blob/main/third_party/proton/test/unittest/Backend/CMakeLists.txt Defines a CMake library for the testing backend, specifying source files and include directories. ```cmake add_library(TestBackend OBJECT TestBackend.cpp) target_include_directories(TestBackend PRIVATE "${PROTON_COMMON_DIR}/include" "${PROTON_SRC_DIR}/include" ) ``` -------------------------------- ### Include Directories for DialectPlugin Source: https://github.com/openai/triton/blob/main/examples/plugins/DialectPlugins/DialectPlugin/include/DialectPlugin/CMakeLists.txt Sets the include directories for the DialectPlugin, essential for the build system to locate header files. ```cmake include_directories(${PROJECT_SOURCE_DIR}/examples/plugins/DialectPlugins/DialectPlugin/include) include_directories(${PROJECT_BINARY_DIR}/examples/plugins/DialectPlugins/DialectPlugin/include) ``` -------------------------------- ### Run All Triton Tests (GPU Required) Source: https://github.com/openai/triton/blob/main/docs/getting-started/installation.rst Execute all available tests for Triton. This command requires a GPU to be present and configured. ```bash make test ``` -------------------------------- ### Run Proton DSL Profiler (Timeline Trace Mode) Source: https://github.com/openai/triton/blob/main/third_party/proton/tutorials/intra_kernel/README.md Executes the Proton DSL profiler in its default timeline trace mode. This is used for native integration with Triton and Gluon DSL kernels. ```bash python3 example_dsl.py ``` -------------------------------- ### Insert Pass into Compiler Pipeline Source: https://github.com/openai/triton/blob/main/examples/plugins/README.md This example shows how to define a hook to insert a custom pass at the end of the 'make_ttir' pipeline. The hook's placement is arbitrary within the Triton pipeline. This functionality can be toggled by commenting out the assignment or setting it to None, avoiding core compiler changes. ```python # loc(#loc) #loc = loc("/home/triton/test.py":13:0) #loc1 = loc("/home/triton/test.py":14:4) knobs.runtime.add_stages_inspection_hook = inspect_stages_hook ``` -------------------------------- ### Custom Out-of-Tree Targets Source: https://github.com/openai/triton/blob/main/docs/meetups/01-06-2026/notes.md Work in progress to enable loading backends from external shared objects at runtime, rather than requiring static linking with Triton at compile time. ```text Alternate path for loading backends (backend is something like AMD, Nvidia, etc.) Current backends are loaded with macros. New way, load from external shared object, don’t need to statically link in with Triton at compile time. ``` -------------------------------- ### Set Triton build with ccache Source: https://github.com/openai/triton/blob/main/README.md Enable ccache for builds by setting the TRITON_BUILD_WITH_CCACHE environment variable. ```shell export TRITON_BUILD_WITH_CCACHE=true ``` -------------------------------- ### Reproduce Compiler Crashes with MLIR Source: https://github.com/openai/triton/blob/main/AGENTS.md This command reproduces compiler crashes locally by running an MLIR reproducer. Save the MLIR output to a file and then execute this command. ```bash triton-opt /tmp/.mlir --run-reproducer ``` -------------------------------- ### Run Proton DSL Profiler (Operation Measurement Mode) Source: https://github.com/openai/triton/blob/main/third_party/proton/tutorials/intra_kernel/README.md Executes the Proton DSL profiler in operation measurement mode. This provides cycle-accurate latency measurements for individual operations. ```bash python3 example_dsl.py --op-measure ``` -------------------------------- ### Create Testing Backend Registry Library Source: https://github.com/openai/triton/blob/main/third_party/proton/test/unittest/Backend/CMakeLists.txt Defines a CMake library for the testing backend registry, including the generated registration code and necessary include directories. ```cmake add_library(TestBackendRegistry OBJECT "${CMAKE_CURRENT_BINARY_DIR}/RegisteredBackends.cpp" ) target_include_directories(TestBackendRegistry PRIVATE "${PROTON_COMMON_DIR}/include" "${PROTON_SRC_DIR}/include" ) ``` -------------------------------- ### Set Triton build with clang and lld Source: https://github.com/openai/triton/blob/main/README.md Use clang and lld for faster builds by setting the TRITON_BUILD_WITH_CLANG_LLD environment variable. ```shell export TRITON_BUILD_WITH_CLANG_LLD=true ``` -------------------------------- ### Minimal Dummy Profiler Implementation Source: https://github.com/openai/triton/blob/main/third_party/proton/README.md A minimal dummy implementation of a profiler, inheriting from a base class. Refer to existing implementations for detailed requirements. ```cpp class MyProfilerImplementation : public GPUProfiler { // Similar to CuptiProfiler or RoctracerProfiler } ``` -------------------------------- ### Run Proton DSL Profiler with High Accuracy Source: https://github.com/openai/triton/blob/main/third_party/proton/tutorials/intra_kernel/README.md Executes the Proton DSL profiler with settings for increased accuracy. This may involve higher overhead but provides more precise measurements. ```bash python3 example_dsl.py --increase-accuracy ``` -------------------------------- ### Building GSan Runtime IR with Custom Command Source: https://github.com/openai/triton/blob/main/third_party/nvidia/CMakeLists.txt A custom CMake command to build the GSan runtime IR. It ensures the output directory exists and then uses clang++ to compile the CUDA source to LLVM IR, with specific flags for a freestanding device runtime. ```cmake add_custom_command( OUTPUT "${GSAN_RUNTIME_IR}" COMMAND "${CMAKE_COMMAND}" -E make_directory "${PROJECT_SOURCE_DIR}/third_party/nvidia/backend/lib" COMMAND "${TRITON_GSAN_CLANGXX}" -x cuda -std=c++17 -O3 -S -emit-llvm --cuda-device-only -nocudainc -nocudalib -nostdinc -fno-exceptions -fcuda-flush-denormals-to-zero --cuda-gpu-arch=sm_80 --cuda-feature='+ptx70' ${TRITON_GSAN_INCLUDE_FLAGS} "${GSAN_RUNTIME_SRC}" -o "${GSAN_RUNTIME_IR}" DEPENDS "${GSAN_RUNTIME_SRC}" ${GSAN_RUNTIME_HDRS} COMMENT "Building GSan runtime" VERBATIM) add_custom_target(TritonNVIDIAGSanRuntime ALL DEPENDS "${GSAN_RUNTIME_IR}") ``` -------------------------------- ### Profile a Python Region Source: https://github.com/openai/triton/blob/main/third_party/proton/README.md Use proton.start to begin profiling a region, proton.deactivate to skip, proton.activate to resume, and proton.finalize to write profile data. ```python session_id = proton.start(name="profile_name", context="python") ... # Skip a region proton.deactivate(session_id) ... # Restart profiling proton.activate(session_id) ... # Write out the profile data and finalize the profiler proton.finalize() ``` -------------------------------- ### Run Proton DSL Profiler with Warp Sampling Source: https://github.com/openai/triton/blob/main/third_party/proton/tutorials/intra_kernel/README.md Executes the Proton DSL profiler with warp sampling enabled, focusing on specific warps and using global memory for buffers. This is useful for reducing overhead while profiling. ```bash python3 example_dsl.py --warp-sampling --warp-ids "0,1,2,3" --gmem_buffer ``` -------------------------------- ### mbarrier Source: https://github.com/openai/triton/blob/main/docs/gluon/api/nvidia.hopper.rst Manages memory barriers for synchronization on NVIDIA Hopper. ```APIDOC ## mbarrier ### Description Manages memory barriers for synchronization. ### Method Not specified (likely a function call within the Triton framework). ### Endpoint Not applicable (SDK function). ### Parameters Not specified in the provided text. ### Response Not specified in the provided text. ``` -------------------------------- ### Profile with Shadow Context and Scopes Source: https://github.com/openai/triton/blob/main/third_party/proton/README.md Utilize the 'shadow' context and proton.scope for annotated regions. Nested scopes are supported. Ensure proton.finalize is called. ```python import triton.profiler as proton session_id = proton.start(name="profile_name", context="shadow") with proton.scope("test0"): with proton.scope("test1"): foo[1,](x, y) with proton.scope("test2"): foo[1,](x, y) ... proton.finalize() ``` -------------------------------- ### Inline Assembly Source: https://github.com/openai/triton/blob/main/docs/python-api/triton.language.rst Allows embedding inline assembly code. ```APIDOC ## Inline Assembly Enables the insertion of inline assembly code for element-wise operations. ### Functions - **inline_asm_elementwise**: Embeds inline assembly for element-wise computations. ``` -------------------------------- ### mbarrier Source: https://github.com/openai/triton/blob/main/docs/gluon/api/nvidia.blackwell.rst Manages memory barriers for synchronization on Blackwell hardware. ```APIDOC ## mbarrier ### Description Manages memory barriers for synchronization on NVIDIA Blackwell hardware. This is crucial for ensuring correct data visibility across asynchronous operations. ### Function Signature `mbarrier()` ### Parameters None explicitly documented. ``` -------------------------------- ### Adding Backend and Plugin Subdirectories Source: https://github.com/openai/triton/blob/main/CMakeLists.txt This section iterates through specified codegen backends and conditionally adds the 'proton' subdirectory. It also manages a list of Triton plugin names, reading their names from configuration files and adding their respective subdirectories. ```cmake foreach(CODEGEN_BACKEND ${TRITON_CODEGEN_BACKENDS}) add_subdirectory(third_party/${CODEGEN_BACKEND}) endforeach() if (TRITON_BUILD_PROTON) add_subdirectory(third_party/proton) endif() # We always build proton dialect list(APPEND TRITON_PLUGIN_NAMES "proton") add_subdirectory(third_party/proton/Dialect) if (DEFINED TRITON_PLUGIN_DIRS) foreach(PLUGIN_DIR ${TRITON_PLUGIN_DIRS}) # Read the plugin name under dir/backend/name.conf cmake_path(APPEND PLUGIN_DIR "backend" "name.conf" OUTPUT_VARIABLE PLUGIN_NAME_PATH) file(READ ${PLUGIN_NAME_PATH} PLUGIN_NAME) string(STRIP ${PLUGIN_NAME} PLUGIN_NAME) list(APPEND TRITON_PLUGIN_NAMES ${PLUGIN_NAME}) # Include the plugin as part of the build, placing the build output under # ${TRITON_BINARY_DIR}/third_party/${PLUGIN_NAME} cmake_path(APPEND TRITON_BINARY_DIR "third_party" ${PLUGIN_NAME} OUTPUT_VARIABLE PLUGIN_DIR_BUILD_OUTPUT) message(STATUS "Building plugin '${PLUGIN_NAME}' from ${PLUGIN_DIR} with output ${PLUGIN_DIR_BUILD_OUTPUT}") add_subdirectory(${PLUGIN_DIR} ${PLUGIN_DIR_BUILD_OUTPUT}) endforeach() endif() ``` -------------------------------- ### Set Compile Options and Output Directory Source: https://github.com/openai/triton/blob/main/examples/plugins/DialectPlugins/DialectPlugin/lib/DialectPlugin/CMakeLists.txt Sets private compile options for visibility and configures the library output directory for the MLIR dialect plugin. ```cmake target_compile_options(MLIRDialectPlugin PRIVATE -fvisibility=hidden) if(DEFINED CMAKE_LIBRARY_OUTPUT_DIRECTORY) set_target_properties(MLIRDialectPlugin PROPERTIES LIBRARY_OUTPUT_DIRECTORY "${CMAKE_LIBRARY_OUTPUT_DIRECTORY}/../plugins") endif(DEFINED CMAKE_LIBRARY_OUTPUT_DIRECTORY) ``` -------------------------------- ### C++ Backend Registration Hook Source: https://github.com/openai/triton/blob/main/third_party/proton/README.md C++ code providing the `registerProtonBackend` hook to register profiler, device, and runtime implementations. ```cpp namespace MyBackendNamespace { proton::BackendRegistration registerProtonBackend() { return { proton::ProfilerRegistration{ "MyBackend", "TritonBackendForThisProfiler", []() -> proton::Profiler * { return &MyProfilerImplementation::instance(); }}, proton::DeviceRegistration{ "MY_DEVICE", proton::DeviceType::MY_DEVICE, [](uint64_t index) { return getMyDevice(index); }}, proton::RuntimeRegistration{ "MY_DEVICE", []() -> proton::Runtime * { return &MyRuntimeImplementaiton::instance(); }}, }; } } ``` -------------------------------- ### Run Triton Tests (No GPU) Source: https://github.com/openai/triton/blob/main/docs/getting-started/installation.rst Execute Triton tests without requiring a GPU. This is useful for testing CPU-bound functionalities or in environments without GPU access. ```bash make test-nogpu ``` -------------------------------- ### Add Dialect Plugins Subdirectory Source: https://github.com/openai/triton/blob/main/examples/plugins/CMakeLists.txt Includes the DialectPlugins subdirectory, allowing for the configuration and building of dialect-specific plugins. ```cmake # Dialect Plugins add_subdirectory(DialectPlugins) ``` -------------------------------- ### buffer_load Source: https://github.com/openai/triton/blob/main/docs/gluon/api/amd.cdna4.rst Loads data from a buffer. ```APIDOC ## buffer_load ### Description Loads data from a buffer. ### Method Not specified (likely an SDK function call) ### Endpoint Not applicable (SDK function) ### Parameters Not specified in the source. ### Request Example Not applicable. ### Response Not specified in the source. ``` -------------------------------- ### Triton Plugin Framework Overview Source: https://github.com/openai/triton/blob/main/docs/meetups/01-06-2026/notes.md The new plugin framework allows for custom extensions, dialects, and operations without needing to fork or recompile the Triton compiler. It enables experimentation and integration of out-of-tree components. ```text New API Enforces layerings Full featured: dialects, custom operations Don’t need to work on a fork! Don’t need to recompile compiler to experiment. See https://github.com/triton-lang/triton/tree/main/plugins for examples. ``` -------------------------------- ### Run Triton Program with Interpreter and PDB Source: https://github.com/openai/triton/blob/main/docs/programming-guide/chapter-3/debugging.rst Use the interpreter and pdb for step-by-step debugging of Triton programs. Set TRITON_INTERPRET to 1 and use pdb commands to set breakpoints and run the script. ```bash TRITON_INTERPRET=1 pdb main.py b main.py: r ``` -------------------------------- ### Add MLIR Documentation for Dialect Source: https://github.com/openai/triton/blob/main/third_party/proton/Dialect/include/Dialect/Proton/IR/CMakeLists.txt Adds a target to generate documentation for the Proton dialect using MLIR's doc generation utility. ```cmake add_mlir_doc(ProtonDialect ProtonDialect dialects/ -gen-dialect-doc -dialect=proton) ``` -------------------------------- ### Define Sources for GPUInstrumentationTestLib Source: https://github.com/openai/triton/blob/main/test/lib/Instrumentation/CMakeLists.txt Specifies the source files for the GPUInstrumentationTestLib. ```cmake set(GPUInstrumentationTestLib_SOURCES GPUHello.cpp ) ``` -------------------------------- ### Configure Dependencies for Build Helpers Source: https://github.com/openai/triton/blob/main/CMakeLists.txt Configures dependencies for build helpers, ensuring that configure outputs are regenerated when helper inputs change. This is typically done at the beginning of the CMakeLists.txt file. ```cmake set_property( DIRECTORY APPEND PROPERTY CMAKE_CONFIGURE_DEPENDS "${CMAKE_CURRENT_SOURCE_DIR}/python/build_helpers.py" "${CMAKE_CURRENT_SOURCE_DIR}/cmake/llvm-info.json" "${CMAKE_CURRENT_SOURCE_DIR}/cmake/json-version.txt" ) find_package(Python3 REQUIRED COMPONENTS Interpreter) set(TRITON_THIRD_PARTY_CMAKE_VARS_FILE "${CMAKE_CURRENT_BINARY_DIR}/triton-third-party-vars.cmake") execute_process( COMMAND ${Python3_EXECUTABLE} "${CMAKE_CURRENT_SOURCE_DIR}/python/build_helpers.py" write_thirdparty_cmake_vars ${TRITON_BUILD_HELPER_COMMON_ARGS} --output "${TRITON_THIRD_PARTY_CMAKE_VARS_FILE}" --packages llvm json WORKING_DIRECTORY "${CMAKE_CURRENT_SOURCE_DIR}" COMMAND_ERROR_IS_FATAL ANY ) include("${TRITON_THIRD_PARTY_CMAKE_VARS_FILE}") ``` -------------------------------- ### Enable Kernel Dumping and Set Dump Directory Source: https://github.com/openai/triton/blob/main/README.md Set TRITON_KERNEL_DUMP to 1 to enable dumping of IR and PTX/AMDGCN. Specify the TRITON_DUMP_DIR to set the save location. ```bash export TRITON_KERNEL_DUMP=1 export TRITON_DUMP_DIR= ``` -------------------------------- ### Compiler Hint Operations Source: https://github.com/openai/triton/blob/main/docs/python-api/triton.language.rst Hints for the Triton compiler to optimize code. ```APIDOC ## Compiler Hint Operations Provides hints to the Triton compiler for optimization purposes. ### Functions - **assume**: Asserts a condition that the compiler can use for optimization. - **debug_barrier**: Inserts a barrier for debugging purposes. - **max_constancy**: Hints at the maximum level of constancy for an operation. - **max_contiguous**: Hints at the maximum contiguity for memory accesses. - **multiple_of**: Indicates that a value is a multiple of a given number. ``` -------------------------------- ### Add ProtonData Library Source: https://github.com/openai/triton/blob/main/third_party/proton/csrc/lib/Data/CMakeLists.txt This snippet shows how to add the ProtonData library using the add_proton_library CMake command, listing all its source files. ```cmake add_proton_library(ProtonData Data.cpp Dump/ChromeTrace.cpp Dump/TraceDataDump.cpp Metric.cpp TraceData.cpp TreeData.cpp ) ``` -------------------------------- ### Add TritonGPU LinearLayoutConversions Unit Test Source: https://github.com/openai/triton/blob/main/unittest/Dialect/TritonGPU/CMakeLists.txt Sets up a unit test for LinearLayoutConversions within TritonGPU, specifying its core library dependencies. ```cmake add_triton_ut( NAME LinearLayoutConversions SRCS LinearLayoutConversionsTest.cpp LIBS TritonGPUIR TritonGPUTransforms TritonNvidiaGPUTransforms ) ``` -------------------------------- ### Define TritonTestProton Library Source: https://github.com/openai/triton/blob/main/test/lib/Proton/CMakeLists.txt Defines a test library named 'TritonTestProton' using the specified C++ source file. ```cmake add_library(TritonTestProton TestScopeIdAllocation.cpp) ``` -------------------------------- ### jit Source: https://github.com/openai/triton/blob/main/docs/gluon/api/runtime.rst Decorator or function to perform Just-In-Time (JIT) compilation for Triton kernels. It allows users to define Python functions that will be compiled into optimized kernels for specific hardware. ```APIDOC ## jit ### Description Decorator or function to perform Just-In-Time (JIT) compilation for Triton kernels. It allows users to define Python functions that will be compiled into optimized kernels for specific hardware. ### Usage Users can apply this as a decorator to Python functions or call it to compile kernels. ### Parameters (No specific parameters are detailed in the source for direct user invocation.) ### Returns (The return type is implied to be a compiled kernel or a callable object representing it.) ``` -------------------------------- ### Build Shared Libraries for Triton Plugins Source: https://github.com/openai/triton/blob/main/examples/plugins/CMakeLists.txt Iterates through the defined plugin passes, creating a shared library for each. It also sets up dependencies and links the necessary Triton libraries. ```cmake foreach( plugin ${TRITON_PLUGIN_PASSES} ) add_library(${plugin} SHARED ${${plugin}_SOURCES}) add_dependencies(${plugin} TritonTableGen TritonGPUTableGen TritonNvidiaGPUTableGen TritonCanonicalizeIncGen TritonPluginsIncGen ) target_link_libraries(${plugin} PRIVATE triton) # CMAKE_LIBRARY_OUTPUT_DIRECTORY is only set during the Python # build. It is empty if building directly from the root # CMakeLists.txt file. Therefore if not building from Python just # use the default CMake shared lib path otherwise this causes a hard # build error if(DEFINED CMAKE_LIBRARY_OUTPUT_DIRECTORY) set_target_properties(${plugin} PROPERTIES LIBRARY_OUTPUT_DIRECTORY "${CMAKE_LIBRARY_OUTPUT_DIRECTORY}/../plugins") endif(DEFINED CMAKE_LIBRARY_OUTPUT_DIRECTORY) target_compile_options(${plugin} PRIVATE -fvisibility=hidden ${TRITON_DISABLE_EH_RTTI_FLAGS}) endforeach() ``` -------------------------------- ### Link TritonTestProton with Dependencies Source: https://github.com/openai/triton/blob/main/test/lib/Proton/CMakeLists.txt Links the 'TritonTestProton' library with public dependencies including MLIRPass and ProtonAnalysis. ```cmake target_link_libraries(TritonTestProton PUBLIC MLIRPass ProtonAnalysis) ``` -------------------------------- ### Original MLIR Module Source: https://github.com/openai/triton/blob/main/examples/plugins/README.md This is the initial MLIR code before the custom out-of-tree pass is applied. It defines a simple function '@foo'. ```MLIR module attributes {"ttg.num-warps" = 4 : i32, ttg.target = "cuda:80"} { tt.func @foo() { tt.return } } ``` -------------------------------- ### Generate Testing Backend Registry Source: https://github.com/openai/triton/blob/main/third_party/proton/test/unittest/Backend/CMakeLists.txt Uses a custom CMake macro 'codegen_proton_backend_registry' to generate C++ code for registering the testing backend. ```cmake codegen_proton_backend_registry( "${CMAKE_CURRENT_BINARY_DIR}/RegisteredBackends.cpp" TestBackend ) ``` -------------------------------- ### Build Triton Plugins Source: https://github.com/openai/triton/blob/main/CMakeLists.txt Adds the Triton plugins subdirectory for building when Triton extensions are enabled. ```cmake if(TRITON_EXT_ENABLED) add_subdirectory(examples/plugins) endif() ``` -------------------------------- ### Add MLIR Operation Documentation Target Source: https://github.com/openai/triton/blob/main/third_party/amd/include/Dialect/TritonAMDGPU/IR/CMakeLists.txt Uses add_mlir_doc to create a documentation target for TritonAMDGPU operations. Requires '-gen-op-doc' flag. ```cmake add_mlir_doc(TritonAMDGPUOps TritonAMDGPUOps dialects/ -gen-op-doc -dialect=amdg) ``` -------------------------------- ### Run Triton Program with Compute-Sanitizer Source: https://github.com/openai/triton/blob/main/docs/programming-guide/chapter-3/debugging.rst Use compute-sanitizer to check for data races and memory access issues when debugging Triton programs on NVIDIA GPUs. Prepend the command to your execution. ```bash compute-sanitizer ```