### Install Triton Python Package from Source Source: https://github.com/triton-lang/triton/blob/main/docs/getting-started/installation.rst Clones the Triton repository, navigates to the Python directory, installs necessary build-time dependencies (ninja, cmake, wheel), and then installs the Triton Python package in editable mode. This setup is suitable for local development and contributions, allowing modifications to be immediately reflected. ```bash git clone https://github.com/triton-lang/triton.git; cd triton/python; pip install ninja cmake wheel; # build-time dependencies pip install -e . ``` -------------------------------- ### Run Triton Unit Tests Source: https://github.com/triton-lang/triton/blob/main/docs/getting-started/installation.rst Installs the required test dependencies for Triton and then executes the unit tests using pytest. This step is crucial for verifying the correctness and integrity of the installed Triton library after a source build. ```bash pip install -e '.[tests]' pytest -vs test/unit/ ``` -------------------------------- ### Install Triton Stable Release via pip Source: https://github.com/triton-lang/triton/blob/main/docs/getting-started/installation.rst Installs the latest stable version of the Triton library using pip. This method provides pre-built binary wheels for CPython 3.8-3.12 and PyPy 3.8-3.9, ensuring a quick and easy setup. ```bash pip install triton ``` -------------------------------- ### Run Triton Benchmarks Source: https://github.com/triton-lang/triton/blob/main/docs/getting-started/installation.rst Navigates to the benchmark directory and runs the Triton benchmarks, generating performance plots and saving results to a specified directory. This allows users to evaluate the performance characteristics of their Triton installation. ```bash cd bench python -m run --with-plots --result-dir /tmp/triton-bench ``` -------------------------------- ### Install Triton Nightly Release via pip Source: https://github.com/triton-lang/triton/blob/main/docs/getting-started/installation.rst Installs the latest nightly build of the Triton library from a specified index URL. This provides access to the most recent development features and bug fixes, ideal for users who need the cutting edge. ```bash pip install -U --index-url https://aiinfra.pkgs.visualstudio.com/PublicPackages/_packaging/Triton-Nightly/pypi/simple/ triton-nightly ``` -------------------------------- ### Install Triton via pip Source: https://github.com/triton-lang/triton/blob/main/README.md Installs the latest stable release of Triton using pip, providing pre-built binary wheels for CPython 3.9-3.13. This is the quickest way to get started with Triton. ```shell pip install triton ``` -------------------------------- ### Install Triton Tutorial Dependencies Source: https://github.com/triton-lang/triton/blob/main/python/tutorials/README.rst This command sequence installs the necessary Python dependencies for running the Triton tutorials. It first navigates into the 'triton' directory and then installs the 'tutorials' extra for the Python package in editable mode. ```bash cd triton pip install -e './python[tutorials]' ``` -------------------------------- ### Setup VSCode IntelliSense for Triton C++ Development Source: https://github.com/triton-lang/triton/blob/main/README.md This section provides a step-by-step guide for configuring VSCode IntelliSense to work with Triton's C++ codebase. It involves performing a local build to generate the `compile_commands.json` file, locating this file, and then configuring the C/C++ extension in VSCode to use it for accurate code completion and navigation. ```shell pip install -e . find ./build -name 'compile_commands.json' | xargs readlink -f ``` -------------------------------- ### Install Proton from Source Source: https://github.com/triton-lang/triton/blob/main/third_party/proton/README.md Instructions to clone the Triton repository and install the Proton profiler from its Python directory using pip. ```bash git clone https://github.com/triton-lang/triton cd triton/python pip install . ``` -------------------------------- ### Install Hatchet for Profile Visualization Source: https://github.com/triton-lang/triton/blob/main/third_party/proton/README.md Instructions to install the `llnl-hatchet` library, which is required to visualize Proton's default JSON profile data. Note that `pip install hatchet` is not the correct command due to API differences. ```bash pip install llnl-hatchet ``` -------------------------------- ### Install Triton from source with virtual environment Source: https://github.com/triton-lang/triton/blob/main/README.md Clones the Triton repository and installs it in editable mode within a dedicated Python virtual environment. This ensures isolated dependencies and a clean development setup, preventing conflicts with system-wide packages. ```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 . ``` -------------------------------- ### Execute Triton Test Suite Source: https://github.com/triton-lang/triton/blob/main/README.md This snippet provides the necessary commands to set up the development environment and run Triton's test suite. It includes commands for a one-time development installation and options to run all tests, both with and without GPU acceleration. ```shell # One-time setup: make dev-install # To run all tests (requires a GPU): make test # Or, to run tests without a GPU: make test-nogpu ``` -------------------------------- ### Install Triton from source Source: https://github.com/triton-lang/triton/blob/main/README.md Clones the Triton repository and installs it in editable mode, including build-time dependencies. This method is suitable for development and allows for local modifications to the Triton source code. ```shell git clone https://github.com/triton-lang/triton.git cd triton pip install -r python/requirements.txt # build-time dependencies pip install -e . ``` -------------------------------- ### Build Triton with custom LLVM using dev-install Source: https://github.com/triton-lang/triton/blob/main/README.md Provides an automated command to build LLVM and install Triton using a custom LLVM version. This simplifies the process for developers who need to work with a specific or modified LLVM build. ```shell make dev-install-llvm ``` -------------------------------- ### Print Autotuning Configuration and Time Source: https://github.com/triton-lang/triton/blob/main/README.md Prints the best autotuning configuration and the total time spent for each kernel after the autotuning process is complete. ```Shell TRITON_PRINT_AUTOTUNING=1 ``` -------------------------------- ### Visualize Sorted Profile Data with proton-viewer Source: https://github.com/triton-lang/triton/blob/main/third_party/proton/README.md This command-line example demonstrates how to use `proton-viewer` with the `--print-sorted` flag to display a sorted list of kernels based on the first specified metric, such as `time/ns`. It helps in quickly identifying performance bottlenecks. ```bash proton-viewer -m time/ns,time/% --print-sorted ``` -------------------------------- ### Manually build LLVM for Triton Source: https://github.com/triton-lang/triton/blob/main/README.md Detailed steps to manually build LLVM from source at a specific revision required by Triton. This is an alternative to the automated `dev-install-llvm` command, useful for fine-grained control over the LLVM build process. ```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" -DLLVM_TARGETS_TO_BUILD="host;NVPTX;AMDGPU" ninja ``` -------------------------------- ### Configure Triton Build Environment Variables and Pip Flags Source: https://github.com/triton-lang/triton/blob/main/README.md This snippet outlines various environment variables and `pip install` flags that can be used to optimize the Triton build process. These include enabling clang/lld for faster builds, utilizing ccache, specifying a custom cache directory, limiting parallel compilation jobs, and speeding up subsequent `pip install` operations. ```shell TRITON_BUILD_WITH_CLANG_LLD=true TRITON_BUILD_WITH_CCACHE=true TRITON_HOME=/some/path # To limit jobs during pip install: MAX_JOBS= pip install -e . # To speed up nop builds: pip install --no-build-isolation -e . ``` -------------------------------- ### Enable LLVM IR Dumping Source: https://github.com/triton-lang/triton/blob/main/README.md Dumps the LLVM Intermediate Representation (IR) before every pass run over the LLVM IR. ```Shell LLVM_IR_ENABLE_DUMP=1 ``` -------------------------------- ### Install Triton with manually built LLVM Source: https://github.com/triton-lang/triton/blob/main/README.md Installs Triton using a manually built LLVM by setting specific environment variables to point to the custom LLVM build directory. This step follows the manual LLVM build process and ensures Triton links against the desired LLVM version. ```shell 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 . ``` -------------------------------- ### Halide/TVM Scheduling Language Example for Matrix Multiplication Source: https://github.com/triton-lang/triton/blob/main/docs/programming-guide/chapter-2/related-work.rst This example demonstrates the separation of concerns in scheduling languages like Halide and TVM. The algorithm for matrix multiplication is defined separately from its optimization schedule, allowing independent maintenance and distribution. The schedule applies transformations such as vectorization, splitting, reordering, and parallelization to optimize performance. ```python // algorithm Var x("x"), y("y"); Func matmul("matmul"); RDom k(0, matrix_size); RVar ki; matmul(x, y) = 0.0f; matmul(x, y) += A(k, y) * B(x, k); // schedule Var xi("xi"), xo("xo"), yo("yo"), yi("yo"), yii("yii"), xii("xii"); matmul.vectorize(x, 8); matmul.update(0) .split(x, x, xi, block_size).split(xi, xi, xii, 8) .split(y, y, yi, block_size).split(yi, yi, yii, 4) .split(k, k, ki, block_size) .reorder(xii, yii, xi, ki, yi, k, x, y) .parallel(y).vectorize(xii).unroll(xi).unroll(yii); ``` -------------------------------- ### Define Triton Unit Tests with add_triton_ut Source: https://github.com/triton-lang/triton/blob/main/unittest/Dialect/TritonGPU/CMakeLists.txt The `add_triton_ut` command is used to register a new unit test within the Triton build system. It takes parameters for the test's `NAME`, its `SRCS` (source files), and a list of `LIBS` (libraries) it depends on. This configuration helps the build system compile and link the test correctly. ```Triton Build Config add_triton_ut( NAME TestSwizzling SRCS SwizzleTest.cpp LIBS TritonTools LLVMSupport MLIRSupport ) ``` ```Triton Build Config add_triton_ut( NAME Dialect SRCS DialectTest.cpp LIBS MLIRParser TritonGPUIR TritonGPUTransforms TritonNvidiaGPUTransforms ) ``` ```Triton Build Config add_triton_ut( NAME LinearLayoutConversions SRCS LinearLayoutConversionsTest.cpp LIBS TritonGPUIR TritonGPUTransforms TritonNvidiaGPUTransforms ) ``` ```Triton Build Config add_triton_ut( NAME DumpLayoutTest SRCS DumpLayoutTest.cpp LIBS TritonGPUIR TritonGPUTransforms TritonNvidiaGPUTransforms ) ``` -------------------------------- ### Enable LLVM Debugging Output Source: https://github.com/triton-lang/triton/blob/main/README.md Passes the `-debug` flag to LLVM, resulting in extensive debugging information printed to stdout. For less verbose output, consider using `TRITON_LLVM_DEBUG_ONLY` or extracting IR and running `opt` standalone with `-debug-only=foo`. ```Shell TRITON_ENABLE_LLVM_DEBUG=1 ``` -------------------------------- ### Triton Project CMake Build Setup Source: https://github.com/triton-lang/triton/blob/main/third_party/nvidia/CMakeLists.txt This snippet configures the core build process for the Triton project using CMake. It defines include directories, adds subdirectories for various components (like 'include', 'lib', 'hopper', 'unittest'), and conditionally compiles a Python module ('TritonNVIDIA') and links it with Python and pybind11 if 'TRITON_BUILD_PYTHON_MODULE' is enabled. It also conditionally includes unit tests if 'TRITON_BUILD_UT' is set. ```CMake include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include) include_directories(${CMAKE_CURRENT_BINARY_DIR}/include) add_subdirectory(include) add_subdirectory(lib) if(TRITON_BUILD_PYTHON_MODULE) add_triton_plugin(TritonNVIDIA ${CMAKE_CURRENT_SOURCE_DIR}/triton_nvidia.cc LINK_LIBS TritonNVIDIAGPUToLLVM NVGPUToLLVM) target_link_libraries(TritonNVIDIA PRIVATE Python3::Module pybind11::headers) endif() if(TRITON_BUILD_UT) add_subdirectory(unittest) endif() add_subdirectory(hopper) ``` -------------------------------- ### Dump LLVM Pass Timing Information Source: https://github.com/triton-lang/triton/blob/main/README.md Dumps timing information for each LLVM pass executed. ```Shell LLVM_ENABLE_TIMING ``` -------------------------------- ### Enable Instruction Sampling with Proton in Python Source: https://github.com/triton-lang/triton/blob/main/third_party/proton/README.md This snippet demonstrates how to initiate instruction sampling using the `triton.profiler` module. It starts a profiling session named 'profile_name' with a 'shadow' context and 'cupti_pcsampling' backend, enabling experimental instruction sampling on NVIDIA GPUs. ```python import triton.profiler as proton proton.start(name="profile_name", context="shadow", backend="cupti_pcsampling") ``` -------------------------------- ### Define TritonGPUIR Library Build Configuration Source: https://github.com/triton-lang/triton/blob/main/lib/Dialect/TritonGPU/IR/CMakeLists.txt This configuration snippet specifies the build process for the `TritonGPUIR` library within the Triton project's build system. It enumerates the C++ source files that comprise the library, lists internal dependencies required for its compilation (e.g., `TritonGPUTableGen`), and declares external libraries it publicly links against, such as `MLIRGPUDialect` and `TritonIR`. This setup ensures the library is correctly compiled and integrated with its necessary components. ```Build System Configuration add_triton_library(TritonGPUIR Dialect.cpp LinearLayoutConversions.cpp LayoutUtility.cpp Ops.cpp Types.cpp DEPENDS TritonGPUTableGen TritonGPUAttrDefsIncGen TritonGPUTypeInterfacesIncGen LINK_LIBS PUBLIC MLIRGPUDialect TritonIR TritonTools ) ``` -------------------------------- ### Triton LLVM/MLIR Build System Setup Source: https://github.com/triton-lang/triton/blob/main/include/triton/Dialect/TritonInstrument/Transforms/CMakeLists.txt This snippet configures the build process for Triton, setting LLVM target definitions, generating MLIR pass declarations using `mlir_tablegen`, and adding a public TableGen target for transformations. These commands are typically found in CMakeLists.txt files within LLVM/MLIR-based projects. ```CMake set(LLVM_TARGET_DEFINITIONS Passes.td) mlir_tablegen(Passes.h.inc -gen-pass-decls -name TritonInstrument) add_public_tablegen_target(TritonInstrumentTransformsIncGen) ``` -------------------------------- ### Specify MLIR Dump Output Path Source: https://github.com/triton-lang/triton/blob/main/README.md Defines the destination for MLIR IR dumps. If unset, dumps to stderr. ```Shell MLIR_DUMP_PATH= ``` -------------------------------- ### Generate MLIR Reproducer File Source: https://github.com/triton-lang/triton/blob/main/README.md Specifies a path to generate an MLIR reproducer file before each MLIR compiler stage. If a stage fails, the file at this path will be a local MLIR reproducer captured just before the failing pass, aiding in debugging. ```Shell TRITON_REPRODUCER_PATH= ``` -------------------------------- ### Dump MLIR Pass Timing Information Source: https://github.com/triton-lang/triton/blob/main/README.md Dumps timing information for each MLIR pass executed. ```Shell MLIR_ENABLE_TIMING ``` -------------------------------- ### C Loop Nest for Polyhedral Iteration Domain Example Source: https://github.com/triton-lang/triton/blob/main/docs/programming-guide/chapter-2/related-work.rst This C code snippet demonstrates a simple loop nest that forms a Static Control Part (SCoP). In polyhedral compilation, such constructs, where loop bounds and array accesses are affine functions of loop indices, define iteration domains that are bounded by affine inequalities, forming polyhedra. ```C for(int i = 0; i < 3; i++) for(int j = i; j < 5; j++) A[i][j] = 0; ``` -------------------------------- ### Define Triton Unit Test for PTX Assembly Format Source: https://github.com/triton-lang/triton/blob/main/third_party/nvidia/unittest/Conversion/TritonGPUToLLVM/CMakeLists.txt This snippet configures a unit test named `TestPtxAsmFormat` within the Triton project. It specifies `PTXAsmFormatTest.cpp` as the source file and links several Triton GPU-related libraries, including `TritonGPUToLLVM`, `TritonNVIDIAGPUToLLVM`, `NVGPUIR`, and `MLIRUBToLLVM`. This setup is crucial for testing the PTX assembly format within the Triton framework. ```Build Script add_triton_ut( NAME TestPtxAsmFormat SRCS PTXAsmFormatTest.cpp LIBS TritonGPUToLLVM TritonNVIDIAGPUToLLVM NVGPUIR MLIRUBToLLVM ) ``` -------------------------------- ### Display proton-viewer Help Options Source: https://github.com/triton-lang/triton/blob/main/third_party/proton/README.md This command shows all available command-line options and flags for the `proton-viewer` utility, providing comprehensive guidance on its usage. ```bash proton-viewer -h ``` -------------------------------- ### Triton Kernel Override Workflow using Bash Source: https://github.com/triton-lang/triton/blob/main/README.md This bash script provides a step-by-step guide to perform a kernel override in Triton. It involves setting specific environment variables to dump kernel intermediate representations (IRs) and then load modified versions for re-execution, enabling advanced debugging and optimization. ```bash export TRITON_ALWAYS_COMPILE=1 export TRITON_KERNEL_DUMP=1 export TRITON_DUMP_DIR= export TRITON_KERNEL_OVERRIDE=1 export TRITON_OVERRIDE_DIR= # Step 1: Run the kernel once to dump kernel's IRs and ptx/amdgcn in $TRITON_DUMP_DIR # Step 2: Copy $TRITON_DUMP_DIR/ to $TRITON_OVERRIDE_DIR # Step 3: Delete the stages that you do not want to override and modify the stage you do want to override # Step 4: Run the kernel again to see the overridden result ``` -------------------------------- ### Install Triton Python Package Without Building Proton Source: https://github.com/triton-lang/triton/blob/main/third_party/proton/README.md How to install Triton's Python package while explicitly disabling the build of the Proton profiler by setting the TRITON_BUILD_PROTON environment variable to OFF. ```bash TRITON_BUILD_PROTON=OFF pip install . ``` -------------------------------- ### Define TritonToTritonGPU Library Source: https://github.com/triton-lang/triton/blob/main/lib/Conversion/TritonToTritonGPU/CMakeLists.txt Defines the `TritonToTritonGPU` library within the Triton project's build system. It lists the C++ source files that compose the library, its build-time dependencies, and the public libraries it links against, such as MLIR components and Triton-specific IRs. ```Build Configuration add_triton_library(TritonToTritonGPU RelayoutTritonGPU.cpp TritonGPUConversion.cpp TritonToTritonGPUPass.cpp DEPENDS TritonConversionPassIncGen LINK_LIBS PUBLIC MLIRIR MLIRPass MLIRTransforms TritonIR ProtonIR TritonGPUIR ) ``` -------------------------------- ### Example of unsupported indirect memory access in Triton interpreter (Python) Source: https://github.com/triton-lang/triton/blob/main/docs/programming-guide/chapter-3/debugging.rst This Python code snippet provides an example of an indirect memory access pattern that is not supported by the Triton interpreter. Such patterns involve loading a pointer value from memory and then using that loaded pointer to access another memory location. ```python ptr = tl.load(ptr) x = tl.load(ptr) ``` -------------------------------- ### Build and Link `triton-tensor-layout` Tool (CMake) Source: https://github.com/triton-lang/triton/blob/main/bin/CMakeLists.txt Configures the `triton-tensor-layout` executable, a tool likely related to tensor layout optimization or analysis within Triton. It adds the executable and links it against Triton-specific libraries, MLIR conversion and dialect libraries, and Triton test components. ```CMake add_llvm_executable(triton-tensor-layout triton-tensor-layout.cpp PARTIAL_SOURCES_INTENDED) target_link_libraries(triton-tensor-layout PRIVATE ${triton_libs} ${conversion_libs} ${dialect_libs} TritonTestAnalysis TritonTestDialect TritonAMDGPUTestAnalysis ) ``` -------------------------------- ### Profile a Python Code Region with Proton Source: https://github.com/triton-lang/triton/blob/main/third_party/proton/README.md Shows how to use `proton.start` to begin profiling a code region, `proton.deactivate` to temporarily skip a region, `proton.activate` to resume, and `proton.finalize` to write out profile data and complete the session. ```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() ``` -------------------------------- ### C Loop Example for Sparse Matrix Computation Source: https://github.com/triton-lang/triton/blob/main/docs/programming-guide/chapter-2/related-work.rst This C code snippet illustrates a nested loop structure for a computation involving irregular iteration spaces, typical in sparse matrix operations. The inner loop's bound `K[i]` and the column access `col[i, k]` depend on outer loop indices, which is problematic for traditional scheduling languages that impose severe constraints on such dependencies. This example highlights a limitation that Triton's block-based representation aims to address by allowing block-structured 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; ``` -------------------------------- ### Override Default Floating Point Fusion Behavior Source: https://github.com/triton-lang/triton/blob/main/README.md Overrides the default behavior regarding floating-point fusion (e.g., `mul+add` to `fma`). ```Shell TRITON_DEFAULT_FP_FUSION ``` -------------------------------- ### Include Utility, Binaries, and Test Subdirectories in CMake Source: https://github.com/triton-lang/triton/blob/main/CMakeLists.txt Adds standard subdirectories for 'f2reduce' (likely a utility library), 'bin' (for executables), and 'test' (for general project tests) to the CMake build system. This organizes the project structure and ensures all components are built. ```CMake add_subdirectory(third_party/f2reduce) add_subdirectory(bin) add_subdirectory(test) ``` -------------------------------- ### Force Kernel Compilation Source: https://github.com/triton-lang/triton/blob/main/README.md Forces Triton to compile kernels regardless of whether a cached version is available, ensuring fresh compilation. ```Shell TRITON_ALWAYS_COMPILE=1 ``` -------------------------------- ### Provide Kernel Metadata with Triton Hook Source: https://github.com/triton-lang/triton/blob/main/third_party/proton/README.md Explains how to use the 'triton' hook to enable Proton to invoke a `launch_metadata` function before launching a GPU kernel. This function returns a dictionary of metadata for the kernel, such as its name and FLOPs. ```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)) ``` -------------------------------- ### Define ProtonContext Library with Source Files Source: https://github.com/triton-lang/triton/blob/main/third_party/proton/csrc/lib/Context/CMakeLists.txt This build system command defines the 'ProtonContext' library and specifies the C++ source files (Context.cpp, Python.cpp, Shadow.cpp) that comprise it. This typically instructs the build system to compile and link these files to create the library. ```Build System add_proton_library(ProtonContext Context.cpp Python.cpp Shadow.cpp ) ``` -------------------------------- ### Enable Triton Interpreter Mode Source: https://github.com/triton-lang/triton/blob/main/README.md Forces Triton to use its interpreter instead of running on the GPU, allowing for Python breakpoints within kernel code for easier debugging. ```Shell TRITON_INTERPRET=1 ``` -------------------------------- ### Build and Link `triton-lsp` Language Server (CMake) Source: https://github.com/triton-lang/triton/blob/main/bin/CMakeLists.txt Configures the `triton-lsp` executable, which provides Language Server Protocol (LSP) support for Triton. It adds the executable, updates its compile flags, and links it against core MLIR LSP server libraries, Triton-specific libraries, and test components to enable IDE features. ```CMake add_llvm_executable(triton-lsp triton-lsp.cpp PARTIAL_SOURCES_INTENDED) llvm_update_compile_flags(triton-lsp) target_link_libraries(triton-lsp PRIVATE ${dialect_libs} ${conversion_libs} ${triton_libs} # tests TritonTestAnalysis TritonTestDialect TritonAMDGPUTestAnalysis # MLIR core MLIRLspServerLib MLIRPass MLIRTransforms ) mlir_check_all_link_libraries(triton-lsp) ``` -------------------------------- ### Customize Call Path with proton.state Source: https://github.com/triton-lang/triton/blob/main/third_party/proton/README.md This Python example demonstrates how to use `proton.state` to customize the call path of GPU operations. Unlike `scope`, `state` is not recursive and acts as a suffix to the original call path, allowing for fine-grained annotation of individual operations. ```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) ``` -------------------------------- ### Define Triton Unit Tests for Layout Components Source: https://github.com/triton-lang/triton/blob/main/unittest/Tools/CMakeLists.txt This snippet uses the `add_triton_ut` macro or function to register a unit test suite within the Triton build system. It takes the `NAME` of the test suite, a list of `SRCS` (source files) to compile, and `LIBS` (libraries) to link against. This ensures that `LinearLayoutTest.cpp` and `LayoutUtilsTest.cpp` are compiled and linked with `TritonTools` to create the `LinearLayout` test executable. ```Triton Build System add_triton_ut( NAME LinearLayout SRCS LayoutUtilsTest.cpp LinearLayoutTest.cpp LIBS TritonTools ) ``` -------------------------------- ### Final Include Directories and Subdirectory Inclusion Source: https://github.com/triton-lang/triton/blob/main/CMakeLists.txt This section adds a comprehensive list of include directories to the project, including current directory, MLIR, LLVM, project source/binary includes, and third-party directories. It also includes the `include` and `lib` subdirectories, which likely contain source code and further CMake configurations. ```CMake include_directories(".") include_directories(${MLIR_INCLUDE_DIRS}) include_directories(${LLVM_INCLUDE_DIRS}) include_directories(${PROJECT_SOURCE_DIR}/include) include_directories(${PROJECT_BINARY_DIR}/include) # Tablegen'd files include_directories(${PROJECT_SOURCE_DIR}/third_party) include_directories(${PROJECT_BINARY_DIR}/third_party) # Tablegen'd files # link_directories(${LLVM_LIBRARY ``` -------------------------------- ### Disable LLVM Optimizations Source: https://github.com/triton-lang/triton/blob/main/README.md Disables LLVM optimizations for `make_llir` and `make_ptx`. Can be set to a boolean `true` to disable all, or a list of flags (e.g., `disable-lsr`) to disable specific optimizations. Loop strength reduction (`disable-lsr`) is noted for causing performance changes. ```Shell DISABLE_LLVM_OPT=true DISABLE_LLVM_OPT="disable-lsr" ``` -------------------------------- ### Measure CPU Time with cpu_timed_scope in Triton Source: https://github.com/triton-lang/triton/blob/main/third_party/proton/README.md This Python example demonstrates how to use `proton.cpu_timed_scope` to measure the CPU time spent within a specific code block. It wraps a `scope` and records `cpu_time` distinct from accelerator `time`, which is useful for analyzing CPU-bound operations. ```python import triton.profiler as proton with proton.cpu_timed_scope("test"): foo[1,](x, y) ``` -------------------------------- ### Generate NVGPU Dialect and Operations Files using MLIR TableGen Source: https://github.com/triton-lang/triton/blob/main/third_party/nvidia/include/Dialect/NVGPU/IR/CMakeLists.txt This configuration block uses `mlir_tablegen` to generate C++ declarations and definitions for the NVGPU dialect and its operations from `.td` files. It also includes commands to generate documentation for the dialect and operations, and defines a public TableGen target. ```CMake set(LLVM_TARGET_DEFINITIONS NVGPUOps.td) mlir_tablegen(Dialect.h.inc -gen-dialect-decls -dialect=nvgpu) mlir_tablegen(Dialect.cpp.inc -gen-dialect-defs -dialect=nvgpu) mlir_tablegen(OpsConversions.inc -gen-llvmir-conversions) mlir_tablegen(Ops.h.inc -gen-op-decls) mlir_tablegen(Ops.cpp.inc -gen-op-defs) mlir_tablegen(OpsEnums.h.inc -gen-enum-decls) mlir_tablegen(OpsEnums.cpp.inc -gen-enum-defs) add_mlir_doc(NVGPUDialect NVGPUDialect dialects/ -gen-dialect-doc) add_mlir_doc(NVGPUOps NVGPUOps dialects/ -gen-op-doc) add_public_tablegen_target(NVGPUTableGen) ``` -------------------------------- ### Build and Link `triton-opt` MLIR Optimizer (CMake) Source: https://github.com/triton-lang/triton/blob/main/bin/CMakeLists.txt Configures the `triton-opt` executable, an MLIR optimizer tool. It adds the executable, updates its compile flags, and links it against core MLIR optimization libraries, Triton-specific libraries, and various test components, ensuring all necessary dependencies are met for compilation and execution. ```CMake add_llvm_executable(triton-opt triton-opt.cpp PARTIAL_SOURCES_INTENDED) llvm_update_compile_flags(triton-opt) target_link_libraries(triton-opt PRIVATE ${dialect_libs} ${conversion_libs} ${triton_libs} # tests TritonTestAnalysis TritonTestDialect TritonAMDGPUTestAnalysis # MLIR core MLIROptLib MLIRPass MLIRTransforms ) mlir_check_all_link_libraries(triton-opt) ``` -------------------------------- ### Limit LLVM Debug Output to Specific Components Source: https://github.com/triton-lang/triton/blob/main/README.md Equivalent to LLVM's `-debug-only` option, this variable limits LLVM debug output to specific pass or component names (defined by `#define DEBUG_TYPE`). Accepts one or more comma-separated values to filter debug information. ```Shell TRITON_LLVM_DEBUG_ONLY="tritongpu-remove-layout-conversions" TRITON_LLVM_DEBUG_ONLY="tritongpu-remove-layout-conversions,regalloc" ``` -------------------------------- ### Enable MLIR IR Dumping for Debugging Source: https://github.com/triton-lang/triton/blob/main/README.md Controls dumping of MLIR Intermediate Representation (IR) before each MLIR pass. Can be set to '1' for all kernels or 'kernelName' for a specific kernel. Note that Triton cache might interfere, requiring cache clearing (`rm -r ~/.triton/cache/*`). ```Shell MLIR_ENABLE_DUMP=1 MLIR_ENABLE_DUMP=kernelName ``` -------------------------------- ### Configure LIT Site Configuration File Source: https://github.com/triton-lang/triton/blob/main/test/CMakeLists.txt Configures a LIT (LLVM Integrated Tester) site configuration file from a template. This process sets up the testing environment for the project, specifying the main configuration file and output location for the generated configuration. ```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_SOURCe_DIR}/lit.cfg.py ) ``` -------------------------------- ### Triton Compiler Environment Variables Reference Source: https://github.com/triton-lang/triton/blob/main/README.md This section documents the environment variables used to configure the Triton compiler's behavior, including IR dumping, kernel overriding, default floating-point precision, and debugging options. These variables primarily interact with the C++ layer of Triton. ```APIDOC TRITON_KERNEL_DUMP: Description: Enables the dumping of the IR from each compilation stage and the final ptx/amdgcn. TRITON_DUMP_DIR: Description: Specifies the directory to save the dumped IR and ptx/amdgcn when TRITON_KERNEL_DUMP is set to 1. TRITON_KERNEL_OVERRIDE: Description: Enables the override of the compiled kernel with a user-specified IR/ptx/amdgcn at the beginning of each compilation stage. TRITON_OVERRIDE_DIR: Description: Specifies the directory from which to load the IR/ptx/amdgcn files when TRITON_KERNEL_OVERRIDE is set to 1. TRITON_F32_DEFAULT: Description: Sets the default input precision of tl.dot when using 32-bit floats, which can be either ieee, tf32, or tf32x3. TRITON_FRONT_END_DEBUGGING=1: Description: Disables exception wrapping when an error occurs in the compiler frontend, allowing the full stack trace to be seen. TRITON_DISABLE_LINE_INFO=1: Description: Removes all line information from the module. ``` -------------------------------- ### Define Triton Library with Source and Dependencies Source: https://github.com/triton-lang/triton/blob/main/third_party/proton/dialect/lib/TritonProtonToLLVM/CMakeLists.txt This snippet illustrates the definition of a Triton-specific library named `TritonProtonToLLVM`. It includes `RecordOpToLLVM.cpp` as a source file and establishes a public linkage dependency on the `ProtonIR` library. ```CMake add_triton_library(TritonProtonToLLVM RecordOpToLLVM.cpp LINK_LIBS PUBLIC ProtonIR ) ``` -------------------------------- ### Reparse IR for Location Information Source: https://github.com/triton-lang/triton/blob/main/README.md Reparses the Intermediate Representation (IR) to use the line number of the IR file (with `.ttir` or `.ttgir` extension) as location information, instead of the Python file. This provides a direct mapping from IR to LLIR/PTX and can be used with performance tools for instruction breakdown. ```Shell USE_IR_LOC={ttir,ttgir} ``` -------------------------------- ### Enable LLVM Address Sanitizer (ASAN) Source: https://github.com/triton-lang/triton/blob/main/README.md Invokes the LLVM address sanitizer for detecting memory leaks and out-of-bounds access. Currently supported only on the AMD backend and requires specific ASAN libraries. It's recommended to disable memory caching strategies in ROCm and PyTorch for optimal detection. ```Shell TRITON_ENABLE_ASAN=1 ``` -------------------------------- ### Define TritonNvidiaGPUIR Library Build Source: https://github.com/triton-lang/triton/blob/main/lib/Dialect/TritonNvidiaGPU/IR/CMakeLists.txt This snippet defines the `TritonNvidiaGPUIR` library's build process. It specifies the C++ source files (`Dialect.cpp`, `Ops.cpp`), lists internal build dependencies (e.g., `TritonNvidiaGPUTableGen`), and declares the public libraries it links against (`TritonIR`, `TritonGPUIR`). This configuration is crucial for the compilation and linking of the GPU-specific components of Triton. ```Build Configuration add_triton_library(TritonNvidiaGPUIR Dialect.cpp Ops.cpp DEPENDS TritonNvidiaGPUTableGen TritonNvidiaGPUAttrDefsIncGen TritonNvidiaGPUOpInterfacesIncGen LINK_LIBS PUBLIC TritonIR TritonGPUIR ) ``` -------------------------------- ### Build and Link `triton-reduce` MLIR Reducer (CMake) Source: https://github.com/triton-lang/triton/blob/main/bin/CMakeLists.txt Configures the `triton-reduce` executable, an MLIR reduction tool used for minimizing test cases. It adds the executable, updates its compile flags, and links it against core MLIR reduction libraries, Triton-specific libraries, and various test components, similar to `triton-opt` but with `MLIRReduceLib`. ```CMake add_llvm_executable(triton-reduce triton-reduce.cpp PARTIAL_SOURCES_INTENDED) mlir_check_all_link_libraries(triton-reduce) llvm_update_compile_flags(triton-reduce) target_link_libraries(triton-reduce PRIVATE ${dialect_libs} ${conversion_libs} ${triton_libs} # tests TritonTestAnalysis TritonTestDialect TritonAMDGPUTestAnalysis # MLIR core MLIRReduceLib MLIRPass MLIRTransforms ) mlir_check_all_link_libraries(triton-reduce) ``` -------------------------------- ### Control MLIR Diagnostic Emission Source: https://github.com/triton-lang/triton/blob/main/README.md Controls the emission of diagnostics in MLIR. Options include `warnings`, `remarks`, `stacktraces`, and `operations`. Multiple options can be comma-separated. By default, only errors are shown. Setting `warnings` includes errors and warnings; `remarks` includes errors, warnings, and remarks. This replaces the deprecated `MLIR_ENABLE_REMARK`. ```Shell MLIR_ENABLE_DIAGNOSTICS=remarks,operations MLIR_ENABLE_DIAGNOSTICS=warnings,stacktraces ``` -------------------------------- ### Define TritonInstrumentIR Library Source: https://github.com/triton-lang/triton/blob/main/lib/Dialect/TritonInstrument/IR/CMakeLists.txt This configuration block defines the `TritonInstrumentIR` library within the Triton project's custom build system. It specifies the source files (`Dialect.cpp`, `Ops.cpp`) that comprise the library, declares a dependency on `TritonInstrumentTableGen`, and lists the public libraries it links against, including `MLIRIR`, `TritonIR`, and `TritonGPUIR`. ```Triton Build Config add_triton_library(TritonInstrumentIR Dialect.cpp Ops.cpp DEPENDS TritonInstrumentTableGen LINK_LIBS PUBLIC MLIRIR TritonIR TritonGPUIR ) ``` -------------------------------- ### Define Triton Unit Test: TestOptimizeLDS Source: https://github.com/triton-lang/triton/blob/main/third_party/amd/unittest/Conversion/CMakeLists.txt This snippet defines a unit test named 'TestOptimizeLDS' for the Triton project. It specifies 'OptimizeLDSTest.cpp' as the source file and lists several Triton-specific libraries required for compilation and linking, including analysis, IR, GPU, and AMD-specific components. ```Triton Build Configuration add_triton_ut( NAME TestOptimizeLDS SRCS OptimizeLDSTest.cpp LIBS TritonAnalysis TritonIR TritonGPUIR TritonAMDGPUToLLVM MLIRUBToLLVM TritonAMDUtils TritonAMDAnalysis TritonAMDGPUTransforms TritonAMDGPUDialectToLLVM ) ``` -------------------------------- ### Configure LLVM/MLIR TableGen for Triton Passes Source: https://github.com/triton-lang/triton/blob/main/include/triton/Conversion/TritonToTritonGPU/CMakeLists.txt This CMake snippet configures the build system for the Triton project. It specifies the LLVM target definition file, generates MLIR pass declarations for 'TritonToTritonGPU' from 'Passes.h.inc', and adds a public tablegen target for 'TritonConversionPassIncGen'. This setup is essential for defining and integrating custom compiler passes within the Triton framework. ```CMake set(LLVM_TARGET_DEFINITIONS Passes.td) mlir_tablegen(Passes.h.inc -gen-pass-decls --name TritonToTritonGPU) add_public_tablegen_target(TritonConversionPassIncGen) ``` -------------------------------- ### Define Triton Unit Test Configuration Source: https://github.com/triton-lang/triton/blob/main/unittest/Analysis/CMakeLists.txt This snippet illustrates how to configure a unit test named 'TestTritonAnalysis' within the Triton project's build system. It specifies 'UtilityTest.cpp' as the source file and links a comprehensive set of Triton libraries, including those for analysis, IR, GPU IR, and GPU transformations. ```Build Configuration add_triton_ut( NAME TestTritonAnalysis SRCS UtilityTest.cpp LIBS TritonAnalysis TritonIR TritonGPUIR TritonGPUTransforms TritonNvidiaGPUTransforms ) ``` -------------------------------- ### Define TritonAMDUtils Library Source: https://github.com/triton-lang/triton/blob/main/third_party/amd/lib/Dialect/TritonAMDGPU/Utility/CMakeLists.txt Configures the `TritonAMDUtils` library, including `CommonUtils.cpp` as a source file and linking against `MLIRLLVMDialect`, `TritonIR`, and `TritonGPUIR`. This specifies the dependencies required for the library to compile and link correctly. ```Build Configuration add_triton_library(TritonAMDUtils CommonUtils.cpp LINK_LIBS PUBLIC MLIRLLVMDialect TritonIR TritonGPUIR ) ``` -------------------------------- ### Triton Project CMake Build Configuration Source: https://github.com/triton-lang/triton/blob/main/lib/Dialect/Triton/IR/CMakeLists.txt This CMake snippet defines the build process for the Triton project. It sets LLVM target definitions, generates MLIR tablegen files for canonicalization, and adds the TritonIR library. The library includes various C++ source files, depends on TritonTableGen and TritonCanonicalizeIncGen, and links against several MLIR dialects. ```CMake set(LLVM_TARGET_DEFINITIONS Canonicalize.td) mlir_tablegen(TritonCanonicalize.inc -gen-rewriters) add_public_tablegen_target(TritonCanonicalizeIncGen) add_triton_library(TritonIR Dialect.cpp DiscardableAttributes.cpp Ops.cpp Traits.cpp Types.cpp OpInterfaces.cpp Utility.cpp DEPENDS TritonTableGen TritonCanonicalizeIncGen LINK_LIBS PUBLIC MLIRIR MLIRArithDialect MLIRMathDialect MLIRSCFDialect ) ``` -------------------------------- ### Configure Triton LLVM/MLIR Build Definitions Source: https://github.com/triton-lang/triton/blob/main/include/triton/Dialect/Triton/Transforms/CMakeLists.txt This CMake snippet configures essential build definitions for the Triton project. It sets the LLVM target definitions, generates MLIR pass declarations using `mlir_tablegen`, and adds a public tablegen target for Triton transformations. This setup is crucial for integrating Triton's custom passes and transformations within an LLVM/MLIR-based compilation framework. ```CMake set(LLVM_TARGET_DEFINITIONS Passes.td) mlir_tablegen(Passes.h.inc -gen-pass-decls -name Triton) add_public_tablegen_target(TritonTransformsIncGen) ``` -------------------------------- ### Configure Triton Project with CMake Source: https://github.com/triton-lang/triton/blob/main/third_party/proton/dialect/CMakeLists.txt This CMake snippet configures the Triton project by including necessary directories and adding subdirectories for `include` and `lib`. It also conditionally builds a Python plugin named `TritonProton` if the `TRITON_BUILD_PYTHON_MODULE` flag is enabled, linking it with `ProtonIR`, `Python3::Module`, and `pybind11::headers`. ```CMake include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include) include_directories(${CMAKE_CURRENT_BINARY_DIR}/include) add_subdirectory(include) add_subdirectory(lib) if(TRITON_BUILD_PYTHON_MODULE) add_triton_plugin(TritonProton ${CMAKE_CURRENT_SOURCE_DIR}/triton_proton.cc) target_link_libraries(TritonProton PRIVATE ProtonIR Python3::Module pybind11::headers) endif() ``` -------------------------------- ### Configure Triton GPU Instrumentation Shared Libraries with CMake Source: https://github.com/triton-lang/triton/blob/main/test/lib/Instrumentation/CMakeLists.txt This CMake script defines and builds shared libraries for GPU instrumentation passes. It iterates through a list of passes, sets their source files, compiles them as shared libraries, links them with LLVM, and configures their output directory and symbol visibility. It handles platform-specific linking and ensures the `llvmGetPassPluginInfo` symbol is visible for plugin loading. ```CMake set(GPU_INSTRUMENTATION_PASSES GPUInstrumentationTestLib ) set(GPUInstrumentationTestLib_SOURCES GPUHello.cpp ) foreach( plugin ${GPU_INSTRUMENTATION_PASSES} ) add_library( ${plugin} SHARED ${${plugin}_SOURCES} ) target_link_libraries( ${plugin} PRIVATE LLVMCore "$<$:-undefined dynamic_lookup>" ) # 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}/../instrumentation") endif(DEFINED CMAKE_LIBRARY_OUTPUT_DIRECTORY) # This is set to -fvisibility=hidden in the top level CMake file # which causes the llvmGetPassPluginInfo symbol to be hidden and # an "entry point not found" error. Reset it just for this target if(NOT MSVC) target_compile_options(${plugin} PRIVATE -fvisibility=default) endif() endforeach() ``` -------------------------------- ### CMake: Including Core Subdirectories in Triton Build Source: https://github.com/triton-lang/triton/blob/main/third_party/amd/lib/CMakeLists.txt This CMake snippet demonstrates the inclusion of essential subdirectories into the Triton project's build system. Each `add_subdirectory` command incorporates a specific component, such as 'Analysis' for static analysis tools, 'Dialect' for custom language dialects, and various 'TritonAMDGPU' modules for LLVM transformations targeting AMD GPUs. This setup is crucial for modularizing the compiler's development and build process. ```CMake add_subdirectory(Analysis) add_subdirectory(Dialect) add_subdirectory(TritonAMDGPUToLLVM) add_subdirectory(TritonAMDGPUDialectToLLVM) add_subdirectory(TritonAMDGPUTransforms) ``` -------------------------------- ### Define NVGPUIR Library Build Configuration Source: https://github.com/triton-lang/triton/blob/main/third_party/nvidia/lib/Dialect/NVGPU/IR/CMakeLists.txt This snippet configures the `NVGPUIR` library within the Triton build system. It specifies `Dialect.cpp` as a source file, declares dependencies on `NVGPUTableGen` and `NVGPUAttrDefsIncGen`, and links `MLIRLLVMDialect` publicly. ```Triton Build System add_triton_library(NVGPUIR Dialect.cpp DEPENDS NVGPUTableGen NVGPUAttrDefsIncGen LINK_LIBS PUBLIC MLIRLLVMDialect ) ``` -------------------------------- ### Profile a Python Function with Proton Source: https://github.com/triton-lang/triton/blob/main/third_party/proton/README.md Demonstrates how to use `proton.profile` to profile a simple Python function, specifying the profile data name and the context method ('python' or 'shadow') for annotating GPU kernels. ```python import triton.profiler as proton # name: The path to the profile data # context: The method used to annotate the context of each GPU kernel. Currently, "shadow" and "python" are supported. session_id = proton.profile(func, name="profile_name", context="python")(args) ```