From 174db187a89065acddcf9336bab331c505afea28 Mon Sep 17 00:00:00 2001 From: cmhhelgeson <62450112+cmhhelgeson@users.noreply.github.com> Date: Mon, 30 Oct 2023 18:11:52 -0700 Subject: [PATCH] Removed pointerEvents from non-interactive gui elements and added atomic account of number of swaps performed during a sort --- src/sample/bitonicSort/atomicToZero.wgsl | 7 ++ src/sample/bitonicSort/computeShader.ts | 13 ++-- src/sample/bitonicSort/main.ts | 95 ++++++++++++++++++++++-- 3 files changed, 104 insertions(+), 11 deletions(-) create mode 100644 src/sample/bitonicSort/atomicToZero.wgsl diff --git a/src/sample/bitonicSort/atomicToZero.wgsl b/src/sample/bitonicSort/atomicToZero.wgsl new file mode 100644 index 00000000..23e7e4ee --- /dev/null +++ b/src/sample/bitonicSort/atomicToZero.wgsl @@ -0,0 +1,7 @@ +@group(0) @binding(3) var counter: atomic; + +@compute @workgroup_size(1, 1, 1) +fn atomicToZero() { + let counterValue = atomicLoad(&counter); + atomicSub(&counter, counterValue); +} \ No newline at end of file diff --git a/src/sample/bitonicSort/computeShader.ts b/src/sample/bitonicSort/computeShader.ts index 5b112245..e0ca6134 100644 --- a/src/sample/bitonicSort/computeShader.ts +++ b/src/sample/bitonicSort/computeShader.ts @@ -15,13 +15,19 @@ struct Uniforms { } // Create local workgroup data that can contain all elements - var local_data: array; +// Define groups (functions refer to this data) +@group(0) @binding(0) var input_data: array; +@group(0) @binding(1) var output_data: array; +@group(0) @binding(2) var uniforms: Uniforms; +@group(0) @binding(3) var counter: atomic; + // Compare and swap values in local_data fn local_compare_and_swap(idx_before: u32, idx_after: u32) { //idx_before should always be < idx_after if (local_data[idx_after] < local_data[idx_before]) { + atomicAdd(&counter, 1); var temp: u32 = local_data[idx_before]; local_data[idx_before] = local_data[idx_after]; local_data[idx_after] = temp; @@ -54,11 +60,6 @@ fn get_disperse_indices(thread_id: u32, block_height: u32) -> vec2 { return idx; } -@group(0) @binding(0) var input_data: array; -@group(0) @binding(1) var output_data: array; -@group(0) @binding(2) var uniforms: Uniforms; - - fn global_compare_and_swap(idx_before: u32, idx_after: u32) { if (input_data[idx_after] < input_data[idx_before]) { output_data[idx_before] = input_data[idx_after]; diff --git a/src/sample/bitonicSort/main.ts b/src/sample/bitonicSort/main.ts index b89d7368..6e0411c9 100644 --- a/src/sample/bitonicSort/main.ts +++ b/src/sample/bitonicSort/main.ts @@ -4,6 +4,7 @@ import BitonicDisplayRenderer from './bitonicDisplay'; import bitonicDisplay from './bitonicDisplay.frag.wgsl'; import { NaiveBitonicCompute } from './computeShader'; import fullscreenTexturedQuad from '../../shaders/fullscreenTexturedQuad.wgsl'; +import atomicToZero from './atomicToZero.wgsl'; // Type of step that will be executed in our shader enum StepEnum { @@ -40,6 +41,7 @@ interface SettingsInterface { 'Next Swap Span': number; 'Total Workgroups': number; 'Display Mode': DisplayType; + 'Total Swaps': number; executeStep: boolean; 'Randomize Values': () => void; 'Execute Sort Step': () => void; @@ -98,6 +100,8 @@ SampleInitFactoryWebGPU( 'Next Swap Span': 2, // Workgroups to dispatch per frame, 'Total Workgroups': maxElements / (maxThreadsX * 2), + // The number of swap operations executed over time + 'Total Swaps': 0, // Whether we will dispatch a workload this frame executeStep: false, 'Display Mode': 'Elements', @@ -138,6 +142,17 @@ SampleInitFactoryWebGPU( usage: GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST, }); + // Initialize atomic swap buffer on GPU and CPU. Counts number of swaps actually performed by + // compute shader (when value at index x is greater than value at index y) + const atomicSwapsOutputBuffer = device.createBuffer({ + size: Uint32Array.BYTES_PER_ELEMENT, + usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC, + }); + const atomicSwapsStagingBuffer = device.createBuffer({ + size: Uint32Array.BYTES_PER_ELEMENT, + usage: GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST, + }); + // Create uniform buffer for compute shader const computeUniformsBuffer = device.createBuffer({ // width, height, blockHeight, algo @@ -146,22 +161,29 @@ SampleInitFactoryWebGPU( }); const computeBGCluster = createBindGroupCluster( - [0, 1, 2], + [0, 1, 2, 3], [ GPUShaderStage.COMPUTE | GPUShaderStage.FRAGMENT, GPUShaderStage.COMPUTE, GPUShaderStage.COMPUTE | GPUShaderStage.FRAGMENT, + GPUShaderStage.COMPUTE, + ], + ['buffer', 'buffer', 'buffer', 'buffer'], + [ + { type: 'read-only-storage' }, + { type: 'storage' }, + { type: 'uniform' }, + { type: 'storage' }, ], - ['buffer', 'buffer', 'buffer'], - [{ type: 'read-only-storage' }, { type: 'storage' }, { type: 'uniform' }], [ [ { buffer: elementsInputBuffer }, { buffer: elementsOutputBuffer }, { buffer: computeUniformsBuffer }, + { buffer: atomicSwapsOutputBuffer }, ], ], - 'NaiveBitonicSort', + 'BitonicSort', device ); @@ -177,6 +199,19 @@ SampleInitFactoryWebGPU( }, }); + // Simple pipeline that zeros out an atomic value at group 0 binding 3 + const atomicToZeroComputePipeline = device.createComputePipeline({ + layout: device.createPipelineLayout({ + bindGroupLayouts: [computeBGCluster.bindGroupLayout], + }), + compute: { + module: device.createShaderModule({ + code: atomicToZero, + }), + entryPoint: 'atomicToZero', + }, + }); + // Create bitonic debug renderer const renderPassDescriptor: GPURenderPassDescriptor = { colorAttachments: [ @@ -230,6 +265,17 @@ SampleInitFactoryWebGPU( // Reset block heights prevBlockHeightController.setValue(0); nextBlockHeightController.setValue(2); + + // Reset Total Swaps by setting atomic value to 0 + const commandEncoder = device.createCommandEncoder(); + const computePassEncoder = commandEncoder.beginComputePass(); + computePassEncoder.setPipeline(atomicToZeroComputePipeline); + computePassEncoder.setBindGroup(0, computeBGCluster.bindGroups[0]); + computePassEncoder.dispatchWorkgroups(1); + computePassEncoder.end(); + device.queue.submit([commandEncoder.finish()]); + totalSwapsController.setValue(0); + highestBlockHeight = 2; }; @@ -400,6 +446,10 @@ SampleInitFactoryWebGPU( settings, 'Next Step' ); + const totalSwapsController = executionInformationFolder.add( + settings, + 'Total Swaps' + ); const prevBlockHeightController = executionInformationFolder.add( settings, 'Prev Swap Span' @@ -435,7 +485,11 @@ SampleInitFactoryWebGPU( }); // Deactivate interaction with select GUI elements + totalWorkgroupsController.domElement.style.pointerEvents = 'none'; + hoveredCellController.domElement.style.pointerEvents = 'none'; + swappedCellController.domElement.style.pointerEvents = 'none'; stepIndexController.domElement.style.pointerEvents = 'none'; + totalStepsController.domElement.style.pointerEvents = 'none'; prevStepController.domElement.style.pointerEvents = 'none'; prevBlockHeightController.domElement.style.pointerEvents = 'none'; nextStepController.domElement.style.pointerEvents = 'none'; @@ -516,6 +570,8 @@ SampleInitFactoryWebGPU( ? nextStepController.setValue('DISPERSE_GLOBAL') : nextStepController.setValue('DISPERSE_LOCAL'); } + + // Copy GPU accessible buffers to CPU accessible buffers commandEncoder.copyBufferToBuffer( elementsOutputBuffer, 0, @@ -523,6 +579,14 @@ SampleInitFactoryWebGPU( 0, elementsBufferSize ); + + commandEncoder.copyBufferToBuffer( + atomicSwapsOutputBuffer, + 0, + atomicSwapsStagingBuffer, + 0, + Uint32Array.BYTES_PER_ELEMENT + ); } device.queue.submit([commandEncoder.finish()]); @@ -537,14 +601,31 @@ SampleInitFactoryWebGPU( 0, elementsBufferSize ); + // Copy atomic swaps data to CPU + await atomicSwapsStagingBuffer.mapAsync( + GPUMapMode.READ, + 0, + Uint32Array.BYTES_PER_ELEMENT + ); + const copySwapsBuffer = atomicSwapsStagingBuffer.getMappedRange( + 0, + Uint32Array.BYTES_PER_ELEMENT + ); // Get correct range of data from CPU copy of GPU Data const elementsData = copyElementsBuffer.slice( 0, Uint32Array.BYTES_PER_ELEMENT * settings['Total Elements'] ); + const swapsData = copySwapsBuffer.slice( + 0, + Uint32Array.BYTES_PER_ELEMENT + ); // Extract data const elementsOutput = new Uint32Array(elementsData); + totalSwapsController.setValue(new Uint32Array(swapsData)[0]); elementsStagingBuffer.unmap(); + atomicSwapsStagingBuffer.unmap(); + // Elements output becomes elements input, swap accumulate elements = elementsOutput; setSwappedCell(); } @@ -578,7 +659,11 @@ const bitonicSortExample: () => JSX.Element = () => }, { name: './bitonicCompute.frag.wgsl', - contents: NaiveBitonicCompute(16), + contents: NaiveBitonicCompute(64), + }, + { + name: './atomicToZero.wgsl', + contents: atomicToZero, }, ], filename: __filename,