环境变量与编译选项
本文汇总 Triton-Ascend 中可由开发者显式控制的行为开关,包括运行前设置的环境变量,以及编译期通过 triton.Config 或 kernel launch meta-parameter 传入的 NPU 编译选项。
环境变量
环境变量用法示例
环境变量需在运行 Python 程序前设置,例如:
export TRITON_DEBUG=1
python run_kernel.py
环境变量参考表
环境变量配置参考下表:
类别 |
环境变量 |
默认值 |
功能说明 |
配置说明 |
变更声明 |
|---|---|---|---|---|---|
调试与日志 |
TRITON_DEBUG |
0 或未设置 |
启用 Triton 的调试输出功能,用于在运行时打印详细的调试信息。这对于排查编译或执行阶段的问题非常有用。 当设置为 1 时,Triton 会输出更多关于编译过程、内核生成和执行的信息。 某些实现中可能支持更细粒度的调试级别(如 2, 3 等),具体取决于 Triton 的版本和实现。 |
0:不启用DEBUG |
|
调试与日志 |
MLIR_ENABLE_DUMP |
0 或未设置 |
在每次 MLIR 优化前转储所有内核的 IR。使用 |
0:不转储 |
Triton 缓存可能干扰转储。如果 |
调试与日志 |
LLVM_IR_ENABLE_DUMP |
0 或未设置 |
在每次 LLVM IR 优化前转储 IR。 |
0:不转储 |
|
调试与日志 |
TRITON_REPRODUCER_PATH |
未设置 |
在每个 MLIR 编译阶段前生成 MLIR 复现文件。如果某阶段失败, |
<reproducer_path>:保存路径 |
|
调试与日志 |
TRITON_INTERPRET |
0 或未设置 |
使用 Triton 解释器而非 GPU 运行,支持在核函数代码中插入 Python 断点 |
0:不支持断点 |
|
调试与日志 |
TRITON_ENABLE_LLVM_DEBUG |
0 或未设置 |
向LLVM 传递 |
0:不传递 |
另一种减少输出干扰的方法是:先设置 |
调试与日志 |
TRITON_LLVM_DEBUG_ONLY |
未设置 |
功能等同于 LLVM 的 |
逗号分隔值:通道或组件名称 |
|
调试与日志 |
USE_IR_LOC |
0 或未设置 |
控制是否在生成的中间表示(IR)中包含位置信息(如文件名、行号等)。这些信息对调试很有帮助,但可能会增加生成的IR的大小。设置为1,会重新解析中间表示(IR),将位置信息映射为具有特定扩展名的IR文件行号(而非Python源文件行号)。这能建立从IR到LLVM IR/PTX的直接映射关系。配合性能分析工具使用时,可实现对IR指令的细粒度性能剖析。 |
0:不包含位置信息 |
|
调试与日志 |
TRITON_PRINT_AUTOTUNING |
0 或未设置 |
在自动调优完成后,输出每个内核的最佳配置及总耗时。 |
0:不输出 |
|
调试与日志 |
MLIR_ENABLE_REMARK |
0 或未设置 |
启用MLIR 编译过程中的备注信息输出,包括以备注形式输出的性能警告。 |
0:不启用 |
|
调试与日志 |
TRITON_KERNEL_DUMP |
0 或未设置 |
启用或禁用 Triton 内核的转储功能,当启用时,Triton 会将生成的内核代码(各编译阶段IR及最终PTX)保存到指定目录。 |
0:不启用 |
|
调试与日志 |
TRITON_DUMP_DIR |
当前工作目录或未设置 |
指定 Triton 内核转储文件的保存目录。当 |
"path":保存路径 |
|
调试与日志 |
TRITON_DEVICE_PRINT |
0 或未设置 |
当设置为 |
0:不启动 |
每个线程的GM缓冲区最大为16KB,超限内容将被丢弃。该值目前固定,后续将通过环境变量调整。 |
调试与日志 |
TRITON_MEMORY_DISPLAY |
0 或未设置 |
控制是否生成内存使用情况的 json 文件。当 |
0:不启用 |
|
编译控制 |
TRITON_ALWAYS_COMPILE |
0 或未设置 |
控制 Triton 是否每次运行都强制重新编译内核,而不是使用已有的缓存版本。 默认情况下,Triton 会对已经编译过的内核进行缓存(基于参数和配置),以提高性能。 设置为 1 后,Triton 将忽略缓存并每次都重新编译内核,这在调试或测试新编译器特性时非常有用。 |
0:不启用 |
|
编译控制 |
DISABLE_LLVM_OPT |
0 或未设置 |
当设置为 1 时,可以禁用 LLVM 编译过程中的优化步骤(make_llir和make_ptx的LLVM优化)。当设置为字符串,解析为要禁用的LLVM优化标志列表。例如使用 |
0:LLVM 的优化是启用状态 |
|
编译控制 |
MLIR_ENABLE_TIMING |
0 或未设置 |
启用或禁用 MLIR 编译过程中的时间统计功能。 |
0:不启用 |
|
编译控制 |
LLVM_ENABLE_TIMING |
0 或未设置 |
启用或禁用 LLVM 编译过程中的时间统计功能。 |
0:不启用 |
|
编译控制 |
TRITON_DEFAULT_FP_FUSION |
1 启用 |
控制是否默认启用浮点运算融合优化,覆盖默认的浮点运算融合行为(如mul+add->fma)。 |
0:不启用 |
|
编译控制 |
TRITON_KERNEL_OVERRIDE |
0 或未设置 |
启用或禁用 Triton 内核覆盖功能,允许在每个编译阶段开始时用用户指定的外部文件(IR/PTX等)覆盖默认生成的内核代码。 |
0:不启用 |
|
编译控制 |
TRITON_OVERRIDE_DIR |
当前工作目录或未设置 |
指定 Triton 内核覆盖文件的查找目录。当 |
"path":保存路径 |
|
编译控制 |
TRITON_ASCEND_COMPILE_SPEED_OPT |
0 或未设置 |
控制JIT编译器在发现内核编译失败后是否跳过后续编译阶段。设为 |
0:继续尝试 |
|
编译控制 |
TRITON_COMPILE_ONLY |
0 或未设置 |
remote_launch时使用,只编译不运行。 |
0:不启用 |
|
编译控制 |
TRITON_DISABLE_FFTS |
0 或未设置 |
是否禁用FFTS。 |
0:启用 |
|
编译控制 |
TRITON_DISABLE_PRECOMPILE |
0 或未设置 |
是否禁用预编译。 |
0:启用预编译 |
|
运行与调度 |
TRITON_ALL_BLOCKS_PARALLEL |
0 或未设置 |
启用或禁用自动根据物理核数优化逻辑核数,仅当逻辑核间可并行时方可启动。当逻辑核数大于物理核数时,启动该优化,则编译器自动调整逻辑核数量为物理核数,减少调度开销;启用后允许grid>65535。限制:triton kernel的逻辑必须对执行顺序不敏感才能开启该选项,否则可能会导致死锁。 |
0:不启用 |
|
运行与调度 |
TRITON_ENABLE_TASKQUEUE |
0 或未设置 |
是否开启task_queue。 |
0:不启用 |
|
运行与调度 |
TRITON_ENABLE_SANITIZER |
0 或未设置 |
是否启用 SANITIZER。 |
0:不启用 |
|
运行与调度 |
ENABLE_PRINT_UB_BITS |
0 或未设置 |
打开后可以获取当前UB占用量,给inductor使用。 |
0:不启用 |
|
其他 |
TRITON_BENCH_METHOD |
未设置 |
使用昇腾NPU时,将 |
"npu":切换为 |
|
其他 |
TRITON_REMOTE_RUN_CONFIG_PATH |
path |
指定远程运行的配置路径。 |
直接给定path |
编译选项
编译选项用于控制单个 Triton kernel 的编译策略,可通过 triton.Config、Autotune 参数或 kernel launch meta-parameter 传入。
编译选项用法示例
例如,可在 kernel launch 时直接传入 multibuffer:
import triton
import triton.language as tl
@triton.jit
def kernel(..., BLOCK_SIZE: tl.constexpr):
...
grid = (triton.cdiv(n_elements, 1024),)
kernel[grid](..., BLOCK_SIZE=1024, multibuffer=True)
编译选项参考表
编译选项配置参考下表:
类别 |
编译选项 |
默认值/可选值 |
功能说明 |
配置说明 |
|---|---|---|---|---|
通用流水 |
|
|
启用或禁用 ping-pong/double buffer 流水。默认开启。 |
|
CV 融合 |
|
|
启用或禁用自动绑定 sub-block。 |
|
CV 融合 |
|
|
启用或禁用自动 CV balance。 |
|
CV 融合/同步 |
|
|
启用或禁用 HIVM 同步求解器。 |
|
同步 |
|
|
Cube 搬出相关同步优化项。 |
|
同步 |
|
|
启用或禁用自动注入 barrier 同步。 |
|
同步 |
|
|
启用或禁用自动注入 block 同步。 |
|
多缓冲范围 |
|
|
限制自动 multi-buffer 只作用于 local buffer。 |
|
多缓冲范围 |
|
|
配置 local buffer 自动 multi-buffer 的 scope。 |
|
Workspace |
|
|
配置 workspace multi-buffer 档位。 |
|
CV 融合 tiling |
|
|
配置 Vector loop 的切分份数。 |
|
CV 融合 tiling |
|
|
配置 Cube loop 的切分份数。 |
|
CV 融合/同步 |
|
|
启用或禁用自动 block sync 注入。 |
|
运行流 |
|
|
指定 NPU stream。 |
launch meta-parameter |
编译 Pass |
|
版本相关 |
启用或禁用 linearization pass。 |
|
CV 融合/layout |
|
默认 |
启用或禁用 Vector 路径上的 ND 到 NZ 布局转换。 |
|
大 grid 优化 |
|
默认 |
启用或禁用 AutoBlockify pass。未设置 |
launch meta-parameter 或 |