llm-catalog-archive

Change

801e001

801e0016cb481c218c95a5f0db83d6d6b615d51f · commit on GitHub

pytorch-blog-feed: changed (366763 bytes, HTTP 200)

raw/pytorch-blog-feed/response.xml modified

Lines added
+700
Lines removed
-355
Stored bytes at this commit
366,763
Timestamp
origin
Raw artifact at this commit
raw/pytorch-blog-feed/response.xml
Recorded headers
observed_at2026-09-17T04:53:29.961Z
origin_date2026-09-17T03:24:12.000Z
status200
final URLhttps://pytorch.org/blog/feed/
etag"2f292fd0a8a26fe1bb5718e606c00d81"
last-modifiedWed, 16 Sep 2026 22:48:18 GMT
dateThu, 17 Sep 2026 04:53:29 GMT
age5357
cache-controlpublic, max-age=60, s-maxage=43200, stale-while-revalidate=86400, stale-if-error=604800
cf-cache-statusnull
content-encodingnull
content-length366763
@@@ -12,7 +12,7 @@
<atom:link href="https://pytorch.org/blog/feed/" rel="self" type="application/rss+xml" />
<link>https://pytorch.org</link>
<description></description>
- <lastBuildDate>Fri, 11 Sep 2026 13:04:23 +0000</lastBuildDate>
+ <lastBuildDate>Wed, 16 Sep 2026 22:48:18 +0000</lastBuildDate>
<language>en-US</language>
<sy:updatePeriod>
hourly </sy:updatePeriod>
@@@ -28,6 +28,692 @@
<height>32</height>
</image>
<item>
+ <title>Open Research, Tooling &#038; Optimization at PyTorch Conference North America 2026</title>
+ <link>https://pytorch.org/blog/open-research-tooling-optimization-at-pytorch-conference-north-america-2026/</link>
+
+ <dc:creator><![CDATA[PyTorch Foundation]]></dc:creator>
+ <pubDate>Wed, 16 Sep 2026 19:00:47 +0000</pubDate>
+ <category><![CDATA[Blog]]></category>
+ <guid isPermaLink="false">https://pytorch.org/?p=167741</guid>
+
+ <description><![CDATA[TL;DR Taking place October 20 to 21 in San Jose, California, PyTorch Conference North America 2026 highlights open research, tooling, and performance optimization across compiler architecture, cross-hardware kernel domain-specific languages,...]]></description>
+ <content:encoded><![CDATA[<h2>TL;DR</h2>
+<p><span style="font-weight: 400;">Taking place October 20 to 21 in San Jose, California, PyTorch Conference North America 2026 highlights open research, tooling, and performance optimization across compiler architecture, cross-hardware kernel domain-specific languages, exascale distributed training
+<h2><span style="font-weight: 400;">Introduction</span></h2>
+<p><span style="font-weight: 400;">Every year, the beating heart of PyTorch Conference is the work that happens <em>underneath</em> the flashy headlines- the compilers, kernels, autotuners, distributed runtimes, profilers, and open-source libraries that make the rest of the ecosystem possible. At Py
+<p><span style="font-weight: 400;">This is where you&#8217;ll find the engineers rewriting torch.compile&#8217;s internals for speed, the teams building DSLs (Helion, CuteDSL, FlyDSL) that let a single kernel target NVIDIA, AMD, Intel, and TPU silicon, the researchers pushing quantization down to FP
+<p><span style="font-weight: 400;">In this blog we have picked out the sessions that fall into this theme &#8211; open source tooling, systems research, and hard-won performance optimization.</span></p>
+<p><a href="https://hubs.ly/Q04tDx8f0">View the full conference schedule</a></p>
+<p><a href="https://hubs.ly/Q04tDx8f0">Register for PyTorch Conference North America 2026</a></p>
+<h2><span style="font-weight: 400;">Compilers, Graph Capture &amp; Dynamic Shapes</span></h2>
+<p><span style="font-weight: 400;">Everything to do with torch.compile, Dynamo, and the machinery that turns eager PyTorch code into a graph format that can be optimized at runtime, or used to generate artifacts for later execution</span></p>
+<p><strong>Beyond Size and Stride: Unleash Performance with Device-Aware Tensor Layouts</strong><br />
+Olivier Tardieu; Matthew Arnold (IBM)<br />
+Tue 20 Oct, 16:20–16:30, LL20AB</p>
+<p>A lightweight tensor layout extension that lets torch.compile (Inductor) automatically adapt tiling and NUMA-aware placement to the target device, without changing existing PyTorch model code.</p>
+<p><b>Nested Graph Breaks: Reducing the Cost of Graph Breaks in torch.compile</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">William Wen (Meta) </span></p>
+<p><span style="font-weight: 400;">Tue 20 Oct, 16:35–16:45, LL20AB </span></p>
+<p><span style="font-weight: 400;">New Dynamo support reduces the cost of a graph break nested O(N) layers deep from O(N) duplicate breaks and O(N²) frame traces down to O(1) duplicate breaks and O(N) frame traces &#8211; yielding larger captured graphs, fewer breaks, and faster, more debuggable com
+<p><b>Static Tensor Shape Checking for PyTorch with Pyrefly</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">Steven Troxler; Avik Chaudhuri (Meta) </span></p>
+<p><span style="font-weight: 400;">Tue 20 Oct, 17:30–17:55, LL20AB </span></p>
+<p><span style="font-weight: 400;">Shape mismatches cause roughly 45% of deep learning program failures and drive slow torch.compile recompilations. Pyrefly brings static, near-instant tensor shape checking to the type checker itself, evaluated across 28 real LLM, vision, recommender, and RL models.
+<p><b>Parametrized Dynamic Shape CUDA Graphs</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">Elias Ellison (Meta); Daniel Galvez (NVIDIA) </span></p>
+<p><span style="font-weight: 400;">Wed 21 Oct, 11:45–12:10, LL21ABC </span></p>
+<p><span style="font-weight: 400;">Building on parametrized CUDA Graphs and torch.compile&#8217;s symbolic tracing and guard infrastructure, this work captures and re-parametrizes a single CUDA Graph across dynamic shapes, eliminating the need for whole-model rewrites, padding, or re-recording acros
+<p><b>From Backed to Unbacked: Sound, Predictable, and Controllable Dynamic Shapes in PyTorch</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">Laith Sakka (Meta) </span></p>
+<p><span style="font-weight: 400;">Wed 21 Oct, 16:20–16:45, LL21ABC </span></p>
+<p><span style="font-weight: 400;">This session explains why explicit graph-capture workflows (vLLM, export, pre-compilation) need unbacked dynamic shapes, which disallow implicit guards and force the compiler to prove general validity. </span></p>
+<p><b>Speeding Up torch.compile: A New FakeTensor</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">Angel Li (Meta) </span></p>
+<p><span style="font-weight: 400;">Wed 21 Oct, 16:55–17:05, LL21ABC </span></p>
+<p><span style="font-weight: 400;">FakeTensor propagation eats roughly 20% of Dynamo&#8217;s tracing time. This talk introduces a new C++ FakeTensor implementation that delivers a 30x speedup on operations like </span><span style="font-weight: 400;">aten.mm</span><span style="font-weight: 400;">, me
+<p><b>Lightweight FX Tracing in PyTorch</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">Richard Zou; Yidi Wu (Meta) </span></p>
+<p><span style="font-weight: 400;">Wed 21 Oct, 17:10–17:20, LL21ABC </span></p>
+<p><span style="font-weight: 400;">A new FX tracer based on </span><span style="font-weight: 400;">make_fx</span><span style="font-weight: 400;">, targeting functionally pure PyTorch code instead of trying to capture full Python semantics like Dynamo. This means trading flexibility for a much simple
+<p><b>Unlocking the Full Potential of TorchDynamo: Accelerating, Comparing, and Debugging ML Systems</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">Yi Pan (UC Berkeley); Megan Frisella (University of Washington); Stephanie Wang (Paul Allen School, University of Washington) </span></p>
+<p><span style="font-weight: 400;">Wed 21 Oct, 17:30–17:55, LL21ABC </span></p>
+<p><span style="font-weight: 400;">Research using TorchDynamo&#8217;s graph-interception capability well beyond torch.compile itself. DynaFlow and Piper for transparent parallelism acceleration, Magneton for automatically comparing equivalent operations across competing systems, and numerical debugg
+<h2><span style="font-weight: 400;">Kernel Engineering &amp; Domain-Specific Languages</span></h2>
+<p><span style="font-weight: 400;">The DSLs, autotuners, and hand- and agent-written kernels that squeeze more performance out of every GPU cycle, spanning Triton, Helion, CUTLASS, FlyDSL, and the agentic pipelines now writing and optimizing kernels themselves.</span></p>
+<p><b>Extending TorchInductor with FlyDSL: A New MLIR-Native Backend for High-Performance GEMMs</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">Liz Li (AMD) </span></p>
+<p><span style="font-weight: 400;">Tue 20 Oct, 11:10–11:35,  210BF </span></p>
+<p><span style="font-weight: 400;">AMD&#8217;s FlyDSL, a Python-native, MLIR-based GPU kernel DSL, gets integrated into TorchInductor&#8217;s GEMM compilation pipeline. The talk covers how FlyDSL plugs into autotuning, coexists with existing backends, and delivers measured speedups over Triton on AM
+<p><b>Helion: CuteDSL and TPU Backends for Heterogeneous Hardware, and Why It Suits Agents</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">Oguz Ulgen, Dunfan Lu, Jason Ansel (Meta) </span></p>
+<p><span style="font-weight: 400;">Tue 20 Oct, 11:45–12:10, 210BF </span></p>
+<p><span style="font-weight: 400;">Helion, PyTorch&#8217;s high-level kernel-authoring DSL, gains two new compiler backends &#8211; CuteDSL for recent NVIDIA GPUs and Pallas for TPUs &#8211; letting one kernel source target very different hardware. The second half explores why Helion&#8217;s abstrac
+<p><b>Practical GPU Programming with Triton for PyTorch Developers</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">Suman Debnath, JanakiRam Goteti (Crusoe AI) </span></p>
+<p><span style="font-weight: 400;">Tue 20 Oct, 11:45–12:10,  LL20CD</span></p>
+<p><span style="font-weight: 400;">A friendly, from-scratch introduction to writing GPU kernels in Triton &#8211; no CUDA or C++ required.  Building from a simple vector-add up to a small matrix multiplication, with an emphasis on reading and understanding what PyTorch already generates for you.</sp
+<p><b>High-Velocity GPU Kernel Authoring with CUTLASS Python</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">Michael Goldfarb, Guray Ozen (NVIDIA) </span></p>
+<p><span style="font-weight: 400;">Tue 20 Oct, 12:20–12:45, 210BF </span></p>
+<p><span style="font-weight: 400;">New Python-first capabilities in CUTLASS&#8217;s CuTe DSL &#8211; high-level building blocks, low-level hardware-instruction primitives, and a zero-cost Resource and Task Scheduler &#8211; that reduce boilerplate for kernel authors while keeping the DSL&#8217;s zer
+<p><b>Smarter Autotuning for Kernels: From Bayesian Optimization to LLM-Guided Search in Helion DSL</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">Jongsok Choi; Ethan Che (Meta) </span></p>
+<p><span style="font-weight: 400;">Tue 20 Oct, 14:15–14:40,  210BF </span></p>
+<p><span style="font-weight: 400;">Two leaps in Helion&#8217;s autotuning: Likelihood-Free Bayesian Optimization (LFBO), which cuts tuning time by 36.5% via an on-the-fly Random Forest filter, and LLM-guided autotuning, which proposes strong kernel configurations in seconds &#8211; up to 10x faster
+<p><b>KernelAgent: Hardware-Guided GPU Kernel Optimization via Multi-Agent Orchestration</b><br />
+<span style="font-weight: 400;">Kaiming Cheng; Laura Wang (Meta) </span></p>
+<p><span style="font-weight: 400;">Tue 20 Oct, 14:15–14:40, LL21DEF </span></p>
+<p><span style="font-weight: 400;">Building on KernelAgent&#8217;s 100% correctness across all KernelBench L1/L2/L3 tasks, this talk adds a hardware-guided optimization layer that feeds GPU performance signals into a closed-loop multi-agent workflow, achieving 1.56x speedup over default torch.compil
+<p><b>From Weeks to Overnight: Autonomous Day-0 Kernel Bring-Up with Agent Pipelines</b><br />
+<span style="font-weight: 400;">Xiaogang Gu; Qun Yang (Intel) </span></p>
+<p><span style="font-weight: 400;">Tue 20 Oct, 17:30–17:40, 210BF </span></p>
+<p><span style="font-weight: 400;">A pipeline-driven autonomous system that encodes GPU kernel optimization (Profile, Analyze, CodeGen, Verify, Fix, Benchmark, Evaluate) as a deterministic workflow with specialized agents per stage and long-term memory, cutting kernel bring-up cycles from weeks to u
+<p><b>Beyond the Brrr: Building a Unified Ecosystem for Optimized Kernels</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">Sayak Paul (Hugging Face) </span></p>
+<p><span style="font-weight: 400;">Wed 21 Oct, 11:10–11:35, 210BF </span></p>
+<p><span style="font-weight: 400;">Generic operators often use only 20–30% of available GPU performance. This talk introduces Hugging Face&#8217;s Kernels library, which makes discovering and swapping in optimized custom kernels as simple as loading a model checkpoint &#8211; already delivering 2–5x
+<p><b>Making vLLM Faster on Intel GPUs with Triton Kernels</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">Whitney Tsang, Artur Fierka (Intel) </span></p>
+<p><span style="font-weight: 400;">Wed 21 Oct, 12:20–12:30, 210BF </span></p>
+<p><span style="font-weight: 400;">Treats Intel GPU architectures like Arc as first-class optimization targets rather than portability afterthoughts, covering Triton kernel strategies for unified attention, fused and batched MoE, and fusion opportunities like QK-norm+RoPE that outperform hand-writte
+<p><b>FlexGEMM: Flexible PyTorch Epilogues</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">Driss Guessous (Meta) </span></p>
+<p><span style="font-weight: 400;">Wed 21 Oct, 12:35–12:45,  210BF </span></p>
+<p><span style="font-weight: 400;">A proposed PyTorch frontend for writing GEMM epilogues &#8211; bias, activation, residual adds, fp8 scaling, SwiGLU-style gating &#8211; as ordinary PyTorch functions, giving the compiler a clear, fail-closed contract for fusing them directly into the GEMM store pa
+<p><b>Native DSL Operators in PyTorch Core</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">Simon Layton (Meta) </span></p>
+<p><span style="font-weight: 400;">Wed 21 Oct, 14:50–15:15, LL21ABC </span></p>
+<p><span style="font-weight: 400;">DSL-authored kernels (as used in FlashAttention) have stayed outside PyTorch&#8217;s core dispatch system. This talk presents ongoing work to make DSL operators first-class citizens, tied into dispatch and testing, enabling new operators and targeted performance fi
+<p><b>Portable Paged Attention: From Triton to Helion</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">Burkhard Ringlein (IBM Research) </span></p>
+<p><span style="font-weight: 400;">Wed 21 Oct, 15:25–15:50, 210BF </span></p>
+<p><span style="font-weight: 400;">Building on the state-of-the-art portable Triton attention backend in vLLM, this team wrote an experimental Helion attention backend for Paged Attention &#8211; vLLM&#8217;s core kernel &#8211; reporting up to 50% lower latency and 10% higher end-to-end throughput
+<h2><span style="font-weight: 400;">Distributed Training, Parallelism &amp; Systems</span></h2>
+<p><span style="font-weight: 400;">Systems and abstractions for training at scale — sharding, single-GPU-to-cluster developer experience, and the frameworks that decide how a model gets split across thousands of devices.</span></p>
+<p><b>Monarch &#8211; Single-Machine DevX at Distributed Scale</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">Marius Eriksen (Meta) </span></p>
+<p><span style="font-weight: 400;">Tue 20 Oct, 11:10–11:35, LL20CD </span></p>
+<p><span style="font-weight: 400;">Monarch lets developers program an entire cluster as if it were a single machine, using actor-based IPC and RDMA so remote logs, filesystems, and debugging behave like they&#8217;re local. The talk demos how it speeds up building large-scale post-training framework
+<p><b>FlexShard: Pluggable Sharding Placements for Matrix Optimizers and Block-wise Quantization</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">Wei Feng; Anshul Sinha; Ailing Zhang (Meta) </span></p>
+<p><span style="font-weight: 400;">Wed 21 Oct, 11:10–11:35, LL20CD </span></p>
+<p><span style="font-weight: 400;">A prototype PyTorch abstraction that decouples parameter sharding from FSDP&#8217;s runtime machinery via a small Placement contract, enabling optimizer-aware sharding for matrix optimizers like Muon and quantization-aligned sharding for block-wise FP8 all-gather.<
+<p><b>DeepSpeed Is Not Just ZeRO: Tensor, Sequence, and Expert Parallelism with ZeRO-Level Usability</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">Masahiro Tanaka (Anyscale) </span></p>
+<p><span style="font-weight: 400;">Wed 21 Oct, 11:10–11:35, LL21DEF </span></p>
+<p><span style="font-weight: 400;">DeepSpeed&#8217;s AutoTP, AutoSP, and AutoEP bring tensor, sequence, and expert parallelism to existing model implementations, including many Hugging Face models, without rewriting code, and can be combined with ZeRO for large dense models, long-context training, a
+<p><b>Precompile for Training</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">Bob Ren; Aaron Orenstein (Meta) </span></p>
+<p><span style="font-weight: 400;">Wed 21 Oct, 12:00–12:10, LL20CD </span></p>
+<p><span style="font-weight: 400;">A new abstraction for compiling distributed SPMD training programs once and reusing the artifact across every rank, eliminating redundant per-rank compilation, preventing divergent compiler decisions that cause NCCL hangs, and guaranteeing identical kernel selectio
+<p><b>From Full DTensor to spmd_types: Making Parallelism Explicit in TorchTitan</b><br />
+<span style="font-weight: 400;">Chien-Chin Huang, Pian Pawakapan (Meta) </span></p>
+<p><span style="font-weight: 400;">Wed 21 Oct, 14:15–14:40, LL21DEF </span></p>
+<p><span style="font-weight: 400;">The evolution from an all-DTensor approach to spmd_types, a lightweight, zero-runtime-overhead type system where sharding is part of a tensor&#8217;s type, collectives are explicit, and backward gradient sharding follows predictably from forward, covering both glob
+<p><b>AutoParallel: Automatic Discovery of Distributed Training Strategies in PyTorch</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">Sanket Jayant Purandare, Francisco Massa (Meta) </span></p>
+<p><span style="font-weight: 400;">Wed 21 Oct, 15:25–15:50, LL21DEF </span></p>
+<p><span style="font-weight: 400;">Given only a device mesh, memory budget, and sample batch, AutoParallel searches for an optimal parallelization strategy directly on the model graph &#8211; no manual FSDP/TP/SP/EP placement &#8211; matching a strong hand-tuned baseline on Llama-3 8B at 12% less me
+<h2><span style="font-weight: 400;">Communication &amp; Collective Libraries</span></h2>
+<p><span style="font-weight: 400;">The collective-communication layer underneath distributed training and inference: symmetric memory, NCCL extensions, and fault-tolerant libraries built to survive failures at exascale.</span></p>
+<p><b>EP-Overlap: Hiding Expert-Parallel Communication in MoE Training</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">Sanket Jayant Purandare (Meta) </span></p>
+<p><span style="font-weight: 400;">Tue 20 Oct, 15:25–15:35,  LL21DEF </span></p>
+<p><span style="font-weight: 400;">A general technique for hiding expert-parallel all-to-all communication behind any available compute,  not just DualPipe pipeline schedules, implemented as graph passes in TorchTitan&#8217;s GraphTrainer, hiding 76–86% of EP communication on DeepSeek-V3 for a ~1.3x
+<p><b>Future of Distributed Communication in PyTorch: New APIs for Fault Tolerance, RDMA and Extensibility</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">Tristan Rice, Kapil Sharma (Meta) </span></p>
+<p><span style="font-weight: 400;">Tue 20 Oct, 16:20–16:45, LL21ABC</span></p>
+<p><span style="font-weight: 400;">New torch.distributed APIs incubated in TorchComms and now upstreaming directly. Live process-group reconfiguration after rank failures, one-sided RDMA put/get windows, composable collective hooks, and pip-installable backend extensibility &#8211; all with zero mig
+<p><b>rocSHMEM Symmetric Memory in PyTorch for AMD GPUs</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">Prachi Gupta (AMD) </span></p>
+<p><span style="font-weight: 400;">Tue 20 Oct, 16:35–16:45, LL21DEF </span></p>
+<p><span style="font-weight: 400;">Extends PyTorch&#8217;s symmetric memory, previously NVIDIA/NVSHMEM-only, to AMD GPUs via rocSHMEM, exposing one-sided put/get and signal/wait primitives as Triton-callable ops so a single Triton kernel can run MoE all-to-all dispatch and combine on either vendor&#
+<p><b>XCCL: Scaling PyTorch Collectives to Exascale on Intel GPUs with TorchComms</b><br />
+<span style="font-weight: 400;">Panagiotis Kourdis, Tanima Dey (Intel) </span></p>
+<p><span style="font-weight: 400;">Tue 20 Oct, 17:10–17:20, LL21ABC </span></p>
+<p><span style="font-weight: 400;">XCCL extends TorchComms with native Intel GPU support built on oneCCL, validated on Argonne&#8217;s Aurora exascale supercomputer with over 90% scaling efficiency across thousands of nodes, plus a stream-ordered async execution model that overlaps computation and c
+<p><b>Beyond Collectives: Building Fused Compute-Communication Kernels with PyTorch Symmetric Memory</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">Ke Wen (NVIDIA), Natalia Gimelshein, Kapil Sharma (Meta) </span></p>
+<p><span style="font-weight: 400;">Wed 21 Oct, 11:45–12:10, 210BF </span></p>
+<p><span style="font-weight: 400;">PyTorch Symmetric Memory gives every rank a direct window into peers&#8217; GPU memory, letting communication happen inline inside a kernel rather than as a black-box collective. This covers compute–communication fusion, CuTe-DSL bindings for NCCL device primitives
+<p><b>Introducing NCCL Extensions: Communication Patterns for Modern AI</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">Sreeram Potluri, Artem Polyakov (NVIDIA)</span></p>
+<p><span style="font-weight: 400;">Wed 21 Oct, 16:55–17:20, 210AE </span></p>
+<p><span style="font-weight: 400;">Debuts NCCL-EP, bringing tuned MoE dispatch/combine natively into NCCL with GPU-initiated RDMA and NVLink, and NCCL-M2N, enabling zero-copy resharding between disjoint device meshes for RL weight rollout. Both are built on NCCL&#8217;s new Device APIs for custom, f
+<p><b>MCCL: Fault-Tolerant Collective Communication for Large-Scale GPU Training</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">Ben Carver (Meta) </span></p>
+<p><span style="font-weight: 400;">Wed 21 Oct, 17:30–17:55, 210AE </span></p>
+<p><span style="font-weight: 400;">The Meta Collective Communication Library (MCCL) treats hardware failure as a first-class design principle at tens-of-thousands-of-GPU scale, enabling dynamic communicator reconfiguration, shrinking and growing as ranks fail or recover, without full job restarts, i
+<p><span style="font-weight: 400;">Quantization &amp; Low-Precision Training</span></p>
+<p><span style="font-weight: 400;">Pushing training and inference into 4- and 8-bit formats — NVFP4, HiFloat, and the recipes that close the accuracy gap with BF16 while unlocking large throughput gains.</span></p>
+<p><b>HiFloat: Democratizing Ultra-Low Precision Training and Inference in PyTorch Ecosystems</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">Yun Zhao, Haonan Zhang (Huawei) </span></p>
+<p><span style="font-weight: 400;">Wed 21 Oct, 14:50–15:00, LL20AB </span></p>
+<p><span style="font-weight: 400;">HiFloat8 and HiFloat4 are tapered-precision formats designed to match neural network weight and gradient distributions without complex delayed scaling, achieving FP16 loss parity with a 1.5x–1.7x GEMM speedup, integrated via PyTorch custom ops, Triton kernels, and
+<p><b>Efficient Pretraining of LLMs in NVFP4</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">Anjulie Agrusa, Ryan Spring, Bruce Zitelli (NVIDIA) </span></p>
+<p><span style="font-weight: 400;">Wed 21 Oct, 15:25–15:50, LL20AB </span></p>
+<p><span style="font-weight: 400;">NVFP4 GEMMs offer 2x–3x higher peak throughput than FP8 on GB200/GB300, but quality gaps persist versus BF16 at smaller model sizes. This talk presents recipe improvements that close that gap, upstreamed into TorchAO and TorchTitan for native PyTorch NVFP4 pretrain
+<h2><span style="font-weight: 400;">Optimizers &amp; Multi-Loss Training</span></h2>
+<p><span style="font-weight: 400;">New training algorithms and loss-aggregation techniques moving beyond AdamW and single-loss optimization.</span></p>
+<p><b>Beyond AdamW: A Practical PyTorch Walkthrough of Muon, Dion, and Orthogonalized Optimizer Variants</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">Jennifer Wei (Bird of Paradise AI) </span></p>
+<p><span style="font-weight: 400;">Wed 21 Oct, 14:50–15:15, LL21DEF </span></p>
+<p><span style="font-weight: 400;">A hands-on walkthrough of Muon&#8217;s matrix-orthogonalization approach to hidden-layer updates &#8211; momentum, Newton-Schulz iteration, parameter grouping, AdamW compatibility &#8211; bridging from single-GPU implementation to distributed training and touching
+<p><b>TorchJD: Training PyTorch Models with Multiple Losses</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">Valérian Rey, Khush Patel (SimplexLab) </span></p>
+<p><span style="font-weight: 400;">Wed 21 Oct, 15:05–15:15, LL20AB </span></p>
+<p><span style="font-weight: 400;">Now part of the PyTorch ecosystem, TorchJD implements Jacobian descent, aggregating the per-loss gradient Jacobian so parameter updates benefit every objective, with applications spanning multitask learning, machine unlearning, federated learning, and adversarial f
+<p><span style="font-weight: 400;">Reinforcement Learning Systems</span></p>
+<p><span style="font-weight: 400;">Open frameworks and infrastructure for agentic and multi-turn RL — from standardized environments to the training frameworks and numerical-consistency fixes that keep large-scale RL stable.</span></p>
+<p><b>Open Source Reinforcement Learning with Agent Harnesses</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">Ben Burtenshaw (Hugging Face) </span></p>
+<p><span style="font-weight: 400;">Tue 20 Oct, 17:30–17:55, 210AE </span></p>
+<p><span style="font-weight: 400;">OpenEnv is an interoperability layer, co-owned by Hugging Face, Meta, Unsloth, Prime Intellect, Modal, NVIDIA, and Mercor,  that standardizes how RL environments are published, deployed, and consumed via a consistent HTTP/WebSocket/MCP API, decoupling environments
+<p><b>Torchtitan RL: A Unified and Extensible Training Framework for Agentic Tasks</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">Felipe Mello, Jiani Wang (Meta) </span></p>
+<p><span style="font-weight: 400;">Wed 21 Oct, 16:20–16:45, LL21DEF </span></p>
+<p><span style="font-weight: 400;">Addresses the two core instabilities in frontier-scale asynchronous RL &#8211; numerical mismatch between training and generation, and infrastructure complexity. This is done through a single shared model definition for both roles and a hackable, forkable stack bui
+<p><b>When Rollout and Training Disagree: Lessons from Building a Mismatch-Free RL Engine</b><span style="font-weight: 400;"> Neiwen Ling (ByteDance), Tianle Zhong (University of Virginia) </span></p>
+<p><span style="font-weight: 400;">Wed 21 Oct, 16:55–17:05, LL21DEF </span></p>
+<p><span style="font-weight: 400;">Introduces VeXact, an open-source rollout backend that aligns rollout log-probabilities with the training path in verl-based RL systems, showing that small token-level numerical disagreements between inference and training stacks can independently destabilize PPO/G
+<p><span style="font-weight: 400;">Inference &amp; Serving Optimization</span></p>
+<p><span style="font-weight: 400;">Tools that make LLM serving faster and cheaper to operate, from checkpoint loading to prompt caching.</span></p>
+<p><b>Faster LLM Serving Startup with fastsafetensors</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">Takeshi Yoshimura (IBM) </span></p>
+<p><span style="font-weight: 400;">Tue 20 Oct, 15:25–15:35, LL20CD </span></p>
+<p><span style="font-weight: 400;">fastsafetensors treats checkpoint loading as a data-movement problem, removing per-tensor copies and host staging to deliver 4.8x–7.5x faster model loading and up to 28 GB/s of NVMe throughput, now a portable, cross-vendor library with ROCm and Windows DirectStorag
+<p><b>LMCache: A Cluster-Wide Open-Source Solution for LLM Prompt Caching</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">Kuntai Du (Tensormesh, Inc.) </span></p>
+<p><span style="font-weight: 400;">Tue 20 Oct, 15:40–15:50, LL20CD </span></p>
+<p><span style="font-weight: 400;">LMCache is one of the most widely adopted open-source prompt-caching solutions, spanning vLLM, SGLang, and TensorRT-LLM inference engines and storage backends from Mooncake to Redis to S3. This talk includes a quick tutorial for deploying it on Kubernetes.</span></
+<h2><span style="font-weight: 400;">Hardware Portability &amp; Accelerator Backends</span></h2>
+<p><span style="font-weight: 400;">The backends and compilers bringing native PyTorch support to TPUs, Trainium, Spyre, dataflow accelerators, Intel XPUs, and edge hardware — the work that keeps PyTorch truly hardware-agnostic.</span></p>
+<p><b>Unlocking PyTorch for Dataflow Accelerators: An Open-Source Kernel Tile IR and Dataflow Scheduler</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">Prasanth Chatarasi, Bardia Mahjour, Viji Srinivasan (IBM) </span></p>
+<p><span style="font-weight: 400;">Tue 20 Oct, 16:55–17:20, 210BF </span></p>
+<p><span style="font-weight: 400;">KTIR, an open-source MLIR-based tile IR, extends data-parallel abstractions with distributed scratchpads and inter-tile communication, paired with an architecture-agnostic dataflow scheduler. Both are open sourced and integrated with PyTorch via TorchInductor to po
+<p><b>Portable PyTorch Across AI Accelerators: A Triton Operator Stack from Eager Mode to vLLM</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">Yonghua Lin (Beijing Academy of Artificial Intelligence) </span></p>
+<p><span style="font-weight: 400;">Tue 20 Oct, 16:55–17:20, LL20CD </span></p>
+<p><span style="font-weight: 400;">FlagOS is an open source Triton-based operator, compiler, and runtime stack (FlagGems, FlagTree, vllm-plugin-fl) tested on 20+ chips and architectures, enabling Day-0 support for frontier open models across NVIDIA GPUs, non-CUDA accelerators, and ARM CPUs.</span></
+<p><b>Integrating the IBM Spyre Accelerator</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">David Grove, Antoni Viros Martin, Avery Blanchard (IBM) </span></p>
+<p><span style="font-weight: 400;">Tue 20 Oct, 16:55–17:20, LL21DEF </span></p>
+<p><span style="font-weight: 400;">Torch-Spyre, an open source PyTorch PrivateUse1 device with an Inductor backend for IBM&#8217;s 32-core Spyre accelerator, now runs thousands of Hugging Face and vLLM models. This talk covers device-specific tensor layouts, scratchpad tiling, and the upstream contr
+<p><b>TorchNeuron: Native PyTorch on AWS Trainium — From Research to Production Without Compromise</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">Yahav Biran (Annapurna Labs) </span></p>
+<p><span style="font-weight: 400;">Tue 20 Oct, 17:30–17:40, LL21ABC</span></p>
+<p><span style="font-weight: 400;">An open-source native PyTorch backend for Trainium via PrivateUse1 — change </span><span style="font-weight: 400;">.to(&#8216;cuda&#8217;)</span><span style="font-weight: 400;"> to </span><span style="font-weight: 400;">.to(&#8216;neuron&#8217;)</span><span style="
+<p><b>TorchTPU: Running PyTorch Natively on Google TPUs</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">Claudio Basile (Google) </span></p>
+<p><span style="font-weight: 400;">Wed 21 Oct, 11:10–11:35, LL21ABC </span></p>
+<p><span style="font-weight: 400;">TorchTPU brings native, high-performance PyTorch to TPUs (previously JAX-only) through an ATen-to-StableHLO lowering path and &#8220;DeferAndFuse&#8221; execution for automated op fusion, integrating with vLLM for serving and TorchTitan for training, and moving to
+<p><b>Integrating TorchTPU into TorchTitan and Optimizing Models for TPUs</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">Aleksey Vlasenko (Google), Will Constable (Meta) </span></p>
+<p><span style="font-weight: 400;">Wed 21 Oct, 17:10–17:20, LL21DEF </span></p>
+<p><span style="font-weight: 400;">A concise look at the process of adding TorchTPU support into TorchTitan and optimizing selected models to run on TPU, with a direct performance comparison against GPU.</span></p>
+<p><b>From PyTorch to the Edge: Agentic Synthesis of Inference Runtimes for Heterogeneous Hardware</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">Thomas Cottenier (Arm) </span></p>
+<p><span style="font-weight: 400;">Wed 21 Oct, 17:30–17:55, LL20AB </span></p>
+<p><span style="font-weight: 400;">An agentic harness that synthesizes bespoke inference runtimes per edge target, routing through torch.export, ExecuTorch backends, torchao quantization, or direct kernel generation as needed. This is then validated on Apple silicon against llama.cpp and MLX, then e
+<p><b>Scaling PyTorch on Intel XPU: Field Notes from AI for Science on Aurora</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">Sam Foreman (Argonne National Laboratory), Panagiotis Kourdis, Tanima Dey (Intel) </span></p>
+<p><span style="font-weight: 400;">Wed 21 Oct, 17:30–17:55,  LL20CD </span></p>
+<p><span style="font-weight: 400;">Field notes from training scientific foundation models in production on 63,000+ Intel XPUs on Argonne&#8217;s Aurora system, covering the migration from Megatron-DeepSpeed to torchtitan + DTensor + FSDP2, and the operational machinery (silent-hang detection, bad-no
+<h2><span style="font-weight: 400;">Profiling, Observability &amp; Performance Diagnostics</span></h2>
+<p><span style="font-weight: 400;">Tools and techniques for actually seeing what your PyTorch code is doing on the hardware &#8211; reading traces, diagnosing bottlenecks, and estimating memory and runtime before you touch a cluster.</span></p>
+<p><b>What You Cannot Profile, You Cannot Optimize: Learning to Read PyTorch Traces</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">Aritra Roy Gosthipaty, Suvaditya Mukherjee (Hugging Face) </span></p>
+<p><span style="font-weight: 400;">Tue 20 Oct, 14:50–15:15, LL20AB </span></p>
+<p><span style="font-weight: 400;">A Socratic, no-prerequisites walkthrough of PyTorch profiler traces, from a simple matmul-plus-bias workload up to a real GeGLU MLP, teaching durable mental models for overhead-bound vs. compute-bound execution and the CPU→ATen→cuBLAS→GPU dispatch chain.</span></p>
+<p><b>Observability Tooling for Cudagraph Workloads</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">Natalia Gimelshein, Driss Guessous (Meta) </span></p>
+<p><span style="font-weight: 400;">Tue 20 Oct, 16:20–16:30, LL21DEF </span></p>
+<p><span style="font-weight: 400;">CUDA graphs help solve CPU overhead on Blackwell-class hardware but make profiles harder to read. This talk covers new PyTorch utilities that make cudagraph profiling and memory monitoring information-rich with near-zero overhead, by cross-referencing capture-time
+<p><b>Five Bottlenecks, Five Fixes: Systematic Performance Diagnosis for PyTorch Distributed Training</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">Paulo Aragao (Amazon Web Services) </span></p>
+<p><span style="font-weight: 400;">Tue 20 Oct, 16:55–17:20, 210AE </span></p>
+<p><span style="font-weight: 400;">Five real case studies &#8211; from an 83x DataLoader speedup to a 6.8x end-to-end gain on an autonomous-driving model &#8211; mapped to a repeatable diagnostic framework for classifying distributed training bottlenecks as compute-, memory-, communication-, or I/O-
+<p><b>TorchInsights: Zero-GPU Memory &amp; Runtime Estimation for Distributed Training and Agentic Research</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">Sanket Jayant Purandare, Aditya Venkataraman (Meta) </span></p>
+<p><span style="font-weight: 400;">Tue 20 Oct, 17:45–17:55, LL21ABC </span></p>
+<p><span style="font-weight: 400;">Using fake tensors and simulated multi-stream GPU execution, TorchInsights predicts peak memory and runtime across FSDP/TP/EP/CP/PP configurations before touching a real cluster, published at ICML 2025 as TorchSim and fully open source under meta-pytorch.</span></p
+<h2><span style="font-weight: 400;">Data Loading, Storage &amp; Media Pipelines</span></h2>
+<p><span style="font-weight: 400;">Feeding hungry GPUs &#8211; high-speed storage protocols, dataloader scaling, and efficient image/video/audio preprocessing.</span></p>
+<p><b>Keeping GPUs Busy: High-Speed Storage for PyTorch via fsspec</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">Ankita Luthra, Trinadh Kotturu (Google) </span></p>
+<p><span style="font-weight: 400;">Tue 20 Oct, 15:40–15:50, LL21DEF </span></p>
+<p><span style="font-weight: 400;">&#8220;Rapid Storage&#8221; brings Google&#8217;s Colossus stateful protocol to PyTorch through the fsspec interface, bypassing REST APIs via persistent gRPC streams to achieve sub-1ms random I/O latency and 6 TB/s aggregate throughput for training and checkpointin
+<p><b>From Bytes to Tensors: Efficient Media Processing with PyTorch in 2026</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">Nicolas Hug, Scott Schneider (Meta) </span></p>
+<p><span style="font-weight: 400;">Wed 21 Oct, 16:20–16:45, 210BF </span></p>
+<p><span style="font-weight: 400;">A tour of how TorchVision, TorchCodec, and TorchAudio together form an efficient end-to-end media preprocessing pipeline, plus techniques from the past year &#8211; CUDA-accelerated decoding, SIMD-optimized transforms, pre-computed frame indices &#8211; that elimin
+<p><b>Scaling PyTorch Dataloaders 10x: A Case Study</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">Ian Stenbit, Christine Cheng, Dylan Doblar (NVIDIA) </span></p>
+<p><span style="font-weight: 400;">Wed 21 Oct, 16:55–17:20, 210BF</span></p>
+<p><span style="font-weight: 400;">Over two years, NVIDIA&#8217;s autonomous vehicles team scaled a research dataloader to full production across thousands of GPUs, sharing lessons on profiling with Nsight and cprofiler, distributed filesystems, CUDA pipelining, and scaling video decoding to hit a 1
+<h2><span style="font-weight: 400;">AI Agents, Open-Source Community &amp; Release Engineering</span></h2>
+<p><span style="font-weight: 400;">How AI agents are reshaping PyTorch&#8217;s own development process, plus the open-source education and release-engineering tooling that keeps the ecosystem healthy.</span></p>
+<p><b>Contributing to PyTorch with AI Agents</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">Edward Yang (Meta) </span></p>
+<p><span style="font-weight: 400;">Tue 20 Oct, 10:35–11:05, Community Expo (Birds of a Feather) </span></p>
+<p><span style="font-weight: 400;">A community discussion on how AI agents can be used productively to contribute to PyTorch itself &#8211; how to get a PR reviewed, what maintainers want to see, and lessons from the project&#8217;s own devlog on AI coding practices.</span></p>
+<p><b>Relay and Reuse: The Dual Engine Behind PyTorch Out-of-Tree Release Readiness</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">Jiahao Chen, Jiahao Tan (Huawei) </span></p>
+<p><span style="font-weight: 400;">Tue 20 Oct, 11:10–11:35, LL21DEF </span></p>
+<p><span style="font-weight: 400;">How device-agnostic test reuse and a Cross-Repo CI Relay (CRCR) let out-of-tree hardware backends ship high-quality releases within 30 days of each PyTorch update, turning 580K+ community tests into a shared safety net.</span></p>
+<p><b>Scaling PyTorch&#8217;s Compatibility Promise: A Tiered Cross-Repository CI Relay for Out-of-Tree Backends</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">Subin George, Jewel K M (Red Hat) </span></p>
+<p><span style="font-weight: 400;">Tue 20 Oct, 12:00–12:10, LL21DEF </span></p>
+<p><span style="font-weight: 400;">A deeper technical dive into CRCR&#8217;s four-tier trust model (event dispatch through blocking merge prerequisites), its DynamoDB/ClickHouse ingestion pipeline, and how it cut breakage detection for out-of-tree backends like Ascend NPU and RISC-V from days to min
+<p><b>Fighting Agents with Agents — Bringing Claude to PyTorch CI, Triage, and PR Review</b><br />
+<span style="font-weight: 400;">Driss Guessous (Meta) </span></p>
+<p><span style="font-weight: 400;">Tue 20 Oct, 12:35–12:45, LL21DEF </span></p>
+<p><span style="font-weight: 400;">How the PyTorch project brought Claude into its own infrastructure: </span><span style="font-weight: 400;">@claude</span><span style="font-weight: 400;"> on issues and PRs, automatic issue triage, reusable onboarding skills, and CI/autorevert investigation &#8211;
+<p><b>From Scratch to PyTorch: Demystifying ML Frameworks by Building Your Own</b><span style="font-weight: 400;"> </span><br />
+<span style="font-weight: 400;">Andrea Mattia Garavagno (University of Genoa), Vijay Janapa Reddi (Harvard University) </span></p>
+<p><span style="font-weight: 400;">Tue 20 Oct, 14:15–14:40, LL20AB </span></p>
+<p><span style="font-weight: 400;">TinyTorch is an open-source CLI teaching tool which guides developers through rebuilding PyTorch&#8217;s foundational components &#8211; tensors, autograd, optimizers, transformers &#8211; in pure Python, closing the gap between using a framework and understanding
+<h2><span style="font-weight: 400;">Conclusion</span></h2>
+<p><span style="font-weight: 400;">Taken together, these sessions tell a consistent story: the PyTorch ecosystem&#8217;s edge doesn&#8217;t just come from bigger models &#8211; it comes from relentless, open, collaborative work on the tools underneath them. Whether it&#8217;s a new DSL that lets one
+<p><span style="font-weight: 400;">If you work anywhere near performance, compilers, distributed systems, or open-source ML infrastructure, this track alone is worth the trip. </span><a href="https://hubs.ly/Q04tDx8f0"><b>Register now for PyTorch Conference North America 2026</b></a><span style="fon
+<p><i>PyTorch Conference is the open source AI community&#8217;s town square. Where what&#8217;s next gets decided.</i></p>
+<p>&nbsp;</p>
+]]></content:encoded>
+
+
+
+ </item>
+ <item>
+ <title>Low Precision Flash Attention 4: End-to-End Block-Scaled Attention for Blackwell</title>
+ <link>https://pytorch.org/blog/low-precision-flash-attention-4-end-to-end-block-scaled-attention-for-blackwell/</link>
+
+ <dc:creator><![CDATA[Dev (Devashish) Shankar, Darren Liu, Chunzhi Yang, Jackie (Jiaqi) Xu,Markus Hoehnerbach, Jason Xie, Santosh Mohan, Han Xu, Rich Zhu, Josh Fromm, Hongtao Yu, Max Leung, and John Bocharov]]></dc:creator>
+ <pubDate>Wed, 16 Sep 2026 18:55:21 +0000</pubDate>
+ <category><![CDATA[Blog]]></category>
+ <guid isPermaLink="false">https://pytorch.org/?p=159902</guid>
+
+ <description><![CDATA[TL;DR We extend FlashAttention-4 [1] with MXFP8 forward and backward, reaching 2.85 PF/s forward and 2 PF/s backward on LLM shapes. On our internal shapes, FA4 MX8 reaches 2.54 PF/s...]]></description>
+ <content:encoded><![CDATA[<h2><span style="font-weight: 400;">TL;DR</span></h2>
+<p><span style="font-weight: 400;">We extend FlashAttention-4 [1] with MXFP8 forward and backward, reaching 2.85 PF/s forward and 2 PF/s backward on LLM shapes. On our internal shapes, FA4 MX8 reaches 2.54 PF/s forward and 1.58 PF/s backward, delivering up to 1.6× and 1.52× gains over BF16. We fuse
+<h2><span style="font-weight: 400;">1. Introduction</span></h2>
+<p><span style="font-weight: 400;">Blackwell&#8217;s tensor cores introduce block-scaled MMA instructions (tcgen05.mma.block_scale) that operate natively on microscaling formats &#8212; MXFP8, MXFP6, MXFP4 &amp; NVFP4 &#8212; delivering 2-4x the throughput of BF16 MMA [4,5]. However, exploiting this
+<p><span style="font-weight: 400;">In this work, we extend the FA4 attention kernel with end-to-end MXFP8 support for both forward and backward passes, and integrate it into a cross-attention module for Ads training with fused producer and output epilogues. The key contributions are: (1) TMEM alloca
+<h2><span style="font-weight: 400;">2. Implementation Details</span></h2>
+<h3><span style="font-weight: 400;">2.1 Attention Forward</span></h3>
+<p><span style="font-weight: 400;">Attention forward consists of the following primary operations:</span></p>
+<p><span style="font-weight: 400;">S = Q @ K.T</span></p>
+<p><span style="font-weight: 400;">P = Softmax(S)</span></p>
+<p><span style="font-weight: 400;">O = P @ V</span></p>
+<p><span style="font-weight: 400;">To enable blockscaled MMA, we follow existing CuTe DSL examples from Quack GEMM kernels and CUTLASS C++ examples [5,6]. We use TMA loads to fetch scale factors from GMEM to SMEM, and copy SFs from SMEM to TMEM before triggering the UMMA. The primary challenge here
+<p><span style="font-weight: 400;">Currently, in the softmax warp, softmax computation happens in FP32, and then the results are converted to BF16 before the PV multiplication. We convert P to MXFP8 while also computing the scales. We deep dive into the PTX optimizations done to achieve this efficie
+<p><span style="font-weight: 400;">One subtle thing to note is that for P.V blockscaled MMA to work, the scales need to be computed along the MMA K-dim. For Q and K, this is the embedding dimension (D) of attention, but for V, the scales and quantization need to be computed along the sequence dimens
+<h4><span style="font-weight: 400;">2.1.1 TMEM Allocation and barrier synchronization</span></h4>
+<p><span style="font-weight: 400;">Blackwell architecture has a fixed TMEM size of 512 column, which is completely utilized for MMA operands and accumulators in existing Blackwell FA kernels. This makes it challenging to add block scaled MMA, since scales also need to be in TMEM.</span></p>
+<p><span style="font-weight: 400;">FA4 forward uses a ping-pong computation between two Q tiles. We load  two Q tiles, Q0 &amp; Q1 of size [128, 128], and loop over K/V tiles (N dimension). The order of GEMMs is:</span></p>
+<table style="border-collapse: collapse; width: 100%; max-width: 520px; font-family: Arial, Helvetica, sans-serif; font-size: 24px; line-height: 1.25; background: #ffffff;">
+<tbody>
+<tr>
+<td style="border: 2px solid #111111; padding: 14px 12px;"><strong>GEMM</strong></td>
+</tr>
+<tr>
+<td style="border: 2px solid #111111; padding: 14px 12px;"><strong>Prologue</strong></td>
+</tr>
+<tr>
+<td style="border: 2px solid #111111; padding: 14px 12px; color: #366fbc;"><strong>S0</strong> = Q0 @ K0</td>
+</tr>
+<tr>
+<td style="border: 2px solid #111111; padding: 14px 12px; color: #366fbc;"><strong>S1</strong> = Q1 @ K0</td>
+</tr>
+<tr>
+<td style="border: 2px solid #111111; padding: 14px 12px;"><strong>Mainloop</strong> (for n in 0 .. N-1)</td>
+</tr>
+<tr>
+<td style="border: 2px solid #111111; padding: 14px 12px; color: #6a4adf;"><strong>O0</strong> = P0_n * V_n</td>
+</tr>
+<tr>
+<td style="border: 2px solid #111111; padding: 14px 12px; color: #366fbc;"><strong>S0</strong> = Q0 * K_{n+1}</td>
+</tr>
+<tr>
+<td style="border: 2px solid #111111; padding: 14px 12px; color: #d7622b;"><strong>O1</strong> = P1_n * V_n</td>
+</tr>
+<tr>
+<td style="border: 2px solid #111111; padding: 14px 12px; color: #366fbc;"><strong>S1</strong> = Q1 * K_{n+1}</td>
+</tr>
+<tr>
+<td style="border: 2px solid #111111; padding: 14px 12px;"><strong>Epilogue</strong></td>
+</tr>
+<tr>
+<td style="border: 2px solid #111111; padding: 14px 12px; color: #6a4adf;"><strong>O0</strong> = P0_N * V_N</td>
+</tr>
+<tr>
+<td style="border: 2px solid #111111; padding: 14px 12px; color: #d7622b;"><strong>O1</strong> = P1_N * V_N</td>
+</tr>
+</tbody>
+</table>
+<p><span style="font-weight: 400;">This is the TMEM Alloc:</span><img fetchpriority="high" decoding="async" class="aligncenter wp-image-159911 size-full" src="https://pytorch.org/wp-content/uploads/2026/08/1.png" alt="" width="1788" height="426" srcset="https://pytorch.org/wp-content/uploads/2026/08
+<ul>
+<li style="font-weight: 400;" aria-level="1"><span style="font-weight: 400;">For the prologue S(i) GEMMs, we can use O(i) for S(i) SFs, as O(i) hasn’t started yet</span></li>
+<li style="font-weight: 400;" aria-level="1"><span style="font-weight: 400;">SFs for O(i) can live overlapped with S(i) &#8211; this is the same strategy used by regular FA which overlaps P(i) with S(i). Thus, there already exists a barrier which ensures that S(i) TMEM is consumed, before we copy O(
+<li style="font-weight: 400;" aria-level="1"><span style="font-weight: 400;">SFs for S(i) live overlap with S(1-i). This requires an additional barrier between the MMA and Softmax warps, since MMA is executed asynchronously, it is possible that there is a write-write conflict between S(1-i) accumula
+</ul>
+<p><span style="font-weight: 400;">Concretely, the GEMM execution order with SF placement is shown below:</span></p>
+<table style="border-collapse: collapse; width: 100%; max-width: 1040px; font-family: Arial, Helvetica, sans-serif; font-size: 24px; line-height: 1.25; background: #ffffff;">
+<colgroup>
+<col style="width: 50%;" />
+<col style="width: 50%;" /> </colgroup>
+<tbody>
+<tr>
+<td style="border: 2px solid #111111; padding: 14px 12px;"><strong>GEMM</strong></td>
+<td style="border: 2px solid #111111; padding: 14px 12px;"><strong>SF TMEM Region</strong></td>
+</tr>
+<tr>
+<td style="border: 2px solid #111111; padding: 14px 12px;" colspan="2"><strong>Prologue</strong></td>
+</tr>
+<tr>
+<td style="border: 2px solid #111111; padding: 14px 12px; color: #366fbc;"><strong>S0</strong> = Q0 @ K0</td>
+<td style="border: 2px solid #111111; padding: 14px 12px; color: #6a4adf;">O0 (free, not started)</td>

Diff display stops at 400 lines. The line counts above are from the whole diff. 70 lines shown here cut at 300 characters. The raw artifact at this commit is linked above.