How GPUs Work: Graphics Pipeline, Shader Cores, Ray Tracing and Compute

A GPU is a throughput processor. It takes a scene described as triangles, runs a small program called a shader on every vertex and every covered pixel, tests each result against a depth buffer and blends it into an image, many millions of times per frame. The same arrays of shader cores that do this also run general purpose compute kernels, which is why GPUs now train and serve AI models.

This guide follows one draw call from the application all the way to the pixels on screen, then opens up the hardware behind each step: unified shader cores, texture units, raster engines, ROPs and the ray tracing units that walk a bounding volume hierarchy. It finishes with the shift from graphics to general purpose compute and with worked numbers for fill rate, texel rate and FLOPS, every one taken from a vendor whitepaper and checked in Python. We read the Vulkan specification, the Direct3D Raytracing and Mesh Shader specifications, the OpenGL 4.6 core specification, NVIDIA's Fermi, Turing, Ampere GA102 and Ada Lovelace whitepapers, the CUDA Programming Guide, AMD's RDNA architecture presentation and Apple's Metal documentation on tile based deferred rendering. Code samples follow the vendor documentation but were not executed: this article was written on a Mac without a discrete GPU, so treat the code as illustration, not as tested software.

Key takeaways

• A frame is built by a pipeline: the application issues draw calls, vertex shaders transform vertices, fixed function hardware assembles and clips triangles, the rasterizer turns each triangle into fragments, fragment shaders compute colors, and the depth test and blending units write the survivors into the framebuffer.
• Most of the silicon is programmable. Since NVIDIA's G80 in November 2006, vertex, geometry, pixel and compute programs all run on one pool of unified shader cores, according to NVIDIA's Fermi whitepaper.
• GPUs hide memory latency by keeping many groups of threads (warps on NVIDIA, waves on AMD) resident and switching between them every cycle, instead of using the large caches and branch predictors a CPU relies on.
• Rasterization is a coverage test: for every pixel sample, three edge functions decide whether the sample lies inside the triangle. The same numbers give the barycentric weights used to interpolate depth and attributes.
• Ray tracing hardware does not replace shaders. On NVIDIA's RT cores and similar units, fixed function logic walks the bounding volume hierarchy and tests ray against box and ray against triangle, then hands hits back to programmable shaders.
• Peak figures are simple products. The GeForce RTX 4090's published 443.5 gigapixels per second is 176 ROPs times a 2.52 GHz boost clock, and its 82.6 FP32 TFLOPS is 16,384 cores times 2 operations times 2.52 GHz. Sustained throughput in a real game or kernel is lower and is not published as a single number.
• The GPU software stack (drivers, shader compilers, CUDA and ROCm toolkits, container runtimes) is a supply chain of its own, with security bulletins and patches. That is where software composition analysis applies. It does not make the GPU faster.

What does a GPU do differently from a CPU?

A CPU is built to finish one stream of instructions as fast as possible. It spends transistors on large caches, deep out of order machinery and branch prediction so that a single thread rarely waits. A GPU is built to finish a huge number of independent, similar jobs per second, and it accepts that each individual job is slow.

The CUDA Programming Guide puts it plainly: a CPU "can execute a few tens" of threads in parallel, while a GPU "is designed to excel at executing thousands of threads in parallel, trading off lower single-thread performance to achieve much greater total throughput." It adds that GPUs devote more transistors to data processing units, while CPUs dedicate more to caching and flow control.

Graphics is the workload that shaped this design. A 3840 by 2160 frame has 8,294,400 pixels. At 60 frames per second that is 497,664,000 pixels per second before any overdraw, and each pixel needs its own lighting and texture work. The pixels are mostly independent of each other, so a design that runs thousands of them at once and tolerates long memory waits wins over a design that runs a few very quickly.

That trade, latency for throughput, runs through every section below.

From draw call to pixels: the whole pipeline at a glance

The Vulkan specification describes the graphics pipeline as a sequence of stages. The input assembler "assembles vertices to form geometric primitives such as points, lines, and triangles." The vertex shader transforms them. Optional tessellation and geometry shaders can create more primitives. Primitives are clipped, the rasterizer "produces a series of fragments associated with a region of the framebuffer," fragment shading computes their values, and framebuffer operations "read and write the color and depth/stencil attachments."

Direct3D, Metal and OpenGL use different names for some stages, but the order is the same. Here is the pipeline with the programmable stages marked.

 Application (CPU)
   | draw call: vkCmdDraw / DrawIndexedInstanced / drawPrimitives
   v
 Driver: validate state, compile shaders to GPU ISA, build command buffer
   |
   v
+----------------------------------------------------------------------+
| GPU front end: command processor fetches and decodes the buffer       |
+----------------------------------------------------------------------+
   |
   v
 [fixed]  Input assembler        read index and vertex buffers
 [PROG]   Vertex shader          object space -> clip space
 [PROG]   Tessellation (opt.)    hull/control, tessellator [fixed], domain/evaluation
 [PROG]   Geometry shader (opt.) emit or discard whole primitives
   |        or, in the mesh path: [PROG] task/amplification -> [PROG] mesh shader
   v
 [fixed]  Primitive assembly, clipping, perspective divide, viewport transform
 [fixed]  Face culling (sign of triangle area)
 [fixed]  Rasterizer             coverage per sample, attribute setup
 [fixed]  Early depth/stencil    (when the shader allows it)
 [PROG]   Fragment/pixel shader  texture fetch, lighting, material math
 [fixed]  Late depth/stencil     (when the shader writes depth or discards)
 [fixed]  Blend / ROP            combine with the framebuffer, MSAA resolve
   |
   v
 Framebuffer in memory -> display engine -> scan out to the monitor

The rest of this section walks through each row.

Step 1: the CPU side, the API and the driver

Before the GPU does anything, the application records work. In Vulkan that means binding a pipeline object (which contains compiled shaders and fixed function state), binding vertex and index buffers and descriptor sets, and recording a draw command into a command buffer. The command buffer is then submitted to a queue.

// Vulkan 1.x: record one indexed draw (illustrative, not executed)
vkCmdBeginRenderPass(cmd, &renderPassInfo, VK_SUBPASS_CONTENTS_INLINE);
vkCmdBindPipeline(cmd, VK_PIPELINE_BIND_POINT_GRAPHICS, pipeline);
VkDeviceSize offset = 0;
vkCmdBindVertexBuffers(cmd, 0, 1, &vertexBuffer, &offset);
vkCmdBindIndexBuffer(cmd, indexBuffer, 0, VK_INDEX_TYPE_UINT32);
vkCmdBindDescriptorSets(cmd, VK_PIPELINE_BIND_POINT_GRAPHICS, layout,
                        0, 1, &descriptorSet, 0, nullptr);
vkCmdDrawIndexed(cmd, indexCount, /*instances*/ 1, 0, 0, 0);
vkCmdEndRenderPass(cmd);
vkEndCommandBuffer(cmd);
// later: vkQueueSubmit(queue, 1, &submitInfo, fence);

