WebGPU: linear_indexing feature

at a glance

why it exists

WebGPU compute shaders dispatch work in a 3D grid of workgroups, each containing a 3D grid of invocations. To process a flat buffer — the common case for data-parallel tasks — developers had to compute a linear index from three separate coordinates every time:

// Old pattern — manual linear index calculation
let linear = global_id.z * (workgroup_size.x * workgroup_size.y * num_workgroups.x * num_workgroups.y)
           + global_id.y * (workgroup_size.x * num_workgroups.x)
           + global_id.x;

This is boilerplate that must be repeated in every compute shader that processes linear data, and it's easy to get wrong. linear_indexing moves this calculation into the language itself.

new built-in values

example

Feature detection and WGSL shader

// JavaScript: detect before using
if (!navigator.gpu.wgslLanguageFeatures.has('linear_indexing')) {
  throw new Error('linear_indexing WGSL extension not available');
}

const shaderCode = `
  requires linear_indexing;

  @group(0) @binding(0) var<storage, read_write> data: array<f32>;

  @compute @workgroup_size(64)
  fn main(@builtin(global_invocation_index) idx: u32) {
    if (idx < arrayLength(&data)) {
      data[idx] = data[idx] * 2.0;
    }
  }
`;

const shaderModule = device.createShaderModule({ code: shaderCode });
// ... create pipeline and dispatch as usual

workgroup_index example

requires linear_indexing;

var<workgroup> workgroup_data: array<f32, 64>;

@compute @workgroup_size(64)
fn accumulate(
  @builtin(workgroup_index) wg_idx: u32,
  @builtin(local_invocation_index) local_idx: u32,
) {
  // wg_idx identifies which workgroup we're in (flat index)
  // useful for indexing per-workgroup output buffers
  workgroup_data[local_idx] = f32(wg_idx);
}

browser support

see also

• W3C WebGPU Shading Language (WGSL) specification
• What's New in WebGPU (Chrome 147-148) — Chrome for Developers
• chromestatus — WebGPU: linear_indexing feature (5071243424432128)
• chrome-platform-showcase — live demo