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_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.histogram (#10923, #11194, #11210, #11390, #11424, #11531, #11641, #11403, #11218, #11751)

Backend & Compiler

LLVM Updates

• LLVM pin: Updated LLVM to llvm/llvm-project@b010a18d. AMD code generation now uses its own LLVM pin; see AMD/HIP Backend (#10739, #11163, #11485)

Multi-CTA / Multicast / TMA

• Multi-CTA kernels: Fixed missing cluster synchronization and proxy fences. Multi-CTA kernels now support local gather from subslices, tl.histogram, and ordered atomics in warp-specialized kernels (#10656, #10735, #10843, #10975, #10559, #11775, #10914)
• Unaligned TMA: Fixed wrong results when per-stage TMA buffers are not 128-byte aligned (#10898)

Code Generation & Analysis

• Atomic ordering: Atomics with acquire, release, or acq_rel semantics now insert CTA-wide barriers so the ordering covers the whole block. Read-modify-write and CAS atomics default to acq_rel; pass sem="relaxed" to avoid the barriers (#10816, #11022)
• Barrier insertion: Barrier insertion was rewritten as a standalone pass that adds missing barriers and fences and removes redundant ones (#11389, #11391, #11396, #11401, #11411, #11414, #11431, #11462, #11468, #11479, #10087, #11324, #11464, #11499, #11597)
• Batched matmul: On NVIDIA, batched MMAv2 dots are up to 40x faster (#11822)
• Miscompile fixes: Fixed wrong results in nested-loop flattening, reductions, scans, tl.histogram, gathers, axis analysis, and optimizations that moved or duplicated loads (#11521, #11688, #11617, #11621, #10744, #11186, #11269, #11682, #11842, #11637, #11981, #11990, #11943, #12011, #11767, #10838, #11546, #11548, #11337, #11314, #11906, #11975, #11240)

Sanitizers & Debugging

• ConSan: ConSan catches more races and reports fewer false positives. It compiles and runs about 2.5x faster on one reported workload (#10668, #10736, #10784, #10864, #10948, #11045, #11010, #11162, #11216, #11244, #11310, #11200, #11202, #11175, #11268, #11436, #11459, #11626, #11852, #11857, #11138, #11446, #11447, #11448, #11449, #11450, #11451, #11452, #11642)
• GSan: GSan supports programmatic dependent launch, CUDA graphs built with explicit dependencies, and multi-CTA kernels, and runs up to 3.8x faster on MoE (#10677, #10944, #11020, #11135, #11199, #11303, #11172, #11187, #11297, #11351, #12023)
• FpSan: FpSan supports tl.clamp, pure inline assembly, and Gluon embedding ops. Narrowing casts now fold discarded bits by default; set TRITON_FPSAN_HOMOMORPHIC_CASTS=1 to keep truncation (#10707, #10627, #10908, #10834, #11005, #11083, #11101, #11159, #11170, #11489, #11627, #11665, #11716, #11974)

Extensions & Tooling

• Loading extensions: Extensions load at runtime through extend_with and extend_dialects_with, so a package can register itself on import (#10775)
• Out-of-tree layouts: Out-of-tree MMA and shared encodings work through MmaEncodingTrait and SharedEncodingTrait (#11025, #11764)

AMD/HIP Backend

gfx1250 / CDNA 5

• GFX1250-strict: Added initial support for the gfx1250-strict target (#11945, #12134)
• Tensor Data Movement: Added Gluon tdm.async_load_fused and pipelining of in-loop descriptor stores and scatters. Fixed rank-reducing loads, multi-CTA gather/scatter, and size-1 block dimensions. Non-zero padding (padding_option="nan") now raises an error (#10628, #10564, #11077, #11833, #11963, #12331, #12441, #10870)
• Scaled conversions: Gluon adds scaled_downcast to MXFP4/MXFP8 on CDNA4 and CDNA5, and scaled_upcast accepts compact scales (#11377, #11442, #10979)

Other Targets

• RDNA: Added RDNA 4m (gfx117x) targets and RDNA3 buffer atomics. bf16 multiplication now rounds to nearest even, and dot_scaled rounds its bf16 downcast toward zero (#10878, #11474, #11535, #11227, #11629)

Code Generation and Runtime

• Separately pinned LLVM: AMD code generation uses its own LLVM pin (triton-lang/llvm-project@6bc4aaf6). For a custom build, set TRITON_AMD_CODEGEN_PATH to a prebuilt libtriton_amd_codegen library (#11485, #11717, #11745, #11912, #11925, #11916, #12003, #12016)
• TheRock: Triton finds and prefers ROCm from TheRock Python wheels (rocm_sdk) unless TRITON_* paths are set. It raises an error if a different HIP runtime is already loaded (#11028, #11398, #11492)

Bug Fixes

• Incorrect results: Fixed NaN handling in FP8 FNUZ upcasts and CDNA3 TF32 tl.dot, CDNA barriers emitted on RDNA, buffer-atomic fence ordering, and miscompiles in range analysis and canonicalization (#10747, #11783, #10906, #11081, #11313, #11246, #11239)
• Hangs and races: tl.load(..., volatile=True) keeps its volatility, which fixes hanging polling loops. Also fixed LDS races, wrong shared-memory addresses on RDNA3/3.5, out-of-bounds XF32 MFMA access, and a deadlock when loading HIP modules in parallel (#11223, #11169, #11765, #11719, #11651, #11566, #11727, #11265)
• Compile cache: Changing an external library now invalidates the compile cache (#11884)

NVIDIA Backend

Matrix Instructions

• MMA performance: TCGen5 MMA reuses A operands across N tiles, and FP4×FP4 kernels on SM103 are 6-15% faster (#11345, #11438)
• SM120: Added batched (rank-3) scaled-dot scale layouts (#11262)
• MMA fixes: Fixed wrong results for FP4 MMA with mixed packing, Blackwell MMAs that share an accumulator across pipeline stages, small-K MMAv2 dots with upcast operands, and Hopper warp_group_dot_wait with num_warps > 4 (#11151, #11245, #10988, #11691, #11056)

TMA / Cluster Launch Control

• CLC for Triton kernels: On SM100 and newer, the clc=True launch option turns a one-tile-per-program kernel into a persistent Cluster Launch Control loop, including with warp_specialize=True (#11124, #11116)
• TMA fixes: Fixed TMA on shared-memory subviews, multicast im2col, device-side tensor-map updates, and loops that create tensor descriptors (#11064, #11136, #11480, #10281)

Code Generation

• ptxas: Hopper (sm_90) and newer use the bundled CUDA 13 ptxas, now 13.4.59. Older targets keep the CUDA 12 ptxas (#12035, #11669)
• Conversions: FP16, BF16, FP8, and FP4 conversions use native instructions on Hopper and Blackwell. Fixed rounding in round-to-nearest-even conversions to float8e5 and float8e4nv (#11340, #11647, #11699, #11924)
• Reductions: Warp reductions use native FP32 min/max on GB300 and Rubin, and redux.sync for 64-bit integers (#11299, #11302, #11537, #11788, #11813, #11816)
• 256-bit memory access: Gluon kernels on Blackwell can use 256-bit global loads and stores (#10861, #11766)

Bug Fixes

• Synchronization and tensor memory: Fixed a cluster_barrier phase race, misplaced barriers around mbarrier initialization and TMEM deallocation, TMEM allocations overwritten while in use, and a dropped MMA operand initializer (#11334, #11498, #11525, #11636, #11887)

Gluon & Layout Improvements

New Features

• Per-thread inline assembly: Added gl.inline_asm, which runs inline assembly per thread and can take shared-memory and tensor-memory descriptors (#11623)
• Cache policies: Loads, stores, and TMA operations accept a cache_policy argument. On NVIDIA it controls L1 and L2 policies, fractional eviction, and L2 prefetch size (#11482, #12000)
• Offline compilation: triton.compile accepts Gluon IR (.glir), so kernels can be compiled for a target without a GPU (#11628, #11806, #11198)
• Memory descriptor views: Shared-memory descriptors support more subslice and reshape patterns, and tensor-memory descriptors can be sliced along M (#10997, #11006, #10856, #11034, #10855, #10862)
• Language additions: Gluon exports atomic_load, atomic_store, atomic_poll, cumsum, and cumprod (#11591, #10723, #12007)
• AMD APIs: Added FP4 scaled_upcast layout helpers, load_shared_fp4_repacked, ttgl.amd.hint.sched_barrier, and a cd_regclass option on mfma/mfma_scaled (#11470, #11657, #11149, #11932, #11792)

Improvements and Fixes

• CDNA5 module: cdna5 is the canonical Gluon module for gfx1250, and gfx1250 remains as an alias (#11183, #11213)
• Fixes: Functions specialized on NVMMASharedLayout or TensorMemoryLinearLayout arguments no longer reuse the wrong compiled body. AMD scaled_dot rejects scale layouts that gave wrong results (#11094, #11201, #11443)

Kernels & Benchmarks

Low-Precision Matmul

• New input combinations: triton_kernels.matmul supports unscaled FP8 E4M3 with FP16 or BF16, and DequantScaleRoundingMode.ROUND_NEAREST for NVFP4 with E4M3 scales (#11652, #11668)
• MX/FP4 layout conversion: Hopper and Blackwell FP4 value and scale conversions are often 4-40x faster and use half or less of the peak memory (#11360, #11387, #11495, #11496, #11563, #11565)
• Ragged matmul scales: Ragged-M matmuls with swizzled Blackwell activation scales must pass the same RaggedTensorMetadata object to BlackwellActMXScaleLayout and a_ragged_metadata, or they raise AssertionError. With gather_indx, the scales must already be in gathered row order (#11928)

Performance and Other Updates

• Cluster Launch Control: triton_kernels matmul adds a clc constraint for SM100+ persistent kernels (#11176)
• Performance: New RDNA3/RDNA4 MXFP4 configs are up to 7x faster for small-M MoE shapes, and top-k no longer recompiles for each batch size (#10810, #10875)
• Top-k: topk supports non-power-of-two widths and widths above 32 (#11445)
• Fixes: Fixed garbage values in padded Blackwell activation scales. Bias epilogues use explicit FMA, so results do not depend on the FP-fusion default (#11030, #12032)
• Empty MX tensors: Converting empty MX tensors to Hopper, Blackwell, CDNA4, or GFX1250 layouts requires PyTorch 2.10 or newer (#10623)

Proton Profiling

• Async and cross-warp events: Added pl.allocate_event, pl.start_event, and pl.end_event to time asynchronous work across loops, functions, and warps (#11122, #11142)
• Trace alignment: CPU and GPU timestamps are aligned, so Chrome traces can be merged with PyTorch profiler or viztracer traces (#10954, #10949, #10969)
• ROCm: The rocprofiler backend supports PC sampling and HIP graph profiling, and no longer deadlocks or aborts at exit (#11040, #10506, #10989, #10999, #11009, #11008, #11559, #11824)
• Instrumentation: Added intra-kernel clocks on gfx1250, and fixed events in warp-specialized partitions (#10813, #11523)

Testing & CI

• Hardware coverage: Added ROCm 10.1 on gfx950, CDNA5 tests on the rocjitsu emulator, and sm120 TMA GEMM tests (#11487, #11843, #12031, #12043, #10839)

Build & Infrastructure

• Python versions and wheels: Triton supports Python 3.11 through 3.15. Python 3.12 and newer use a single cp312-abi3 wheel, and source builds enable it with TRITON_STABLE_ABI (#12006, #10958)
• nanobind: The Python bindings moved from pybind11 to nanobind (nanobind==2.10.2), which slightly lowers launch overhead (#10283)
• Toolchain: Triton builds as C++20, and the bundled libdevice is updated to CUDA 12.9 (#11920, #11189)
• Extension builds: With TRITON_EXT_ENABLED=1, the wheel includes Triton's C++ headers and TableGen definitions (#10847)

Documentation

• Language reference: Documented do_bench* return units, bf16x3/bf16x6 tl.dot precision, and type promotion for tl.sum, division, and modulo (#10729, #10798, #10893, #10968, #10883)
• Release process: RELEASE.md documents the release-branch process and ties Triton feature releases to PyTorch minor releases (#11713, #10932)
• Tutorials: Fixed wrong results in the FP8 fused-attention tutorial (#11461)

Deprecations

• Floating-point fusion default: Triton 3.10 will turn off implicit floating-point fusion (multiply-add to FMA) by default. In 3.9, fusion stays on, and Triton warns once per process when it compiles a kernel that does not set enable_fp_fusion. Set TRITON_DEFAULT_FP_FUSION=0 to adopt the 3.10 default now, or =1 to keep fusion and silence the warning. Kernels whose numerics depend on fusion should call tl.fma explicitly. The new TRITON_FORCE_DISABLE_FP_FUSION=1 disables implicit fusion even for kernels that request it (#11270, #12387, #12433)
• can_reorder=True on tl.cat, tl.reshape, and tl.ravel: The argument still works in 3.9 but warns once per process. Triton 3.10 will remove it, and these operations will always preserve element order. Drop can_reorder=True now, and use dim= with tl.cat (#11293, #12393, #12433)
• tl.softmax(keep_dims=...): keep_dims has no effect, because softmax always keeps the input shape. Passing a value warns; remove the argument (#11430)
• Gluon amd.gfx1250 module: amd.gfx1250 re-exports amd.cdna5. Import cdna5 in new code (#11183)

Breaking Changes

Language and Runtime

• Block pointers removed: tl.make_block_ptr raises NotImplementedError, tl.advance was removed, and tl.load/tl.store no longer accept boundary_check or padding_option. Switch to tensor descriptors (tl.make_tensor_descriptor or the host-side TensorDescriptor), which provide bounds checking and zero padding (#10833)
• Keyword-only arguments: cache_modifier and eviction_policy on tl.load/tl.store (and tensor.store), volatile on tl.load, and keep_dims/ieee_rounding on tl.softmax are now keyword-only. Passing them positionally fails to compile (#10833, #11418, #11430)
• Division and modulo dtype: / and % return float32 when an operand is a float narrower than 32 bits, including FP8 with Python scalars. Cast the result to keep the old dtype (#10883, #11121)
• Scalar and bfloat16 promotion: A Python scalar no longer widens the result of tl.minimum, tl.maximum, comparisons, tl.clamp, tl.fma, or binary math functions, so tl.maximum(x_fp16, 0.0) returns float16. tl.minimum, tl.maximum, tl.clamp, tl.cumsum, tl.cumprod, and tl.min/tl.max with return_indices=True no longer upcast bfloat16 to float32. To keep float32, cast the inputs or pass dtype=tl.float32 to cumsum/cumprod. tl.cumprod gains dtype and, like cumsum, now upcasts integers narrower than 32 bits by default (#11472, #11913, #11473)
• Tensor math method dtype checks: Method forms such as x.exp() and x.sqrt() apply the same dtype checks as the tl.* functions and fail to compile on fp16 inputs. Cast first, for example x.to(tl.float32).exp() (#11218)
• Stricter validation: triton.autotune raises ValueError if key, reset_to_zero, or restore_value names something that is not a kernel argument. Indexing with more : slices than the tensor has dimensions fails to compile. tl.make_tensor_descriptor rejects constant strides that are not 16-byte aligned, which previously gave wrong results (#11602, #10429, #10931)
• Pointer address spaces: tl.pointer_type(..., address_space=...) takes a string ("global", "descriptor", or "constant"), and integers no longer work. tl.const pointer arguments use "constant". In IR, use named forms such as !tt.ptr<T, "descriptor">; integer forms such as !tt.ptr<T, 0> are rejected (#11385, #11294)

Platforms and Toolchain

• Python 3.10 dropped: Triton requires Python 3.11 or newer (>=3.11,<3.16) and no longer ships Python 3.10 wheels (#12006)
• RDNA1/RDNA2 removed: The AMD backend no longer supports gfx10 targets (gfx101x, gfx103x). Use Triton 3.8 for those GPUs (#12042)
• Hopper ptxas selection: sm_90 now compiles with the CUDA 13 ptxas, whose path is set with TRITON_PTXAS_BLACKWELL_PATH. TRITON_PTXAS_PATH applies only below sm_90, so move any Hopper override (#12035)
• Building from source: nanobind replaced pybind11, so PYBIND11_SYSPATH is no longer used, and out-of-tree code built against Triton's binding headers must port to nanobind. AMD code generation ignores LLVM_SYSPATH; set TRITON_AMD_CODEGEN_PATH to a prebuilt libtriton_amd_codegen library (#10283, #11485)

Gluon

• Async-copy aliases removed: async_copy_global_to_shared, async_copy_global_to_shared_im2col, and async_copy_shared_to_global raise an error. Use async_load, async_load_im2col, and async_store (#11112)
• Multi-CTA primitives: Removed cluster.arrive, cluster.wait, and mbarrier.fence_init_release_cluster. Use cluster.barrier(); the compiler inserts the mbarrier-init fence (#10960)
• Descriptor reinterpret: _reinterpret on shared- and tensor-memory descriptors is renamed to reinterpret (#11131)
• Blocked-layout linear form: Converting a blocked layout to a linear layout drops zero (broadcast) register bases. Compare layouts by equivalence if your code relies on exact bases (#11541)
• Barrier layouts: mbarrier and CLC operands must use the default MBarrierLayout without subviews, apart from indexing along the buffering dimension (#11501)

Out-of-Tree Backends, Extensions, and IR

• Extension loading: The Python bindings no longer read TRITON_PLUGIN_PATHS; use extend_with/extend_dialects_with. make install and the CMake install components were removed; build the wheel with TRITON_EXT_ENABLED=1. Plugin passes and custom-op builders get add_/create_ prefixes automatically (#10775, #11157, #10847, #10737, #10754)
• Cache-policy IR: tt.load, tt.store, tt.descriptor_load, nvws.descriptor_load, the async and TMA global-to-local copies, and the AMD buffer and async-copy ops replaced their cache/evict attributes with one cachePolicy attribute. The C++ builders and the Python ir.builder load, store, and descriptor-load methods now take a cache-policy attribute (see builder.get_cache_policy). The language-level cache_modifier/eviction_policy arguments are unchanged (#11482, #12000)
• IR operation changes: Removed ttng.cluster_arrive/ttng.cluster_wait (use ttng.cluster_barrier). AMD buffer ops print the base-pointer type (: !tt.ptr<f32> -> tensor<...>), and amdg.local_load_packed_tranposed is renamed amdg.local_load_packed_transposed (#11051, #11379, #11145)
• add_plan_cta removed: Use add_assign_cga_layouts and add_optimize_cta_locality (#10325)
• Passes moved out of add_to_llvmir: Custom pipelines must add barrier insertion before add_to_llvmir, with nvidia.passes.ttgpuir.add_membar(pm, capability, ptx_version) or amd.passes.ttgpuir.add_membar(pm, arch). NVIDIA add_to_llvmir also no longer takes a ConSan flag. Add nvidia.passes.ttnvgpuir.add_tmem_wait_insertion, passes.ttgpuir.add_concurrency_sanitizer, nvidia.passes.ttnvgpuir.add_cluster_barrier_mbar_allocator, and passes.ttgpuir.add_allocate_global_scratch_memory explicitly (#11597)
• AMD target triple: amd.TARGET_TRIPLE was replaced by amd.get_target_triple(arch), and amd.attach_target_triple takes an architecture argument (#11532)
• AMD LLVM pass plugins: With LLVM_PASS_PLUGIN_PATH set, the AMDGPU attributor runs before plugin passes. Instrumentation plugins that add reads of implicit kernel inputs must remove amdgpu-no-* attributes from the functions they instrument (#10849)

Contributors

This release includes contributions from engineers at:

• Meta
• AMD
• NVIDIA
• OpenAI
• Intel
• Google
• And many individual contributors

Special thanks to all contributors who submitted bug reports, feature requests, and code improvements!