28617818631

Committed 02 Jul 2026 07:57PM UTC coverage: 71.941% (+0.4%) from 71.566%

Build # 28617818631

Build Type

push

github

Committed by

web-flow

Commit Message

[codex] Document GPU picking and command-graph roadmap (#2711)

Coverage Stats

10762 of 16860 branches covered (63.83%)

Branch coverage included in aggregate %.

21074 of 27393 relevant lines covered (76.93%)

5474.44 hits per line

Source File
Press 'n' to go to next uncovered line, 'b' for previous

84.62

/modules/experimental/src/gpu-primitives/gpu-compaction.ts

// luma.gl
// SPDX-License-Identifier: MIT
// Copyright (c) vis.gl contributors

import {Computation} from '@luma.gl/engine';
import {GPUCommandGraph, type GraphBufferView} from './gpu-command-graph';
import {GPUScan} from './gpu-scan';
import {
  createTransientView,
  getViewBinding,
  getViewElementOffset,
  validatePackedUint32View
} from './graph-buffer-view-utils';

const COMPACTION_WORKGROUP_SIZE = 256;

export type GPUCompactionProps = {
  id?: string;
  input: GraphBufferView<'uint32'>;
  flags: GraphBufferView<'uint32'>;
  output: GraphBufferView<'uint32'>;
  count: GraphBufferView<'uint32'>;
};

/** Stable graph-composed compaction of uint32 values selected by 0/1 flags. */
export class GPUCompaction {
  readonly id: string;
  readonly input: GraphBufferView<'uint32'>;
  readonly flags: GraphBufferView<'uint32'>;
  readonly output: GraphBufferView<'uint32'>;
  readonly count: GraphBufferView<'uint32'>;

  constructor(props: GPUCompactionProps) {
    this.id = props.id ?? 'gpu-compaction';
    this.input = props.input;
    this.flags = props.flags;
    this.output = props.output;
    this.count = props.count;
    validatePackedUint32View(this.input, `${this.id} input`);
    validatePackedUint32View(this.flags, `${this.id} flags`);
    validatePackedUint32View(this.output, `${this.id} output`);
    validatePackedUint32View(this.count, `${this.id} count`);
    if (this.flags.length !== this.input.length) {
      throw new Error(`${this.id} flags.length must equal input.length`);
    }
    if (this.output.length < this.input.length) {
      throw new Error(`${this.id} output must contain at least input.length rows`);
    }
    if (this.count.length < 1) {
      throw new Error(`${this.id} count must contain one uint32 row`);
    }
  }

  /** Adds scan and scatter passes to the target graph. */
  addToGraph<Parameters>(graph: GPUCommandGraph<Parameters>): void {
    for (const view of [this.input, this.flags, this.output, this.count]) {
      if (view.buffer.graph !== graph) {
        throw new Error(`${this.id} views must belong to the target graph`);
      }
    }

    if (this.input.length === 0) {
      addClearCountPass(graph, this.id, this.count);
      return;
    }

    const offsets = createTransientView(graph, `${this.id}-offsets`, 'uint32', this.input.length);
    new GPUScan({
      id: `${this.id}-scan`,
      input: this.flags,
      output: offsets
    }).addToGraph(graph);
    addScatterPass(graph, {
      id: `${this.id}-scatter`,
      input: this.input,
      flags: this.flags,
      offsets,
      output: this.output,
      count: this.count
    });
  }
}

function addClearCountPass<Parameters>(
  graph: GPUCommandGraph<Parameters>,
  id: string,
  count: GraphBufferView<'uint32'>
): void {
  const passId = `${id}-clear-count`;
  graph.addComputePass({
    id: passId,
    resources: [{buffer: count, usage: 'storage-write'}],
    compile: ({device}) => {
      const computation = new Computation(device, {
        id: passId,
        source: `const COUNT_OFFSET: u32 = ${getViewElementOffset(count)}u;
@group(0) @binding(0) var<storage, read_write> outputCount: array<u32>;
@compute @workgroup_size(1) fn main() { outputCount[COUNT_OFFSET] = 0u; }`,
        shaderLayout: {
          bindings: [{name: 'outputCount', type: 'storage', group: 0, location: 0}]
        }
      });
      return {
        encode: ({computePass, getBuffer}) => {
          computation.setBindings({outputCount: getViewBinding(count, getBuffer)});
          computation.dispatch(computePass, 1);
        },
        destroy: () => computation.destroy()
      };
    }
  });
}

function addScatterPass<Parameters>(
  graph: GPUCommandGraph<Parameters>,
  props: {
    id: string;
    input: GraphBufferView<'uint32'>;
    flags: GraphBufferView<'uint32'>;
    offsets: GraphBufferView<'uint32'>;
    output: GraphBufferView<'uint32'>;
    count: GraphBufferView<'uint32'>;
  }
): void {
  const length = props.input.length;
  const source = /* wgsl */ `
const ELEMENT_COUNT: u32 = ${length}u;
const INPUT_OFFSET: u32 = ${getViewElementOffset(props.input)}u;
const FLAGS_OFFSET: u32 = ${getViewElementOffset(props.flags)}u;
const OFFSETS_OFFSET: u32 = ${getViewElementOffset(props.offsets)}u;
const OUTPUT_OFFSET: u32 = ${getViewElementOffset(props.output)}u;
const COUNT_OFFSET: u32 = ${getViewElementOffset(props.count)}u;
@group(0) @binding(0) var<storage, read> inputValues: array<u32>;
@group(0) @binding(1) var<storage, read> flags: array<u32>;
@group(0) @binding(2) var<storage, read> offsets: array<u32>;
@group(0) @binding(3) var<storage, read_write> outputValues: array<u32>;
@group(0) @binding(4) var<storage, read_write> outputCount: array<u32>;

@compute @workgroup_size(${COMPACTION_WORKGROUP_SIZE}) fn main(
  @builtin(global_invocation_id) globalId: vec3<u32>
) {
  let index = globalId.x;
  if (index >= ELEMENT_COUNT) { return; }
  let flag = min(flags[FLAGS_OFFSET + index], 1u);
  if (flag != 0u) {
    outputValues[OUTPUT_OFFSET + offsets[OFFSETS_OFFSET + index]] = inputValues[INPUT_OFFSET + index];
  }
  if (index == ELEMENT_COUNT - 1u) {
    outputCount[COUNT_OFFSET] = offsets[OFFSETS_OFFSET + index] + flag;
  }
}`;
  graph.addComputePass({
    id: props.id,
    resources: [
      {buffer: props.input, usage: 'storage-read'},
      {buffer: props.flags, usage: 'storage-read'},
      {buffer: props.offsets, usage: 'storage-read'},
      {buffer: props.output, usage: 'storage-write'},
      {buffer: props.count, usage: 'storage-write'}
    ],
    compile: ({device}) => {
      const computation = new Computation(device, {
        id: props.id,
        source,
        shaderLayout: {
          bindings: [
            {name: 'inputValues', type: 'storage', group: 0, location: 0},
            {name: 'flags', type: 'storage', group: 0, location: 1},
            {name: 'offsets', type: 'storage', group: 0, location: 2},
            {name: 'outputValues', type: 'storage', group: 0, location: 3},
            {name: 'outputCount', type: 'storage', group: 0, location: 4}
          ]
        }
      });
      return {
        encode: ({computePass, getBuffer}) => {
          computation.setBindings({
            inputValues: getViewBinding(props.input, getBuffer),
            flags: getViewBinding(props.flags, getBuffer),
            offsets: getViewBinding(props.offsets, getBuffer),
            outputValues: getViewBinding(props.output, getBuffer),
            outputCount: getViewBinding(props.count, getBuffer)
          });
          computation.dispatch(computePass, Math.ceil(length / COMPACTION_WORKGROUP_SIZE));
        },
        destroy: () => computation.destroy()
      };
    }
  });
}

1	// luma.gl
2	// SPDX-License-Identifier: MIT
3	// Copyright (c) vis.gl contributors
4
5	import {Computation} from '@luma.gl/engine';
6	import {GPUCommandGraph, type GraphBufferView} from './gpu-command-graph';
7	import {GPUScan} from './gpu-scan';
8	import {
9	createTransientView,
10	getViewBinding,
11	getViewElementOffset,
12	validatePackedUint32View
13	} from './graph-buffer-view-utils';
14
15	const COMPACTION_WORKGROUP_SIZE = 256;	11✔
16
17	export type GPUCompactionProps = {
18	id?: string;
19	input: GraphBufferView<'uint32'>;
20	flags: GraphBufferView<'uint32'>;
21	output: GraphBufferView<'uint32'>;
22	count: GraphBufferView<'uint32'>;
23	};
24
25	/** Stable graph-composed compaction of uint32 values selected by 0/1 flags. */
26	export class GPUCompaction {
27	readonly id: string;
28	readonly input: GraphBufferView<'uint32'>;
29	readonly flags: GraphBufferView<'uint32'>;
30	readonly output: GraphBufferView<'uint32'>;
31	readonly count: GraphBufferView<'uint32'>;
32
33	constructor(props: GPUCompactionProps) {
34	this.id = props.id ?? 'gpu-compaction';	6✔
35	this.input = props.input;	6✔
36	this.flags = props.flags;	6✔
37	this.output = props.output;	6✔
38	this.count = props.count;	6✔
39	validatePackedUint32View(this.input, `${this.id} input`);	6✔
40	validatePackedUint32View(this.flags, `${this.id} flags`);	6✔
41	validatePackedUint32View(this.output, `${this.id} output`);	6✔
42	validatePackedUint32View(this.count, `${this.id} count`);	6✔
43	if (this.flags.length !== this.input.length) {	6!
44	throw new Error(`${this.id} flags.length must equal input.length`);	×
45	}
46	if (this.output.length < this.input.length) {	6!
47	throw new Error(`${this.id} output must contain at least input.length rows`);	×
48	}
49	if (this.count.length < 1) {	6!
50	throw new Error(`${this.id} count must contain one uint32 row`);	×
51	}
52	}
53
54	/** Adds scan and scatter passes to the target graph. */
55	addToGraph<Parameters>(graph: GPUCommandGraph<Parameters>): void {
56	for (const view of [this.input, this.flags, this.output, this.count]) {	6✔
57	if (view.buffer.graph !== graph) {	24!
58	throw new Error(`${this.id} views must belong to the target graph`);	×
59	}
60	}
61
62	if (this.input.length === 0) {	6✔
63	addClearCountPass(graph, this.id, this.count);	1✔
64	return;	1✔
65	}
66
67	const offsets = createTransientView(graph, `${this.id}-offsets`, 'uint32', this.input.length);	5✔
68	new GPUScan({	5✔
69	id: `${this.id}-scan`,
70	input: this.flags,
71	output: offsets
72	}).addToGraph(graph);
73	addScatterPass(graph, {	5✔
74	id: `${this.id}-scatter`,
75	input: this.input,
76	flags: this.flags,
77	offsets,
78	output: this.output,
79	count: this.count
80	});
81	}
82	}
83
84	function addClearCountPass<Parameters>(
85	graph: GPUCommandGraph<Parameters>,
86	id: string,
87	count: GraphBufferView<'uint32'>
88	): void {
89	const passId = `${id}-clear-count`;	1✔
90	graph.addComputePass({	1✔
91	id: passId,
92	resources: [{buffer: count, usage: 'storage-write'}],
93	compile: ({device}) => {
94	const computation = new Computation(device, {	1✔
95	id: passId,
96	source: `const COUNT_OFFSET: u32 = ${getViewElementOffset(count)}u;
97	@group(0) @binding(0) var<storage, read_write> outputCount: array<u32>;
98	@compute @workgroup_size(1) fn main() { outputCount[COUNT_OFFSET] = 0u; }`,
99	shaderLayout: {
100	bindings: [{name: 'outputCount', type: 'storage', group: 0, location: 0}]
101	}
102	});
103	return {	1✔
104	encode: ({computePass, getBuffer}) => {
105	computation.setBindings({outputCount: getViewBinding(count, getBuffer)});	1✔
106	computation.dispatch(computePass, 1);	1✔
107	},
108	destroy: () => computation.destroy()	1✔
109	};
110	}
111	});
112	}
113
114	function addScatterPass<Parameters>(
115	graph: GPUCommandGraph<Parameters>,
116	props: {
117	id: string;
118	input: GraphBufferView<'uint32'>;
119	flags: GraphBufferView<'uint32'>;
120	offsets: GraphBufferView<'uint32'>;
121	output: GraphBufferView<'uint32'>;
122	count: GraphBufferView<'uint32'>;
123	}
124	): void {
125	const length = props.input.length;	5✔
126	const source = /* wgsl */ `	5✔
127	const ELEMENT_COUNT: u32 = ${length}u;
128	const INPUT_OFFSET: u32 = ${getViewElementOffset(props.input)}u;
129	const FLAGS_OFFSET: u32 = ${getViewElementOffset(props.flags)}u;
130	const OFFSETS_OFFSET: u32 = ${getViewElementOffset(props.offsets)}u;
131	const OUTPUT_OFFSET: u32 = ${getViewElementOffset(props.output)}u;
132	const COUNT_OFFSET: u32 = ${getViewElementOffset(props.count)}u;
133	@group(0) @binding(0) var<storage, read> inputValues: array<u32>;
134	@group(0) @binding(1) var<storage, read> flags: array<u32>;
135	@group(0) @binding(2) var<storage, read> offsets: array<u32>;
136	@group(0) @binding(3) var<storage, read_write> outputValues: array<u32>;
137	@group(0) @binding(4) var<storage, read_write> outputCount: array<u32>;
138
139	@compute @workgroup_size(${COMPACTION_WORKGROUP_SIZE}) fn main(
140	@builtin(global_invocation_id) globalId: vec3<u32>
141	) {
142	let index = globalId.x;
143	if (index >= ELEMENT_COUNT) { return; }
144	let flag = min(flags[FLAGS_OFFSET + index], 1u);
145	if (flag != 0u) {
146	outputValues[OUTPUT_OFFSET + offsets[OFFSETS_OFFSET + index]] = inputValues[INPUT_OFFSET + index];
147	}
148	if (index == ELEMENT_COUNT - 1u) {
149	outputCount[COUNT_OFFSET] = offsets[OFFSETS_OFFSET + index] + flag;
150	}
151	}`;
152	graph.addComputePass({	5✔
153	id: props.id,
154	resources: [
155	{buffer: props.input, usage: 'storage-read'},
156	{buffer: props.flags, usage: 'storage-read'},
157	{buffer: props.offsets, usage: 'storage-read'},
158	{buffer: props.output, usage: 'storage-write'},
159	{buffer: props.count, usage: 'storage-write'}
160	],
161	compile: ({device}) => {
162	const computation = new Computation(device, {	5✔
163	id: props.id,
164	source,
165	shaderLayout: {
166	bindings: [
167	{name: 'inputValues', type: 'storage', group: 0, location: 0},
168	{name: 'flags', type: 'storage', group: 0, location: 1},
169	{name: 'offsets', type: 'storage', group: 0, location: 2},
170	{name: 'outputValues', type: 'storage', group: 0, location: 3},
171	{name: 'outputCount', type: 'storage', group: 0, location: 4}
172	]
173	}
174	});
175	return {	5✔
176	encode: ({computePass, getBuffer}) => {
177	computation.setBindings({	5✔
178	inputValues: getViewBinding(props.input, getBuffer),
179	flags: getViewBinding(props.flags, getBuffer),
180	offsets: getViewBinding(props.offsets, getBuffer),
181	outputValues: getViewBinding(props.output, getBuffer),
182	outputCount: getViewBinding(props.count, getBuffer)
183	});
184	computation.dispatch(computePass, Math.ceil(length / COMPACTION_WORKGROUP_SIZE));	5✔
185	},
186	destroy: () => computation.destroy()	5✔
187	};
188	}
189	});
190	}

visgl / luma.gl / 28617818631

Source File Press 'n' to go to next uncovered line, 'b' for previous

Source File
Press 'n' to go to next uncovered line, 'b' for previous