Bind Groups

Bind Groups

Creating and Reusing Bind Groups

device.createBindGroup({ layout, entries }) binds real resources to a layout's slots, each entry a binding and a resource: a buffer (or { buffer, offset, size } for a window of one), a sampler, a texture view or an external texture. Bind groups are immutable, so build them once, outside the frame loop; changing a buffer's contents needs no new group. With hasDynamicOffset, one group serves many windows of a buffer, the offset passed at bind time:

One bind group, three sale windows through a dynamic offsetHTMLLive

<script type="module">
const device = await (await navigator.gpu.requestAdapter()).requestDevice();
const { UNIFORM, STORAGE, COPY_DST, COPY_SRC, MAP_READ } = GPUBufferUsage;
const module = device.createShaderModule({ code: /* wgsl */ `
  @group(0) @binding(0) var<uniform> percentOff: f32;
  @group(0) @binding(1) var<storage, read_write> out: array<f32, 3>;
  @compute @workgroup_size(1) fn main() {
    out[u32(percentOff) / 10 - 1] = 39.50 * (1 - percentOff / 100);   // book 2 on sale
  }` });
const layout = device.createBindGroupLayout({ entries: [
  { binding: 0, visibility: GPUShaderStage.COMPUTE, buffer: { hasDynamicOffset: true } },
  { binding: 1, visibility: GPUShaderStage.COMPUTE, buffer: { type: 'storage' } }] });
const pipeline = device.createComputePipeline({ compute: { module },
  layout: device.createPipelineLayout({ bindGroupLayouts: [layout] }) });
const sales = device.createBuffer({ size: 3 * 256, usage: UNIFORM | COPY_DST });
for (const [i, percent] of [10, 20, 30].entries()) {        // 256 bytes apart
  device.queue.writeBuffer(sales, i * 256, new Float32Array([percent]));
}
const out = device.createBuffer({ size: 12, usage: STORAGE | COPY_SRC });
const read = device.createBuffer({ size: 12, usage: COPY_DST | MAP_READ });
const group = device.createBindGroup({ layout, entries: [   // one group, three windows
  { binding: 0, resource: { buffer: sales, size: 4 } }, { binding: 1, resource: out }] });
const encoder = device.createCommandEncoder(), pass = encoder.beginComputePass();
pass.setPipeline(pipeline);
for (const offset of [0, 256, 512]) {
  pass.setBindGroup(0, group, [offset]), pass.dispatchWorkgroups(1);
}
pass.end(), encoder.copyBufferToBuffer(out, read), device.queue.submit([encoder.finish()]);
await read.mapAsync(GPUMapMode.READ);
console.log([...new Float32Array(read.getMappedRange())].map((v) => v.toFixed(2)).join(', '));
</script>

35.55, 31.60, 27.65

Three dispatches read 10, 20 and 30 percent off "Patterns of the Deep Web" through one group; an offset of 100 failed with "Dynamic Offset[0] (100) is not 256 byte aligned". Uniforms and Instancing uses this for per-draw uniforms.

MENU