26178691674

Committed 20 May 2026 05:24PM UTC coverage: 74.942% (+0.09%) from 74.85%

Build # 26178691674

Build Type

push

github

Committed by

web-flow

Commit Message

feat(gpgpu): Consolidate arithmetic operations (#2630)

Coverage Stats

7494 of 11282 branches covered (66.42%)

Branch coverage included in aggregate %.

281 of 344 new or added lines in 30 files covered. (81.69%)

26 existing lines in 5 files now uncovered.

16166 of 20289 relevant lines covered (79.68%)

1206.12 hits per line

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

69.0

/modules/gpgpu/src/operations/webgpu/common/row-transform.ts

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

import {Buffer, SignedDataType} from '@luma.gl/core';
import {Computation} from '@luma.gl/engine';
import {ShaderModule} from '@luma.gl/shadertools';
import {GPUTableEvaluator} from '../../../operation/gpu-table';
import {getLiteralValue, getWGSLType, getZeroValue} from './helper';

const WORKGROUP_SIZE = 64;
const GPGPU_OPERATION_STATS = 'GPGPU Operation Counts';
const COMPUTATION_RUNS = 'Computation Runs';

export function runRowComputation({
  module,
  elementWise = false,
  expression,
  inputs,
  output,
  operationType = output.type,
  outputBuffer
}: {
  module: ShaderModule;
  elementWise?: boolean;
  expression?: (laneIndex: number) => string;
  inputs: {[name: string]: GPUTableEvaluator};
  output: GPUTableEvaluator;
  operationType?: SignedDataType;
  outputBuffer: Buffer;
}): void {
  if (!module.source) {
    throw new Error(`WebGPU computation ${module.name} requires WGSL source`);
  }

  const bindings = Object.entries(inputs).map(([name, input]) => ({name, input}));
  const storageBindings = bindings
    .filter(({input}) => !input.isConstant)
    .map((binding, index) => ({...binding, index}));
  const castToType = getWGSLType(operationType);
  const outputType = getWGSLType(output.type);
  const defines: Record<string, string> = {
    TYPE: castToType,
    RESULT_LEN: output.size.toString()
  };

  for (const name in inputs) {
    const input = inputs[name];
    defines[`${name.toUpperCase()}_LEN`] = input.size.toString();
  }

  const source = /* wgsl */ `
${preprocess(module.source, defines)}
${storageBindings.map(({name, input, index}) => getInputBinding(name, input, index)).join('\n')}
${bindings.map(({name, input}) => getInputAccessor(name, input, operationType)).join('\n')}
${getOutputBinding(output, storageBindings.length)}
${getOutputWriter(output)}

@compute @workgroup_size(${WORKGROUP_SIZE}) fn main(
  @builtin(global_invocation_id) id: vec3<u32>
) {
  let rowIndex = id.x;
  if (rowIndex >= ${output.length}u) {
    return;
  }

${bindings.map(({name}) => `  let ${name} = read_${name}(rowIndex);`).join('\n')}
  var result: array<${outputType}, ${output.size}>;
${getComputeBlock(module.name, inputs, output, elementWise, expression)}
  write_result(rowIndex, result);
}
`;

  const computation = new Computation(outputBuffer.device, {
    source,
    shaderLayout: {
      bindings: [
        ...storageBindings.map(({name}, index) => ({
          name,
          type: 'storage' as const,
          group: 0,
          location: index
        })),
        {name: 'result', type: 'storage' as const, group: 0, location: storageBindings.length}
      ]
    }
  });

  const computationBindings: Record<string, Buffer> = Object.fromEntries(
    storageBindings.map(({name, input}) => [name, input.buffer])
  );
  computationBindings['result'] = outputBuffer;
  computation.setBindings(computationBindings);

  const computePass = outputBuffer.device.beginComputePass({});
  outputBuffer.device.statsManager
    .getStats(GPGPU_OPERATION_STATS)
    .get(COMPUTATION_RUNS)
    .incrementCount();
  computation.dispatch(computePass, Math.ceil(output.length / WORKGROUP_SIZE));
  computePass.end();
  outputBuffer.device.submit();
  computation.destroy();
}

function getInputBinding(name: string, input: GPUTableEvaluator, index: number): string {
  if (input.isConstant) {
    return '';
  }
  const inputType = getWGSLType(input.type);
  return `@group(0) @binding(${index}) var<storage, read> ${name}: array<${inputType}>;`;
}

function getInputAccessor(name: string, input: GPUTableEvaluator, asType: SignedDataType): string {
  const type = getWGSLType(asType);
  const castToType = input.type === asType ? '' : type;
  const stride = input.stride / input.ValueType.BYTES_PER_ELEMENT;
  const offset = input.offset / input.ValueType.BYTES_PER_ELEMENT;

  if (input.isConstant) {
    return `fn read_${name}(_rowIndex: u32) -> array<${type}, ${input.size}> {
  return array<${type}, ${input.size}>(${getConstantValues(input, castToType)});
}`;
  }

  return `fn read_${name}(rowIndex: u32) -> array<${type}, ${input.size}> {
  var value: array<${type}, ${input.size}>;
  let rowOffset = ${offset}u + rowIndex * ${stride}u;
${Array.from({length: input.size}, (_, elementIndex) =>
  castToType
    ? `  value[${elementIndex}] = ${castToType}(${name}[rowOffset + ${elementIndex}u]);`
    : `  value[${elementIndex}] = ${name}[rowOffset + ${elementIndex}u];`
).join('\n')}
  return value;
}`;
}

function getOutputBinding(output: GPUTableEvaluator, bindingIndex: number): string {
  const type = getWGSLType(output.type);
  return `@group(0) @binding(${bindingIndex}) var<storage, read_write> result: array<${type}>;`;
}

function getOutputWriter(output: GPUTableEvaluator): string {
  const stride = output.stride / output.ValueType.BYTES_PER_ELEMENT;
  const offset = output.offset / output.ValueType.BYTES_PER_ELEMENT;
  const type = getWGSLType(output.type);
  return `fn write_result(rowIndex: u32, value: array<${type}, ${output.size}>) {
  let rowOffset = ${offset}u + rowIndex * ${stride}u;
${Array.from({length: output.size}, (_, elementIndex) => `  result[rowOffset + ${elementIndex}u] = value[${elementIndex}];`).join('\n')}
}`;
}

function getComputeBlock(
  operationName: string,
  inputs: {[name: string]: GPUTableEvaluator},
  output: GPUTableEvaluator,
  elementWise: boolean,
  expression?: (laneIndex: number) => string
): string {
  let result = '';

  if (expression) {
    for (let elementIndex = 0; elementIndex < output.size; elementIndex++) {
      result += `  result[${elementIndex}] = ${expression(elementIndex)};\n`;
    }
  } else if (elementWise) {
    const zero = getZeroValue(output.type);
    const outputType = getWGSLType(output.type);

    for (let elementIndex = 0; elementIndex < output.size; elementIndex++) {
      const elementInputs = Object.keys(inputs).map(name => {
        if (elementIndex < inputs[name].size) {
          const inputType = getWGSLType(inputs[name].type);
          if (inputType === outputType) {
            return `${name}[${elementIndex}]`;
          }
          return `${outputType}(${name}[${elementIndex}])`;
        }
        return zero;
      });
      result += `  result[${elementIndex}] = ${operationName}(${elementInputs.join(', ')});\n`;
    }
  } else {
    result += `result = ${operationName}(${Object.keys(inputs).join(', ')});`;
  }

  return result.trimEnd();
}

function getConstantValues(input: GPUTableEvaluator, asType: string): string {
  const values = input.value;
  if (!values) {
    throw new Error(`Constant input ${input} is missing CPU values`);
  }
  return Array.from({length: input.size}, (_, index) =>
    getLiteralValue(asType, values[index] ?? 0)
  ).join(', ');
}

function preprocess(source: string, defines: Record<string, string>) {
  for (const key in defines) {
    source = source.replaceAll(`{${key}}`, defines[key]);
  }
  return source;
}

1	// luma.gl
2	// SPDX-License-Identifier: MIT
3	// Copyright (c) vis.gl contributors
4
5	import {Buffer, SignedDataType} from '@luma.gl/core';
6	import {Computation} from '@luma.gl/engine';
7	import {ShaderModule} from '@luma.gl/shadertools';
8	import {GPUTableEvaluator} from '../../../operation/gpu-table';
9	import {getLiteralValue, getWGSLType, getZeroValue} from './helper';
10
11	const WORKGROUP_SIZE = 64;	11✔
12	const GPGPU_OPERATION_STATS = 'GPGPU Operation Counts';	11✔
13	const COMPUTATION_RUNS = 'Computation Runs';	11✔
14
15	export function runRowComputation({
16	module,
17	elementWise = false,	27✔
18	expression,
19	inputs,
20	output,
21	operationType = output.type,	27✔
22	outputBuffer
23	}: {
24	module: ShaderModule;
25	elementWise?: boolean;
26	expression?: (laneIndex: number) => string;
27	inputs: {[name: string]: GPUTableEvaluator};
28	output: GPUTableEvaluator;
29	operationType?: SignedDataType;
30	outputBuffer: Buffer;
31	}): void {
32	if (!module.source) {	27!
33	throw new Error(`WebGPU computation ${module.name} requires WGSL source`);	×
34	}
35
36	const bindings = Object.entries(inputs).map(([name, input]) => ({name, input}));	46✔
37	const storageBindings = bindings	27✔
38	.filter(({input}) => !input.isConstant)	46✔
39	.map((binding, index) => ({...binding, index}));	46✔
40	const castToType = getWGSLType(operationType);	27✔
41	const outputType = getWGSLType(output.type);	27✔
42	const defines: Record<string, string> = {	27✔
43	TYPE: castToType,
44	RESULT_LEN: output.size.toString()
45	};
46
47	for (const name in inputs) {	27✔
48	const input = inputs[name];	46✔
49	defines[`${name.toUpperCase()}_LEN`] = input.size.toString();	46✔
50	}
51
52	const source = /* wgsl */ `	27✔
53	${preprocess(module.source, defines)}
54	${storageBindings.map(({name, input, index}) => getInputBinding(name, input, index)).join('\n')}	46✔
55	${bindings.map(({name, input}) => getInputAccessor(name, input, operationType)).join('\n')}	46✔
56	${getOutputBinding(output, storageBindings.length)}
57	${getOutputWriter(output)}
58
59	@compute @workgroup_size(${WORKGROUP_SIZE}) fn main(
60	@builtin(global_invocation_id) id: vec3<u32>
61	) {
62	let rowIndex = id.x;
63	if (rowIndex >= ${output.length}u) {
64	return;
65	}
66
67	${bindings.map(({name}) => ` let ${name} = read_${name}(rowIndex);`).join('\n')}	46✔
68	var result: array<${outputType}, ${output.size}>;
69	${getComputeBlock(module.name, inputs, output, elementWise, expression)}
70	write_result(rowIndex, result);
71	}
72	`;
73
74	const computation = new Computation(outputBuffer.device, {	27✔
75	source,
76	shaderLayout: {
77	bindings: [
78	...storageBindings.map(({name}, index) => ({	46✔
79	name,
80	type: 'storage' as const,
81	group: 0,
82	location: index
83	})),
84	{name: 'result', type: 'storage' as const, group: 0, location: storageBindings.length}
85	]
86	}
87	});
88
89	const computationBindings: Record<string, Buffer> = Object.fromEntries(	27✔
90	storageBindings.map(({name, input}) => [name, input.buffer])	46✔
91	);
92	computationBindings['result'] = outputBuffer;	27✔
93	computation.setBindings(computationBindings);	27✔
94
95	const computePass = outputBuffer.device.beginComputePass({});	27✔
96	outputBuffer.device.statsManager	27✔
97	.getStats(GPGPU_OPERATION_STATS)
98	.get(COMPUTATION_RUNS)
99	.incrementCount();
100	computation.dispatch(computePass, Math.ceil(output.length / WORKGROUP_SIZE));	27✔
101	computePass.end();	27✔
102	outputBuffer.device.submit();	27✔
103	computation.destroy();	27✔
104	}
105
106	function getInputBinding(name: string, input: GPUTableEvaluator, index: number): string {
107	if (input.isConstant) {	46!
108	return '';	×
109	}
110	const inputType = getWGSLType(input.type);	46✔
111	return `@group(0) @binding(${index}) var<storage, read> ${name}: array<${inputType}>;`;	46✔
112	}
113
114	function getInputAccessor(name: string, input: GPUTableEvaluator, asType: SignedDataType): string {
115	const type = getWGSLType(asType);	46✔
116	const castToType = input.type === asType ? '' : type;	46✔
117	const stride = input.stride / input.ValueType.BYTES_PER_ELEMENT;	46✔
118	const offset = input.offset / input.ValueType.BYTES_PER_ELEMENT;	46✔
119
120	if (input.isConstant) {	46!
UNCOV 121	return `fn read_${name}(_rowIndex: u32) -> array<${type}, ${input.size}> {	×
122	return array<${type}, ${input.size}>(${getConstantValues(input, castToType)});
123	}`;
124	}
125
126	return `fn read_${name}(rowIndex: u32) -> array<${type}, ${input.size}> {	46✔
127	var value: array<${type}, ${input.size}>;
128	let rowOffset = ${offset}u + rowIndex * ${stride}u;
129	${Array.from({length: input.size}, (_, elementIndex) =>
130	castToType	137✔
131	? ` value[${elementIndex}] = ${castToType}(${name}[rowOffset + ${elementIndex}u]);`
132	: ` value[${elementIndex}] = ${name}[rowOffset + ${elementIndex}u];`
133	).join('\n')}
134	return value;
135	}`;
136	}
137
138	function getOutputBinding(output: GPUTableEvaluator, bindingIndex: number): string {
139	const type = getWGSLType(output.type);	27✔
140	return `@group(0) @binding(${bindingIndex}) var<storage, read_write> result: array<${type}>;`;	27✔
141	}
142
143	function getOutputWriter(output: GPUTableEvaluator): string {
144	const stride = output.stride / output.ValueType.BYTES_PER_ELEMENT;	27✔
145	const offset = output.offset / output.ValueType.BYTES_PER_ELEMENT;	27✔
146	const type = getWGSLType(output.type);	27✔
147	return `fn write_result(rowIndex: u32, value: array<${type}, ${output.size}>) {	27✔
148	let rowOffset = ${offset}u + rowIndex * ${stride}u;
149	${Array.from({length: output.size}, (_, elementIndex) => ` result[rowOffset + ${elementIndex}u] = value[${elementIndex}];`).join('\n')}	94✔
150	}`;
151	}
152
153	function getComputeBlock(
154	operationName: string,
155	inputs: {[name: string]: GPUTableEvaluator},
156	output: GPUTableEvaluator,
157	elementWise: boolean,
158	expression?: (laneIndex: number) => string
159	): string {
160	let result = '';	27✔
161
162	if (expression) {	27✔
163	for (let elementIndex = 0; elementIndex < output.size; elementIndex++) {	19✔
164	result += ` result[${elementIndex}] = ${expression(elementIndex)};\n`;	61✔
165	}
166	} else if (elementWise) {	8!
UNCOV 167	const zero = getZeroValue(output.type);	×
UNCOV 168	const outputType = getWGSLType(output.type);	×
169
UNCOV 170	for (let elementIndex = 0; elementIndex < output.size; elementIndex++) {	×
UNCOV 171	const elementInputs = Object.keys(inputs).map(name => {	×
UNCOV 172	if (elementIndex < inputs[name].size) {	×
UNCOV 173	const inputType = getWGSLType(inputs[name].type);	×
UNCOV 174	if (inputType === outputType) {	×
UNCOV 175	return `${name}[${elementIndex}]`;	×
176	}
177	return `${outputType}(${name}[${elementIndex}])`;	×
178	}
UNCOV 179	return zero;	×
180	});
UNCOV 181	result += ` result[${elementIndex}] = ${operationName}(${elementInputs.join(', ')});\n`;	×
182	}
183	} else {
184	result += `result = ${operationName}(${Object.keys(inputs).join(', ')});`;	8✔
185	}
186
187	return result.trimEnd();	27✔
188	}
189
190	function getConstantValues(input: GPUTableEvaluator, asType: string): string {
UNCOV 191	const values = input.value;	×
UNCOV 192	if (!values) {	×
193	throw new Error(`Constant input ${input} is missing CPU values`);	×
194	}
UNCOV 195	return Array.from({length: input.size}, (_, index) =>	×
UNCOV 196	getLiteralValue(asType, values[index] ?? 0)	×
197	).join(', ');
198	}
199
200	function preprocess(source: string, defines: Record<string, string>) {
201	for (const key in defines) {	27✔
202	source = source.replaceAll(`{${key}}`, defines[key]);	100✔
203	}
204	return source;	27✔
205	}

visgl / luma.gl / 26178691674

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