Two things happen on the CPU that matter later. First, shaders written in GLSL, HLSL, Metal Shading Language or Slang are compiled to an intermediate form such as SPIR-V or DXIL, and the driver then compiles that intermediate form into the GPU's own instruction set. That second compiler lives inside the driver. Second, the driver translates API state into hardware register writes and packets that the GPU's command processor understands. A modern GPU front end fetches those packets from memory, decodes them, and launches work on the shader arrays.

This is also why drivers are large, frequently updated pieces of software: they contain a compiler, a memory manager and a scheduler for a highly parallel machine. We come back to that in the supply chain section.

Step 2: vertex processing

The input assembler reads indices and fetches vertex attributes from buffers: positions, normals, texture coordinates. Each vertex is handed to a vertex shader invocation. The vertex shader's one mandatory job is to output a position in clip space, a four component homogeneous coordinate (x, y, z, w).

A typical vertex shader multiplies the object space position by model, view and projection matrices:

// GLSL vertex shader (Vulkan flavour), illustrative
#version 450
layout(set = 0, binding = 0) uniform Camera { mat4 view; mat4 proj; } cam;
layout(push_constant) uniform Push { mat4 model; } pc;
layout(location = 0) in vec3 inPos;
layout(location = 1) in vec2 inUV;
layout(location = 0) out vec2 vUV;
void main() {
    vUV = inUV;
    gl_Position = cam.proj * cam.view * pc.model * vec4(inPos, 1.0);
}

Worked example: one vertex through the projection. Take a camera with a 60 degree vertical field of view, a 16:9 aspect ratio, a near plane at 0.1 and a far plane at 100, using the Vulkan depth range of 0 to 1. A point in view space at (1.0, 0.5, -5.0) (five units in front of the camera) becomes, in clip space, approximately (0.974, 0.866, 4.905, 5.0). The hardware divides by w = 5.0 to get normalised device coordinates of about (0.195, 0.173, 0.981). The viewport transform maps x and y from the range of minus one to one onto a 1920 by 1080 framebuffer: x becomes about 1147.1 and y about 633.5 (in Vulkan, framebuffer y grows downward).

Look at the depth value. The point is only 5 units into a 100 unit view range, yet its depth is already 0.981. Perspective projection packs most of the depth range close to the camera. This is why engines often use a reversed depth buffer (far at 0, near at 1) with a floating point format: it spreads precision more evenly. The arithmetic was checked in Python with the standard perspective matrix.

Vertex shaders run on exactly the same shader cores as pixel shaders. A vertex is just another thread.

Tessellation, geometry shaders and mesh shaders

Optional stages can add geometry. Tessellation splits patches into many small triangles, with a fixed function tessellator between two programmable stages. Geometry shaders take a whole primitive and can emit zero or more new primitives.

Mesh shading replaces the front of the pipeline altogether. The Vulkan specification describes it this way: "input primitives are not assembled implicitly, but explicitly through the (Mesh Shader)." An optional task shader (called the amplification shader in Direct3D) launches mesh shader workgroups, and each workgroup outputs a small mesh "with a variable-sized number of primitives." NVIDIA's Turing whitepaper (2018) introduced mesh shading as a new shader model for the vertex, tessellation and geometry stages, saying it moves "the key performance bottleneck of object list processing off of the CPU" and enables new approaches to object level of detail. The Direct3D Mesh Shader specification exposes the same idea through amplification and mesh shaders.

The reason this matters for hardware is that mesh shaders behave like compute shaders. They use workgroups, shared memory and cooperative threads, which pulls the geometry front end closer to the general purpose compute model described later.

Step 3: primitive assembly, clipping and culling

After vertex processing, fixed function hardware groups vertices into primitives, clips them against the view volume, performs the perspective divide and applies the viewport transform. Triangles that cross the near plane are cut so that no vertex ends up behind the camera.

Then comes face culling. The Vulkan specification says the first step of polygon rasterization is to decide whether a triangle is front facing or back facing from "the sign of the (clipped or unclipped) polygon's area computed in framebuffer coordinates." The area is computed with the shoelace formula over the vertices. A closed, opaque mesh viewed from outside shows only its front faces, so discarding back faces before rasterization removes a large share of the triangles before any pixel work is done. Vulkan treats any triangle that is not front facing as back facing, "including zero-area triangles."

Step 4: rasterization

The Vulkan specification defines rasterization as "the process by which a primitive is converted to a two-dimensional image." For each primitive, the rasterizer determines which squares of the framebuffer grid it occupies. A grid square with its coordinates, depth and associated data "is called a fragment." The test is done at sample locations, which are fixed offsets inside each pixel: with no multisampling there is one sample at the pixel centre.

The standard way to test coverage in hardware comes from Juan Pineda's 1988 SIGGRAPH paper "A Parallel Algorithm for Polygon Rasterization". For a triangle with vertices V0, V1, V2, define an edge function for each edge. For the edge from A to B and a sample point P:

E_AB(P) = (B.x - A.x) * (P.y - A.y) - (B.y - A.y) * (P.x - A.x)

The sign tells you which side of the edge the sample lies on. A sample is inside the triangle when all three edge functions have the same sign. The edge function is linear in x and y, so stepping one pixel to the right adds a constant, which makes it cheap to evaluate for a whole block of pixels in parallel. That property is exactly why it suits hardware.

Worked example: rasterize one triangle on an 8 by 8 grid. Take vertices V0 = (1, 1), V1 = (7, 2), V2 = (3, 7) in framebuffer coordinates, with y increasing downward as in Vulkan. Sample each pixel at its centre (x + 0.5, y + 0.5). Twice the signed area is E_01(V2) = 34, so the triangle's area is 17 square pixels. Testing all 64 centres with the three edge functions gives this coverage, where # marks a covered pixel:

   x: 0 1 2 3 4 5 6 7
y=0   . . . . . . . .
y=1   . # # # . . . .
y=2   . # # # # # # .
y=3   . . # # # # . .
y=4   . . # # # . . .
y=5   . . # # . . . .
y=6   . . . . . . . .
y=7   . . . . . . . .

Eighteen pixel centres are covered, close to the triangle's 17 square pixel area, as you would expect. Real hardware adds tie breaking rules so that a sample exactly on an edge shared by two triangles is counted once, never twice and never zero times.

The same numbers give interpolation weights. For pixel (3, 3), sampled at (3.5, 3.5), the three edge function values are 11.5, 10.0 and 12.5. Dividing each by the doubled area of 34 gives barycentric weights of about 0.338, 0.294 and 0.368 for V0, V1 and V2. They sum to one. If the three vertices have depths 0.2, 0.5 and 0.8, the interpolated depth at that pixel is about 0.509. Every attribute (texture coordinates, normals, colors) is interpolated with these weights. For attributes that should look correct under perspective, hardware divides by the interpolated 1/w as well, which is called perspective correct interpolation. The Vulkan specification has a full section on barycentric interpolation. The coverage grid and weights above were computed in Python.

