triton

triton-langMIT活跃维护 有专家点评

Triton 是一门面向 GPU 编程的类 Python 语言及其编译器,让开发者不用写 CUDA 就能写出接近手写性能的高效内核。

Development repository for the Triton language and compiler

v3.9.0发布于 2026-10-27 18:56 UTC20,551star +18 / 7 天3,396fork同步于 2026-11-10 23:25 UTC

专家点评

上手难度:进阶更新于 2026-11-11

用 Python 写 GPU kernel 并编译到 NVPTX/AMDGPU 的编译器,顶替手写 CUDA 和 CUTLASS 模板那一层

适用场景

  • 在线推理里要写一个融合算子(比如 attention 变体、fused norm+激活),用 CUDA 写要处理 warp 调度和 shared memory 布局,用 Triton 几十行能出第一版,代价是极端调优后仍可能落后手写 kernel
  • 离线批量跑算子搜索:autotune 把 block size、num_warps 这些配置扫一遍,比手写多套模板再人工挑快得多,但搜索本身要占 GPU 时间,得和训练任务错峰
  • 给 PyTorch 补一个自定义 backward:Triton kernel 能直接挂进 autograd,不用写 C++ extension 和 setup.py,迭代周期从分钟级降到秒级
  • 在 AMD 卡上做算子开发:同一份 kernel 走 AMDGPU 后端,比从 HIP 重写一遍省事,但后端成熟度不如 NVPTX,复杂 kernel 要留调试余量

踩坑

  • 从源码构建时 LLVM 版本被 llvm-info.json 的 hash 锁死,想用系统自带 LLVM 或自己打的补丁版本会直接编译失败,只有 pip 安装的预编译包能绕开
  • 调试时 dump 出来的 IR/PTX 受 ~/.triton/cache 影响,改了代码但看到旧结果,得先清缓存再复现,否则会误判成编译器 bug
  • 无 GPU 环境只能靠 TRITON_INTERPRET=1 跑解释器模式验证语义,性能特征和真实后端不一致,数值对不代表上卡就对;完整测试套件仍然要 GPU
  • Windows 支持在 issue 里被反复提出(191 反应),说明官方路径长期不在 Windows 上,跨平台团队要提前规划 WSL 或 Linux 容器

同类对比

要极致性能且愿意维护模板代码选 CUTLASS/CUDA;要跨厂商可移植的 kernel 语言可以看 SYCL 或 Mojo 生态,但工具链成熟度和 PyTorch 集成度目前不如 Triton

怎么上手

由 README 提炼,执行前请确认
pip install triton
硬件要求按 README 说明

README 未明确列出最低显存要求;支持 NVIDIA(NVPTX)与 AMD(AMDGPU)GPU,也可编译到 CPU;设置 TRITON_INTERPRET=1 可在无 GPU 的解释器模式下运行。

纯 CPU 可运行需要 CUDA支持 ROCm / AMD
下载:源码 zip·GitHub Releases(安装包)本站不镜像安装包

项目说明

由 README 自动翻译提炼

用 Python 写高效 GPU 深度学习原语的编译器,比 CUDA 更好写、比现有 DSL 更灵活

Triton 是一门面向 GPU 编程的类 Python 语言及其编译器,让开发者不用写 CUDA 就能写出接近手写性能的高效内核。

核心能力

  • 用类 Python 语法编写自定义 GPU kernel
  • 基于 MLIR/LLVM 编译到 GPU 与 CPU 后端
  • 支持 NVPTX 与 AMDGPU 目标
  • 内置 autotune 自动搜索最优配置
  • 提供解释器模式,无 GPU 也能调试 kernel

容易踩的坑

从源码构建会下载预编译 LLVM,且必须匹配 cmake/llvm-info.json 中的 llvm_hash,任意 LLVM 版本无法编译;运行完整测试需 GPU,无 GPU 时用 make test-nogpu;调试 dump 受 Triton 缓存干扰,需清理 ~/.triton/cache。

内容由模型根据 README 生成,可能不完整或不准确,以官方文档为准。看原文 README

最近发版

