Merge branch 'main' into gh-pages

sona1111 · Dec 9, 2023 · 7b868dd · 7b868dd
2 parents 71a8227 + 570ebce
commit 7b868dd
Show file tree

Hide file tree

Showing 26 changed files with 355 additions and 332 deletions.
diff --git a/README.md b/README.md
@@ -3,9 +3,8 @@ WebGPU Image Super Resolution
 
 **University of Pennsylvania, CIS 565: GPU Programming and Architecture, Final Project**
 
-* Paul (San) Jewell
-  * [LinkedIn](https://www.linkedin.com/in/paul-jewell-2aba7379), [work website](
-    https://www.biociphers.org/paul-jewell-lab-member), [personal website](https://gitlab.com/inklabapp), [twitter](https://twitter.com/inklabapp), etc.
+* Sanura Jewell
+  * [LinkedIn](https://www.linkedin.com/in/san-jewell-2aba7379/)
 * Tested on: Linux pop-os 5.11.0-7614-generic, i7-9750H CPU @ 2.60GHz 32GB, GeForce GTX 1650 Mobile / Max-Q 4GB
 
 * Yuxuan Zhu

diff --git a/past_iterations/cnn.js b/past_iterations/cnn.js
@@ -74,24 +74,25 @@
 
     const shaderModule = device.createShaderModule({
       code: `
-      [[block]] struct Matrix {
+      struct Matrix {
         size : vec4<f32>; // batch_size , channel_size , height , width
-        numbers: array<f32>;
+        numbers: array<f32>,
     };
     
-    [[block]] struct Array {
+    struct Array {
         size : f32; // channel_size
-        numbers : array<f32>; 
+        numbers : array<f32>,
     };
     
-    [[group(0), binding(0)]] var<storage, read> inputImage : Matrix;
-    [[group(0), binding(1)]] var<storage, read> inputKernel : Matrix;
-    [[group(0), binding(2)]] var<storage, read> inputBias : Array;
+    @group(0) @binding(0) var<storage, read> inputImage : Matrix;
+    @group(0) @binding(1) var<storage, read> inputKernel : Matrix;
+    @group(0) @binding(2) var<storage, read> inputBias : Array;
     [[group(0), binding(3)]] var<storage, write> resultImage : Matrix;
     
     
-    [[stage(compute), workgroup_size(4, 4, 4)]]
-    fn main([[builtin(global_invocation_id)]] global_id : vec3<u32>) {
+    @compute
+@workgroup_size(4, 4, 4)
+    fn main(@builtin(global_invocation_id) global_id: vec3u) {
         // Guard against out-of-bounds work group sizes.
         if (global_id.x >= u32(inputImage.size.w) || global_id.y >= u32(inputImage.size.z) || global_id.z >= u32(inputKernel.size.x)) {
             return;

diff --git a/past_iterations/interpolateTest.js b/past_iterations/interpolateTest.js
@@ -44,22 +44,23 @@
 
     const shaderModule = device.createShaderModule({
       code: `
-      [[block]] struct Matrix {
+      struct Matrix {
         size : vec4<f32>; // batch_size , channel_size , height , width
-        numbers: array<f32>;
+        numbers: array<f32>,
     };
     
-    [[block]] struct Array {
+    struct Array {
         size : f32; // channel_size
-        numbers : array<f32>; 
+        numbers : array<f32>,
     };
     
-    [[group(0), binding(0)]] var<storage, read> inputImage : Matrix;
-    [[group(0), binding(1)]] var<storage, write> resultImage : Matrix;
+    @group(0) @binding(0) var<storage, read> inputImage : Matrix;
+    @group(0) @binding(1) var<storage, write> resultImage : Matrix;
     
     
-    [[stage(compute), workgroup_size(4, 4, 4)]]
-    fn main([[builtin(global_invocation_id)]] global_id : vec3<u32>) {
+    @compute
+@workgroup_size(4, 4, 4)
+    fn main(@builtin(global_invocation_id) global_id: vec3u) {
         // Guard against out-of-bounds work group sizes.
         if (global_id.x >= u32(inputImage.size.w) || global_id.y >= u32(inputImage.size.z) || global_id.z >= u32(inputImage.size.y)) {
             return;

diff --git a/past_iterations/reluTest.js b/past_iterations/reluTest.js
@@ -34,17 +34,18 @@
 
     const shaderModule = device.createShaderModule({
       code: `
-      [[block]] struct Matrix {
+      struct Matrix {
         size : vec4<f32>; // batch_size , channel_size , height , width
-        numbers: array<f32>;
+        numbers: array<f32>,
     };
     
     
-    [[group(0), binding(0)]] var<storage, read_write> inputImage : Matrix;
+    @group(0) @binding(0) var<storage, read_write> inputImage : Matrix;
     
     
-    [[stage(compute), workgroup_size(4, 4, 4)]]
-    fn main([[builtin(global_invocation_id)]] global_id : vec3<u32>) {
+    @compute
+@workgroup_size(4, 4, 4)
+    fn main(@builtin(global_invocation_id) global_id: vec3u) {
         // Guard against out-of-bounds work group sizes.
         if (global_id.x >= u32(inputImage.size.w) || global_id.y >= u32(inputImage.size.z) || global_id.z >= u32(inputImage.size.y)) {
             return;

diff --git a/rrbdnetC_v7.js b/rrbdnetC_v7.js
@@ -247,11 +247,11 @@ async function run_nn(input_elem, output_elem, status_elem, gpumem_elem, progres
         const passEncoder = commandEncoder.beginComputePass();
         passEncoder.setPipeline(computePipeline);
         passEncoder.setBindGroup(0, bindGroup);
-        const x = Math.ceil(inputshape[1] / 4); // X dimension of the grid of workgroups to dispatch.
-        const y = Math.ceil(inputshape[2] / 4); // Y dimension of the grid of workgroups to dispatch.
+        const x = Math.ceil(inputshape[1] / 4); // X dimension of the grid of workgroups to dispatchWorkgroups.
+        const y = Math.ceil(inputshape[2] / 4); // Y dimension of the grid of workgroups to dispatchWorkgroups.
         const z = Math.ceil(weightshape[0] / 4); 
-        passEncoder.dispatch(x, y, z);
-        passEncoder.endPass();
+        passEncoder.dispatchWorkgroups(x, y, z);
+        passEncoder.end();
 
         // Submit GPU commands.
         const gpuCommands = commandEncoder.finish();
@@ -369,11 +369,11 @@ async function run_nn(input_elem, output_elem, status_elem, gpumem_elem, progres
         const passEncoder = commandEncoder.beginComputePass();
         passEncoder.setPipeline(computePipeline);
         passEncoder.setBindGroup(0, bindGroup);
-        const x = Math.ceil(inputshape[1] / 4); // X dimension of the grid of workgroups to dispatch.
-        const y = Math.ceil(inputshape[2] / 4); // Y dimension of the grid of workgroups to dispatch.
+        const x = Math.ceil(inputshape[1] / 4); // X dimension of the grid of workgroups to dispatchWorkgroups.
+        const y = Math.ceil(inputshape[2] / 4); // Y dimension of the grid of workgroups to dispatchWorkgroups.
         const z = Math.ceil(weightshape[0] / 4); 
-        passEncoder.dispatch(x, y, z);
-        passEncoder.endPass();
+        passEncoder.dispatchWorkgroups(x, y, z);
+        passEncoder.end();
 
         // Submit GPU commands.
         const gpuCommands = commandEncoder.finish();
@@ -543,11 +543,11 @@ async function run_nn(input_elem, output_elem, status_elem, gpumem_elem, progres
         const passEncoder = commandEncoder.beginComputePass();
         passEncoder.setPipeline(computePipeline);
         passEncoder.setBindGroup(0, bindGroup);
-        const x = Math.ceil(inputSize[2] / 4); // X dimension of the grid of workgroups to dispatch.
-        const y = Math.ceil(inputSize[1] / 4); // Y dimension of the grid of workgroups to dispatch.
+        const x = Math.ceil(inputSize[2] / 4); // X dimension of the grid of workgroups to dispatchWorkgroups.
+        const y = Math.ceil(inputSize[1] / 4); // Y dimension of the grid of workgroups to dispatchWorkgroups.
         const z = Math.ceil(inputSize[0] / 4);
-        passEncoder.dispatch(x, y, z);
-        passEncoder.endPass();
+        passEncoder.dispatchWorkgroups(x, y, z);
+        passEncoder.end();
         const gpuCommands = commandEncoder.finish();
         device.queue.submit([gpuCommands]);
         freeGPUArray(unifbuffer);
@@ -627,11 +627,11 @@ async function run_nn(input_elem, output_elem, status_elem, gpumem_elem, progres
         const passEncoder = commandEncoder.beginComputePass();
         passEncoder.setPipeline(computePipeline);
         passEncoder.setBindGroup(0, bindGroup);
-        const x = Math.ceil(matrixSize[2] / 4); // X dimension of the grid of workgroups to dispatch.
-        const y = Math.ceil(matrixSize[1] / 4); // Y dimension of the grid of workgroups to dispatch.
+        const x = Math.ceil(matrixSize[2] / 4); // X dimension of the grid of workgroups to dispatchWorkgroups.
+        const y = Math.ceil(matrixSize[1] / 4); // Y dimension of the grid of workgroups to dispatchWorkgroups.
         const z = Math.ceil(matrixSize[0] / 4); 
-        passEncoder.dispatch(x, y, z);
-        passEncoder.endPass();
+        passEncoder.dispatchWorkgroups(x, y, z);
+        passEncoder.end();
         const gpuCommands = commandEncoder.finish();
         device.queue.submit([gpuCommands]);
         freeGPUArray(unifbuffer);
@@ -713,11 +713,11 @@ async function run_nn(input_elem, output_elem, status_elem, gpumem_elem, progres
         const passEncoder = commandEncoder.beginComputePass();
         passEncoder.setPipeline(computePipeline);
         passEncoder.setBindGroup(0, bindGroup);
-        const x = Math.ceil(matrixSize[2] / 4); // X dimension of the grid of workgroups to dispatch.
-        const y = Math.ceil(matrixSize[1] / 4); // Y dimension of the grid of workgroups to dispatch.
+        const x = Math.ceil(matrixSize[2] / 4); // X dimension of the grid of workgroups to dispatchWorkgroups.
+        const y = Math.ceil(matrixSize[1] / 4); // Y dimension of the grid of workgroups to dispatchWorkgroups.
         const z = Math.ceil(matrixSize[0] / 4);
-        passEncoder.dispatch(x, y, z);
-        passEncoder.endPass();
+        passEncoder.dispatchWorkgroups(x, y, z);
+        passEncoder.end();
         const gpuCommands = commandEncoder.finish();
         device.queue.submit([gpuCommands]);
 

diff --git a/shaders_f32/addbias.wgsl b/shaders_f32/addbias.wgsl
@@ -1,27 +1,27 @@
 
-[[block]] struct Array {
+struct Array {
 
-    numbers : array<f32>; 
+    numbers : array<f32>,
 };
 
 
-[[block]] struct UBO {
-  channelIdxs: array<i32, 2>;
-  outputSizes: array<i32, 2>;
-  inputSizes: array<i32, 3>;
-  kernSizes: array<i32, 4>;
+struct UBO {
+  channelIdxs: array<i32, 2>,
+  outputSizes: array<i32, 2>,
+  inputSizes: array<i32, 3>,
+  kernSizes: array<i32, 4>,
 };
 
-[[group(0), binding(0)]] var<storage, read> inputImage : Array;
-[[group(0), binding(1)]] var<storage, read> inputKernel : Array;
-[[group(0), binding(2)]] var<storage, read> inputBias : Array;
+@group(0) @binding(0) var<storage, read> inputImage : Array;
+@group(0) @binding(1) var<storage, read> inputKernel : Array;
+@group(0) @binding(2) var<storage, read> inputBias : Array;
 [[group(0), binding(3)]] var<storage, read_write> resultImage : Array;
 [[group(0), binding(4)]] var<storage, read> ufs : UBO;
 
 
 
 [[stage(compute), workgroup_size(4, 4)]]
-fn main([[builtin(global_invocation_id)]] global_id : vec3<u32>) {
+fn main(@builtin(global_invocation_id) global_id: vec3u) {
     // Guard against out-of-bounds work group sizes.
 
     if (global_id.x >= u32(ufs.outputSizes[0]) || global_id.y >= u32(ufs.outputSizes[1])) {

diff --git a/shaders_f32/addition.wgsl b/shaders_f32/addition.wgsl
@@ -1,21 +1,22 @@
-[[block]] struct Matrix {
-    numbers: array<f32>;
+struct Matrix {
+    numbers: array<f32>,
 };
 
-[[block]] struct UBO {
-    inputSizesX: u32;
-  inputSizesY: u32;
-  inputSizesZ: u32;
+struct UBO {
+    inputSizesX: u32,
+  inputSizesY: u32,
+  inputSizesZ: u32,
 };
 
 
-[[group(0), binding(0)]] var<storage, read> inputImage : Matrix;
-[[group(0), binding(1)]] var<storage, read_write> outputImage : Matrix;
-[[group(0), binding(2)]] var<storage, read> ufs : UBO;
+@group(0) @binding(0) var<storage, read> inputImage : Matrix;
+@group(0) @binding(1) var<storage, read_write> outputImage : Matrix;
+@group(0) @binding(2) var<storage, read> ufs : UBO;
 
 
-[[stage(compute), workgroup_size(4, 4, 4)]]
-fn main([[builtin(global_invocation_id)]] global_id : vec3<u32>) {
+@compute
+@workgroup_size(4, 4, 4)
+fn main(@builtin(global_invocation_id) global_id: vec3u) {
     // Guard against out-of-bounds work group sizes.
     if (global_id.x >= ufs.inputSizesZ || global_id.y >= ufs.inputSizesY || global_id.z >= ufs.inputSizesX) {
         return;

diff --git a/shaders_f32/conv2d.wgsl b/shaders_f32/conv2d.wgsl
@@ -1,21 +1,22 @@
-[[block]] struct Matrix {
+struct Matrix {
     size : vec4<f32>; // batch_size , channel_size , height , width
-    numbers: array<f32>;
+    numbers: array<f32>,
 };
 
-[[block]] struct Array {
+struct Array {
     size : f32; // channel_size
-    numbers : array<f32>; 
+    numbers : array<f32>,
 };
 
-[[group(0), binding(0)]] var<storage, read> inputImage : Matrix;
-[[group(0), binding(1)]] var<storage, read> inputKernel : Matrix;
-[[group(0), binding(2)]] var<storage, read> inputBias : Array;
-[[group(0), binding(3)]] var<storage, read_write> resultImage : Matrix;
+@group(0) @binding(0) var<storage, read> inputImage : Matrix;
+@group(0) @binding(1) var<storage, read> inputKernel : Matrix;
+@group(0) @binding(2) var<storage, read> inputBias : Array;
+@group(0) @binding(3) var<storage, read_write> resultImage : Matrix;
 
 
-[[stage(compute), workgroup_size(4, 4, 4)]]
-fn main([[builtin(global_invocation_id)]] global_id : vec3<u32>) {
+@compute
+@workgroup_size(4, 4, 4)
+fn main(@builtin(global_invocation_id) global_id: vec3u) {
     // Guard against out-of-bounds work group sizes.
     if (global_id.x >= u32(inputImage.size.w) || global_id.y >= u32(inputImage.size.z) || global_id.z >= u32(inputKernel.size.x)) {
         return;

diff --git a/shaders_f32/conv2d_allch.wgsl b/shaders_f32/conv2d_allch.wgsl
@@ -1,28 +1,29 @@
 
-[[block]] struct Array {
+struct Array {
 
-    numbers : array<f32>; 
+    numbers : array<f32>,
 };
 
 
-[[block]] struct UBO {
-  outputSizes: array<i32, 2>;
-  inputSizes: array<i32, 3>;
-  kernSizes: array<i32, 4>;
-  offsetw: u32;
-  offsetb: i32;
+struct UBO {
+  outputSizes: array<i32, 2>,
+  inputSizes: array<i32, 3>,
+  kernSizes: array<i32, 4>,
+  offsetw: u32,
+  offsetb: i32,
 };
 
-[[group(0), binding(0)]] var<storage, read> inputImage : Array;
-[[group(0), binding(1)]] var<storage, read> inputKernel : Array;
-[[group(0), binding(2)]] var<storage, read> inputBias : Array;
+@group(0) @binding(0) var<storage, read> inputImage : Array;
+@group(0) @binding(1) var<storage, read> inputKernel : Array;
+@group(0) @binding(2) var<storage, read> inputBias : Array;
 [[group(0), binding(3)]] var<storage, read_write> resultImage : Array;
 [[group(0), binding(4)]] var<storage, read> ufs : UBO;
 
 
 
-[[stage(compute), workgroup_size(4, 4, 4)]]
-fn main([[builtin(global_invocation_id)]] global_id : vec3<u32>) {
+@compute
+@workgroup_size(4, 4, 4)
+fn main(@builtin(global_invocation_id) global_id: vec3u) {
     // Guard against out-of-bounds work group sizes.
 
     if (global_id.x >= u32(ufs.outputSizes[0]) || global_id.y >= u32(ufs.outputSizes[1]) || global_id.z >= u32(ufs.kernSizes[0])) {

diff --git a/shaders_f32/conv2d_allch_rrdb.wgsl b/shaders_f32/conv2d_allch_rrdb.wgsl
@@ -1,29 +1,30 @@
 
-[[block]] struct Array {
+struct Array {
 
-    numbers : array<f32>; 
+    numbers : array<f32>,
 };
 
 
-[[block]] struct UBO {
-  outputSizes: array<i32, 2>;
-  inputSizes: array<i32, 3>;
-  kernSizes: array<i32, 4>;
-  offsetw: u32;
-  offsetb: i32;
-  inputOffset: u32;
-  outputOffset: i32;
+struct UBO {
+  outputSizes: array<i32, 2>,
+  inputSizes: array<i32, 3>,
+  kernSizes: array<i32, 4>,
+  offsetw: u32,
+  offsetb: i32,
+  inputOffset: u32,
+  outputOffset: i32,
 };
 
-[[group(0), binding(0)]] var<storage, read> inputKernel : Array;
-[[group(0), binding(1)]] var<storage, read> inputBias : Array;
-[[group(0), binding(2)]] var<storage, read_write> resultImage : Array;
-[[group(0), binding(3)]] var<storage, read> ufs : UBO;
+@group(0) @binding(0) var<storage, read> inputKernel : Array;
+@group(0) @binding(1) var<storage, read> inputBias : Array;
+@group(0) @binding(2) var<storage, read_write> resultImage : Array;
+@group(0) @binding(3) var<storage, read> ufs : UBO;
 
 
 
-[[stage(compute), workgroup_size(4, 4, 4)]]
-fn main([[builtin(global_invocation_id)]] global_id : vec3<u32>) {
+@compute
+@workgroup_size(4, 4, 4)
+fn main(@builtin(global_invocation_id) global_id: vec3u) {
     // Guard against out-of-bounds work group sizes.
 
     if (global_id.x >= u32(ufs.outputSizes[0]) || global_id.y >= u32(ufs.outputSizes[1]) || global_id.z >= u32(ufs.kernSizes[0])) {