feat: bring back WebGPU

Cargo.toml

@@ -24,6 +24,7 @@ members = [
   "ext/node",
   "ext/url",
   "ext/web",
+  "ext/webgpu",
   "ext/webidl",
   "ext/websocket",
   "ext/webstorage",
@@ -65,6 +66,7 @@ deno_kv = { version = "0.28.0", path = "./ext/kv" }
 deno_tls = { version = "0.107.0", path = "./ext/tls" }
 deno_url = { version = "0.120.0", path = "./ext/url" }
 deno_web = { version = "0.151.0", path = "./ext/web" }
+deno_webgpu = { version = "0.94.0", path = "./ext/webgpu" }
 deno_webidl = { version = "0.120.0", path = "./ext/webidl" }
 deno_websocket = { version = "0.125.0", path = "./ext/websocket" }
 deno_webstorage = { version = "0.115.0", path = "./ext/webstorage" }
@@ -158,6 +160,11 @@ p384 = { version = "0.13.0", features = ["ecdh"] }
 rsa = { version = "0.7.0", default-features = false, features = ["std", "pem", "hazmat"] } # hazmat needed for PrehashSigner in ext/node
 hkdf = "0.12.3"
 
+# webgpu
+raw-window-handle = "0.5.0"
+wgpu-core = "0.17"
+wgpu-types = "0.17"
+
 # macros
 proc-macro2 = "1"
 quote = "1"

cli/args/flags.rs

@@ -755,6 +755,7 @@ static ENV_VARIABLES_HELP: &str = r#"ENVIRONMENT VARIABLES:
     DENO_NO_UPDATE_CHECK Set to disable checking if a newer Deno version is
                          available
     DENO_V8_FLAGS        Set V8 command line options
+    DENO_WEBGPU_TRACE    Directory to use for wgpu traces
     DENO_JOBS            Number of parallel workers used for the --parallel
                          flag with the test subcommand. Defaults to number
                          of available CPUs.

cli/build.rs

@@ -164,6 +164,7 @@ mod ts {
     op_crate_libs.insert("deno.url", deno_url::get_declaration());
     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.webstorage", deno_webstorage::get_declaration());
     op_crate_libs.insert("deno.crypto", deno_crypto::get_declaration());
@@ -532,3 +533,11 @@ fn main() {
     res.compile().unwrap();
   }
 }
+
+fn deno_webgpu_get_declaration() -> PathBuf {
+  let manifest_dir = std::path::Path::new(env!("CARGO_MANIFEST_DIR"));
+  manifest_dir
+    .join("tsc")
+    .join("dts")
+    .join("lib.deno_webgpu.d.ts")
+}

cli/js/40_testing.js

@@ -141,6 +141,9 @@ const OP_DETAILS = {
   "op_utime_async": ["change file timestamps", "awaiting the result of a `Deno.utime` call"],
   "op_host_recv_message": ["receive a message from a web worker", "terminating a `Worker`"],
   "op_host_recv_ctrl": ["receive a message from a web worker", "terminating a `Worker`"],
+  "op_webgpu_buffer_get_map_async": ["map a WebGPU buffer", "awaiting the result of a `GPUBuffer#mapAsync` call"],
+  "op_webgpu_request_adapter": ["request a WebGPU adapter", "awaiting the result of a `navigator.gpu.requestAdapter` call"],
+  "op_webgpu_request_device": ["request a WebGPU device", "awaiting the result of a `GPUAdapter#requestDevice` call"],
   "op_ws_close": ["close a WebSocket", "awaiting until the `close` event is emitted on a `WebSocket`, or the `WebSocketStream#closed` promise resolves"],
   "op_ws_create": ["create a WebSocket", "awaiting until the `open` event is emitted on a `WebSocket`, or the result of a `WebSocketStream#connection` promise"],
   "op_ws_next_event": ["receive the next message on a WebSocket", "closing a `WebSocket` or `WebSocketStream`"],

cli/tests/testdata/webgpu/computepass_shader.wgsl

@@ -0,0 +1,38 @@
+@group(0)
+@binding(0)
+var<storage, read_write> v_indices: array<u32>; // this is used as both input and output for convenience
+
+// 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 {
+            // Overflow? (i.e. 3*n + 1 > 0xffffffffu?)
+            if (n >= 1431655765u) {   // 0x55555555u
+                return 4294967295u;   // 0xffffffffu
+            }
+
+            n = 3u * n + 1u;
+        }
+        i = i + 1u;
+    }
+    return i;
+}
+
+@compute
+@workgroup_size(1)
+fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
+    v_indices[global_id.x] = collatz_iterations(v_indices[global_id.x]);
+}

cli/tests/testdata/webgpu/hellotriangle_shader.wgsl

@@ -0,0 +1,11 @@
+@vertex
+fn vs_main(@builtin(vertex_index) in_vertex_index: u32) -> @builtin(position) vec4<f32> {
+    let x = f32(i32(in_vertex_index) - 1);
+    let y = f32(i32(in_vertex_index & 1u) * 2 - 1);
+    return vec4<f32>(x, y, 0.0, 1.0);
+}
+
+@fragment
+fn fs_main() -> @location(0) vec4<f32> {
+    return vec4<f32>(1.0, 0.0, 0.0, 1.0);
+}

