feat: WebGPU API (#7977)

Co-authored-by: Luca Casonato <[email protected]>

cli/build.rs

@@ -11,6 +11,7 @@ use deno_core::RuntimeOptions;
 use deno_runtime::deno_crypto;
 use deno_runtime::deno_fetch;
 use deno_runtime::deno_web;
+use deno_runtime::deno_webgpu;
 use deno_runtime::deno_websocket;
 use regex::Regex;
 use std::collections::HashMap;
@@ -62,6 +63,7 @@ fn create_compiler_snapshot(
   let mut op_crate_libs = HashMap::new();
   op_crate_libs.insert("deno.web", deno_web::get_declaration());
   op_crate_libs.insert("deno.fetch", deno_fetch::get_declaration());
+  op_crate_libs.insert("deno.webgpu", deno_webgpu::get_declaration());
   op_crate_libs.insert("deno.websocket", deno_websocket::get_declaration());
   op_crate_libs.insert("deno.crypto", deno_crypto::get_declaration());
 
@@ -260,6 +262,10 @@ fn main() {
     "cargo:rustc-env=DENO_FETCH_LIB_PATH={}",
     deno_fetch::get_declaration().display()
   );
+  println!(
+    "cargo:rustc-env=DENO_WEBGPU_LIB_PATH={}",
+    deno_webgpu::get_declaration().display()
+  );
   println!(
     "cargo:rustc-env=DENO_WEBSOCKET_LIB_PATH={}",
     deno_websocket::get_declaration().display()

cli/dts/lib.deno.window.d.ts

@@ -3,6 +3,7 @@
 /// <reference no-default-lib="true" />
 /// <reference lib="deno.ns" />
 /// <reference lib="deno.shared_globals" />
+/// <reference lib="deno.webgpu" />
 /// <reference lib="esnext" />
 
 declare class Window extends EventTarget {
@@ -17,12 +18,18 @@ declare class Window extends EventTarget {
   confirm: (message?: string) => boolean;
   prompt: (message?: string, defaultValue?: string) => string | null;
   Deno: typeof Deno;
+  navigator: Navigator;
 }
 
 declare var window: Window & typeof globalThis;
 declare var self: Window & typeof globalThis;
 declare var onload: ((this: Window, ev: Event) => any) | null;
 declare var onunload: ((this: Window, ev: Event) => any) | null;
+declare var navigator: Navigator;
+
+declare interface Navigator {
+  readonly gpu: GPU;
+}
 
 /**
  * Shows the given message and waits for the enter key pressed.

cli/dts/lib.deno.worker.d.ts

@@ -3,6 +3,7 @@
 /// <reference no-default-lib="true" />
 /// <reference lib="deno.ns" />
 /// <reference lib="deno.shared_globals" />
+/// <reference lib="deno.webgpu" />
 /// <reference lib="esnext" />
 
 declare class WorkerGlobalScope {
@@ -29,6 +30,13 @@ declare class WorkerGlobalScope {
   close: () => void;
   postMessage: (message: any) => void;
   Deno: typeof Deno;
+  navigator: WorkerNavigator;
+}
+
+declare var navigator: WorkerNavigator;
+
+declare interface WorkerNavigator {
+  readonly gpu: GPU;
 }
 
 declare class DedicatedWorkerGlobalScope extends WorkerGlobalScope {

cli/flags.rs

@@ -234,6 +234,7 @@ static ENV_VARIABLES_HELP: &str = r#"ENVIRONMENT VARIABLES:
     DENO_DIR             Set the cache directory
     DENO_INSTALL_ROOT    Set deno install's output directory
                          (defaults to $HOME/.deno/bin)
+    DENO_WEBGPU_TRACE    Directory to use for wgpu traces
     HTTP_PROXY           Proxy address for HTTP requests
                          (module downloads, fetch)
     HTTPS_PROXY          Proxy address for HTTPS requests

cli/main.rs

@@ -278,10 +278,11 @@ fn print_cache_info(
 
 pub fn get_types(unstable: bool) -> String {
   let mut types = format!(
-    "{}\n{}\n{}\n{}\n{}\n{}",
+    "{}\n{}\n{}\n{}\n{}\n{}\n{}",
     crate::tsc::DENO_NS_LIB,
     crate::tsc::DENO_WEB_LIB,
     crate::tsc::DENO_FETCH_LIB,
+    crate::tsc::DENO_WEBGPU_LIB,
     crate::tsc::DENO_WEBSOCKET_LIB,
     crate::tsc::SHARED_GLOBALS_LIB,
     crate::tsc::WINDOW_LIB,
@@ -1022,6 +1023,8 @@ fn init_logger(maybe_level: Option<Level>) {
   )
   // https://github.com/denoland/deno/issues/6641
   .filter_module("rustyline", LevelFilter::Off)
+  // wgpu backend crates (gfx_backend), have a lot of useless INFO and WARN logs
+  .filter_module("gfx", LevelFilter::Error)
   .format(|buf, record| {
     let mut target = record.target().to_string();
     if let Some(line_no) = record.line() {

cli/tests/unit/unit_tests.ts

@@ -76,3 +76,4 @@ import "./write_text_file_test.ts";
 import "./performance_test.ts";
 import "./version_test.ts";
 import "./websocket_test.ts";
+import "./webgpu_test.ts";

cli/tests/unit/webgpu_test.ts

@@ -0,0 +1,225 @@
+// TODO(lucacasonato): remove when GPUBufferUsage and friends are added to dlint
+// deno-lint-ignore-file no-undef
+
+import { assert, assertEquals, unitTest } from "./test_util.ts";
+
+let isCI: boolean;
+try {
+  isCI = (Deno.env.get("CI")?.length ?? 0) > 0;
+} catch {
+  isCI = true;
+}
+
+// Skip this test on linux CI, because the vulkan emulator is not good enough
+// yet, and skip on macOS because these do not have virtual GPUs.
+unitTest({
+  perms: { read: true, env: true },
+  ignore: (Deno.build.os === "linux" || Deno.build.os === "darwin") && isCI,
+}, async function webgpuComputePass() {
+  const adapter = await navigator.gpu.requestAdapter();
+  assert(adapter);
+
+  const numbers = [1, 4, 3, 295];
+
+  const device = await adapter.requestDevice();
+  assert(device);
+
+  const shaderCode = await Deno.readTextFile(
+    "cli/tests/webgpu_computepass_shader.wgsl",
+  );
+
+  const shaderModule = device.createShaderModule({
+    code: shaderCode,
+  });
+
+  const size = new Uint32Array(numbers).byteLength;
+
+  const stagingBuffer = device.createBuffer({
+    size: size,
+    usage: 1 | 8,
+  });
+
+  const storageBuffer = device.createBuffer({
+    label: "Storage Buffer",
+    size: size,
+    usage: 0x80 | 8 | 4,
+    mappedAtCreation: true,
+  });
+
+  const buf = new Uint32Array(storageBuffer.getMappedRange());
+
+  buf.set(numbers);
+
+  storageBuffer.unmap();
+
+  const bindGroupLayout = device.createBindGroupLayout({
+    entries: [
+      {
+        binding: 0,
+        visibility: 4,
+        buffer: {
+          type: "storage",
+          minBindingSize: 4,
+        },
+      },
+    ],
+  });
+
+  const bindGroup = device.createBindGroup({
+    layout: bindGroupLayout,
+    entries: [
+      {
+        binding: 0,
+        resource: {
+          buffer: storageBuffer,
+        },
+      },
+    ],
+  });
+
+  const pipelineLayout = device.createPipelineLayout({
+    bindGroupLayouts: [bindGroupLayout],
+  });
+
+  const computePipeline = device.createComputePipeline({
+    layout: pipelineLayout,
+    compute: {
+      module: shaderModule,
+      entryPoint: "main",
+    },
+  });
+
+  const encoder = device.createCommandEncoder();
+
+  const computePass = encoder.beginComputePass();
+  computePass.setPipeline(computePipeline);
+  computePass.setBindGroup(0, bindGroup);
+  computePass.insertDebugMarker("compute collatz iterations");
+  computePass.dispatch(numbers.length);
+  computePass.endPass();
+
+  encoder.copyBufferToBuffer(storageBuffer, 0, stagingBuffer, 0, size);
+
+  device.queue.submit([encoder.finish()]);
+
+  await stagingBuffer.mapAsync(1);
+
+  const data = stagingBuffer.getMappedRange();
+
+  assertEquals(new Uint32Array(data), new Uint32Array([0, 2, 7, 55]));
+
+  stagingBuffer.unmap();
+
+  device.destroy();
+
+  // TODO(lucacasonato): webgpu spec should add a explicit destroy method for
+  // adapters.
+  const resources = Object.keys(Deno.resources());
+  Deno.close(Number(resources[resources.length - 1]));
+});
+
+// Skip this test on linux CI, because the vulkan emulator is not good enough
+// yet, and skip on macOS because these do not have virtual GPUs.
+unitTest({
+  perms: { read: true, env: true },
+  ignore: (Deno.build.os === "linux" || Deno.build.os === "darwin") && isCI,
+}, async function webgpuHelloTriangle() {
+  const adapter = await navigator.gpu.requestAdapter();
+  assert(adapter);
+
+  const device = await adapter.requestDevice();
+  assert(device);
+
+  const shaderCode = await Deno.readTextFile(
+    "cli/tests/webgpu_hellotriangle_shader.wgsl",
+  );
+
+  const shaderModule = device.createShaderModule({
+    code: shaderCode,
+  });
+
+  const pipelineLayout = device.createPipelineLayout({
+    bindGroupLayouts: [],
+  });
+
+  const renderPipeline = device.createRenderPipeline({
+    layout: pipelineLayout,
+    vertex: {
+      module: shaderModule,
+      entryPoint: "vs_main",
+    },
+    fragment: {
+      module: shaderModule,
+      entryPoint: "fs_main",
+      targets: [
+        {
+          format: "rgba8unorm-srgb",
+        },
+      ],
+    },
+  });
+
+  const dimensions = {
+    width: 200,
+    height: 200,
+  };
+  const unpaddedBytesPerRow = dimensions.width * 4;
+  const align = 256;
+  const paddedBytesPerRowPadding = (align - unpaddedBytesPerRow % align) %
+    align;
+  const paddedBytesPerRow = unpaddedBytesPerRow + paddedBytesPerRowPadding;
+
+  const outputBuffer = device.createBuffer({
+    label: "Capture",
+    size: paddedBytesPerRow * dimensions.height,
+    usage: GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST,
+  });
+  const texture = device.createTexture({
+    label: "Capture",
+    size: dimensions,
+    format: "rgba8unorm-srgb",
+    usage: GPUTextureUsage.RENDER_ATTACHMENT | GPUTextureUsage.COPY_SRC,
+  });
+
+  const encoder = device.createCommandEncoder();
+  const renderPass = encoder.beginRenderPass({
+    colorAttachments: [
+      {
+        view: texture.createView(),
+        storeOp: "store",
+        loadValue: [0, 1, 0, 1],
+      },
+    ],
+  });
+  renderPass.setPipeline(renderPipeline);
+  renderPass.draw(3, 1);
+  renderPass.endPass();
+
+  encoder.copyTextureToBuffer(
+    {
+      texture,
+    },
+    {
+      buffer: outputBuffer,
+      bytesPerRow: paddedBytesPerRow,
+      rowsPerImage: 0,
+    },
+    dimensions,
+  );
+
+  device.queue.submit([encoder.finish()]);
+
+  await outputBuffer.mapAsync(1);
+  const data = new Uint8Array(outputBuffer.getMappedRange());
+
+  assertEquals(data, await Deno.readFile("cli/tests/webgpu_hellotriangle.out"));
+
+  outputBuffer.unmap();
+
+  device.destroy();
+
+  // TODO(lucacasonato): webgpu spec should add a explicit destroy method for
+  // adapters.
+  const resources = Object.keys(Deno.resources());
+  Deno.close(Number(resources[resources.length - 1]));
+});

cli/tests/webgpu_computepass_shader.wgsl

@@ -0,0 +1,39 @@
+[[builtin(global_invocation_id)]]
+var global_id: vec3<u32>;
+
+[[block]]
+struct PrimeIndices {
+    data: [[stride(4)]] array<u32>;
+}; // this is used as both input and output for convenience
+
+[[group(0), binding(0)]]
+var<storage> v_indices: [[access(read_write)]] PrimeIndices;
+
+// The Collatz Conjecture states that for any integer n:
+// If n is even, n = n/2
+// If n is odd, n = 3n+1
+// And repeat this process for each new n, you will always eventually reach 1.
+// Though the conjecture has not been proven, no counterexample has ever been found.
+// This function returns how many times this recurrence needs to be applied to reach 1.
+fn collatz_iterations(n_base: u32) -> u32{
+    var n: u32 = n_base;
+    var i: u32 = 0u;
+    loop {
+        if (n <= 1u) {
+            break;
+        }
+        if (n % 2u == 0u) {
+            n = n / 2u;
+        }
+        else {
+            n = 3u * n + 1u;
+        }
+        i = i + 1u;
+    }
+    return i;
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+    v_indices.data[global_id.x] = collatz_iterations(v_indices.data[global_id.x]);
+}

cli/tests/webgpu_hellotriangle_shader.wgsl

@@ -0,0 +1,19 @@
+[[builtin(vertex_index)]]
+var<in> in_vertex_index: u32;
+[[builtin(position)]]
+var<out> out_pos: vec4<f32>;
+
+[[stage(vertex)]]
+fn vs_main() {
+    var x: f32 = f32(i32(in_vertex_index) - 1);
+    var y: f32 = f32(i32(in_vertex_index & 1) * 2 - 1);
+    out_pos = vec4<f32>(x, y, 0.0, 1.0);
+}
+
+[[location(0)]]
+var<out> out_color: vec4<f32>;
+
+[[stage(fragment)]]
+fn fs_main() {
+    out_color = vec4<f32>(1.0, 0.0, 0.0, 1.0);
+}