How hardware organises this. GPUs do not walk pixels one by one. They test blocks of pixels at once, discard blocks that a triangle cannot touch, and emit fragments in groups of 2 by 2 pixels called quads. Quads exist because fragment shaders need screen space derivatives: to choose a texture mip level, the hardware compares a texture coordinate with its neighbours to the right and below. Pixels at a triangle edge that are not covered still run as "helper" invocations inside a partly covered quad so that derivatives are defined. Their results are thrown away. This is one reason why very small triangles are expensive: a triangle that covers one pixel still pays for a whole quad.

Multisample anti-aliasing (MSAA) tests coverage at several sample positions per pixel, typically 4 or 8, but runs the fragment shader once per pixel and stores the color in each covered sample. A coverage mask records which samples the triangle hit. When the frame is finished, the samples are averaged in a resolve step. This is why MSAA multiplies depth and color storage and ROP work, but not shader work.

On NVIDIA's Ada generation the rasterization hardware sits in the Graphics Processing Cluster. The Ada whitepaper says each GPC includes "a dedicated Raster Engine," and the full AD102 chip has 12 GPCs, so a frame's triangles are spread across a dozen rasterizers working in parallel.

Step 5: fragment shading

Each fragment runs the fragment shader (pixel shader in Direct3D). This is where most of a frame's arithmetic usually goes: sampling textures, evaluating lighting, applying shadows and material models.

// GLSL fragment shader, illustrative
#version 450
layout(set = 0, binding = 1) uniform sampler2D albedoTex;
layout(location = 0) in vec2 vUV;
layout(location = 0) out vec4 outColor;
void main() {
    vec3 albedo = texture(albedoTex, vUV).rgb;   // filtered texture fetch
    vec3 lightDir = normalize(vec3(0.3, 0.8, 0.5));
    float ndotl = 0.6;                            // placeholder lighting term
    outColor = vec4(albedo * ndotl, 1.0);
}

The call to texture() does not run as ordinary shader instructions. It is sent to a texture unit, a fixed function block that computes addresses, fetches texels from memory through a cache and filters them.

Texture units and filtering

A texture is an image stored with a chain of smaller copies called mipmaps, each half the width and height of the one before. The full chain costs one third more memory than the base level alone: a 4096 by 4096 RGBA8 texture uses 64 MiB, and about 85.3 MiB with its mip chain, because 1/4 + 1/16 + 1/64 and so on sums to 1/3.

The OpenGL 4.6 core specification defines how a sampler picks a mip level. It computes a scale factor rho from how fast the texture coordinates change across the screen, then a level of detail lambda that is, before biases and clamps, the base 2 logarithm of rho. If one screen pixel spans 4 texels, lambda is 2 and the sampler reads mip level 2. If a pixel spans about 2.83 texels, lambda is about 1.5 and the sampler sits between levels 1 and 2.

The filter modes then decide how many texels are read:

For anisotropic filtering, the OpenGL specification says "the particular scheme for anisotropic texture filtering is implementation dependent," and then describes one possible approach using two scale factors, one per screen axis. This is a case where vendors deliberately do not publish their exact method.

Texture units are counted per SM on NVIDIA parts. The Ada whitepaper lists four texture units per SM. The 128 SMs enabled on the GeForce RTX 4090 therefore give 512 texture units, which matches the 512 in the whitepaper's specification table.

Step 6: depth, stencil, blending and ROPs

After shading, each fragment must be tested and merged into the framebuffer.

Depth test. The depth buffer stores, for each sample, the depth of the closest surface drawn so far. A new fragment passes if its depth compares favourably (usually "less than") with the stored value. If it passes, it may overwrite the stored depth and color. This is how the GPU resolves visibility without sorting triangles.

Early and late depth. Running the shader and then throwing the result away is wasted work. If the shader does not write depth, does not discard fragments and has no other side effects that forbid it, the GPU can run the depth test before the fragment shader. Vulkan calls these "early fragment tests." When a shader writes gl_FragDepth or uses discard, the hardware may have to fall back to testing after the shader. This is a common, invisible performance cliff.

Stencil. A small integer per sample (8 bits in common formats such as Vulkan's VK_FORMAT_S8_UINT) that fragments can test and update, used for masking, portals and outlines.

Blending. For transparent surfaces the new color is combined with the color already in the framebuffer. The most common equation is "source over":

result = src.rgb * src.a + dst.rgb * (1 - src.a)

A half transparent red fragment (1, 0, 0) with alpha 0.5 drawn over a blue pixel (0, 0, 1) produces (0.5, 0, 0.5). Blending needs the existing color, so it reads framebuffer memory as well as writing it. Blending results also depend on the order fragments arrive in, which is why the API guarantees that blending follows primitive order even though shading runs in parallel.

ROPs. The units that perform depth, stencil and blend operations and write samples to memory are called raster operation units, or ROPs, on NVIDIA parts and render backends on AMD parts. The Ada whitepaper says each GPC has "two Raster Operations (ROPs) partitions, with each partition containing eight individual ROP units." Their throughput sets the pixel fill rate, which we calculate below.

Immediate mode versus tile based deferred rendering

Everything above describes the logical pipeline. Two physical strategies implement it.

Desktop and data center GPUs are usually described as immediate mode renderers: triangles flow through the pipeline roughly in submission order, and fragments are written to framebuffer memory as they are produced, with large on chip caches catching much of the traffic.

Most mobile GPUs use tile based rendering. Apple's Metal documentation explains that Apple silicon GPUs use tile based deferred rendering (TBDR): the GPU "breaks up the render destination into a grid of smaller regions, called tiles," processes all the geometry for a render pass first, and "defers, or postpones, the rendering phase for each tile until after it evaluates all the geometry for that tile." Each tile is shaded into tile memory on the chip, and only the final result goes out to device memory. Apple lists the advantages of tile memory as bandwidth many times higher, latency many times lower and energy use much lower than device memory.

Immediate mode                         Tile based deferred
-----------------                      ---------------------
for each triangle:                     pass 1, for each triangle:
  rasterize                              transform, bin into tiles it touches
  shade fragments                      pass 2, for each tile (on chip):
  depth test, blend -> DRAM                 hidden surface removal
                                            shade only visible fragments
                                            blend in tile memory
                                            write finished tile -> DRAM once

The trade is extra work and storage to bin geometry in exchange for far fewer trips to off chip memory. On a phone, where memory traffic costs battery, that trade usually wins. Apple also notes that a TBDR GPU can start the vertex stage of the next render pass while it finishes the fragment stage of the current one, since the two use different hardware.

Practical consequence for developers: on tile based GPUs, telling the API that you do not need to load or store an attachment (for example a depth buffer used only within one pass) lets the GPU skip moving it to and from memory. In Metal this is MTLLoadActionDontCare and MTLStoreActionDontCare; in Vulkan it is the load and store operations on a render pass attachment.

The hardware behind the pipeline: unified shader cores

Early graphics chips had separate hardware for vertex work and pixel work. If a scene was heavy on geometry, the pixel units idled, and the reverse.

NVIDIA's Fermi whitepaper describes the change. The G80 based GeForce 8800, introduced in November 2006, "was the first GPU to replace the separate vertex and pixel pipelines with a single, unified processor that executed vertex, geometry, pixel, and computing programs." The same whitepaper says G80 introduced "the single-instruction multiple-thread (SIMT) execution model" and shared memory with barrier synchronisation. The architecture is described in detail in Lindholm, Nickolls, Oberman and Montrym, "NVIDIA Tesla: A Unified Graphics and Computing Architecture," IEEE Micro, March 2008. (These "first" claims are NVIDIA's own wording.)

