feat: WebGPU API (denoland#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);
+}