28670509643

Committed 03 Jul 2026 03:41PM UTC coverage: 72.312% (+0.02%) from 72.29%

Build # 28670509643

Build Type

Pull #2723

github

Committed by

web-flow

Commit Message

Merge da8b95de7 into aeb41a2cd

Pull Request Pull Request #2723: feat(experimental) Add multi-chunk GPU histogram

Coverage Stats

10852 of 16893 branches covered (64.24%)

Branch coverage included in aggregate %.

12 of 13 new or added lines in 1 file covered. (92.31%)

1 existing line in 1 file now uncovered.

21219 of 27458 relevant lines covered (77.28%)

5661.26 hits per line

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

87.93

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

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

import {type Binding} from '@luma.gl/core';
import {Computation} from '@luma.gl/engine';
import {
  GPUCommandGraph,
  GraphVectorView,
  type GraphBufferUse,
  type GraphDataView
} from './gpu-command-graph';
import {
  createTransientView,
  getViewBinding,
  getViewElementOffset,
  type GPUScalarFormat,
  validatePackedUint32View,
  validatePackedView
} from './graph-data-view-utils';
import {GPUReduction} from './gpu-reduction';

const HISTOGRAM_WORKGROUP_SIZE = 256;
const MAXIMUM_LOCAL_BIN_COUNT = 256;
const SCALAR_FORMATS = ['uint32', 'sint32', 'float32'] as const;

/** Domain accepted by {@link GPUHistogram}. */
export type GPUHistogramDomain<T extends GPUScalarFormat = GPUScalarFormat> =
  | readonly [number, number]
  | GraphDataView<T>
  | 'auto';

/** One scalar chunk or an ordered scalar vector counted by {@link GPUHistogram}. */
export type GPUHistogramInput<T extends GPUScalarFormat = GPUScalarFormat> =
  | GraphDataView<T>
  | GraphVectorView<T>;

/** Properties for graph-native scalar histogram counting. */
export type GPUHistogramProps<T extends GPUScalarFormat = GPUScalarFormat> = {
  id?: string;
  input: GPUHistogramInput<T>;
  output: GraphDataView<'uint32'>;
  domain: GPUHistogramDomain<T>;
};

/** Graph-native histogram counting for packed 32-bit scalar values. */
export class GPUHistogram<T extends GPUScalarFormat = GPUScalarFormat> {
  readonly id: string;
  readonly input: GPUHistogramInput<T>;
  readonly output: GraphDataView<'uint32'>;
  readonly domain: GPUHistogramDomain<T>;

  constructor(props: GPUHistogramProps<T>) {
    this.id = props.id ?? 'gpu-histogram';
    this.input = props.input;
    this.output = props.output;
    this.domain = props.domain;
    for (const [chunkIndex, input] of getHistogramInputs(this.input).entries()) {
      const name = this.input instanceof GraphVectorView ? ` input chunk ${chunkIndex}` : ' input';
      validatePackedView(input, SCALAR_FORMATS, `${this.id}${name}`);
      if (input.format !== this.input.format) {
        throw new Error(`${this.id}${name} format must match the input format`);
      }
    }
    validatePackedUint32View(this.output, `${this.id} output`);
    if (this.output.length === 0) {
      throw new Error(`${this.id} output must contain at least one bin`);
    }
    if (getHistogramInputs(this.input).some(input => input.buffer === this.output.buffer)) {
      throw new Error(`${this.id} input and output must use separate buffers`);
    }
    if (isGPUHistogramDomainView(this.domain)) {
      validatePackedView(this.domain, SCALAR_FORMATS, `${this.id} domain`);
      if (this.domain.format !== this.input.format || this.domain.length !== 2) {
        throw new Error(`${this.id} GPU domain must contain two ${this.input.format} rows`);
      }
      if (this.domain.buffer === this.output.buffer) {
        throw new Error(`${this.id} domain and output must use separate buffers`);
      }
    } else if (Array.isArray(this.domain)) {
      validateLiteralDomain(this.domain, this.input.format, this.id);
    }
  }

