[{"url":"https://api.github.com/repos/triton-lang/triton/releases/341462708","assets_url":"https://api.github.com/repos/triton-lang/triton/releases/341462708/assets","upload_url":"https://uploads.github.com/repos/triton-lang/triton/releases/341462708/assets{?name,label}","html_url":"https://github.com/triton-lang/triton/releases/tag/v3.7.1","id":341462708,"author":{"login":"atalman","id":7563158,"node_id":"MDQ6VXNlcjc1NjMxNTg=","avatar_url":"https://avatars.githubusercontent.com/u/7563158?v=4","gravatar_id":"","url":"https://api.github.com/users/atalman","html_url":"https://github.com/atalman","followers_url":"https://api.github.com/users/atalman/followers","following_url":"https://api.github.com/users/atalman/following{/other_user}","gists_url":"https://api.github.com/users/atalman/gists{/gist_id}","starred_url":"https://api.github.com/users/atalman/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/atalman/subscriptions","organizations_url":"https://api.github.com/users/atalman/orgs","repos_url":"https://api.github.com/users/atalman/repos","events_url":"https://api.github.com/users/atalman/events{/privacy}","received_events_url":"https://api.github.com/users/atalman/received_events","type":"User","user_view_type":"public","site_admin":false},"node_id":"RE_kwDOAWaHXs4UWk60","tag_name":"v3.7.1","target_commitish":"main","name":"Triton 3.7.1 Release Notes","draft":false,"immutable":false,"prerelease":false,"created_at":"2026-06-17T00:31:10Z","updated_at":"2026-06-18T14:39:14Z","published_at":"2026-06-18T14:38:51Z","assets":[{"url":"https://api.github.com/repos/triton-lang/triton/releases/assets/451294588","id":451294588,"node_id":"RA_kwDOAWaHXs4a5jV8","name":"triton-3.7.1.tar.gz","label":"","uploader":{"login":"github-actions[bot]","id":41898282,"node_id":"MDM6Qm90NDE4OTgyODI=","avatar_url":"https://avatars.githubusercontent.com/in/15368?v=4","gravatar_id":"","url":"https://api.github.com/users/github-actions%5Bbot%5D","html_url":"https://github.com/apps/github-actions","followers_url":"https://api.github.com/users/github-actions%5Bbot%5D/followers","following_url":"https://api.github.com/users/github-actions%5Bbot%5D/following{/other_user}","gists_url":"https://api.github.com/users/github-actions%5Bbot%5D/gists{/gist_id}","starred_url":"https://api.github.com/users/github-actions%5Bbot%5D/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/github-actions%5Bbot%5D/subscriptions","organizations_url":"https://api.github.com/users/github-actions%5Bbot%5D/orgs","repos_url":"https://api.github.com/users/github-actions%5Bbot%5D/repos","events_url":"https://api.github.com/users/github-actions%5Bbot%5D/events{/privacy}","received_events_url":"https://api.github.com/users/github-actions%5Bbot%5D/received_events","type":"Bot","user_view_type":"public","site_admin":false},"content_type":"application/gzip","state":"uploaded","size":6673650,"digest":"sha256:21cab714d4fc9579b728f4d597660c9598fbbd52c1154896c71d2d42f9b61626","download_count":332,"created_at":"2026-06-18T14:39:13Z","updated_at":"2026-06-18T14:39:14Z","browser_download_url":"https://github.com/triton-lang/triton/releases/download/v3.7.1/triton-3.7.1.tar.gz"}],"tarball_url":"https://api.github.com/repos/triton-lang/triton/tarball/v3.7.1","zipball_url":"https://api.github.com/repos/triton-lang/triton/zipball/v3.7.1","body":"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.\r\n\r\n## Regression fixes\r\n\r\n- 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)\r\n- [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)\r\n"},{"url":"https://api.github.com/repos/triton-lang/triton/releases/319213104","assets_url":"https://api.github.com/repos/triton-lang/triton/releases/319213104/assets","upload_url":"https://uploads.github.com/repos/triton-lang/triton/releases/319213104/assets{?name,label}","html_url":"https://github.com/triton-lang/triton/releases/tag/v3.7.0","id":319213104,"author":{"login":"atalman","id":7563158,"node_id":"MDQ6VXNlcjc1NjMxNTg=","avatar_url":"https://avatars.githubusercontent.com/u/7563158?v=4","gravatar_id":"","url":"https://api.github.com/users/atalman","html_url":"https://github.com/atalman","followers_url":"https://api.github.com/users/atalman/followers","following_url":"https://api.github.com/users/atalman/following{/other_user}","gists_url":"https://api.github.com/users/atalman/gists{/gist_id}","starred_url":"https://api.github.com/users/atalman/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/atalman/subscriptions","organizations_url":"https://api.github.com/users/atalman/orgs","repos_url":"https://api.github.com/users/atalman/repos","events_url":"https://api.github.com/users/atalman/events{/privacy}","received_events_url":"https://api.github.com/users/atalman/received_events","type":"User","user_view_type":"public","site_admin":false},"node_id":"RE_kwDOAWaHXs4TBs4w","tag_name":"v3.7.0","target_commitish":"main","name":"Triton 3.7.0 Release Notes","draft":false,"immutable":false,"prerelease":false,"created_at":"2026-05-07T16:50:13Z","updated_at":"2026-07-10T22:21:06Z","published_at":"2026-05-07T22:19:06Z","assets":[{"url":"https://api.github.com/repos/triton-lang/triton/releases/assets/426360199","id":426360199,"node_id":"RA_kwDOAWaHXs4Zab2H","name":"triton-3.7.0.tar.gz","label":null,"uploader":{"login":"atalman","id":7563158,"node_id":"MDQ6VXNlcjc1NjMxNTg=","avatar_url":"https://avatars.githubusercontent.com/u/7563158?v=4","gravatar_id":"","url":"https://api.github.com/users/atalman","html_url":"https://github.com/atalman","followers_url":"https://api.github.com/users/atalman/followers","following_url":"https://api.github.com/users/atalman/following{/other_user}","gists_url":"https://api.github.com/users/atalman/gists{/gist_id}","starred_url":"https://api.github.com/users/atalman/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/atalman/subscriptions","organizations_url":"https://api.github.com/users/atalman/orgs","repos_url":"https://api.github.com/users/atalman/repos","events_url":"https://api.github.com/users/atalman/events{/privacy}","received_events_url":"https://api.github.com/users/atalman/received_events","type":"User","user_view_type":"public","site_admin":false},"content_type":"application/x-gzip","state":"uploaded","size":6698909,"digest":"sha256:980397efada8ed2de303a95ad918e7454df13d7df5941840646e8004d4bd0462","download_count":324,"created_at":"2026-05-21T18:00:06Z","updated_at":"2026-05-21T18:00:09Z","browser_download_url":"https://github.com/triton-lang/triton/releases/download/v3.7.0/triton-3.7.0.tar.gz"}],"tarball_url":"https://api.github.com/repos/triton-lang/triton/tarball/v3.7.0","zipball_url":"https://api.github.com/repos/triton-lang/triton/zipball/v3.7.0","body":"## Table of Contents\r\n- [Dialect & Frontend](#dialect--frontend)\r\n- [Backend & Compiler](#backend--compiler)\r\n- [AMD/HIP Backend](#amdhip-backend)\r\n- [NVIDIA Backend](#nvidia-backend)\r\n- [Gluon & Layout Improvements](#gluon--layout-improvements)\r\n- [Kernels & Benchmarks](#kernels--benchmarks)\r\n- [Proton Profiling](#proton-profiling)\r\n- [Testing & CI](#testing--ci)\r\n- [Build & Infrastructure](#build--infrastructure)\r\n- [Documentation](#documentation)\r\n- [Breaking Changes](#breaking-changes)\r\n- [Contributors](#contributors)\r\n\r\n---\r\n\r\n## Dialect & Frontend\r\n\r\n### New Features\r\n- **`tl.squeeze` / `tl.unsqueeze`**: Added `tl.squeeze` and `tl.unsqueeze` operations to the standard library (#8924)\r\n- **Scaled BMM**: Added support for scaled batched matmul in the frontend (#9000)\r\n- **FP8 Constants**: Frontend can now create FP8 constants directly (#8882)\r\n- **Returning Constexpr from JIT**: Functions can return `constexpr` values from JIT-compiled code (#8785)\r\n- **`get_int_attr` for Out-of-Tree Walk**: Added `get_int_attr` to `Operation` to support out-of-tree IR walks (#8892)\r\n- **Optional Device Arg to `preload`**: Added optional device argument to `preload` and guardrails for cross-target preload (#8951, #8952, #9234)\r\n- **`tl.cat(can_reorder=False)`**: Added a non-reordering variant of `tl.cat` with broadcast support (#9312, #9163)\r\n- **Round f32→tf32 in Descriptor**: Added option to round f32 to tf32 inside tensor descriptors (#9295)\r\n- **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)\r\n\r\n### Bug Fixes\r\n- **`desc.shape` for FP4 Padded**: Fixed `desc.shape` values for fp4-padded tensor descriptors (#9012)\r\n- **Setting Attr on Constexpr Argument**: Fixed setting attributes on constexpr arguments (#9053)\r\n- **Named Tuples in Constexpr Functions**: Preserved named tuples through `constexpr_functions` (#8876)\r\n- **`must_use_result` for Methods**: Fixed `must_use_result` check for methods (#8902)\r\n- **`_semantic` Default to None**: Defaulted `_semantic` parameter to None (#8909)\r\n- **`make_tensor_descriptor` Error Typo**: Fixed typo in `make_tensor_descriptor` error message (#8912)\r\n- **`tl.cat` Determinism**: Made `tl.cat` deterministic via permute+reshape+join, then reverted (#9312, #8854, #8878)\r\n- **Deprecation Warning for `make_block_ptr`**: Emitted a deprecation warning when `make_block_ptr` is used (#9667)\r\n\r\n### Improvements\r\n- **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)\r\n- **Revert Deep Copy on Scope Entry**: Removed deep copy when entering a new scope (#8832)\r\n- **Default 32-bit Dot Precision Change**: Briefly changed default 32-bit dot precision to TF32x3, then reverted (#9080, #9090)\r\n- **Tutorial Updates** (#8565, #8853, #8982)\r\n- **Interpreter Cleanups**: Typing and efficiency cleanups in the interpreter (#9072)\r\n\r\n---\r\n\r\n## Backend & Compiler\r\n\r\n### LLVM Updates\r\n- **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)\r\n- **llvm-head Merge**: Merged changes from llvm-head (#8842)\r\n- **Infinite Rewrite Loop in Latest LLVM**: Fixed an infinite rewrite loop introduced by a newer LLVM revision (#9249)\r\n\r\n### 2CTA / Multicast / TMA\r\n- **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)\r\n- **TMA + Multicast**: Backend support for TMA with multicast (#9005)\r\n- **`tcgen05.mma` + Multicast**: Added multicast support for `tcgen05.mma` (#9071)\r\n- **TMA Index Translation**: Moved TMA index translation from mid-end to lowering (#9082)\r\n- **`tcgen05.mma` Verifier & Errors**: Throw a clear error instead of miscompiling very large `tcgen05.mma` along N (#8915)\r\n- **MMAv5 Illegal Instruction Fix**: Fixed illegal instruction in MMAv5 lowering (#8910)\r\n\r\n### Warp Specialization\r\n- **Nested Loops**: Nested-loop support in warp specialization (#8687)\r\n- **Partition Scheduling**: Improved partition scheduling pass; correct stage/cluster annotations for block-arg producers (#7312, #8883)\r\n- **WS Lowering Hardening**: Variable naming fix in `LowerAref`, per-partition `asyncOp` storage, explicit captures to `WarpSpecializePartitionsOp`, skip `InsertTmemAref` when WS isn't used (#8978, #9007, #9023, #9133, #9212)\r\n- **`RegionBranchInterface`**: Made `WarpSpecializePartitionsOp` implement `RegionBranchInterface` (#8799)\r\n- **Mixed TMA / non-TMA Loads**: Fixed AutoWS when mixing TMA and non-TMA loads (#9111)\r\n- **`aref.get` Filtering**: `aref.get` creation now filters results not in the scheduled loop (#9114)\r\n- **Multibuffering Acc Logic**: Improved multibuffering accumulator logic in WS (#8950)\r\n\r\n### Code Generation & Analysis\r\n- **`tt.scan` Layout Fixes**: Fixed `tt.scan` with broadcasted layouts and additional scan layout issues (#9185, #9189)\r\n- **Reduce/Scan Verifier**: Verify reduce/scan op axis values (#9061)\r\n- **Pipelined Loops Skip Asserts/Prints**: Loops containing `assert` or `print` are no longer pipelined (#9055)\r\n- **Async Op Semantics**: Added explicit semantics for async ops (#8966)\r\n- **WGMMA Wait Delay**: Delay `wgmma wait(0)` to first use of the accumulator (#9021, #9179)\r\n- **WGMMA Register Pipelining**: Added missing waits in WGMMA RHS register pipelining (#8964, #8970, #8997)\r\n- **WGMMA RS Split Limit**: Limit RS-dot splitting to two splits (#9152)\r\n- **Layout Hoisting Fix**: Fixed handling of conflicting layouts when hoisting convert into conditionals (#9083)\r\n- **Rematerialization Cost**: Consider rematerialisation cost when hoisting over `ext`; improved robustness of ext slice rematerialization (#9194, #9019)\r\n- **AxisInfo Improvements**: Enhanced divisibility handling in `AxisInfo` for add/sub; reland of unvisited-operand handling (#9297, #8758)\r\n- **Layout Picker for Small `async_cp`**: Pick better layouts for small `async_cp` (#9183)\r\n- **Skip Conversion-Backward-Slice Cycle**: Skip values with existing conversions in `getConvertBackwardSlice` (#8291)\r\n- **Membar Improvements**: Consider `memdesc_slice` in Membar; extended membar with third-party ops via traits; AMD-aware `membarFilter` (#8755, #8798, #9265)\r\n- **Reduce Op Lowering**: Improvements to `ReduceOp` lowering, later reverted on release branch (#9192, #9214)\r\n- **Clamp on Scalars**: Support clamp optimization on scalars (#8796)\r\n- **`kReg` smem Padding**: Separated additive `kReg` shared-memory padding contribution (#9286)\r\n- **`tcgen05.mma` + multicast support** and Generalized Encodings: continued generalization of TMEM and shared-memory layouts (#9071)\r\n- **`SwizzledShared` Layout**, **uniform hint** on `ttg.warp_id`, and CGAEncoding rename (#9286, #9073, #8850, #9040, #9125)\r\n- **Pipelining Barrier Location**: Fixed barrier placement in loop lowering for MMA ops with non-pipelined operands (#8732)\r\n- **Properly Async wgmma Loop Detection**: Fixed `dotCanBeProperlyAsync` when wgmma is not yielded by the loop and an associated infinite loop (#9274, #9282)\r\n- **`FuncOpToLLVM` Refactors**: Moved `handleArgPtrDatatype` to `Utility.h`; support for LLVM struct/array types in `DITypeAttr` (#9120, #9124)\r\n- **Cache Robustness**: Handle corrupted on-disk cache (#8923)\r\n- **Async Sentinel**: Added a sentinel when async-compiling (#9251)\r\n- **`JITFunction` in `preload`**: Support `JITFunction` in `preload` (#8794)\r\n\r\n### CONSAN (Concurrency Sanitizer) & Debug\r\n- Buffer-region analysis, aliasing support, false-positive deadlock fix, overflow-check disable, compile-time optimization, reduced coverage configurations, TMEM allocation handling, and removal of TMEM size verification (#8837, #8939, #9046, #8940, #9240, #9294, #8787, #8782)\r\n- **Debug Info**: Fixed missing kernel arguments in LLVM debug info; fixed address-sanitizer stack-use-after-scope (#9002, #9088)\r\n\r\n---\r\n\r\n## AMD/HIP Backend\r\n\r\n> 3.7 is heavy on **gfx1250 (RDNA4)** maturation, **warp specialization on AMD**, **Tensor Data Movement (TDM)**, and a new **warp-pipeline** path.\r\n\r\n### Warp Specialization & Warp Pipelining on AMD\r\n- **Warp-Pipeline Support**: New AMD warp-pipeline path with Gluon and LLVM lowering (#8586, #8975, #8980)\r\n- **Warp Specialization on gfx1250** (#8947, #8968)\r\n- **Warp-Pipeline Fixes**: Priority hints and Gluon fixes for the new pipeline (#9301)\r\n- **`ttg.warp_id` and AMD Conversions** (#8659)\r\n\r\n### gfx1250 / RDNA4 Maturation\r\n- **Mixed-Precision Scaled Dot**: Enabled mixed-precision (scaled) dot in Triton on gfx1250 (#8938)\r\n- **4-Warp / 8-Warp MXFP GEMM**: 4-warp scheduling and 8-warp pingpong + MXGEMM refactor (#9031, #9356)\r\n- **Persistent WS f16 GEMM**: Persistent variant and persistent subtiled variant for WS f16 GEMM (#8990, #9052)\r\n- **F16 GEMM Examples Updates**: Updated MXFP FA example and f16 GEMM examples (#9326, #8972)\r\n- **Buffer Atomics for RDNA4**: Enabled buffer atomics on RDNA4 (#8778)\r\n- **`v_permlane16_swap`**: Enabled for `convert_layout` and `reduceOp` on GFX1250 (#8724)\r\n- **Extended FP Conversion**: Including RTZ rounding fixes for GFX1250 (#8821, #8965)\r\n- **libdevice for ROCm 7.1**: Updated libdevice bitcode files (#8807)\r\n- **Cluster Loads / Multi-CTA**: Multi-CTA GEMM example for gfx1250, multi-CTA support for `AMDWmmaEncodingAttr`, scalar-pointer cluster-load avoidance (#9342, #9340, #9129)\r\n- **Gluon `AMDWMMALayout` Rank Consistency** (#9127)\r\n- **WMMA Database Additions**: Added `i8xi8xi32` v3, missing `f64.16x16x4.f64`, and clamp operand on WMMA int intrinsic (#9267, #9271, #9291, #9359)\r\n- **Wavefront Scheduling**: Fixed waitcnt for gfx1250 (#8835)\r\n- **Gluon Stream-K**: 4- and 8-warp stream-k Gluon kernels for gfx1250 (#9370)\r\n- **Roll-up Updates**: Bundled small gfx1250 fixes (#9365)\r\n\r\n### Tensor Data Movement (TDM)\r\n- **Multi-CTA & Multicast for TDM** (#8790)\r\n- **Host-Side TDM Descriptor**: 1D-5D support on gfx1250 (#8977)\r\n- **TDM L2 Prefetch**: Backend and Gluon exposure (#9086, #9148)\r\n- **TDM Predicate**: Use TDM predicate in f16 GEMM variants (#9054)\r\n- **TDM Async Wait**: Support TDM `AsyncWait` in `UpdateAsyncWaitCount` (#9352)\r\n- **TDM Padding in Store**: Support padding when interval equals the inner dimension (#9360)\r\n- **TDM Async Scatter/Gather**: Tensor async scatter/gather support and fixed OOB handling (#9299, #9313, #9371)\r\n- **TDM Shape Adjustment**: Account for CGA offset in TDM shape adjustment (#9341)\r\n- **4D+ TDM Bug Fix**: Fixed TDM behavior when `dim > 2` (#8994)\r\n- **Some TDM Features Enabled** (#9283)\r\n\r\n### Async Copy / LDS\r\n- **AsyncCopy Default On**: Enabled `AsyncCopy` by default for gfx950 and gfx1250 — later reverted on release/3.7.x (#9445, #9087)\r\n- **Async Copy Block Dim Duplication**: Allow async load global-to-load block-dim duplication (#8788)\r\n- **Direct-to-LDS Refactors**: Fixed shared-order selection on GFX9, refactored coalescing checks, contiguity hints, vector-size fixes for padded encodings (#9028, #9041, #9048, #9089, #9149)\r\n- **`v_perm` for `convert_layout`** (#9014)\r\n- **Padded Layout Heuristic**: Relaxed heuristics for smaller block sizes (#9074)\r\n\r\n### Reorder / Pipelining Cleanup\r\n- **`ReorderInstructions`**: Removed `sinkSecondLoad`, `sinkDotConversion`, and `moveUpTranspose` optimizations (#9119, #9139, #9204, #9229)\r\n- **Replace `ReorderInstructions` with `MoveUpPrologueLoads`** (#9328)\r\n- **`UpdateAsyncWaitCount`**: Support single-block `execute` regions (#9126)\r\n- **`OptimizeLDSUsage` Removal** (#8282)\r\n\r\n### libdevice / Layouts / Misc\r\n- **`finite`/`isfinited`**, **`rint`**, **`clampf` via `v_med3`**: libdevice and codegen additions (#9097, #9166, #9256)\r\n- **`BlockPingpong` Improvements**: Debug messages and dot-dominates-predecessors fix (#8804, #9027)\r\n- **`kWidth` mandatory for WMMA v3** (#8783)\r\n- **`copysign` Replacement**: Replaced LLVM `copysign` intrinsic (#8789)\r\n- **WMMA Layout CTA Fields**: Generalized (#8946)\r\n- **TDM with `CanonicalizePointers`**: Support `MakeTensorDescOp` in `CanonicalizePointers` (#9228)\r\n- **`PartitionedSharedEncodingAttr`**: Introduced and reverted (#9314, #9367)\r\n- **`scf.if` Combining**: Added `PrepareIfCombining` pass (#9253)\r\n- **Fine-Grained Cluster Barrier**: New AMD cluster barrier exposed to Gluon (#9206)\r\n- **`SinkLayoutConversions` Pass** (#9168)\r\n- **MIR Swap**: Option to swap MIR; `addOccurrence` for proper LLVM-option disabling; `ScopedNoAliasAAWrapperPass` in MIR swap pipeline (#8711, #9311, #9309)\r\n\r\n### AMD Bug Fixes (selected)\r\n- **`atomic_cas` Fixes**: Wrong struct index for atomic-CAS pattern, ignored sem/scope, and atomic-CAS for non-int types (#8867, #9042, #9116)\r\n- **Atomic-RMW Mask Vectorization**: Fixed wrong vectorization width for masked atomic-RMW (#9142)\r\n- **BroadcastedRegisters in Compilation**: Fixed compilation crash (#8828)\r\n- **`uniformSum` Crash**: Fixed null `uniformSum` in `CanonicalizePointers` (#8991)\r\n- **Cooperative Groups Support**: Driver check (#8935)\r\n- **FP8/BF8 WMMA Selection** on release/3.7.x: Fixed mixed FP8 promotion / instruction selection (#9567, #9581)\r\n- **True16 on gfx11**: Disabled True16 for assembler on gfx11 (#9447, #9476)\r\n- **`RangeAnalysis` `tripCount`**: Fixed trip-count calculation (#9383, #9944)\r\n- **Padded-Layout Async Copy OOM**: Fixed OOM in pipelining with padded async copy on GFX950 (#9442, #9945)\r\n- **`BlockPingpong` for non-MFMA dot** (#9618, #9948)\r\n- **`CanonicalizePointers` Different Bases** (#9541, #9950)\r\n- **Backend cherry-pick dance** (#9487, #9502, #9673, #9675)\r\n- **FP4 Matmul Tests Skipping**: Skip tests packed along M/N for gfx1250 (#9176)\r\n\r\n---\r\n\r\n## NVIDIA Backend\r\n\r\n### Blackwell & Newer SMs\r\n- **`tcgen05` MMA on sm110 (Jetson Thor)** (#9160)\r\n- **`tcgen05.ld.red` on sm103**: Implemented in Gluon (#9151)\r\n- **x Scale Swizzling for Blackwell + Batched Matmul** (#8863)\r\n- **Block-Scaled Matmul Baselining**: mxfp8/nvfp4 block-scaled cuBLAS baselines (#9044)\r\n- **ptxas for Blackwell**: Repeated ptxas-version uprev/revert; final state on release/3.7.x cherry-picks the GB300/Spark/THOR-required commits (#8941, #9011, #9016, #8983, #9363, #9621)\r\n- **NVMMA Variadic CUDA Launcher**: Variadic-argument pre-compiled CUDA launcher (#6788)\r\n- **`NVIDIA::canSkipBarSync`**: Resurrected (#9246)\r\n\r\n### TMA\r\n- **TMA im2col Mode**: End-to-end im2col TMA support — `AsyncTMACopyGlobalToLocalOp`, tensor-descriptor support, fix for `tma load`, and driver support (#9202, #9225, #9303, #9305)\r\n- **TMA Encoding Verification**: Verify encodings on TMA ops (#8886)\r\n- **TMA Descriptor Mitigation**: Mitigation against potential TMA descriptor creation errors (#9235)\r\n\r\n### Hopper / WS\r\n- **`tt.split`/`join` in WS Data Partition**: Hopper WS support for `tt.split`/`tt.join` (#456, #9147)\r\n- **mx8 `w_scale` Mask**: Fixed Hopper mask (#8974)\r\n- **Small-Batch Hopper**: Bench fixes for small batches on Hopper (#8877)\r\n- **SM89 ptxas Workaround Reverted**: Removed the older workaround for the SM89 ptxas bug now that it is unnecessary (#9756)\r\n\r\n---\r\n\r\n## Gluon & Layout Improvements\r\n\r\n### New Features\r\n- **Local Scatter/Gather**: Added local scatter/gather support to Gluon (#8480)\r\n- **`get_view()`**: Added `get_view()` for Gluon layouts (#9270)\r\n- **Finer Cluster Fences**: Exposed finer-grained cluster fences (#9076)\r\n- **Multi-CTA Refactor of `PaddedSharedLayouts`** (#9336)\r\n- **\"Illegal Instruction\" Sanitize Mode**: Tightened TMA op verifiers and added an \"illegal instruction\" sanitize mode (#9112)\r\n- **Verifier Improvements**: Tightened Gluon dialect verifiers and moved checks into C++ (#8981, #9018, #9033)\r\n- **TensorMemory in `to_linear_layout`**: Allow TM layouts in `to_linear_layout` for printing (#8682)\r\n- **More Blackwell Tutorials** (#8982)\r\n\r\n### Layouts & Shared Encodings\r\n- **LinearEncoding Tightening**: Tightened LinearEncoding checks (#9215)\r\n- **`SharedLinearEncoding`**: Continued lowering generalization (carry-over from 3.6 with backend updates).\r\n\r\n---\r\n\r\n## Kernels & Benchmarks\r\n\r\n### Persistent Matmul\r\n- **Persistent Matmul Heuristics**: Fixed and refined heuristics (#8791, #8813)\r\n- **Hopper HBM Swizzling**: Persistent matmul now supports Hopper HBM swizzling (#8917)\r\n- **Hopper FP4 Swizzled, num_warps=4** (#9029)\r\n- **Don't Flatten Mixed-Precision Hopper Persistent Matmul** (#9279)\r\n- **High-Occupancy Persistent Matmul**: Re-enabled (#9248)\r\n- **4-Warp Persistent Kernel**: Re-enabled after fixes (#9331)\r\n- **Strided Layout Handling for Persistent**: Fixed when setting `requires_persistent` (#9198)\r\n- **Mxfp Non-Persistent Strided Layout**: Allow non-persistent mx matmul with strided layout (#8808)\r\n\r\n### Triton Kernels Refactor\r\n- **Matrix-Multiplication Refactor**: Major refactor of triton_kernels matmul (#8765)\r\n- **Tensor/Layout/Distributed Refactor**: Reland of the tensor/layout/distributed refactor; small follow-ups (#9134, #9140, #9186, #9187, #9213)\r\n- **Closure-Based Output Mapping**: For peer shards (#8999)\r\n- **Distributed Tests**: Distributed routing kernels test fix (#9258)\r\n- **Device Descriptor Allocator**: Keep a pool to fix descriptor allocator behavior (#9259)\r\n- **Reduce Kernel**: Unfuse FMA for numeric stability, unpadded batch handling, global scale (#9320, #9332, #9372)\r\n- **`Tensor.clone`**: Briefly added `clone` for `triton_kernels.tensor.Tensor`, then reverted (#9178, #9208)\r\n\r\n### MXFP / Scaled-Dot Kernels\r\n- **Force `mxfp4→bf16` Conversion via `mul.bf16x2`** (#8967)\r\n- **Hopper mxfp4 Swizzled, num_warps=4** (#9029)\r\n- **swiglu Optimizations**: Save instructions, then partial revert; later use of `ex2.approx.ftz` for swiglu (#8801, #8905, #9164)\r\n- **matmul Output mxfp Format Fixes** (#8865)\r\n- **Symmetric Memory in Bench**: Release symmetric memory between runs (#8900)\r\n- **`distributed.py` / `bench_utils.py`**: Extracted common code from `bench_mlp.py` and `distributed.py` (#8866)\r\n- **`num_stages` Adjustment**: For bf16/fp16 × mxfp (#8773)\r\n\r\n### Other\r\n- **X Scale Swizzling for Ragged** (#8897)\r\n- **`reduce_forward` Metadata**: Improved performance (#9068)\r\n- **TF32 Rounding in MoE** (#9296)\r\n- **`p_matmul` Asserts & Fixes** (#9376)\r\n- **Distributed `symm_mem_pool` by Argument** (#9092, #9155)\r\n\r\n---\r\n\r\n## Proton Profiling\r\n\r\n### Highlights\r\n- **Hardware Trace on Blackwell**: Enabled low-overhead hardware trace (#9307)\r\n- **Significant `deactivate` / `get_data` Overhead Reduction**: Especially for CUDA-graph profiling; exposed `get_data_msgpack` (#9030)\r\n- **Periodic Dumping**: Periodic profile dumping; metadata profiling with periodic flushing (#9150, #9236)\r\n- **Multi-Device Metric Profiling**: Fixed metric buffer deadlock and added multi-device support (#8943)\r\n- **Capture-on-Error**: Capture traces even when code exits with an error (#8955)\r\n- **Vector Metrics**: New vector metric type (#9329)\r\n\r\n### API & Internals\r\n- **`get_data` API**: Export profile data directly in Python (#8928)\r\n- **`clear_data` API**: Remove pre-deactivation data (#8971)\r\n- **`finalize` Cleanup**: Clean up context source after teardown (#9069)\r\n- **Fewer Locks**: Further reduce unnecessary locks (#9257)\r\n- **Runtime/Metric Correlation**: Simplified to reduce overhead (#9132)\r\n- **Selective Kernel Metadata**: Allow Proton to record metadata for selective kernels (#9158)\r\n- **Metric Type Restrictions**: Restrict frontend metric types (#8858)\r\n- **Init/Final Timestamps**: Added to Chrome trace (#8870)\r\n- **`GlobalScratchAllocOp` Deprecation**: Deprecated Proton's own op in favor of TritonGPU's, with a custom backend (#8976)\r\n- **Drop Invalid-Time Kernels** (#8961)\r\n- **Ignore Metric-Kernel Timing** (#9058)\r\n- **Documented Experimental APIs** (#9056)\r\n- **HW Trace Default Fix**: Fixed default value for `TRITON_ENABLE_HW_TRACE` in `CuptiProfiler` (#9324)\r\n- **Tensor Descriptor & 2-CTA Tests** (#9070)\r\n- **AMD Proton Test Fixes** (#8763)\r\n\r\n---\r\n\r\n## Testing & CI\r\n- **Gluon TMA + MMA Hopper/Blackwell Test** (#8873)\r\n- **AMD Shadow CI**: New AMD runner setup, then reverted (#9032, #9049)\r\n- **`fresh_knobs` Default Behavior** (#9184)\r\n- **`tl.dot` BF16xN Nondeterminism** (#8818)\r\n- **Disable Stack Traces in Performance Remarks** (#8884)\r\n- **Pin pandas < 3.0** (#9273)\r\n- **Fix pytorch Deprecation Warning in CI** (#8857)\r\n- **Reduce Wheel Size, Pin `DOCKER_API_VERSION` (release/3.7.x)** (#10244)\r\n- **Increase Release Wheel Timeout** (#10250)\r\n- **Skip Tests for RDNA / gfx1250**: Various AMD test skips and enables (#9210, #9176, #9177, #9232, #9343, #9095)\r\n- **Triton's `assert_close`**: Propagate `err_msg` to numpy (#9170)\r\n- **Float8 × MX Tolerance** (#9316, #9338)\r\n- **NumPy 2.4 Compatibility**: Explicit numpy-array-to-scalar conversion (#9172)\r\n- **`test_line_info_ir_source` Flake Fix** (#9161)\r\n\r\n---\r\n\r\n## Build & Infrastructure\r\n- **`CMAKE_LIBRARY_OUTPUT_DIRECTORY`**: Fixed build with empty directory (#8810)\r\n- **`llvm_update_compile_flags` Removal** (#9167)\r\n- **`LLVM_BUILD_SHARED_LIBS` Canonicalization** (#8933)\r\n- **actions/checkout v5 → v6** (#8826)\r\n- **Version Bumps**: 3.5.0 → 3.6.0 ; 3.6.0 → 3.7.0 (#8836, #9885, #9888)\r\n- **`TRITON_EXT_ENABLED` for Wheels** (#9935, #9959)\r\n- **`nvidia-toolchain-version.json` Update**.\r\n- **`TRITON_DEFAULT_BACKEND`**: Control `driver.active` via this env var (#9144)\r\n- **`TRITON_PTXAS_BLACKWELL_PATH`**: Allow override of `ptxas-blackwell` binary (#8945)\r\n- **Release to PyPI** (#10251)\r\n- **`topk` in Plugin Example**: Increment index in plugin example (#9315)\r\n- **HIP Support in `link.py`** (#9084)\r\n\r\n---\r\n\r\n## Documentation\r\n- **Divisibility Reset Logic**: Clarified for contiguous dimensions in `AxisInfo` (#9266)\r\n- **`topk` Operation**: Added to language documentation (#9345)\r\n- **Plugin Example README**: Added a second pass-plugin README example (#8815)\r\n- **Conference Materials**: Updated README (#9009)\r\n- **Community Meetup Notes**: Added 2026-01-06 meetup notes (#9288)\r\n- **`warp_specialize` Docs**: Updated `gl.warp_specialize` docs (#8553)\r\n- **`LinearLayout` Output Matrix Comment**: Doc fix (#9243)\r\n\r\n---\r\n\r\n## Breaking Changes\r\n- **`triton_kernels` matmul refactor (BC-breaking)**: The matrix-multiplication refactor introduces a backwards-incompatible API surface. Downstream users of `triton_kernels.matmul_*` should review call sites (#8765)\r\n- **`tcgen05.cp` Lowering Generalization & `tcgen05.mma` Encoding Acceptance**: Continued from 3.6, with new verifier behavior and stricter encoding checks.\r\n- **Proton `GlobalScratchAllocOp` Deprecated**: Replaced with TritonGPU's `GlobalScratchAllocOp` + custom backend. Out-of-tree consumers must migrate (#8976)\r\n- **`make_block_ptr` Deprecated**: A deprecation warning is now emitted; users should migrate to tensor descriptors (#9667)\r\n- **Default 32-bit Dot Precision Reverted**: Default 32-bit dot precision was briefly TF32x3 — the default in 3.7 remains as in 3.6. Note the new \"round f32→tf32 in descriptor\" option (#9080, #9090, #9295)\r\n- **AsyncCopy Default for gfx950 / gfx1250**: Was enabled by default and then reverted on the release branch. Users must opt in explicitly in 3.7 (#9087, #9445)\r\n- **SM89 ptxas Workaround Reverted**: The ptxas workaround introduced earlier is removed on release/3.7.x (#9756, #7067)\r\n\r\n---\r\n\r\n## Contributors\r\n\r\nThis release includes contributions from engineers at:\r\n\r\n- Meta\r\n- AMD\r\n- NVIDIA\r\n- OpenAI\r\n- Intel\r\n- Google\r\n- And many individual contributors\r\n\r\nSpecial thanks to all contributors who submitted bug reports, feature requests, and code improvements!\r\n\r\n","reactions":{"url":"https://api.github.com/repos/triton-lang/triton/releases/319213104/reactions","total_count":14,"+1":7,"-1":0,"laugh":0,"hooray":0,"confused":0,"heart":0,"rocket":7,"eyes":0}},{"url":"https://api.github.com/repos/triton-lang/triton/releases/278767719","assets_url":"https://api.github.com/repos/triton-lang/triton/releases/278767719/assets","upload_url":"https://uploads.github.com/repos/triton-lang/triton/releases/278767719/assets{?name,label}","html_url":"https://github.com/triton-lang/triton/releases/tag/v3.6.0","id":278767719,"author":{"login":"atalman","id":7563158,"node_id":"MDQ6VXNlcjc1NjMxNTg=","avatar_url":"https://avatars.githubusercontent.com/u/7563158?v=4","gravatar_id":"","url":"https://api.github.com/users/atalman","html_url":"https://github.com/atalman","followers_url":"https://api.github.com/users/atalman/followers","following_url":"https://api.github.com/users/atalman/following{/other_user}","gists_url":"https://api.github.com/users/atalman/gists{/gist_id}","starred_url":"https://api.github.com/users/atalman/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/atalman/subscriptions","organizations_url":"https://api.github.com/users/atalman/orgs","repos_url":"https://api.github.com/users/atalman/repos","events_url":"https://api.github.com/users/atalman/events{/privacy}","received_events_url":"https://api.github.com/users/atalman/received_events","type":"User","user_view_type":"public","site_admin":false},"node_id":"RE_kwDOAWaHXs4Qnahn","tag_name":"v3.6.0","target_commitish":"main","name":"Triton 3.6.0 release","draft":false,"immutable":false,"prerelease":false,"created_at":"2026-01-20T14:17:56Z","updated_at":"2026-01-21T23:27:58Z","published_at":"2026-01-21T23:27:22Z","assets":[{"url":"https://api.github.com/repos/triton-lang/triton/releases/assets/344008379","id":344008379,"node_id":"RA_kwDOAWaHXs4UgSa7","name":"triton-3.6.0.tar.gz","label":"","uploader":{"login":"github-actions[bot]","id":41898282,"node_id":"MDM6Qm90NDE4OTgyODI=","avatar_url":"https://avatars.githubusercontent.com/in/15368?v=4","gravatar_id":"","url":"https://api.github.com/users/github-actions%5Bbot%5D","html_url":"https://github.com/apps/github-actions","followers_url":"https://api.github.com/users/github-actions%5Bbot%5D/followers","following_url":"https://api.github.com/users/github-actions%5Bbot%5D/following{/other_user}","gists_url":"https://api.github.com/users/github-actions%5Bbot%5D/gists{/gist_id}","starred_url":"https://api.github.com/users/github-actions%5Bbot%5D/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/github-actions%5Bbot%5D/subscriptions","organizations_url":"https://api.github.com/users/github-actions%5Bbot%5D/orgs","repos_url":"https://api.github.com/users/github-actions%5Bbot%5D/repos","events_url":"https://api.github.com/users/github-actions%5Bbot%5D/events{/privacy}","received_events_url":"https://api.github.com/users/github-actions%5Bbot%5D/received_events","type":"Bot","user_view_type":"public","site_admin":false},"content_type":"application/gzip","state":"uploaded","size":6424965,"digest":"sha256:2251e067e4229e4a379259292e284232369977cb3ba1fb2424457df80115ec0a","download_count":1223,"created_at":"2026-01-21T23:27:57Z","updated_at":"2026-01-21T23:27:58Z","browser_download_url":"https://github.com/triton-lang/triton/releases/download/v3.6.0/triton-3.6.0.tar.gz"}],"tarball_url":"https://api.github.com/repos/triton-lang/triton/tarball/v3.6.0","zipball_url":"https://api.github.com/repos/triton-lang/triton/zipball/v3.6.0","body":"# Triton 3.6 Release Notes\r\n\r\n## Table of Contents\r\n- [Dialect & Frontend](#dialect--frontend)\r\n- [Backend & Compiler](#backend--compiler)\r\n- [AMD/HIP Backend](#amdhip-backend)\r\n- [NVIDIA Backend](#nvidia-backend)\r\n- [Gluon & Layout Improvements](#gluon--layout-improvements)\r\n- [Kernels & Benchmarks](#kernels--benchmarks)\r\n- [Proton Profiling](#proton-profiling)\r\n- [Testing & CI](#testing--ci)\r\n- [Build & Infrastructure](#build--infrastructure)\r\n- [Documentation](#documentation)\r\n- [Breaking Changes](#breaking-changes)\r\n\r\n---\r\n\r\n## Dialect & Frontend\r\n\r\n### New Features\r\n- **Multidimensional Batch Support** (#8542): Added support for multidimensional batches in `tl.trans` and `tl.dot` operations\r\n- **Ragged TMA Atomic Add** (#8238): Added atomic add support for ragged TMA operations\r\n- **Integer Range Utility** (#8753): Exposed an integer-range utility from AMD range analysis code for broader use\r\n- **Constexpr Through Min/Max** (#8733): Propagate constexpr through builtin min/max functions (BC-breaking)\r\n- **Scales Dimension Checks** (#8564): Added dimension checks for scales in `dot_scaled` operations\r\n- **Loop Bounds Verification** (#8243): Added verification that loop bounds are scalars\r\n\r\n### Bug Fixes\r\n- **For Loop Induction Variable** (#8750): Fixed modification of for loop induction variable handling\r\n- **Store Broadcasting** (#8661): Fixed broadcasting issues in store operations\r\n- **Missing `dot_scaled` Handling** (#8658): Fixed missing handling for None acc in `dot_scaled`\r\n- **AugAssign Line Information** (#8703): Attached proper line number information to AugAssign nodes\r\n- **Starred Argument Handling** (#8686): Made starred argument handling more robust\r\n- **Saved Exception Cloning** (#8115): Fixed clone of saved exception before raising\r\n- **Tuple Mangling** (#8060): Fixed mangling for tuples in JIT compilation\r\n\r\n### Improvements\r\n- **Optimized `tl.cdiv`** (#8669): Optimized `tl.cdiv` for common case of 32-bit divisors\r\n- **Un-deprecated min/max** (#8734): Un-deprecated min/max on scalar tensors\r\n- **Warmup in KernelInterface** (#8757): Moved warmup functionality into KernelInterface\r\n- **Verification with Diagnostics** (#8074): Frontend always verifies with diagnostics enabled\r\n- **Constexpr with do_not_specialize Error** (#8275): Added error when constexpr is combined with do_not_specialize\r\n- **Deprecated ast.Num Replacement** (#8698): Replaced usage of deprecated `ast.Num`\r\n\r\n---\r\n\r\n## Backend & Compiler\r\n\r\n### LLVM Updates\r\n- **LLVM Bump** (#8299): Bumped to llvm/llvm-project@f6ded0be897e\r\n- **LLVM Head Merge** (#8612): Merged back changes from llvm-head with updated APIs\r\n- **Inliner Import** (#8152): Import inliner in triton-opt for better optimization\r\n\r\n### Code Generation\r\n- **CTALayout as LinearLayout** (#8770): Made CTALayout an honest-to-goodness LinearLayout for better representation\r\n- **Shared Layout Rank Check** (#8772): Added check that Shared layouts have rank equal to the tensor or one less\r\n- **Backward Propagation Fix Point** (#8776): Run remove backward prop until fix point for correctness\r\n- **Generic `tcgen05.cp` Lowering** (#8225): Implemented generic lowering for `tcgen05.cp`\r\n- **Generic Matrix Descriptors** (#8321): Implemented shmem matrix descriptors generically\r\n- **LinearSharedEncoding Support** (#8116): Added support for LinearSharedEncoding\r\n- **BF16x3 Trick** (#7592): Implemented BF16x3 trick for improved performance\r\n- **Padded Shared Linear Remapping** (#7929): Added linear remapping to padded shared layout\r\n\r\n### Optimizations\r\n- **Compilation Time Improvement** (#8689): Improved compilation time in constant sanitizer pass\r\n- **AxisInfo Loop Removal** (#8679): Removed unnecessary loop over roots in AxisInfo analysis\r\n- **Constant Analysis** (#8502): Improved constant analysis in AxisInfo\r\n- **Combinatory Explosion Prevention** (#8477): Prevented combinatory explosion when checking tmem_load uses\r\n- **Layout Conversion Vectorization** (#8655): Fixed vectorization for convert_layout with ldmatrix and stmatrix\r\n- **Maybeduplicate Generalization** (#8492): Generalized maybeDeduplicate to all layouts\r\n\r\n### Bug Fixes\r\n- **cp_async Alignment** (#8752): Fixed cp_async used in pipeliner when alignment info gets lost\r\n- **While Op Layout Propagation** (#8751): Prevented backward layout propagation through while op\r\n- **AxisInfo Handling** (#8723, #8754): Fixed handling of unvisited operands in AxisInfoAnalysis\r\n- **64-bit Atomic CAS** (#8105): Fixed 64-bit `atomic_cas` operation\r\n- **Memdesc of Pointers** (#8515): Fixed memdesc handling for pointer types\r\n- **Alloc Shape Reset** (#8537): Reset alloc_shape when doing memdesc_index\r\n- **Denorm Flushing** (#8557): Don't flush denorms for precise div/sqrt\r\n- **Local Load Reordering** (#8423): Prevented reordering local_load across side-effecting operations\r\n- **Pattern Reordering** (#8266): Restricted pattern re-ordering of alloc and reshape\r\n- **Poison Op AxisInfo** (#8489): Fixed AxisInfo handling of PoisonOp producing MemDesc\r\n\r\n### Analysis Improvements\r\n- **Trans Contiguity** (#8226): Added tt.trans contiguity analysis support\r\n- **Hint Analysis** (#5254): Fixed hint analysis in axis info\r\n- **Topological Sort Deprecation** (#8596): Deprecated triton's custom topological sort in favor of MLIR's\r\n\r\n---\r\n\r\n## AMD/HIP Backend\r\n\r\n### GFX1250 (RDNA4) Support\r\n- **Initial Skeleton** (#8131): Added gfx1250 skeleton support\r\n- **WMMA Support** (#8174, #8283, #8312): Added initial and scaled WMMA support for gfx1250\r\n- **TDM Support** (#8333, #8392, #8479): Added Tensor Data Movement (TDM) load/store support\r\n- **Async Copy** (#8509, #8510, #8621, #8622): Added async copy and async wait support\r\n- **Buffer Ops** (#8130, #8532): Enabled buffer atomics and exposed buffer ops\r\n- **Multicast Loads** (#8719, #8759): Added async load to LDS multicast and multicast in `tt.LoadOp`\r\n- **ds_read_tr** (#8461): Added gfx1250 support for ds_read_tr\r\n- **LDS Memory Barriers** (#8681): Added support for LDS memory barriers\r\n- **Shared Memory Size** (#8517): Updated shared memory size from TargetInfo\r\n- **num_cta > 1** (#8718): Support launches with num_cta > 1 on gfx1250\r\n- **Scale Preshuffling** (#8576): Implemented scale preshuffling and opSel\r\n\r\n### MXFP & Scaled Dot\r\n- **Scale Preshuffling in Decomposed Dot** (#8170): Support scale preshuffling in decomposed scaled dot\r\n- **Pipeline Scale via LDS** (#8258): Pipeline scale in decomposed scaled dot via LDS\r\n- **Scaled Upcast Ops** (#8088): Introduced scaled upcast ops for hardware upcasting\r\n- **FP4->BF16 Optimized Conversion** (#8145): Added optimized fp4->bf16 conversion for MI300\r\n- **Scaled Dot Decomposition for GFX950** (#7839): Enabled f16 * mxfp scaled dot decomposition\r\n\r\n### Layout & Memory Optimizations\r\n- **Permlane Swap** (#7947): Use permlane_swap for layout conversions between dot operations\r\n- **Padded Shared with AsyncCopy** (#8365): Use PaddedLayout with AsyncCopy on gfx950 when pipelining\r\n- **LDS Layout Selection Redesign** (#8053): Redesigned stream pipeliner LDS layout selection logic\r\n- **Padded Encoding Restrictions** (#8583): Relaxed padded encoding block size restrictions\r\n- **Direct-to-LDS with Padded** (#8185): Coalesce direct-to-lds loads with padded encodings\r\n- **Contiguity Hint for Direct-to-LDS** (#8761): Use contiguity hint for direct-to-lds ops\r\n- **BypassLDS Feature** (#7968): Added bypassLDS feature to StreamPipeline\r\n\r\n### Code Generation\r\n- **ds_read_tr with Linear Layout** (#8235): Use linear layout to infer and emit ds_read_tr\r\n- **ds_read_tr Restrictions Lifted** (#8442): Lift unneeded ds_read_tr lowering restrictions\r\n- **ds_read_tr Vec Size Limit** (#8377): Limit vec size for ds_read_tr + padded layouts by min interval\r\n- **Wave ID Optimization** (#8601): Optimized gfx9 wave id code generation\r\n- **MFMA Layout Refactor** (#8213): Refactored MFMA layout implementation\r\n- **MFMA Select Replacement** (#8320): Replaced mfma select in LLVM conversion\r\n- **FP8/BF8 WMMA Instruction Selection** (#8649): Fixed instruction selection for fp8/bf8 wmma\r\n- **Chained WMMA Optimization** (#7374): Optimized chained multiplications for WMMA\r\n- **BF16 v_dot** (#8444): Use v_dot for bf16 multiplication on gfx11/gfx12\r\n\r\n### Build & Driver\r\n- **ROCm 7 Docker Image** (#8224): Switched to use official ROCm 7 docker image\r\n- **HIP v6 Requirement** (#8748): Only require HIP v6 which is necessary\r\n- **HIP Header Update** (#8709): Updated HIP header files to 7.1\r\n- **Optional Symbols Support** (#8729): Support optional symbols in driver.py\r\n- **Uniform Workgroup Size** (#8720): Indicate uniform workgroup size to LLVM\r\n- **MIR Dump Option** (#8663): Added option to dump MIR\r\n- **Custom LLVM Scheduler** (#8326, #8700): Added schedule hint for custom LLVM scheduler\r\n\r\n### Bug Fixes\r\n- **Pointer Canonicalization** (#8465, #8276): Fixed ptr-canonicalization segfault and assertion\r\n- **Large Tensor Pointer Canonicalization** (#8359): Disabled pointer-canonicalization for large tensors\r\n- **Padded Shared Local Load** (#8683): Fixed padded shared when lowering local load\r\n- **Nondeterministic Atomic Tests** (#8633): Fixed nondeterministic atomic tests failure on RDNA\r\n- **Buffer Cache Swizzling** (#8264): Turned off buffer op cache swizzling temporarily\r\n- **Direct-to-LDS on CDNA1/2** (#8280): Disabled direct-to-lds loads on CDNA1 and CDNA2\r\n- **Floating-point Upcasting Rounding** (#8268): Skip rounding mode for floating-point upcasting\r\n- **TilesPerWarp Boundary Cases** (#8467): Fixed deduceTilesPerWarp boundary cases\r\n- **fast_tanhf Overflow** (#8551): Reimplemented fast_tanhf() to avoid overflow\r\n- **MFMA Small K Selection** (#8278): Avoid selecting MFMA with smaller K than problem size\r\n\r\n---\r\n\r\n## NVIDIA Backend\r\n\r\n### Blackwell Features\r\n- **TMEM Bitwidth** (#8136): Added bitwidth to TMEM encoding for better representation\r\n- **TMEM Layout Broadcasting** (#8148): Represent broadcasting in TensorMemoryLayouts\r\n- **TMEM Layout Construction** (#8202): Simplified TMEM layout construction and row/col computation\r\n- **Generic tcgen05.ld/st Layouts** (#8421, #8495): Generate distributed layouts for `tcgen05.ld/st` generically\r\n- **tcgen05.mma Generalization** (#8386): Generalized `tcgen05.mma` to accept `SharedLinearEncodingAttr`\r\n- **tcgen05.cp Generic Lowering** (#8102, #8338): Towards a generic tcgen05.cp lowering via matrix descriptors\r\n- **tcgen05.mma Verifier** (#8725): Fixed missing case in tcgen05.mma verifier\r\n- **Explicit Commit Merge** (#8026): Added rewrite pattern to merge explicit commit ops into MMAv5\r\n- **2CTA Mode Support** (#8644, #8653): Initial support for 2CTA mode in Gluon with global flag\r\n- **reqnctapercluster Emission** (#8645): Emit reqnctapercluster for better cluster sizing\r\n\r\n### SM120 Features\r\n- **Native FP4 Scaled Dot** (#8494): Added native FP4 scaled_dot for SM120\r\n- **Native MXFP FP8 Scaled Dot** (#7918, #8029, #8129): Added native MXFP FP8 scaled_dot for SM120\r\n- **TMA Gather4** (#8498): Enabled TMA gather4 on sm_120 and sm_121\r\n- **DotScaledScaleLayout Rewrite** (#8482): Rewrote getSM120DotScaledScaleLayout and refactored MMAv2\r\n\r\n### Warp Specialization\r\n- **E2E Aref** (#8262): Enabled end-to-end aref for warp specialization\r\n- **TMA Load Aref Insertion** (#7826): Use aref for TMA load pipelining and lowering\r\n- **TMEM Aref Insertion Pass** (#8009): Added aref tmem insertion pass\r\n- **Partition Representation Rework** (#8123): Reworked partition representation\r\n- **Assign Partitions to All Ops** (#8534): Assign partitions to all ops for consistency\r\n- **Nested Loop Recognition** (#8451): Recognize warp-specialized nested loops in AssignLatencies\r\n- **Scalar Ops Across Partition** (#8061): Support scalar ops across partition boundaries\r\n- **Stage/Phase Assignment** (#8329): Assign stage-phase only to partitions that need it\r\n- **Partition Scheduler Annotations** (#8215): Partition-scheduler annotates all ops with fixes\r\n- **Control Flow Support** (#8651): Support ops annotations outside tt.ws loops\r\n- **Then/Else Heuristic Patch** (#8656): Patched partitioner then/else heuristic\r\n- **Fence After Local Store** (#8317): Added missing fence after local_store for MMAv5 consumers\r\n\r\n### Other Enhancements\r\n- **Descriptor Bit 46** (#8032): Turn on bit 46 for descriptors in mmav5\r\n- **Matrix Descriptor No-Swizzle** (#8027): Fixed matrix descriptor for no-swizzle case\r\n- **WGMMA Wait Op CVT** (#8579): Fixed unnecessary cvt caused by wgmma wait op\r\n- **Enable Reflect FTZ Flag** (#8762): Added enable_reflect_ftz flag to NVIDIA backend\r\n- **libcuda.so.1 Usage** (#8668): Modified NVIDIA backend driver to use libcuda.so.1\r\n- **Padded Shared in MemDescSubslice** (#7944): Support padded shared in MemDescSubsliceOp\r\n- **Ptxas Workaround** (#8155): Fixed ptxas workaround in convert_layout\r\n- **ldmatrix/stmatrix.b8.trans** (#7542): Added support in local_load/store for ldmatrix/stmatrix.b8.trans\r\n\r\n---\r\n\r\n## Gluon & Layout Improvements\r\n\r\n### Gluon Language Features\r\n- **Warp Specialize API Change** (#8527): Changed `gl.warp_specialize` API for better usability\r\n- **Multi-CTA Support** (#8468, #8587, #8644): Basic multi-cta support with initial implementation\r\n- **num_ctas Implementation** (#8602): Implemented `num_ctas` in Gluon\r\n- **Device-Side TMA** (#8505): Added device-side TMA support\r\n- **Coalesced Layout** (#8604): Added coalesced layout support\r\n- **get_num_warps** (#8133): Added `ttgl.get_num_warps` metafunction\r\n- **gather Integration** (#8018): Integrated `gather` and its layout tests\r\n- **reduce with No Axis** (#8396): Added support for reduce with no axis\r\n- **assume Support** (#8394): Added support for assume operation\r\n- **cat Remapping** (#8715): Remap more `tl` functions into gluon and expose `cat`\r\n- **Type Verifiers** (#8007): Added type verifiers for many methods\r\n\r\n### Layout System\r\n- **bank_conflicts Exposure** (#8181): Exposed bank_conflicts and to_linear_layout\r\n- **Linear Layout Python Interface** (#8521): Added LL Python Interface\r\n- **Layout Check Message** (#8456): Improved layout check error messages\r\n- **Tensor Rank Verification** (#8242): Verify tensor rank and layout rank match\r\n- **MemDesc Trans/Reshape** (#8251): Have MemDesc{Trans,Reshape} accept equivalent layouts\r\n- **Fp4ToFp Backward Propagation** (#8438): Fixed backwards propagation for Fp4ToFp\r\n- **ResolveAutoEncodings Print** (#8228): Print encoding in ResolveAutoEncodings\r\n\r\n### Gluon AMD Support\r\n- **Host-Side TDM Descriptor** (#8722): Initial support for host-side tdm descriptor\r\n- **TDM 1D-5D Support** (#8743): Support TDM load/store for 1D-5D tensors\r\n- **TDM Pred Exposure** (#8767): Expose pred for TDM load\r\n- **Scale Layout Selection** (#8673): Turn select scale layout into constexpr function\r\n- **WMMA/MFMA Scale Layout** (#8496): Expose get wmma/mfma scale layout\r\n- **AMDWMMALayout Exposure** (#8090): Exposed AMDWMMALayout\r\n- **WMMA for RDNA3/RDNA4** (#8111): Exposed WMMA for RDNA3 and RDNA4\r\n- **Buffer Ops Exposure** (#8532): Expose buffer ops to gfx1250\r\n- **buffer_atomic_rmw API** (#8325): Refactored buffer_atomic_rmw API\r\n- **async_copy for gfx1250** (#8622): Added `async_copy` to Gluon for gfx1250\r\n- **Async Wait Groups** (#8605): Wait outstanding async commit groups instead of instructions\r\n\r\n### Gluon NVIDIA Support\r\n- **tcgen05 mma scaled** (#8393): Added tcgen05 mma scaled support\r\n- **MMAv2 and Dot FMA** (#8227): Exposed MMAv2 and Dot FMA\r\n- **Float2 API** (#8209): Added proper float2 API for Blackwell\r\n- **warp_specialize Docs** (#8553): Updated gl.warp_specialize docs\r\n\r\n### Bug Fixes\r\n- **Translator Fixes** (#8569): Fixed several things in the translator\r\n- **SwizzledSharedLayout** (#8003): Fixed getting layout from a SwizzledSharedLayout\r\n- **Bank Conflict Computation** (#8200): Fixed bank_conflict computation with shmem broadcasting\r\n- **Trans Alloc Optimization** (#8193): Simplified and fixed trans(alloc) optimization\r\n- **TMem Alloc/Store Pattern** (#8192): Fixed pattern combining tmem_alloc and store\r\n- **Constant CSE** (#8323): Disabled constant CSE before auto layout propagation\r\n\r\n---\r\n\r\n## Kernels & Benchmarks\r\n\r\n### MXFP Improvements\r\n- **MXFP Conversions Speedup** (#8610): Significant speedup for mxfp conversions\r\n- **FP32 MXFP Support** (#8672 from 3.5): Added quant/dequant from/to fp32\r\n- **MXFP4 Hopper Layout on A100** (#8474): Apply MXFP4 Hopper layout on A100\r\n- **A100 MXFP4 Upcasting** (#8428): Support A100 upcasting for mxfp4\r\n- **MXFP8 X Support** (#8062): Support mxfp8 `x` in triton_kernels\r\n- **BF16 x MXFP4 Bug Fix** (#8478): Fixed bf16 x mxfp4 bug with SUBTILE_FACTOR > 1\r\n- **EXPT_IS_INNER Support** (#8385): Support EXPT_IS_INNER for MX\r\n- **w_scale Swizzle Handling** (#8652): Handle w_scale without swizzle correctly\r\n- **Max Value Handling** (#8356): Handle values close to max correctly without overflow\r\n- **x_scale OOB Fix** (#8369): Fixed x_scale out-of-bounds access\r\n- **Round-to-Nearest-Even** (#8110): Use round-to-nearest-even mxfp4 quant for consistency\r\n\r\n### Matmul Optimizations\r\n- **Batched Block Sizes** (#7897, #8084): Improved block sizes for batched matmul_ogs with small m/n/k\r\n- **Ragged Matmul DW** (#8256): Added support for ragged matmul dw\r\n- **Split-K Fixes** (#8252): Two small split-k fixes\r\n- **Batched Split-K** (#8327): Fixed and enabled batched matmul with split-k\r\n- **Split-K Constraint** (#8404): Added constraint on `split_k` on `m * n`\r\n- **Launch Metadata** (#8429): Fixed launch metadata computations for matmul_ogs\r\n- **Transposed X Fix** (#8156): Fixed _p_matmul_ogs when x is transposed\r\n- **MX Scale Mask** (#8161): Fixed mx scale mask update\r\n\r\n### Expert Parallelism & MoE\r\n- **Basic Expert Parallelism** (#8448): Basic expert parallelism implementation\r\n- **EP Sharding** (#8493): Incorporated EP sharding and deprecated legacy communication\r\n- **CUDA Graph Tracing** (#8563): vllm compatible version of CUDA Graph tracing for expert parallelism\r\n- **Fused Matmul + Comms** (#8340): Fused matmul_ogs + communications\r\n- **Split-K Decoupling** (#8483): Decoupled split-k reduction from inter-expert reductions\r\n- **Small Batch MoE Tuning** (#8206): Tuning for small batch MoE\r\n- **BitmatrixMetadata** (#8375): Added `BitmatrixMetadata` and `RaggedTensorMetadata`; deprecated triton_kernels.routing\r\n- **BitMatrix Fix** (#8599): Fixed BitmatrixMetadata col/row_sorted_indx\r\n- **y_indx Support** (#8472): Support `y_indx` and uniform distribution\r\n\r\n### Benchmarks\r\n- **Roofline Plotting** (#8244): Fixed roofline plotting\r\n- **HipBlas Roofline** (#8216): Integrated hipblas in roofline measurement\r\n- **GFX950 BF16 x MXFP4 MoE** (#8176): Updated parameters for bf16 x mxfp4 MoE kernel\r\n- **MLP Benchmark Fix** (#8699): Added missing `reduction_n=2` to `bench_mlp.py`\r\n- **tl.clamp Usage** (#8728): Use tl.clamp whenever possible in triton_kernels\r\n\r\n### Other Improvements\r\n- **Redundant Reduce Removal** (#8647): Removed redundant reduce for topk=1\r\n- **Split-K with Fused Scatter** (#8618): Forbid use of `split_k > 1` with fused scatter\r\n- **Layout Dataclasses** (#8690): Made layout classes dataclasses (NFC)\r\n- **HopperValue Padding** (#8677): Pad tensors in `HopperValue` layout\r\n- **A100 Default Layout Revert** (#8549): Reverted a100 default layout change\r\n- **opt_flags Reset** (#8453): Added function to reset opt_flags\r\n\r\n---\r\n\r\n## Proton Profiling\r\n\r\n### New Features\r\n- **Global Memory Support** (#8641): Global memory support for proton intra kernel profiler\r\n- **Global Timestamps** (#7729): Capture global timestamps for consistent cross-CTA timeline\r\n- **Intra Kernel Call Stack** (#8071): Added kernel call stack to intra kernel events\r\n- **NVTX/ROCTX Support** (#8095): Init NVTX/ROCTX support for external profilers\r\n- **Graph Profiling** (#8676): Improved graph profiling part-1\r\n- **Disable Flag** (#8293): Added flag to disable proton to use other profilers\r\n\r\n### Improvements\r\n- **Scope ID Allocation Refactor** (#8613): Refactored scope id allocation to allow flexible annotations\r\n- **Concrete Line Info** (#8614): Attached concrete line info to proton operations\r\n- **FinalizeOp Refactor** (#8635): Refactored finalizeOp to reduce buffer write overhead\r\n- **Buffer Size Description** (#8650): Improved default buffer size description\r\n- **Profile Allocator** (#8730): Made profile allocator a global var\r\n- **Backend Lib Settings** (#8246): Simplified backend lib settings\r\n- **Python Frame Representation** (#8241): Unified python frame representation\r\n\r\n### Bug Fixes\r\n- **Dominance Analysis** (#8712): Fixed dominance analysis in Proton\r\n- **Function Metadata Cleanup** (#8713): Do not clean up function metadata at finalize\r\n- **Memory Leak Fix** (#8692): Fixed memory leak and removed unused variables\r\n- **Buffer Overflow Warning** (#8109): Fixed proton intra kernel profiling buffer overflow warning\r\n- **Concurrent Profiling** (#8210): Do not allow concurrent profiling with different modes\r\n- **Triton Function Filtering** (#8021): Filter out all intrinsics when counting triton functions\r\n- **Global Time Trace Precision** (#8309): Fixed global time trace precision\r\n\r\n### Testing\r\n- **Internal Testing Utility** (#8204): Use more internal testing utility\r\n- **Proton Tests Conditional** (#8237): Conditionally include Proton tests\r\n- **AMD Proton Tests** (#8388): Simplified proton tests on AMD\r\n- **Skip AMD Overhead Tests** (#8665): Skip hip overhead tests\r\n- **Globaltime GFX950** (#8627): Disabled test_globaltime on gfx950\r\n\r\n---\r\n\r\n## Concurrency Sanitizer (ConSan)\r\n\r\n### New Features\r\n- **Deadlock Detection** (#8285): Added deadlock detection capability\r\n- **Warp Specialization Support** (#8189, #8265): Added support for WarpSpecialization with fixes\r\n- **TMA Store Validation** (#8672): Support for TMA store validation\r\n\r\n### Improvements\r\n- **Function Call Opcodes** (#8559): Converted consan instrumentation opcodes to function calls\r\n- **Compilation Time** (#8689): Improved compilation time\r\n- **Cache Invalidation** (#8332, #8342): ConSan env var should be cache invalidating\r\n\r\n---\r\n\r\n## Testing & CI\r\n\r\n### Test Infrastructure\r\n- **Frontend Tests for test-nogpu** (#8771): Added frontend tests to make test-nogpu\r\n- **Device Fixture Usage** (#8512): Using device fixture instead of cuda in tensor descriptor tests\r\n- **tb=short in CI** (#8440): Added tb=short to CI for shorter tracebacks\r\n- **Subprocess Removal** (#8350): Removed subprocess usage from test_triton_debuginfo_on\r\n- **SmallVector Crash Fix** (#8544): Fixed SmallVector crash issue of AxisInfoAnalysis\r\n\r\n### AMD Testing\r\n- **GFX950 CI Fixes** (#8741, #8760): Avoid gfx950 runner failing others, fix continue-on-error\r\n- **GFX1250 Tests** (#8680): Updated gfx1250 Gluon tests\r\n- **Padded Layout Lit Tests** (#8399): Added lit tests for pipelining with padded layouts on gfx950\r\n- **CDNA2 Atomic CAS** (#8376): Disabled flaky atomic cas test on CDNA2\r\n\r\n### NVIDIA Testing\r\n- **Warp Specialization Tests**: Enabled WS tests for various features\r\n- **GB200 Error Handling**: Continue running CI when GB200 errors out\r\n\r\n### Lit Tests\r\n- **Redundant CTALayout Removal** (#8704): Removed all redundant CTALayout information from LIT tests\r\n- **ASAN Fix** (#8117): Fixed ASAN initialization-order-fiasco issue in tensor_layout_print.mlir test\r\n- **MMA Support Check** (#8640): Perform supportMMA check during IR verification\r\n\r\n---\r\n\r\n## Build & Infrastructure\r\n\r\n### Build System\r\n- **Python 3.9 Support Removal** (#8222): Cleaned up Python 3.9 related code/docs\r\n- **Python 3.10 Minimum** (#8167): Updated MIN_PYTHON version to 3.10\r\n- **Python 3.14 Wheels** (#7695 from 3.5): Python 3.14 wheel build support\r\n- **Python 3.13 Fix** (#8403): Fixed Python 3.13 compatibility issues\r\n- **CentOS 7 Removal** (#8191): Removed CentOS 7 build\r\n- **Actions Updates** (#8347, #8361, #8187): Bumped actions/setup-python to v6, tj-actions/changed-files to v47\r\n- **TarFile Deprecation** (#8337): Fixed deprecation warning from TarFile.extractall\r\n- **Unused CMake Removal** (#8408, #8362): Removed unused include(ExternalProject) and find_library\r\n\r\n### Compilation & Runtime\r\n- **Native Specialize** (#7771): Native specialize for improved launch latency\r\n- **AsyncCompile Error Option** (#8756): Added option to ignore errors in AsyncCompile\r\n- **JIT Functions to Kernels** (#8721): Added test that jit functions can be passed to kernels safely\r\n- **JIT Specialization Serialization** (#8639): Fixed JIT specialization data (de)serialization for tuples and constexprs\r\n- **Aggregate Cache Keys** (#8528, #8568): Made sure aggregate members are added to the cache key\r\n- **Interpreter Mode Cache** (#8499): Disabled cache when interpreter is enabled\r\n- **Backend Detection** (#8046): Added env var to speed up backend detection in tree\r\n\r\n### Compiler Pipeline\r\n- **Configurable Pass Pipeline** (#8137): Added hook for configurable/overridable compiler pass pipeline\r\n- **MLIR Reproducer Retention** (#8113): Retain mlir reproducer temporaries from prior run pass pipelines\r\n- **MLIR Multithreading Disable** (#8255): Disabled MLIR multithreading\r\n- **SCF to CF Inliner** (#8017): Run the inliner after scf-to-cf\r\n\r\n### CUDA Updates\r\n- **PTXAS Upgrade** (#8476): Upgraded ptxas to 12.9.86 for Blackwell\r\n- **CUDA 13 CRT Headers** (#8336): Fixed crt header download location for CUDA >= 13\r\n- **ptxas_options Knobs** (#8121): Updated ptxas_options knobs default value\r\n\r\n### AOT Compilation\r\n- **Gluon Kernel Compilation** (#8660): Support compile gluon kernels in compile.py\r\n\r\n### Interpreter\r\n- **TRITON_INTERPRET Cleanup** (#8735, #8736): Made TRITON_INTERPRET cleanup after itself with improvements\r\n- **Tensor Descriptor Stride Validation** (#8670): Fixed tensor descriptor stride validation\r\n- **Histogram Silent Corruption** (#8550): Fixed silent data corruption in histogram\r\n- **TensorHandle Dtype Validation** (#8594): Validated TensorHandle np/tl dtypes size\r\n- **Pre-run Hooks** (#8573): Enabled pre-run hooks in interpreter mode\r\n\r\n---\r\n\r\n## Documentation\r\n\r\n### Community Meetup Notes\r\n- **2025-09-03** (#8178): Adding meeting notes for 2025-09-03 community meetup\r\n- **2025-11-05** (#8727): Added meeting notes for 2025-11-05 community meetup\r\n\r\n### Technical Documentation\r\n- **dot_scaled Requirements** (#8433): Clarified lhs_scale and rhs_scale requirements in dot_scaled\r\n- **Install Command Fix** (#8271): Fixed install command in tutorials README.rst\r\n- **Gluon Tutorial Fix** (#8593): Fixed gluon tutorial example\r\n- **Gluon Layout Explanation** (#8020): Fixed description in layout explanation in gluon tutorial\r\n- **Proton README** (#8319): Updated Proton README\r\n- **Proton Tutorial** (#8334): Intra kernel profiling tutorial and examples\r\n- **Tutorial Units** (#8631): Added units to result tables in tutorials\r\n- **AMD Scaled Matmul Tutorial** (#8099): Added AMD GPUs in scaled matmul tutorial\r\n\r\n### README Updates\r\n- **Triton Conference 2025** (#8186): Added Triton Conference 2025 details to README\r\n- **Conference Registration** (#8114): Added conference registration link\r\n\r\n---\r\n\r\n## Breaking Changes\r\n\r\n### API Changes\r\n- **Constexpr Through min/max** (#8733): BC-breaking propagation of constexpr through builtin min/max\r\n- **Aggregate Cache Keys** (#8568): Aggregate members are now added to the cache key\r\n- **warp_specialize Argument Tuples** (#8368): Required warp_specialize default_args and worker_args to be tuples\r\n- **warp_specialize API Change** (#8527): Changed `gl.warp_specialize` API\r\n\r\n### Proton Changes\r\n- **Metric ValueId Types** (#7979): BC-break - Prevent updating the same metric valueId with different types\r\n\r\n### Removed Features\r\n- **Python 3.9 Support** (#8222): Removed Python 3.9 support, minimum is now 3.10\r\n- **CentOS 7 Build** (#8191): Removed CentOS 7 build support\r\n- **GlobalPrefetch/LocalPrefetch Knobs** (#8295): Removed GlobalPrefetch and LocalPrefetch Knobs for AMD\r\n\r\n### Deprecations\r\n- **triton_kernels.routing** (#8375): Deprecated triton_kernels.routing in favor of BitmatrixMetadata\r\n- **Custom Topological Sort** (#8596): Deprecated triton's custom topological sort\r\n\r\n---\r\n\r\n## Performance Improvements\r\n\r\n### Compilation Performance\r\n- **Native Specialization** (#7771): Significant launch latency improvements through native specialize\r\n- **ConSan Compilation Time** (#8689): Improved compilation time in constant sanitizer\r\n\r\n### Runtime Performance\r\n- **MXFP Conversions** (#8610): Speedup for mxfp conversions\r\n- **FP4->BF16 Conversion** (#8145): Optimized fp4->bf16 conversion for MI300\r\n- **Permlane Swap** (#7947): Use permlane_swap for efficient layout conversions\r\n- **Chained WMMA** (#7374): Optimized chained multiplications for WMMA\r\n- **Expert Parallelism** (#8448): New expert parallelism implementation\r\n\r\n### Memory Optimizations\r\n- **BypassLDS** (#7968): Added bypassLDS feature to skip LDS when possible\r\n- **Padded Layout Selection** (#8053): Redesigned stream pipeliner LDS layout selection\r\n\r\n---\r\n\r\n## Notable Bug Fixes\r\n\r\n### Correctness Issues\r\n- **Loop Induction Variable** (#8750): Fixed modification of for loop induction variable\r\n- **Store Broadcasting** (#8661): Fixed broadcasting in store operations\r\n- **64-bit Atomic CAS** (#8105): Fixed 64-bit atomic_cas\r\n- **Histogram Corruption** (#8550): Fixed silent data corruption in histogram\r\n- **MXFP Overflow** (#8356): Handle values close to max correctly without overflow\r\n\r\n### Crash Fixes\r\n- **Pointer Canonicalization** (#8465): Fixed ptr-canonicalization segmentation fault\r\n- **SmallVector Crash** (#8544): Fixed SmallVector crash issue in AxisInfoAnalysis\r\n- **ASAN Issues** (#8117): Fixed ASAN initialization-order-fiasco\r\n\r\n### Regression Fixes\r\n- **Batched Block Sizes Reapply** (#8084): Reapplied improved block sizes after fixes\r\n- **Native MXFP FP8 Reapply** (#8129): Reapplied native MXFP FP8 scaled_dot for SM120\r\n\r\n---\r\n\r\n## Experimental Triton to Gluon Translator\r\n\r\n- **Translator Tool** (#8417): Added experimental translator from Triton to Gluon for easier migration\r\n\r\n---\r\n\r\n## Contributors\r\n\r\nThis release includes contributions from engineers at:\r\n- Meta\r\n- AMD\r\n- NVIDIA\r\n- OpenAI\r\n- Intel\r\n- Google\r\n- And many individual contributors\r\n\r\nSpecial thanks to all contributors who submitted bug reports, feature requests, and code improvements!\r\n","reactions":{"url":"https://api.github.com/repos/triton-lang/triton/releases/278767719/reactions","total_count":12,"+1":0,"-1":0,"laugh":0,"hooray":2,"confused":0,"heart":9,"rocket":1,"eyes":0}},{"url":"https://api.github.com/repos/triton-lang/triton/releases/261919265","assets_url":"https://api.github.com/repos/triton-lang/triton/releases/261919265/assets","upload_url":"https://uploads.github.com/repos/triton-lang/triton/releases/261919265/assets{?name,label}","html_url":"https://github.com/triton-lang/triton/releases/tag/v3.5.1","id":261919265,"author":{"login":"atalman","id":7563158,"node_id":"MDQ6VXNlcjc1NjMxNTg=","avatar_url":"https://avatars.githubusercontent.com/u/7563158?v=4","gravatar_id":"","url":"https://api.github.com/users/atalman","html_url":"https://github.com/atalman","followers_url":"https://api.github.com/users/atalman/followers","following_url":"https://api.github.com/users/atalman/following{/other_user}","gists_url":"https://api.github.com/users/atalman/gists{/gist_id}","starred_url":"https://api.github.com/users/atalman/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/atalman/subscriptions","organizations_url":"https://api.github.com/users/atalman/orgs","repos_url":"https://api.github.com/users/atalman/repos","events_url":"https://api.github.com/users/atalman/events{/privacy}","received_events_url":"https://api.github.com/users/atalman/received_events","type":"User","user_view_type":"public","site_admin":false},"node_id":"RE_kwDOAWaHXs4PnJIh","tag_name":"v3.5.1","target_commitish":"main","name":"Triton 3.5.1 release, bug fix release","draft":false,"immutable":false,"prerelease":false,"created_at":"2025-11-11T16:12:26Z","updated_at":"2025-11-12T22:45:34Z","published_at":"2025-11-12T22:45:12Z","assets":[{"url":"https://api.github.com/repos/triton-lang/triton/releases/assets/315704165","id":315704165,"node_id":"RA_kwDOAWaHXs4S0UNl","name":"triton-3.5.1.tar.gz","label":"","uploader":{"login":"github-actions[bot]","id":41898282,"node_id":"MDM6Qm90NDE4OTgyODI=","avatar_url":"https://avatars.githubusercontent.com/in/15368?v=4","gravatar_id":"","url":"https://api.github.com/users/github-actions%5Bbot%5D","html_url":"https://github.com/apps/github-actions","followers_url":"https://api.github.com/users/github-actions%5Bbot%5D/followers","following_url":"https://api.github.com/users/github-actions%5Bbot%5D/following{/other_user}","gists_url":"https://api.github.com/users/github-actions%5Bbot%5D/gists{/gist_id}","starred_url":"https://api.github.com/users/github-actions%5Bbot%5D/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/github-actions%5Bbot%5D/subscriptions","organizations_url":"https://api.github.com/users/github-actions%5Bbot%5D/orgs","repos_url":"https://api.github.com/users/github-actions%5Bbot%5D/repos","events_url":"https://api.github.com/users/github-actions%5Bbot%5D/events{/privacy}","received_events_url":"https://api.github.com/users/github-actions%5Bbot%5D/received_events","type":"Bot","user_view_type":"public","site_admin":false},"content_type":"application/gzip","state":"uploaded","size":6243015,"digest":"sha256:4afb7e2af6ab4585ea923eb5054ea1748989d3a4b5b2e035d239723924c8a7fa","download_count":940,"created_at":"2025-11-12T22:45:33Z","updated_at":"2025-11-12T22:45:34Z","browser_download_url":"https://github.com/triton-lang/triton/releases/download/v3.5.1/triton-3.5.1.tar.gz"}],"tarball_url":"https://api.github.com/repos/triton-lang/triton/tarball/v3.5.1","zipball_url":"https://api.github.com/repos/triton-lang/triton/zipball/v3.5.1","body":"This release is meant to fix the following issue:\r\n\r\nFix sm103 (GB300) support broken by Triton 3.5.0 release (https://github.com/triton-lang/triton/pull/8045)","reactions":{"url":"https://api.github.com/repos/triton-lang/triton/releases/261919265/reactions","total_count":10,"+1":10,"-1":0,"laugh":0,"hooray":0,"confused":0,"heart":0,"rocket":0,"eyes":0}},{"url":"https://api.github.com/repos/triton-lang/triton/releases/255856323","assets_url":"https://api.github.com/repos/triton-lang/triton/releases/255856323/assets","upload_url":"https://uploads.github.com/repos/triton-lang/triton/releases/255856323/assets{?name,label}","html_url":"https://github.com/triton-lang/triton/releases/tag/v3.5.0","id":255856323,"author":{"login":"atalman","id":7563158,"node_id":"MDQ6VXNlcjc1NjMxNTg=","avatar_url":"https://avatars.githubusercontent.com/u/7563158?v=4","gravatar_id":"","url":"https://api.github.com/users/atalman","html_url":"https://github.com/atalman","followers_url":"https://api.github.com/users/atalman/followers","following_url":"https://api.github.com/users/atalman/following{/other_user}","gists_url":"https://api.github.com/users/atalman/gists{/gist_id}","starred_url":"https://api.github.com/users/atalman/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/atalman/subscriptions","organizations_url":"https://api.github.com/users/atalman/orgs","repos_url":"https://api.github.com/users/atalman/repos","events_url":"https://api.github.com/users/atalman/events{/privacy}","received_events_url":"https://api.github.com/users/atalman/received_events","type":"User","user_view_type":"public","site_admin":false},"node_id":"RE_kwDOAWaHXs4PQA7D","tag_name":"v3.5.0","target_commitish":"main","name":"Triton 3.5.0 release","draft":false,"immutable":false,"prerelease":false,"created_at":"2025-10-13T15:10:54Z","updated_at":"2025-10-21T10:02:53Z","published_at":"2025-10-21T10:00:38Z","assets":[{"url":"https://api.github.com/repos/triton-lang/triton/releases/assets/306835712","id":306835712,"node_id":"RA_kwDOAWaHXs4SSfEA","name":"triton-3.5.0.tar.gz","label":"","uploader":{"login":"github-actions[bot]","id":41898282,"node_id":"MDM6Qm90NDE4OTgyODI=","avatar_url":"https://avatars.githubusercontent.com/in/15368?v=4","gravatar_id":"","url":"https://api.github.com/users/github-actions%5Bbot%5D","html_url":"https://github.com/apps/github-actions","followers_url":"https://api.github.com/users/github-actions%5Bbot%5D/followers","following_url":"https://api.github.com/users/github-actions%5Bbot%5D/following{/other_user}","gists_url":"https://api.github.com/users/github-actions%5Bbot%5D/gists{/gist_id}","starred_url":"https://api.github.com/users/github-actions%5Bbot%5D/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/github-actions%5Bbot%5D/subscriptions","organizations_url":"https://api.github.com/users/github-actions%5Bbot%5D/orgs","repos_url":"https://api.github.com/users/github-actions%5Bbot%5D/repos","events_url":"https://api.github.com/users/github-actions%5Bbot%5D/events{/privacy}","received_events_url":"https://api.github.com/users/github-actions%5Bbot%5D/received_events","type":"Bot","user_view_type":"public","site_admin":false},"content_type":"application/gzip","state":"uploaded","size":6243115,"digest":"sha256:a812f387de09277623ccd58af477f2ae9fc6e028ccfadf4b13919af441b59928","download_count":404,"created_at":"2025-10-21T10:01:05Z","updated_at":"2025-10-21T10:01:05Z","browser_download_url":"https://github.com/triton-lang/triton/releases/download/v3.5.0/triton-3.5.0.tar.gz"}],"tarball_url":"https://api.github.com/repos/triton-lang/triton/tarball/v3.5.0","zipball_url":"https://api.github.com/repos/triton-lang/triton/zipball/v3.5.0","body":"# Triton Release Notes\r\n\r\n## Table of Contents\r\n- [Dialect & Frontend](#dialect--frontend)\r\n- [Backend & Compiler](#backend--compiler)\r\n- [AMD/HIP Backend](#amdhip-backend)\r\n- [NVIDIA Backend](#nvidia-backend)\r\n- [Gluon & Layout Improvements](#gluon--layout-improvements)\r\n- [Kernels & Benchmarks](#kernels--benchmarks)\r\n- [Testing & CI](#testing--ci)\r\n- [Build & Infrastructure](#build--infrastructure)\r\n- [Documentation](#documentation)\r\n- [Breaking Changes](#breaking-changes)\r\n\r\n---\r\n\r\n## Dialect & Frontend\r\n\r\n### New Features\r\n- **Warp Specialization Enhancements** (#8005): Made warp specialization require at least 4 warps with proper error messaging to prevent compiler crashes\r\n- **Ragged TMA Support** (#7792, #7783): Added support for write-only and general ragged TMAs with automatic bounds checking using higher-dimensional TMA descriptors\r\n- **Device Assert Mask Support** (#7905): Added `mask` parameter to `tl.device_assert` for easier debugging with masked operations\r\n- **Padding Option for TMA Loads** (#7993): Added support for padding option (including NaN) in TMA descriptor creation and fallback paths\r\n- **Implicit Downcast in TMA Descriptor Store** (#6236): Fixed missing implicit downcast when storing blocks through TMA descriptors\r\n- **Mutations Disallowed** (#7762): Disabled all mutations to address semantic issues in the language\r\n- **Specialized Recursion** (#7468): Enabled functions to recurse on specialized versions of themselves\r\n- **Constexpr Function Cache Invalidation** (#7802): Reworked `constexpr_function` to support cache invalidation and capability checks\r\n\r\n### Bug Fixes\r\n- **Floating Point Argument Passing** (#7439): Fixed floating point argument passing for `tl.float16` and other FP types\r\n- **Non-Associative Reduce Rematerialization** (#7272): Avoided rematerialization for non-associative reduce operations to prevent data consistency issues\r\n- **PDL Issue Fix** (#7379): Fixed PDL-related issues in the frontend\r\n- **Constexpr in Tuples** (#7442): Improved handling of constexpr in tuples, fixing type mismatches and in-place mutations\r\n- **Loop Carry Detection** (#7200): Improved detection of loop carries when `@builtin` or `@core.extern` functions modify their arguments\r\n- **Liveouts in Conditionals** (#7318): Fixed detection of liveouts in conditional blocks\r\n\r\n### Improvements\r\n- **MLIR Verifier After Parsing** (#7999): Run MLIR verifier after parsing to catch errors early\r\n- **Better Error for num_cta > 1 on sm < 90** (#7812): Improved error messaging for unsupported configurations\r\n- **Extern Elementwise Type Handling** (#7930): Fixed mismatched type handling for `core.extern_elementwise`\r\n- **Libdevice Exposure in Gluon** (#7890): Exposed libdevice functions with improved layout propagation\r\n\r\n---\r\n\r\n## Backend & Compiler\r\n\r\n### LLVM Updates\r\n- **LLVM Bump** (#7881): Updated to llvm/llvm-project@bc773632355b with multiple API changes including:\r\n  - Switched `Constant{Int|Float}Op` type and value order\r\n  - Provided triple for `TargetLibraryInfoImpl`\r\n  - Fixed atomic sync scope for NVIDIA\r\n  - Updated MLIR lib names and ops\r\n\r\n### Code Generation\r\n- **Generic Swizzling for convert_layout** (#6982, #7565): Implemented generalized swizzling algorithm for `convert_layout` that:\r\n  - Finds optimal shared memory layout maximizing read/write vectorization\r\n  - Minimizes bank conflicts\r\n  - Supports `ldmatrix/stmatrix` and transpose versions\r\n  - Uses columns and diagonals for better performance\r\n- **Warp-Local Layout Conversion** (#7558): Improved warp-local layout conversion algorithm using shuffles with:\r\n  - Better handling of broadcasting in layouts\r\n  - Fewer `select` and `shuffle` instructions\r\n  - Register packing for sub-32-bit data types\r\n- **Byte Permutes in Intra-Warp Conversion** (#7809): Used byte permute instructions for better performance in layout conversions\r\n- **Tmem Alloc Hoisting** (#7568): Hoisted tmem alloc outside of if statements to reduce register pressure\r\n- **CP.Async Lowering Improvements** (#7314): Moved cp.async to better lowering sequence reusing previous optimizations\r\n\r\n### Optimizations\r\n- **Simpler Codegen for Linear Layouts** (#7201): Simplified code generation for linear layouts\r\n- **Vectorization Fixes** (#7845): Fixed vectorization for `PaddedSharedEncoding` with non-default order\r\n- **XOR Trick Refactoring** (#7397): Refactored XOR trick into helper function for better code reuse\r\n- **Shared Memory Offset Fixes** (#7949): Fixed various issues with smem base offsets\r\n- **Min/Max Redux Optimization for Blackwell** (#7465): Implemented new redux.sync optimization\r\n\r\n### Bug Fixes\r\n- **Atomic RMW Broadcasting** (#7460): Fixed atomic rmw ops to broadcast results when necessary\r\n- **TMA Load with Multiple Users** (#7398): Fixed lowering of TMA load when users have differing encodings\r\n- **Subview Padding** (#7404): Fixed subview padding for PaddedSharedEncoding\r\n- **Memdesc Subview Fixes** (#7480, #7515): Properly handled memdesc_subview with slicing and offsets\r\n- **FP16 to FP32 Conversion** (#7585): Fixed fp16 to fp32 conversion issues\r\n- **Barrier Synchronization** (#7993): Added bar.sync before deallocating tmem to prevent race conditions\r\n\r\n---\r\n\r\n## AMD/HIP Backend\r\n\r\n### New Features\r\n- **GFX950 (MI350) Support**: Added comprehensive support for AMD's latest architecture including:\r\n  - MFMA scale support (#7799)\r\n  - Scale preshuffling (#7603, #7836)\r\n  - OpSel implementation for scaled MFMA\r\n  - Buffer load/store operations (#7738)\r\n  - Improved register usage in Float8 conversions (#7527)\r\n- **ChainedDot Schedule** (#7601, #7638): Added new scheduling variant for loops with 2 chained dots\r\n- **Ping-Pong Transformation** (#7638, #7458): Added ping-pong support for:\r\n  - Chained dot schedules\r\n  - Async load with num_stages=3\r\n  - MXFP types\r\n- **Buffer Atomic CAS** (#7292): Added support for buffer atomic compare-and-swap\r\n- **FP64 MFMA Support** (#7461): Added support for fp64 dot operations using MFMA intrinsics\r\n\r\n### Layout & Memory Optimizations\r\n- **General Swizzling Support** (#7482, #7606): Enabled ConvertLayoutOp general swizzling\r\n- **Padded vs Swizzled Allocation** (#7328, #7750): Introduced specialized allocation pass with proper layout selection strategy\r\n- **Improved LDS Usage** (#7750, #7813): Optimized LDS usage by:\r\n  - Preferring swizzle layouts when LDS limits allow\r\n  - Using single LDS for both transposed and non-transposed access\r\n  - Better layout selection in optimize-lds-usage pass\r\n- **TilesPerWarp Parameter** (#7283): Added tilesPerWarp parameter to MFMA layout for contiguous tile computation\r\n- **Extract Slice Rewrite** (#7128): Refactored extract_slice to support:\r\n  - Arbitrary tensor ranks\r\n  - Relaxed layout constraints\r\n  - CTA tile boundary alignment\r\n\r\n### Code Generation Improvements\r\n- **PermlaneSwap Pattern** (#7825, #7861): Added general permlane_swap pattern for ConvertLayoutOp\r\n- **Register Broadcast** (#7407): Added support for register broadcast in slice/concat ops\r\n- **Shared Memory Ops for FP4** (#7626): Added support for M/N packed FP4 with transposition\r\n- **Direct-to-LDS Loads** (#7829): Refactored lowering via common `lowerLdSt` path\r\n- **Local Load/Store Lowering** (#7355): Enabled common code path for local_load/store operations\r\n\r\n### FP8 & Numeric Support\r\n- **FP8 Variant Support**:\r\n  - Software emulation for non-gfx942 architectures (#7401)\r\n  - Improved conversions with proper clamping (#7337, #7361, #7363)\r\n  - BF16 to OCP FP8 conversion on CDNA3 (#7469)\r\n  - Float8E4M3FN emulation on CDNA3 and below (#7186)\r\n- **Dot Scaled Support**: Enabled on gfx11 (#7954) and gfx12 (#7644) with emulation via decomposition\r\n- **True16 Handling**: Disabled on gfx11 due to test failures (#7953)\r\n\r\n### Stream Pipeliner Enhancements\r\n- **Refactoring** (#7526, #7556): Refactored to use more common pipeliner functionality\r\n- **Async Wait Handling** (#7577): Restricted merging async_wait when pipelining with num_stages=3\r\n- **Mask Operation Support** (#7620): Added ttg.mask handling in stream pipeliner\r\n\r\n### Build & Driver\r\n- **LLD Library API** (#7548): Replaced shell-out to lld with direct library API calls\r\n- **hipGetProcAddress** (#7350): Switched to using hipGetProcAddress for querying HIP symbols\r\n- **Driver Version Check** (#7501): Added runtime driver version check with descriptive errors\r\n- **AOT Compilation** (#7007): Added HIP AOT compilation support to compile.py tool\r\n\r\n### Bug Fixes\r\n- **Pointer Canonicalizer** (#7242): Fixed attribute propagation when ranks don't match\r\n- **Global Atomic Optimization** (#7496): Optimized global atomic operations following memory model semantics\r\n- **FP32/FP16 to OCP FP8** (#7382): Fixed conversion for subnormal numbers\r\n- **Async Copy Vectorization** (#7250): Fixed async load pipeline for less than 32-bit loads\r\n- **OptimizeLDSUtility Crash** (#7434): Fixed nullptr crash in createTmpLayout\r\n- **Memrealtime on GFX11/12** (#7357): Added proper support using s_sendmsg_rtn_b64\r\n\r\n---\r\n\r\n## NVIDIA Backend\r\n\r\n### Hopper/Blackwell Features\r\n- **Warp Specialization**:\r\n  - Enable for persistent matmul and FA (#7642, #7623)\r\n  - Assign final try_wait to partition (#7757)\r\n  - Tightened user critical section with accumulator (#7509)\r\n  - Fixed rematerialization bug in partitioner (#7427)\r\n  - Optimized partitioning by hoisting above broadcasts (#7692)\r\n  - Enabled 1 buffer for SSA partition dependencies (#7686)\r\n  - Control flow support in TMEM allocation (#7698)\r\n- **WGMMA Support in Gluon** (#7300, #7313): Added Hopper WGMMA with async wait support\r\n- **Aref Operations** (#7479, #7561, #7645): Updated aref ops and lower_aref pass with:\r\n  - Multi-consumer support\r\n  - Stage/cluster attribute passing\r\n  - TMA load aref insertion\r\n  - Control flow handling\r\n- **Partition Loops Rewrite** (#7415): Reimplemented supporting general control flow using mutual recursion\r\n\r\n### Blackwell-Specific\r\n- **TMEM Support**:\r\n  - Fixed codegen for Nx1xf32 (#7234)\r\n  - Fixed tmem_subslice for packed layouts (#7207)\r\n  - Allowed splitting block_m=64 along N (#7589)\r\n  - Tcgen05_copy exposure (#7936)\r\n  - Generic lowering for tcgen05.ld/st (#7831, #7874)\r\n- **Tcgen05.commit Op** (#7335): Added separate commit op for better persistent kernel support\r\n- **Subtile QK TMEM Load** (#7655): Improved non-causal fp8 performance by 40-50 TFLOPS\r\n\r\n### MMA Improvements\r\n- **FP64 SIMT FMA** (#7310): Added fp64 simt fma support and fp64 mma for SM80/SM90\r\n- **FP8 MMAv2** (#7409): Don't promote fp8 MMAv2 dot inputs for sm120 (~1.9x speedup)\r\n- **Min Dot Sizes Update** (#7411): Reduced minimum dot sizes (e.g., N=16 to lower values)\r\n- **NVVM Op Migration** (#7420, #7471, #7512): Replaced inline assembly with NVVM ops for:\r\n  - WGMMAFenceOp, WGMMACommitGroupOp, ClusterWaitOp\r\n  - ClusterCTAIdOp conversion\r\n  - Better optimization opportunities\r\n\r\n### Other Enhancements\r\n- **Bar.warp.sync for 1 Warp** (#7336): Emit more efficient bar.warp.sync for single-warp barriers\r\n- **L2 Cache Hints** (#7219): Limited L2 cache hints to sm >= 80\r\n- **Cublas.gemm Exposure** (#7656): Exposed cublas.gemm for performance testing\r\n- **PTX Workarounds**:\r\n  - Matrix descriptor arithmetic (#7197)\r\n  - TMA device-side descriptor race condition (#7293)\r\n  - Byte permutes ptxas bug (#7933)\r\n\r\n---\r\n\r\n## Gluon & Layout Improvements\r\n\r\n### Gluon Language Features\r\n- **AutoLayout** (#7447, #7466): Added AutoLayout for backward layout inference with:\r\n  - Custom layout inference interface\r\n  - Propagation through operations\r\n  - Conflict detection and error reporting\r\n  - `assert_trivial` flag for performance validation\r\n- **Docstrings** (#7323): Added comprehensive documentation for public Gluon API\r\n- **API Improvements**:\r\n  - Added `numel` and `nbytes` properties (#7507)\r\n  - Added `map_elementwise` (#7564)\r\n  - Fixed `tensor.sum` (#7617)\r\n  - Fixed `splat` returning auto encoding (#7490)\r\n  - Fixed auto encoding inconsistencies (#7726)\r\n  - Added constexpr_function and static_range (#7531)\r\n\r\n### Layout System\r\n- **Padded Shared Layout** (#7212): Added new shared memory layout for padding\r\n- **Slice Encoding for SplitOp** (#7247): Improved slice encoding inference\r\n- **LinearLayout Improvements**:\r\n  - Implemented toLinearLayout for TensorMemoryEncodingAttr (#7748)\r\n  - Fixed split op backward propagation (#7340)\r\n  - Generalized getShapePerCTA (#7580)\r\n  - Fixed memdesc reshape encoding inference (#7544)\r\n- **NVIDIA Shared Layout Improvements**:\r\n  - Added NVMMASharedLayout constructor with default swizzle (#7534)\r\n  - Fixed handling of non-default order (#7845)\r\n- **DotOperandLayout Exposure** (#7730): Exposed for WGMMA with LHS in registers\r\n\r\n### Tutorial & Examples\r\n- **Attention Kernels** (#7009, #7298, #7488): Implemented complete attention for d64 and d128 with:\r\n  - Persistent kernel support\r\n  - Causal masking optimization\r\n  - FADD2 for row_sum computation (D64)\r\n  - FFMA2 for QK scale\r\n  - Turnstile over exp2 to control MFU access\r\n  - Subtiling optimizations\r\n  - 100-120 TFLOPS improvement for D64\r\n- **Tutorials Added** (#7657): Comprehensive set covering basic to advanced optimized techniques\r\n\r\n---\r\n\r\n## Kernels & Benchmarks\r\n\r\n### MXFP Support\r\n- **Naming Fixes** (#7870): Changed dequantize to quantize, matched arg names\r\n- **FP32 Support** (#7672): Added quant/dequant from/to fp32\r\n- **Transposed Weight Support** (#7795): Handle both transposed and non-transposed mxfp weights\r\n- **Act In/Out Matmul** (#7598): Added mxfp act input and output support\r\n- **Blackwell Value Padding** (#7958): Fixed mxfp value padding for Blackwell\r\n- **Test Coverage** (#7591): Added missing mxfp4 tests\r\n- **MXFP_BLOCK_SIZE Constant** (#7567): Added constant for better readability\r\n- **Empty Tensor Handling** (#7579): Fixed handling of empty tensors in downcast_to_mxfp\r\n\r\n### Matmul Optimizations\r\n- **Significant Cleanup** (#7882): Major refactoring of matmul_ogs.py with:\r\n  - More efficient post-processing\r\n  - Better intelligibility\r\n  - Improved documentation\r\n- **Heuristics Improvements** (#7664, #7632): Tweaked block sizes and heuristics\r\n- **Block Size Improvements** (#7897): Better block sizes for batched matmul_ogs with small m/n/k\r\n- **Swizzling Fixes** (#7582, #7587): Fixed swizzling numerics and contiguity\r\n- **Host TMA Usage** (#7182): Increased use of host TMA for X, W, Mx scales\r\n- **Bias Subtiling** (#7232): Added then reverted bias subtiling changes due to regression\r\n- **Zero Elements Support** (#7808): Added support for inputs with 0 elements\r\n- **Index Casting** (#7794): Cast index to int64 to avoid overflow\r\n\r\n### MoE & Multi-GPU\r\n- **Routing Improvements** (#7369): 30% performance improvement through:\r\n  - Liberal kernel fusion (7 launches → 4 launches)\r\n  - Specific kernel optimizations\r\n  - FP32 logits: 22.8us → 18.0us\r\n  - FP16 logits: 17.3us → 12.2us\r\n- **BitMatrix for Routing** (#7789): Used bitmatrix for distributed routing\r\n- **Simple Multi-GPU MoE** (#7352): Initialized baseline implementation\r\n\r\n### Benchmarks & Tests\r\n- **Launch Overhead** (#7849): Added microbenchmark to track dispatch overhead\r\n- **Total Time Computation** (#7752): Added total kernel time computation\r\n- **Roofline Fixes** (#7670): Fixed roofline plots for compute-bound kernels\r\n- **MLP Fixes** (#7926): Fixed bench_mlp.py for various issues\r\n\r\n---\r\n\r\n## Testing & CI\r\n\r\n### Test Infrastructure\r\n- **Fresh Knobs Usage** (#7687): Use fresh_knobs when touching triton.knobs\r\n- **Environment Variable Restore** (#7807): Fixed monkey patching for proper cleanup\r\n- **Test Cleanup** (#7801): Various cleanups on test_core.py including:\r\n  - Changed kernel launch to warmup\r\n  - Moved tuple tests to test_tuple.py\r\n  - Improved error handling\r\n- **Reduce Test Overflow Fix** (#7470): Limited integer range to avoid overflow\r\n- **Input Generation Consolidation** (#7477): Consolidated input generation for reduce tests\r\n\r\n### AMD Testing\r\n- **GFX950 CI** (#7189): Enabled CI for GFX950\r\n- **GPU Isolation** (#7650): Added env-file for better GPU isolation in CI\r\n- **Passing Tests** (#7363, #7365, #7236, #7183): Enabled many passing tests for AMD GFX942 and GFX12\r\n- **Test Skipping**: Properly skipped flaky or unsupported tests (globaltimer, True16, etc.)\r\n\r\n### Lit Tests\r\n- **Subfolder Test Fixes** (#7966): Fixed lit tests failing when run via ninja check-triton-lit-tests-<folder>\r\n- **LLD Configuration** (#7992): Added llc to lit tool configuration\r\n- **Test Updates**: Updated numerous lit tests for new features and bug fixes\r\n\r\n### NVIDIA Testing\r\n- **GB200 Error Handling** (#7537): Continue running CI when GB200 errors out\r\n- **Warp Specialization Tests** (#7623, #7642): Enabled WS tests for Hopper\r\n\r\n---\r\n\r\n## Build & Infrastructure\r\n\r\n### Build System\r\n- **Out-of-Tree Build** (#7347, #7871): Enabled complete out-of-tree build with TRITON_BUILD_DIR\r\n- **Compile Commands Symlink** (#7305, #7341): Symlink compile_commands.json to root for better IDE support\r\n- **Elapsed Time for MacOS** (#7559): Added elapsed time logging to MacOS builds\r\n- **LLD for MacOS** (#7559): Enabled LLD for macOS build to reduce time\r\n- **Clang Warning Fixes** (#7868): Fixed warnings to build triton with clang\r\n- **Debug Build Default** (#7872): Changed default LLVM build to release\r\n\r\n### Dependencies & Environment\r\n- **Custom LLVM Build** (#7279, #6709): Mentioned `make dev-install-llvm` in README\r\n- **Python 3.14 Wheels** (#7695): Added Python 3.14 wheel build support\r\n- **Setuptools Removal** (#7983): Removed setuptools requirement from setup.py\r\n- **New CUDA Versions** (#7384): Automatically handle newer CUDA versions\r\n- **LLVM System Suffix** (#7430): Added TRITON_LLVM_SYSTEM_SUFFIX for user-specified prebuilt LLVM\r\n\r\n### Runtime & Compilation\r\n- **Async Compile Mode** (#7306): Added AsyncCompileMode to build multiple kernels in parallel\r\n- **Thread-Safe Allocator** (#7685): Made set_allocator thread-safe using ContextVar\r\n- **AsyncCompileMode Thread Safety** (#7701): Made AsyncCompileMode thread safe\r\n- **Kernel Caching**:\r\n  - Cache invalidation for constexpr_function (#7802)\r\n  - Fixed cache key computation thread safety (#7974)\r\n  - Include constexprs in cache keys (#7348)\r\n  - NvidiaTool.from_path caching (#7569)\r\n\r\n### Driver & Backend\r\n- **Host Compiler Flags** (#7659): Allow backend-provided runtime host compiler flags\r\n- **NVIDIA Driver Improvements** (#7769): Slightly improved NVIDIA driver backend with string caching\r\n- **HIP Driver Updates**:\r\n  - Fixed hipError discard (#7832)\r\n  - Fixed multiple compiler warnings (#7838)\r\n  - Fixed undefined behavior (#7806)\r\n\r\n---\r\n\r\n## Documentation\r\n\r\n### New Documentation\r\n- **Community Meetup Notes**:\r\n  - 2025-03-12 (#7255)\r\n  - 2025-05-01 (#7256)\r\n  - 2025-07-09 (#7788)\r\n- **Moderators Guide** (#7787): Updated with YouTube info and event creation\r\n- **Installation Instructions** (#7572): Updated install instructions in docs\r\n- **README Improvements** (#7368): Improved readability and fixed minor issues\r\n- **Running Meetups** (#7103): Added documentation for running Triton Community Meetups\r\n\r\n### Code Quality\r\n- **CODEOWNERS Updates**:\r\n  - Linear Layouts (#7754)\r\n  - Gluon section (#7744)\r\n  - Proton Backend (#7782)\r\n- **NFC Refactorings**: Multiple no-functional-change refactorings for better code organization\r\n- **Unused Code Removal** (#7703): Killed unused functions throughout codebase\r\n\r\n---\r\n\r\n## Breaking Changes\r\n\r\n### API Changes\r\n- **Mutations Disallowed** (#7762): All mutations are now disabled in the language\r\n- **Min Dot Sizes** (#7411, #7451): Relaxed minimum dot size requirements (may affect autotuning)\r\n- **Constexpr Handling**: Changed how constexprs are included in cache keys (#7348)\r\n- **Environment Variables**:\r\n  - Only check `TRITON_DEBUG` at import time (#7767)\r\n  - Removed getattr overhead from DriverConfig and CompiledKernel (#7770)\r\n- **Hook System Changes**:\r\n  - Renamed and changed signature for kernel load hooks (#7834)\r\n  - More generalized hooking system (#7866)\r\n\r\n### Removed Features\r\n- **Python 3.9 Support** (#8222, #8287): Cleaned up Python 3.9 related code\r\n- **Nightly Installation**: Removed from documentation\r\n- **Local Prefetch Schedule** (#7395): Retired AMD local prefetch schedule hint variant\r\n\r\n### Deprecations\r\n- **Warp Size**: Removed hardcoded warp size assumptions (#7253)\r\n- **GetShapePerCTA**: Moving toward elimination in AMD backend (#7740)\r\n\r\n---\r\n\r\n## Performance Improvements\r\n\r\n### Measured Improvements\r\n- **Attention Kernels**: Up to 785 TFLOPS for D64, 1230 TFLOPS for D128\r\n- **MoE Routing**: 30% faster (17.3us → 12.2us for fp16)\r\n- **FP8 on Blackwell**: ~1.9x speedup for large matmuls\r\n- **Launch Overhead**: Reduced by various optimizations (DriverConfig cleanup, etc.)\r\n- **Compile Time**: ~20% savings by skipping link_extern_libs when unnecessary (#7570)\r\n\r\n### Optimization Techniques\r\n- **Register Pressure**: Better management through tmem alloc hoisting\r\n- **Vectorization**: Improved through generic swizzling and layout optimizations\r\n- **Bank Conflicts**: Minimized through optimized shared memory layouts\r\n- **Instruction Scheduling**: Better code generation for linear layouts\r\n\r\n---\r\n\r\n## Notable Bug Fixes\r\n\r\n### Correctness Issues\r\n- **Non-Associative Reduce** (#7272): Fixed rematerialization causing incorrect results\r\n- **Atomic Operations** (#7460): Fixed broadcasting for atomic_cas and rmw operations\r\n- **Memory Model**: Multiple fixes for proper fence insertion and synchronization\r\n- **FP8 Conversions**: Fixed numerous rounding and clamping issues\r\n- **TMA Operations**: Fixed various edge cases in TMA load/store\r\n\r\n### Crash Fixes\r\n- **Warp Specialization**: Fixed iterator invalidation and use-after-free issues\r\n- **AMD OptimizeLDS**: Fixed nullptr crash\r\n- **Memory Leaks**: Fixed in TritonNvidiaGPU InterleaveTMem.cpp (#7924)\r\n- **Nullptr Access**: Fixed in AMD pingpong ChainedDot (#7694)\r\n\r\n### Regression Fixes\r\n- **Block Size Logic Revert** (#7971): Reverted fp8 matmul issues\r\n- **Byte Permutes Revert** (#7899): Reverted due to functional regression, then relanded with fix (#7933)\r\n- **Diagonal Iteration Partial Revert** (#7245): Addressed internal regressions\r\n\r\n---\r\n\r\n## Contributors\r\n\r\nThis release includes contributions from engineers at:\r\n- OpenAI\r\n- Meta\r\n- AMD\r\n- NVIDIA\r\n- Intel\r\n- Google\r\n- And many individual contributors\r\n\r\nSpecial thanks to all contributors who submitted bug reports, feature requests, and code improvements!","reactions":{"url":"https://api.github.com/repos/triton-lang/triton/releases/255856323/reactions","total_count":16,"+1":0,"-1":0,"laugh":0,"hooray":11,"confused":0,"heart":2,"rocket":3,"eyes":0}},{"url":"https://api.github.com/repos/triton-lang/triton/releases/235969377","assets_url":"https://api.github.com/repos/triton-lang/triton/releases/235969377/assets","upload_url":"https://uploads.github.com/repos/triton-lang/triton/releases/235969377/assets{?name,label}","html_url":"https://github.com/triton-lang/triton/releases/tag/v3.4.0","id":235969377,"author":{"login":"atalman","id":7563158,"node_id":"MDQ6VXNlcjc1NjMxNTg=","avatar_url":"https://avatars.githubusercontent.com/u/7563158?v=4","gravatar_id":"","url":"https://api.github.com/users/atalman","html_url":"https://github.com/atalman","followers_url":"https://api.github.com/users/atalman/followers","following_url":"https://api.github.com/users/atalman/following{/other_user}","gists_url":"https://api.github.com/users/atalman/gists{/gist_id}","starred_url":"https://api.github.com/users/atalman/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/atalman/subscriptions","organizations_url":"https://api.github.com/users/atalman/orgs","repos_url":"https://api.github.com/users/atalman/repos","events_url":"https://api.github.com/users/atalman/events{/privacy}","received_events_url":"https://api.github.com/users/atalman/received_events","type":"User","user_view_type":"public","site_admin":false},"node_id":"RE_kwDOAWaHXs4OEJth","tag_name":"v3.4.0","target_commitish":"main","name":"Triton 3.4.0 Release","draft":false,"immutable":false,"prerelease":false,"created_at":"2025-07-30T18:47:33Z","updated_at":"2025-07-30T20:48:05Z","published_at":"2025-07-30T20:47:15Z","assets":[{"url":"https://api.github.com/repos/triton-lang/triton/releases/assets/277876325","id":277876325,"node_id":"RA_kwDOAWaHXs4QkA5l","name":"triton-3.4.0.tar.gz","label":"","uploader":{"login":"github-actions[bot]","id":41898282,"node_id":"MDM6Qm90NDE4OTgyODI=","avatar_url":"https://avatars.githubusercontent.com/in/15368?v=4","gravatar_id":"","url":"https://api.github.com/users/github-actions%5Bbot%5D","html_url":"https://github.com/apps/github-actions","followers_url":"https://api.github.com/users/github-actions%5Bbot%5D/followers","following_url":"https://api.github.com/users/github-actions%5Bbot%5D/following{/other_user}","gists_url":"https://api.github.com/users/github-actions%5Bbot%5D/gists{/gist_id}","starred_url":"https://api.github.com/users/github-actions%5Bbot%5D/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/github-actions%5Bbot%5D/subscriptions","organizations_url":"https://api.github.com/users/github-actions%5Bbot%5D/orgs","repos_url":"https://api.github.com/users/github-actions%5Bbot%5D/repos","events_url":"https://api.github.com/users/github-actions%5Bbot%5D/events{/privacy}","received_events_url":"https://api.github.com/users/github-actions%5Bbot%5D/received_events","type":"Bot","user_view_type":"public","site_admin":false},"content_type":"application/gzip","state":"uploaded","size":5978956,"digest":"sha256:019221ceede5ed482b1856c7e47dc8c88eb7a6d1c187ffa70c0875d86c62ed17","download_count":1394,"created_at":"2025-07-30T20:47:38Z","updated_at":"2025-07-30T20:47:39Z","browser_download_url":"https://github.com/triton-lang/triton/releases/download/v3.4.0/triton-3.4.0.tar.gz"}],"tarball_url":"https://api.github.com/repos/triton-lang/triton/tarball/v3.4.0","zipball_url":"https://api.github.com/repos/triton-lang/triton/zipball/v3.4.0","body":"## Highlights\r\n\r\n### Gluon Framework Comprehensive Enhancement\r\nThe Gluon framework has received major enhancements across all areas including new APIs, tensor memory management, layout operations, and synchronization primitives. Key additions include `static_assert` functionality, TensorDescriptor kernel arguments, async TMA operations, tensor memory implementation, thread synchronization barriers, and comprehensive tensor operations like split/join/reshape and reductions. ([#7172](https://github.com/triton-lang/triton/pull/7172), [#7168](https://github.com/triton-lang/triton/pull/7168), [#7165](https://github.com/triton-lang/triton/pull/7165), [#7160](https://github.com/triton-lang/triton/pull/7160), [#7152](https://github.com/triton-lang/triton/pull/7152), [#7151](https://github.com/triton-lang/triton/pull/7151), [#7149](https://github.com/triton-lang/triton/pull/7149), [#7145](https://github.com/triton-lang/triton/pull/7145), [#7142](https://github.com/triton-lang/triton/pull/7142), [#7122](https://github.com/triton-lang/triton/pull/7122), [#7121](https://github.com/triton-lang/triton/pull/7121), [#7120](https://github.com/triton-lang/triton/pull/7120), [#7115](https://github.com/triton-lang/triton/pull/7115), [#7114](https://github.com/triton-lang/triton/pull/7114), [#7106](https://github.com/triton-lang/triton/pull/7106), [#7102](https://github.com/triton-lang/triton/pull/7102), [#7099](https://github.com/triton-lang/triton/pull/7099), [#7097](https://github.com/triton-lang/triton/pull/7097), [#7091](https://github.com/triton-lang/triton/pull/7091), [#7089](https://github.com/triton-lang/triton/pull/7089), [#7080](https://github.com/triton-lang/triton/pull/7080), [#7061](https://github.com/triton-lang/triton/pull/7061), [#7057](https://github.com/triton-lang/triton/pull/7057), [#7022](https://github.com/triton-lang/triton/pull/7022), [#7020](https://github.com/triton-lang/triton/pull/7020), [#7009](https://github.com/triton-lang/triton/pull/7009), [#7006](https://github.com/triton-lang/triton/pull/7006), [#7004](https://github.com/triton-lang/triton/pull/7004), [#7001](https://github.com/triton-lang/triton/pull/7001), [#6998](https://github.com/triton-lang/triton/pull/6998), [#6997](https://github.com/triton-lang/triton/pull/6997), [#6994](https://github.com/triton-lang/triton/pull/6994), [#6992](https://github.com/triton-lang/triton/pull/6992), [#6989](https://github.com/triton-lang/triton/pull/6989), [#6985](https://github.com/triton-lang/triton/pull/6985), [#6971](https://github.com/triton-lang/triton/pull/6971), [#6950](https://github.com/triton-lang/triton/pull/6950))\r\n\r\n### Hardware Support Expansion\r\n- **AMD GFX950 Architecture Support** - Comprehensive support for GFX950 including WMMA operations, performance optimizations, and architectural-specific features ([#7175](https://github.com/triton-lang/triton/pull/7175), [#7171](https://github.com/triton-lang/triton/pull/7171), [#7127](https://github.com/triton-lang/triton/pull/7127), [#6744](https://github.com/triton-lang/triton/pull/6744), [#6594](https://github.com/triton-lang/triton/pull/6594))\r\n- **Blackwell Enhanced TMEM Support** - Improved tensor memory operations with better register usage and performance optimizations ([#7160](https://github.com/triton-lang/triton/pull/7160), [#7079](https://github.com/triton-lang/triton/pull/7079), [#6817](https://github.com/triton-lang/triton/pull/6817))\r\n- **Hopper WGMMA Improvements** - Enhanced matrix multiplication with subtiling and prefetching optimizations ([#7136](https://github.com/triton-lang/triton/pull/7136), [#6130](https://github.com/triton-lang/triton/pull/6130))\r\n\r\n### Performance Optimizations\r\n- **Automatic Warp Specialization** - Introduced automatic warp specialization optimization for enhanced kernel performance on NVIDIA GPUs ([#6289](https://github.com/triton-lang/triton/pull/6289), [#6246](https://github.com/triton-lang/triton/pull/6246), [#6217](https://github.com/triton-lang/triton/pull/6217))\r\n- **MMAv5 Pipelining** - Re-enabled and improved MMAv5 pipelining with better performance and scheduling ([#6732](https://github.com/triton-lang/triton/pull/6732), [#6613](https://github.com/triton-lang/triton/pull/6613), [#6256](https://github.com/triton-lang/triton/pull/6256))\r\n- **TMA Operations Enhancement** - Improved tensor memory access with better layout support and reduced register pressure ([#6725](https://github.com/triton-lang/triton/pull/6725), [#6238](https://github.com/triton-lang/triton/pull/6238), [#6580](https://github.com/triton-lang/triton/pull/6580))\r\n\r\n## New Features\r\n\r\n### Language and Frontend\r\n- **Aggregate Type Support** - Added `@tl.aggregate` decorator for autogenerating Triton types from Python classes ([#6970](https://github.com/triton-lang/triton/pull/6970))\r\n- **JITFunction Constexpr Support** - Enhanced constexpr support for function lists and improved JIT functionality ([#6988](https://github.com/triton-lang/triton/pull/6988), [#6963](https://github.com/triton-lang/triton/pull/6963), [#7105](https://github.com/triton-lang/triton/pull/7105))\r\n- **Enhanced Boolean Operations** - Improved handling of boolean operators and scalars with chained operations ([#6769](https://github.com/triton-lang/triton/pull/6769))\r\n- **Bitonic Top-k and Sorting** - Added support for bitonic top-k operations and improved sort implementations ([#6461](https://github.com/triton-lang/triton/pull/6461), [#6486](https://github.com/triton-lang/triton/pull/6486))\r\n- **Masked Histograms** - Added support for masked histogram operations ([#6695](https://github.com/triton-lang/triton/pull/6695))\r\n- **Syntactic Sugar Additions** - Added `.item()` as syntactic sugar for `.reshape([])` ([#6873](https://github.com/triton-lang/triton/pull/6873))\r\n\r\n### Backend and Compilation\r\n- **Generic Swizzling Implementation** - Implemented generic swizzling algorithm for convert_layout lowering ([#6982](https://github.com/triton-lang/triton/pull/6982))\r\n- **Enhanced Register Allocation** - Improved dynamic register reallocation for warp specialization ([#6877](https://github.com/triton-lang/triton/pull/6877), [#6694](https://github.com/triton-lang/triton/pull/6694), [#6407](https://github.com/triton-lang/triton/pull/6407))\r\n- **TMA Reduce Operations** - Added TMA reduce operations for descriptor-based reducing stores ([#6580](https://github.com/triton-lang/triton/pull/6580))\r\n- **Improved Subtiling** - Enhanced subtiling code generation for tensor memory loading ([#6415](https://github.com/triton-lang/triton/pull/6415))\r\n- **BF16 Atomic Operations** - Added support for BF16 atomic add operations ([#6519](https://github.com/triton-lang/triton/pull/6519))\r\n- **Stmatrix Support** - Added comprehensive stmatrix support including transpose operations ([#6910](https://github.com/triton-lang/triton/pull/6910), [#6899](https://github.com/triton-lang/triton/pull/6899))\r\n\r\n### Hardware-Specific Features\r\n- **AMD AsyncCopy Optimizations** - Enhanced AsyncCopy support in StreamPipeliner with improved memory operations ([#6270](https://github.com/triton-lang/triton/pull/6270), [#6639](https://github.com/triton-lang/triton/pull/6639), [#6382](https://github.com/triton-lang/triton/pull/6382))\r\n- **AMD Buffer Operations** - Comprehensive improvements to buffer operations with better vectorization and alignment ([#6126](https://github.com/triton-lang/triton/pull/6126), [#6145](https://github.com/triton-lang/triton/pull/6145), [#6329](https://github.com/triton-lang/triton/pull/6329))\r\n- **AMD Ping-pong Scheduler** - Enhanced ping-pong scheduler for better memory operation handling ([#6254](https://github.com/triton-lang/triton/pull/6254), [#6301](https://github.com/triton-lang/triton/pull/6301), [#6198](https://github.com/triton-lang/triton/pull/6198))\r\n- **NVIDIA PDL Support** - Enabled Programmatic Dependent Launch for overlapping kernel execution ([#6394](https://github.com/triton-lang/triton/pull/6394))\r\n- **AMD HIP AOT Support** - Added HIP Ahead-of-Time compilation support ([#7007](https://github.com/triton-lang/triton/pull/7007))\r\n\r\n## Improvements\r\n\r\n### Performance\r\n- **Routing Kernel Optimizations** - Multiple performance improvements achieving up to 5% runtime reduction ([#6866](https://github.com/triton-lang/triton/pull/6866), [#6546](https://github.com/triton-lang/triton/pull/6546), [#7040](https://github.com/triton-lang/triton/pull/7040))\r\n- **Matrix Multiplication Enhancements** - Enhanced persistent TMA matmul with epilogue subtiling and metadata alignment ([#6724](https://github.com/triton-lang/triton/pull/6724), [#6882](https://github.com/triton-lang/triton/pull/6882), [#7123](https://github.com/triton-lang/triton/pull/7123))\r\n- **SwiGLU Optimizations** - Improved SwiGLU kernel performance and fused activation functions ([#6797](https://github.com/triton-lang/triton/pull/6797), [#6553](https://github.com/triton-lang/triton/pull/6553))\r\n- **Attention Kernel Fixes** - Fixed and optimized attention tutorials with better performance metrics ([#7037](https://github.com/triton-lang/triton/pull/7037), [#6839](https://github.com/triton-lang/triton/pull/6839))\r\n\r\n### Developer Experience\r\n- **Enhanced CI/CD** - Improved continuous integration with better caching and timeout handling ([#6815](https://github.com/triton-lang/triton/pull/6815), [#6816](https://github.com/triton-lang/triton/pull/6816), [#6582](https://github.com/triton-lang/triton/pull/6582))\r\n- **Testing Infrastructure** - Enhanced test coverage and organization ([#7109](https://github.com/triton-lang/triton/pull/7109), [#6867](https://github.com/triton-lang/triton/pull/6867))\r\n- **Documentation Updates** - Improved documentation for installation and new features ([#7103](https://github.com/triton-lang/triton/pull/7103), [#6778](https://github.com/triton-lang/triton/pull/6778), [#6235](https://github.com/triton-lang/triton/pull/6235))\r\n- **Build System Improvements** - Better CMake support and dependency management ([#6330](https://github.com/triton-lang/triton/pull/6330), [#6903](https://github.com/triton-lang/triton/pull/6903))\r\n\r\n### Code Quality\r\n- **Type System Enhancements** - Improved type checking with mypy integration ([#6596](https://github.com/triton-lang/triton/pull/6596), [#6704](https://github.com/triton-lang/triton/pull/6704))\r\n- **Layout System Improvements** - Better layout handling with LinearLayout-based implementations ([#6252](https://github.com/triton-lang/triton/pull/6252), [#6169](https://github.com/triton-lang/triton/pull/6169), [#6170](https://github.com/triton-lang/triton/pull/6170))\r\n- **Code Organization** - Extensive refactoring and cleanup for better maintainability ([#6500](https://github.com/triton-lang/triton/pull/6500), [#6285](https://github.com/triton-lang/triton/pull/6285))\r\n\r\n## Bug Fixes\r\n\r\n### Critical Fixes\r\n- **AST Parsing Regression** - Fixed parsing failures for float(\"inf\") and float(\"-inf\") expressions ([#6344](https://github.com/triton-lang/triton/pull/6344))\r\n- **Memory Allocation Issues** - Fixed tensor memory allocation boundary collisions and use-after-free errors ([#6318](https://github.com/triton-lang/triton/pull/6318), [#6433](https://github.com/triton-lang/triton/pull/6433))\r\n- **TMA Layout Consistency** - Fixed layout assignment from rank-reducing loads ([#6362](https://github.com/triton-lang/triton/pull/6362))\r\n- **Dot Operation Fixes** - Fixed bug where passing None as accumulator caused errors ([#7130](https://github.com/triton-lang/triton/pull/7130))\r\n- **Version Detection** - Fixed version detection when using source tarballs ([#7164](https://github.com/triton-lang/triton/pull/7164), [#6381](https://github.com/triton-lang/triton/pull/6381))\r\n\r\n### Hardware-Specific Fixes\r\n- **AMD Range Analysis** - Improved range analysis for persistent kernels and loop bounds ([#6390](https://github.com/triton-lang/triton/pull/6390), [#6133](https://github.com/triton-lang/triton/pull/6133))\r\n- **AMD Buffer Operations** - Fixed vector size computation and alignment issues ([#6114](https://github.com/triton-lang/triton/pull/6114), [#6126](https://github.com/triton-lang/triton/pull/6126))\r\n- **AMD Atomic Operations** - Fixed f16/bf16 buffer atomic operations ([#6090](https://github.com/triton-lang/triton/pull/6090), [#6139](https://github.com/triton-lang/triton/pull/6139))\r\n- **NVIDIA Register Pressure** - Fixed register allocation issues in warp specialization ([#6403](https://github.com/triton-lang/triton/pull/6403))\r\n- **NVIDIA TMEM Operations** - Fixed various tensor memory access issues ([#6888](https://github.com/triton-lang/triton/pull/6888))\r\n\r\n### Stability Improvements\r\n- **Test Reliability** - Resolved intermittent test failures across various components ([#6861](https://github.com/triton-lang/triton/pull/6861), [#6889](https://github.com/triton-lang/triton/pull/6889))\r\n- **Memory Usage** - Fixed memory leaks and reduced peak memory consumption ([#6796](https://github.com/triton-lang/triton/pull/6796))\r\n- **Error Handling** - Improved error messages and crash prevention ([#6865](https://github.com/triton-lang/triton/pull/6865))\r\n\r\n## Deprecations and Breaking Changes\r\n\r\n### Breaking Changes\r\n- **Cumsum Type Promotion** - Upcast boolean inputs in cumsum to uint32_t for correct results ([#6927](https://github.com/triton-lang/triton/pull/6927))\r\n- **Experimental API Cleanup** - Removed outdated experimental descriptor APIs ([#6488](https://github.com/triton-lang/triton/pull/6488))\r\n- **Python Support** - Dropped Python 3.8 support, minimum version now 3.9 ([#6649](https://github.com/triton-lang/triton/pull/6649))\r\n- **Tensor Descriptor APIs** - Removed experimental prefix from tensor descriptor operations ([#6194](https://github.com/triton-lang/triton/pull/6194))\r\n- **Register Spilling Performance Regression** - Bad interaction between new LLVM changes and PTXAS optimizations can cause increased register spilling in some kernels ([#7138](https://github.com/triton-lang/triton/pull/7138))\r\n\r\n### Deprecations\r\n- **FP8 Format Warnings** - Enhanced warnings for deprecated FP8 formats ([#6931](https://github.com/triton-lang/triton/pull/6931))\r\n- **Configuration Module** - Renamed config.py to knobs.py to avoid confusion ([#6641](https://github.com/triton-lang/triton/pull/6641))\r\n\r\n## Performance\r\n\r\n### Benchmark Results\r\n- **Matrix Multiplication** - Up to 15% speedup in dense 8k x 8k x 8k operations ([#6804](https://github.com/triton-lang/triton/pull/6804))\r\n- **Attention Kernels** - Achieved 700+ TFLOPS on DHEAD=64, 960-1080 TFLOPS on DHEAD=128 ([#6660](https://github.com/triton-lang/triton/pull/6660))\r\n- **Routing Operations** - 5% runtime reduction with optimized kernels ([#6866](https://github.com/triton-lang/triton/pull/6866))\r\n- **MoE Kernels** - Up to 30% performance boost with optimized TMA layouts ([#7123](https://github.com/triton-lang/triton/pull/7123))\r\n\r\n### Memory Optimizations\r\n- **Register Usage** - Reduced register pressure in various operations ([#6817](https://github.com/triton-lang/triton/pull/6817))\r\n- **Shared Memory** - Improved shared memory utilization with better swizzling ([#6982](https://github.com/triton-lang/triton/pull/6982))\r\n- **Cache Efficiency** - Enhanced cache utilization with L2 cache hints ([#6278](https://github.com/triton-lang/triton/pull/6278))\r\n\r\n## Documentation\r\n\r\n### New Guides\r\n- **Community Meetups** - Added documentation for running Triton Community Meetups ([#7103](https://github.com/triton-lang/triton/pull/7103))\r\n- **Installation Instructions** - Updated with better memory management guidance ([#6235](https://github.com/triton-lang/triton/pull/6235))\r\n- **Hardware Support** - Updated PyTorch installation for Blackwell support ([#6778](https://github.com/triton-lang/triton/pull/6778))\r\n\r\n### API Documentation\r\n- **Tensor Descriptors** - Comprehensive documentation for tensor descriptor APIs ([#6911](https://github.com/triton-lang/triton/pull/6911), [#7028](https://github.com/triton-lang/triton/pull/7028))\r\n- **Cache Modifiers** - Updated tl.load documentation with correct cache modifier usage ([#6214](https://github.com/triton-lang/triton/pull/6214))\r\n- **Scan Operations** - Enhanced docstrings with appropriate parameters ([#6946](https://github.com/triton-lang/triton/pull/6946))\r\n\r\n## Developers\r\n\r\n### Build System\r\n- **LLVM Integration** - Multiple LLVM version bumps with latest upstream changes ([#7138](https://github.com/triton-lang/triton/pull/7138), [#7129](https://github.com/triton-lang/triton/pull/7129), [#6754](https://github.com/triton-lang/triton/pull/6754), [#6361](https://github.com/triton-lang/triton/pull/6361))\r\n- **CMake Updates** - Improved build configuration and parallel building support ([#6830](https://github.com/triton-lang/triton/pull/6830), [#6953](https://github.com/triton-lang/triton/pull/6953))\r\n- **Dependency Management** - Better handling of external dependencies ([#7078](https://github.com/triton-lang/triton/pull/7078))\r\n\r\n### Testing Infrastructure\r\n- **Lit Tests** - Enhanced lit test coverage and organization ([#6855](https://github.com/triton-lang/triton/pull/6855), [#6661](https://github.com/triton-lang/triton/pull/6661))\r\n- **Benchmarking** - Enhanced benchmarking infrastructure with roofline analysis ([#6703](https://github.com/triton-lang/triton/pull/6703))\r\n- **CI/CD Improvements** - Better hardware support and workflow organization ([#6582](https://github.com/triton-lang/triton/pull/6582))\r\n\r\n### Code Organization\r\n- **Module Structure** - Better organization of modules and passes ([#6500](https://github.com/triton-lang/triton/pull/6500))\r\n- **Type System** - Enhanced type checking and inference ([#6285](https://github.com/triton-lang/triton/pull/6285), [#6231](https://github.com/triton-lang/triton/pull/6231))\r\n- **Error Handling** - Improved error messages and debugging support throughout the codebase\r\n","reactions":{"url":"https://api.github.com/repos/triton-lang/triton/releases/235969377/reactions","total_count":26,"+1":0,"-1":0,"laugh":0,"hooray":14,"confused":0,"heart":0,"rocket":8,"eyes":4}}]