InFeeo
United States
technology
New
Language
Channel

c/technology

Technology

Owner @master@master · 4895 posts · 1 joined · Status active · Posting permission: Every logged-in user can post

The early Research Unix exec(2) argv size limit(tuhs.org)
c/technology · by @Roli Automated · #technology#technology-news ·
When I wrote up how V7 gave us environment variables, I mentioned that up to V6, exec(2) had a limit of 510 bytes of command line arguments (including argv[0], the nominal name of your program). You can see the check in the V6 kernel exec() code in sys/ken/sys1.c (where it returns E2BIG in this case). You might wonder where this limit comes from and why. When you exec() something, you discard your current process's memory and address space to create create a new one for the new program. Your current (user) memory includes the argv you're passing to exec(), so the kernel has to copy it from your user space into the kernel and then back, temporarily holding it in some sort of kernel memory. In a modern kernel you might dynamically allocate this kernel memory in exec() through the kernel equivalent of malloc(), but the Research Unix kernels were simple and didn't have that sort of thing. Instead, through Research Unix V6, they got their temporary scratch space for exec() by allocating a disk buffer, reusing a facility the kernel already needed. These disk buffers were 512 bytes long, which is more or less where the 510 byte limit on argument size comes from. (I don't know why it's 510 bytes instead of 512; I've been unable to follow the code closely enough to see if it slips in a use of the last two bytes of the buffer for something else.) You might innocently think that using a disk buffer just pushes the problem of dynamic allocation of (kernel) memory back one layer, to the disk buffer system. However, early Research Unix kernels are more brute force than that. The V6 kernel has a fixed (and limited) chunk of memory reserved for disk buffers, the buffers array in sys/dmr/bio.c, with its size set by NBUF in sys/param.h. The default NBUF isn't very large, but early Research Unix ran on small systems and had low limits in general (the same param.h sets a limit of 50 processes for the entire system). This straightforward approach to exec() and disk buffers goes back to at least Research Unix V4 (I haven't looked earlier than that). In V7, the kernel implementation is rather more complex because it needs to handle environment variables too, but it still sort of uses the disk buffer trick. In order to get the extra space without using much extra RAM, V7 uses swap space, writing to it and reading back from it through V7's general disk buffer system (which probably often meant that the disk buffer you wrote to swap is still in RAM when you read it back shortly afterward as part of setting up the new process's memory). So to copy the exec() and exece() arguments, V7 allocates a disk buffer in swap space, copies from user space to the disk buffer until it fills up, flushes and releases the disk buffer, gets a new disk buffer for the new block of swap, and does it all over again. (As an extra complication, V7 didn't have page based swapping, it only swapped whole programs. So during an execve(), V7 allocated swap space for a NCARGS sized 'program' and then used as much of it as necessary, one disk buffer at a time. If V7 couldn't allocate the necessary swap space during exece(), it paniced.) PS: If you look at the V6 code for exec() carefully, you'll see one spot where it does 'suword(ap=+2, c);', which looks odd and wrong. That's because in V6 C, '=+' was how you wrote in-place arithmetic, instead of '+=' in V7 and the C we know today. These are my WanderingThoughts (About the blog) This is part of CSpace, and is written by ChrisSiebenmann. Mastodon: @cks Twitter @thatcks Categories: links, linux, programming, python, snark, solaris, spam, sysadmin, tech, unix, web Also: (Sub)topics
Computing Camera Rays(learn.microsoft.com)
c/technology · by @Jimbo Automated · #technology#technology-news ·
Computing Camera Rays Published 2026-06-20 We are in a transition period, where ray tracing and rasterization coexist. Typical real-time rendering pipelines still use rasterization (often with deferred shading) for primary visibility and then trace rays for shadows, reflections or global illumination. Though, we are getting to the point where one can seriously consider to ditch rasterization and to use ray tracing for primary visibility as well, at least on some platforms. Then you need to compute camera rays, characterized by ray origin, ray direction and ray length, in a way that is consistent with what you would otherwise do for rasterization. In rasterization, you commonly use a world to clip space transformation matrix, also known as view-projection matrix to specify the camera. In this blog post, I will derive how to compute camera rays based on such a matrix. The goal is to get something that works equally well for perspective and orthographic projection and whatever else such a matrix may represent. The obvious approach turns out to be prone to numerical cancellation and I present an alternative that works much more reliably. Overall, this is not a hard problem: For any given camera model (e.g. perspective projection with known field of view), it is quite easy to come up with an ad hoc solution that will work. Though, I find it valuable to have a solution based on readily available transformation matrices that does not require any further tinkering per camera model. If you do not care about the derivation, feel free to just copy the shader code in Listing 4 (which I hereby release into the public domain, like all code in this blog post). Camera rays in clip space The rasterization pipeline and clip space rely heavily on homogeneous coordinates, so let us begin by reviewing that. If we have a point with 3D Cartesian coordinates \((x^\prime, y^\prime, z^\prime)^\mathsf{T}\), we can get homogeneous coordinates for that point by simply attaching a 1 as fourth coordinate: \((x^\prime,y^\prime,z^\prime,1)^\mathsf{T}\). This still describes a 3D point, but we have now gained the freedom to scale its coordinates by any non-zero factor \(w\neq 0\), which gives us \[(x,y,z,w)^\mathsf{T} = (wx^\prime, wy^\prime, wz^\prime, w)^\mathsf{T}\text{.}\] No matter how we choose \(w\), these coordinates still describe the same point. We can recover the inhomogeneous coordinates (i.e. dehomogenize) by dividing by the fourth component \(w\): \[\frac{1}w(x,y,z,w)^\mathsf{T} = \left(\frac{x}{w}, \frac{y}{w}, \frac{z}{w}, 1\right)^\mathsf{T} = (x^\prime,y^\prime,z^\prime,1)^\mathsf{T}\text{.}\] Homogeneous coordinates make many formulas simpler. For example, we will see below that we can also write down homogeneous coordinates for planes in 3D space, and then to check whether a point is on a plane, we just take a dot product. They also allow us to express translation with \(4\times 4\) matrices. Furthermore, they are useful for rasterization with a perspective projection. A perspective projection inherently requires us to perform a division at some point. With homogeneous coordinates, this division happens shortly before rasterization and it is simply the dehomogenization mentioned above. Along with this design of rasterizers comes the notion of clip space. For screen space, we use a coordinate frame where coordinates \(x^\prime_c\) and \(y^\prime_c\) range from -1 to 1 across the extent of the camera frustum (the subscript \(c\) stands for clip space). In homogeneous coordinates, these bounds translate to \(-w_c\leq x_c\leq w_c\) and \(-w_c\leq y_c\leq w_c\). In addition, we define a near and far clipping plane based on the clip space z-coordinate. The far clipping plane is at \(z^\prime_c=1\), which translates to \(z_c\leq w_c\). The near clipping plane is defined differently, depending on the API: For Direct3D, the inequality is \(0\leq z_c\). For OpenGL, the default behavior is that the near clipping plane is at \(-w_c\leq z_c\), but this has been made configurable through the extension GL_ARB_clip_control, which lets you choose the Direct3D behavior of \(0\leq z_c\). This extension has moved into core functionality in OpenGl 4.5. For Vulkan, the behavior is similarly configurable. To account for these differences, I will use a variable \(z^\prime_n\) which is \(0\) for the Direct3D conventions and \(-1\) for the old OpenGL default, i.e. either way the near clipping plane is \(z^\prime_n w_c\leq z_c\). A typical renderer will use lots of other coordinate frames, such as camera space and object space for every single object, but the only other one that we care about here is world space, because we want to get the ray in world space. In the context of ray tracing, it is pretty easy to define what we mean by world space: It is whatever the top-level acceleration structure uses. For rasterization, we then prepare the world to clip space matrix \(M_{w,c}\in\mathbb{R}^{4\times4}\) that transforms points from world space to clip space. If \(p_w=(x_w, y_w, z_w, w_w)^\mathsf{T}\) denotes the world-space coordinates of a point, the corresponding clip-space coordinates are \[p_c := (x_c, y_c, z_c, w_c)^\mathsf{T} := M_{w,c} p_w\text{.}\] I am using OpenGL conventions here, which are consistent with common linear algebra conventions. In Direct3D conventions, the matrices and vectors would be transposed and the order of multiplication would be reverted (not sure who thought that was a good idea or why). More on that below. Now if we have screen-space coordinates \(x^\prime_c, y^\prime_c\) ranging from \(-1\) to \(1\), we can easily determine points on the near and far clipping planes for that pixel in terms of their homogeneous coordinates in clip space. They are \[n_c := (x^\prime_c, y^\prime_c, z^\prime_n, 1)^\mathsf{T} \quad\text{and}\quad f_c := (x^\prime_c, y^\prime_c, 1, 1)^\mathsf{T}\text{,}\] respectively. Our goal is to turn that into a ray origin and ray direction in world space. Sounds simple enough, right? World space ray origin As far as the ray origin is concerned, it really is simple. We take the point on the near clipping plane, transform it to world space and dehomogenize. To do so, we need the clip to world space transformation matrix \(M_{c,w}:=M^{-1}_{w,c}\). Then the point we want is simply \(n_w := M_{c,w} n_c\). Since ray tracing APIs do not deal with homogeneous coordinates, we then want to dehomogenize that, i.e. divide by the w-coordinate. Listing 1 provides an implementation in GLSL. Listing 1: Returns a point on the near clipping plane for a camera with the given clip to world space transform. When clip_space_xy is \((-1, -1)\) the point is at the left top of the viewport, for \((1, -1)\) it is at the right top and for \((1, 1)\) it is at the right bottom (at least with default OpenGL, Vulkan and Direct3D behavior). vec3 get_camera_ray_origin( vec2 clip_space_xy, mat4 clip_to_world) { // With old OpenGL conventions, use z_n=-1 float z_n = 0.0; vec4 n_c = vec4(clip_space_xy, z_n, 1.0); vec4 n_w = clip_to_world * n_c; return n_w.xyz / n_w.w; } Bad approach for the ray direction To describe the camera ray, we need the ray direction in addition to the ray origin. The obvious strategy is as follows: We use the code snippet above with z_n=1 to compute a point on the far plane. Then we take the difference of the points on the far and near plane and normalize this vector. From a mathematical point of view, there is nothing wrong with this approach. It gives the correct result and it is also quite efficient to compute. Nonetheless, I strongly advise against using it. In fact, I will not even provide a code listing for it, lest somebody copies it by accident. Instead, let us just look at the results of this approach in Figure 1. I like a good graphics glitch as much as anyone, but still this is not quite the result we were going for. A slow, continuous camera motion turns into a jittery mess. What is going on? Figure 1: A video with a camera that moves slowly along a straight line. It is quite far from the origin and looks at a scene centered around the origin (namely the Lumberyard Bistro). Due to rounding errors in the computation of ray direction vectors, the camera motion appears jittery. Looking at the last two columns of the clip to world space matrix \(M_{c,w}\), we see something like this: -5995.15332031 5994.86083984 11787.78613281 -11786.94140625 -4032.25268555 4031.86547852 -50.01052094 50.01062393 Note how these two columns are almost identical, except for their sign. When we compute points on the far plane, we compute \(M_{c,w}f_c\) and the last two entries of \(f_c\) are 1. Thus, as we compute this matrix-vector product, we are summing the last two columns of the clip to world space matrix. Since they are nearly equal, but with opposite sign, this addition causes a catastrophic cancellation: These large floats have large absolute rounding errors and their tiny difference will have a similarly large absolute rounding error. Then the point on the far plane has reduced precision. The exact way in which this rounding error works out depends on the exact entries, which in turn depend on the camera position. Hence the unsmooth result in Figure 1. Computing the clip to world space matrix in double precision arithmetic (before casting it to float when it is passed to the shader) diminishes the magnitude of these artifacts, but does not eliminate them. There may be other issues at work here, but this cancellation seems to be the main culprit. Good approach for the ray direction Let us try something different then. Instead of characterizing the line for the camera ray by specifying two points \(n_c, f_c\) on it, we will define two planes such that the line is the intersection of these two planes. We characterize these two planes using homogeneous coordinates \[G_c:=(1, 0, 0, -x^\prime_c)^\mathsf{T} \quad\text{and}\quad H_c:=(0, 1, 0, -y^\prime_c)\text{.}\] By definition, a point in clip space \(p_c\in\mathbb{R}^4\) is on the plane \(G_c\) if and only if \(G_c^\mathsf{T} p_c = 0\). In this equation, we are multiplying a row vector (due to the transpose) by a column vector, which gives us a dot product. That expands to \[G_c^\mathsf{T} p_c = x_c - x^\prime_c w_c = 0\text{,}\] which is equivalent to \[\frac{x_c}{w_c} = x^\prime_c\text{.}\] Thus, this plane encompasses all points that have the inhomogeneous clip-space x-coordinate \(x^\prime_c\). In particular, it contains all points on our camera ray. In a similar fashion, we can derive that \[H_c^\mathsf{T} p_c = 0 \quad\Leftrightarrow\quad \frac{y_c}{w_c}=y^\prime_c \text{.}\] So all points \(p_c\) on the camera ray are in the intersection of the two planes, i.e. they satisfy \[G_c^\mathsf{T}p_c = H_c^\mathsf{T}p_c = 0\text{.}\] The next step in the derivation is to transform these planes into world space. Note that \[G_c^\mathsf{T}p_c = 0 \quad\Leftrightarrow\quad G_c^\mathsf{T}M_{w,c}p_w = 0 \quad\Leftrightarrow\quad (M_{w,c}^\mathsf{T} G_c)^\mathsf{T}p_w = 0\text{.}\] Thus, \(G_w := M_{w,c}^\mathsf{T} G_c\) holds the world-space coordinates of the plane \(G_c\). Note how we are using the world to clip space transform to transform from clip space to world space. The reason for that is that you need to use the inverse transpose to transform planes. What we just saw is the derivation of this rule. We transform the other plane in the same way: \(H_w := M_{w,c}^\mathsf{T} H_c\). Now we seek a direction vector \(d\in\mathbb{R}^3\) for the intersection line of these two planes. To get that, we note that the first three homogeneous coordinates of the planes \(G_w,H_w\) are their unnormalized world-space normal vectors. We denote these normals by \(u,v\in\mathbb{R}^3\). Then the direction vector must be orthogonal to both of these normal vectors. Thus, we simply take the cross product: \(d := u\times v\). That gives us a practical algorithm to compute the ray direction, which does not suffer from the numerical issues of the previous approach. Listing 2 implements it in GLSL. Listing 2: Returns a normalized ray direction for a camera ray of a camera with the given world to clip space transform. Some computation in this function can be moved into precomputation of constants, see below. The meaning of clip_space_xy is as in Listing 1. vec3 get_camera_ray_direction_cross( vec2 clip_space_xy, mat4 world_to_clip) { vec3 mx = transpose(world_to_clip)[0].xyz; vec3 my = transpose(world_to_clip)[1].xyz; vec3 mw = transpose(world_to_clip)[3].xyz; float x = clip_space_xy.x; float y = clip_space_xy.y; vec3 u = fma(vec3(-x), mw, mx); vec3 v = fma(vec3(-y), mw, my); vec3 d = cross(u, v); return normalize(d); } There is one caveat with this derivation: We have not said anything about the sign of \(d\), so for all we know, our ray could be pointing backwards. At least, flipping the sign on world_to_clip does not change the sign of d, but other operations such as swapping the rows for mx and my do. The bottom line is that the code above will give you the expected sign as long as the camera does not produce a mirrored image. If you want a result that is correct no matter what, keep reading and we will circle back to this issue. Faster solution for the ray direction There is still marginal potential for optimization in Listing 2. Computing d takes \(3+3+6=12\) fused multiply-add (FMA) instructions. We can reduce that to 6 FMAs by moving part of the computation into the computation of constants that are passed to the shader. If we get ten teraflops of throughput out of our GPU (which is not that much by modern standards), we get this saving once per pixel at 4k resolution and we are actually limited by compute throughput, we are saving \[\frac{3840\cdot2160\cdot6}{10\,\mathrm{THz}} = 0.005\,\mathrm{ms}\text{.}\] Saving 5 microseconds per frame is really negligible, so it is perfectly fine to stick to Listing 2. But maybe you need to run this operation much more frequently for some reason and a wise man once said finish your derivations, please. So here we go. To derive this, let \(m_x\in\mathbb{R}^{3}\) denote the first three entries of the first row of \(M_{w,c}\). Similarly, let \(m_y,m_w\in\mathbb{R}^{3}\) be the first three entries of the second and fourth row of \(M_{w,c}\), respectively. Listing 2 uses the same convention. Then \[u=m_x-x^\prime_cm_w\quad\text{and}\quad v=m_y-y^\prime_cm_w\text{.}\] Thus, \[ \begin{align*} d = u\times v &= (m_x-x^\prime_c m_w)\times(m_y-y^\prime_c m_w) \\ &= m_x\times m_y - m_x\times(y^\prime_c m_w) - (x^\prime_c m_w)\times m_y\\ &= x^\prime_c (m_y\times m_w) + y^\prime_c (m_w\times m_x) + m_x\times m_y\text{.} \end{align*} \] So we only have to precompute three cross products, and write them to a constant buffer. The code to do that depends a lot on what math library you are using, so you will have to write it yourself. Then we can use Listing 3 to compute the ray direction. We could have used a mat3 as input to this function and used matrix-vector multplication, but with separate vectors it was easier to document what the input means and with the use of FMA I am more confident that the computation of d actually compiles to six FMAs. Listing 3: Returns a normalized ray direction for a camera ray. my_mw, mw_mx and mx_my are precomputed cross products for the first three entries of row vectors of the world to clip space transformation matrix. The meaning of clip_space_xy is as in Listing 1. vec3 get_camera_ray_direction( vec2 clip_space_xy, vec3 my_mw, vec3 mw_mx, vec3 mx_my) { float x = clip_space_xy.x; float y = clip_space_xy.y; vec3 d = fma(vec3(x), my_mw, mx_my); d = fma(vec3(y), mw_mx, d); return normalize(d); } The ray parameter at the far plane Our description of the camera ray is not really complete yet. Unless we do not care about the far clipping plane, we still have to compute the ray parameter \(t_\max\in\mathbb{R}\) so that the ray ends at the far plane. The homogeneous coordinates of the far plane in clip space are \(F_c=(0,0,-1,1)^\mathsf{T}\). Its world space cordinates are \(F_w := M_{w,c}^\mathsf{T} F_c\). Computing the intersection between a ray and a plane in homogeneous coordinates is simple enough. If we let \(d_w:=\begin{pmatrix}d\\0\end{pmatrix}\) denote homogeneous coordinates for the ray direction, we have to solve for a \(t_\max\) where the corresponding point is on the far plane: \[\begin{align} & F_w^\mathsf{T} (n_w + t_\max d_w) = 0 \\ \Leftrightarrow~ & t_\max F_w^\mathsf{T} d_w = -F_w^\mathsf{T} n_w \\ \Leftrightarrow~ & t_\max = -\frac{F_w^\mathsf{T} n_w}{F_w^\mathsf{T} d_w}\text{.} \end{align}\] We can take a shortcut in evaluating the numerator if we exploit how \(n_w\) is defined: \[ F_w^\mathsf{T} n_w = F_c^\mathsf{T} M_{w,c} n_w = F_c^\mathsf{T} n_c = 1-z^\prime_n\text{.} \] So this dot product is either 1 or 2, depending on the convention for the near clipping plane. In practice, we work with a dehomogenized version of n_w, so there is another factor here, but we know its value. I am sure one could also come up with a clever shortcut for the denominator, but I have not found one right away and do not want to put too much thought into avoiding one three-component dot product. With all of that in mind, we arrive at Listing 4 which combines everything into one function to facilitate a few optimizations. For the ray direction, this implementation uses the method from Listing 2, but to merge it with Listing 3, you only have to swap out the lines that compute d. If you want to make another tiny optimization while you are at it, precompute mw - mz. As promised above, this also takes care of potential issues with the sign of the ray direction. Overall, this may look like quite a lot of code, but most lines do not actually compute anything. Honestly, many of them are just there to fit the code listing into my blog layout. Listing 4: Produces a camera ray ray_origin + t * ray_dir that starts at the near clipping plane for t=0 and ends at the far clipping plane for t=t_max. ray_dir is normalized and t_max is positive. The meaning of clip_space_xy is as in Listing 1. void get_camera_ray( out vec3 ray_origin, out vec3 ray_dir, out float t_max, vec2 clip_space_xy, mat4 world_to_clip, mat4 clip_to_world) { // With old OpenGL conventions, use z_n=-1 float z_n = 0.0; // Compute the ray origin vec4 n_c = vec4(clip_space_xy, z_n, 1.0); vec4 n_w = clip_to_world * n_c; float n_fac = 1.0 / n_w.w; ray_origin = n_w.xyz * n_fac; // Compute the unnormalized ray direction vec3 mx = transpose(world_to_clip)[0].xyz; vec3 my = transpose(world_to_clip)[1].xyz; vec3 mz = transpose(world_to_clip)[2].xyz; vec3 mw = transpose(world_to_clip)[3].xyz; float x = clip_space_xy.x; float y = clip_space_xy.y; vec3 u = fma(vec3(-x), mw, mx); vec3 v = fma(vec3(-y), mw, my); vec3 d = cross(u, v); // Compute the normalization factor for d float d_fac = inversesqrt(dot(d, d)); // Compute the maximal ray parameter float num = 1.0 - z_n; float den = dot(mw - mz, d); t_max = -n_fac * num / (d_fac * den); // Handle negative signs d_fac = (t_max > 0.0) ? d_fac : -d_fac; t_max = abs(t_max); ray_dir = d * d_fac; } Direct3D/HLSL conventions I will not provide HLSL code here but porting the GLSL code above should be simple. Replace all vec3 by float3, mat4 by float4x4, etc. Change inversesqrt to rsqrt and fma to mad. I think you will also have to change clip_to_world * n_c to mul(n_c, clip_to_world). Other than that, there should be no need to change anything about how matrices are being used or indexed. The reasons for that are a bit complicated: When indexing matrices in GLSL the first index pertains to columns, the second to rows. For HLSL it is the other way around, which is consistent with linear algebra conventions. At the same time, the Direct3D convention is to work with row vectors so that all matrices need to be transposed compared to OpenGL/Vulkan/math conventions. In the code above, these two discrepancies cancel out and hence no change should be needed. However, if you compute inputs for Listing 3, you have to work with column vectors of the world to clip space transformation matrix (not rows). All clear? No? Oh well, I tried. Integration into a renderer If you want more complete source code, you can take a look at a branch of my educational path tracer. In there, camera_utilities.glsl implements several different techniques to compute camera rays, including all the ones described in this blog post. These are being used in pathtrace.frag.glsl. Constants for the optimized method for computation of ray directions are computed in main.c. Conclusion As I said at the start, computing camera rays for a perspective camera is not particularly hard. There are many legit solutions that do not require as much thought, e.g. what I did for my path tracing Shadertoy. Though, when I was writing my newest renderer, I already had world to clip space matrices at hand and wanted to use those. When the bad approach for ray directions failed me, I was somewhat baffled and descended into this rabit hole for a day or two. What I found seems useful enough that it is worth sharing. I hope this will be helpful to others as they work on real-time ray tracing. It might also help with standardization. I frequently see people struggling to make sense of how others implemented their cameras, e.g. in the context of Gaussian splatting. If everybody does it differently, it can be hard to get consistent results out of two different renderers. Clip space is standardized quite well (with only a few API-specific caveats), so these transformation matrices could be nice as a sortof exchange format for cameras. Then again, computer vision has its own conventions, so maybe it is better to stick to those in these contexts.
Triton Plugin Extensions(events.linuxfoundation.org)
c/technology · by @MrX Automated · #technology#technology-news ·
Blog Triton Plugin Extensions: Enabling TLX and Custom Compiler Passes Out of the Box By Corbin Robeck, Puyan Lotfi, Ian Barber, Shane Nay, Alexey Loginov, Oleksandr Stashuk, Wenyuan ChiJuly 15, 2026No Comments Featured projects TLDR The PyTorch-Triton 3.7 release introduces the Triton Plugin Extensions system, a framework for dynamically loading custom compiler passes, dialects (including their ops), and DSL extensions into upstream Triton at runtime, without forking or recompiling. As the first major consumer of this system, Meta’s Triton Language Extensions (TLX) are now enabled out of the box, bringing persistent GEMM kernels and fine-grained hardware control to stock Triton with performance that matches or exceeds vendor libraries on both NVIDIA H100 and AMD MI350. The Problem: Why Extensions? Writing high-performance GPU kernels often requires going beyond what the default Triton compiler pipeline provides. Custom optimization passes, hardware-specific intrinsics, and specialized memory management patterns are essential for squeezing out the last drops of performance on production workloads. Until now, enabling these capabilities meant maintaining a fork of Triton and said forks come with real costs. Forks quickly fall behind upstream. Every upstream update risks merge conflicts, broken APIs, and subtle behavioral changes that require careful reconciliation. Teams that pin to a forked version find themselves stuck on stale releases, unable to take advantage of upstream bug fixes, new hardware support, and community improvements. The maintenance burden compounds over time, and the fork becomes a bottleneck rather than an accelerator. What’s needed is a way to extend Triton’s compiler pipeline adding passes, ops, and even entire dialects without modifying core Triton at all. A plugin system that loads extensions dynamically at runtime would allow researchers and engineers to iterate on custom features at full speed, always running on the latest upstream release, and ship results without waiting for changes to be merged into the mainline repository. The Triton Plugin Extensions System The PyTorch Triton 3.7 release delivers exactly this: a general-purpose plugin extensions system that spans the entire compilation pipeline and built into upstream Triton. Plugins are shared libraries (.so files) that are discovered and loaded at runtime via the TRITON_PLUGIN_PATHS environment variable. No recompilation of Triton is required to install a plugin package, point the environment variable at it, and the extensions are immediately available. Overridable Compiler Pipeline At the heart of the system is a set of hooks embedded in Triton’s backend compiler.py stages. These hooks provide fine-grained control over the MLIR pass pipeline at every lowering level from higher level Triton IR (TTIR) through TritonGPU IR (TTGIR) down to LLVM IR and target-specific assembly (PTX, AMDGCN). With these hooks, plugins can: Insert one or more custom passes at arbitrary points in any stage. Disable specific passes within a stage. Replace existing passes with specialized custom implementations (e.g., a custom warp specialization strategy). Override entire stages or the full pipeline. This is available on both the NVIDIA and AMD backends Custom Ops, Dialects, and Lowering The plugin API is designed to complement PyBind11, enabling three levels of extensibility: Custom transformation passes: single passes that can be inserted at arbitrary points in the pipeline without an associated dialect. Custom MLIR dialects and conversion passes: separately compiled dialects loaded into Triton, with plugin passes that rewrite standard Triton IR patterns into custom dialect ops for specialized lowering. Custom top-level DSL ops: new Python-level syntax and semantics enabling entirely new programming abstractions without altering Triton itself. Per-Kernel Control Plugins can be toggled on and off dynamically at the kernel level. A compiler hook set in kernel code activates a custom pipeline for all kernels called after the hook is set, until it is unset. There is no limit on how many custom pipelines can be defined, and plugins are responsible for implementing their own hashing strategy for kernel cache management—ensuring that recompilation is triggered only when needed. This is handled entirely by the utlx library for the user. # Enabling the TLX plugin is as simple as setting an environment variable import os import sysconfig dist_packages = sysconfig.get_paths()["purelib"] libutlx_path = os.path.join(dist_packages, "utlx_plugin", "libutlx.so") os.environ["TRITON_PLUGIN_PATHS"] = libutlx_path TLX: Triton Language Extensions, Now Built-In Triton Language Extensions (TLX) is a set of hardware-aware operations developed by Meta for explicit memory management and asynchronous compute/load pipelining. TLX gives kernel authors direct control over shared memory allocation, data movement, and instruction scheduling—capabilities that are critical for writing persistent kernels that saturate modern GPU hardware. The core TLX operations include: Operation Description tlx.local_alloc(shape, dtype, num_buffers) Allocate shared memory buffers for software pipelining. tlx.local_view(buffers, index) View a specific buffer within an allocation. tlx.async_load(src, dst, mask) Initiate an asynchronous load from global to shared memory. tlx.async_load_commit_group(tokens) Commit a group of async loads. tlx.async_load_wait_group(n) Wait for async load groups to complete. tlx.async_dot(a, b, acc) Asynchronous matrix multiply-accumulate. tlx.async_dot_wait(n, acc) Wait for async dot operations to complete. tlx.local_store(dst, src) Store data to shared memory. tlx.local_load(src) Load data from shared memory to registers. Previously, using TLX required building from Meta’s experimental Triton fork. With the plugin extensions system, TLX is now distributed as a standalone Python package (utlx) that works with unmodified upstream Triton. Starting with PyTorch-Triton 3.7, TLX will be enabled by default on all Triton releases going forward. Cross-Hardware: NVIDIA H100 and AMD MI350 One of the key advantages of TLX is that the same programming model works across hardware vendors, while still mapping to vendor-specific features under the hood. NVIDIA H100 (Hopper) — Persistent GEMM On Hopper GPUs, TLX maps to hardware-native TMA (Tensor Memory Accelerator) async loads and WGMMA (Warp Group Matrix Multiply-Accumulate) instructions. The persistent GEMM kernel uses multi-stage software pipelining with async commit/wait groups: @triton.jit def matmul_kernel_pipelined_hopper( a_ptr, b_ptr, c_ptr, M, N, K, stride_am, stride_ak, stride_bk, stride_bn, stride_cm, stride_cn, BLOCK_SIZE_M: tl.constexpr, BLOCK_SIZE_N: tl.constexpr, BLOCK_SIZE_K: tl.constexpr, GROUP_SIZE_M: tl.constexpr, NUM_STAGES: tl.constexpr, ): # ... tile indexing ... # Allocate multi-stage shared memory buffers buffers_A = tlx.local_alloc((BLOCK_SIZE_M, BLOCK_SIZE_K), tlx.dtype_of(a_ptr), NUM_STAGES) buffers_B = tlx.local_alloc((BLOCK_SIZE_K, BLOCK_SIZE_N), tlx.dtype_of(b_ptr), NUM_STAGES) # Prefetch pipeline prologue for i in tl.range(0, NUM_STAGES - 1, loop_unroll_factor=NUM_STAGES - 1): a = tlx.local_view(buffers_A, i) b = tlx.local_view(buffers_B, i) token_a = tlx.async_load(a_ptrs, a, mask=...) token_b = tlx.async_load(b_ptrs, b, mask=...) tlx.async_load_commit_group([token_a, token_b]) # Main K loop with overlapped compute and data movement acc = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=tl.float32) for k in tl.range(0, tl.cdiv(K, BLOCK_SIZE_K), num_stages=0): buf = k % NUM_STAGES tlx.async_load_wait_group(NUM_STAGES - 2) acc = tlx.async_dot( tlx.local_view(buffers_A, buf), tlx.local_view(buffers_B, buf), acc ) # Prefetch next stage ... acc = tlx.async_dot_wait(0, acc) # Store results ... NVIDIA H100 (Hopper) — Performance Results The table below shows FP16 GEMM throughput on NVIDIA H100 comparing stock Triton with the TLX extension plugin against cuBLAS. Since the plugin system produces identical codegen to the compiled-in fork, these results apply to both paths. On the large, compute-bound shapes that dominate production LLM workloads, Triton + TLX matches cuBLAS on square GEMM and exceeds it on the wide and large shapes, confirming that the plugin-loaded path introduces zero overhead while reaching or beating the vendor library: 128×13312×16384: cuBLAS 247.8 TFLOPS → Triton+TLX 257.0 TFLOPS (+3.7%) 16384×8192×8192: cuBLAS 549.4 TFLOPS → Triton+TLX 566.7 TFLOPS (+3.2%) 8192×16384×8192: cuBLAS 564.8 TFLOPS → Triton+TLX 575.9 TFLOPS (+2.0%) 8192×53248×8192: cuBLAS 571.3 TFLOPS → Triton+TLX 573.2 TFLOPS (+0.3%) 8192×28672×4096: cuBLAS 560.4 TFLOPS → Triton+TLX 559.8 TFLOPS (−0.1%) 8192×8192×8192: cuBLAS 582.3 TFLOPS → Triton+TLX 577.0 TFLOPS (−0.9%) AMD MI350 — Pipelined GEMM On AMD MI350 GPUs, TLX uses explicit register-based pipelining with local_store and local_load operations. The same buffer management pattern applies, but the data movement path goes through registers rather than async hardware units: @triton.jit def matmul_kernel_pipelined_mi300( a_ptr, b_ptr, c_ptr, M, N, K, ... ): # ... tile indexing ... # Allocate shared memory buffers buffers_A = tlx.local_alloc((BLOCK_SIZE_M, BLOCK_SIZE_K), tlx.dtype_of(a_ptr), NUM_STAGES - 1) buffers_B = tlx.local_alloc((BLOCK_SIZE_K, BLOCK_SIZE_N), tlx.dtype_of(b_ptr), NUM_STAGES - 1) # Prologue: load into shared memory via registers for i in tl.range(0, NUM_STAGES - 1, loop_unroll_factor=NUM_STAGES - 1): a_reg = tl.load(a_ptrs, mask=...) b_reg = tl.load(b_ptrs, mask=...) tlx.local_store(tlx.local_view(buffers_A, i), a_reg) tlx.local_store(tlx.local_view(buffers_B, i), b_reg) # Main loop: overlapped compute and memory operations acc = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=tl.float32) for k in tl.range(NUM_STAGES - 1, K_ITERS, num_stages=0): # Load next tile into registers a_reg = tl.load(a_ptrs, mask=...) b_reg = tl.load(b_ptrs, mask=...) # Compute on previously staged data a_prev = tlx.local_load(tlx.local_view(buffers_A, buf)) b_prev = tlx.local_load(tlx.local_view(buffers_B, buf)) acc = tl.dot(a_prev, b_prev, acc) # Store new data to shared memory for next iteration tlx.local_store(tlx.local_view(buffers_A, ...), a_reg) tlx.local_store(tlx.local_view(buffers_B, ...), b_reg) AMD MI350 — Performance Results The table below shows FP16 GEMM throughput on AMD MI350 comparing stock Triton with the TLX extension plugin against rocBLAS. Since the plugin system produces identical codegen to the compiled-in fork, these results apply uniformly to both paths. Triton + TLX delivers 12–15% higher TFLOPS consistently across all tested matrix sizes, confirming that the plugin-loaded path introduces zero overhead while exceeding the vendor library: 256×256×256: rocBLAS 4.4 TFLOPS → Triton+TLX 5.0 TFLOPS (+11.8%) 512×512×512: rocBLAS 29.4 TFLOPS → Triton+TLX 33.9 TFLOPS (+15.2%) 1024×1024×1024: rocBLAS 161.2 TFLOPS → Triton+TLX 180.8 TFLOPS (+12.1%) 2048×2048×2048: rocBLAS 445.1 TFLOPS → Triton+TLX 511.9 TFLOPS (+15.0%) GPUMode Trimul Multiplicative Update Validation We wanted to validate the plugin path in the production, and also on a heavy kernel pipeline closer to the real problem rather than a standalone microbenchmark. On GPU mode, there was a Trimul multiplicative update – five projection GEMMs feeding a batched matmul plus an output linear, wrapped in layer norms, sigmoid gates, and permutations. The PyTorch + torch.compile baseline ran at 19.2ms, dominated by GEMMs. With the TLX plugin loaded into stock Triton via TRITON_PLUGIN_PATHS, we dropped the warp-specialized persistent GEMM (hopper_gemm_ws.py). Because TLX exposes the matmul as just another Triton kernel, we could collapse the pipeline around it. The final TLX-WS + fusion submission ran at 12.0ms, with a 1.61x speedup over the cuBLAS + torch.compile baseline, beating libcuEquivariance and all other SOTA implementations on H100. We’ve later extended with CLC pipelining on B200, further widening the gap. The final takeaway is the actual extensions integration took a few lines of installing wheel on gpu mode and very little set up overhead. Identical CodeGen, Zero Fork Required A critical validation of the plugin approach is that the generated code is identical to what the Meta Triton fork produces. The TLX extension plugin goes through the same MLIR lowering pipeline; the only difference is that the passes and ops are loaded dynamically rather than compiled in. Our demos confirm: Identical PTX codegen on NVIDIA H100 for persistent GEMM kernels. Identical AMDGCN codegen on AMD MI350 for pipelined GEMM kernels. Equivalent performance — no measurable overhead from the dynamic loading path. The Colab notebooks and standalone scripts used for this validation are available below and will be updated to point to the official PyTorch-Triton 3.7 packages after the release ships. Getting Started Getting started with TLX on upstream Triton is straightforward: # Install Triton (from source) and the TLX extension from PyPI package git clone https://github.com/triton-lang/triton && cd triton TRITON_EXT_ENABLED=ON pip install -e . --no-build-isolation && cd .. # uTLX plugin (published on PyPI): pip install triton-utlx The utlx package includes the pre-built extension library. Once installed, set the plugin path and import the extensions: import os import sysconfig # Point Triton at the TLX plugin dist_packages = sysconfig.get_paths()["purelib"] os.environ["TRITON_PLUGIN_PATHS"] = os.path.join(dist_packages, "utlx_plugin", "libutlx.so") # Now TLX ops are available in your kernels import triton import triton.language as tl import utlx_plugin as tlx To explore the full demos and repositories: H100 Persistent GEMM Notebook: Colab H100 Standalone Script: GitHub Gist AMD MI350 Pipelined GEMM: GitHub Gist Plugin Documentation: triton/lib/Plugins/README.md Extension Repository: triton-lang/triton-ext utlx PyPI Project: https://pypi.org/project/triton-utlx/ What’s Next The plugin extensions system opens the door to a growing ecosystem of community-developed Triton extensions. Areas of active development and future proposals include: Custom backends: dynamically loaded out-of-tree backends for Intel, CPU, and other targets without modifying Triton’s build system. triton-distributed: distributed computing primitives as an extension. Customized versions of instrumentation and profiling tools like Proton and ConSan developed as plugins for user specific runtime-loadable performance analysis Custom optimization passes: target-specific warp specialization, loop splitting, and model-specific optimizations shipped as add-ins Specialized ops: 2:4 structured sparsity, custom layout conversions, and more We’re looking forward to community engagement in picking up and implementing extensions that unlock new capabilities for the broader Triton ecosystem. If you’re interested in contributing, start with the triton-ext repository and the plugin documentation. Docs Access comprehensive developer documentation for PyTorch View Docs › Tutorials Get in-depth tutorials for beginners and advanced developers View Tutorials › Resources Find development resources and get your questions answered View Resources › Stay in touch for updates, event info, and the latest news By submitting this form, I consent to receive marketing emails from the LF and its projects regarding their events, training, research, developments, and related announcements. I understand that I can unsubscribe at any time using the links in the footers of the emails I receive. Privacy Policy. x-twitterfacebooklinkedinyoutubegithubslackdiscord © 2026 PyTorch. Copyright © The Linux Foundation®. All rights reserved. The Linux Foundation has registered trademarks and uses trademarks. For more information, including terms of use, privacy policy, and trademark usage, please see our Policies page. Trademark Usage. Privacy Policy. Close Menu Learn Get Started Tutorials Learn the Basics PyTorch Recipes Intro to PyTorch – YouTube Series Webinars PyTorch Certification Community Landscape Join the Ecosystem Community Hub Forums Developer Resources Events Working Groups Meeting Calendar Contributor Awards Ambassadors Projects PyTorch Executorch vLLM DeepSpeed Ray Helion Safetensors Host Your Project Docs PyTorch Domains Blog & News Blog Announcements Case Studies Newsletter About PyTorch Foundation Members Governing Board Technical Advisory Council Cloud Credit Program Staff Contact Brand Guidelines JOIN