Thread Coordinates

When launching a compute shader, execution is divided into a three-level hierarchy: Invocations (Threads), Workgroups, and the Grid.

Understanding the coordinates of each executing thread is critical for GPGPU programming. WGSL provides several built-in input variables to locate an invocation within this hierarchy.

Thread Coordinates Built-ins

Every execution thread has access to the following built-ins in its entry point:

1. @builtin(global_invocation_id) (vec3<u32>): The unique 3D coordinate of the current thread across the entire dispatch grid.
2. @builtin(local_invocation_id) (vec3<u32>): The 3D coordinate of the thread within its parent workgroup.
3. @builtin(workgroup_id) (vec3<u32>): The 3D coordinate of the workgroup within the dispatch grid.
4. @builtin(local_invocation_index) (u32): A linearized (1D) index of the thread within its workgroup, starting at 0 and ending at (size_x * size_y * size_z) - 1.

Mathematical Coordination Mapping

The relationship between these coordinates is given by:

\[ \text{global_invocation_id} = \text{workgroup_id} \times \text{workgroup_size} + \text{local_invocation_id} \]

Or in WGSL component form:

let global_x = workgroup_id.x * workgroup_size_x + local_id.x;

Similarly, the linearized index local_invocation_index is computed from the 3D local coordinate as:

\[ \text{local_invocation_index} = \text{local_id.z} \times (\text{size_x} \times \text{size_y}) + \text{local_id.y} \times \text{size_x} + \text{local_id.x} \]

Inspecting Coordinates

In the shader code on the right, we capture these coordinate values. The visualizer evaluates them to show the layout mappings!

/*
 * Copyright ©2026 Michael R. Bernstein. Licensed under Apache 2.0.
 * See root README.md for global project-wide upstream attributions.
 */

struct Coords {
  global_id: vec3<u32>,
  local_id: vec3<u32>,
  workgroup_id: vec3<u32>,
  local_idx: u32,
}

// We demonstrate a mock dispatch with:
//   workgroup_size = (8, 8, 1)
//   workgroup_id   = (2, 3, 0)
//   local_id       = (4, 5, 0)
//
// Let's compute the thread coordinates mathematically:
//   global_id = workgroup_id * workgroup_size + local_id
//             = (2, 3, 0) * (8, 8, 1) + (4, 5, 0)
//             = (16, 24, 0) + (4, 5, 0)
//             = (20, 29, 0)
//
//   local_idx = local_id.z * (size_x * size_y) + local_id.y * size_x + local_id.x
//             = 0 * 64 + 5 * 8 + 4
//             = 44
const computed_coords = Coords(
  vec3<u32>(20u, 29u, 0u), // global_invocation_id
  vec3<u32>(4u, 5u, 0u),   // local_invocation_id
  vec3<u32>(2u, 3u, 0u),   // workgroup_id
  44u                      // local_invocation_index
);

// This helper function shows how thread coordinates are mathematically related in WGSL.
fn get_global_id(w_id: vec3<u32>, w_size: vec3<u32>, l_id: vec3<u32>) -> vec3<u32> {
  return w_id * w_size + l_id;
}

fn get_local_index(l_id: vec3<u32>, w_size: vec3<u32>) -> u32 {
  return l_id.z * (w_size.x * w_size.y) + l_id.y * w_size.x + l_id.x;
}