  /** Adds domain inference, output clearing, and accumulation nodes to a graph. */
  addToGraph<Parameters>(graph: GPUCommandGraph<Parameters>): void {
    const inputs = getHistogramInputs(this.input);
    if (inputs.some(input => input.buffer.graph !== graph) || this.output.buffer.graph !== graph) {
      throw new Error(`${this.id} views must belong to the target graph`);
    }
    let domain = this.domain;
    if (isGPUHistogramDomainView(domain) && domain.buffer.graph !== graph) {
      throw new Error(`${this.id} domain must belong to the target graph`);
    }
    if (domain === 'auto') {
      const inferredDomain = createTransientView(
        graph,
        `${this.id}-auto-domain`,
        this.input.format,
        2
      ) as GraphDataView<T>;
      new GPUReduction({
        id: `${this.id}-extent`,
        input: this.input,
        output: inferredDomain,
        operation: 'extent'
      }).addToGraph(graph);
      domain = inferredDomain;
    }
    addClearHistogramPass(graph, this.id, this.output);
    const accumulationPath = this.output.length <= MAXIMUM_LOCAL_BIN_COUNT ? 'local' : 'global';
    inputs.forEach((input, chunkIndex) => {
      if (input.length === 0) {
        return;
      }
      addHistogramPass(graph, {
        id:
          this.input instanceof GraphVectorView
            ? `${this.id}-chunk-${chunkIndex}-${accumulationPath}`
            : `${this.id}-${accumulationPath}`,
        input,
        output: this.output,
        domain
      });
    });
  }
}

function getHistogramInputs<T extends GPUScalarFormat>(
  input: GPUHistogramInput<T>
): readonly GraphDataView<T>[] {
  return input instanceof GraphVectorView ? input.data : [input];
}

function addClearHistogramPass<Parameters>(
  graph: GPUCommandGraph<Parameters>,
  id: string,
  output: GraphDataView<'uint32'>
): void {
  const passId = `${id}-clear`;
  const source = /* wgsl */ `
const BIN_COUNT: u32 = ${output.length}u;
const OUTPUT_OFFSET: u32 = ${getViewElementOffset(output)}u;
@group(0) @binding(0) var<storage, read_write> outputCounts: array<atomic<u32>>;
@compute @workgroup_size(${HISTOGRAM_WORKGROUP_SIZE}) fn main(
  @builtin(global_invocation_id) globalId: vec3<u32>
) {
  if (globalId.x < BIN_COUNT) { atomicStore(&outputCounts[OUTPUT_OFFSET + globalId.x], 0u); }
}`;
  addComputationPass(graph, {
    id: passId,
    source,
    resources: [{buffer: output, usage: 'storage-write'}],
    bindings: {outputCounts: output},
    dispatchCount: Math.ceil(output.length / HISTOGRAM_WORKGROUP_SIZE)
  });
}

