diff --git a/examples/experimental/gpu-trace-viewer/app.ts b/examples/experimental/gpu-trace-viewer/app.ts index dc6e517a92..636f6815cf 100644 --- a/examples/experimental/gpu-trace-viewer/app.ts +++ b/examples/experimental/gpu-trace-viewer/app.ts @@ -40,6 +40,7 @@ import { getTraceDependencyCapacityOptions, isTraceDensityMode, makeTraceDataset, + makeTraceSpanChunks, TRACE_COLLAPSED_STATE, TRACE_DENSITY_BIN_COUNT, TRACE_DEPENDENCY_BATCH_CAPACITY, @@ -56,6 +57,8 @@ import { TRACE_LANES_PER_THREAD, TRACE_PROCESS_COUNT, TRACE_SPAN_BATCH_CAPACITY, + TRACE_SPAN_CHUNK_TARGET_BYTE_LENGTH, + TRACE_SPAN_RECORD_WORD_LENGTH, TRACE_STATUS_COUNT, TRACE_THREAD_COUNT, TRACE_THREADS_PER_PROCESS, @@ -130,6 +133,24 @@ type TraceGroupResources = { firstSpanIndex: number; }; +type TraceSpanChunkResources = { + buffer: Buffer; + uniforms: Buffer; + chunkIndex: number; + firstSpanIndex: number; + spanCount: number; + firstBatchIndex: number; + batchCount: number; +}; + +type TraceSpanDrawResources = { + commandIndex: number; + groupIndex: number; + chunkIndex: number; + firstBatchIndex: number; + batchCount: number; +}; + type PickPosition = { time: number; lane: number; @@ -147,7 +168,10 @@ type TraceGraphResources = { readbackRing: GPUReadbackRing; renderBundle: RenderBundle; groups: TraceGroupResources[]; - spans: Buffer; + spanChunks: TraceSpanChunkResources[]; + spanDraws: TraceSpanDrawResources[]; + dependencyDrawCommandIndex: number; + densityDrawCommandIndex: number; spanBatchIndex: Buffer; candidateBatchIds: Buffer; visibleSpanIds: Buffer; @@ -186,7 +210,7 @@ function getTraceResourceBuffers(resources: TraceGraphResources): Array<{byteLen resources.densityCandidateDispatchCommands.buffer, resources.pickCandidateDispatchCommands.buffer, resources.candidateDependencyDispatchCommands.buffer, - resources.spans, + ...resources.spanChunks.flatMap(chunk => [chunk.buffer, chunk.uniforms]), resources.spanBatchIndex, resources.candidateBatchIds, resources.visibleSpanIds, @@ -237,6 +261,7 @@ export default class GPUTraceViewerAnimationLoopTemplate extends AnimationLoopTe readonly graphInspector = new GPUCommandGraphInspector({maxSamples: 90}); readonly capacityOptions: number[]; readonly dependencyCapacityOptions: number[]; + private readonly spanChunkByteLength: number; private resources: TraceGraphResources | null = null; private graphObservation: GPUCommandGraphInspectorObservation | null = null; @@ -286,13 +311,22 @@ export default class GPUTraceViewerAnimationLoopTemplate extends AnimationLoopTe constructor({ device, traceCapacity = DEFAULT_CAPACITY, - dependencyCapacity = DEFAULT_DEPENDENCY_CAPACITY - }: AnimationProps & {traceCapacity?: number; dependencyCapacity?: number}) { + dependencyCapacity = DEFAULT_DEPENDENCY_CAPACITY, + spanChunkByteLength = TRACE_SPAN_CHUNK_TARGET_BYTE_LENGTH + }: AnimationProps & { + traceCapacity?: number; + dependencyCapacity?: number; + spanChunkByteLength?: number; + }) { super(); if (device.type !== 'webgpu') { throw new Error('GPU Hierarchical Trace Viewer requires WebGPU'); } this.device = device; + if (!Number.isSafeInteger(spanChunkByteLength) || spanChunkByteLength < 1) { + throw new RangeError('Trace span chunk byte length must be a positive safe integer'); + } + this.spanChunkByteLength = spanChunkByteLength; this.capacityOptions = getTraceCapacityOptions( device.limits.maxStorageBufferBindingSize, device.limits.maxBufferSize @@ -449,7 +483,8 @@ export default class GPUTraceViewerAnimationLoopTemplate extends AnimationLoopTe {name: 'threadOffsets', type: 'read-only-storage', group: 0, location: 2}, {name: 'threadStates', type: 'read-only-storage', group: 0, location: 3}, {name: 'reachedSpans', type: 'read-only-storage', group: 0, location: 4}, - {name: 'viewUniforms', type: 'uniform', group: 0, location: 5} + {name: 'viewUniforms', type: 'uniform', group: 0, location: 5}, + {name: 'spanChunk', type: 'uniform', group: 0, location: 6} ] }, parameters: makeTraceBlendParameters() @@ -504,9 +539,9 @@ export default class GPUTraceViewerAnimationLoopTemplate extends AnimationLoopTe const started = performance.now(); this.destroyResources(); this.spanCapacity = spanCapacity; - this.dependencyCapacity = dependencyCapacity; + this.dependencyCapacity = this.getSupportedDependencyCapacity(spanCapacity, dependencyCapacity); this.selectedSpanIndex = INVALID_SPAN_INDEX; - const dataset = makeTraceDataset(spanCapacity, dependencyCapacity); + const dataset = makeTraceDataset(spanCapacity, this.dependencyCapacity); const resources = this.createResources(dataset); resources.renderBundle = this.createRenderBundle(resources); resources.compiled = this.createGraph(resources, dataset); @@ -526,6 +561,16 @@ export default class GPUTraceViewerAnimationLoopTemplate extends AnimationLoopTe this.updateInspector(); } + private getSupportedDependencyCapacity(spanCapacity: number, dependencyCapacity: number): number { + const spanByteLength = + spanCapacity * TRACE_SPAN_RECORD_WORD_LENGTH * Uint32Array.BYTES_PER_ELEMENT; + const maximumMonolithicSpanByteLength = Math.min( + this.device.limits.maxStorageBufferBindingSize, + this.device.limits.maxBufferSize + ); + return spanByteLength <= maximumMonolithicSpanByteLength ? dependencyCapacity : 0; + } + /** Uploads each canonical source or mutable interaction allocation exactly once. */ private createResources(dataset: TraceDatasetData): TraceGraphResources { const groups = dataset.groups.map(group => ({ @@ -539,11 +584,51 @@ export default class GPUTraceViewerAnimationLoopTemplate extends AnimationLoopTe const topologyChunkLengths = getTopologyChunkLengths(dataset.spanCount); const outgoingOffsets = makePartitionedOffsets(dataset.outgoing.offsets, topologyChunkLengths); const incomingOffsets = makePartitionedOffsets(dataset.incoming.offsets, topologyChunkLengths); + const maximumDirectSpanByteLength = Math.min( + this.device.limits.maxStorageBufferBindingSize, + this.device.limits.maxBufferSize + ); + const maximumSpanChunkByteLength = + dataset.dependencyCount > 0 + ? maximumDirectSpanByteLength + : Math.min(maximumDirectSpanByteLength, this.spanChunkByteLength); + const spanChunkData = makeTraceSpanChunks( + dataset.spans, + dataset.spanBatches, + maximumSpanChunkByteLength + ); + // Dependency endpoint rendering currently requires one directly addressable span allocation. + if (spanChunkData.length > 1 && dataset.dependencyCount > 0) { + throw new Error(); + } + const spanDraws: TraceSpanDrawResources[] = spanChunkData.flatMap(chunk => + groups.flatMap((_, groupIndex) => { + const chunkBatches = dataset.spanBatches + .slice(chunk.firstBatchIndex, chunk.firstBatchIndex + chunk.batchCount) + .filter(batch => batch.groupIndex === groupIndex); + return chunkBatches.length > 0 + ? [ + { + commandIndex: 0, + groupIndex, + chunkIndex: chunk.chunkIndex, + firstBatchIndex: chunkBatches[0].batchIndex, + batchCount: chunkBatches.length + } + ] + : []; + }) + ); + spanDraws.forEach((draw, commandIndex) => { + draw.commandIndex = commandIndex; + }); + const dependencyDrawCommandIndex = spanDraws.length; + const densityDrawCommandIndex = dependencyDrawCommandIndex + 1; const drawCommands = new DrawCommandBuffer(this.device, { id: 'gpu-trace-draw-commands', type: 'draw', commands: [ - ...groups.map(() => ({vertexCount: 6, instanceCount: 0})), + ...spanDraws.map(() => ({vertexCount: 6, instanceCount: 0})), {vertexCount: 2, instanceCount: 0}, {vertexCount: 6, instanceCount: densityBinCount} ] @@ -584,7 +669,31 @@ export default class GPUTraceViewerAnimationLoopTemplate extends AnimationLoopTe candidateDependencyDispatchCommands, readbackRing, groups, - spans: this.createDataBuffer('gpu-trace-spans', dataset.spans), + spanChunks: spanChunkData.map(chunk => { + const buffer = this.createDataBuffer(`gpu-trace-spans-${chunk.chunkIndex}`, chunk.data); + const uniforms = this.createDataBuffer( + `gpu-trace-span-chunk-uniforms-${chunk.chunkIndex}`, + Uint32Array.of( + chunk.firstSpanIndex, + chunk.spanCount, + chunk.firstBatchIndex, + chunk.batchCount + ), + Buffer.UNIFORM + ); + return { + buffer, + uniforms, + chunkIndex: chunk.chunkIndex, + firstSpanIndex: chunk.firstSpanIndex, + spanCount: chunk.spanCount, + firstBatchIndex: chunk.firstBatchIndex, + batchCount: chunk.batchCount + }; + }), + spanDraws, + dependencyDrawCommandIndex, + densityDrawCommandIndex, spanBatchIndex: this.createDataBuffer('gpu-trace-span-batch-index', dataset.spanBatchIndex), candidateBatchIds: this.createStorageBuffer( 'gpu-trace-candidate-batch-ids', @@ -666,11 +775,11 @@ export default class GPUTraceViewerAnimationLoopTemplate extends AnimationLoopTe }; } - private createDataBuffer(id: string, data: Uint32Array): Buffer { + private createDataBuffer(id: string, data: Uint32Array, additionalUsage = 0): Buffer { return this.device.createBuffer({ id, data: data.length > 0 ? data : new Uint32Array(1), - usage: Buffer.STORAGE | Buffer.COPY_DST + usage: Buffer.STORAGE | Buffer.COPY_DST | additionalUsage }); } @@ -692,7 +801,15 @@ export default class GPUTraceViewerAnimationLoopTemplate extends AnimationLoopTe }); const handles = { uniforms: importTraceBuffer(graph, 'view-uniforms', this.viewUniformBuffer), - spans: importTraceBuffer(graph, 'spans', resources.spans), + spanChunks: resources.spanChunks.map(chunk => ({ + ...chunk, + spans: importTraceBuffer(graph, `spans-${chunk.chunkIndex}`, chunk.buffer), + uniforms: importTraceBuffer( + graph, + `span-chunk-uniforms-${chunk.chunkIndex}`, + chunk.uniforms + ) + })), spanBatchIndex: importTraceBuffer(graph, 'span-batch-index', resources.spanBatchIndex), candidateBatchIds: importTraceBuffer( graph, @@ -1001,11 +1118,12 @@ export default class GPUTraceViewerAnimationLoopTemplate extends AnimationLoopTe length: 1, workgroupSize: 1 }); + for (const chunk of handles.spanChunks) { addTraceIndirectComputePass(graph, { - id: 'trace-candidate-span-visibility', - source: getCandidateVisibilityShader(), + id: `trace-candidate-span-visibility-${chunk.chunkIndex}`, + source: getCandidateVisibilityShader(chunk), bindings: [ - storageRead('spans', handles.spans), + storageRead('spans', chunk.spans), storageRead('spanBatches', handles.spanBatchIndex), storageRead('candidateBatchIds', handles.candidateBatchIds), uniformBinding('viewUniforms', handles.uniforms), @@ -1017,6 +1135,7 @@ export default class GPUTraceViewerAnimationLoopTemplate extends AnimationLoopTe ], dispatchBuffer: handles.exactCandidateDispatchCommands }); + } addTraceComputePass(graph, { id: 'trace-clear-density', source: getDensityClearShader(), @@ -1024,11 +1143,12 @@ export default class GPUTraceViewerAnimationLoopTemplate extends AnimationLoopTe length: TRACE_LANE_COUNT * TRACE_DENSITY_BIN_COUNT, workgroupSize: TRACE_WORKGROUP_SIZE }); + for (const chunk of handles.spanChunks) { addTraceIndirectComputePass(graph, { - id: 'trace-candidate-density', - source: getCandidateDensityShader(), + id: `trace-candidate-density-${chunk.chunkIndex}`, + source: getCandidateDensityShader(chunk), bindings: [ - storageRead('spans', handles.spans), + storageRead('spans', chunk.spans), storageRead('spanBatches', handles.spanBatchIndex), storageRead('candidateBatchIds', handles.candidateBatchIds), uniformBinding('viewUniforms', handles.uniforms), @@ -1041,10 +1161,10 @@ export default class GPUTraceViewerAnimationLoopTemplate extends AnimationLoopTe dispatchBuffer: handles.densityCandidateDispatchCommands }); addTraceIndirectComputePass(graph, { - id: 'trace-candidate-pick', - source: getCandidatePickShader(), + id: `trace-candidate-pick-${chunk.chunkIndex}`, + source: getCandidatePickShader(chunk), bindings: [ - storageRead('spans', handles.spans), + storageRead('spans', chunk.spans), storageRead('spanBatches', handles.spanBatchIndex), storageRead('candidateBatchIds', handles.candidateBatchIds), uniformBinding('viewUniforms', handles.uniforms), @@ -1055,6 +1175,7 @@ export default class GPUTraceViewerAnimationLoopTemplate extends AnimationLoopTe ], dispatchBuffer: handles.pickCandidateDispatchCommands }); + } const visibleSpanCountBuffer = graph.createTransientBuffer({ id: 'trace-visible-span-count', byteLength: UINT32_BYTE_LENGTH, @@ -1084,27 +1205,21 @@ export default class GPUTraceViewerAnimationLoopTemplate extends AnimationLoopTe }), count: graph.createDataView(visibleSpanCountBuffer, {format: 'uint32', length: 1}) }).addToGraph(graph); - const groupBatchRanges = resources.groups.map((_, groupIndex) => { - const firstBatchIndex = dataset.spanBatches.findIndex( - batch => batch.groupIndex === groupIndex - ); - const batchCount = dataset.spanBatches.filter( - batch => batch.groupIndex === groupIndex - ).length; - return {firstBatchIndex, batchCount}; - }); addTraceComputePass(graph, { id: 'trace-publish-span-draw-commands', - source: getTraceDrawCommandsShader(groupBatchRanges), + source: getTraceDrawCommandsShader(resources.spanDraws), bindings: [ storageRead('rangeCounts', rangeCompaction.rangeCounts.buffer), storageRead('rangeOffsets', rangeCompaction.rangeOffsets.buffer), storageWrite('drawCommands', handles.drawCommands) ], - length: resources.groups.length + length: resources.spanDraws.length }); const renderResources: GraphBufferUse[] = [ - {buffer: handles.spans, usage: 'storage-read'}, + ...handles.spanChunks.flatMap(chunk => [ + {buffer: chunk.spans, usage: 'storage-read'} as const, + {buffer: chunk.uniforms, usage: 'uniform'} as const + ]), {buffer: handles.visibleSpanIds, usage: 'storage-read'}, {buffer: handles.dependencies, usage: 'storage-read'}, {buffer: handles.processStates, usage: 'storage-read'}, @@ -1194,7 +1309,9 @@ export default class GPUTraceViewerAnimationLoopTemplate extends AnimationLoopTe count: graph.createDataView(handles.drawCommands, { format: 'uint32', length: 1, - byteOffset: resources.drawCommands.getInstanceCountByteOffset(resources.groups.length) + byteOffset: resources.drawCommands.getInstanceCountByteOffset( + resources.dependencyDrawCommandIndex + ) }) }).addToGraph(graph); } @@ -1225,31 +1342,35 @@ export default class GPUTraceViewerAnimationLoopTemplate extends AnimationLoopTe }); encoder.setPipeline(this.model.pipeline); encoder.setVertexArray(this.model.vertexArray); - for (const [groupIndex] of resources.groups.entries()) { + for (const draw of resources.spanDraws) { + const chunk = resources.spanChunks[draw.chunkIndex]; encoder.setBindings({ - spans: resources.spans, + spans: chunk.buffer, visibleIds: resources.visibleSpanIds, threadOffsets: resources.threadOffsets, threadStates: resources.threadStates, reachedSpans: resources.reachedSpans, - viewUniforms: this.viewUniformBuffer + viewUniforms: this.viewUniformBuffer, + spanChunk: chunk.uniforms }); - resources.drawCommands.draw(encoder, groupIndex); + resources.drawCommands.draw(encoder, draw.commandIndex); } + if (resources.dependencyCount > 0) { encoder.setPipeline(this.dependencyModel.pipeline); encoder.setVertexArray(this.dependencyModel.vertexArray); encoder.setBindings({ dependencies: resources.dependencies, visibleDependencyIds: resources.visibleDependencyIds, - spans: resources.spans, + spans: resources.spanChunks[0].buffer, processStates: resources.processStates, threadStates: resources.threadStates, threadOffsets: resources.threadOffsets, dependencyResults: resources.dependencyResults, viewUniforms: this.viewUniformBuffer }); - resources.drawCommands.draw(encoder, resources.groups.length); + resources.drawCommands.draw(encoder, resources.dependencyDrawCommandIndex); + } encoder.setPipeline(this.densityModel.pipeline); encoder.setVertexArray(this.densityModel.vertexArray); @@ -1257,7 +1378,7 @@ export default class GPUTraceViewerAnimationLoopTemplate extends AnimationLoopTe densityBins: resources.densityBins, viewUniforms: this.viewUniformBuffer }); - resources.drawCommands.draw(encoder, resources.groups.length + 1); + resources.drawCommands.draw(encoder, resources.densityDrawCommandIndex); return encoder.finish(); } @@ -1302,10 +1423,12 @@ export default class GPUTraceViewerAnimationLoopTemplate extends AnimationLoopTe return; } const values = new Uint32Array(bytes.buffer, bytes.byteOffset, bytes.byteLength / 4); - this.sampledVisibleCounts = resources.groups.map( - (_, groupIndex) => values[groupIndex * 4 + 1] ?? 0 + this.sampledVisibleCounts = resources.groups.map((_, groupIndex) => + resources.spanDraws + .filter(draw => draw.groupIndex === groupIndex) + .reduce((count, draw) => count + (values[draw.commandIndex * 4 + 1] ?? 0), 0) ); - this.sampledDependencyCount = values[resources.groups.length * 4 + 1] ?? 0; + this.sampledDependencyCount = values[resources.dependencyDrawCommandIndex * 4 + 1] ?? 0; this.recordWorkloadCounters(); this.updateInspector(); } catch { @@ -1409,7 +1532,7 @@ export default class GPUTraceViewerAnimationLoopTemplate extends AnimationLoopTe resources.candidateDependencyDispatchCommands.destroy(); resources.readbackRing.destroy(); for (const buffer of [ - resources.spans, + ...resources.spanChunks.flatMap(chunk => [chunk.buffer, chunk.uniforms]), resources.spanBatchIndex, resources.candidateBatchIds, resources.visibleSpanIds, @@ -1513,6 +1636,7 @@ export default class GPUTraceViewerAnimationLoopTemplate extends AnimationLoopTe `` ) .join('')} + Dependencies automatically switch off when the selected span source requires multiple GPU chunks.
Span groups${groupControls}
Status${statusControls}
@@ -1564,8 +1688,15 @@ export default class GPUTraceViewerAnimationLoopTemplate extends AnimationLoopTe } if (target.matches('[data-span-capacity]')) { this.rebuild(Number(target.value), this.dependencyCapacity); + const dependencyCapacity = root.querySelector( + '[data-dependency-capacity]' + ); + if (dependencyCapacity) { + dependencyCapacity.value = String(this.dependencyCapacity); + } } else if (target.matches('[data-dependency-capacity]')) { this.rebuild(this.spanCapacity, Number(target.value)); + target.value = String(this.dependencyCapacity); } else if (target instanceof HTMLInputElement && target.dataset.group !== undefined) { const group = Number(target.dataset.group); this.enabledMask = setBit(this.enabledMask, group, target.checked); @@ -1719,7 +1850,7 @@ export default class GPUTraceViewerAnimationLoopTemplate extends AnimationLoopTe this.dependencyCapacity, this.device.limits ); - this.capacityElement.innerHTML = `${formatCount(this.spanCapacity)} spans · ${formatCount(resources.spanBatchCount)} batches · ${formatCount(resources.dependencyCount)}/${formatCount(this.dependencyCapacity)} dependencies · graph compile #${this.compileCount} (${this.compileTimeMilliseconds.toFixed(1)} ms)
${formatBytes(this.allocationStats.persistentByteLength)} persistent in ${formatCount(this.allocationStats.bufferCount)} buffers · largest ${formatBytes(this.allocationStats.largestBufferByteLength)} · source contract ${capacityContract.fitsDeviceLimits ? 'fits device' : 'exceeds device'}`; + this.capacityElement.innerHTML = `${formatCount(this.spanCapacity)} spans · ${formatCount(resources.spanBatchCount)} batches · ${formatCount(resources.dependencyCount)}/${formatCount(this.dependencyCapacity)} dependencies · graph compile #${this.compileCount} (${this.compileTimeMilliseconds.toFixed(1)} ms)
${resources.spanChunks.length} span chunk${resources.spanChunks.length === 1 ? '' : 's'} · ${formatBytes(this.allocationStats.persistentByteLength)} persistent in ${formatCount(this.allocationStats.bufferCount)} buffers · largest ${formatBytes(this.allocationStats.largestBufferByteLength)} · chunked contract ${capacityContract.fitsChunkedDeviceLimits ? 'fits device' : 'exceeds device'}`; } if (this.selectionElement) { this.selectionElement.textContent = diff --git a/examples/experimental/gpu-trace-viewer/trace-benchmark.ts b/examples/experimental/gpu-trace-viewer/trace-benchmark.ts index b9a9eb6121..afcdff9a8f 100644 --- a/examples/experimental/gpu-trace-viewer/trace-benchmark.ts +++ b/examples/experimental/gpu-trace-viewer/trace-benchmark.ts @@ -2,7 +2,11 @@ // SPDX-License-Identifier: MIT // SPDX-FileCopyrightText: Copyright (c) vis.gl contributors -import {TRACE_DEPENDENCY_RECORD_WORD_LENGTH, TRACE_SPAN_RECORD_WORD_LENGTH} from './trace-data'; +import { + TRACE_DEPENDENCY_RECORD_WORD_LENGTH, + TRACE_SPAN_CHUNK_TARGET_BYTE_LENGTH, + TRACE_SPAN_RECORD_WORD_LENGTH +} from './trace-data'; const UINT32_BYTE_LENGTH = Uint32Array.BYTES_PER_ELEMENT; @@ -69,6 +73,9 @@ export type TraceCapacityContract = { fitsStorageBufferBindingSize: boolean; fitsMaxBufferSize: boolean; fitsDeviceLimits: boolean; + spanChunkCount: number; + largestSpanChunkByteLength: number; + fitsChunkedDeviceLimits: boolean; }; /** Persistent GPU-buffer accounting independent of command-graph transient allocations. */ @@ -112,6 +119,22 @@ export function getTraceCapacityContract( const fitsStorageBufferBindingSize = largestSourceBufferByteLength <= limits.maxStorageBufferBindingSize; const fitsMaxBufferSize = largestSourceBufferByteLength <= limits.maxBufferSize; + const maximumSpanChunkByteLength = Math.min( + TRACE_SPAN_CHUNK_TARGET_BYTE_LENGTH, + limits.maxStorageBufferBindingSize, + limits.maxBufferSize + ); + const spanChunkCount = + maximumSpanChunkByteLength > 0 + ? Math.max(1, Math.ceil(spanBufferByteLength / maximumSpanChunkByteLength)) + : 0; + const largestSpanChunkByteLength = + maximumSpanChunkByteLength > 0 ? Math.min(spanBufferByteLength, maximumSpanChunkByteLength) : 0; + const fitsChunkedDeviceLimits = + maximumSpanChunkByteLength >= TRACE_SPAN_RECORD_WORD_LENGTH * UINT32_BYTE_LENGTH && + dependencyBufferByteLength <= limits.maxStorageBufferBindingSize && + dependencyBufferByteLength <= limits.maxBufferSize && + (dependencyCapacity === 0 || (fitsStorageBufferBindingSize && fitsMaxBufferSize)); return Object.freeze({ spanCapacity, dependencyCapacity, @@ -122,7 +145,10 @@ export function getTraceCapacityContract( maxBufferSize: limits.maxBufferSize, fitsStorageBufferBindingSize, fitsMaxBufferSize, - fitsDeviceLimits: fitsStorageBufferBindingSize && fitsMaxBufferSize + fitsDeviceLimits: fitsStorageBufferBindingSize && fitsMaxBufferSize, + spanChunkCount, + largestSpanChunkByteLength, + fitsChunkedDeviceLimits }); } diff --git a/examples/experimental/gpu-trace-viewer/trace-data.ts b/examples/experimental/gpu-trace-viewer/trace-data.ts index e756875d5f..b2116e8f37 100644 --- a/examples/experimental/gpu-trace-viewer/trace-data.ts +++ b/examples/experimental/gpu-trace-viewer/trace-data.ts @@ -20,6 +20,8 @@ export const TRACE_SPAN_BATCH_RECORD_WORD_LENGTH = 8; export const TRACE_DEPENDENCY_BATCH_RECORD_WORD_LENGTH = 6; // One span batch maps to one portable WebGPU workgroup for candidate-driven local compaction. export const TRACE_SPAN_BATCH_CAPACITY = 256; +/** Keeps each chunk comfortably below portable storage-binding and allocation ceilings. */ +export const TRACE_SPAN_CHUNK_TARGET_BYTE_LENGTH = 64 * 1024 * 1024; // One dependency batch maps to one portable WebGPU workgroup for candidate-driven compaction. export const TRACE_DEPENDENCY_BATCH_CAPACITY = 128; const TRACE_DEMONSTRATION_CAPACITIES = [250_000, 1_000_000, 4_000_000, 10_000_000]; @@ -50,16 +52,19 @@ export const TRACE_COLLAPSED_STATE = 0; export const TRACE_EXPANDED_STATE = 1; export const TRACE_INVALID_SPAN_INDEX = 0xffffffff; -/** Returns useful demonstration sizes that fit in one span storage-buffer binding. */ +/** Returns useful demonstration sizes when one portable span batch fits in a storage chunk. */ export function getTraceCapacityOptions( maxStorageBufferBindingSize: number, maxBufferSize: number ): number[] { const spanRecordByteLength = TRACE_SPAN_RECORD_WORD_LENGTH * Uint32Array.BYTES_PER_ELEMENT; - const maximumSpanCapacity = Math.floor( - Math.min(maxStorageBufferBindingSize, maxBufferSize) / spanRecordByteLength + const maximumChunkSpanCount = Math.floor( + Math.min(maxStorageBufferBindingSize, maxBufferSize, TRACE_SPAN_CHUNK_TARGET_BYTE_LENGTH) / + spanRecordByteLength ); - return TRACE_DEMONSTRATION_CAPACITIES.filter(capacity => capacity <= maximumSpanCapacity); + return maximumChunkSpanCount >= TRACE_SPAN_BATCH_CAPACITY + ? [...TRACE_DEMONSTRATION_CAPACITIES] + : TRACE_DEMONSTRATION_CAPACITIES.filter(capacity => capacity <= maximumChunkSpanCount); } /** Returns useful dependency limits that fit in one dependency storage-buffer binding. */ @@ -72,7 +77,10 @@ export function getTraceDependencyCapacityOptions( const maximumDependencyCapacity = Math.floor( Math.min(maxStorageBufferBindingSize, maxBufferSize) / dependencyRecordByteLength ); - return TRACE_DEMONSTRATION_CAPACITIES.filter(capacity => capacity <= maximumDependencyCapacity); + return [ + 0, + ...TRACE_DEMONSTRATION_CAPACITIES.filter(capacity => capacity <= maximumDependencyCapacity) + ]; } /** Matches the GPU's adaptive exact-span versus density-rendering decision. */ @@ -111,6 +119,16 @@ export type TraceSpanBatchData = { data: Uint32Array; }; +/** Borrowed contiguous span rows grouped on complete candidate-batch boundaries. */ +export type TraceSpanChunkData = { + chunkIndex: number; + firstSpanIndex: number; + spanCount: number; + firstBatchIndex: number; + batchCount: number; + data: Uint32Array; +}; + /** Stable dependency range with a conservative endpoint and ancestor time envelope. */ export type TraceDependencyBatchData = { batchIndex: number; @@ -390,6 +408,49 @@ export function makeTraceSpanBatches( return {spanBatches, spanBatchIndex}; } +/** Splits canonical spans without repacking rows or cutting candidate batches. */ +export function makeTraceSpanChunks( + spans: Uint32Array, + spanBatches: readonly TraceSpanBatchData[], + maximumChunkByteLength: number +): TraceSpanChunkData[] { + const spanRecordByteLength = TRACE_SPAN_RECORD_WORD_LENGTH * Uint32Array.BYTES_PER_ELEMENT; + const maximumChunkSpanCount = Math.floor(maximumChunkByteLength / spanRecordByteLength); + // Candidate shaders require every portable batch to remain wholly addressable in one chunk. + if (maximumChunkSpanCount < TRACE_SPAN_BATCH_CAPACITY) { + throw new RangeError(); + } + const chunks: TraceSpanChunkData[] = []; + let firstBatchIndex = 0; + while (firstBatchIndex < spanBatches.length) { + const firstBatch = spanBatches[firstBatchIndex]; + let batchCount = 0; + let spanCount = 0; + while (firstBatchIndex + batchCount < spanBatches.length) { + const batch = spanBatches[firstBatchIndex + batchCount]; + if (batchCount > 0 && spanCount + batch.count > maximumChunkSpanCount) { + break; + } + spanCount += batch.count; + batchCount++; + } + const firstSpanIndex = firstBatch.firstSpanIndex; + chunks.push({ + chunkIndex: chunks.length, + firstSpanIndex, + spanCount, + firstBatchIndex, + batchCount, + data: spans.subarray( + firstSpanIndex * TRACE_SPAN_RECORD_WORD_LENGTH, + (firstSpanIndex + spanCount) * TRACE_SPAN_RECORD_WORD_LENGTH + ) + }); + firstBatchIndex += batchCount; + } + return chunks; +} + /** Maintains the original example API while retaining canonical source references. */ export function makeTraceGroups(totalSpanCount: number): TraceGroupData[] { return makeTraceDataset(totalSpanCount).groups; diff --git a/examples/experimental/gpu-trace-viewer/trace-shaders.ts b/examples/experimental/gpu-trace-viewer/trace-shaders.ts index 0f73762859..1b7683e3b3 100644 --- a/examples/experimental/gpu-trace-viewer/trace-shaders.ts +++ b/examples/experimental/gpu-trace-viewer/trace-shaders.ts @@ -26,6 +26,20 @@ import { const TRACE_WORKGROUP_SIZE = 256; export const TRACE_FOCUS_FRONTIER_WORKGROUP_SIZE = 64; +export type TraceSpanChunkShaderProps = { + firstSpanIndex: number; + spanCount: number; + firstBatchIndex: number; + batchCount: number; +}; + +function getSpanChunkDeclarations(props: TraceSpanChunkShaderProps): string { + return `const CHUNK_FIRST_SPAN_INDEX: u32 = ${props.firstSpanIndex}u; +const CHUNK_SPAN_COUNT: u32 = ${props.spanCount}u; +const CHUNK_FIRST_BATCH_INDEX: u32 = ${props.firstBatchIndex}u; +const CHUNK_BATCH_COUNT: u32 = ${props.batchCount}u;`; +} + const TRACE_SHADER_DECLARATIONS = /* wgsl */ ` struct TraceSpan { start: f32, @@ -137,6 +151,13 @@ ${TRACE_SHADER_DECLARATIONS} @group(0) @binding(3) var threadStates: array; @group(0) @binding(4) var reachedSpans: array; @group(0) @binding(5) var viewUniforms: ViewUniforms; +struct SpanChunkUniforms { + firstSpanIndex: u32, + spanCount: u32, + firstBatchIndex: u32, + batchCount: u32, +}; +@group(0) @binding(6) var spanChunk: SpanChunkUniforms; struct VertexOutput { @builtin(position) position: vec4, @@ -149,7 +170,7 @@ struct VertexOutput { @builtin(instance_index) instanceIndex: u32 ) -> VertexOutput { let sourceIndex = visibleIds[instanceIndex]; - let span = spans[sourceIndex]; + let span = spans[sourceIndex - spanChunk.firstSpanIndex]; let corner = getCorner(vertexIndex); let timeRange = max(viewUniforms.timeMax - viewUniforms.timeMin, 0.0001); let laneRange = max(viewUniforms.laneMax - viewUniforms.laneMin, 1.0); @@ -523,9 +544,10 @@ fn main(@builtin(global_invocation_id) globalId: vec3) { } /** Classifies and aggregates focused candidate density without span-sized intermediate keys. */ -export function getCandidateDensityShader(): string { +export function getCandidateDensityShader(props: TraceSpanChunkShaderProps): string { return /* wgsl */ ` ${TRACE_SHADER_DECLARATIONS} +${getSpanChunkDeclarations(props)} struct TraceSpanBatch { firstSpanIndex: u32, spanCount: u32, @@ -552,13 +574,20 @@ fn main( @builtin(global_invocation_id) globalId: vec3, @builtin(workgroup_id) workgroupId: vec3 ) { - let batch = spanBatches[candidateBatchIds[workgroupId.y]]; + let batchIndex = candidateBatchIds[workgroupId.y]; + if ( + batchIndex < CHUNK_FIRST_BATCH_INDEX || + batchIndex >= CHUNK_FIRST_BATCH_INDEX + CHUNK_BATCH_COUNT + ) { + return; + } + let batch = spanBatches[batchIndex]; let batchRowIndex = globalId.x; if (batchRowIndex >= batch.spanCount) { return; } let sourceIndex = batch.firstSpanIndex + batchRowIndex; - let span = spans[sourceIndex]; + let span = spans[sourceIndex - CHUNK_FIRST_SPAN_INDEX]; let processExpanded = processStates[span.processIndex] != 0u; let localLane = select(0u, span.lane % LANES_PER_THREAD, threadStates[span.threadIndex] != 0u); let expandedLane = threadOffsets[span.threadIndex] + localLane; @@ -580,41 +609,42 @@ fn main( }`; } -/** Publishes stable global visible-ID slices into the per-group indirect draw commands. */ +/** Publishes stable visible-ID slices into the per-group, per-chunk indirect draw commands. */ export function getTraceDrawCommandsShader( - groupBatchRanges: readonly {firstBatchIndex: number; batchCount: number}[] + drawBatchRanges: readonly {firstBatchIndex: number; batchCount: number}[] ): string { - const firstBatchIndices = groupBatchRanges.map(range => `${range.firstBatchIndex}u`).join(', '); - const lastBatchIndices = groupBatchRanges + const firstBatchIndices = drawBatchRanges.map(range => `${range.firstBatchIndex}u`).join(', '); + const lastBatchIndices = drawBatchRanges .map(range => `${range.firstBatchIndex + range.batchCount - 1}u`) .join(', '); return /* wgsl */ ` -const GROUP_COUNT: u32 = ${groupBatchRanges.length}u; -const FIRST_BATCH_INDICES = array(${firstBatchIndices}); -const LAST_BATCH_INDICES = array(${lastBatchIndices}); +const DRAW_COUNT: u32 = ${drawBatchRanges.length}u; +const FIRST_BATCH_INDICES = array(${firstBatchIndices}); +const LAST_BATCH_INDICES = array(${lastBatchIndices}); @group(0) @binding(0) var rangeCounts: array; @group(0) @binding(1) var rangeOffsets: array; @group(0) @binding(2) var drawCommands: array; @compute @workgroup_size(${TRACE_WORKGROUP_SIZE}) fn main(@builtin(global_invocation_id) globalId: vec3) { - let groupIndex = globalId.x; - if (groupIndex >= GROUP_COUNT) { + let drawIndex = globalId.x; + if (drawIndex >= DRAW_COUNT) { return; } - let firstBatchIndex = FIRST_BATCH_INDICES[groupIndex]; - let lastBatchIndex = LAST_BATCH_INDICES[groupIndex]; + let firstBatchIndex = FIRST_BATCH_INDICES[drawIndex]; + let lastBatchIndex = LAST_BATCH_INDICES[drawIndex]; let firstInstance = rangeOffsets[firstBatchIndex]; let endInstance = rangeOffsets[lastBatchIndex] + rangeCounts[lastBatchIndex]; - let commandOffset = groupIndex * 4u; + let commandOffset = drawIndex * 4u; drawCommands[commandOffset + 1u] = endInstance - firstInstance; drawCommands[commandOffset + 3u] = firstInstance; }`; } /** Publishes focused, generation-tagged exact visibility for candidate spans. */ -export function getCandidateVisibilityShader(): string { +export function getCandidateVisibilityShader(props: TraceSpanChunkShaderProps): string { return /* wgsl */ ` ${TRACE_SHADER_DECLARATIONS} +${getSpanChunkDeclarations(props)} struct TraceSpanBatch { firstSpanIndex: u32, spanCount: u32, @@ -641,13 +671,20 @@ fn main( @builtin(global_invocation_id) globalId: vec3, @builtin(workgroup_id) workgroupId: vec3 ) { - let batch = spanBatches[candidateBatchIds[workgroupId.y]]; + let batchIndex = candidateBatchIds[workgroupId.y]; + if ( + batchIndex < CHUNK_FIRST_BATCH_INDEX || + batchIndex >= CHUNK_FIRST_BATCH_INDEX + CHUNK_BATCH_COUNT + ) { + return; + } + let batch = spanBatches[batchIndex]; let batchRowIndex = globalId.x; if (batchRowIndex >= batch.spanCount) { return; } let sourceIndex = batch.firstSpanIndex + batchRowIndex; - let span = spans[sourceIndex]; + let span = spans[sourceIndex - CHUNK_FIRST_SPAN_INDEX]; let processExpanded = processStates[span.processIndex] != 0u; let localLane = select(0u, span.lane % LANES_PER_THREAD, threadStates[span.threadIndex] != 0u); let expandedLane = threadOffsets[span.threadIndex] + localLane; @@ -668,9 +705,10 @@ fn main( } /** Resolves explicit picking inside the same compacted candidate batches as classification. */ -export function getCandidatePickShader(): string { +export function getCandidatePickShader(props: TraceSpanChunkShaderProps): string { return /* wgsl */ ` ${TRACE_SHADER_DECLARATIONS} +${getSpanChunkDeclarations(props)} struct TraceSpanBatch { firstSpanIndex: u32, spanCount: u32, @@ -699,13 +737,20 @@ fn main( if (viewUniforms.pickLane < 0.0) { return; } - let batch = spanBatches[candidateBatchIds[workgroupId.y]]; + let batchIndex = candidateBatchIds[workgroupId.y]; + if ( + batchIndex < CHUNK_FIRST_BATCH_INDEX || + batchIndex >= CHUNK_FIRST_BATCH_INDEX + CHUNK_BATCH_COUNT + ) { + return; + } + let batch = spanBatches[batchIndex]; let batchRowIndex = globalId.x; if (batchRowIndex >= batch.spanCount) { return; } let sourceIndex = batch.firstSpanIndex + batchRowIndex; - let span = spans[sourceIndex]; + let span = spans[sourceIndex - CHUNK_FIRST_SPAN_INDEX]; let processExpanded = processStates[span.processIndex] != 0u; let localLane = select(0u, span.lane % LANES_PER_THREAD, threadStates[span.threadIndex] != 0u); let expandedLane = threadOffsets[span.threadIndex] + localLane; diff --git a/test/examples/gpu-trace-viewer.node.spec.ts b/test/examples/gpu-trace-viewer.node.spec.ts index 3dfde6f1b8..0c7a38a4ba 100644 --- a/test/examples/gpu-trace-viewer.node.spec.ts +++ b/test/examples/gpu-trace-viewer.node.spec.ts @@ -20,6 +20,7 @@ import { makeTraceDependencyBatches, makeTraceGroups, makeTraceSpanBatches, + makeTraceSpanChunks, TRACE_CROSS_PROCESS_DEPENDENCY, TRACE_DEPENDENCY_BATCH_RECORD_WORD_LENGTH, TRACE_DEPENDENCY_RECORD_WORD_LENGTH, @@ -30,6 +31,7 @@ import { TRACE_PROCESS_COUNT, TRACE_SAME_PROCESS_DEPENDENCY, TRACE_SPAN_BATCH_RECORD_WORD_LENGTH, + TRACE_SPAN_CHUNK_TARGET_BYTE_LENGTH, TRACE_SPAN_RECORD_WORD_LENGTH, TRACE_THREAD_COUNT, TRACE_THREADS_PER_PROCESS @@ -89,18 +91,18 @@ test('deck GPU trace copies complete canonical span records', t => { test('GPU trace capacity options adapt to negotiated WebGPU buffer limits', t => { t.deepEqual( getTraceCapacityOptions(128 * 1024 * 1024, 256 * 1024 * 1024), - [250_000, 1_000_000, 4_000_000], - 'portable limits retain the four-million-span ceiling' + [250_000, 1_000_000, 4_000_000, 10_000_000], + 'portable limits expose ten million spans through chunked source storage' ); t.deepEqual( getTraceDependencyCapacityOptions(128 * 1024 * 1024, 256 * 1024 * 1024), - [250_000, 1_000_000, 4_000_000], + [0, 250_000, 1_000_000, 4_000_000], 'dependency options use their smaller fixed-width record size' ); t.deepEqual( getTraceCapacityOptions(256 * 1024 * 1024, 1024 * 1024 * 1024), - [250_000, 1_000_000, 4_000_000], - 'a 256 MiB storage binding remains below the ten-million-span requirement' + [250_000, 1_000_000, 4_000_000, 10_000_000], + 'chunking removes the single-binding ceiling' ); t.deepEqual( getTraceCapacityOptions(1024 * 1024 * 1024, 1024 * 1024 * 1024), @@ -110,6 +112,39 @@ test('GPU trace capacity options adapt to negotiated WebGPU buffer limits', t => t.end(); }); +test('GPU trace span chunks preserve complete candidate batches and borrowed source rows', t => { + const dataset = makeTraceDataset(2048, 0); + const chunks = makeTraceSpanChunks( + dataset.spans, + dataset.spanBatches, + 300 * TRACE_SPAN_RECORD_WORD_LENGTH * Uint32Array.BYTES_PER_ELEMENT + ); + t.ok(chunks.length > 1, 'small synthetic chunk limit creates multiple source buffers'); + t.equal( + chunks.reduce((count, chunk) => count + chunk.spanCount, 0), + dataset.spanCount, + 'chunks cover every source span exactly once' + ); + t.ok( + chunks.every(chunk => chunk.data.buffer === dataset.spans.buffer), + 'chunk rows remain borrowed views of canonical source storage' + ); + t.ok( + chunks.every(chunk => + dataset.spanBatches + .slice(chunk.firstBatchIndex, chunk.firstBatchIndex + chunk.batchCount) + .every( + batch => + batch.firstSpanIndex >= chunk.firstSpanIndex && + batch.firstSpanIndex + batch.count <= chunk.firstSpanIndex + chunk.spanCount + ) + ), + 'no candidate batch crosses a chunk boundary' + ); + t.equal(TRACE_SPAN_CHUNK_TARGET_BYTE_LENGTH, 64 * 1024 * 1024, 'target stays portable'); + t.end(); +}); + test('GPU trace supremacy contract exposes standard scales and interaction scenarios', t => { t.deepEqual( TRACE_BENCHMARK_CAPACITIES, @@ -148,6 +183,23 @@ test('GPU trace capacity contract makes the monolithic 10M limit explicit', t => '10M dependencies require a 160 MB source buffer' ); t.equal(portable.fitsDeviceLimits, false, 'portable limits reject the monolithic 10M source'); + t.equal(portable.spanChunkCount, 5, 'portable chunk target splits the 320 MB source five ways'); + t.equal( + getTraceCapacityContract(10_000_000, 0, { + maxStorageBufferBindingSize: 128 * 1024 * 1024, + maxBufferSize: 256 * 1024 * 1024 + }).fitsChunkedDeviceLimits, + true, + 'portable limits admit ten million spans when dependency endpoint lookup is disabled' + ); + t.equal( + getTraceCapacityContract(10_000_000, 250_000, { + maxStorageBufferBindingSize: 128 * 1024 * 1024, + maxBufferSize: 256 * 1024 * 1024 + }).fitsChunkedDeviceLimits, + false, + 'dependency endpoint lookup keeps its monolithic span requirement explicit' + ); const maximum = getTraceCapacityContract(10_000_000, 10_000_000, { maxStorageBufferBindingSize: 1024 * 1024 * 1024, @@ -207,17 +259,23 @@ test('GPU trace LOD switches at a stable trace-time-per-pixel threshold', t => { }); test('GPU trace adaptive LOD shaders parse as WGSL', t => { + const spanChunk = { + firstSpanIndex: 0, + spanCount: 11, + firstBatchIndex: 0, + batchCount: 3 + }; const shaders = [ TRACE_RENDER_SHADER, TRACE_DEPENDENCY_RENDER_SHADER, TRACE_DENSITY_RENDER_SHADER, getBatchVisibilityShader(3), getPickClearShader(), - getCandidateVisibilityShader(), + getCandidateVisibilityShader(spanChunk), getCandidatePassDispatchShader(), getDensityClearShader(), - getCandidateDensityShader(), - getCandidatePickShader(), + getCandidateDensityShader(spanChunk), + getCandidatePickShader(spanChunk), getTraceDrawCommandsShader([ {firstBatchIndex: 0, batchCount: 2}, {firstBatchIndex: 2, batchCount: 1} diff --git a/test/examples/gpu-trace-viewer.spec.ts b/test/examples/gpu-trace-viewer.spec.ts index 6e791f5b9a..7f4f7e9853 100644 --- a/test/examples/gpu-trace-viewer.spec.ts +++ b/test/examples/gpu-trace-viewer.spec.ts @@ -12,6 +12,8 @@ import { TRACE_FILTER_HIDE_OVERLAPPING_CHILDREN, TRACE_FILTER_HIDE_SIMILAR_DURATION_PARENTS, TRACE_PROCESS_COUNT, + TRACE_SPAN_BATCH_CAPACITY, + TRACE_SPAN_RECORD_WORD_LENGTH, TRACE_THREAD_COUNT } from '../../examples/experimental/gpu-trace-viewer/trace-data'; @@ -373,4 +375,74 @@ describe('GPU hierarchical trace viewer', () => { host.remove(); } }, 30_000); + + test('renders exact spans from multiple bounded source chunks', async () => { + const device = await getWebGPUTestDevice('core'); + if (!device) { + return; + } + if ( + device.info.gpu === 'software' || + device.info.gpuType === 'cpu' || + Boolean(device.info.fallback) + ) { + return; + } + + const host = document.createElement('div'); + host.id = 'example-panel-host'; + document.body.append(host); + const spanChunkByteLength = + TRACE_SPAN_BATCH_CAPACITY * TRACE_SPAN_RECORD_WORD_LENGTH * Uint32Array.BYTES_PER_ELEMENT; + let viewer: GPUTraceViewerAnimationLoopTemplate | null = null; + try { + viewer = new GPUTraceViewerAnimationLoopTemplate({ + device, + traceCapacity: 4096, + dependencyCapacity: 0, + spanChunkByteLength + } as AnimationProps & { + traceCapacity: number; + dependencyCapacity: number; + spanChunkByteLength: number; + }); + const state = viewer as unknown as { + resources: { + drawCommands: {buffer: {readAsync: () => Promise}}; + spanChunks: Array<{buffer: {byteLength: number}}>; + spanDraws: Array<{commandIndex: number; chunkIndex: number}>; + dependencyCount: number; + }; + }; + + expect(state.resources.dependencyCount).toBe(0); + expect(state.resources.spanChunks.length).toBeGreaterThan(1); + expect(state.resources.spanDraws.length).toBeGreaterThan(3); + expect( + state.resources.spanChunks.every(chunk => chunk.buffer.byteLength <= spanChunkByteLength) + ).toBe(true); + + viewer.onRender({ + device, + time: 6000, + width: 2048, + height: 1 + } as AnimationProps); + device.submit(); + const drawCommandBytes = await state.resources.drawCommands.buffer.readAsync(); + const drawCommands = new Uint32Array( + drawCommandBytes.buffer, + drawCommandBytes.byteOffset, + drawCommandBytes.byteLength / Uint32Array.BYTES_PER_ELEMENT + ); + const visibleSpanCount = state.resources.spanDraws.reduce( + (count, draw) => count + drawCommands[draw.commandIndex * 4 + 1], + 0 + ); + expect(visibleSpanCount).toBeGreaterThan(0); + } finally { + viewer?.onFinalize(); + host.remove(); + } + }, 30_000); });