Papers › Fast and Stable Triangular Inversion for Delta-Rule Linear Transformers

Fast and Stable Triangular Inversion for Delta-Rule Linear Transformers

20 May 2026arXiv:2605.21325added by Syntology

Aleksandros Sobczyk, Gioele Gottardo, Christos K. Matzoros, Mirko De Vita, Filip Skogh, Anastasios Zouzias, Jiawei Zhuang

Title, abstract, authors and date from arXiv's metadata (CC0); this paper is not in the Papers with Code archive (frozen 2025-07-28).

Linear attention has emerged as a cornerstone for efficient long-context architectures, as evidenced by its integration into state-of-the-art open-source models including Qwen3.5/3.6, Kimi Linear, and RWKV-7. Models that incorporate linear attention layers with the so-called Delta-Rule involve the inversion of triangular matrices as a core sub-routine. This operation often forms a performance bottleneck, and, due to its high-sensitivity to numerical errors, it can significantly deteriorate end-to-end model accuracy if it is not carefully implemented. This work provides a systematic analysis of both direct and iterative triangular inversion algorithms, targeting methods that are rich in matrix products, and, therefore, have the potential to efficiently utilize modern hardware. To that end, our analysis covers a broad spectrum of mathematical and practical aspects, with a heavy focus on numerical stability, computational complexity, and, ultimately, hardware efficiency and practical considerations. We provide a rigorous experimental evaluation to verify these properties in practical scenarios, and in low-precision floating-point representations, highlighting the strengths and limitations of each method. Performance benchmarks on NPUs reveal up to 4.3× speed-up against the state-of-the-art implementations of SGLang for triangular matrix inversion, leading to significant performance improvements on the entire layer level, while maintaining full end-to-end model accuracy.

PaperPDFCode Syntology ran

In Syntology Open this paper in Syntology's Atlas, the map of the papers in Syntology's graph and their citations.

For agents, Syntology's MCP tool lists every function and class Syntology harvested from this paper and whether it ran (how to connect): get_harvested_code_for_paper(arxiv_id="2605.21325")

Code

Syntology Ran 15 of 28 code samples harvested from 3 repositories linked to this paper; 13 have no recorded run. Of those that ran: 15 ran with no contract checked.

By repository: found in paper text by Syntology: 28 samples from 3 repositories, 15 ran. The run record, sample by sample. “Ran” means executed on a synthesized input, not that the code is correct or reproduces the paper.

huawei-csl/pto-kernels found in paper text by Syntology report
huawei-csl/gdn-tri-inverse found in paper text by Syntology report
huawei-csl/megagdn-pto found in paper text by Syntology report

Reachability, where shown, is from one Syntology probe window (2026-09-16 to 2026-09-18); repositories not probed show nothing. GitHub stars are not tracked.

Code Syntology ran Syntology

28 samples harvested; 15 ran; 0 honoured the contract we drafted; 13 have no recorded run. Read from Syntology's graph 2026-09-24; that is when this build read the record, not when the samples ran.

15ran
13unverified

Licence: 28 of the 28 samples are pointer only, meaning Syntology does not serve that copy's text. This page shows no code text for any sample; each one links to its file in the repository.

Harvested from 3 repositories linked to this paper, official or community; each sample names its own and says which. “Ran” means the sample executed on a synthesized input. It does not mean the output is correct, and nothing here reproduces the paper's results. “Honoured” and “violated” refer to a contract Syntology drafted from the code itself; “our draft was wrong” and “fixture could not drive it” are failures of Syntology's instrument, not of the code.

Each sample ends with its code_sha256, Syntology's identity for that exact code. An agent fetches the stored sample with Syntology's MCP tool get_code(code_sha256="…") (how to connect); click an identity to copy that call.

