GPGPU Compute Concepts

General Purpose Graphics Programming Unit (GPGPU) computing is an effective way of using all of the computational resources your computer has to offer by offloading work to the GPU, making it an easy way to access heterogeneous parallel computing. Even though the "G" in GPU stands for graphics, it's more accurate to call modern GPUs highly parallel general processors capable of processing large quantities of data quickly thanks to having a large number of simple microprocessor cores, called Stream Processors (AMD) or Compute Unified Device Architecture (CUDA) cores (NVIDIA), or just GPU cores in other architectures like the Apple M series. These microprocessor cores are designed to not preform well with latency or branching, but instead focus on throughput [Hwu et al. 2022]. This architecture is designed around processing large amounts of data with a single instruction, otherwise known as a Single Instruction Multiple Data (SIMD) or more accurately Single Instruction Multiple Thread (SIMT) architecture, where a group of threads process a single instruction in lock-step.

On the AMD RDNA3 architecture in the AMD RX 7900 XTX, when processing a compute shader, every scheduled stream processor (96 Compute Units (CUs) with 64 shading units each for a maximum of 6144 stream prcessors) receives the same instructions, with the only differences being what the current thread ID is for that a single thread, which waves are on which instruction, and what data is accessible between threads of the same wave/group:

; API shader hash:                  0x00000000000001009731533EAB7158CE
; API PSO hash:                     0xB4116117903F261D
; Driver internal pipeline hash:    0x9337B876E9519DBAD84A08F55EB136FC
;
; Vector registers:                 56 (64 allocated)
; Scalar registers:                 52 (128 allocated)

_amdgpu_cs_main:
 s_version 0x4004                 ; 000000000000: B0804004
 s_inst_prefetch 0x3              ; 000000000004: BFA00003
 s_getpc_b64 s[0:1]               ; 000000000008: BE801F80
 s_mov_b32 s4, s2                 ; 00000000000C: BE840302
 s_mov_b32 s5, s1                 ; 000000000010: BE850301
 s_mov_b32 s7, s1                 ; 000000000014: BE870301
 s_movk_i32 s10, 0x1000           ; 000000000018: B00A1000
    ; ...

Groups of threads are executed together in hardware in what's called a Wave (DirectX), Subgroup (Vulkan), Wavefront (AMD), SIMD Group (Apple), or Warp (NVIDIA CUDA). On both recent AMD and NVIDIA hardware, waves have 32 lanes, with each lane executing one thread. Older AMD hardware traditionally uses 64 threads per wave.

Since this term is different for every hardware vendor and graphics API, we'll stick with the word wave for all general mentions of the term unless working in a specific graphics API.

GPUs can be asked to perform local arithmetic operations, transcendental operations that cost 4 cycles such as square roots, cosine/sine, load/store operations, or hardware accelerated ray tracing traversal or machine learning convolution steps, making them very similar to ASICs or FPGAs.

Compute shaders (or kernels) are programs that execute computations on the GPU, and are useful for a variety of tasks within and outside of computer graphics:

• Postprocessing - traditional fragment based post-processing effects can be easily rewritten as compute shaders and executed more efficiently and even at the same time.

• Geometric Processing - every aspect of ray tracing, from the routines to trace rays, build bounding volume hierarchies (BVHs), and updating them, can be done by the GPU, which can process large quantities of triangle data in real time. Physics simulations such as rigid bodies can be simulated by processing vertices.

• Encoding/Compression - Video encoding, while traditionally hardware accelerated, can be done on compute. Compressing mesh data, scene data such as your BVH or KD trees, compressing frame buffer attachments such as 3D normals to 2D view space normals in a preprocessing step on the GPU, etc.

• Materials and Lighting - Some games opt to use compute shaders in tandem with lookup structured buffers to add a level of abstraction from rasterization, with Nanite using software rendering in combination with meshlets; Mortal Kombat 11 indexed material IDs from a given fragment for its lighting passes.

Let's review GPGPU computing, the terminology, and techniques involved in writing GPGPU programs in different graphics and compute APIs.

Dispatch

Dispatch in WebGPU
// 🚄 Dispatch Compute Work
let width = 1920;
let height = 1080;
let workgroupSize = { x: 8, y: 4 };
let workgroupCount = {
    x: Math.ciel(width / workgroupSize.x),
    y: Math.ciel(height / workgroupSize.y),
    z: 1
};
passEncoder.dispatchWorkgroups(workgroupCount.x, workgroupCount.y, workgroupCount.z);

Dispatch in C++ DirectX 12
// 🚄 Dispatch Compute Work
unsigned width = 1920;
unsigned height = 1080;
unsigned numThreadsX = 8;
unsigned numThreadsY = 4;
unsigned threadgroupX = static_cast<unsigned>(ceil(static_cast<float>(width) / static_cast<float>(numThreadsX)));
unsigned threadgroupY = static_cast<unsigned>(ceil(static_cast<float>(height) / static_cast<float>(numThreadsY)));
unsigned threadgroupZ = 1;
commandList->Dispatch(threadgroupX, threadgroupY, threadgroupZ);