The unified design is what made general purpose GPU computing practical: once all shader stages ran on one pool of programmable cores, exposing that pool directly to C programs was a smaller step. (For how the silicon, memory stacks and boards behind these cores are manufactured, see how a GPU is made.)

Inside a modern GPU: the NVIDIA Ada example

The Ada Lovelace whitepaper gives a complete hierarchy for its largest chip, AD102:

• 12 GPCs, 72 Texture Processing Clusters (TPCs), 144 SMs, and a 384 bit memory interface built from twelve 32 bit memory controllers.
• Each GPC has a raster engine, two ROP partitions of eight ROPs each, and six TPCs.
• Each TPC has one PolyMorph Engine (NVIDIA's name for its per TPC geometry front end) and two SMs.
• Each SM has 128 CUDA cores, one third generation RT core, four fourth generation Tensor Cores, four texture units, a 256 KB register file and 128 KB of L1 data cache and shared memory that can be split between the two uses.
• The full chip therefore has 18,432 CUDA cores, 144 RT cores, 576 Tensor Cores and 576 texture units. The whitepaper also notes 2 FP64 cores per SM, giving FP64 at one sixty fourth of the FP32 rate.

AD102 (full chip, per Ada whitepaper)
+---------------------------------------------------------------+
| GigaThread engine / command front end                         |
+---------------------------------------------------------------+
| GPC x 12                                                      |
|  +---------------------------------------------------------+  |
|  | Raster engine   |  ROP partition (8)  ROP partition (8)  |  |
|  +---------------------------------------------------------+  |
|  | TPC x 6                                                 |  |
|  |  +------------------+  +---------------------------+    |  |
|  |  | PolyMorph engine |  | SM            | SM        |    |  |
|  |  +------------------+  | 128 CUDA cores| ...       |    |  |
|  |                        | 4 Tensor Cores|           |    |  |
|  |                        | 1 RT core     |           |    |  |
|  |                        | 4 tex units   |           |    |  |
|  |                        | 256 KB regs   |           |    |  |
|  |                        | 128 KB L1/SMEM|           |    |  |
|  |                        +---------------------------+    |  |
|  +---------------------------------------------------------+  |
+---------------------------------------------------------------+
| L2 cache (72 MB on RTX 4090, 96 MB on the AD102 based L40)     |
+---------------------------------------------------------------+
| 12 x 32-bit GDDR6X memory controllers                         |
+---------------------------------------------------------------+

The GeForce RTX 4090 uses a cut down AD102 with 11 GPCs, 64 TPCs and 128 SMs enabled, according to the whitepaper's appendix: 16,384 CUDA cores, 128 RT cores, 512 Tensor Cores, 512 texture units, 176 ROPs and a 73,728 KB (72 MB) L2 cache. The arithmetic is consistent: 128 SMs times 128 cores is 16,384, and 128 SMs times 4 texture units is 512.

The Ampere GA102 whitepaper (2020) explains the jump in cores per SM between generations. In Turing, each of the four SM partitions had one FP32 datapath and one integer datapath. In GA102, one datapath does FP32 only and the other can do either FP32 or INT32, so "all four SM partitions combined can execute 128 FP32 operations per clock, which is double the FP32 rate of the Turing SM," or 64 FP32 and 64 INT32 operations. That is why the whitepapers list 64 CUDA cores per SM for Turing and 128 for Ampere and Ada. It also means the doubled FP32 peak is only reached when the instruction mix has few integer operations.

The AMD RDNA example: workgroup processors and wave32

AMD's RDNA architecture presentation (2019) describes a different organisation with the same goals. The building block is the Workgroup Processor (WGP), which combines the resources of two Compute Units so that one workgroup can use both. Each SIMD in RDNA is 32 lanes wide. AMD lists "Native Wave32 with support for Wave64 via dual-issue," and "single-cycle instruction issue."

The presentation contrasts this with the older GCN design. In GCN, each 64 thread wave ran on a 16 lane SIMD, which "issues 1 instruction every 4 cycles." In RDNA each wave is assigned to one SIMD32, "up to 20 waves per SIMD32," and a SIMD32 "issues 1 instruction every cycle," with 5 cycles of latency exposed that other waves can fill.

Which wave size is used is a compiler decision. AMD's presentation says "compute and vertex shaders usually as Wave32, pixel shaders usually as Wave64," and gives the trade: Wave32 offers lower latency per wave and is more efficient for partly filled waves; Wave64 allows more threads per lane and is more efficient for attribute interpolation.

How shader cores hide latency

A texture fetch or a global memory load can take hundreds of cycles to return. A CPU core hides that with caches and out of order execution. A GPU core mostly hides it by having other work ready.

Each SM or compute unit keeps many warps or waves resident at once, each with its own registers already loaded. When one warp issues a load and has to wait, the scheduler issues an instruction from another warp that is ready, in the very next cycle, at no switching cost, because nothing has to be saved or restored. The register file is huge for this reason: Ada's is 256 KB per SM, which is larger than the 128 KB of L1 and shared memory in the same SM.

Warp scheduler timeline, one SM sub-partition (illustrative)

cycle:   1    2    3    4    5    6    7    8    9   10   ...  400  401
warp 0: LD   .    .    .    .    .    .    .    .    .   ...  data  FMA
warp 1:      FMA  LD   .    .    .    .    .    .    .   ...
warp 2:                FMA  FMA  LD   .    .    .    .   ...
warp 3:                               FMA  FMA  FMA  LD  ...
  ...
Each cycle the scheduler picks a warp whose next instruction has its inputs.
With enough warps, the arithmetic units never sit idle while loads are in flight.

How much work must be in flight is a direct consequence of Little's law: data in flight equals bandwidth times latency. As an illustrative calculation only, if a GPU with 1,008 GB/s of memory bandwidth (the RTX 4090's published figure) saw an average memory latency of 500 ns, it would need about 504 KB of loads in flight across the chip to keep the memory system busy. The 500 ns figure is an assumption for the example, not a vendor number; NVIDIA does not publish a single memory latency figure for these parts. The point is the shape of the relationship, not the exact value. More bandwidth or more latency means more threads must be resident, which is why registers and occupancy matter so much.

SIMT and divergence

The CUDA Programming Guide describes SIMT this way: threads are grouped into warps of 32, and "all threads in the warp are executing the same kernel code, but each thread may follow different branches through the code." If threads in a warp disagree about a branch, "the threads which do not follow the branch will be masked off while the threads which follow the branch are executed." The guide calls this warp divergence and notes that utilisation is highest "when threads within a warp follow the same control flow path."