Repository labels, per sample. official repository: The archive marks this repository official for the paper. named in the paper: The archive records that the paper mentions this repository; it is not marked official. community (archive-listed): In the archive's code links for this paper, not marked official and not recorded as mentioned in the paper. found in paper text by Syntology: Syntology found this repository in the paper's own text; whether it is the authors' implementation is not asserted. community: Not in the archive's code links for this paper; a community repository Syntology harvested. Samples from a repository marked official are listed first. Licence labels name the repository's licence as recorded at harvest. “Pointer only” means Syntology does not serve that copy's text, for one of four reasons: no licence file was found; the licence was not identified; the licence is recorded as permissive but that copy's record is not marked cleared; or the licence is outside the permissive list Syntology serves text under (MIT, Apache-2.0, BSD and similar). Some licences outside that list permit redistribution, such as WTFPL, and GPL-3.0 under its conditions; they are simply not on the list. Hover a licence label for the reason. File links open the file on GitHub at the default branch, which may have changed since the harvest.

bind_triton huawei-csl/megagdn-pto/vllm_patch/chunk_kda_pto.py found in paper text by Syntology ran no licence file found · pointer only · 3be586d513227377 · report
chunk_gated_delta_rule_native huawei-csl/gdn-tri-inverse/src/gdn_tri_inverse/triton_chunk_gdn/fla/chunk.py found in paper text by Syntology ran BSD-3-Clause-Clear · pointer only · 8c58a327ad934418 · report
chunk_kda_pto huawei-csl/megagdn-pto/vllm_patch/chunk_kda_pto.py found in paper text by Syntology ran no licence file found · pointer only · 88c21f1f458ccf5c · report
compile_kernel huawei-csl/pto-kernels/.skills/testing-pto-kernels/reference/pto_demo_utils.py found in paper text by Syntology ran BSD-3-Clause-Clear · pointer only · 49cd9155190e49b2 · report
get_tp_context huawei-csl/megagdn-pto/docker/overrides/model_runner_v1.py found in paper text by Syntology ran no licence file found · pointer only · 179b4ecb6419827a · report
inv_tril_inplace huawei-csl/gdn-tri-inverse/src/gdn_tri_inverse/triton_chunk_gdn/fla/chunk.py found in paper text by Syntology ran fingerprinted BSD-3-Clause-Clear · pointer only · 65c98279ed930932 · report
npu_from_cpu huawei-csl/pto-kernels/.skills/testing-pto-kernels/reference/dynamic_multi_core/a5/run_kernel_ctypes.py found in paper text by Syntology ran BSD-3-Clause-Clear · pointer only · 29854c7c42411c2b · report
npu_tensor huawei-csl/pto-kernels/.skills/testing-pto-kernels/reference/dynamic_multi_core/a5/run_kernel_pybind.py found in paper text by Syntology ran BSD-3-Clause-Clear · pointer only · 43739c7f7a9f008f · report
summarize huawei-csl/pto-kernels/.skills/testing-pto-kernels/reference/dynamic_multi_core/a2a3/benchmark.py found in paper text by Syntology ran fingerprinted BSD-3-Clause-Clear · pointer only · 7140e33bf7bef21f · report
tensor_ptr huawei-csl/pto-kernels/.skills/testing-pto-kernels/reference/pto_demo_utils.py found in paper text by Syntology ran BSD-3-Clause-Clear · pointer only · 6c19edb11fb6754c · report
total_chunks huawei-csl/megagdn-pto/megagdn_pto/kernel_libs.py found in paper text by Syntology ran no licence file found · pointer only · e4e5891b8b4488b1 · report
transpose_beta huawei-csl/megagdn-pto/megagdn_pto/kernel_libs.py found in paper text by Syntology ran no licence file found · pointer only · 33a4536ecf79ed25 · report
transpose_gates huawei-csl/megagdn-pto/megagdn_pto/kernel_libs.py found in paper text by Syntology ran no licence file found · pointer only · e0779707b00e56f4 · report
tri_inv_qwen3_next_default huawei-csl/gdn-tri-inverse/src/gdn_tri_inverse/linalg.py found in paper text by Syntology ran fingerprinted BSD-3-Clause-Clear · pointer only · 27aa399aa0968a73 · report
workspace huawei-csl/pto-kernels/.skills/testing-pto-kernels/reference/dynamic_multi_core/a2a3/run_mix_ctypes.py found in paper text by Syntology ran BSD-3-Clause-Clear · pointer only · bf1d857bd30bc516 · report
add_reference_to_path huawei-csl/pto-kernels/.skills/testing-pto-kernels/reference/pto_demo_utils.py found in paper text by Syntology unverified BSD-3-Clause-Clear · pointer only · b26eb7a5b59d8e1f · report
build_kernel huawei-csl/pto-kernels/.skills/testing-pto-kernels/reference/dynamic_multi_core/a2a3/run_kernel_pybind.py found in paper text by Syntology unverified BSD-3-Clause-Clear · pointer only · bcc14bcd0eec3750 · report
compile_chunk_kernel huawei-csl/megagdn-pto/megagdn_pto/compile.py found in paper text by Syntology unverified no licence file found · pointer only · 1a65bd82810630c4 · report
compile_tri_inverse huawei-csl/megagdn-pto/megagdn_pto/compile.py found in paper text by Syntology unverified no licence file found · pointer only · 8991055d8785b3a2 · report
event_bench huawei-csl/pto-kernels/.skills/testing-pto-kernels/reference/dynamic_multi_core/a2a3/benchmark.py found in paper text by Syntology unverified BSD-3-Clause-Clear · pointer only · b0bfd6df1362ea99 · report
fast_inv_tril huawei-csl/gdn-tri-inverse/src/gdn_tri_inverse/triton_chunk_gdn/fla/chunk.py found in paper text by Syntology unverified BSD-3-Clause-Clear · pointer only · 3872a6d7e9a842a6 · report
recurrent_gated_delta_rule_ref huawei-csl/gdn-tri-inverse/src/gdn_tri_inverse/core.py found in paper text by Syntology unverified BSD-3-Clause-Clear · pointer only · 94a3d31751316112 · report
ref_chunk_cumsum huawei-csl/gdn-tri-inverse/src/gdn_tri_inverse/tilelang_chunk_gdn/opt_gdn_chunk_cumsum.py found in paper text by Syntology unverified BSD-3-Clause-Clear · pointer only · 1902ecb472db2fcc · report
ref_chunk_h huawei-csl/gdn-tri-inverse/src/gdn_tri_inverse/tilelang_chunk_gdn/opt_gdn_chunk_h.py found in paper text by Syntology unverified BSD-3-Clause-Clear · pointer only · 2996cf978e93d4a1 · report
ref_chunk_o huawei-csl/gdn-tri-inverse/src/gdn_tri_inverse/tilelang_chunk_gdn/opt_gdn_chunk_o.py found in paper text by Syntology unverified BSD-3-Clause-Clear · pointer only · 7136a894cedf2d88 · report
ref_seq_gdn huawei-csl/gdn-tri-inverse/src/gdn_tri_inverse/tilelang_chunk_gdn/opt_gdn_full.py found in paper text by Syntology unverified BSD-3-Clause-Clear · pointer only · f3d36efe796d3c18 · report
ref_wy_fast huawei-csl/gdn-tri-inverse/src/gdn_tri_inverse/tilelang_chunk_gdn/opt_gdn_wy_fast.py found in paper text by Syntology unverified BSD-3-Clause-Clear · pointer only · 18bceabcfb10e5ac · report
staged_separate_launch huawei-csl/megagdn-pto/benchmarks/kernel/bench_pto_pipeline_latency.py found in paper text by Syntology unverified no licence file found · pointer only · 04b68677ccfb33a0 · report

Results from the paper

The Papers with Code archive ends with its 2025-07-28 snapshot. This paper's arXiv identifier, 2605.21325, was issued in May 2026, after that date, so the archive has no leaderboard rows for it.

Report a problem or propose a change · a person checks every report against the paper or source before anything changes; decisions are listed on /corrections