- 在线推理里要写一个融合算子(比如 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 要留调试余量
专家点评
上手难度:进阶更新于 2026-11-11用 Python 写 GPU kernel 并编译到 NVPTX/AMDGPU 的编译器,顶替手写 CUDA 和 CUTLASS 模板那一层
适用场景
踩坑
- 从源码构建时 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
pip install tritonREADME 未明确列出最低显存要求;支持 NVIDIA(NVPTX)与 AMD(AMDGPU)GPU,也可编译到 CPU;设置 TRITON_INTERPRET=1 可在无 GPU 的解释器模式下运行。
项目说明
由 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。
- 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
- Backend & Compiler
- AMD/HIP Backend
- NVIDIA Backend
- Gluon & Layout Improvements
- Kernels & Benchmarks
- Proton Profiling
- Testing & CI
- Build & Infrastructure
- Documentation
- Deprecations
- Breaking Changes
- Contributors
Dialect & Frontend
New Features
- Atomic loads and stores: Added
tl.atomic_loadandtl.atomic_store, also available in Gluon and the interpreter. Loads default toacquireand stores default torelease(#11591, #11650, #11676, #11821) tl.atomic_poll: Addedtl.atomic_poll, which waits until the value at a pointer equals an expected value and returnsFalseon 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.clampaccepts integer tensors (#12005)
Improvements
- Earlier errors: Mismatched tuple unpacking, invalid
tl.topkkvalues,tl.catshape mismatches, and non-32-bittl.histograminputs 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.softmaxwithdimother than 0, integer stores to boolean pointers, list-comprehensioniffilters, negative-zero constants, and stale compilations when a captured constexpr shares a global's name (#11409, #11853, #10888, #11426, #11224) - E8M0 scale decoding: E8M0
0x00now decodes to2^-127and0xFFto 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=1now matches compiled kernels for integer overflow,tl.fma,tl.umulhi, and `tl.histo
- 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
- Backend & Compiler
- AMD/HIP Backend
- NVIDIA Backend
- Gluon & Layout Improvements
- Kernels & Benchmarks
- Proton Profiling
- Testing & CI
- Build & Infrastructure
- Documentation
- Breaking Changes
- Contributors
Dialect & Frontend
New Features
- Aggregate types:
@triton.aggregateand@gluon.aggregateare now public APIs. Aggregates support inherited fields, default values, generated constructors, immutable instances, andaggregate_replace()(#10095, #9572) tl.topk: Added adescendingargument. Setdescending=Falseto 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-pointatomic_minnow returns a value-typed result (#10074, #10485) - Interpreter NaN handling:
argmin,argmax,minimum,maximum, andclampnow 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_waitnow accepts aread_onlyargument. The default remainsTrue; useread_only=Falsewhen the store must reach global memory before a release operation (#10415, [#10419](https://github.com/triton-lang/trit
- 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)
- v3.7.0Triton 3.7.0 Release Notes2026-05-07 · 2026-05-07 22:19 UTC
查看更新说明
Table of Contents
- Dialect & Frontend
- Backend & Compiler
- AMD/HIP Backend
- NVIDIA Backend
- Gluon & Layout Improvements
- Kernels & Benchmarks
- Proton Profiling
- Testing & CI
- Build & Infrastructure
- Documentation
- Breaking Changes
- Contributors
Dialect & Frontend
New Features
tl.squeeze/tl.unsqueeze: Addedtl.squeezeandtl.unsqueezeoperations 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
constexprvalues from JIT-compiled code (#8785) get_int_attrfor Out-of-Tree Walk: Addedget_int_attrtoOperationto support out-of-tree IR walks (#8892)- Optional Device Arg to
preload: Added optional device argument topreloadand guardrails for cross-target preload (#8951, #8952, #9234) tl.cat(can_reorder=False): Added a non-reordering variant oftl.catwith 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.shapefor FP4 Padded: Fixeddesc.shapevalues 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_resultfor Methods: Fixedmust_use_resultcheck for methods (#8902)_semanticDefault to None: Defaulted_semanticparameter to None (#8909)make_tensor_descriptorError Typo: Fixed typo inmake_tensor_descriptorerror message (#8912)tl.catDeterminism: Madetl.catdeterministic via permute+reshape+join, then reverted (#9312, #8854, #8878)- Deprecation Warning for
make_block_ptr: Emitted a deprecation warning whenmake_block_ptris used (#9667)
Improvements
- Frontend Performance: Pre-computed
inspect.signaturefor builtins, lazily computed tuple type names, avoidedfind_paths_ifandinspect.getclosurevars, removed outdatedcatch_warningsblocks — 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 fortcgen05.mma(#9071)- TMA Index Translation: Moved TMA index translation from mid-end to lowering (#9082)
- **`tcgen05.mma
- 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
- Backend & Compiler
- AMD/HIP Backend
- NVIDIA Backend
- Gluon & Layout Improvements
- Kernels & Benchmarks
- Proton Profiling
- Testing & CI
- Build & Infrastructure
- Documentation
- Breaking Changes
Dialect & Frontend
New Features
- Multidimensional Batch Support (#8542): Added support for multidimensional batches in
tl.transandtl.dotoperations - 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_scaledoperations - 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_scaledHandling (#8658): Fixed missing handling for None acc indot_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): Optimizedtl.cdivfor 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.cpLowering (#8225): Implemented generic lowering fortcgen05.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
用户点评
登录用户的真实使用记录,可按版本和硬件筛选还没有人点评。第一个说出真实体验的人最有价值。
Documentation | Nightly Wheels |
|---|---|
Triton Conference 2026
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
-
Find the version of LLVM that Triton builds against. Check the
llvm_hashfield incmake/llvm-info.jsonto 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.
-
git checkoutLLVM at this revision. Optionally, make additional modifications to LLVM. -
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 -
Grab a snack, this will take a while.
-
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=trueas an environment variable to use clang and lld. lld in particular results in faster builds. -
Set
TRITON_BUILD_WITH_CCACHE=trueto build with ccache. -
Set
TRITON_HOME=/some/pathto change the location of the.tritondirectory 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_JOBSenvironment variable (to thepip install -e .command) to limit the number of jobs. -
Pass
--no-build-isolationtopip installto make nop builds faster. Without this, every invocation ofpip installuses a different symlink to cmake, and this forces ninja to rebuild most of the.afiles. -
The build system creates a
compile_commands.jsonfile 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.jsonfile 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 + Pon Mac, orShift + Ctrl + Pon Windows/Linux) and openC/C++: Edit Configurations (UI). - Open "Advanced Settings" and paste the full path to
compile_commands.jsoninto the "Compile Commands" textbox.
- Do a local build. Run command
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=1dumps the IR before every MLIR pass Triton runs, for all kernels. UseMLIR_ENABLE_DUMP=kernelNameto dump for a specific kernel only.- Triton cache can interfere with the dump. In cases where
MLIR_ENABLE_DUMP=1does not work, try cleaning your triton cache:rm -r ~/.triton/cache/*.
- Triton cache can interfere with the dump. In cases where
-
MLIR_DUMP_PATHspecifies whereMLIR_ENABLE_DUMPwill dump to. If unset will dump to stderr. -
LLVM_IR_ENABLE_DUMP=1dumps 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=1uses the Triton interpreter instead of running on the GPU. You can insert Python breakpoints in your kernel code! -
TRITON_ENABLE_LLVM_DEBUG=1passes-debugto LLVM, printing a lot of debugging information to stdout. If this is too noisy, run with justTRITON_LLVM_DEBUG_ONLYinstead 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'soptstandalone, perhaps passing-debug-only=fooon the command line.
- An alternative way to reduce output noisiness is running with
-
TRITON_LLVM_DEBUG_ONLY=<comma-separated>is the equivalent of LLVM's-debug-onlycommand-line option. This limits the LLVM debug output to specific pass or component names (which are specified using#define DEBUG_TYPEthroughout LLVM and Triton) in order to allow the debug output to be less noisy.TRITON_LLVM_DEBUG_ONLYallows for one or more comma separated values to be specified (egTRITON_LLVM_DEBUG_ONLY="tritongpu-remove-layout-conversions"orTRITON_LLVM_DEBUG_ONLY="tritongpu-remove-layout-conversions,regalloc"). -
TRITON_ENABLE_ASAN=1invokes 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=1prints out the best autotuning config and total time spent for each kernel after autotuning is complete. -
DISABLE_LLVM_OPTwill 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 isDISABLE_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=1forces to compile kernels regardless of cache hit. -
MLIR_ENABLE_TIMINGdumps the timing information for each MLIR pass. -
LLVM_ENABLE_TIMINGdumps the timing information for each LLVM pass. -
TRITON_DEFAULT_FP_FUSIONoverrides the default behavior of disabling fp fusion (mul+add->fma). -
TRITON_ALLOW_FP_FUSION=1allows kernels to control implicit fp fusion throughenable_fp_fusion. Without this opt-in, implicit fusion is disabled. Explicit fused operations such astl.fmaare 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,operationsenables remarks and IR operations, whileMLIR_ENABLE_DIAGNOSTICS=warnings,stacktracesenables warnings with stacktraces. By default, only errors are shown. Settingwarningsincludes errors and warnings;remarksincludes errors, warnings, and remarks. -
MLIR_ENABLE_REMARKis deprecated. Please useMLIR_ENABLE_DIAGNOSTICS=remarks. -
TRITON_KERNEL_DUMPenables the dumping of the IR from each compilation stage and the final ptx/amdgcn. -
TRITON_DUMP_DIRspecifies the directory to save the dumped IR and ptx/amdgcn whenTRITON_KERNEL_DUMPis set to 1. -
TRITON_KERNEL_OVERRIDEenables the override of the compiled kernel with a user-specified IR/ptx/amdgcn at the beginning of each compilation stage. -
TRITON_OVERRIDE_DIRspecifies the directory from which to load the IR/ptx/amdgcn files whenTRITON_KERNEL_OVERRIDEis set to 1. -
TRITON_F32_DEFAULTsets the default input precision oftl.dotwhen using 32-bit floats, which can be eitherieee,tf32, ortf32x3. -
TRITON_FRONT_END_DEBUGGING=1disables exception wrapping when an error occurs in the compiler frontend, allowing the full stack trace to be seen. -
TRITON_DISABLE_LINE_INFO=1removes all line information from the module. -
PTXAS_OPTIONSpasses additional command-line options to the PTX assemblerptxas(only on NVIDIA). -
LLVM_EXTRACT_DI_LOCAL_VARIABLESemit 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.