function addHistogramPass<Parameters, T extends GPUScalarFormat>(
  graph: GPUCommandGraph<Parameters>,
  props: {
    id: string;
    input: GraphDataView<T>;
    output: GraphDataView<'uint32'>;
    domain: readonly [number, number] | GraphDataView<T>;
  }
): void {
  const local = props.output.length <= MAXIMUM_LOCAL_BIN_COUNT;
  const shaderType = getShaderType(props.input.format);
  const gpuDomain = isGPUHistogramDomainView(props.domain);
  const literalDomain = props.domain as readonly [number, number];
  const domainBinding = gpuDomain
    ? '@group(0) @binding(1) var<storage, read> domainValues: array<' + shaderType + '>;'
    : '';
  const outputBinding = gpuDomain ? 2 : 1;
  const domainInitialization = gpuDomain
    ? `let minimum: ${shaderType} = domainValues[DOMAIN_OFFSET];
  let maximum: ${shaderType} = domainValues[DOMAIN_OFFSET + 1u];`
    : `let minimum: ${shaderType} = ${getLiteral(literalDomain[0], props.input.format)};
  let maximum: ${shaderType} = ${getLiteral(literalDomain[1], props.input.format)};`;
  const finiteCondition =
    props.input.format === 'float32' ? 'value == value && abs(value) <= 3.402823466e+38' : 'true';
  const integerMultiplierBitCount = 32 - Math.clz32(props.output.length);
  const integerBinningFunction =
    props.input.format === 'float32'
      ? ''
      : `fn multiplyDivideFloor(numerator: u32, multiplier: u32, denominator: u32) -> u32 {
  var quotient = 0u;
  var remainder = 0u;
  var bitIndex = ${integerMultiplierBitCount}u;
  loop {
    bitIndex = bitIndex - 1u;
    // Double and add modulo denominator without overflowing u32.
    let doubledThreshold = denominator - remainder;
    if (remainder >= doubledThreshold) {
      remainder = remainder - doubledThreshold;
      quotient = quotient * 2u + 1u;
    } else {
      remainder = remainder * 2u;
      quotient = quotient * 2u;
    }
    if (((multiplier >> bitIndex) & 1u) != 0u) {
      let additionThreshold = denominator - numerator;
      if (remainder >= additionThreshold) {
        remainder = remainder - additionThreshold;
        quotient = quotient + 1u;
      } else {
        remainder = remainder + numerator;
      }
    }
    if (bitIndex == 0u) { break; }
  }
  return quotient;
}`;
  const binCalculation =
    props.input.format === 'float32'
      ? `let ratio = (value - minimum) / (maximum - minimum);
          binIndex = min(u32(ratio * f32(BIN_COUNT)), BIN_COUNT - 1u);`
      : props.input.format === 'uint32'
        ? 'binIndex = multiplyDivideFloor(value - minimum, BIN_COUNT, maximum - minimum);'
        : `let orderedValue = bitcast<u32>(value) ^ 0x80000000u;
          let orderedMinimum = bitcast<u32>(minimum) ^ 0x80000000u;
          let orderedMaximum = bitcast<u32>(maximum) ^ 0x80000000u;
          binIndex = multiplyDivideFloor(
            orderedValue - orderedMinimum,
            BIN_COUNT,
            orderedMaximum - orderedMinimum
          );`;
  const accumulation = local
    ? `if (accepted) { atomicAdd(&localCounts[binIndex], 1u); }
  workgroupBarrier();
  if (lane < BIN_COUNT) {
    atomicAdd(&outputCounts[OUTPUT_OFFSET + lane], atomicLoad(&localCounts[lane]));
  }`
    : 'if (accepted) { atomicAdd(&outputCounts[OUTPUT_OFFSET + binIndex], 1u); }';
  const source = /* wgsl */ `
const ELEMENT_COUNT: u32 = ${props.input.length}u;
const BIN_COUNT: u32 = ${props.output.length}u;
const INPUT_OFFSET: u32 = ${getViewElementOffset(props.input)}u;
${gpuDomain ? `const DOMAIN_OFFSET: u32 = ${getViewElementOffset(props.domain as GraphDataView)}u;` : ''}
const OUTPUT_OFFSET: u32 = ${getViewElementOffset(props.output)}u;
@group(0) @binding(0) var<storage, read> inputValues: array<${shaderType}>;
${domainBinding}
@group(0) @binding(${outputBinding}) var<storage, read_write> outputCounts: array<atomic<u32>>;
${local ? `var<workgroup> localCounts: array<atomic<u32>, ${props.output.length}>;` : ''}
${integerBinningFunction}

@compute @workgroup_size(${HISTOGRAM_WORKGROUP_SIZE}) fn main(
  @builtin(global_invocation_id) globalId: vec3<u32>,
  @builtin(local_invocation_id) localId: vec3<u32>
) {
  let index = globalId.x;
  let lane = localId.x;
  ${domainInitialization}
  ${local ? 'if (lane < BIN_COUNT) { atomicStore(&localCounts[lane], 0u); }\n  workgroupBarrier();' : ''}
  var accepted = false;
  var binIndex = 0u;
  if (index < ELEMENT_COUNT && maximum >= minimum) {
    let value = inputValues[INPUT_OFFSET + index];
    if (${finiteCondition} && value >= minimum && value <= maximum) {
      if (maximum == minimum) {
        accepted = value == minimum;
      } else {
        accepted = true;
        if (value == maximum) {
          binIndex = BIN_COUNT - 1u;
        } else {
          ${binCalculation}
        }
      }
    }
  }
  ${accumulation}
}`;
  const resources: GraphBufferUse[] = [
    {buffer: props.input, usage: 'storage-read'},
    ...(gpuDomain
      ? ([{buffer: props.domain as GraphDataView, usage: 'storage-read'}] as GraphBufferUse[])
      : []),
    {buffer: props.output, usage: 'storage-read-write'}
  ];
  addComputationPass(graph, {
    id: props.id,
    source,
    resources,
    bindings: {
      inputValues: props.input,
      ...(gpuDomain ? {domainValues: props.domain as GraphDataView} : {}),
      outputCounts: props.output
    },
    dispatchCount: Math.ceil(props.input.length / HISTOGRAM_WORKGROUP_SIZE)
  });
}