5 个版本v3.9.0
  1. v3.9.0Triton 3.9.0 Release Notes最新正式版2026-10-27 · 2026-10-27 18:56 UTC

    Triton 3.9.0 Release Notes

    Table of Contents


    Dialect & Frontend

    New Features

    • Atomic loads and stores: Added tl.atomic_load and tl.atomic_store, also available in Gluon and the interpreter. Loads default to acquire and stores default to release (#11591, #11650, #11676, #11821)
    • tl.atomic_poll: Added tl.atomic_poll, which waits until the value at a pointer equals an expected value and returns False on timeout (#10723, #10755, #11540)
    • Approximate division: Added tl.fdiv(..., approx=True) for fast FP32 division (#11900)
    • Language syntax: Kernels support chained comparisons such as 0 <= x <= 8, keyword-only parameters, tuple unpacking in list comprehensions, and operator methods on aggregate types (#12002, #11418, #11076, #10990)
    • Integer tl.clamp: tl.clamp accepts integer tensors (#12005)

    Improvements

    • Earlier errors: Mismatched tuple unpacking, invalid tl.topk k values, tl.cat shape mismatches, and non-32-bit tl.histogram inputs now fail at compile time (#11053, #11074, #11353, #10891)
    • Hints through helpers: Alignment and contiguity hints on values returned from inlined helpers are now kept (#11150)
    • Compile time: Deeply nested loops compile about 10x faster at nesting depth 8 (#11369)

    Bug Fixes

    • Wrong results: Fixed tl.softmax with dim other than 0, integer stores to boolean pointers, list-comprehension if filters, negative-zero constants, and stale compilations when a captured constexpr shares a global's name (#11409, #11853, #10888, #11426, #11224)
    • E8M0 scale decoding: E8M0 0x00 now decodes to 2^-127 and 0xFF to NaN (previously zero and infinity), which changes results for kernels that use those values (#11624)
    • Forked processes: A forked child process no longer unloads GPU modules that belong to its parent (#11128)
    • Interpreter: TRITON_INTERPRET=1 now matches compiled kernels for integer overflow, tl.fma, tl.umulhi, and `tl.histo
  2. v3.8.0Triton 3.8.0 Release Notes2026-08-28 · 2026-08-28 18:25 UTC
    查看更新说明

    Triton 3.8.0 Release Notes

    Table of Contents


    Dialect & Frontend

    New Features

    • Aggregate types: @triton.aggregate and @gluon.aggregate are now public APIs. Aggregates support inherited fields, default values, generated constructors, immutable instances, and aggregate_replace() (#10095, #9572)
    • tl.topk: Added a descending argument. Set descending=False to return the smallest values (#9355)
    • Tensor descriptors: Tensor descriptors can be passed inside tuple-valued kernel arguments (#9422)
    • Interpreter: Added support for tl.dot_scaled (#10311)

    Improvements

    • Autotuning listener: Added a listener that reports the selected configuration, measured timings, tuning duration, and disk-cache status (#10125)
    • JIT cache keys: Dependency cache keys are now generated deterministically (#10494)

    Bug Fixes

    • Division and atomics: tl.fdiv(..., ieee_rounding=True) now emits IEEE-rounded division, and floating-point atomic_min now returns a value-typed result (#10074, #10485)
    • Interpreter NaN handling: argmin, argmax, minimum, maximum, and clamp now match compiled behavior more closely when inputs contain NaNs (#10298, #10333, #10699)
    • Block-pointer padding: Fixed zero padding for block-pointer loads (#11252)
    • Python 3.14 annotations: Updated annotation handling for PEP 649, including aggregate field discovery (#10581)

    Backend & Compiler

    LLVM Updates

    • Correctness fixes: Updated the pinned LLVM revision with fixes for a GFX950 BF16 miscompilation and SLP-vectorizer issues (#10719, #11356)

    Multi-CTA / Multicast / TMA

    • Generic multi-CTA support: Extended multi-CTA support to layout conversion, reductions, local gather/scatter, TMA gather/scatter, and multicast, with corresponding updates to barrier insertion and memory analysis (#9317, #9221, #9977, #10472, #9318, #9615)
    • TMA store waits: tma.store_wait now accepts a read_only argument. The default remains True; use read_only=False when the store must reach global memory before a release operation (#10415, [#10419](https://github.com/triton-lang/trit
  3. v3.7.1Triton 3.7.1 Release Notes2026-06-18 · 2026-06-18 14:38 UTC
    查看更新说明

    Triton 3.7.1 is a patch release on top of 3.7.0. It fixes the following 2 regressions and contains no new features or API changes.

    Regression fixes

    • Add async read dependencies to FenceAsync — a missing fence between a shared-memory store (st.shared) and an async copy_local_to_global could let the async copy read shared memory before the store completed, producing incorrect results. FenceAsync now inserts the required fence. (triton-lang/triton#9610)
    • [InstCombine] Shrink added constant using LHS known zeros — fixes an LLVM InstCombine miscompilation where add simplification used known-zero bits only from the RHS, mishandling the symmetric case where the LHS has known zeros and the low bits are unused. Picked up by Triton through its pinned LLVM. (llvm/llvm-project#174380)
  4. v3.7.0Triton 3.7.0 Release Notes2026-05-07 · 2026-05-07 22:19 UTC
    查看更新说明

    Table of Contents


    Dialect & Frontend

    New Features

    • tl.squeeze / tl.unsqueeze: Added tl.squeeze and tl.unsqueeze operations to the standard library (#8924)
    • Scaled BMM: Added support for scaled batched matmul in the frontend (#9000)
    • FP8 Constants: Frontend can now create FP8 constants directly (#8882)
    • Returning Constexpr from JIT: Functions can return constexpr values from JIT-compiled code (#8785)
    • get_int_attr for Out-of-Tree Walk: Added get_int_attr to Operation to support out-of-tree IR walks (#8892)
    • Optional Device Arg to preload: Added optional device argument to preload and guardrails for cross-target preload (#8951, #8952, #9234)
    • tl.cat(can_reorder=False): Added a non-reordering variant of tl.cat with broadcast support (#9312, #9163)
    • Round f32→tf32 in Descriptor: Added option to round f32 to tf32 inside tensor descriptors (#9295)
    • Plugin Hooks & Out-of-Tree Dialects: Added support for out-of-tree TTIR/TTGIR passes and Triton Dialect Plugins, with example documentation (#8401, #8523, #8815)

    Bug Fixes

    • desc.shape for FP4 Padded: Fixed desc.shape values for fp4-padded tensor descriptors (#9012)
    • Setting Attr on Constexpr Argument: Fixed setting attributes on constexpr arguments (#9053)
    • Named Tuples in Constexpr Functions: Preserved named tuples through constexpr_functions (#8876)
    • must_use_result for Methods: Fixed must_use_result check for methods (#8902)
    • _semantic Default to None: Defaulted _semantic parameter to None (#8909)
    • make_tensor_descriptor Error Typo: Fixed typo in make_tensor_descriptor error message (#8912)
    • tl.cat Determinism: Made tl.cat deterministic via permute+reshape+join, then reverted (#9312, #8854, #8878)
    • Deprecation Warning for make_block_ptr: Emitted a deprecation warning when make_block_ptr is used (#9667)

    Improvements

    • Frontend Performance: Pre-computed inspect.signature for builtins, lazily computed tuple type names, avoided find_paths_if and inspect.getclosurevars, removed outdated catch_warnings blocks — all to reduce JIT overhead (#8843, #8844, #8846, #8845, #8881)
    • Revert Deep Copy on Scope Entry: Removed deep copy when entering a new scope (#8832)
    • Default 32-bit Dot Precision Change: Briefly changed default 32-bit dot precision to TF32x3, then reverted (#9080, #9090)
    • Tutorial Updates (#8565, #8853, #8982)
    • Interpreter Cleanups: Typing and efficiency cleanups in the interpreter (#9072)

    Backend & Compiler

    LLVM Updates

    • LLVM Bumps: Multiple LLVM uprevs through the cycle, with one bump reverted on the release branch for stability (#8766, #8840, #8919, #8987, #9264, #9333, #9431, #9942)
    • llvm-head Merge: Merged changes from llvm-head (#8842)
    • Infinite Rewrite Loop in Latest LLVM: Fixed an infinite rewrite loop introduced by a newer LLVM revision (#9249)

    2CTA / Multicast / TMA

    • 2CTA Mode End-to-End: Gluon multi-cta + 2CTA support, M=64 2CTA mode, removed unnecessary synchronization in 2CTA MMA, and proper TMEM deallocation timing (#8684, #8874, #8922, #8986)
    • TMA + Multicast: Backend support for TMA with multicast (#9005)
    • tcgen05.mma + Multicast: Added multicast support for tcgen05.mma (#9071)
    • TMA Index Translation: Moved TMA index translation from mid-end to lowering (#9082)
    • **`tcgen05.mma
  5. v3.6.0Triton 3.6.0 release2026-01-21 · 2026-01-21 23:27 UTC
    查看更新说明

    Triton 3.6 Release Notes

    Table of Contents


    Dialect & Frontend

    New Features

    • Multidimensional Batch Support (#8542): Added support for multidimensional batches in tl.trans and tl.dot operations
    • Ragged TMA Atomic Add (#8238): Added atomic add support for ragged TMA operations
    • Integer Range Utility (#8753): Exposed an integer-range utility from AMD range analysis code for broader use
    • Constexpr Through Min/Max (#8733): Propagate constexpr through builtin min/max functions (BC-breaking)
    • Scales Dimension Checks (#8564): Added dimension checks for scales in dot_scaled operations
    • Loop Bounds Verification (#8243): Added verification that loop bounds are scalars

    Bug Fixes

    • For Loop Induction Variable (#8750): Fixed modification of for loop induction variable handling
    • Store Broadcasting (#8661): Fixed broadcasting issues in store operations
    • Missing dot_scaled Handling (#8658): Fixed missing handling for None acc in dot_scaled
    • AugAssign Line Information (#8703): Attached proper line number information to AugAssign nodes
    • Starred Argument Handling (#8686): Made starred argument handling more robust
    • Saved Exception Cloning (#8115): Fixed clone of saved exception before raising
    • Tuple Mangling (#8060): Fixed mangling for tuples in JIT compilation

    Improvements

    • Optimized tl.cdiv (#8669): Optimized tl.cdiv for common case of 32-bit divisors
    • Un-deprecated min/max (#8734): Un-deprecated min/max on scalar tensors
    • Warmup in KernelInterface (#8757): Moved warmup functionality into KernelInterface
    • Verification with Diagnostics (#8074): Frontend always verifies with diagnostics enabled
    • Constexpr with do_not_specialize Error (#8275): Added error when constexpr is combined with do_not_specialize
    • Deprecated ast.Num Replacement (#8698): Replaced usage of deprecated ast.Num

    Backend & Compiler

    LLVM Updates

    • LLVM Bump (#8299): Bumped to llvm/llvm-project@f6ded0be897e
    • LLVM Head Merge (#8612): Merged back changes from llvm-head with updated APIs
    • Inliner Import (#8152): Import inliner in triton-opt for better optimization

    Code Generation

    • CTALayout as LinearLayout (#8770): Made CTALayout an honest-to-goodness LinearLayout for better representation
    • Shared Layout Rank Check (#8772): Added check that Shared layouts have rank equal to the tensor or one less
    • Backward Propagation Fix Point (#8776): Run remove backward prop until fix point for correctness
    • Generic tcgen05.cp Lowering (#8225): Implemented generic lowering for tcgen05.cp
    • Generic Matrix Descriptors (#8321): Implemented shmem matrix descriptors generically
    • LinearSharedEncoding Support (#8116): Added support for LinearSharedEncoding
    • BF16x3 Trick (#7592): Implemented BF16x3 trick for improved performance
    • Padded Shared Linear Remapping (#7929): Added linear remapping to padded shared layout

    Optimizations

    • Compilation Time Improvement (#8689): Improved compilation time in constant sanitizer pass
    • AxisInfo Loop Removal (#8679): Removed unnecessary loop over roots in AxisInfo analysis
    • Constant Analysis (#8502): Improved constant analysis in AxisInfo
    • Combinatory Explosion Prevention (#8477): Prevented combinatory explosion when checking tmem_load uses
    • Layout Conversion Vectorization (#8655): Fixed vectorization for conver

用户点评

登录用户的真实使用记录,可按版本和硬件筛选
共 0 条

    还没有人点评。第一个说出真实体验的人最有价值。

    README 原文

    未翻译
    DocumentationNightly Wheels
    DocumentationWheels

    Triton Conference 2026

    Triton Banner

    Triton Developer Conference registration is open: https://tritonconference2026.splashthat.com/

    Folks, please join us at the Triton Developer Conference on October 19th at the San Jose Convention Center, just before PTC at the same venue! We are finalizing an agenda with keynotes, technical talks, lightning talks, and posters.

    Registration and attendance are free. Breakfast, lunch, as well as a happy hour organized by our Nvidia colleagues, are included! Learn something new, share your experience, mingle with colleagues, and make new friends!

    To allow us to prepare correctly, please register and reserve your spot at your earliest convenience: https://tritonconference2026.splashthat.com/ You will see more info, including a draft agenda there soon, as well.

    For previous conference materials, see:

    Triton

    This is the development repository of Triton, a language and compiler for writing highly efficient custom Deep-Learning primitives. The aim of Triton is to provide an open-source environment to write fast code at higher productivity than CUDA, but also with higher flexibility than other existing DSLs.

    The foundations of this project are described in the following MAPL2019 publication: Triton: An Intermediate Language and Compiler for Tiled Neural Network Computations. Please consider citing this work if you use Triton!

    The official documentation contains installation instructions and tutorials. See also these third-party Triton puzzles, which can all be run using the Triton interpreter -- no GPU required.

    Quick Installation

    You can install the latest stable release of Triton from pip:

    pip install triton
    

    Binary wheels are available for CPython 3.10-3.14.

    Install from source

    git clone https://github.com/triton-lang/triton.git
    cd triton
    
    pip install -r python/requirements.txt # build-time dependencies
    pip install -e .
    

    Or with a virtualenv:

    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 .
    

    Building with a custom LLVM

    Triton uses LLVM to generate code for GPUs and CPUs. Normally, the Triton build downloads a prebuilt LLVM, but you can also build and use LLVM from source.

    LLVM does not have a stable API, so the Triton build will not work at an arbitrary LLVM version.

    For convenience, use the following command to build LLVM and install Triton with the custom LLVM:

    make dev-install-llvm                 # Core LLVM (default)
    make dev-install-llvm LLVM_BUILD=amd  # AMD LLVM only
    make dev-install-llvm LLVM_BUILD=all
    
    1. Find the version of LLVM that Triton builds against. Check the llvm_hash field in cmake/llvm-info.json to see the current version. For example, if it says: "llvm_hash": "49af6502c6dcb4a7f7520178bd14df396f78240c"

      This means that the version of Triton you have builds against LLVM 49af6502.

    2. git checkout LLVM at this revision. Optionally, make additional modifications to LLVM.

    3. Build LLVM. For example, you might run:

      $ 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
      
    4. Grab a snack, this will take a while.

    5. Build Triton as above, but set the following environment variables:

      # Modify as appropriate to point to your LLVM build.
      $ export LLVM_BUILD_DIR=$HOME/llvm-project/build
      
      $ cd <triton install>
      $ LLVM_INCLUDE_DIRS=$LLVM_BUILD_DIR/include \
        LLVM_LIBRARY_DIR=$LLVM_BUILD_DIR/lib \
        LLVM_SYSPATH=$LLVM_BUILD_DIR \
        pip install -e .
      

    Tips for building

    • Set TRITON_BUILD_WITH_CLANG_LLD=true as an environment variable to use clang and lld. lld in particular results in faster builds.

    • Set TRITON_BUILD_WITH_CCACHE=true to build with ccache.

    • Set TRITON_HOME=/some/path to change the location of the .triton directory where Triton's cache is located and downloads are stored during the build. By default, this is the user's home directory. It can be changed anytime.

    • If you're running out of memory when building Triton, specify the MAX_JOBS environment variable (to the pip install -e . command) to limit the number of jobs.

    • Pass --no-build-isolation to pip install to make nop builds faster. Without this, every invocation of pip install uses a different symlink to cmake, and this forces ninja to rebuild most of the .a files.

    • The build system creates a compile_commands.json file under the Triton repo directory. This file is used by VSCode IntelliSense and clangd to provide code completion and other features for C++ code.

      If IntelliSense does not work, you can try the following steps:

      • Do a local build. Run command pip install -e ..
      • Get the full path to the compile_commands.json file produced by the build: find ./build -name 'compile_commands.json' | xargs readlink -f. You might get a full path similar to /Users/{username}/triton/build/cmake.macosx-11.1-arm64-cpython-3.12/compile_commands.json.
      • In VSCode, install the C/C++ extension, then open the command palette (Shift + Command + P on Mac, or Shift + Ctrl + P on Windows/Linux) and open C/C++: Edit Configurations (UI).
      • Open "Advanced Settings" and paste the full path to compile_commands.json into the "Compile Commands" textbox.

    Running tests

    There currently isn't a turnkey way to run all the Triton tests, but you can follow the following recipe:

    # One-time setup.  Note this will reinstall local Triton because torch
    # overwrites it with the public version.
    $ make dev-install
    
    # To run all tests (requires a GPU)
    $ make test
    
    # Or, to run tests without a gpu
    $ make test-nogpu
    

    Tips for hacking

    For detailed instructions on how to debug Triton's frontend, please refer to this tutorial. The following includes additional tips for hacking on Triton's backend.

    Configuration knobs

    See python/triton/knobs.py for the full list of configuration knobs. You can set those knobs directly in python or use environment variables to control them. Below are some of the environment variables you can specify (see knobs.py for the full list):

    • MLIR_ENABLE_DUMP=1 dumps the IR before every MLIR pass Triton runs, for all kernels. Use MLIR_ENABLE_DUMP=kernelName to dump for a specific kernel only.

      • Triton cache can interfere with the dump. In cases where MLIR_ENABLE_DUMP=1 does not work, try cleaning your triton cache: rm -r ~/.triton/cache/*.
    • MLIR_DUMP_PATH specifies where MLIR_ENABLE_DUMP will dump to. If unset will dump to stderr.

    • LLVM_IR_ENABLE_DUMP=1 dumps the IR before every pass run over the LLVM IR.

    • TRITON_REPRODUCER_PATH=<reproducer_path> will generate an MLIR reproducer file at <reproducer_path> before each MLIR compiler stage. If any of the stages fail, <reproducer_path> will be a local MLIR reproducer captured right before the failing pass.

    • TRITON_INTERPRET=1 uses the Triton interpreter instead of running on the GPU. You can insert Python breakpoints in your kernel code!

    • TRITON_ENABLE_LLVM_DEBUG=1 passes -debug to LLVM, printing a lot of debugging information to stdout. If this is too noisy, run with just TRITON_LLVM_DEBUG_ONLY instead to limit the output.

      • An alternative way to reduce output noisiness is running with LLVM_IR_ENABLE_DUMP=1, extract the IR before the LLVM pass of interest, and then run LLVM's opt standalone, perhaps passing -debug-only=foo on the command line.
    • TRITON_LLVM_DEBUG_ONLY=<comma-separated> is the equivalent of LLVM's -debug-only command-line option. This limits the LLVM debug output to specific pass or component names (which are specified using #define DEBUG_TYPE throughout LLVM and Triton) in order to allow the debug output to be less noisy. TRITON_LLVM_DEBUG_ONLY allows for one or more comma separated values to be specified (eg TRITON_LLVM_DEBUG_ONLY="tritongpu-remove-layout-conversions" or TRITON_LLVM_DEBUG_ONLY="tritongpu-remove-layout-conversions,regalloc").

    • TRITON_ENABLE_ASAN=1 invokes the LLVM address sanitizer for memory leak and out of bounds access detection. Currently only supported on the AMD backend. This must be run using the ASAN libraries documented here.

      • When enabling the address sanitizer it is recommended to disable various memory caching strategies both within the ROCm stack and PyTorch. This will give the address sanitizer the best chance at finding the memory fault where it originates. See this test for more details.
    • USE_IR_LOC={ttir,ttgir} reparses the IR such that the location information will be the line number of the IR file with that particular extension, instead of line number of the python file. This can provide a direct mapping from the IR to llir/ptx. When used with performance tools, it can provide a breakdown on IR instructions.

    • TRITON_PRINT_AUTOTUNING=1 prints out the best autotuning config and total time spent for each kernel after autotuning is complete.

    • DISABLE_LLVM_OPT will disable llvm optimizations for make_llir and make_ptx if its value is true when parsing as Bool. Otherwise, it will be parsed as a list of flags to disable llvm optimizations. One usage case is DISABLE_LLVM_OPT="disable-lsr" Loop strength reduction is known to cause up to 10% performance changes for certain kernels with register pressure.

    • TRITON_ALWAYS_COMPILE=1 forces to compile kernels regardless of cache hit.

    • MLIR_ENABLE_TIMING dumps the timing information for each MLIR pass.

    • LLVM_ENABLE_TIMING dumps the timing information for each LLVM pass.

    • TRITON_DEFAULT_FP_FUSION overrides the default behavior of disabling fp fusion (mul+add->fma).

    • TRITON_ALLOW_FP_FUSION=1 allows kernels to control implicit fp fusion through enable_fp_fusion. Without this opt-in, implicit fusion is disabled. Explicit fused operations such as tl.fma are unchanged.

    • MLIR_ENABLE_DIAGNOSTICS=<comma-separated> controls diagnostic emission in MLIR. Options are: warnings, remarks, stacktraces, operations. Use comma-separated values to customize output. For example, MLIR_ENABLE_DIAGNOSTICS=remarks,operations enables remarks and IR operations, while MLIR_ENABLE_DIAGNOSTICS=warnings,stacktraces enables warnings with stacktraces. By default, only errors are shown. Setting warnings includes errors and warnings; remarks includes errors, warnings, and remarks.

    • MLIR_ENABLE_REMARK is deprecated. Please use MLIR_ENABLE_DIAGNOSTICS=remarks.

    • TRITON_KERNEL_DUMP enables the dumping of the IR from each compilation stage and the final ptx/amdgcn.

    • TRITON_DUMP_DIR specifies the directory to save the dumped IR and ptx/amdgcn when TRITON_KERNEL_DUMP is set to 1.

    • TRITON_KERNEL_OVERRIDE enables the override of the compiled kernel with a user-specified IR/ptx/amdgcn at the beginning of each compilation stage.

    • TRITON_OVERRIDE_DIR specifies the directory from which to load the IR/ptx/amdgcn files when TRITON_KERNEL_OVERRIDE is set to 1.

    • TRITON_F32_DEFAULT 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 disables exception wrapping when an error occurs in the compiler frontend, allowing the full stack trace to be seen.

    • TRITON_DISABLE_LINE_INFO=1 removes all line information from the module.

    • PTXAS_OPTIONS passes additional command-line options to the PTX assembler ptxas (only on NVIDIA).

    • LLVM_EXTRACT_DI_LOCAL_VARIABLES emit full debug info, allowing for eval of values in gpu debuggers (ie cuda-gdb, rocm-gdb etc)

    • TRITON_DEFAULT_BACKEND=<backend> optionally sets the default backend used by Triton when constructing the active driver (i.e., triton.runtime.driver.active).

    [!NOTE] Some of these environment variables don't have a knob in knobs.py-- those are only relevant to the C++ layer(s), hence they don't exist in the python layer.

    Kernel Override Steps

    export TRITON_ALWAYS_COMPILE=1
    export TRITON_KERNEL_DUMP=1
    export TRITON_DUMP_DIR=<dump_dir>
    export TRITON_KERNEL_OVERRIDE=1
    export TRITON_OVERRIDE_DIR=<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/<kernel_hash> 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
    

    Compiler Pipeline Inspection Steps To introspect the pipeline add_stages, before running your kernels, simply set the add_stages_inspection_hook like so:

    def inspect_stages(_self, stages, options, language, capability):
        # inspect or modify add_stages here
    triton.knobs.runtime.add_stages_inspection_hook = inspect_stages
    

    Examples of how to use this for out of tree plugin passes is here

    Changelog

    Version 2.0 is out! New features include:

    • Many, many bug fixes
    • Performance improvements
    • Backend rewritten to use MLIR
    • Support for kernels that contain back-to-back matmuls (e.g., flash attention)

    Contributing

    Community contributions are more than welcome, whether it be to fix bugs or to add new features at github. For more detailed instructions, please visit our contributor's guide.

    Compatibility

    Supported Platforms:

    • Linux

    Supported Hardware:

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

    Development Container (Dev Container)

    Dev Containers for the Triton project are available from the triton-dev-containers repository.

    Key Benefits:

    • Consistency: All developers can work with the same development environment, ensuring uniform behavior across different systems.
    • Isolation: The container prevents potential conflicts with software installed on your local machine.
    • Portability: Easily share the development environment with team members, minimizing onboarding time and setup issues.

    How to Use the Dev Container:

    For detailed instructions on how to use the dev containers, please see the dev container user guide.