]> zoso.dev Git - libnemo.git/commitdiff
Refactor WebGPU pow to align with MDN compute shader example which was clearer to...
authorChris Duncan <chris@zoso.dev>
Tue, 31 Dec 2024 22:36:42 +0000 (14:36 -0800)
committerChris Duncan <chris@zoso.dev>
Tue, 31 Dec 2024 22:36:42 +0000 (14:36 -0800)
src/lib/workers/powgpu.ts

index a52cd5987d328f65cf9f172bec40f05edb7b01a6..81bdc57c10c98361014f51331e54043eb657f7dc 100644 (file)
@@ -4,10 +4,11 @@
 // https://github.com/numtel/nano-webgl-pow
 /// <reference types="@webgpu/types" />
 import { WorkerInterface } from '../pool.js'
+import powgl from './powgl.js'
 
-export class Pow extends WorkerInterface {
+export class PowGpu extends WorkerInterface {
        static {
-               Pow.listen()
+               PowGpu.listen()
        }
 
        /**
@@ -32,7 +33,7 @@ export class Pow extends WorkerInterface {
        /**
        * Finds a nonce that satisfies the Nano proof-of-work requirements.
        *
-       * @param {string} hashHex - Hexadecimal hash of previous block, or public key for new accounts
+       * @param {string} hash - Hexadecimal hash of previous block, or public key for new accounts
        * @param {number} [threshold=0xfffffff8] - Difficulty of proof-of-work calculation
        */
        static async find (hash: string, threshold: number = 0xfffffff8): Promise<string> {
@@ -41,34 +42,19 @@ export class Pow extends WorkerInterface {
                })
        }
 
-
-       /** Used to set canvas size. Must be a multiple of 256. */
-       static #WORKLOAD: number = 256 * Math.max(1, Math.floor(navigator.hardwareConcurrency))
-
-       static #hexify (arr: number[] | Uint8Array): string {
-               let out = ''
-               for (let i = arr.length - 1; i >= 0; i--) {
-                       out += arr[i].toString(16).padStart(2, '0')
-               }
-               return out
-       }
-
-       static #device: GPUDevice | null = null;
-       static #pipeline: GPUComputePipeline | null = null;
-       static #workgroupSize: number = 256; // Must align with shader
-       static #bindGroupLayout: GPUBindGroupLayout | null = null;
-
+       static #workDispatchSize: number = Math.max(1, Math.floor(navigator.hardwareConcurrency))
+       static #workgroupSize: number = 256 // Must align with shader
+       static #workload: number = this.#workDispatchSize * this.#workgroupSize
 
        // WebGPU Compute Shader
-       static #computeShader = `
-               @group(0) @binding(0) var<uniform> ubo: UBO;
-               @group(0) @binding(1) var<storage, read_write> work: array<vec4<u32>, 2>;
-
+       static #shader = `
                struct UBO {
                        blockhash: array<vec4<u32>, 2>,
                        threshold: u32,
-                       workload: f32,
+                       workload: u32
                };
+               @group(0) @binding(0) var<uniform> ubo: UBO;
+               @group(0) @binding(1) var<storage, read_write> work: array<vec2<u32>>;
 
                /**
                * Defined separately from uint v[32] below as the original value is required
@@ -156,31 +142,18 @@ export class Pow extends WorkerInterface {
                /**
                * Main compute function
                */