function addComputationPass<Parameters>(
  graph: GPUCommandGraph<Parameters>,
  props: {
    id: string;
    source: string;
    resources: GraphBufferUse[];
    bindings: Record<string, GraphDataView>;
    dispatchCount: number;
  }
): void {
  graph.addComputePass({
    id: props.id,
    resources: props.resources,
    compile: ({device}) => {
      const computation = new Computation(device, {
        id: props.id,
        source: props.source,
        shaderLayout: {
          bindings: Object.keys(props.bindings).map((name, location) => ({
            name,
            type: 'storage' as const,
            group: 0,
            location
          }))
        }
      });
      return {
        encode: ({computePass, getBuffer}) => {
          const bindings: Record<string, Binding> = {};
          for (const [name, view] of Object.entries(props.bindings)) {
            bindings[name] = getViewBinding(view, getBuffer);
          }
          computation.setBindings(bindings);
          computation.dispatch(computePass, props.dispatchCount);
        },
        destroy: () => computation.destroy()
      };
    }
  });
}

function validateLiteralDomain(
  domain: readonly number[],
  format: GPUScalarFormat,
  id: string
): void {
  if (domain.length !== 2 || !domain.every(Number.isFinite) || domain[0] > domain[1]) {
    throw new Error(`${id} literal domain must be a finite [min, max] pair`);
  }
  if (format !== 'float32') {
    const minimum = format === 'uint32' ? 0 : -0x80000000;
    const maximum = format === 'uint32' ? 0xffffffff : 0x7fffffff;
    if (!domain.every(value => Number.isInteger(value) && value >= minimum && value <= maximum)) {
      throw new Error(`${id} literal domain values must fit ${format}`);
    }
  }
}

function isGPUHistogramDomainView<T extends GPUScalarFormat>(
  domain: GPUHistogramDomain<T>
): domain is GraphDataView<T> {
  return domain !== 'auto' && !Array.isArray(domain);
}

function getShaderType(format: GPUScalarFormat): 'u32' | 'i32' | 'f32' {
  return format === 'uint32' ? 'u32' : format === 'sint32' ? 'i32' : 'f32';
}

function getLiteral(value: number, format: GPUScalarFormat): string {
  if (format === 'uint32') return `${value}u`;
  if (format === 'sint32') return `${value}`;
  return Number.isInteger(value) ? `${value}.0` : `${value}`;
}