cli/tests/unit/webgpu_test.ts

@@ -0,0 +1,242 @@
+// Copyright 2018-2023 the Deno authors. All rights reserved. MIT license.
+
+import { assert, assertEquals } from "./test_util.ts";
+
+let isCI: boolean;
+try {
+  isCI = (Deno.env.get("CI")?.length ?? 0) > 0;
+} catch {
+  isCI = true;
+}
+
+// Skip these tests on linux CI, because the vulkan emulator is not good enough
+// yet, and skip on macOS CI because these do not have virtual GPUs.
+const isLinuxOrMacCI =
+  (Deno.build.os === "linux" || Deno.build.os === "darwin") && isCI;
+// Skip these tests in WSL because it doesn't have good GPU support.
+const isWsl = await checkIsWsl();
+
+Deno.test({
+  permissions: { read: true, env: true },
+  ignore: isWsl || isLinuxOrMacCI,
+}, 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/testdata/webgpu/computepass_shader.wgsl",
+  );
+
+  const shaderModule = device.createShaderModule({
+    code: shaderCode,
+  });
+
+  const size = new Uint32Array(numbers).byteLength;
+
+  const stagingBuffer = device.createBuffer({
+    size: size,
+    usage: GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST,
+  });
+
+  const storageBuffer = device.createBuffer({
+    label: "Storage Buffer",
+    size: size,
+    usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST |
+      GPUBufferUsage.COPY_SRC,
+    mappedAtCreation: true,
+  });
+
+  const buf = new Uint32Array(storageBuffer.getMappedRange());
+
+  buf.set(numbers);
+
+  storageBuffer.unmap();
+
+  const computePipeline = device.createComputePipeline({
+    layout: "auto",
+    compute: {
+      module: shaderModule,
+      entryPoint: "main",
+    },
+  });
+  const bindGroupLayout = computePipeline.getBindGroupLayout(0);
+
+  const bindGroup = device.createBindGroup({
+    layout: bindGroupLayout,
+    entries: [
+      {
+        binding: 0,
+        resource: {
+          buffer: storageBuffer,
+        },
+      },
+    ],
+  });
+
+  const encoder = device.createCommandEncoder();
+
+  const computePass = encoder.beginComputePass();
+  computePass.setPipeline(computePipeline);
+  computePass.setBindGroup(0, bindGroup);
+  computePass.insertDebugMarker("compute collatz iterations");
+  computePass.dispatchWorkgroups(numbers.length);
+  computePass.end();
+
+  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]));
+});
+
+Deno.test({
+  permissions: { read: true, env: true },
+  ignore: isWsl || isLinuxOrMacCI,
+}, 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/testdata/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 view = texture.createView();
+  const renderPass = encoder.beginRenderPass({
+    colorAttachments: [
+      {
+        view,
+        storeOp: "store",
+        loadOp: "clear",
+        clearValue: [0, 1, 0, 1],
+      },
+    ],
+  });
+  renderPass.setPipeline(renderPipeline);
+  renderPass.draw(3, 1);
+  renderPass.end();
+
+  encoder.copyTextureToBuffer(
+    {
+      texture,
+    },
+    {
+      buffer: outputBuffer,
+      bytesPerRow: paddedBytesPerRow,
+      rowsPerImage: 0,
+    },
+    dimensions,
+  );
+
+  const bundle = encoder.finish();
+  device.queue.submit([bundle]);
+
+  await outputBuffer.mapAsync(1);
+  const data = new Uint8Array(outputBuffer.getMappedRange());
+
+  assertEquals(
+    data,
+    await Deno.readFile("cli/tests/testdata/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]));
+});
+
+Deno.test({
+  ignore: isWsl || isLinuxOrMacCI,
+}, async function webgpuAdapterHasFeatures() {
+  const adapter = await navigator.gpu.requestAdapter();
+  assert(adapter);
+  assert(adapter.features);
+  const resources = Object.keys(Deno.resources());
+  Deno.close(Number(resources[resources.length - 1]));
+});
+
+async function checkIsWsl() {
+  return Deno.build.os === "linux" && await hasMicrosoftProcVersion();
+
+  async function hasMicrosoftProcVersion() {
+    // https://github.com/microsoft/WSL/issues/423#issuecomment-221627364
+    try {
+      const procVersion = await Deno.readTextFile("/proc/version");
+      return /microsoft/i.test(procVersion);
+    } catch {
+      return false;
+    }
+  }
+}

cli/tsc/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="deno.webstorage" />
 /// <reference lib="esnext" />
 /// <reference lib="deno.cache" />
@@ -101,6 +102,7 @@ declare var caches: CacheStorage;
 
 /** @category Web APIs */
 declare interface Navigator {
+  readonly gpu: GPU;
   readonly hardwareConcurrency: number;
   readonly userAgent: string;
   readonly language: string;

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

@@ -62,6 +62,7 @@ declare var WorkerGlobalScope: {
 
 /** @category Web APIs */
 declare interface WorkerNavigator {
+  readonly gpu: GPU;
   readonly hardwareConcurrency: number;
   readonly userAgent: string;
   readonly language: string;