Worked example. Suppose half the threads of a 32 thread warp take an if branch that costs 10 instructions and the other half take the else branch that also costs 10. The warp issues 20 instructions in total, but in each one only 16 lanes do useful work. Useful lane slots are 16 times 10 plus 16 times 10, which is 320, out of 32 times 20, which is 640. Efficiency is 50 percent. If all 32 threads agreed, the same work would take 10 issue slots.

// CUDA C++: a divergent branch inside one warp (illustrative, not executed)
__global__ void shade(const float* in, float* out, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i >= n) return;
    float v = in[i];
    if (i % 2 == 0) {          // even lanes: path A
        v = v * 1.5f + 0.25f;
    } else {                   // odd lanes: path B, runs while path A lanes are masked
        v = sqrtf(fabsf(v));
    }
    out[i] = v;
}

The same thing happens in pixel shaders. A material shader that branches on a per pixel value pays for both sides wherever neighbouring pixels disagree. Graphics programmers therefore try to make branches coherent across screen regions.

The CUDA guide adds that "hardware execution may differ" from strict lockstep, points to its sections on independent thread execution, and says that exploiting knowledge of how warp execution maps to real hardware "is discouraged." The rule of thumb still holds: keep threads in a warp on the same path.

Ray tracing hardware: walking the BVH

Rasterization asks, for each triangle, which pixels it covers. Ray tracing asks, for each ray, which triangle it hits first. That second question is natural for reflections, refractions, soft shadows and global illumination, which rasterization can only approximate. NVIDIA's Turing whitepaper lists those approximations and their artifacts as the motivation for hardware ray tracing.

The API model

Microsoft's DirectX Raytracing (DXR) specification describes the model that Vulkan's ray tracing extensions also follow.

• Acceleration structures. Scene geometry is described in two levels. "Bottom-level acceleration structures each consist of a set of geometries," either triangle meshes or procedural primitives described by axis aligned bounding boxes. "A top-level acceleration structure represents a set of instances of bottom-level acceleration structures," each with a transform and an ID. The application asks the driver to build them; their internal format is opaque.
• Shaders. DispatchRays() "invokes a grid of ray generation shader invocations," each of which can call TraceRay(). When a ray hits geometry, an any hit shader can accept or ignore the hit (for alpha tested foliage, for example), an intersection shader handles custom procedural primitives, the closest hit shader runs for the nearest accepted hit, and a miss shader runs if nothing is hit.
• Rays. The DXR specification defines a ray as "an origin, direction and parametric interval (TMin, TMax)," and says the system reduces TMax as closer intersections are found, so that at the end TMax represents the closest hit.
• Inline ray tracing. DXR Tier 1.1 also allows a shader of any stage, including compute and graphics shaders, to trace rays inline with RayQuery, without the separate shader table machinery.

// HLSL ray generation shader (DXR), illustrative, not executed
RaytracingAccelerationStructure Scene : register(t0);
RWTexture2D<float4> Output : register(u0);

struct Payload { float3 color; };

[shader("raygeneration")]
void RayGen() {
    uint2 pix  = DispatchRaysIndex().xy;
    uint2 size = DispatchRaysDimensions().xy;
    float2 ndc = (float2(pix) + 0.5) / float2(size) * 2.0 - 1.0;

    RayDesc ray;
    ray.Origin    = float3(0, 0, -5);
    ray.Direction = normalize(float3(ndc.x, -ndc.y, 1));
    ray.TMin      = 0.001;
    ray.TMax      = 10000.0;

    Payload p = { float3(0, 0, 0) };
    TraceRay(Scene, RAY_FLAG_NONE, 0xFF, 0, 1, 0, ray, p);
    Output[pix] = float4(p.color, 1);
}

What a BVH is

A bounding volume hierarchy is a tree. The root box encloses the whole scene. Each child box encloses part of it, and so on down to leaves that hold a handful of triangles. To find what a ray hits, you test the ray against the root box; if it misses, you are done. If it hits, you test the child boxes, and only descend into those the ray enters.

                [ scene box ]
               /             \
        [ box A ]           [ box B ]
        /      \            /      \
   [A1]       [A2]     [B1]       [B2]
   tri 0-3    tri 4-7  tri 8-11   tri 12-15

Ray misses box B  -> skip 8 triangles with one test.
Ray hits A, misses A1, hits A2 -> test only triangles 4 to 7.

For a well built tree over N triangles, a ray visits a number of nodes that grows roughly with the logarithm of N rather than with N, which is what makes ray tracing scenes with millions of triangles feasible.

Here is the core loop in plain C++, the way a software ray tracer or a shader without hardware support would do it:

// Software BVH traversal sketch (stack based), illustrative, not executed
struct Ray { float3 o, d, invD; float tMin, tMax; };
struct Node { AABB box; int left, right, firstTri, triCount; };

bool slabTest(const AABB& b, const Ray& r, float tMax) {
    float3 t0 = (b.min - r.o) * r.invD;
    float3 t1 = (b.max - r.o) * r.invD;
    float tNear = max3(min(t0, t1));
    float tFar  = min3(max(t0, t1));
    return tNear <= tFar && tFar >= r.tMin && tNear <= tMax;
}

int trace(const Node* nodes, const Tri* tris, Ray& r) {
    int stack[64], sp = 0, hit = -1;
    stack[sp++] = 0;                          // root
    while (sp) {
        const Node& n = nodes[stack[--sp]];
        if (!slabTest(n.box, r, r.tMax)) continue;
        if (n.triCount) {                     // leaf: ray-triangle tests
            for (int i = 0; i != n.triCount; ++i) {
                float t;
                if (intersect(tris[n.firstTri + i], r, t) && t < r.tMax) {
                    r.tMax = t; hit = n.firstTri + i;   // shrink TMax, as DXR describes
                }
            }
        } else {
            stack[sp++] = n.left;
            stack[sp++] = n.right;
        }
    }
    return hit;
}

What RT cores accelerate

Every iteration of that loop is a box test or a triangle test, each with a data dependent branch and a memory fetch. On ordinary shader cores, NVIDIA's Turing whitepaper says, this "would need to be performed by shader operations and take thousands of instruction slots per ray cast." Divergence makes it worse: neighbouring rays in a warp quickly visit different nodes.

Turing added an RT core to each SM. According to the whitepaper, the RT core "includes two specialized units. The first unit does bounding box tests, and the second unit does ray-triangle intersection tests." The SM launches a ray probe, the RT core traverses the BVH "autonomously," and returns hit or no hit, while the SM carries on with "other vertex, pixel, and compute shading work." Building and refitting the BVH is done by the driver, and ray generation and shading stay in programmable shaders.

 SM (shader cores)                     RT core (fixed function)
 -----------------                     -------------------------
 ray generation shader
   TraceRay(ray) -------------------> box intersection engine
   (warp may switch to other work)      walk BVH nodes
                                      triangle intersection engine
                                        test leaf triangles
   <---------------------------------- closest hit / any hit / miss result
 hit or miss shader runs on SM
 (material, lighting, maybe more rays)