-               @compute @workgroup_size(${this.#workgroupSize})
-               fn main(@builtin(global_invocation_id) global_id : vec3<u32>) {
-                       let uv_x = global_id.x;
-                       let uv_y = global_id.y;
-
+               @compute @workgroup_size(${this.#workgroupSize}, ${this.#workgroupSize})
+               fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
                        // Check bounds, may be unnecessary with proper dispatch size
-                       if (uv_x >= u32(ubo.workload) || uv_y >= u32(ubo.workload) ) {
+                       if (global_id.x >= ubo.workload || global_id.y >= ubo.workload) {
                                return;
                        }
 
                        var m: array<u32, 16>;
-                       var u_work0: vec4<u32> = work[0u];
-                       var u_work1: vec4<u32> = work[1u];
-                       let x_pos = uv_x % 256u;
-                       let y_pos = uv_y % 256u;
-                       let x_index = (uv_x - x_pos) / 256u;
-                       let y_index = (uv_y - y_pos) / 256u;
 
-                       // First 2 work bytes are the x,y pos within the 256x256 area, the next
-                       // two bytes are modified from the random generated value, XOR'd with
-                       // the x,y area index of where this pixel is located
-                       m[0u] = (x_pos ^ (y_pos << 8u) ^ ((u_work0.b ^ x_index) << 16u) ^ ((u_work0.a ^ y_index) << 24u));
-
-                       // Remaining bytes are un-modified from the random generated value
-                       m[1u] = (u_work1.r ^ (u_work1.g << 8u) ^ (u_work1.b << 16u) ^ (u_work1.a << 24u));
+                       // Workgroup index defines work value for this calculation point
+                       m[0u] = global_id.x;
+                       m[1u] = global_id.y;
 
                        // Block hash
                        m[2u] = ubo.blockhash[0u].x;
@@ -227,98 +200,117 @@ export class Pow extends WorkerInterface {
                        }
 
                        // Store the result directly into work array
-                       if ((BLAKE2B_IV32_1 ^ v[1u] ^ v[17u]) > ubo.threshold) {
-                               work[0u].x = x_index + 1u;
-                               work[0u].y = y_index + 1u;
-                               work[0u].z = x_pos;
-                               work[0u].w = y_pos;
-                       }
+                       // if ((BLAKE2B_IV32_1 ^ v[1u] ^ v[17u]) > 0) {
+                               work[global_id.x].x = 1u ^ (2u << 8u);
+                               work[global_id.x].y = 3u ^ (4u << 8u);
+                       // }
                }
        `;
 
-       // Initialize WebGPU
-       static async #initializeWebGPU (): Promise<void> {
-               if (!navigator.gpu) {
-                       throw new Error("WebGPU is not supported.")
-               }
-
-               const adapter = await navigator.gpu.requestAdapter()
-               if (!adapter) {
-                       throw new Error("No suitable WebGPU adapter found.")
+       static #hexify (arr: number[] | Uint8Array): string {
+               let out = ''
+               for (let i = arr.length - 1; i >= 0; i--) {
+                       out += arr[i].toString(16).padStart(2, '0')
                }
+               return out
+       }
 
-               this.#device = await adapter.requestDevice()
-               this.#bindGroupLayout = this.#device.createBindGroupLayout({
-                       entries: [
-                               {
-                                       binding: 0,
-                                       visibility: GPUShaderStage.COMPUTE,
-                                       buffer: { type: 'uniform', },
-                               },
-                               {
-                                       binding: 1,
-                                       visibility: GPUShaderStage.COMPUTE,
-                                       buffer: { type: 'storage', },
-                               },
-                       ],
-               })
+       // Initialize WebGPU
+       static #device: GPUDevice | null = null
+       static #gpuBuffer: GPUBuffer
+       static #cpuBuffer: GPUBuffer
+       static #bindGroupLayout: GPUBindGroupLayout
+       static #pipeline: GPUComputePipeline
 
-               this.#pipeline = this.#device.createComputePipeline({
-                       layout: this.#device.createPipelineLayout({
-                               bindGroupLayouts: [this.#bindGroupLayout],
-                       }),
-                       compute: {
-                               module: this.#device.createShaderModule({
-                                       code: this.#computeShader,
-                               }),
-                               entryPoint: 'main',
-                       },
-               })
+       // Initialize WebGPU
+       static {
+               // Request device and adapter
+               if (navigator.gpu == null) {
+                       throw new Error('WebGPU is not supported in this browser.')
+               }
+               navigator.gpu.requestAdapter()
+                       .then(adapter => {
+                               if (adapter == null) {
+                                       throw new Error('WebGPU adapter refused by browser.')
+                               }
+                               adapter.requestDevice()
+                                       .then(device => {
+                                               this.#device = device
+
+                                               // Create buffers for writing GPU calculations and reading from Javascript
+                                               this.#gpuBuffer = this.#device.createBuffer({
+                                                       size: this.#workload,
+                                                       usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST | GPUBufferUsage.COPY_SRC
+                                               })
+                                               this.#cpuBuffer = this.#device.createBuffer({
+                                                       size: this.#workload,
+                                                       usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ
+                                               })
+
+                                               // Create binding group data structure and use it later once UBO is known
+                                               this.#bindGroupLayout = this.#device.createBindGroupLayout({
+                                                       entries: [
+                                                               {
+                                                                       binding: 0,
+                                                                       visibility: GPUShaderStage.COMPUTE,
+                                                                       buffer: {
+                                                                               type: 'uniform'
+                                                                       },
+                                                               },
+                                                               {
+                                                                       binding: 1,
+                                                                       visibility: GPUShaderStage.COMPUTE,
+                                                                       buffer: {
+                                                                               type: 'storage'
+                                                                       },
+                                                               }
+                                                       ]
+                                               })
+
+                                               // Create pipeline to connect compute shader to binding layout
+                                               this.#pipeline = this.#device.createComputePipeline({
+                                                       layout: this.#device.createPipelineLayout({
+                                                               bindGroupLayouts: [this.#bindGroupLayout]
+                                                       }),
+                                                       compute: {
+                                                               entryPoint: 'main',
+                                                               module: this.#device.createShaderModule({
+                                                                       code: this.#shader
+                                                               })
+                                                       }
+                                               })
+                                       })
+                       })
+                       .catch(err => { throw new Error(err.message) })
        }
 
-
        static async #calculate (hashHex: string, callback: (nonce: string | PromiseLike<string>) => any, threshold: number): Promise<void> {
-
-               if (!/^[A-F-a-f0-9]{64}$/.test(hashHex)) throw new Error(`invalid_hash ${hashHex}`)
+               if (!/^[A-Fa-f0-9]{64}$/.test(hashHex)) throw new Error(`Invalid hash ${hashHex}`)
                if (typeof threshold !== 'number') throw new TypeError(`Invalid threshold ${threshold}`)
 
-               // Ensure WebGPU is initialized. Call this once
-               if (!this.#device) {
-                       this.#initializeWebGPU()
-                               .then(() => {
-                                       this.#calculate(hashHex, callback, threshold) // restart calculation
-                               })
-                               .catch((error) => {
-                                       console.error("Failed to initalize WebGPU:", error)
-                               })
-                       return // Stop execution until WebGPU is initalized
+               // Ensure WebGPU is initialized else restart calculation
+               if (PowGpu.#device == null) {
+                       setTimeout(async () => { await this.#calculate(hashHex, callback, threshold) })
+                       return
                }
 
                // Set up uniform buffer object
-               const uboView = new DataView(new ArrayBuffer(144))
+               const uboView = new DataView(new ArrayBuffer(48))
                for (let i = 0; i < 64; i += 8) {
                        const uint32 = hashHex.slice(i, i + 8)
-                       uboView.setUint32(i * 2, parseInt(uint32, 16))
+                       uboView.setUint32(i / 2, parseInt(uint32, 16))
                }
-               uboView.setUint32(128, threshold, true)
-               uboView.setFloat32(132, Pow.#WORKLOAD - 1, true)
-               const uboBuffer = this.#device.createBuffer({
+               uboView.setUint32(32, threshold, true)
+               uboView.setUint32(40, PowGpu.#workload, true)
+               const uboBuffer = PowGpu.#device.createBuffer({
                        size: uboView.byteLength,
                        usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST,
                })
-               this.#device.queue.writeBuffer(uboBuffer, 0, uboView)
+               PowGpu.#device.queue.writeBuffer(uboBuffer, 0, uboView)
 
                // Work buffer
-               const work = new Uint8Array(8)
-               crypto.getRandomValues(work)
-               const workBuffer = this.#device.createBuffer({
-                       size: work.byteLength,
-                       usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST | GPUBufferUsage.COPY_SRC,
-               })
-               this.#device.queue.writeBuffer(workBuffer, 0, work)
-
-               const bindGroup = this.#device.createBindGroup({
-                       layout: this.#bindGroupLayout!,
+               const bindGroup = PowGpu.#device.createBindGroup({
+                       layout: PowGpu.#bindGroupLayout,
                        entries: [
                                {
                                        binding: 0,
@@ -329,53 +321,55 @@ export class Pow extends WorkerInterface {
                                {
                                        binding: 1,
                                        resource: {
-                                               buffer: workBuffer,
+                                               buffer: PowGpu.#gpuBuffer,
                                        },
                                },
                        ],
                })
 
-
-               const commandEncoder = this.#device.createCommandEncoder()
+               // Create command encoder to issue commands to GPU and initiate render pass
+               const commandEncoder = PowGpu.#device.createCommandEncoder()
                const passEncoder = commandEncoder.beginComputePass()
-               passEncoder.setPipeline(this.#pipeline!)
-               passEncoder.setBindGroup(0, bindGroup)
-
-               const dispatchSizeX = Math.ceil(this.#WORKLOAD / this.#workgroupSize)
-               const dispatchSizeY = Math.ceil(this.#WORKLOAD / this.#workgroupSize)
 
-               passEncoder.dispatchWorkgroups(dispatchSizeX, dispatchSizeY)
+               // Issue commands and end render pass
+               passEncoder.setPipeline(PowGpu.#pipeline)
+               passEncoder.setBindGroup(0, bindGroup)
+               passEncoder.dispatchWorkgroups(PowGpu.#workDispatchSize)
                passEncoder.end()
 
-               // Get result
-               const resultStagingBuffer = this.#device.createBuffer({
-                       size: workBuffer.size,
-                       usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ
-               })
+               // Copy result from GPU buffer to CPU buffer
                commandEncoder.copyBufferToBuffer(
-                       workBuffer,
+                       PowGpu.#gpuBuffer,
                        0,
-                       resultStagingBuffer,
+                       PowGpu.#cpuBuffer,
                        0,
-                       workBuffer.size
+                       PowGpu.#workload
                )
-               this.#device.queue.submit([commandEncoder.finish()])
-
-               // Get result
-               await resultStagingBuffer.mapAsync(GPUMapMode.READ)
-
-               const arrayBuffer = resultStagingBuffer.getMappedRange()
-               const result = new Uint32Array(arrayBuffer)
-               resultStagingBuffer.unmap() // Unmap after reading
-
-               for (let i = 0; i < result.length; i += 4) {
-                       if (result[i] !== 0) {
-                               const hex = this.#hexify(work.subarray(4, 8)) + this.#hexify([
-                                       result[i + 2],
-                                       result[i + 3],
-                                       work[2] ^ (result[i] - 1),
-                                       work[3] ^ (result[i + 1] - 1)
-                               ])
+
+               // End frame by passing array of command buffers to command queue for execution
+               PowGpu.#device.queue.submit([commandEncoder.finish()])
+
+               // Read results back to Javascript and then unmap buffer after reading
+               await PowGpu.#cpuBuffer.mapAsync(GPUMapMode.READ)
+               const result = new Uint32Array(PowGpu.#cpuBuffer.getMappedRange())
+               PowGpu.#cpuBuffer.unmap()
+
+               console.log(`result`)
+               console.dir(result)
+               console.log(`result?`)
+               for (let i = 0; i < result.length; i += 2) {
+                       const work = new Uint8Array([result[i], result[i + 1]])
+                       console.log(`result[${i}]: ${result[i]}`)
+                       console.log(`result[${i + 1}]: ${result[i + 1]}`)
+                       console.log(`work: ${work}`)
+                       if (result[i] !== 0 || result[i + 1] !== 0) {
+                               // const hex = PowGpu.#hexify(work.subarray(4, 8)) + PowGpu.#hexify([
+                               //      result[i + 2],
+                               //      result[i + 3],
+                               //      work[2] ^ (result[i] - 1),
+                               //      work[3] ^ (result[i + 1] - 1)
+                               // ])
+                               const hex = PowGpu.#hexify(work)
                                typeof callback === 'function' && callback(hex)
                                return
                        }
@@ -387,5 +381,5 @@ export class Pow extends WorkerInterface {
 
 export default `
        const WorkerInterface = ${WorkerInterface}
-       const Pow = ${Pow}
+       const Pow = ${PowGpu}
 `