From 11fdf1eea9dfed48439e3df2ead77ea4db35cb03 Mon Sep 17 00:00:00 2001 From: Jin Wei Tan <48111396+CarrotzRule123@users.noreply.github.com> Date: Sat, 20 Aug 2022 19:58:12 +0800 Subject: [PATCH] fix: Update WGPU (#26) --- backend/webgpu/backend.ts | 5 +- backend/webgpu/shaders/binary.ts | 46 ++++----- backend/webgpu/shaders/binary.wgsl | 12 +-- backend/webgpu/shaders/matmul.ts | 76 +++++++-------- backend/webgpu/shaders/matmul.wgsl | 32 +++--- backend/webgpu/shaders/pad.ts | 62 ++++++------ backend/webgpu/shaders/pad.wgsl | 26 ++--- backend/webgpu/shaders/perlin.ts | 130 ++++++++++++------------- backend/webgpu/shaders/perlin.wgsl | 42 ++++---- backend/webgpu/shaders/transpose.ts | 52 +++++----- backend/webgpu/shaders/transpose.wgsl | 24 ++--- backend/webgpu/shaders/unary.ts | 42 ++++---- backend/webgpu/shaders/unary.wgsl | 10 +- examples/webgpu_binary.ts | 30 ------ examples/webgpu_matmul.ts | 22 ----- examples/webgpu_pad.ts | 29 ------ examples/webgpu_perlin.ts | 135 -------------------------- examples/webgpu_reduce.ts | 19 ---- examples/webgpu_transpose.ts | 27 ------ examples/webgpu_unary.ts | 15 --- tests/deps.ts | 1 - tests/webgpu_binary_test.ts | 59 ----------- tests/webgpu_matmul_test.ts | 29 ------ tests/webgpu_pad_test.ts | 42 -------- tests/webgpu_transpose_test.ts | 39 -------- tests/webgpu_unary_test.ts | 23 ----- 26 files changed, 280 insertions(+), 749 deletions(-) delete mode 100644 examples/webgpu_binary.ts delete mode 100644 examples/webgpu_matmul.ts delete mode 100644 examples/webgpu_pad.ts delete mode 100644 examples/webgpu_perlin.ts delete mode 100644 examples/webgpu_reduce.ts delete mode 100644 examples/webgpu_transpose.ts delete mode 100644 examples/webgpu_unary.ts delete mode 100644 tests/deps.ts delete mode 100644 tests/webgpu_binary_test.ts delete mode 100644 tests/webgpu_matmul_test.ts delete mode 100644 tests/webgpu_pad_test.ts delete mode 100644 tests/webgpu_transpose_test.ts delete mode 100644 tests/webgpu_unary_test.ts diff --git a/backend/webgpu/backend.ts b/backend/webgpu/backend.ts index 74f948c..1c1f9ee 100644 --- a/backend/webgpu/backend.ts +++ b/backend/webgpu/backend.ts @@ -47,6 +47,7 @@ export class WebGPUBackend implements Backend { const module = this.device.createShaderModule({ code }); const pipeline = await this.device.createComputePipelineAsync({ compute: { module, entryPoint: "main" }, + layout: "auto", }); const layout = pipeline.getBindGroupLayout(0); @@ -75,8 +76,8 @@ export class WebGPUBackend implements Backend { const passEncoder = commandEncoder.beginComputePass(); passEncoder.setBindGroup(0, bindgroup); passEncoder.setPipeline(pipeline); - passEncoder.dispatch(...workgroups as [number, number, number]); - passEncoder.endPass(); + passEncoder.dispatchWorkgroups(...workgroups as [number, number, number]); + passEncoder.end(); this.device.queue.submit([commandEncoder.finish()]); } diff --git a/backend/webgpu/shaders/binary.ts b/backend/webgpu/shaders/binary.ts index 1a10e83..22ced95 100644 --- a/backend/webgpu/shaders/binary.ts +++ b/backend/webgpu/shaders/binary.ts @@ -6,34 +6,34 @@ import { DataType } from "../../types/data.ts"; export default (type: DataType, expr: string) => ` // #import "prelude.wgsl" - var e: f32 = 2.718281828459045; - var pi: f32 = 3.141592653589793; - var tau: f32 = 6.283185307179586; - var phi: f32 = 1.618033988749895; - var feigd: f32 = 4.66920160910299; - var feiga: f32 = -2.5029078750958926; +var e: f32 = 2.718281828459045; +var pi: f32 = 3.141592653589793; +var tau: f32 = 6.283185307179586; +var phi: f32 = 1.618033988749895; +var feigd: f32 = 4.66920160910299; +var feiga: f32 = -2.5029078750958926; var gauss: f32 = 0.8346268416740732; - // #input type: DataType - // #input expr: string +// #input type: DataType +// #input expr: string - struct Data { - values: array<${type}>; - }; +struct Data { + values: array<${type}>, +}; - [[group(0), binding(0)]] - var a_data: Data; - [[group(0), binding(1)]] - var b_data: Data; - [[group(0), binding(2)]] - var c_data: Data; +@group(0) @binding(0) +var a_data: Data; +@group(0) @binding(1) +var b_data: Data; +@group(0) @binding(2) +var c_data: Data; - fn binary(a: ${type}, b: ${type}) -> ${type} { - ${expr} - } +fn binary(a: ${type}, b: ${type}) -> ${type} { + ${expr} +} - [[stage(compute), workgroup_size(128)]] - fn main([[builtin(global_invocation_id)]] global_id: vec3) { - c_data.values[global_id.x] = binary(a_data.values[global_id.x], b_data.values[global_id.x]); +@compute @workgroup_size(128) +fn main(@builtin(global_invocation_id) global_id: vec3) { + c_data.values[global_id.x] = binary(a_data.values[global_id.x], b_data.values[global_id.x]); } `; diff --git a/backend/webgpu/shaders/binary.wgsl b/backend/webgpu/shaders/binary.wgsl index 11bc8ab..c315bd9 100644 --- a/backend/webgpu/shaders/binary.wgsl +++ b/backend/webgpu/shaders/binary.wgsl @@ -3,21 +3,21 @@ #input expr: string struct Data { - values: array; + values: array, }; -[[group(0), binding(0)]] +@group(0) @binding(0) var a_data: Data; -[[group(0), binding(1)]] +@group(0) @binding(1) var b_data: Data; -[[group(0), binding(2)]] +@group(0) @binding(2) var c_data: Data; fn binary(a: type, b: type) -> type { expr } -[[stage(compute), workgroup_size(128)]] -fn main([[builtin(global_invocation_id)]] global_id: vec3) { +@compute @workgroup_size(128) +fn main(@builtin(global_invocation_id) global_id: vec3) { c_data.values[global_id.x] = binary(a_data.values[global_id.x], b_data.values[global_id.x]); } diff --git a/backend/webgpu/shaders/matmul.ts b/backend/webgpu/shaders/matmul.ts index ed6737e..ad82b0e 100644 --- a/backend/webgpu/shaders/matmul.ts +++ b/backend/webgpu/shaders/matmul.ts @@ -6,45 +6,45 @@ import { DataType } from "../../types/data.ts"; export default (type: DataType) => ` // #import "prelude.wgsl" - var e: f32 = 2.718281828459045; - var pi: f32 = 3.141592653589793; - var tau: f32 = 6.283185307179586; - var phi: f32 = 1.618033988749895; - var feigd: f32 = 4.66920160910299; - var feiga: f32 = -2.5029078750958926; +var e: f32 = 2.718281828459045; +var pi: f32 = 3.141592653589793; +var tau: f32 = 6.283185307179586; +var phi: f32 = 1.618033988749895; +var feigd: f32 = 4.66920160910299; +var feiga: f32 = -2.5029078750958926; var gauss: f32 = 0.8346268416740732; - // #input type: DataType - - struct Uniform { - m: u32; - n: u32; - k: u32; - }; - - struct Data { - values: array<${type}>; - }; - - [[group(0), binding(0)]] - var a: Data; - [[group(0), binding(1)]] - var b: Data; - [[group(0), binding(2)]] - var c: Data; - [[group(0), binding(3)]] - var meta: Uniform; - - [[stage(compute), workgroup_size(8, 8, 1)]] - fn main([[builtin(global_invocation_id)]] global_id: vec3) { - if (global_id.x >= meta.n || global_id.y >= meta.m) { - return; - } - - var sum = ${type}(0); - for (var k = 0u; k < meta.k; k = k + 1u) { - sum = sum + a.values[global_id.y * meta.k + k] * b.values[k * meta.n + global_id.x]; - } - c.values[global_id.x + global_id.y * meta.n] = sum; +// #input type: DataType + +struct Uniforms { + m: u32, + n: u32, + k: u32, +}; + +struct Data { + values: array<${type}>, +}; + +@group(0) @binding(0) +var a: Data; +@group(0) @binding(1) +var b: Data; +@group(0) @binding(2) +var c: Data; +@group(0) @binding(3) +var uniforms: Uniforms; + +@compute @workgroup_size(8, 8, 1) +fn main(@builtin(global_invocation_id) global_id: vec3) { + if (global_id.x >= uniforms.n || global_id.y >= uniforms.m) { + return; + } + + var sum = ${type}(0); + for (var k = 0u; k < uniforms.k; k = k + 1u) { + sum = sum + a.values[global_id.y * uniforms.k + k] * b.values[k * uniforms.n + global_id.x]; + } + c.values[global_id.x + global_id.y * uniforms.n] = sum; } `; diff --git a/backend/webgpu/shaders/matmul.wgsl b/backend/webgpu/shaders/matmul.wgsl index 2a20a00..f77981b 100644 --- a/backend/webgpu/shaders/matmul.wgsl +++ b/backend/webgpu/shaders/matmul.wgsl @@ -1,34 +1,34 @@ #import "prelude.wgsl" #input type: DataType -struct Uniform { - m: u32; - n: u32; - k: u32; +struct Uniforms { + m: u32, + n: u32, + k: u32, }; struct Data { - values: array; + values: array, }; -[[group(0), binding(0)]] +@group(0) @binding(0) var a: Data; -[[group(0), binding(1)]] +@group(0) @binding(1) var b: Data; -[[group(0), binding(2)]] +@group(0) @binding(2) var c: Data; -[[group(0), binding(3)]] -var meta: Uniform; +@group(0) @binding(3) +var uniforms: Uniforms; -[[stage(compute), workgroup_size(8, 8, 1)]] -fn main([[builtin(global_invocation_id)]] global_id: vec3) { - if (global_id.x >= meta.n || global_id.y >= meta.m) { +@compute @workgroup_size(8, 8, 1) +fn main(@builtin(global_invocation_id) global_id: vec3) { + if (global_id.x >= uniforms.n || global_id.y >= uniforms.m) { return; } var sum = type(0); - for (var k = 0u; k < meta.k; k = k + 1u) { - sum = sum + a.values[global_id.y * meta.k + k] * b.values[k * meta.n + global_id.x]; + for (var k = 0u; k < uniforms.k; k = k + 1u) { + sum = sum + a.values[global_id.y * uniforms.k + k] * b.values[k * uniforms.n + global_id.x]; } - c.values[global_id.x + global_id.y * meta.n] = sum; + c.values[global_id.x + global_id.y * uniforms.n] = sum; } diff --git a/backend/webgpu/shaders/pad.ts b/backend/webgpu/shaders/pad.ts index 9f628cc..e42d14b 100644 --- a/backend/webgpu/shaders/pad.ts +++ b/backend/webgpu/shaders/pad.ts @@ -6,39 +6,39 @@ import { DataType } from "../../types/data.ts"; export default (type: DataType) => ` // #import "prelude.wgsl" - var e: f32 = 2.718281828459045; - var pi: f32 = 3.141592653589793; - var tau: f32 = 6.283185307179586; - var phi: f32 = 1.618033988749895; - var feigd: f32 = 4.66920160910299; - var feiga: f32 = -2.5029078750958926; +var e: f32 = 2.718281828459045; +var pi: f32 = 3.141592653589793; +var tau: f32 = 6.283185307179586; +var phi: f32 = 1.618033988749895; +var feigd: f32 = 4.66920160910299; +var feiga: f32 = -2.5029078750958926; var gauss: f32 = 0.8346268416740732; - // #input type: DataType - - struct Uniform { - w: u32; - h: u32; - n: u32; - }; - - struct Data { - values: array<${type}>; - }; - - [[group(0), binding(0)]] - var a: Data; - [[group(0), binding(1)]] - var b: Data; - [[group(0), binding(2)]] - var meta: Uniform; - - [[stage(compute), workgroup_size(8, 8, 1)]] - fn main([[builtin(global_invocation_id)]] global_id: vec3) { - if (global_id.x >= meta.w || global_id.y >= meta.h) { - return; - } +// #input type: DataType + +struct Uniforms { + w: u32, + h: u32, + n: u32, +}; + +struct Data { + values: array<${type}>, +}; + +@group(0) @binding(0) +var a: Data; +@group(0) @binding(1) +var b: Data; +@group(0) @binding(2) +var uniforms: Uniforms; + +@compute @workgroup_size(8, 8, 1) +fn main(@builtin(global_invocation_id) global_id: vec3) { + if (global_id.x >= uniforms.w || global_id.y >= uniforms.h) { + return; + } - b.values[global_id.x + global_id.y * meta.n] = a.values[global_id.x + global_id.y * meta.w]; + b.values[global_id.x + global_id.y * uniforms.n] = a.values[global_id.x + global_id.y * uniforms.w]; } `; diff --git a/backend/webgpu/shaders/pad.wgsl b/backend/webgpu/shaders/pad.wgsl index 30e7904..c0f85a3 100644 --- a/backend/webgpu/shaders/pad.wgsl +++ b/backend/webgpu/shaders/pad.wgsl @@ -1,28 +1,28 @@ #import "prelude.wgsl" #input type: DataType -struct Uniform { - w: u32; - h: u32; - n: u32; +struct Uniforms { + w: u32, + h: u32, + n: u32, }; struct Data { - values: array; + values: array, }; -[[group(0), binding(0)]] +@group(0) @binding(0) var a: Data; -[[group(0), binding(1)]] +@group(0) @binding(1) var b: Data; -[[group(0), binding(2)]] -var meta: Uniform; +@group(0) @binding(2) +var uniforms: Uniforms; -[[stage(compute), workgroup_size(8, 8, 1)]] -fn main([[builtin(global_invocation_id)]] global_id: vec3) { - if (global_id.x >= meta.w || global_id.y >= meta.h) { +@compute @workgroup_size(8, 8, 1) +fn main(@builtin(global_invocation_id) global_id: vec3) { + if (global_id.x >= uniforms.w || global_id.y >= uniforms.h) { return; } - b.values[global_id.x + global_id.y * meta.n] = a.values[global_id.x + global_id.y * meta.w]; + b.values[global_id.x + global_id.y * uniforms.n] = a.values[global_id.x + global_id.y * uniforms.w]; } diff --git a/backend/webgpu/shaders/perlin.ts b/backend/webgpu/shaders/perlin.ts index c44d395..b266c07 100644 --- a/backend/webgpu/shaders/perlin.ts +++ b/backend/webgpu/shaders/perlin.ts @@ -3,84 +3,84 @@ // Make changes to `perlin.wgsl` instead and rebuild export default () => ` - struct Meta { - m: u32; - n: u32; - k: u32; - }; +struct Uniforms { + m: u32, + n: u32, + k: u32, +}; - struct Data { - values: array; - }; +struct Data { + values: array, +}; - struct Data2 { - values: array>; - }; +struct Data2 { + values: array>, +}; - struct Data3 { - values: array; - }; +struct Data3 { + values: array, +}; - [[group(0), binding(0)]] - var a: Data; - [[group(0), binding(1)]] - var b: Data2; - [[group(0), binding(2)]] - var c: Data3; - [[group(0), binding(3)]] - var meta: Meta; +@group(0) @binding(0) +var a: Data; +@group(0) @binding(1) +var b: Data2; +@group(0) @binding(2) +var c: Data3; +@group(0) @binding(3) +var uniforms: Uniforms; - fn lerp(a: f32, b: f32, x: f32) -> f32 { - return a + (b - a) * x; - }; +fn lerp(a: f32, b: f32, x: f32) -> f32 { + return a + (b - a) * x; +}; - fn fade(t: f32) -> f32 { - return ((6.0 * t - 15.0) * t + 10.0) * t * t * t; - }; +fn fade(t: f32) -> f32 { + return ((6.0 * t - 15.0) * t + 10.0) * t * t * t; +}; - fn perlin2D(posX: f32, posY: f32) -> f32 { - let batchX = i32(floor(posX / f32(meta.k))); - let batchY = i32(floor(posY / f32(meta.k))); - - let trp = a.values[a.values[batchX + 1] + batchY + 1]; - let tlp = a.values[a.values[batchX] + batchY + 1]; - let brp = a.values[a.values[batchX + 1] + batchY]; - let blp = a.values[a.values[batchX] + batchY]; +fn perlin2D(posX: f32, posY: f32) -> f32 { + let batchX = i32(floor(posX / f32(uniforms.k))); + let batchY = i32(floor(posY / f32(uniforms.k))); + + let trp = a.values[a.values[batchX + 1] + batchY + 1]; + let tlp = a.values[a.values[batchX] + batchY + 1]; + let brp = a.values[a.values[batchX + 1] + batchY]; + let blp = a.values[a.values[batchX] + batchY]; - let trv = b.values[u32(trp % 3)]; - let tlv = b.values[u32(tlp % 3)]; - let brv = b.values[u32(brp % 3)]; - let blv = b.values[u32(blp % 3)]; + let trv = b.values[u32(trp % 3)]; + let tlv = b.values[u32(tlp % 3)]; + let brv = b.values[u32(brp % 3)]; + let blv = b.values[u32(blp % 3)]; - let x = posX / f32(meta.k) - f32(batchX); - let y = posY / f32(meta.k) - f32(batchY); + let x = posX / f32(uniforms.k) - f32(batchX); + let y = posY / f32(uniforms.k) - f32(batchY); - let tr = vec2(x - 1.0, y - 1.0); - let tl = vec2(x, y - 1.0); - let br = vec2(x - 1.0, y); - let bl = vec2(x, y); + let tr = vec2(x - 1.0, y - 1.0); + let tl = vec2(x, y - 1.0); + let br = vec2(x - 1.0, y); + let bl = vec2(x, y); - let trd = dot(trv, tr); - let tld = dot(tlv, tl); - let brd = dot(brv, br); - let bld = dot(blv, bl); + let trd = dot(trv, tr); + let tld = dot(tlv, tl); + let brd = dot(brv, br); + let bld = dot(blv, bl); - let u = fade(x); - let v = fade(y); + let u = fade(x); + let v = fade(y); - let left = lerp(bld, tld, v); - let right = lerp(brd, trd, v); - return lerp(left, right, u); - }; + let left = lerp(bld, tld, v); + let right = lerp(brd, trd, v); + return lerp(left, right, u); +}; - [[stage(compute), workgroup_size(8, 8, 1)]] - fn main([[builtin(global_invocation_id)]] global_id: vec3) { - if (global_id.x >= meta.n || global_id.y >= meta.m) { - return; - } - - let idx = global_id.x + global_id.y * meta.n; - let perlin = perlin2D(f32(global_id.y), f32(global_id.x)); - c.values[idx] = perlin; +@compute @workgroup_size(8, 8, 1) +fn main(@builtin(global_invocation_id) global_id: vec3) { + if (global_id.x >= uniforms.n || global_id.y >= uniforms.m) { + return; + } + + let idx = global_id.x + global_id.y * uniforms.n; + let perlin = perlin2D(f32(global_id.y) @f32(global_id.x)); + c.values[idx] = perlin; } `; diff --git a/backend/webgpu/shaders/perlin.wgsl b/backend/webgpu/shaders/perlin.wgsl index 4246c5f..d22b251 100644 --- a/backend/webgpu/shaders/perlin.wgsl +++ b/backend/webgpu/shaders/perlin.wgsl @@ -1,29 +1,29 @@ -struct Meta { - m: u32; - n: u32; - k: u32; +struct Uniforms { + m: u32, + n: u32, + k: u32, }; struct Data { - values: array; + values: array, }; struct Data2 { - values: array>; + values: array>, }; struct Data3 { - values: array; + values: array, }; -[[group(0), binding(0)]] +@group(0) @binding(0) var a: Data; -[[group(0), binding(1)]] +@group(0) @binding(1) var b: Data2; -[[group(0), binding(2)]] +@group(0) @binding(2) var c: Data3; -[[group(0), binding(3)]] -var meta: Meta; +@group(0) @binding(3) +var uniforms: Uniforms; fn lerp(a: f32, b: f32, x: f32) -> f32 { return a + (b - a) * x; @@ -34,8 +34,8 @@ fn fade(t: f32) -> f32 { }; fn perlin2D(posX: f32, posY: f32) -> f32 { - let batchX = i32(floor(posX / f32(meta.k))); - let batchY = i32(floor(posY / f32(meta.k))); + let batchX = i32(floor(posX / f32(uniforms.k))); + let batchY = i32(floor(posY / f32(uniforms.k))); let trp = a.values[a.values[batchX + 1] + batchY + 1]; let tlp = a.values[a.values[batchX] + batchY + 1]; @@ -47,8 +47,8 @@ fn perlin2D(posX: f32, posY: f32) -> f32 { let brv = b.values[u32(brp % 3)]; let blv = b.values[u32(blp % 3)]; - let x = posX / f32(meta.k) - f32(batchX); - let y = posY / f32(meta.k) - f32(batchY); + let x = posX / f32(uniforms.k) - f32(batchX); + let y = posY / f32(uniforms.k) - f32(batchY); let tr = vec2(x - 1.0, y - 1.0); let tl = vec2(x, y - 1.0); @@ -68,13 +68,13 @@ fn perlin2D(posX: f32, posY: f32) -> f32 { return lerp(left, right, u); }; -[[stage(compute), workgroup_size(8, 8, 1)]] -fn main([[builtin(global_invocation_id)]] global_id: vec3) { - if (global_id.x >= meta.n || global_id.y >= meta.m) { +@compute @workgroup_size(8, 8, 1) +fn main(@builtin(global_invocation_id) global_id: vec3) { + if (global_id.x >= uniforms.n || global_id.y >= uniforms.m) { return; } - let idx = global_id.x + global_id.y * meta.n; - let perlin = perlin2D(f32(global_id.y), f32(global_id.x)); + let idx = global_id.x + global_id.y * uniforms.n; + let perlin = perlin2D(f32(global_id.y) @f32(global_id.x)); c.values[idx] = perlin; } diff --git a/backend/webgpu/shaders/transpose.ts b/backend/webgpu/shaders/transpose.ts index 6aa5834..731d9ea 100644 --- a/backend/webgpu/shaders/transpose.ts +++ b/backend/webgpu/shaders/transpose.ts @@ -6,38 +6,38 @@ import { DataType } from "../../types/data.ts"; export default (type: DataType) => ` // #import "prelude.wgsl" - var e: f32 = 2.718281828459045; - var pi: f32 = 3.141592653589793; - var tau: f32 = 6.283185307179586; - var phi: f32 = 1.618033988749895; - var feigd: f32 = 4.66920160910299; - var feiga: f32 = -2.5029078750958926; +var e: f32 = 2.718281828459045; +var pi: f32 = 3.141592653589793; +var tau: f32 = 6.283185307179586; +var phi: f32 = 1.618033988749895; +var feigd: f32 = 4.66920160910299; +var feiga: f32 = -2.5029078750958926; var gauss: f32 = 0.8346268416740732; - // #input type: DataType +// #input type: DataType - struct Uniform { - w: u32; - h: u32; - }; +struct Uniforms { + w: u32, + h: u32, +}; - struct Data { - values: array<${type}>; - }; +struct Data { + values: array<${type}>, +}; - [[group(0), binding(0)]] - var a: Data; - [[group(0), binding(1)]] - var b: Data; - [[group(0), binding(2)]] - var meta: Uniform; +@group(0) @binding(0) +var a: Data; +@group(0) @binding(1) +var b: Data; +@group(0) @binding(2) +var uniforms: Uniforms; - [[stage(compute), workgroup_size(8, 8, 1)]] - fn main([[builtin(global_invocation_id)]] global_id: vec3) { - if (global_id.x >= meta.w || global_id.y >= meta.h) { - return; - } +@compute @workgroup_size(8, 8, 1) +fn main(@builtin(global_invocation_id) global_id: vec3) { + if (global_id.x >= uniforms.w || global_id.y >= uniforms.h) { + return; + } - b.values[global_id.y + global_id.x * meta.h] = a.values[global_id.x + global_id.y * meta.w]; + b.values[global_id.y + global_id.x * uniforms.h] = a.values[global_id.x + global_id.y * uniforms.w]; } `; diff --git a/backend/webgpu/shaders/transpose.wgsl b/backend/webgpu/shaders/transpose.wgsl index 69a8f25..0acf985 100644 --- a/backend/webgpu/shaders/transpose.wgsl +++ b/backend/webgpu/shaders/transpose.wgsl @@ -1,27 +1,27 @@ #import "prelude.wgsl" #input type: DataType -struct Uniform { - w: u32; - h: u32; +struct Uniforms { + w: u32, + h: u32, }; struct Data { - values: array; + values: array, }; -[[group(0), binding(0)]] +@group(0) @binding(0) var a: Data; -[[group(0), binding(1)]] +@group(0) @binding(1) var b: Data; -[[group(0), binding(2)]] -var meta: Uniform; +@group(0) @binding(2) +var uniforms: Uniforms; -[[stage(compute), workgroup_size(8, 8, 1)]] -fn main([[builtin(global_invocation_id)]] global_id: vec3) { - if (global_id.x >= meta.w || global_id.y >= meta.h) { +@compute @workgroup_size(8, 8, 1) +fn main(@builtin(global_invocation_id) global_id: vec3) { + if (global_id.x >= uniforms.w || global_id.y >= uniforms.h) { return; } - b.values[global_id.y + global_id.x * meta.h] = a.values[global_id.x + global_id.y * meta.w]; + b.values[global_id.y + global_id.x * uniforms.h] = a.values[global_id.x + global_id.y * uniforms.w]; } diff --git a/backend/webgpu/shaders/unary.ts b/backend/webgpu/shaders/unary.ts index 41ff36b..69d6bb9 100644 --- a/backend/webgpu/shaders/unary.ts +++ b/backend/webgpu/shaders/unary.ts @@ -6,32 +6,32 @@ import { DataType } from "../../types/data.ts"; export default (type: DataType, expr: string) => ` // #import "prelude.wgsl" - var e: f32 = 2.718281828459045; - var pi: f32 = 3.141592653589793; - var tau: f32 = 6.283185307179586; - var phi: f32 = 1.618033988749895; - var feigd: f32 = 4.66920160910299; - var feiga: f32 = -2.5029078750958926; +var e: f32 = 2.718281828459045; +var pi: f32 = 3.141592653589793; +var tau: f32 = 6.283185307179586; +var phi: f32 = 1.618033988749895; +var feigd: f32 = 4.66920160910299; +var feiga: f32 = -2.5029078750958926; var gauss: f32 = 0.8346268416740732; - // #input type: DataType - // #input expr: string +// #input type: DataType +// #input expr: string - struct Data { - values: array<${type}>; - }; +struct Data { + values: array<${type}>, +}; - [[group(0), binding(0)]] - var a_data: Data; - [[group(0), binding(1)]] - var b_data: Data; +@group(0) @binding(0) +var a_data: Data; +@group(0) @binding(1) +var b_data: Data; - fn unary(a: ${type}) -> ${type} { - ${expr} - } +fn unary(a: ${type}) -> ${type} { + ${expr} +} - [[stage(compute), workgroup_size(128)]] - fn main([[builtin(global_invocation_id)]] global_id: vec3) { - b_data.values[global_id.x] = unary(a_data.values[global_id.x]); +@compute @workgroup_size(128) +fn main(@builtin(global_invocation_id) global_id: vec3) { + b_data.values[global_id.x] = unary(a_data.values[global_id.x]); } `; diff --git a/backend/webgpu/shaders/unary.wgsl b/backend/webgpu/shaders/unary.wgsl index ab8b9bd..2004e7d 100644 --- a/backend/webgpu/shaders/unary.wgsl +++ b/backend/webgpu/shaders/unary.wgsl @@ -3,19 +3,19 @@ #input expr: string struct Data { - values: array; + values: array, }; -[[group(0), binding(0)]] +@group(0) @binding(0) var a_data: Data; -[[group(0), binding(1)]] +@group(0) @binding(1) var b_data: Data; fn unary(a: type) -> type { expr } -[[stage(compute), workgroup_size(128)]] -fn main([[builtin(global_invocation_id)]] global_id: vec3) { +@compute @workgroup_size(128) +fn main(@builtin(global_invocation_id) global_id: vec3) { b_data.values[global_id.x] = unary(a_data.values[global_id.x]); } diff --git a/examples/webgpu_binary.ts b/examples/webgpu_binary.ts deleted file mode 100644 index 0b605d9..0000000 --- a/examples/webgpu_binary.ts +++ /dev/null @@ -1,30 +0,0 @@ -import { WebGPUBackend } from "../backend/webgpu/backend.ts"; -import { WebGPUData } from "../backend/webgpu/data.ts"; -import { add, div, mul, sub } from "../backend/webgpu/operators/binary.ts"; - -const backend = new WebGPUBackend(); -await backend.initialize(); - -const a = await WebGPUData.from( - backend, - new Float32Array(16).fill(4), - "vec2", -); -const b = await WebGPUData.from( - backend, - new Float32Array(16).fill(4), - "vec2", -); -const c = new WebGPUData(backend, "vec2", 16); - -await add(backend, a, b, c); -console.log(await c.get()); - -await sub(backend, a, b, c); -console.log(await c.get()); - -await mul(backend, a, b, c); -console.log(await c.get()); - -await div(backend, a, b, c); -console.log(await c.get()); diff --git a/examples/webgpu_matmul.ts b/examples/webgpu_matmul.ts deleted file mode 100644 index fc1f3d0..0000000 --- a/examples/webgpu_matmul.ts +++ /dev/null @@ -1,22 +0,0 @@ -import { WebGPUBackend } from "../backend/webgpu/backend.ts"; -import { WebGPUData } from "../backend/webgpu/data.ts"; -import { matmul } from "../backend/webgpu/operators/matmul.ts"; - -const backend = new WebGPUBackend(); -await backend.initialize(); - -const uniform = { m: 2, n: 2, k: 2, alpha: 1 }; - -const a = await WebGPUData.from( - backend, - new Float32Array(uniform.m * uniform.k).fill(2), -); -const b = await WebGPUData.from( - backend, - new Float32Array(uniform.n * uniform.k).fill(2), -); -const c = new WebGPUData(backend, "f32", uniform.m * uniform.n); - -await matmul(backend, a, b, c, uniform); - -console.log(await c.get()); diff --git a/examples/webgpu_pad.ts b/examples/webgpu_pad.ts deleted file mode 100644 index 4febb22..0000000 --- a/examples/webgpu_pad.ts +++ /dev/null @@ -1,29 +0,0 @@ -import { WebGPUBackend } from "../backend/webgpu/backend.ts"; -import { WebGPUData } from "../backend/webgpu/data.ts"; -import { pad } from "../backend/webgpu/operators/pad.ts"; - -const backend = new WebGPUBackend(); -await backend.initialize(); - -// deno-fmt-ignore -const a = await WebGPUData.from<"f32">( - backend, - new Float32Array([ - 1, 2, 3, - 4, 5, 6, - 7, 8, 9 - ]), -); -const b = new WebGPUData(backend, "f32", 16); - -await pad(backend, a, b, { w: 3, h: 3, t: 4 }); -console.log(await b.get()); - -// 1 2 3 -// 4 5 6 -// 7 8 9 - -// 1 2 3 0 -// 4 5 6 0 -// 7 8 9 0 -// 0 0 0 0 diff --git a/examples/webgpu_perlin.ts b/examples/webgpu_perlin.ts deleted file mode 100644 index aad1c19..0000000 --- a/examples/webgpu_perlin.ts +++ /dev/null @@ -1,135 +0,0 @@ -import { WebGPUBackend } from "../backend/webgpu/backend.ts"; -import { WebGPUData } from "../backend/webgpu/data.ts"; -import { perlin } from "../backend/webgpu/operators/perlin.ts"; -import { PlotWindow } from "https://deno.land/x/plotsaur@v0.2.1/mod.ts"; - -function shuffle(tab: number[]) { - for (let e = tab.length - 1; e > 0; e--) { - const index = Math.round(Math.random() * (e - 1)), - temp = tab[e]; - - tab[e] = tab[index]; - tab[index] = temp; - } -} - -function makePermutation() { - const P = []; - for (let i = 0; i < 256; i++) { - P.push(i); - } - shuffle(P); - for (let i = 0; i < 256; i++) { - P.push(P[i]); - } - - return P; -} - -const perms = makePermutation(); -const vecs: number[] = []; -const constants = [ - { x: 1, y: 1 }, - { x: -1, y: 1 }, - { x: -1, y: -1 }, - { x: 1, y: -1 }, -]; -constants.map((v) => vecs.push(v.x, v.y)); -const meta = { - m: 100, - n: 100, - k: 20, -}; - -const backend = new WebGPUBackend(); -await backend.initialize(); - -const a = await WebGPUData.from(backend, new Int32Array(perms), "i32"); -const b = await WebGPUData.from(backend, new Float32Array(vecs), "vec2"); -const c = new WebGPUData(backend, "f32", meta.m * meta.n); - -await perlin(backend, a, b, c, meta); -const result = await c.get(); -// console.log(result); - -// CPU calculation - -function _calculate() { - const vec = []; - for (let i = 0; i < meta.k; i++) { - for (let j = 0; j < meta.k; j++) { - vec.push(perlin2D(i, j)); - } - } - return vec; -} - -function perlin2D(posX: number, posY: number) { - const batchX = Math.floor(posX / meta.k); - const batchY = Math.floor(posY / meta.k); - const trv = constants[perms[perms[batchX + 1] + batchY + 1] % 3]; - const tlv = constants[perms[perms[batchX] + batchY + 1] % 3]; - const brv = constants[perms[perms[batchX + 1] + batchY] % 3]; - const blv = constants[perms[perms[batchX] + batchY] % 3]; - - const x = posX / meta.k - batchX; - const y = posY / meta.k - batchY; - const tr = { x: x - 1, y: y - 1 }; - const tl = { x: x, y: y - 1 }; - const br = { x: x - 1, y: y }; - const bl = { x: x, y: y }; - - const trd = Dot(trv, tr); - const tld = Dot(tlv, tl); - const brd = Dot(brv, br); - const bld = Dot(blv, bl); - - const u = Fade(x); - const v = Fade(y); - - const left = Lerp(bld, tld, v); - const right = Lerp(brd, trd, v); - return Lerp(left, right, u); -} - -type Vector = { - x: number; - y: number; -}; - -function Dot(a: Vector, b: Vector) { - return a.x * b.x + a.y * b.y; -} - -function Lerp(a: number, b: number, x: number) { - return a + (b - a) * x; -} - -function Fade(t: number) { - return ((6 * t - 15) * t + 10) * t * t * t; -} -// console.log(calculate()) - -const plot = new PlotWindow("Webgpu Perlin", 600, 600); -plot.addPlot({ - xLabelAreaSize: 0, - yLabelAreaSize: 0, - mesh: undefined, - seriesLabel: undefined, -}); -plot.cartesian2D({ - type: "ranged", - x_axis: { start: -1, end: 1 }, - y_axis: { start: -1, end: 1 }, -}); -for (let i = 0; i < meta.m; i++) { - for (let j = 0; j < meta.m; j++) { - const idx = i + j * meta.n; - plot.drawRect({ - style: { r: 0, g: 0, b: 0, a: result[idx] + 0.5 }, - points: [{ x: i * 5, y: j * 5 }, { x: (i + 1) * 5, y: (j + 1) * 5 }], - filled: true, - }); - } -} -plot.show(); diff --git a/examples/webgpu_reduce.ts b/examples/webgpu_reduce.ts deleted file mode 100644 index 425abba..0000000 --- a/examples/webgpu_reduce.ts +++ /dev/null @@ -1,19 +0,0 @@ -// Todo! - -// import { WebGPUBackend } from "../backend/webgpu/backend.ts"; -// import { WebGPUData } from "../backend/webgpu/data.ts"; -// import { reduce } from "../backend/webgpu/operators/reduce.ts"; - -// const backend = new WebGPUBackend(); -// await backend.initialize(); - -// // deno-fmt-ignore -// const a = await WebGPUData.from<"f32">( -// backend, -// new Float32Array(1238).fill(123), -// ); - -// console.log(await reduce("return p + c;")(backend, a)); - -// // what about https://www.w3.org/TR/WGSL/#atomic-type -// // Probably a good idea, but then again, no idea how to use it lol diff --git a/examples/webgpu_transpose.ts b/examples/webgpu_transpose.ts deleted file mode 100644 index 87e335d..0000000 --- a/examples/webgpu_transpose.ts +++ /dev/null @@ -1,27 +0,0 @@ -import { WebGPUBackend } from "../backend/webgpu/backend.ts"; -import { WebGPUData } from "../backend/webgpu/data.ts"; -import { transpose } from "../backend/webgpu/operators/transpose.ts"; - -const backend = new WebGPUBackend(); -await backend.initialize(); - -// deno-fmt-ignore -const a = await WebGPUData.from<"f32">( - backend, - new Float32Array([ - 1, 2, 3, - 4, 5, 6, - 7, 8, 9 - ]), -); -const b = new WebGPUData(backend, "f32", 9); - -await transpose(backend, a, b, { w: 3, h: 3 }); -console.log(await b.get()); - -// 1 2 3 -// 4 5 6 - -// 1 4 -// 2 5 -// 3 6 diff --git a/examples/webgpu_unary.ts b/examples/webgpu_unary.ts deleted file mode 100644 index 6c06f19..0000000 --- a/examples/webgpu_unary.ts +++ /dev/null @@ -1,15 +0,0 @@ -import { WebGPUBackend } from "../backend/webgpu/backend.ts"; -import { WebGPUData } from "../backend/webgpu/data.ts"; -import { relu } from "../backend/webgpu/operators/unary.ts"; - -const backend = new WebGPUBackend(); -await backend.initialize(); - -const a = await WebGPUData.from<"f32">( - backend, - new Float32Array(10).fill(-1), -); -const b = new WebGPUData(backend, "f32", 10); - -await relu(backend, a, b); -console.log(await b.get()); diff --git a/tests/deps.ts b/tests/deps.ts deleted file mode 100644 index e4e80ee..0000000 --- a/tests/deps.ts +++ /dev/null @@ -1 +0,0 @@ -export { assertEquals } from "https://deno.land/std@0.125.0/testing/asserts.ts"; diff --git a/tests/webgpu_binary_test.ts b/tests/webgpu_binary_test.ts deleted file mode 100644 index 0ada0dd..0000000 --- a/tests/webgpu_binary_test.ts +++ /dev/null @@ -1,59 +0,0 @@ -import { WebGPUBackend } from "../backend/webgpu/backend.ts"; -import { WebGPUData } from "../backend/webgpu/data.ts"; -import { add, div, mul, sub } from "../backend/webgpu/operators/binary.ts"; -import { assertEquals } from "./deps.ts"; - -const backend = new WebGPUBackend(); -await backend.initialize(); - -const a = await WebGPUData.from( - backend, - new Float32Array(16).fill(4), - "vec2", -); -const b = await WebGPUData.from( - backend, - new Float32Array(16).fill(4), - "vec2", -); -const c = new WebGPUData(backend, "vec2", 16); - -Deno.test({ - name: "Binary Add", - async fn() { - await add(backend, a, b, c); - const expected = new Float32Array(16).fill(8); - assertEquals(await c.get(), expected); - }, - sanitizeResources: false, -}); - -Deno.test({ - name: "Binary Subtract", - async fn() { - await sub(backend, a, b, c); - const expected = new Float32Array(16).fill(0); - assertEquals(await c.get(), expected); - }, - sanitizeResources: false, -}); - -Deno.test({ - name: "Binary Multiply", - async fn() { - await mul(backend, a, b, c); - const expected = new Float32Array(16).fill(16); - assertEquals(await c.get(), expected); - }, - sanitizeResources: false, -}); - -Deno.test({ - name: "Binary Divide", - async fn() { - await div(backend, a, b, c); - const expected = new Float32Array(16).fill(1); - assertEquals(await c.get(), expected); - }, - sanitizeResources: false, -}); diff --git a/tests/webgpu_matmul_test.ts b/tests/webgpu_matmul_test.ts deleted file mode 100644 index a6d149c..0000000 --- a/tests/webgpu_matmul_test.ts +++ /dev/null @@ -1,29 +0,0 @@ -import { WebGPUBackend } from "../backend/webgpu/backend.ts"; -import { WebGPUData } from "../backend/webgpu/data.ts"; -import { matmul } from "../backend/webgpu/operators/matmul.ts"; -import { assertEquals } from "./deps.ts"; - -const backend = new WebGPUBackend(); -await backend.initialize(); - -const uniform = { m: 2, n: 2, k: 2, alpha: 1 }; - -const a = await WebGPUData.from( - backend, - new Float32Array(uniform.m * uniform.k).fill(2), -); -const b = await WebGPUData.from( - backend, - new Float32Array(uniform.n * uniform.k).fill(2), -); -const c = new WebGPUData(backend, "f32", uniform.m * uniform.n); - -Deno.test({ - name: "Matrix Multiply", - async fn() { - await matmul(backend, a, b, c, uniform); - const expected = new Float32Array(4).fill(8); - assertEquals(await c.get(), expected); - }, - sanitizeResources: false, -}); diff --git a/tests/webgpu_pad_test.ts b/tests/webgpu_pad_test.ts deleted file mode 100644 index 1180177..0000000 --- a/tests/webgpu_pad_test.ts +++ /dev/null @@ -1,42 +0,0 @@ -import { WebGPUBackend } from "../backend/webgpu/backend.ts"; -import { WebGPUData } from "../backend/webgpu/data.ts"; -import { pad } from "../backend/webgpu/operators/pad.ts"; -import { assertEquals } from "./deps.ts"; - -const backend = new WebGPUBackend(); -await backend.initialize(); - -// deno-fmt-ignore -const a = await WebGPUData.from<"f32">( - backend, - new Float32Array([ - 1, 2, 3, - 4, 5, 6, - 7, 8, 9 - ]), -); -const b = new WebGPUData(backend, "f32", 16); - -// deno-fmt-ignore -Deno.test({ - name: "Matrix Pad", - async fn() { - await pad(backend, a, b, { w: 3, h: 3, t: 4 }); - const expected = new Float32Array([ - 1, 2, 3, 0, 4, 5, - 6, 0, 7, 8, 9, 0, - 0, 0, 0, 0 - ]) - assertEquals(await b.get(), expected); - }, - sanitizeResources: false -}); - -// 1 2 3 -// 4 5 6 -// 7 8 9 - -// 1 2 3 0 -// 4 5 6 0 -// 7 8 9 0 -// 0 0 0 0 diff --git a/tests/webgpu_transpose_test.ts b/tests/webgpu_transpose_test.ts deleted file mode 100644 index 5982976..0000000 --- a/tests/webgpu_transpose_test.ts +++ /dev/null @@ -1,39 +0,0 @@ -import { WebGPUBackend } from "../backend/webgpu/backend.ts"; -import { WebGPUData } from "../backend/webgpu/data.ts"; -import { transpose } from "../backend/webgpu/operators/transpose.ts"; -import { assertEquals } from "./deps.ts"; - -const backend = new WebGPUBackend(); -await backend.initialize(); - -// deno-fmt-ignore -const a = await WebGPUData.from<"f32">( - backend, - new Float32Array([ - 1, 2, 3, - 4, 5, 6, - 7, 8, 9 - ]), -); -const b = new WebGPUData(backend, "f32", 9); - -// deno-fmt-ignore -Deno.test({ - name: "Matrix Transpose", - async fn() { - await transpose(backend, a, b, { w: 3, h: 3 }); - const expected = new Float32Array([ - 1, 4, 7, 2, 5, - 8, 3, 6, 9 - ]) - assertEquals(await b.get(), expected); - }, - sanitizeResources: false -}); - -// 1 2 3 -// 4 5 6 - -// 1 4 -// 2 5 -// 3 6 diff --git a/tests/webgpu_unary_test.ts b/tests/webgpu_unary_test.ts deleted file mode 100644 index 00beba3..0000000 --- a/tests/webgpu_unary_test.ts +++ /dev/null @@ -1,23 +0,0 @@ -import { WebGPUBackend } from "../backend/webgpu/backend.ts"; -import { WebGPUData } from "../backend/webgpu/data.ts"; -import { relu } from "../backend/webgpu/operators/unary.ts"; -import { assertEquals } from "./deps.ts"; - -const backend = new WebGPUBackend(); -await backend.initialize(); - -const a = await WebGPUData.from<"f32">( - backend, - new Float32Array(10).fill(-1), -); -const b = new WebGPUData(backend, "f32", 10); - -Deno.test({ - name: "Unary ReLU", - async fn() { - await relu(backend, a, b); - const expected = new Float32Array(10).fill(0); - assertEquals(await b.get(), expected); - }, - sanitizeResources: false, -});