The Ada whitepaper describes the third generation RT core: the same box and triangle engines plus an Opacity Micromap Engine (for alpha tested geometry such as leaves: the whitepaper says the RT core can then resolve many rays "without invoking any shader code," and reports a measured doubling of traversal performance for alpha tested geometry that depends heavily on usage) and a Displaced Micro-Mesh Engine (for very finely detailed geometry stored compactly).

Worked numbers, with care. The Turing whitepaper's own table lists the GeForce RTX 2080 Ti at 10 Giga Rays per second. At 3840 by 2160 and 60 frames per second (497.7 million pixels per second), that is about 20 rays per pixel per frame at the published peak. NVIDIA presents this as a headline figure; it is not a sustained number for an arbitrary scene, and real ray throughput depends heavily on how coherent the rays are and how deep the BVH is. Later whitepapers report "RT TFLOPS" instead (191 for the RTX 4090 in the Ada whitepaper), which is a different metric and cannot be converted back into rays per second without information NVIDIA has not published.

Hybrid rendering. Real time renderers do not ray trace everything. The Turing whitepaper describes a "hybrid rendering" approach in which "rasterization is used where it is most effective, and ray tracing is used where it provides the most visual benefit," such as reflections, refractions and shadows, combined with denoising filters that reconstruct a clean image from few rays per pixel. That is still the standard approach.

From graphics to general purpose compute

Once vertex and pixel shaders ran on one pool of programmable cores, researchers began expressing non graphics problems as pixel shaders: store data in textures, draw a full screen rectangle, read the result back. It worked, but it was awkward. G80, introduced in November 2006 according to NVIDIA's Fermi whitepaper, removed the graphics disguise: the whitepaper says G80 "was the first GPU to support C," and Nickolls, Buck, Garland and Skadron described the programming model in "Scalable Parallel Programming with CUDA" in ACM Queue in March 2008.

Why throughput beats latency for these jobs

The workloads that moved to GPUs share the properties that pixels have: lots of independent elements, the same operation on each, and arithmetic that can be arranged to reuse data. Dense linear algebra is the clearest case. Multiplying two N by N matrices takes about 2N³ floating point operations on 3N² numbers, so the arithmetic per byte grows with N. That lets a GPU keep its arithmetic units busy rather than waiting on memory.

Graphics and compute now share the same hardware pipeline entry points. Vulkan describes the compute pipeline as "a separate pipeline from the graphics pipeline, which operates on one-, two-, or three-dimensional workgroups." The hardware is the same SMs or compute units; only the front end differs. A compute dispatch skips the input assembler, rasterizer and ROPs and launches thread groups directly.

The same idea in three languages

A simple operation, scaling a vector and adding another (SAXPY), shows how the graphics and compute models line up.

// CUDA C++ (illustrative, not executed)
__global__ void saxpy(int n, float a, const float* x, float* y) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) y[i] = a * x[i] + y[i];
}
// host: saxpy<<<(n + 255) / 256, 256>>>(n, 2.0f, d_x, d_y);

// GLSL compute shader for Vulkan (illustrative, not executed)
#version 450
layout(local_size_x = 256) in;
layout(set = 0, binding = 0) readonly buffer X { float x[]; };
layout(set = 0, binding = 1) buffer Y { float y[]; };
layout(push_constant) uniform P { float a; uint n; } p;
void main() {
    uint i = gl_GlobalInvocationID.x;
    if (i < p.n) y[i] = p.a * x[i] + y[i];
}

// Metal Shading Language kernel (illustrative, not executed)
#include <metal_stdlib>
using namespace metal;
kernel void saxpy(device const float* x [[buffer(0)]],
                  device float*       y [[buffer(1)]],
                  constant float&     a [[buffer(2)]],
                  constant uint&      n [[buffer(3)]],
                  uint i [[thread_position_in_grid]]) {
    if (i < n) y[i] = a * x[i] + y[i];
}

In all three, one thread handles one element, threads are grouped (256 per block or workgroup here), and the hardware schedules the groups across its cores. The mapping is the same one used for pixels.

SAXPY itself is a poor fit for showing off a GPU, and that is worth seeing in numbers. Each element reads 8 bytes (x and y), writes 4 bytes, and performs 2 floating point operations: 2 operations per 12 bytes, about 0.17 FLOP per byte. As we calculate in the next section, an RTX 4090 needs about 82 FLOP per byte to be limited by arithmetic rather than by memory. SAXPY is therefore entirely bandwidth bound: it will run near memory speed, and the arithmetic units will mostly wait.

Tensor cores and the AI turn

