{"id":24984,"date":"2026-11-10T08:59:31","date_gmt":"2026-11-10T16:59:31","guid":{"rendered":"https:\/\/engineering.fb.com\/?p=24984"},"modified":"2026-11-09T15:33:59","modified_gmt":"2026-11-09T23:33:59","slug":"tlx-a-hardware-native-gpu-compiler-for-production-ml","status":"publish","type":"post","link":"https:\/\/engineering.fb.com\/2026\/11\/10\/open-source\/tlx-a-hardware-native-gpu-compiler-for-production-ml\/","title":{"rendered":"How We Built TLX: A Hardware-Native GPU Compiler for Production ML at Scale"},"content":{"rendered":"<ul>\n<li style=\"font-weight: 400;\" aria-level=\"1\"><a href=\"https:\/\/arxiv.org\/abs\/2605.10905\" target=\"_blank\" rel=\"noopener\"><span style=\"font-weight: 400;\">Triton Low-Level Extensions (TLX)<\/span><\/a><span style=\"font-weight: 400;\">, is an extension to Triton that helps engineers optimize GPU performance while preserving Triton&#8217;s productive programming model.\u00a0<\/span><\/li>\n<li style=\"font-weight: 400;\" aria-level=\"1\"><span style=\"font-weight: 400;\">TLX introduces explicit warp-group and cluster-level orchestration, letting developers adopt new hardware capabilities without sacrificing velocity.\u00a0<\/span><\/li>\n<li style=\"font-weight: 400;\" aria-level=\"1\"><span style=\"font-weight: 400;\">TLX has demonstrated performance improvements in RecSys training, RecSys inference, and large-scale LLM pretraining.<\/span><\/li>\n<\/ul>\n<p><span style=\"font-weight: 400;\">Modern GPU architectures increasingly rely on execution-oriented features such as warp specialization, asynchronous memory movement, tensor memory, and cluster-level execution to optimize their performance. But these same features also expose a growing gap between low-level programming models for control and high-level programming models for productivity. To bridge this gap, we built Triton Low-Level Extensions (TLX), an extension <\/span><span style=\"font-weight: 400;\">that gives engineers explicit control over GPU execution while preserving Triton&#8217;s productive programming model.\u00a0<\/span><\/p>\n<p><span style=\"font-weight: 400;\">Today, TLX powers production workloads across Meta, from AI training to large-scale LLM pretraining, enabling engineers to adopt new hardware capabilities quickly without sacrificing developer velocity. Across NVIDIA and AMD GPUs, TLX is pushing the frontier of GPU performance \u2014 combining state-of-the-art kernels with explicit execution orchestration and model-aware fusion to deliver end-to-end improvements beyond what standalone libraries can achieve.\u00a0<\/span><\/p>\n<h2><span style=\"font-weight: 400;\">From AI Models to GPU Kernels<\/span><\/h2>\n<p><span style=\"font-weight: 400;\">To understand where TLX fits, it helps to trace the path from an AI model to the GPU instructions that ultimately execute it. At Meta, machine learning (ML) workloads are built on a common software stack centered around PyTorch. Model code is transformed through graph capture, optimization, compilation, and lowering before eventually becoming GPU kernels running across large accelerator fleets. Over the past several years, Triton has evolved from an experimental kernel language into a foundational component of this stack. A significant share of production GPU kernels \u2014 both compiler-generated and hand-optimized \u2014 are now implemented in Triton. Improvements at this layer can therefore propagate broadly across training and inference workloads, making kernel performance an increasingly important lever for improving end-to-end system efficiency.<\/span><\/p>\n<figure id=\"attachment_24988\" aria-describedby=\"caption-attachment-24988\" style=\"width: 1999px\" class=\"wp-caption alignnone\"><img loading=\"lazy\" decoding=\"async\" class=\"size-full wp-image-24988\" src=\"https:\/\/engineering.fb.com\/wp-content\/uploads\/2026\/11\/Meta-TLX-Figure-1.png\" alt=\"\" width=\"1999\" height=\"1016\" srcset=\"https:\/\/engineering.fb.com\/wp-content\/uploads\/2026\/11\/Meta-TLX-Figure-1.png 1999w, https:\/\/engineering.fb.com\/wp-content\/uploads\/2026\/11\/Meta-TLX-Figure-1.png?resize=916,466 916w, https:\/\/engineering.fb.com\/wp-content\/uploads\/2026\/11\/Meta-TLX-Figure-1.png?resize=768,390 768w, https:\/\/engineering.fb.com\/wp-content\/uploads\/2026\/11\/Meta-TLX-Figure-1.png?resize=1024,520 1024w, https:\/\/engineering.fb.com\/wp-content\/uploads\/2026\/11\/Meta-TLX-Figure-1.png?resize=1536,781 1536w, https:\/\/engineering.fb.com\/wp-content\/uploads\/2026\/11\/Meta-TLX-Figure-1.png?resize=96,49 96w, https:\/\/engineering.fb.com\/wp-content\/uploads\/2026\/11\/Meta-TLX-Figure-1.png?resize=192,98 192w\" sizes=\"auto, (max-width: 992px) 100vw, 62vw\" \/><figcaption id=\"caption-attachment-24988\" class=\"wp-caption-text\">Figure 1: Normalized kernel count growth at Meta over a 10-month period (start = 1.0 in month 1). Over this period, Triton kernel count grew ~3.8x compared with ~2.0x for CUDA, reflecting Triton&#8217;s rapidly expanding role in Meta&#8217;s GPU software stack.<\/figcaption><\/figure>\n<p><span style=\"font-weight: 400;\">This shift has also changed who writes kernels. Kernel optimization is no longer the exclusive domain of GPU specialists. ML engineers routinely profile workloads, identify bottlenecks, and optimize the hottest kernels through autotuning, fusion, memory-layout improvements, or algorithmic restructuring. Triton&#8217;s productivity allows many of these iterations to happen in hours rather than days, enabling engineers to stay close to the model-development workflow. As a result, GPU kernel development has increasingly become part of the standard optimization toolkit for ML engineers.<\/span><\/p>\n<p><span style=\"font-weight: 400;\">Triton already provides a productive foundation for many memory-bound workloads through fusion and compiler-managed memory-access optimizations. Building on these strengths, TLX addresses a complementary challenge: coordinating execution groups, asynchronous data movement, and shared resources on modern GPUs.<\/span><\/p>\n<h2><span style=\"font-weight: 400;\">The Problem: When Hardware Evolves Faster Than Programming Models<\/span><\/h2>\n<p><span style=\"font-weight: 400;\">While Triton dramatically improved developer productivity, modern GPU architectures have evolved in ways that place increasing emphasis on execution orchestration rather than computation alone. NVIDIA\u2019s Hopper introduced warp specialization and the Tensor Memory Accelerator (TMA). NVIDIA\u2019s Blackwell GPU expanded this trend further with Tensor Memory (TMEM), cluster-level execution, paired-CTA execution, and increasingly asynchronous execution mechanisms. These capabilities are not simply new instructions that can be hidden behind compiler optimizations, they fundamentally change how high-performance kernels are organized.<\/span><\/p>\n<p><span style=\"font-weight: 400;\">For example, a modern FlashAttention kernel may dedicate one warp group to issuing TMA transfers while another executes tensor-core computation. A Blackwell GEMM kernel may coordinate multiple warp groups, TMEM allocations, and cluster resources simultaneously. In both cases, performance depends not only on the computation being performed, but also on how multiple execution agents cooperate.\u00a0<\/span><\/p>\n<p><span style=\"font-weight: 400;\">As we adopted Hopper and Blackwell across Meta&#8217;s ML infrastructure, we repeatedly encountered the same challenge: CUDA could express these execution structures, but doing so often required large amounts of architecture-specific orchestration code. Triton provided a dramatically more productive programming model, but many of the execution relationships required by modern hardware remained hidden behind compiler abstractions.<\/span><\/p>\n<p><span style=\"font-weight: 400;\">The result was a growing compiler catch-up problem. Each new hardware generation introduced new execution capabilities, while programmers either waited for compiler support to arrive or dropped into increasingly specialized implementations. We realized that the challenge was not simply missing compiler features, it was the absence of an appropriate programming abstraction.<\/span><\/p>\n<h2><span style=\"font-weight: 400;\">The Missing Middle: Why SIMT and SIMB Weren&#8217;t Enough<\/span><\/h2>\n<p><span style=\"font-weight: 400;\">The capabilities that matter most on recent GPUs operate neither at thread granularity nor at block granularity. Warp specialization partitions work across cooperating warp groups. TMA introduces dedicated agents for memory movement. TMEM enables new communication patterns between execution groups. Cluster execution lets multiple CTAs cooperate within a shared execution context. Organizing a kernel around these features means reasoning about how several execution agents cooperate, communicate, and synchronize \u2014 a level of structure that sits between CUDA and Triton.<\/span><\/p>\n<p><span style=\"font-weight: 400;\">This exposed what we came to view as a missing abstraction. CUDA&#8217;s SIMT model operates at too fine a granularity: It provides complete control, but requires developers to express every execution relationship manually. Triton&#8217;s blocked programming model operates at too coarse a granularity: It improves productivity, but turns many execution structures into implicit compiler decisions. Modern hardware increasingly requires something in between.<\/span><\/p>\n<p><span style=\"font-weight: 400;\">We refer to this model as<\/span> <span style=\"font-weight: 400;\">Multiple Instructions, Multiple Warps (MIMW). MIMW elevates the warp group to the fundamental unit of orchestration. Rather than reasoning about individual threads or entire blocks, programmers express how multiple execution agents cooperate, communicate, synchronize, and share resources. This is the natural level at which many Hopper and Blackwell optimizations are expressed, and it is the abstraction layer TLX implements.<\/span><\/p>\n<p><span style=\"font-weight: 400;\">Figure 2 illustrates the evolution of GPU programming models. CUDA&#8217;s SIMT model exposed execution directly at thread granularity, giving developers complete control over scheduling, synchronization, and memory movement. Triton moved toward a SIMB model, raising the abstraction to blocked tensor computations and improving productivity. MIMW occupies the space between these two extremes. Rather than treating threads or tensor blocks as the primary execution abstraction, it elevates warp groups to first-class execution agents. <\/span><\/p>\n<figure id=\"attachment_24985\" aria-describedby=\"caption-attachment-24985\" style=\"width: 1999px\" class=\"wp-caption alignnone\"><img loading=\"lazy\" decoding=\"async\" class=\"size-full wp-image-24985\" src=\"https:\/\/engineering.fb.com\/wp-content\/uploads\/2026\/11\/Meta-TLX-Figure-2.png\" alt=\"\" width=\"1999\" height=\"646\" srcset=\"https:\/\/engineering.fb.com\/wp-content\/uploads\/2026\/11\/Meta-TLX-Figure-2.png 1999w, https:\/\/engineering.fb.com\/wp-content\/uploads\/2026\/11\/Meta-TLX-Figure-2.png?resize=916,296 916w, https:\/\/engineering.fb.com\/wp-content\/uploads\/2026\/11\/Meta-TLX-Figure-2.png?resize=768,248 768w, https:\/\/engineering.fb.com\/wp-content\/uploads\/2026\/11\/Meta-TLX-Figure-2.png?resize=1024,331 1024w, https:\/\/engineering.fb.com\/wp-content\/uploads\/2026\/11\/Meta-TLX-Figure-2.png?resize=1536,496 1536w, https:\/\/engineering.fb.com\/wp-content\/uploads\/2026\/11\/Meta-TLX-Figure-2.png?resize=96,31 96w, https:\/\/engineering.fb.com\/wp-content\/uploads\/2026\/11\/Meta-TLX-Figure-2.png?resize=192,62 192w\" sizes=\"auto, (max-width: 992px) 100vw, 62vw\" \/><figcaption id=\"caption-attachment-24985\" class=\"wp-caption-text\">Figure 2: GPU architecture and execution model evolution. Recent NVIDIA GPUs evolved from mostly synchronous SIMT-style execution toward increasingly specialized and asynchronous execution, including Tensor Cores, Async Copy, TMA, WGMMA, TMEM, register reallocation, and paired-CTA MMA. This hardware evolution makes warp-group orchestration a programming-model concern rather than merely a compiler optimization.<\/figcaption><\/figure>\n<p><span style=\"font-weight: 400;\">The emergence of MIMW is not a software-design preference but a consequence of how hardware has evolved. Volta and Ampere exposed tensor cores through thread- and warp-level programming. Hopper made asynchronous data movement (enabled byTMA) and warp-group execution (enabled by WGMMA) central to high-performance kernels. Blackwell extends this trend further with TMEM, register reallocation, cluster execution, and paired-CTA MMA. Each generation pushes more cooperation across warp groups \u2014 and sometimes across CTAs \u2014 into territory that the programming model, not just the compiler, must address.<\/span><\/p>\n<figure id=\"attachment_24987\" aria-describedby=\"caption-attachment-24987\" style=\"width: 600px\" class=\"wp-caption alignnone\"><img loading=\"lazy\" decoding=\"async\" class=\"wp-image-24987\" src=\"https:\/\/engineering.fb.com\/wp-content\/uploads\/2026\/11\/Meta-TLX-Figure-3.png\" alt=\"\" width=\"600\" height=\"391\" srcset=\"https:\/\/engineering.fb.com\/wp-content\/uploads\/2026\/11\/Meta-TLX-Figure-3.png 1206w, https:\/\/engineering.fb.com\/wp-content\/uploads\/2026\/11\/Meta-TLX-Figure-3.png?resize=916,597 916w, https:\/\/engineering.fb.com\/wp-content\/uploads\/2026\/11\/Meta-TLX-Figure-3.png?resize=768,501 768w, https:\/\/engineering.fb.com\/wp-content\/uploads\/2026\/11\/Meta-TLX-Figure-3.png?resize=1024,667 1024w, https:\/\/engineering.fb.com\/wp-content\/uploads\/2026\/11\/Meta-TLX-Figure-3.png?resize=96,63 96w, https:\/\/engineering.fb.com\/wp-content\/uploads\/2026\/11\/Meta-TLX-Figure-3.png?resize=192,125 192w\" sizes=\"auto, (max-width: 992px) 100vw, 62vw\" \/><figcaption id=\"caption-attachment-24987\" class=\"wp-caption-text\">Figure 3: MIMW occupies the missing middle between SIMT&#8217;s fine-grained thread control and SIMB&#8217;s block-level productivity. TLX implements MIMW by adding explicit warp-group orchestration to Triton&#8217;s productive base.<\/figcaption><\/figure>\n<p><span style=\"font-weight: 400;\">Making MIMW practical requires expressing cooperating warp or wave groups as first-class execution tasks, with explicit control over shared-memory buffer access, synchronization, and handoff between them. TLX brings these capabilities into Triton\u2019s existing tensor programming model, letting developers combine explicit execution control with familiar high-level tensor operations.<\/span><\/p>\n<h2><span style=\"font-weight: 400;\">Introducing TLX: The Missing Middle, Implemented<\/span><\/h2>\n<p><span style=\"font-weight: 400;\">We wanted\u00a0 to avoid creating yet another standalone low-level language with TLX. Instead, TLX extends Triton with a second layer of execution-oriented constructs. Programmers continue to express computation using familiar Triton tensor operations, while TLX is used only where explicit control over execution, memory movement, or synchronization is required. In other words, Triton answers <\/span><i><span style=\"font-weight: 400;\">what<\/span><\/i><span style=\"font-weight: 400;\"> computation should be performed, while TLX answers <\/span><i><span style=\"font-weight: 400;\">how<\/span><\/i><span style=\"font-weight: 400;\"> that computation should be orchestrated. The difference is easiest to see through a simple tiled matrix multiplication. In standard Triton, execution follows a synchronous load-compute loop: Each tile is loaded, consumed, and then discarded before the next tile begins.<\/span><\/p>\n<pre class=\"line-numbers\"><code class=\"language-none\"># Standard Triton: synchronous load-compute loop semantics, \r\n# with compiler-managed software pipelining  \r\n@triton.jit\r\ndef matmul_kernel(a_ptr, b_ptr, c_ptr, M, N, K, ...):\r\n    pid_m = tl.program_id(0)\r\n    pid_n = tl.program_id(1)\r\n    acc = tl.zeros((BLOCK_M, BLOCK_N), dtype=tl.float32)\r\n    for k in range(0, K, BLOCK_K):\r\n        a = tl.load(a_ptr + ...)   # Load A tile from global memory\r\n        b = tl.load(b_ptr + ...)   # Load B tile from global memory\r\n        acc += tl.dot(a, b)        # Compute matmul \r\n    tl.store(c_ptr + ..., acc)\r\n<\/code><\/pre>\n<p><span style=\"font-weight: 400;\">TLX preserves the same computation but makes the execution structure explicit. Producer warp groups move data asynchronously, consumer warp groups perform tensor-core computation, and synchronization is expressed directly through lightweight barriers. The computation itself remains Triton; only the orchestration layer changes. The result is a deeply pipelined kernel that overlaps memory movement, computation, and output processing without leaving Python or sacrificing Triton&#8217;s autotuning infrastructure. More importantly, the execution structure is visible in the source code rather than hidden inside compiler heuristics.<\/span><\/p>\n<pre class=\"line-numbers\"><code class=\"language-none\"># TLX on Blackwell: explicit warp-group orchestration and pipelining \r\n@triton.jit\r\ndef matmul_kernel(desc_a, desc_b, desc_c, M, N, K, ...):\r\n    buffers_a = tlx.local_alloc((BLOCK_M, BLOCK_K), tl.float16, NUM_STAGES)\r\n    buffers_b = tlx.local_alloc((BLOCK_K, BLOCK_N), tl.float16, NUM_STAGES)\r\n    acc = tlx.local_alloc((BLOCK_M, BLOCK_N), tl.float32, 1, tlx.storage_kind.tmem)\r\n    load_bars = tlx.alloc_barriers(num_barriers=NUM_STAGES, arrive_count=1)\r\n    dot_bars  = tlx.alloc_barriers(num_barriers=NUM_STAGES, arrive_count=1)\r\n    with tlx.async_tasks():\r\n        # Warp group 0: PRODUCER \u2014 async TMA loads into shared memory\r\n        with tlx.async_task(\"default\"):\r\n             # Wait for last matmul\r\n            tlx.barrier_wait(dot_bars[num_k_tiles - 1], phase)  \r\n            result = tlx.local_load(acc[0])\r\n            desc_c.store(offsets, result.to(tl.float16))\r\n            \r\n        # Warp group 1: CONSUMER \u2014 tensor-core matmul from shared memory\r\n        with tlx.async_task(num_warps=1):\r\n            for k in range(num_k_tiles):\r\n                buf = k % NUM_STAGES\r\n                 # Wait for TMA load to complete\r\n                tlx.barrier_wait(load_bars[buf], phase)      \r\n                tlx.async_dot(buffers_a[buf], buffers_b[buf], acc[0],\r\n                              use_acc=(k &gt; 0), mBarriers=[dot_bars[buf]])\r\n         \r\n        # Warp group 2: EPILOGUE \u2014 store result to global memory\r\n        with tlx.async_task(num_warps=1):\r\n            for k in range(num_k_tiles):\r\n                buf = k % NUM_STAGES\r\n                 # Wait for consumer to release buffer\r\n                tlx.barrier_wait(dot_bars[buf], phase)       \r\n                tlx.barrier_expect_bytes(load_bars[buf], ...)\r\n                tlx.async_descriptor_load(desc_a, buffers_a[buf], offsets, load_bars[buf])\r\n                tlx.async_descriptor_load(desc_b, buffers_b[buf], offsets, load_bars[buf])<\/code><\/pre>\n<p><span style=\"font-weight: 400;\">TLX extends Triton in five key areas:\u00a0<\/span><\/p>\n<p><b>Warp-group orchestration:<\/b><span style=\"font-weight: 400;\"> TLX treats warp groups as first-class execution agents. Programmers can assign explicit roles to cooperating warp groups \u2014 for example, producer warp groups that move data, consumer warp groups that execute tensor-core computation, and epilogue warp groups that process results.\u00a0<\/span><\/p>\n<p><b>Synchronization: <\/b><span style=\"font-weight: 400;\">TLX exposes synchronization at the warp-group level rather than the thread level. Programmers express dependencies between cooperating execution groups \u2014 for example, producer-consumer relationships in a software pipeline \u2014 using hardware primitives such as mbarriers. Fine-grained thread-level synchronization remains handled by the underlying execution model and compiler, letting developers focus on coordinating larger execution structures that increasingly determine performance on modern GPUs.<\/span><\/p>\n<p><b>Memory orchestration:<\/b><span style=\"font-weight: 400;\"> TLX provides direct access to hardware capabilities such as TMA, asynchronous tensor-core operations, and TMEM. These primitives allow memory movement and computation to overlap naturally within the execution structure defined by the program, while avoiding much of the low-level setup and bookkeeping typically required in CUDA.<\/span><\/p>\n<p><b>Cluster execution:<\/b><span style=\"font-weight: 400;\"> TLX exposes cluster-level execution as a first-class programming construct. Multiple CTAs can cooperate through distributed shared memory and cluster-wide synchronization primitives, enabling on-chip communication patterns that would otherwise require global-memory round trips.<\/span><\/p>\n<p><b>Optional layout: <\/b><span style=\"font-weight: 400;\">Unlike CUDA and CUTLASS, TLX does not require programmers to explicitly specify layouts before computation can be expressed. The compiler infers reasonable defaults automatically, and programmers only specify layouts when they want to override them for performance reasons. A correct kernel can therefore be written without any layout annotations at all, with layout tuning introduced incrementally during optimization.<\/span><\/p>\n<p><span style=\"font-weight: 400;\">An important consequence is that kernel fusion becomes significantly easier. Fusion is one of the most effective GPU optimizations because it eliminates intermediate memory traffic and kernel launch overhead. In traditional CUDA-style workflows, however, fusion often becomes dominated by layout management. Each stage may assume a different tiling, swizzling, or shared-memory mapping, forcing programmers to manually reconcile layouts across operation boundaries.\u00a0<\/span><\/p>\n<p><span style=\"font-weight: 400;\">Because TLX infers layouts automatically, operations can often be fused without manually reconciling intermediate layouts. The compiler resolves layout compatibility where possible, and programmers intervene only when a particular layout choice improves performance. This significantly lowers the engineering resources needed for building complex fused kernels and enables multi-stage pipelines that would otherwise require multiple kernels or substantial layout-specific infrastructure.<\/span><\/p>\n<p><span style=\"font-weight: 400;\">Taken together, these capabilities allow TLX to occupy the space between Triton&#8217;s productivity and CUDA&#8217;s control. Developers retain the advantages of Python, autotuning, and blocked tensor programming while gaining direct access to the execution structures that increasingly determine performance on modern GPUs.<\/span><\/p>\n<h2><span style=\"font-weight: 400;\">TLX in Production: Real-World Impact<\/span><\/h2>\n<p><span style=\"font-weight: 400;\">We&#8217;ve been running TLX in production across Meta&#8217;s ML infrastructure for the past several months. Three deployments \u2014 spanning RecSys training, RecSys inference, and large-scale LLM pretraining \u2014 highlight different aspects of what TLX enables.<\/span><\/p>\n<table border=\"1\">\n<tbody>\n<tr>\n<td><b>Case study<\/b><\/td>\n<td><b>Hardware<\/b><\/td>\n<td><b>Problem TLX Addresses<\/b><\/td>\n<td><b>Headline Results<\/b><\/td>\n<td><b>Deployment<\/b><\/td>\n<\/tr>\n<tr>\n<td><a href=\"https:\/\/engineering.fb.com\/2025\/11\/10\/ml-applications\/metas-generative-ads-model-gem-the-central-brain-accelerating-ads-recommendation-ai-innovation\/\"><b>GEM<\/b><\/a><b> Model Training<\/b><\/td>\n<td><span style=\"font-weight: 400;\">NVIDIA Blackwell<\/span><\/td>\n<td><span style=\"font-weight: 400;\">Reaching roofline MFU on heterogeneous RecSys foundation models needs warp specialization and explicit tensor-memory control beyond standard Triton.<\/span><\/td>\n<td><span style=\"font-weight: 400;\">20-25% E2E MFU on thousands of GPUs<\/span><\/td>\n<td><span style=\"font-weight: 400;\">GEM training stack.\u00a0<\/span><\/td>\n<\/tr>\n<tr>\n<td><b>In-Kernel Broadcast Optimization (IKBO)<\/b><\/td>\n<td><span style=\"font-weight: 400;\">NVIDIA Hopper<\/span><\/td>\n<td><span style=\"font-weight: 400;\">User features replicated many times per candidate batch, wasting bandwidth and compute; eliminating it needs warp-specialized multi-stage fusion.<\/span><\/td>\n<td><span style=\"font-weight: 400;\">Up to 2\/3 net-latency reduction in compute-intensive latency for co-designed models.<\/span><\/td>\n<td><span style=\"font-weight: 400;\">RecSys inference, early\u2192late ranking, GPU.<\/span><\/td>\n<\/tr>\n<tr>\n<td><b>Distributed Grouped GEMM for LLM Pretraining<\/b><\/td>\n<td><span style=\"font-weight: 400;\">NVIDIA Blackwell<\/span><\/td>\n<td><span style=\"font-weight: 400;\">Expert-parallel MoE layers require grouped GEMM with inter-node communication; existing implementations leave performance on the table at scale.<\/span><\/td>\n<td><span style=\"font-weight: 400;\">Production-validated at large-fleet scale over an extended period.<\/span><\/td>\n<td><span style=\"font-weight: 400;\">LLM pretraining.<\/span><\/td>\n<\/tr>\n<\/tbody>\n<\/table>\n<h3><\/h3>\n<h3><span style=\"font-weight: 400;\">Case Study 1: GEM Model Training<\/span><\/h3>\n<p><span style=\"font-weight: 400;\">Training Meta\u2019s <\/span><a href=\"https:\/\/engineering.fb.com\/2025\/11\/10\/ml-applications\/metas-generative-ads-model-gem-the-central-brain-accelerating-ads-recommendation-ai-innovation\/\" target=\"_blank\" rel=\"noopener\"><span style=\"font-weight: 400;\">Generative Ads Model (GEM)<\/span><\/a><span style=\"font-weight: 400;\"> is more difficult to optimize than LLM training because smaller input shapes<\/span><span style=\"font-weight: 400;\">, 10+ norms per layer, heterogeneous module types, and architectural diversity all limit MFU<\/span><span style=\"font-weight: 400;\">. <\/span><span style=\"font-weight: 400;\">Through a systematic model-kernel co-design approach spanning model architecture, custom kernels, compiler optimizations, and hardware-specific tuning, GEM training in its latest version reached O(10^24) total training flops on several thousand of NVIDIA Blackwell GPUs with 20-25% E2E MFU.<\/span><\/p>\n<h4><span style=\"font-weight: 400;\">TLX&#8217;s Role in the Multi-DSL Strategy\u00a0<\/span><\/h4>\n<p><span style=\"font-weight: 400;\">We evaluated three programming models on Blackwell \u2014 CUTLASS, CuTeDSL, and TLX \u2014 and found that TLX provides the best balance between productivity and performance for most kernels. It combines Triton-like simplicity, autotuning, and composability with direct access to Blackwell-specific capabilities such as <\/span><span style=\"font-weight: 400; font-family: 'courier new', courier; color: #339966;\">async_dot<\/span><span style=\"font-weight: 400;\">, tensor memory operations, and mbarrier synchronization.<\/span><\/p>\n<p><span style=\"font-weight: 400;\">For Generalized Dot Product Attention (GDPA), a kernel in GEM, TLX achieved near-optimal performance with a fraction of the implementation complexity, speeding up training by targeting at unifying the models&#8217; many attention-like operations into a single GPU kernel, delivering up to 2x forward speedup. While different kernels impose different requirements, TLX proved to be the right abstraction for the majority of GEM workloads where both developer velocity and performance matter.<\/span><\/p>\n<p><span style=\"font-weight: 400;\">TLX enabled us to move beyond optimizing individual kernels toward kernelizing and fusing larger portions of the model. For GEM workloads, reducing memory movement often provides greater benefits than optimizing compute alone, making composability and fusion first-class concerns. To accelerate this process, we built Agent for Compute Efficiency (ACE), an AI-assisted optimization framework that automates profiling, analysis, and tuning. ACE reduced the optimization cycle for a kernel, allowing engineers to focus on higher-level optimization opportunities rather than low-level performance debugging.<\/span><\/p>\n<h3><span style=\"font-weight: 400;\">Case Study 2: In-Kernel Broadcast Optimization for RecSys<\/span><\/h3>\n<p><span style=\"font-weight: 400;\">In RecSys inference, user features are identical across all scored candidates, yet they are typically replicated dozens of times per batch to match the candidate dimension. Validated across a broad spectrum of candidate-to-user broadcast ratios, this broadcast becomes a significant source of wasted memory bandwidth and compute, with overhead that scales linearly with the number of candidates.<\/span><\/p>\n<p><span style=\"font-weight: 400;\">In-Kernel Broadcast Optimization (IKBO) eliminates this inefficiency by moving the broadcast into the kernel itself. Rather than materializing replicated user features, kernels consume user and candidate inputs at their natural, mismatched batch sizes and perform the broadcast implicitly during computation. This approach is enabled through a co-design of the model architecture, inference runtime, and kernel implementation.<\/span><\/p>\n<h4><span style=\"font-weight: 400;\">TLX Provides a Building Block for the Final Stage of This Optimization<\/span><\/h4>\n<p><span style=\"font-weight: 400;\">Consider our linear compression kernel. After a series of optimizations \u2014 including matmul decomposition, memory alignment, and broadcast fusion \u2014 the remaining bottleneck is addressed through a warp-specialized, multi-stage TLX kernel that delivers a 4x speedup on the NVIDIA H100. The final implementation collapses the entire computation into a single persistent mega-kernel. Producer and consumer warp groups execute concurrently using ping-pong scheduling, while user and candidate GEMMs share circular buffers and synchronize through mbarriers without draining the pipeline between stages. Tile-level release-acquire synchronization further enables coordination across CTAs.<\/span><\/p>\n<p><span style=\"font-weight: 400;\">For FlashAttention, IKBO co-design shifts the kernel from I\/O-bound to compute-bound by amortizing K\/V reads across candidates sharing the same user context. TLX then layers on FA3-style warp specialization with persistence, eliminating cross-CTA launch stalls that dominate at RecSys&#8217;s short query lengths and delivering up to a two-thirds reduction in compute-intensive net latency for co-designed models. Expressing this execution structure directly in CUDA would require substantial manual orchestration; TLX allows it to be expressed concisely within Triton&#8217;s Python programming model and autotuning framework.<\/span><\/p>\n<p><span style=\"font-weight: 400;\">Today, IKBO helps optimize <\/span><a href=\"https:\/\/engineering.fb.com\/2026\/03\/31\/ml-applications\/meta-adaptive-ranking-model-bending-the-inference-scaling-curve-to-serve-llm-scale-models-for-ads\/\" target=\"_blank\" rel=\"noopener\"><span style=\"font-weight: 400;\">Meta&#8217;s Adaptive Ranking Model<\/span><\/a><span style=\"font-weight: 400;\"> across both GPUs and the <\/span><a href=\"https:\/\/ai.meta.com\/blog\/meta-mtia-scale-ai-chips-for-billions\/\" target=\"_blank\" rel=\"noopener\"><span style=\"font-weight: 400;\">Meta Training and Inference Accelerator (MTIA)<\/span><\/a><span style=\"font-weight: 400;\">. More broadly, it illustrates a shift in how we approach RecSys inference. Rather than optimizing around broadcast overhead, we eliminate it at the computational primitive itself, enabling dense interaction quality regardless of candidate count.<\/span><\/p>\n<h3><span style=\"font-weight: 400;\">Case Study 3: <\/span><span style=\"font-weight: 400;\">Distributed Grouped GEMM for LLM Pretraining<\/span><\/h3>\n<p><span style=\"font-weight: 400;\">TLX&#8217;s production impact extends beyond recommendation workloads. A TLX-based distributed grouped GEMM kernel is currently deployed in Meta&#8217;s large-scale LLM pretraining infrastructure, where it has been operating continuously across a large GPU fleet for an extended training period. The kernel accelerates the grouped matrix multiplications required by mixture-of-experts (MoE) models, while tightly integrating computation and communication to reduce end-to-end latency.<\/span><\/p>\n<p><span style=\"font-weight: 400;\">At this scale, peak throughput is only one part of the equation. Reliability, debuggability, and long-term performance stability are equally important. Production kernels must sustain performance over weeks of continuous operation while remaining maintainable as models, hardware, and system software evolve.<\/span><\/p>\n<p><span style=\"font-weight: 400;\">TLX proved particularly well suited for this environment. Its support for warp-specialized execution, explicit memory orchestration, and asynchronous pipelines enabled the performance characteristics required by large-scale distributed training. At the same time, its integration with Triton&#8217;s Python programming model preserved the developer productivity needed to iterate, validate, and deploy complex kernels at production velocity.<\/span><\/p>\n<p><span style=\"font-weight: 400;\">This deployment highlights an important aspect of TLX&#8217;s design philosophy: It is not intended solely for achieving peak kernel performance, but for making advanced hardware capabilities practical to adopt in real production systems.<\/span><\/p>\n<h2><span style=\"font-weight: 400;\">Take a Deeper Look Into TLX Kernels<\/span><\/h2>\n<p><span style=\"font-weight: 400;\">For deeper technical dives into TLX kernels, see our companion PyTorch blog posts:<\/span><\/p>\n<ul>\n<li style=\"font-weight: 400;\" aria-level=\"1\"><a href=\"https:\/\/pytorch.org\/blog\/in-kernel-broadcast-optimization-co-designing-kernels-for-recsys-inference\" target=\"_blank\" rel=\"noopener\"><span style=\"font-weight: 400;\">In-Kernel Broadcast Optimization: Co-Designing Kernels for RecSys Inference<\/span><\/a><span style=\"font-weight: 400;\">.<\/span><\/li>\n<li style=\"font-weight: 400;\" aria-level=\"1\"><a href=\"https:\/\/pytorch.org\/blog\/tlx-block-attention-a-warp-specialized-blackwell-kernel-for-fixed-block-sparse-self-attention\" target=\"_blank\" rel=\"noopener\"><span style=\"font-weight: 400;\">TLX Block Attention: A Warp-Specialized Blackwell Kernel for Fixed-Block Sparse Self-Attention<\/span><\/a><span style=\"font-weight: 400;\">.<\/span><\/li>\n<li style=\"font-weight: 400;\" aria-level=\"1\"><a href=\"https:\/\/pytorch.org\/blog\/fast-2-simplicial-attention-hardware-efficient-kernels-in-tlx\/\" target=\"_blank\" rel=\"noopener\"><span style=\"font-weight: 400;\">Fast 2-Simplicial Attention: Hardware-Efficient Kernels in TLX<\/span><\/a><span style=\"font-weight: 400;\">.<\/span><\/li>\n<li style=\"font-weight: 400;\" aria-level=\"1\"><a href=\"https:\/\/pytorch.org\/blog\/enabling-cluster-launch-control-with-tlx\/\" target=\"_blank\" rel=\"noopener\"><span style=\"font-weight: 400;\">Enabling Cluster Launch Control with TLX<\/span><\/a><span style=\"font-weight: 400;\">.<\/span><\/li>\n<li style=\"font-weight: 400;\" aria-level=\"1\"><a href=\"https:\/\/pytorch.org\/blog\/towards-free-normalization-fusing-normalization-into-gemm-and-attention-kernels\/\" target=\"_blank\" rel=\"noopener\"><span style=\"font-weight: 400;\">Achieving Free Normalization: Fusing Norms into GEMM and Attention via TLX<\/span><\/a><span style=\"font-weight: 400;\">.<\/span><\/li>\n<li style=\"font-weight: 400;\" aria-level=\"1\"><a href=\"https:\/\/pytorch.org\/blog\/optimizing-jagged-flash-attention-with-tlx-the-road-toward-sota-fa4-on-blackwell\/%20\" target=\"_blank\" rel=\"noopener\"><span style=\"font-weight: 400;\">Optimizing Jagged Flash Attention with TLX: The Road Toward SOTA FA4 on Blackwell<\/span><\/a><span style=\"font-weight: 400;\">.<\/span><\/li>\n<li style=\"font-weight: 400;\" aria-level=\"1\"><a href=\"https:\/\/www.amd.com\/en\/developer\/resources\/technical-articles\/2026\/optimizing-gemm-with-tlx.html\" target=\"_blank\" rel=\"noopener\"><span style=\"font-weight: 400;\">Optimizing GEMM + Activation on CDNA4 with TLX<\/span><\/a><span style=\"font-weight: 400;\">.<\/span><\/li>\n<\/ul>\n<h2><span style=\"font-weight: 400;\">What We Learned: Advice for Building GPU Compilers<\/span><\/h2>\n<h3><span style=\"font-weight: 400;\">Start With a Real Workload, Not a Benchmark<\/span><\/h3>\n<p><span style=\"font-weight: 400;\">We built TLX to solve actual production problems at Meta rather than to win benchmark contests. Both the GEM MFU project and IKBO illustrate this principle. Each began with a concrete pain point (low MFU in GEM, broadcast overhead in RecSys), and the tooling followed. This focus kept us oriented toward real-world performance rather than synthetic scores.<\/span><\/p>\n<h3><span style=\"font-weight: 400;\">Productivity and Performance Are Multipliers, Not Trade-Offs<\/span><\/h3>\n<p><span style=\"font-weight: 400;\">TLX typically delivers\u00a0 performance on-par with hand-tuned CUDA performance, but engineers can iterate significantly faster. A kernel that is marginally faster but takes three times as long to write is a net loss at fleet scale. With GEM, the multi-DSL approach let us select TLX for rapid iteration; with IKBO, the TLX-based implementation took weeks rather than months.<\/span><\/p>\n<h3><span style=\"font-weight: 400;\">Hardware Features Are Only Useful if They Are Accessible\u00a0<\/span><\/h3>\n<p><span style=\"font-weight: 400;\">WGMMA, TMA, and TMEM are powerful capabilities, but they saw little use because the programming model was too complex. TLX makes them accessible through concise annotations. The warp-specialized fusion in IKBO and the multi-DSL kernels in GEM would have been prohibitively complex in raw CUDA.<\/span><\/p>\n<h3><span style=\"font-weight: 400;\">The Compiler Is Part of the User Interface\u00a0<\/span><\/h3>\n<p><span style=\"font-weight: 400;\">Error messages, debugging support, and profiling integration matter as much as code-generation quality. We invested as much in the developer experience as in the optimizer. When GEM engineers debugged MFU issues or IKBO engineers encountered warp-group synchronization problems, they required clear feedback rather than obscure compiler errors.<\/span><\/p>\n<h3><span style=\"font-weight: 400;\">Evolution Outperforms Revolution<\/span><\/h3>\n<p><span style=\"font-weight: 400;\">By building TLX as an extension to Triton, we leveraged the existing ecosystem, tooling, and expertise. The GEM team could use TLX alongside CUTLASS and CuTeDSL; the IKBO team, already familiar with Triton, needed only to learn the TLX annotations for performance-critical sections. A clean-slate design might have been more elegant, but might never have shipped.<\/span><\/p>\n<h2><span style=\"font-weight: 400;\">What&#8217;s Next for TLX<\/span><\/h2>\n<p><span style=\"font-weight: 400;\">TLX is going further. Here&#8217;s what we&#8217;re working on:<\/span><\/p>\n<ul>\n<li style=\"font-weight: 400;\" aria-level=\"1\"><b>More Hardware<\/b><span style=\"font-weight: 400;\">: Right now, TLX targets NVIDIA Hopper and Blackwell, and AMD Mi300 and Mi350x. We&#8217;re working on support for more architectures, including new generations from NVIDIA and AMD, and custom accelerators like MTIA. The MIMW model is hardware-agnostic \u2014 the challenge is mapping it efficiently to different ISAs. GEM already runs on both GPU and MTIA, and IKBO spans GPU and MTIA.<\/span><\/li>\n<li style=\"font-weight: 400;\" aria-level=\"1\"><b>Higher-Level Abstractions<\/b><span style=\"font-weight: 400;\">: TLX gives developers control over warp groups. Next, we want to raise the abstraction level again \u2014 automatically inferring warp-group orchestration from higher-level patterns. Think &#8220;auto-TLX&#8221; that writes the annotations for you. The patterns we learned from GEM (multi-DSL selection) and IKBO (three-stage warp specialization) are candidates for automation.<\/span><\/li>\n<li style=\"font-weight: 400;\" aria-level=\"1\"><b>Integration with PyTorch 2.x<\/b><span style=\"font-weight: 400;\">: We&#8217;re working on making TLX a backend for torch.compile. The vision is to write PyTorch code and get TLX-optimized kernels automatically, with no manual annotation required for common patterns. GEM shows what&#8217;s possible with systematic co-design, and IKBO shows the value of pattern-specific optimization. We want to make both automatic.<\/span><\/li>\n<\/ul>\n<p><span style=\"font-weight: 400;\">The missing middle is real, and it&#8217;s not just for GPUs. As accelerators get more complex, we need programming models that expose the right level of control without sacrificing productivity. TLX is our first step in that direction.<\/span><\/p>\n<h2><span style=\"font-weight: 400;\">Acknowledgements\u00a0<\/span><\/h2>\n<p><i><span style=\"font-weight: 400;\">TLX is a collaboration across the Triton, PyTorch, MSL, Ads Training Infra, Ads Ranking, RecSys, and AI Compilers teams at Meta. We\u2019d like to thank the following people for their contributions:<\/span><\/i><\/p>\n<p><b><i>Triton:<\/i><\/b><i><span style=\"font-weight: 400;\"> Shane Nay, Manman Ren, Oleksandr Stashuk, Karthik Manivannan, Zhijing Li, Yuanwei Fang, Shunting Zhang, Arda Unal, Yueming Hao, Taylor Robie, Puyan Lotfi, Corbin Robeck, Bert Maher, Alexey Loginov, Ying Liu, Bill Yoshimi.<\/span><\/i><\/p>\n<p><b><i>PG:<\/i><\/b><i><span style=\"font-weight: 400;\"> Jackie Xu, Jacky Zhou, Han Xu, Hao Yan, Jake Siso, Devashish Shankar, Shikai Li, Jianyu Huang, Andrew Tulloch, Boda Li, Jiao Jian, Zhuoran Zhao, Ao Cai, Shuyao Bi, Yuanjuan Yao, Ruizhen Guo, Doris Xiang, Jason Lu, Marvin Fan, Hai Zheng, Zhaokang Ke, Albert Chen, Mingyu Guan<\/span><\/i><\/p>\n<p><b><i>UCSD:<\/i><\/b><i><span style=\"font-weight: 400;\"> Yufei Ding and Yue Guan.<\/span><\/i><\/p>\n","protected":false},"excerpt":{"rendered":"<p>Triton Low-Level Extensions (TLX), is an extension to Triton that helps engineers optimize GPU performance while preserving Triton&#8217;s productive programming model.\u00a0 TLX introduces explicit warp-group and cluster-level orchestration, letting developers adopt new hardware capabilities without sacrificing velocity.\u00a0 TLX has demonstrated performance improvements in RecSys training, RecSys inference, and large-scale LLM pretraining. Modern GPU architectures increasingly [&#8230;]<\/p>\n<p><a class=\"btn btn-secondary understrap-read-more-link\" href=\"https:\/\/engineering.fb.com\/2026\/11\/10\/open-source\/tlx-a-hardware-native-gpu-compiler-for-production-ml\/\">Read More&#8230;<\/a><\/p>\n","protected":false},"author":51,"featured_media":24986,"comment_status":"closed","ping_status":"closed","sticky":false,"template":"","format":"standard","meta":{"_jetpack_newsletter_access":"","_jetpack_dont_email_post_to_subs":false,"_jetpack_newsletter_tier_id":0,"_jetpack_memberships_contains_paywalled_content":false,"_jetpack_hide_featured_image":false,"_jetpack_featured_image_duplicate_dismissed":0,"_jetpack_feature_clip_id":0,"_jetpack_memberships_contains_paid_content":false,"footnotes":"","jetpack_post_was_ever_published":false},"categories":[66,72,174,67],"tags":[],"coauthors":[3073,3074,3075,3076,3077,1944],"class_list":["post-24984","post","type-post","status-publish","format-standard","has-post-thumbnail","hentry","category-developer-tools","category-ml-applications","category-open-source","category-production-engineering","fb_content_type-article"],"yoast_head":"<!-- This site is optimized with the Yoast SEO Premium plugin v19.3 (Yoast SEO v27.3) - https:\/\/yoast.com\/product\/yoast-seo-premium-wordpress\/ -->\n<title>How We Built TLX: A Hardware-Native GPU Compiler for Production ML at Scale - Engineering at Meta<\/title>\n<meta name=\"robots\" content=\"index, follow, max-snippet:-1, max-image-preview:large, max-video-preview:-1\" \/>\n<link rel=\"canonical\" href=\"https:\/\/engineering.fb.com\/2026\/11\/10\/open-source\/tlx-a-hardware-native-gpu-compiler-for-production-ml\/\" \/>\n<meta name=\"twitter:label1\" content=\"Written by\" \/>\n\t<meta name=\"twitter:data1\" content=\"Hongtao Yu, Peng Chen, Daohang Shi, Nick Riasanovsky, Lei Wang, Partha Kanuparthy\" \/>\n\t<meta name=\"twitter:label2\" content=\"Est. reading time\" \/>\n\t<meta name=\"twitter:data2\" content=\"19 minutes\" \/>\n<script type=\"application\/ld+json\" class=\"yoast-schema-graph\">{\"@context\":\"https:\\\/\\\/schema.org\",\"@graph\":[{\"@type\":\"Article\",\"@id\":\"https:\\\/\\\/engineering.fb.com\\\/2026\\\/11\\\/10\\\/open-source\\\/tlx-a-hardware-native-gpu-compiler-for-production-ml\\\/#article\",\"isPartOf\":{\"@id\":\"https:\\\/\\\/engineering.fb.com\\\/2026\\\/11\\\/10\\\/open-source\\\/tlx-a-hardware-native-gpu-compiler-for-production-ml\\\/\"},\"author\":{\"@id\":\"https:\\\/\\\/engineering.fb.com\\\/2026\\\/11\\\/10\\\/open-source\\\/tlx-a-hardware-native-gpu-compiler-for-production-ml\\\/#author\",\"name\":\"\"},\"headline\":\"How We Built TLX: A Hardware-Native GPU Compiler for Production ML at Scale\",\"datePublished\":\"2026-11-10T16:59:31+00:00\",\"mainEntityOfPage\":{\"@id\":\"https:\\\/\\\/engineering.fb.com\\\/2026\\\/11\\\/10\\\/open-source\\\/tlx-a-hardware-native-gpu-compiler-for-production-ml\\\/\"},\"wordCount\":3591,\"publisher\":{\"@id\":\"https:\\\/\\\/engineering.fb.com\\\/#organization\"},\"image\":{\"@id\":\"https:\\\/\\\/engineering.fb.com\\\/2026\\\/11\\\/10\\\/open-source\\\/tlx-a-hardware-native-gpu-compiler-for-production-ml\\\/#primaryimage\"},\"thumbnailUrl\":\"https:\\\/\\\/engineering.fb.com\\\/wp-content\\\/uploads\\\/2026\\\/11\\\/Meta-TLX-hero.png\",\"articleSection\":[\"DevInfra\",\"ML Applications\",\"Open Source\",\"Production Engineering\"],\"inLanguage\":\"en-US\"},{\"@type\":\"WebPage\",\"@id\":\"https:\\\/\\\/engineering.fb.com\\\/2026\\\/11\\\/10\\\/open-source\\\/tlx-a-hardware-native-gpu-compiler-for-production-ml\\\/\",\"url\":\"https:\\\/\\\/engineering.fb.com\\\/2026\\\/11\\\/10\\\/open-source\\\/tlx-a-hardware-native-gpu-compiler-for-production-ml\\\/\",\"name\":\"How We Built TLX: A Hardware-Native GPU Compiler for Production ML at Scale - Engineering at Meta\",\"isPartOf\":{\"@id\":\"https:\\\/\\\/engineering.fb.com\\\/#website\"},\"primaryImageOfPage\":{\"@id\":\"https:\\\/\\\/engineering.fb.com\\\/2026\\\/11\\\/10\\\/open-source\\\/tlx-a-hardware-native-gpu-compiler-for-production-ml\\\/#primaryimage\"},\"image\":{\"@id\":\"https:\\\/\\\/engineering.fb.com\\\/2026\\\/11\\\/10\\\/open-source\\\/tlx-a-hardware-native-gpu-compiler-for-production-ml\\\/#primaryimage\"},\"thumbnailUrl\":\"https:\\\/\\\/engineering.fb.com\\\/wp-content\\\/uploads\\\/2026\\\/11\\\/Meta-TLX-hero.png\",\"datePublished\":\"2026-11-10T16:59:31+00:00\",\"breadcrumb\":{\"@id\":\"https:\\\/\\\/engineering.fb.com\\\/2026\\\/11\\\/10\\\/open-source\\\/tlx-a-hardware-native-gpu-compiler-for-production-ml\\\/#breadcrumb\"},\"inLanguage\":\"en-US\",\"potentialAction\":[{\"@type\":\"ReadAction\",\"target\":[\"https:\\\/\\\/engineering.fb.com\\\/2026\\\/11\\\/10\\\/open-source\\\/tlx-a-hardware-native-gpu-compiler-for-production-ml\\\/\"]}]},{\"@type\":\"ImageObject\",\"inLanguage\":\"en-US\",\"@id\":\"https:\\\/\\\/engineering.fb.com\\\/2026\\\/11\\\/10\\\/open-source\\\/tlx-a-hardware-native-gpu-compiler-for-production-ml\\\/#primaryimage\",\"url\":\"https:\\\/\\\/engineering.fb.com\\\/wp-content\\\/uploads\\\/2026\\\/11\\\/Meta-TLX-hero.png\",\"contentUrl\":\"https:\\\/\\\/engineering.fb.com\\\/wp-content\\\/uploads\\\/2026\\\/11\\\/Meta-TLX-hero.png\",\"width\":1376,\"height\":768},{\"@type\":\"BreadcrumbList\",\"@id\":\"https:\\\/\\\/engineering.fb.com\\\/2026\\\/11\\\/10\\\/open-source\\\/tlx-a-hardware-native-gpu-compiler-for-production-ml\\\/#breadcrumb\",\"itemListElement\":[{\"@type\":\"ListItem\",\"position\":1,\"name\":\"Home\",\"item\":\"https:\\\/\\\/engineering.fb.com\\\/\"},{\"@type\":\"ListItem\",\"position\":2,\"name\":\"How We Built TLX: A Hardware-Native GPU Compiler for Production ML at Scale\"}]},{\"@type\":\"WebSite\",\"@id\":\"https:\\\/\\\/engineering.fb.com\\\/#website\",\"url\":\"https:\\\/\\\/engineering.fb.com\\\/\",\"name\":\"Engineering at Meta\",\"description\":\"Engineering at Meta Blog\",\"publisher\":{\"@id\":\"https:\\\/\\\/engineering.fb.com\\\/#organization\"},\"potentialAction\":[{\"@type\":\"SearchAction\",\"target\":{\"@type\":\"EntryPoint\",\"urlTemplate\":\"https:\\\/\\\/engineering.fb.com\\\/?s={search_term_string}\"},\"query-input\":{\"@type\":\"PropertyValueSpecification\",\"valueRequired\":true,\"valueName\":\"search_term_string\"}}],\"inLanguage\":\"en-US\"},{\"@type\":\"Organization\",\"@id\":\"https:\\\/\\\/engineering.fb.com\\\/#organization\",\"name\":\"Meta\",\"url\":\"https:\\\/\\\/engineering.fb.com\\\/\",\"logo\":{\"@type\":\"ImageObject\",\"inLanguage\":\"en-US\",\"@id\":\"https:\\\/\\\/engineering.fb.com\\\/#\\\/schema\\\/logo\\\/image\\\/\",\"url\":\"https:\\\/\\\/engineering.fb.com\\\/wp-content\\\/uploads\\\/2023\\\/08\\\/Meta_lockup_positive-primary_RGB.jpg\",\"contentUrl\":\"https:\\\/\\\/engineering.fb.com\\\/wp-content\\\/uploads\\\/2023\\\/08\\\/Meta_lockup_positive-primary_RGB.jpg\",\"width\":29011,\"height\":12501,\"caption\":\"Meta\"},\"image\":{\"@id\":\"https:\\\/\\\/engineering.fb.com\\\/#\\\/schema\\\/logo\\\/image\\\/\"},\"sameAs\":[\"https:\\\/\\\/www.facebook.com\\\/Engineering\\\/\",\"https:\\\/\\\/x.com\\\/fb_engineering\"]},[]]}<\/script>\n<!-- \/ Yoast SEO Premium plugin. -->","yoast_head_json":{"title":"How We Built TLX: A Hardware-Native GPU Compiler for Production ML at Scale - Engineering at Meta","robots":{"index":"index","follow":"follow","max-snippet":"max-snippet:-1","max-image-preview":"max-image-preview:large","max-video-preview":"max-video-preview:-1"},"canonical":"https:\/\/engineering.fb.com\/2026\/11\/10\/open-source\/tlx-a-hardware-native-gpu-compiler-for-production-ml\/","twitter_misc":{"Written by":"Hongtao Yu, Peng Chen, Daohang Shi, Nick Riasanovsky, Lei Wang, Partha Kanuparthy","Est. reading time":"19 minutes"},"schema":{"@context":"https:\/\/schema.org","@graph":[{"@type":"Article","@id":"https:\/\/engineering.fb.com\/2026\/11\/10\/open-source\/tlx-a-hardware-native-gpu-compiler-for-production-ml\/#article","isPartOf":{"@id":"https:\/\/engineering.fb.com\/2026\/11\/10\/open-source\/tlx-a-hardware-native-gpu-compiler-for-production-ml\/"},"author":{"@id":"https:\/\/engineering.fb.com\/2026\/11\/10\/open-source\/tlx-a-hardware-native-gpu-compiler-for-production-ml\/#author","name":""},"headline":"How We Built TLX: A Hardware-Native GPU Compiler for Production ML at Scale","datePublished":"2026-11-10T16:59:31+00:00","mainEntityOfPage":{"@id":"https:\/\/engineering.fb.com\/2026\/11\/10\/open-source\/tlx-a-hardware-native-gpu-compiler-for-production-ml\/"},"wordCount":3591,"publisher":{"@id":"https:\/\/engineering.fb.com\/#organization"},"image":{"@id":"https:\/\/engineering.fb.com\/2026\/11\/10\/open-source\/tlx-a-hardware-native-gpu-compiler-for-production-ml\/#primaryimage"},"thumbnailUrl":"https:\/\/engineering.fb.com\/wp-content\/uploads\/2026\/11\/Meta-TLX-hero.png","articleSection":["DevInfra","ML Applications","Open Source","Production Engineering"],"inLanguage":"en-US"},{"@type":"WebPage","@id":"https:\/\/engineering.fb.com\/2026\/11\/10\/open-source\/tlx-a-hardware-native-gpu-compiler-for-production-ml\/","url":"https:\/\/engineering.fb.com\/2026\/11\/10\/open-source\/tlx-a-hardware-native-gpu-compiler-for-production-ml\/","name":"How We Built TLX: A Hardware-Native GPU Compiler for Production ML at Scale - Engineering at Meta","isPartOf":{"@id":"https:\/\/engineering.fb.com\/#website"},"primaryImageOfPage":{"@id":"https:\/\/engineering.fb.com\/2026\/11\/10\/open-source\/tlx-a-hardware-native-gpu-compiler-for-production-ml\/#primaryimage"},"image":{"@id":"https:\/\/engineering.fb.com\/2026\/11\/10\/open-source\/tlx-a-hardware-native-gpu-compiler-for-production-ml\/#primaryimage"},"thumbnailUrl":"https:\/\/engineering.fb.com\/wp-content\/uploads\/2026\/11\/Meta-TLX-hero.png","datePublished":"2026-11-10T16:59:31+00:00","breadcrumb":{"@id":"https:\/\/engineering.fb.com\/2026\/11\/10\/open-source\/tlx-a-hardware-native-gpu-compiler-for-production-ml\/#breadcrumb"},"inLanguage":"en-US","potentialAction":[{"@type":"ReadAction","target":["https:\/\/engineering.fb.com\/2026\/11\/10\/open-source\/tlx-a-hardware-native-gpu-compiler-for-production-ml\/"]}]},{"@type":"ImageObject","inLanguage":"en-US","@id":"https:\/\/engineering.fb.com\/2026\/11\/10\/open-source\/tlx-a-hardware-native-gpu-compiler-for-production-ml\/#primaryimage","url":"https:\/\/engineering.fb.com\/wp-content\/uploads\/2026\/11\/Meta-TLX-hero.png","contentUrl":"https:\/\/engineering.fb.com\/wp-content\/uploads\/2026\/11\/Meta-TLX-hero.png","width":1376,"height":768},{"@type":"BreadcrumbList","@id":"https:\/\/engineering.fb.com\/2026\/11\/10\/open-source\/tlx-a-hardware-native-gpu-compiler-for-production-ml\/#breadcrumb","itemListElement":[{"@type":"ListItem","position":1,"name":"Home","item":"https:\/\/engineering.fb.com\/"},{"@type":"ListItem","position":2,"name":"How We Built TLX: A Hardware-Native GPU Compiler for Production ML at Scale"}]},{"@type":"WebSite","@id":"https:\/\/engineering.fb.com\/#website","url":"https:\/\/engineering.fb.com\/","name":"Engineering at Meta","description":"Engineering at Meta Blog","publisher":{"@id":"https:\/\/engineering.fb.com\/#organization"},"potentialAction":[{"@type":"SearchAction","target":{"@type":"EntryPoint","urlTemplate":"https:\/\/engineering.fb.com\/?s={search_term_string}"},"query-input":{"@type":"PropertyValueSpecification","valueRequired":true,"valueName":"search_term_string"}}],"inLanguage":"en-US"},{"@type":"Organization","@id":"https:\/\/engineering.fb.com\/#organization","name":"Meta","url":"https:\/\/engineering.fb.com\/","logo":{"@type":"ImageObject","inLanguage":"en-US","@id":"https:\/\/engineering.fb.com\/#\/schema\/logo\/image\/","url":"https:\/\/engineering.fb.com\/wp-content\/uploads\/2023\/08\/Meta_lockup_positive-primary_RGB.jpg","contentUrl":"https:\/\/engineering.fb.com\/wp-content\/uploads\/2023\/08\/Meta_lockup_positive-primary_RGB.jpg","width":29011,"height":12501,"caption":"Meta"},"image":{"@id":"https:\/\/engineering.fb.com\/#\/schema\/logo\/image\/"},"sameAs":["https:\/\/www.facebook.com\/Engineering\/","https:\/\/x.com\/fb_engineering"]},[]]}},"jetpack_shortlink":"https:\/\/wp.me\/pa0Lhq-6uY","jetpack_sharing_enabled":true,"jetpack_featured_media_url":"https:\/\/engineering.fb.com\/wp-content\/uploads\/2026\/11\/Meta-TLX-hero.png","_links":{"self":[{"href":"https:\/\/engineering.fb.com\/wp-json\/wp\/v2\/posts\/24984","targetHints":{"allow":["GET"]}}],"collection":[{"href":"https:\/\/engineering.fb.com\/wp-json\/wp\/v2\/posts"}],"about":[{"href":"https:\/\/engineering.fb.com\/wp-json\/wp\/v2\/types\/post"}],"author":[{"embeddable":true,"href":"https:\/\/engineering.fb.com\/wp-json\/wp\/v2\/users\/51"}],"replies":[{"embeddable":true,"href":"https:\/\/engineering.fb.com\/wp-json\/wp\/v2\/comments?post=24984"}],"version-history":[{"count":6,"href":"https:\/\/engineering.fb.com\/wp-json\/wp\/v2\/posts\/24984\/revisions"}],"predecessor-version":[{"id":24999,"href":"https:\/\/engineering.fb.com\/wp-json\/wp\/v2\/posts\/24984\/revisions\/24999"}],"wp:featuredmedia":[{"embeddable":true,"href":"https:\/\/engineering.fb.com\/wp-json\/wp\/v2\/media\/24986"}],"wp:attachment":[{"href":"https:\/\/engineering.fb.com\/wp-json\/wp\/v2\/media?parent=24984"}],"wp:term":[{"taxonomy":"category","embeddable":true,"href":"https:\/\/engineering.fb.com\/wp-json\/wp\/v2\/categories?post=24984"},{"taxonomy":"post_tag","embeddable":true,"href":"https:\/\/engineering.fb.com\/wp-json\/wp\/v2\/tags?post=24984"},{"taxonomy":"author","embeddable":true,"href":"https:\/\/engineering.fb.com\/wp-json\/wp\/v2\/coauthors?post=24984"}],"curies":[{"name":"wp","href":"https:\/\/api.w.org\/{rel}","templated":true}]}}