A dispatch call in a given graphics API can process groups of work (aptly called workgroups or threadgroups), normally a set of threads that can number anywhere between 8-256, though generally you'll want to stick with groups of 32 or 64 threads, the size of a wave on most hardware, unless your workload would benefit from more data locality over large groups such as a radix sort. In most graphics APIs, the shader itself describes the thread counts for each of these thread groups, Metal is the only API that doesn't have this requirement and lets you set the number of threads per group at dispatch time. Other APIs could do the same however with shader permutations.

Compute Shader

// 🫂 Local Group Thread ID
@builtin(local_invocation_id) localInvocationID: vec3<u32>

// 👶 Group ID
@builtin(workgroup_id) workgroupID: vec3

// 🖼️ Dispatch Thread ID
@builtin(global_invocation_id) globalInvocationID: vec3<u32>

// 🗂️ Group Index
@builtin(local_invocation_index) localInvocationIndex: u32

// 🔢 Number of Thread Groups
@builtin(num_workgroups) numWorkgroups: vec3<u32>

// 🫂 Local Group Thread ID
uint3 groupThreadID : SV_GroupThreadID

// 👶 Group ID
uint3 groupID : SV_GroupID

// 🖼️ Dispatch Thread ID
uint3 dispatchThreadID: SV_DispatchThreadID

// 🗂️ Group Index
uint groupIndex : SV_GroupIndex

// 🔢 Number of Thread Groups
// HLSL Has no builtin for the number of workgroups.

// 🫂 Local Group Thread ID
gl_LocalInvocationID

// 👶 Group ID
gl_WorkGroupID

// 🖼️ Dispatch Thread ID
gl_GlobalInvocationID

// 🗂️ Group Index
gl_LocalInvocationIndex

// 🔢 Number of Thread Groups
gl_NumWorkGroups

// 🫂 Local Group Thread ID
uint3 threadPositionInThreadgroup [[thread_position_in_threadgroup]]

// 👶 Group ID
uint3 threadPositionInGrid [[threadgroup_position_in_grid]]

// 🖼️ Dispatch Thread ID
uint3 dispatchThreadID [[thread_position_in_grid]]

// 🗂️ Group Index
uint threadIndex [[thread_index_in_threadgroup]]

// 🔢 Number of Thread Groups
uint3 tpg [[threads_per_grid]]

• Group Thread ID - What thread you're currently working on within a workgroup.

• Group ID - What workgroup you're currently exectuing in, useful for designating certain workgroups for the final parts of an algorithm such as prefix sum.

• Dispatch Thread ID - The actual thread that's currently being executed by the hardware, the "pixel" in the case of a screen space dispatch.

• Group Index - a linear representation of the current thread index.

• Number of Thread Groups - What's passed as arguments to your dispatch, the number of thread groups currently executing for this shader kernel.

Group Local Data

var<workgroup> local_data: array<i32,32*16>;

groupshared uint localData[32*16];

Each workgroup can allocate memory within that group that can be shared between threads of that group. This is particularly useful for peforming cross-lane operations and wave operations such as determining the prefix sum of a given value in every thread, radix binning, scans, sweeps, scatter, and compaction.

Shader Model 6.0 introduces wave intrinsics to HLSL, though this is exclusive to both HLSL and GLSL as subgroup intrinsics at the moment. There is a proposal to bring this to WGSL as well.

In the RDNA3 hardware this looks nearly identical to the shader itself, with that data shared within a given compute unit's L1 cache, known as the local data share (LDS). Level cache rules apply to GPUs as well, the more local your data the faster your algorithm can execute, so try to keep data as close to a thread as possible. This means using local variables (which become scalar or vector registers in hardware) more often than global memory such as your strctured buffers, textures, etc.

Atomic Operations

// https://www.w3.org/TR/WGSL/#atomic-builtin-functions
fn atomicLoad(atomic_ptr: ptr<AS, atomic<T>, read_write>) -> T
fn atomicStore(atomic_ptr: ptr<AS, atomic<T>, read_write>, v: T)

fn atomicAdd(atomic_ptr: ptr<AS, atomic<T>, read_write>, v: T) -> T
fn atomicSub(atomic_ptr: ptr<AS, atomic<T>, read_write>, v: T) -> T
fn atomicMax(atomic_ptr: ptr<AS, atomic<T>, read_write>, v: T) -> T
fn atomicMin(atomic_ptr: ptr<AS, atomic<T>, read_write>, v: T) -> T
fn atomicAnd(atomic_ptr: ptr<AS, atomic<T>, read_write>, v: T) -> T
fn atomicOr(atomic_ptr: ptr<AS, atomic<T>, read_write>, v: T) -> T
fn atomicXor(atomic_ptr: ptr<AS, atomic<T>, read_write>, v: T) -> T