1	// luma.gl
2	// SPDX-License-Identifier: MIT
3	// Copyright (c) vis.gl contributors
4
5	import {type Binding} from '@luma.gl/core';
6	import {Computation} from '@luma.gl/engine';
7	import {
8	GPUCommandGraph,
9	GraphVectorView,
10	type GraphBufferUse,
11	type GraphDataView
12	} from './gpu-command-graph';
13	import {
14	createTransientView,
15	getViewBinding,
16	getViewElementOffset,
17	type GPUScalarFormat,
18	validatePackedUint32View,
19	validatePackedView
20	} from './graph-data-view-utils';
21	import {GPUReduction} from './gpu-reduction';
22
23	const HISTOGRAM_WORKGROUP_SIZE = 256;	11✔
24	const MAXIMUM_LOCAL_BIN_COUNT = 256;	11✔
25	const SCALAR_FORMATS = ['uint32', 'sint32', 'float32'] as const;	11✔
26
27	/** Domain accepted by {@link GPUHistogram}. */
28	export type GPUHistogramDomain<T extends GPUScalarFormat = GPUScalarFormat> =
29	\| readonly [number, number]
30	\| GraphDataView<T>
31	\| 'auto';
32
33	/** One scalar chunk or an ordered scalar vector counted by {@link GPUHistogram}. */
34	export type GPUHistogramInput<T extends GPUScalarFormat = GPUScalarFormat> =
35	\| GraphDataView<T>
36	\| GraphVectorView<T>;
37
38	/** Properties for graph-native scalar histogram counting. */
39	export type GPUHistogramProps<T extends GPUScalarFormat = GPUScalarFormat> = {
40	id?: string;
41	input: GPUHistogramInput<T>;
42	output: GraphDataView<'uint32'>;
43	domain: GPUHistogramDomain<T>;
44	};
45
46	/** Graph-native histogram counting for packed 32-bit scalar values. */
47	export class GPUHistogram<T extends GPUScalarFormat = GPUScalarFormat> {
48	readonly id: string;
49	readonly input: GPUHistogramInput<T>;
50	readonly output: GraphDataView<'uint32'>;
51	readonly domain: GPUHistogramDomain<T>;
52
53	constructor(props: GPUHistogramProps<T>) {
54	this.id = props.id ?? 'gpu-histogram';	14✔
55	this.input = props.input;	14✔
56	this.output = props.output;	14✔
57	this.domain = props.domain;	14✔
58	for (const [chunkIndex, input] of getHistogramInputs(this.input).entries()) {	14✔
59	const name = this.input instanceof GraphVectorView ? ` input chunk ${chunkIndex}` : ' input';	19✔
60	validatePackedView(input, SCALAR_FORMATS, `${this.id}${name}`);	19✔
61	if (input.format !== this.input.format) {	19!
NEW 62	throw new Error(`${this.id}${name} format must match the input format`);	×
63	}
64	}
65	validatePackedUint32View(this.output, `${this.id} output`);	14✔
66	if (this.output.length === 0) {	14!
67	throw new Error(`${this.id} output must contain at least one bin`);	×
68	}
69	if (getHistogramInputs(this.input).some(input => input.buffer === this.output.buffer)) {	19!
70	throw new Error(`${this.id} input and output must use separate buffers`);	×
71	}
72	if (isGPUHistogramDomainView(this.domain)) {	14✔
73	validatePackedView(this.domain, SCALAR_FORMATS, `${this.id} domain`);	1✔
74	if (this.domain.format !== this.input.format \|\| this.domain.length !== 2) {	1!
75	throw new Error(`${this.id} GPU domain must contain two ${this.input.format} rows`);	×
76	}
77	if (this.domain.buffer === this.output.buffer) {	1!
78	throw new Error(`${this.id} domain and output must use separate buffers`);	×
79	}
80	} else if (Array.isArray(this.domain)) {	13✔
81	validateLiteralDomain(this.domain, this.input.format, this.id);	10✔
82	}
83	}
84
85	/** Adds domain inference, output clearing, and accumulation nodes to a graph. */
86	addToGraph<Parameters>(graph: GPUCommandGraph<Parameters>): void {
87	const inputs = getHistogramInputs(this.input);	13✔
88	if (inputs.some(input => input.buffer.graph !== graph) \|\| this.output.buffer.graph !== graph) {	18!
UNCOV 89	throw new Error(`${this.id} views must belong to the target graph`);	×
90	}
91	let domain = this.domain;	13✔
92	if (isGPUHistogramDomainView(domain) && domain.buffer.graph !== graph) {	13!
93	throw new Error(`${this.id} domain must belong to the target graph`);	×
94	}
95	if (domain === 'auto') {	13✔
96	const inferredDomain = createTransientView(	3✔
97	graph,
98	`${this.id}-auto-domain`,
99	this.input.format,
100	2
101	) as GraphDataView<T>;
102	new GPUReduction({	3✔
103	id: `${this.id}-extent`,
104	input: this.input,
105	output: inferredDomain,
106	operation: 'extent'
107	}).addToGraph(graph);
108	domain = inferredDomain;	3✔
109	}
110	addClearHistogramPass(graph, this.id, this.output);	13✔
111	const accumulationPath = this.output.length <= MAXIMUM_LOCAL_BIN_COUNT ? 'local' : 'global';	13✔
112	inputs.forEach((input, chunkIndex) => {	13✔
113	if (input.length === 0) {	18✔
114	return;	2✔
115	}
116	addHistogramPass(graph, {	16✔
117	id:
118	this.input instanceof GraphVectorView	16✔
119	? `${this.id}-chunk-${chunkIndex}-${accumulationPath}`
120	: `${this.id}-${accumulationPath}`,
121	input,
122	output: this.output,
123	domain
124	});
125	});
126	}
127	}
128
129	function getHistogramInputs<T extends GPUScalarFormat>(
130	input: GPUHistogramInput<T>
131	): readonly GraphDataView<T>[] {
132	return input instanceof GraphVectorView ? input.data : [input];	41✔
133	}
134
135	function addClearHistogramPass<Parameters>(
136	graph: GPUCommandGraph<Parameters>,
137	id: string,
138	output: GraphDataView<'uint32'>
139	): void {
140	const passId = `${id}-clear`;	13✔
141	const source = /* wgsl */ `	13✔
142	const BIN_COUNT: u32 = ${output.length}u;
143	const OUTPUT_OFFSET: u32 = ${getViewElementOffset(output)}u;
144	@group(0) @binding(0) var<storage, read_write> outputCounts: array<atomic<u32>>;
145	@compute @workgroup_size(${HISTOGRAM_WORKGROUP_SIZE}) fn main(
146	@builtin(global_invocation_id) globalId: vec3<u32>
147	) {
148	if (globalId.x < BIN_COUNT) { atomicStore(&outputCounts[OUTPUT_OFFSET + globalId.x], 0u); }
149	}`;
150	addComputationPass(graph, {	13✔
151	id: passId,
152	source,
153	resources: [{buffer: output, usage: 'storage-write'}],
154	bindings: {outputCounts: output},
155	dispatchCount: Math.ceil(output.length / HISTOGRAM_WORKGROUP_SIZE)
156	});
157	}
158
159	function addHistogramPass<Parameters, T extends GPUScalarFormat>(
160	graph: GPUCommandGraph<Parameters>,
161	props: {
162	id: string;
163	input: GraphDataView<T>;
164	output: GraphDataView<'uint32'>;
165	domain: readonly [number, number] \| GraphDataView<T>;
166	}
167	): void {
168	const local = props.output.length <= MAXIMUM_LOCAL_BIN_COUNT;	16✔
169	const shaderType = getShaderType(props.input.format);	16✔
170	const gpuDomain = isGPUHistogramDomainView(props.domain);	16✔
171	const literalDomain = props.domain as readonly [number, number];	16✔
172	const domainBinding = gpuDomain	16✔
173	? '@group(0) @binding(1) var<storage, read> domainValues: array<' + shaderType + '>;'
174	: '';
175	const outputBinding = gpuDomain ? 2 : 1;	16✔
176	const domainInitialization = gpuDomain	16✔
177	? `let minimum: ${shaderType} = domainValues[DOMAIN_OFFSET];
178	let maximum: ${shaderType} = domainValues[DOMAIN_OFFSET + 1u];`
179	: `let minimum: ${shaderType} = ${getLiteral(literalDomain[0], props.input.format)};
180	let maximum: ${shaderType} = ${getLiteral(literalDomain[1], props.input.format)};`;
181	const finiteCondition =
182	props.input.format === 'float32' ? 'value == value && abs(value) <= 3.402823466e+38' : 'true';	16✔
183	const integerMultiplierBitCount = 32 - Math.clz32(props.output.length);	16✔
184	const integerBinningFunction =
185	props.input.format === 'float32'	16✔
186	? ''
187	: `fn multiplyDivideFloor(numerator: u32, multiplier: u32, denominator: u32) -> u32 {
188	var quotient = 0u;
189	var remainder = 0u;
190	var bitIndex = ${integerMultiplierBitCount}u;
191	loop {
192	bitIndex = bitIndex - 1u;
193	// Double and add modulo denominator without overflowing u32.
194	let doubledThreshold = denominator - remainder;
195	if (remainder >= doubledThreshold) {
196	remainder = remainder - doubledThreshold;
197	quotient = quotient * 2u + 1u;
198	} else {
199	remainder = remainder * 2u;
200	quotient = quotient * 2u;
201	}
202	if (((multiplier >> bitIndex) & 1u) != 0u) {
203	let additionThreshold = denominator - numerator;
204	if (remainder >= additionThreshold) {
205	remainder = remainder - additionThreshold;
206	quotient = quotient + 1u;
207	} else {
208	remainder = remainder + numerator;
209	}
210	}
211	if (bitIndex == 0u) { break; }
212	}
213	return quotient;
214	}`;
215	const binCalculation =
216	props.input.format === 'float32'	16✔
217	? `let ratio = (value - minimum) / (maximum - minimum);
218	binIndex = min(u32(ratio * f32(BIN_COUNT)), BIN_COUNT - 1u);`
219	: props.input.format === 'uint32'	12✔
220	? 'binIndex = multiplyDivideFloor(value - minimum, BIN_COUNT, maximum - minimum);'
221	: `let orderedValue = bitcast<u32>(value) ^ 0x80000000u;
222	let orderedMinimum = bitcast<u32>(minimum) ^ 0x80000000u;
223	let orderedMaximum = bitcast<u32>(maximum) ^ 0x80000000u;
224	binIndex = multiplyDivideFloor(
225	orderedValue - orderedMinimum,
226	BIN_COUNT,
227	orderedMaximum - orderedMinimum
228	);`;
229	const accumulation = local	16✔
230	? `if (accepted) { atomicAdd(&localCounts[binIndex], 1u); }
231	workgroupBarrier();
232	if (lane < BIN_COUNT) {
233	atomicAdd(&outputCounts[OUTPUT_OFFSET + lane], atomicLoad(&localCounts[lane]));
234	}`
235	: 'if (accepted) { atomicAdd(&outputCounts[OUTPUT_OFFSET + binIndex], 1u); }';
236	const source = /* wgsl */ `	16✔
237	const ELEMENT_COUNT: u32 = ${props.input.length}u;
238	const BIN_COUNT: u32 = ${props.output.length}u;
239	const INPUT_OFFSET: u32 = ${getViewElementOffset(props.input)}u;
240	${gpuDomain ? `const DOMAIN_OFFSET: u32 = ${getViewElementOffset(props.domain as GraphDataView)}u;` : ''}	16✔
241	const OUTPUT_OFFSET: u32 = ${getViewElementOffset(props.output)}u;
242	@group(0) @binding(0) var<storage, read> inputValues: array<${shaderType}>;
243	${domainBinding}
244	@group(0) @binding(${outputBinding}) var<storage, read_write> outputCounts: array<atomic<u32>>;
245	${local ? `var<workgroup> localCounts: array<atomic<u32>, ${props.output.length}>;` : ''}	16✔
246	${integerBinningFunction}
247
248	@compute @workgroup_size(${HISTOGRAM_WORKGROUP_SIZE}) fn main(
249	@builtin(global_invocation_id) globalId: vec3<u32>,
250	@builtin(local_invocation_id) localId: vec3<u32>
251	) {
252	let index = globalId.x;
253	let lane = localId.x;
254	${domainInitialization}
255	${local ? 'if (lane < BIN_COUNT) { atomicStore(&localCounts[lane], 0u); }\n workgroupBarrier();' : ''}	16✔
256	var accepted = false;
257	var binIndex = 0u;
258	if (index < ELEMENT_COUNT && maximum >= minimum) {
259	let value = inputValues[INPUT_OFFSET + index];
260	if (${finiteCondition} && value >= minimum && value <= maximum) {
261	if (maximum == minimum) {
262	accepted = value == minimum;
263	} else {
264	accepted = true;
265	if (value == maximum) {
266	binIndex = BIN_COUNT - 1u;
267	} else {
268	${binCalculation}
269	}
270	}
271	}
272	}
273	${accumulation}
274	}`;
275	const resources: GraphBufferUse[] = [	16✔
276	{buffer: props.input, usage: 'storage-read'},
277	...(gpuDomain	16✔
278	? ([{buffer: props.domain as GraphDataView, usage: 'storage-read'}] as GraphBufferUse[])
279	: []),
280	{buffer: props.output, usage: 'storage-read-write'}
281	];
282	addComputationPass(graph, {	16✔
283	id: props.id,
284	source,
285	resources,
286	bindings: {
287	inputValues: props.input,
288	...(gpuDomain ? {domainValues: props.domain as GraphDataView} : {}),	16✔
289	outputCounts: props.output
290	},
291	dispatchCount: Math.ceil(props.input.length / HISTOGRAM_WORKGROUP_SIZE)
292	});
293	}
294
295	function addComputationPass<Parameters>(
296	graph: GPUCommandGraph<Parameters>,
297	props: {
298	id: string;
299	source: string;
300	resources: GraphBufferUse[];
301	bindings: Record<string, GraphDataView>;
302	dispatchCount: number;
303	}
304	): void {
305	graph.addComputePass({	29✔
306	id: props.id,
307	resources: props.resources,
308	compile: ({device}) => {
309	const computation = new Computation(device, {	29✔
310	id: props.id,
311	source: props.source,
312	shaderLayout: {
313	bindings: Object.keys(props.bindings).map((name, location) => ({	50✔
314	name,
315	type: 'storage' as const,
316	group: 0,
317	location
318	}))
319	}
320	});
321	return {	29✔
322	encode: ({computePass, getBuffer}) => {
323	const bindings: Record<string, Binding> = {};	31✔
324	for (const [name, view] of Object.entries(props.bindings)) {	31✔
325	bindings[name] = getViewBinding(view, getBuffer);	53✔
326	}
327	computation.setBindings(bindings);	31✔
328	computation.dispatch(computePass, props.dispatchCount);	31✔
329	},
330	destroy: () => computation.destroy()	29✔
331	};
332	}
333	});
334	}
335
336	function validateLiteralDomain(
337	domain: readonly number[],
338	format: GPUScalarFormat,
339	id: string
340	): void {
341	if (domain.length !== 2 \|\| !domain.every(Number.isFinite) \|\| domain[0] > domain[1]) {	10✔
342	throw new Error(`${id} literal domain must be a finite [min, max] pair`);	1✔
343	}
344	if (format !== 'float32') {	9!
345	const minimum = format === 'uint32' ? 0 : -0x80000000;	9✔
346	const maximum = format === 'uint32' ? 0xffffffff : 0x7fffffff;	9✔
347	if (!domain.every(value => Number.isInteger(value) && value >= minimum && value <= maximum)) {	18!
348	throw new Error(`${id} literal domain values must fit ${format}`);	×
349	}
350	}
351	}
352
353	function isGPUHistogramDomainView<T extends GPUScalarFormat>(
354	domain: GPUHistogramDomain<T>
355	): domain is GraphDataView<T> {
356	return domain !== 'auto' && !Array.isArray(domain);	43✔
357	}
358
359	function getShaderType(format: GPUScalarFormat): 'u32' \| 'i32' \| 'f32' {
360	return format === 'uint32' ? 'u32' : format === 'sint32' ? 'i32' : 'f32';	16✔
361	}
362
363	function getLiteral(value: number, format: GPUScalarFormat): string {
364	if (format === 'uint32') return `${value}u`;	22✔
365	if (format === 'sint32') return `${value}`;	2!
366	return Number.isInteger(value) ? `${value}.0` : `${value}`;	×
367	}

visgl / luma.gl / 28670509643

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