Matrix multiplication became important enough that GPUs added dedicated units for it. The Ada whitepaper lists four Tensor Cores per SM, and its specification table gives the RTX 4090 a peak of 165.2 dense FP16 Tensor TFLOPS with FP32 accumulate (330.4 with NVIDIA's sparsity feature), against 82.6 TFLOPS for ordinary FP32 shader math. Those are peak rates at the boost clock, per the whitepaper's footnote. Accelerators built for AI push this balance much further. Our posts on what an AI accelerator is and how it works and AI chip architecture go into that design space, and what a GPU costs per million tokens when you self host turns peak figures like these into serving economics.

Worked numbers: fill rate, texel rate and FLOPS

Peak GPU figures in vendor tables are products of unit counts and clock speed. Reproducing them is a good way to understand what each unit does. The NVIDIA Ada whitepaper's appendix compares three GeForce cards side by side, all using the GPU Boost clock, which the whitepaper states in a footnote is the basis for peak rates. We recomputed every figure in Python.

Inputs from the Ada whitepaper appendix:

The formulas:

pixel fill rate (Gpixel/s)  = ROPs           x boost clock (GHz)
texel fill rate (Gtexel/s)  = texture units  x boost clock (GHz)
FP32 peak (TFLOPS)          = CUDA cores x 2 x boost clock (GHz) / 1000
                              (2 because one fused multiply add counts as 2 operations)
memory bandwidth (GB/s)     = bus width (bits) x data rate (Gbps) / 8

Results (computed, then compared with the published table):

All three match the whitepaper. That tells you what the published numbers mean: one pixel per ROP per clock, one filtered texel per texture unit per clock, and one fused multiply add per CUDA core per clock, all at the boost clock. They are peaks. A real frame never keeps every ROP, every texture unit and every core busy at once, clocks vary with power and temperature, and NVIDIA does not publish sustained figures for games. No benchmark results are claimed here.

What the peaks mean for a 4K frame

Take the RTX 4090 and a 3840 by 2160 frame (8,294,400 pixels).

The first column after the frame rate looks enormous: hundreds of pixel writes per pixel. In practice that headroom is consumed by overdraw, multiple render targets, shadow maps, post processing passes and MSAA samples. A single full screen RGBA8 render target at 4K is 33.2 MB, an RGBA16F target is 66.4 MB, and a 32 bit depth buffer is another 33.2 MB. A deferred renderer that writes four such targets, reads them back for lighting, and then runs several post processing passes moves many hundreds of megabytes per frame before textures are counted. The per frame memory budget, not the fill rate, is usually what limits high resolution rendering, and that is one reason the RTX 4090's L2 cache is 72 MB rather than the 6 MB of the RTX 3090 Ti.

A roofline view

The roofline model compares a kernel's arithmetic intensity (operations per byte moved from memory) with the machine's balance point, peak FLOPS divided by peak bandwidth.

attainable
 TFLOPS
   82.6 |                         ______________________  compute roof (RTX 4090)
        |                       /
        |                     /
        |                   /   memory roof: 1,008 GB/s x intensity
        |                 /
        |               /
        |  SAXPY      /
        |   x       /
        |_________/__________________________________________
           0.17    81.9                       FLOP per byte (log scale)

A kernel below about 82 FLOP per byte on an RTX 4090 cannot reach peak FP32 no matter how well it is written; it is limited by memory. SAXPY at about 0.17 FLOP per byte sits far to the left. A blocked matrix multiply that reuses tiles from shared memory or registers can sit far to the right. Notice also the trend across generations: compute grew about 5.8 times from the 2080 Ti to the 4090 while bandwidth grew about 1.6 times, so the balance point moved right. Each generation, more work becomes memory bound unless caches catch the reuse.

These balance points use DRAM bandwidth. Data served from L2 or shared memory has a much higher effective bandwidth, which is why tiling and caching matter so much in both shaders and compute kernels. NVIDIA does not publish L2 bandwidth figures for these GeForce parts in the whitepapers we read.

A short dated timeline

Common mistakes

Treating peak figures as delivered performance. Fill rate, texel rate and TFLOPS in vendor tables are products of unit counts and boost clock. They describe a ceiling. Two GPUs with the same TFLOPS can perform very differently because of memory bandwidth, cache size or how well a workload fits the instruction mix (recall that Ampere's doubled FP32 rate depends on few integer instructions).

Comparing TFLOPS across precisions or sparsity modes. A Tensor Core FP16 figure with sparsity is not comparable with an FP32 shader figure. The Ada whitepaper marks its sparsity figures with a footnote; read it.

Ignoring the memory roof. Many shaders and kernels are bandwidth bound. Adding arithmetic optimisations to a kernel at 0.2 FLOP per byte achieves nothing; reducing bytes moved does.

Breaking early depth testing by accident. Writing depth from a fragment shader, or discarding fragments, can force the depth test to run after shading. Keep discard out of shaders for opaque geometry where you can, and use a depth prepass if you must.

Drawing huge numbers of tiny triangles. Rasterizers emit 2 by 2 quads. A triangle smaller than a pixel still pays for a quad of shader invocations, three of which may be helpers whose output is discarded.

Divergent branches in hot shaders. A per pixel branch that alternates between neighbours halves throughput in the worst case, as the worked example showed.

Assuming every GPU is immediate mode. On tile based mobile GPUs, unnecessary attachment loads and stores, or reading back a render target in the middle of a pass, force tiles out to memory and erase the bandwidth advantage.

Calling ray tracing a replacement for rasterization. Shipping real time renderers use hybrid pipelines. RT units accelerate traversal and intersection; shading still runs on the shader cores.

Leaving the GPU software stack out of inventory. Drivers, compilers and runtime libraries change often and carry security fixes. Teams that track application dependencies carefully often have no record of which driver or CUDA runtime is inside a container image.

FAQ

How does a GPU render a frame, step by step?

The application records draw calls; the driver compiles shaders and builds command buffers; the GPU fetches vertices and runs vertex shaders; fixed function hardware assembles, clips and culls triangles; the rasterizer finds which pixel samples each triangle covers; fragment shaders compute colors using texture units; and the depth test, stencil test and blending units (ROPs) write the visible results into the framebuffer, which the display engine scans out.

What is the difference between a vertex shader and a fragment shader?

A vertex shader runs once per vertex and outputs its clip space position plus attributes such as texture coordinates. A fragment (pixel) shader runs once per covered fragment and outputs a color, using attributes the rasterizer interpolated from the three vertices. On modern GPUs both run on the same unified shader cores.

What is rasterization in simple terms?

Rasterization converts a triangle into the list of pixels it covers. Hardware evaluates three edge functions at each pixel's sample point; if all three agree on sign, the sample is inside. The same values give barycentric weights to interpolate depth, color and texture coordinates.

What does a ROP do in a GPU?

A ROP (raster operation unit) performs depth and stencil tests, blends new colors with existing framebuffer values, handles multisample data, and writes results to memory. On NVIDIA's Ada parts there are eight ROPs per partition and two partitions per GPC, per the Ada whitepaper. ROP count times clock gives the published pixel fill rate.

How do RT cores speed up ray tracing?

They take over the inner loop: walking the bounding volume hierarchy with ray and box tests and running ray and triangle tests at the leaves. NVIDIA's Turing whitepaper says that without such hardware this takes "thousands of instruction slots per ray" on shader cores. Ray generation and shading still run as programmable shaders.

Why are GPUs good for AI and other non graphics work?

Because the unified shader cores are general purpose programmable processors built for throughput, and many workloads, matrix multiplication above all, are large numbers of similar independent operations with data reuse. Dedicated Tensor Cores then raise the matrix math peak well above the ordinary FP32 rate.

What is a warp, and what is warp divergence?

On NVIDIA GPUs a warp is a group of 32 threads that issue the same instruction together (SIMT). If threads in a warp take different branches, the hardware runs each path with the non participating threads masked off, so the warp pays for both paths. AMD's equivalent is the wavefront, 32 or 64 wide on RDNA.

What is tile based deferred rendering?

A strategy used by Apple silicon GPUs and many mobile GPUs: the GPU first bins all of a render pass's geometry into screen tiles, then shades each tile in fast on chip tile memory, removing hidden surfaces before shading and writing each finished tile to device memory once. It saves memory bandwidth and energy.

How is GPU fill rate calculated?

Pixel fill rate is the number of ROPs times the clock; texel rate is the number of texture units times the clock. For the GeForce RTX 4090, 176 ROPs at 2.52 GHz gives 443.5 gigapixels per second and 512 texture units gives 1,290.2 gigatexels per second, matching the Ada whitepaper. These are peaks, not sustained rates.

Do GPUs have security vulnerabilities?

The GPU software stack does. Vendors publish security bulletins for drivers and GPU software on their product security pages, and libraries, runtimes and container hooks that ship inside images need the same patching as other dependencies. Track them in an SBOM and update them like any other component.

Where Safeguard fits

Safeguard does not design, program or speed up GPUs, and nothing in this article depends on it. The honest connection is the software around the GPU. A rendering or AI workload depends on a graphics or compute driver, a user space runtime (CUDA, ROCm, Vulkan loaders, Mesa), math and communication libraries such as cuDNN and NCCL, a container runtime hook such as the NVIDIA Container Toolkit, and an application stack on top. NVIDIA and AMD both run product security pages that publish bulletins for their GPU software, so these components need the same inventory and patch discipline as any other dependency.

Safeguard's generally available tooling applies to the part you build and ship. It generates and analyses SBOMs in SPDX and CycloneDX, scans repositories (safeguard scan --dir .), container images (safeguard scan --image <name>) and existing SBOM files (safeguard scan --sbom <file>), matches components against vulnerability databases, ranks findings with reachability analysis, and applies policies and gates. The platform shows which projects contain a given component version, which helps when a bulletin lands for a library baked into many inference or rendering images. What it does not do: it does not inspect GPU firmware, it does not patch a kernel mode driver installed on a host outside your images, and it does not change how fast a shader or kernel runs.

A GPU workload is only as current as the driver, runtime libraries and container image it runs in, and those images pull in open source components that receive security fixes on their own schedules. The GPU itself does not change that, so the patch discipline for the software around it is the same as for any other service.

More in this series

Other posts in our GPUs, CPUs, QPUs and lithography series:

• AMD GPU Architecture Explained: GCN, RDNA, CDNA, Infinity Fabric and HIP
• CPU vs GPU vs QPU Explained: How Each Works and Why They Coexist
• CUDA Programming Model Explained: Grids, Blocks, Warps and Memory
• GPU Memory and Interconnects: HBM, GDDR, NVLink, PCIe, CXL
• CUDA vs HIP vs SYCL vs OpenCL vs Metal vs Vulkan vs Triton Compared
• GPU Software Supply Chain: Drivers, CUDA, cuDNN, NCCL and Your SBOM
• How a QPU Works: Qubit Chips, Fridges, Control Pulses and Readout
• How CPUs Work: Pipelines, Out-of-Order, Caches, Branches, SIMD
• How EUV Lithography Machines Work: Tin Plasma, Mirrors and High-NA
• Inside an NVIDIA SM: Warp Schedulers, Registers, Tensor Cores
• Intel GPU Architecture: Xe Cores, XMX, Ponte Vecchio and oneAPI SYCL
• NVIDIA GPU Architecture History: From Tesla and Fermi to Blackwell
• NVIDIA vs AMD vs Intel vs Apple GPUs: Architectures Compared
• Optimizing a CUDA Matrix Multiply Step by Step, From Naive to Tensor Cores
• Photolithography Explained: Contact Printing to DUV Immersion and Overlay
• Tile-Based GPUs Explained: Apple, Mali, Adreno, PowerVR and TBDR

Fix it without waiting for upstream

Whatever the vulnerability, most of the risk sits in two gaps. The first is the time between a fix being published and that fix actually running in your production builds. The second, and worse, is the time before any fix exists at all, while a maintainer works out a patch and a release schedule. Waiting through both gaps is a habit, not a requirement. Safeguard gives teams three ways to close them.

1. Autonomous remediation: pull requests, not tickets

When a fixed version exists, Safeguard's Griffin models open a tested, policy gated pull request that moves the affected dependency to the safe version. You can run it on one repository or in bulk across every project that carries the vulnerable version. Remediation covers eight categories, including dependencies, licenses, security posture and findings. Reachability analysis ranks the vulnerable code paths your application actually calls first, so the first pull requests you merge are the ones that matter most. Autonomous remediation ships switched off, with every category set to Manual, so your team decides how much runs on its own and what still needs a human approval.

2. Gold Open Source: the safe version in seconds, free

gold.safeguard.sh is free and needs no login. Search any CVE or package to see affected and fixed versions, CVSS, EPSS, CISA KEV status and a recommended safe version, or use the "Am I affected?" checker. The same data is available through a free read API, README status badges and RSS threat feeds. The safeguard-sh/gold-scan-action GitHub Action can fail a build on a known exploited vulnerability, on any known vulnerability, or never, so a version you have just removed cannot quietly come back in next week's merge.

3. Gold Registry: start from components that are already fixed

registry.safeguard.sh provides hardened, drop in replacements for upstream packages and container images. Every Gold artifact is held continuously to zero exploitable critical or high CVEs, zero malware and a completed license check, with SLSA L3 or higher provenance, a Sigstore signature and a public evidence page you can hand to an auditor. An artifact loses its Gold badge if it regresses. Teams can point their own package directory at the registry so hardened equivalents resolve without changing how developers install anything.

When upstream has no fix yet, or will not ship one, Safeguard first works to get the fix accepted upstream. If the issue is severe or upstream will not move, Safeguard builds a fork, fixes the vulnerability there and publishes a hardened Gold build, so you are not stuck on a maintainer's release schedule. You can also request a hardened, compatibility tested build of a specific package. This fork and patch flow is in early access today, so ask which packages it covers before you plan around it.

One honest limit: when the vulnerable code lives in a vendor appliance, a hosted service or an operating system runtime you do not build yourself, the fix still has to come from that vendor or distribution. The steps above apply to the open source dependencies and container images you build and ship, which is where most teams carry the largest and least tracked share of their exposure.

See how Safeguard would help

These demos show the software side only. None of them touches GPU hardware, firmware or shader performance.

The first demo is a real Safeguard CLI session. You ask it to fix a lodash CVE; it reads the finding and the advisory, asks before any command that changes something, edits package.json, runs npm install and npm test with your approval, and pushes nothing. The same flow applies to a Python inference service whose requirements pin a vulnerable library next to its GPU runtime.

Click through it at your own pace, or open it full screen.

The second demo sets up the local runner on your own machine, approves it with a device code and then scans from the terminal. That suits teams who build rendering or training images on internal build hosts and want to scan them where they are built.

Click through it at your own pace, or open it full screen.

The third demo connects Claude to the Safeguard MCP server with OAuth and chooses which tools Claude may use. An engineer reading a driver or runtime security bulletin can then ask, in the same conversation, which projects still carry the affected component.

Click through it at your own pace, or open it full screen.

To try the CLI yourself, download Safeguard and follow the CLI documentation.

Quick check

Sources

• Vulkan specification (Khronos registry, latest)
• OpenGL 4.6 core profile specification (Khronos, May 2022)
• DirectX Raytracing (DXR) functional specification (Microsoft)
• Mesh Shader specification (Microsoft DirectX Specs)
• Direct3D 11 graphics pipeline overview (Microsoft Learn)
• NVIDIA Ada GPU architecture whitepaper
• NVIDIA Ampere GA102 GPU architecture whitepaper
• NVIDIA Turing GPU architecture whitepaper
• NVIDIA Fermi compute architecture whitepaper
• CUDA Programming Guide: introduction
• CUDA Programming Guide: programming model, warps and SIMT
• AMD RDNA architecture presentation (GPUOpen, 2019)
• AMD GPUOpen RDNA resources
• Apple Metal: tailor your apps for Apple GPUs and tile based deferred rendering
• Lindholm, Nickolls, Oberman, Montrym: NVIDIA Tesla, a unified graphics and computing architecture, IEEE Micro, 2008 (DOI 10.1109/MM.2008.31)
• Pineda: A parallel algorithm for polygon rasterization, SIGGRAPH 1988 (DOI 10.1145/54852.378457)
• Nickolls, Buck, Garland, Skadron: Scalable parallel programming with CUDA, ACM Queue, 2008 (DOI 10.1145/1365490.1365500)
• NVIDIA product security
• AMD product security
• NVIDIA Container Toolkit documentation
• Mesa documentation

Last reviewed: 25 October 2026.