fn atomicExchange(atomic_ptr: ptr<AS, atomic<T>, read_write>, v: T) -> T
fn atomicCompareExchangeWeak(atomic_ptr: ptr<AS, atomic<T>, read_write>, cmp: T, v: T) -> __atomic_compare_exchange_result<T>
struct __atomic_compare_exchange_result<T> {
  old_value : T; // old value stored in the atomic
  exchanged : bool; // true if the exchange was done
}

Atomic operations allow for threads to perform logic on a large dataset and write the result to either group shared or global memory. This is useful for performing a variety of tasks such as atomicAdds for prefix sums, atomic min/max for finding the range of points, or atomicXors for compression.

WebGPU is unique in that atomics are strictly typed, whereas with HLSL atomic operations can work on regular memory locations.

Code Example

Here's a simple example of GPGPU computing in a variety of different shader languages:

🌐 WGSL Compute Shader
struct Constants
{
    time: f32
};
 
@group(0) @binding(0) var<uniform> constants: Constants;
@group(0) @binding(1) var tOutput: texture_storage_2d<rgba8unorm,write>;
 
fn foo() -> f32
{
    return 4.0;
}
 
@compute @workgroup_size(8, 4, 1)
fn main(@builtin(local_invocation_id) localInvocationID: vec3<u32>,
        @builtin(workgroup_id) workgroupID: vec3<u32>,
        @builtin(local_invocation_index) localInvocationIndex: u32,
        @builtin(global_invocation_id) globalInvocationID: vec3<u32>) {
    
    var color = vec4<f32>(
        (f32(localInvocationID.x) / 8.0) + sin((f32(globalInvocationID.x) / 160.0) + (constants.time * 0.01)), 
        f32(localInvocationID.y) / foo(),
        f32(globalInvocationID.y) / 360.0, 
        255
    );
 
    textureStore(tOutput, globalInvocationID.xy, color);
 
}

❎ HLSL Compute Shader
struct Constants
{
  float time : packoffset(c0);
};
 
ConstantBuffer<Constants> constants : register(b0);
Texture2D<float4> albedoTex : register(t1);
Texture2D<float4> pbrLutTex : register(t2);
SamplerState gSampler : register(s0);
 
RWTexture2D<float4> tOutput : register(u0);
 
groupshared float groupData[4];
 
namespace random
{
    float foo()
    {
        return 16.0;
    }
}
 
[numthreads(8, 4, 1)]
void main(uint3 groupThreadID : SV_GroupThreadID,
         uint3 groupID : SV_GroupID,
         uint  groupIndex : SV_GroupIndex,
         uint3 dispatchThreadID: SV_DispatchThreadID)
{
    tOutput[dispatchThreadID.xy] = float4( float(groupThreadID.x) / random::foo(), float(groupThreadID.y) / 16.0, dispatchThreadID.x / 1280.0, 1.0);
}

🌋 GLSL Compute Shader
#version 450
layout(local_size_x = 8, local_size_y = 4, local_size_z = 1) in;
 
layout(binding = 0, std140) uniform Constants
{
    float time;
} constants;
 
layout(binding = 0, rgba32f) uniform readonly writeonly image2D tOutput;
 
shared float groupData[4];
 
float foo()
{
    return 16.0;
}
 
void main()
{
    imageStore(tOutput, gl_GlobalInvocationID.xy,
               vec4(float(gl_SubgroupInvocationID.x) / foo(),
                    float(gl_SubgroupInvocationID.y) / 16.0,
                    gl_GlobalInvocationID.x / 1280.0, 1.0));
}

🤖 MSL Compute Shader
#include <metal_stdlib>
#include <simd/simd.h>
 
using namespace metal;
 
struct Constants
{
    float time;
};
 
float foo()
{
    return 16.0;
}
 
kernel void main(
    uint3 groupThreadID [[ thread_position_in_threadgroup ]],
    uint3 groupID [[threadgroup_position_in_grid]],
    uint groupIndex [[thread_index_in_threadgroup]],
    uint3 dispatchThreadID [[thread_position_in_grid]],
    constant Constants& constants [[buffer(0)]],
    texture2d<float, access::read_write> tOutput [[texture(0)]])
{
    threadgroup float groupData[4];
    tOutput[dispatchThreadID.xy] = float4( float(groupThreadID.x) / foo(), float(groupThreadID.y) / 16.0, dispatchThreadID.x / 1280.0, 1.0);
}

Conclusion

GPGPU computing is an effective way to solve problems with large datasets, which is why you find it most often used in computer graphics, machine learning, and simulations. The SIMT architecture of the GPU requires that you conceptualize the solution to a problem in a different way, mapping the problem to a large number of groups of threads.

More Resources

• RDNA Whitepaper

• RDNA3 Instruction Set Architecture

• NVIDIA Volta Whitepaper

• NVIDIA CUDA C++ Programming Guide

• GPU Gems 2 - GPGPU Primer

• GPU Programming Primitives for Computer Graphics - Siggraph Asia 2024

• Lou Kramer (Technology Developer Engineer @ AMD) released a talk titled Compute Shaders.

• Intel Implicit SPMD Program Compiler