From 9878d7aa428450ec86ce0f37012f97e87989d7da Mon Sep 17 00:00:00 2001 From: Chris Duncan Date: Tue, 31 Dec 2024 14:36:42 -0800 Subject: [PATCH] Refactor WebGPU pow to align with MDN compute shader example which was clearer to follow and to improve alignment with intended work generation algorithm. --- src/lib/workers/powgpu.ts | 292 +++++++++++++++++++------------------- 1 file changed, 143 insertions(+), 149 deletions(-) diff --git a/src/lib/workers/powgpu.ts b/src/lib/workers/powgpu.ts index a52cd59..81bdc57 100644 --- a/src/lib/workers/powgpu.ts +++ b/src/lib/workers/powgpu.ts @@ -4,10 +4,11 @@ // https://github.com/numtel/nano-webgl-pow /// 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 { @@ -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 ubo: UBO; - @group(0) @binding(1) var work: array, 2>; - + static #shader = ` struct UBO { blockhash: array, 2>, threshold: u32, - workload: f32, + workload: u32 }; + @group(0) @binding(0) var ubo: UBO; + @group(0) @binding(1) var work: array>; /** * 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) { - 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) { // 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; - var u_work0: vec4 = work[0u]; - var u_work1: vec4 = 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 { - 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) => any, threshold: number): Promise { - - 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} ` -- 2.34.1