summaryrefslogtreecommitdiff
path: root/cli
diff options
context:
space:
mode:
Diffstat (limited to 'cli')
-rw-r--r--cli/build.rs6
-rw-r--r--cli/dts/lib.deno.window.d.ts7
-rw-r--r--cli/dts/lib.deno.worker.d.ts8
-rw-r--r--cli/flags.rs1
-rw-r--r--cli/main.rs5
-rw-r--r--cli/tests/unit/unit_tests.ts1
-rw-r--r--cli/tests/unit/webgpu_test.ts225
-rw-r--r--cli/tests/webgpu_computepass_shader.wgsl39
-rw-r--r--cli/tests/webgpu_hellotriangle.outbin0 -> 204800 bytes
-rw-r--r--cli/tests/webgpu_hellotriangle_shader.wgsl19
-rw-r--r--cli/tsc.rs1
11 files changed, 311 insertions, 1 deletions
diff --git a/cli/build.rs b/cli/build.rs
index a54bf962c..236097cbf 100644
--- a/cli/build.rs
+++ b/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());
@@ -261,6 +263,10 @@ fn main() {
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()
);
diff --git a/cli/dts/lib.deno.window.d.ts b/cli/dts/lib.deno.window.d.ts
index 721c89af2..2d46e5fe0 100644
--- a/cli/dts/lib.deno.window.d.ts
+++ b/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.
diff --git a/cli/dts/lib.deno.worker.d.ts b/cli/dts/lib.deno.worker.d.ts
index a8f42f336..93549ecd5 100644
--- a/cli/dts/lib.deno.worker.d.ts
+++ b/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 {
diff --git a/cli/flags.rs b/cli/flags.rs
index a803ff63f..6add2df87 100644
--- a/cli/flags.rs
+++ b/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
diff --git a/cli/main.rs b/cli/main.rs
index 7871fe882..0fa63a138 100644
--- a/cli/main.rs
+++ b/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() {
diff --git a/cli/tests/unit/unit_tests.ts b/cli/tests/unit/unit_tests.ts
index 8af28c1bb..2664a9ab0 100644
--- a/cli/tests/unit/unit_tests.ts
+++ b/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";
diff --git a/cli/tests/unit/webgpu_test.ts b/cli/tests/unit/webgpu_test.ts
new file mode 100644
index 000000000..628a3c8db
--- /dev/null
+++ b/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]));
+});
diff --git a/cli/tests/webgpu_computepass_shader.wgsl b/cli/tests/webgpu_computepass_shader.wgsl
new file mode 100644
index 000000000..2433f3243
--- /dev/null
+++ b/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]);
+}
diff --git a/cli/tests/webgpu_hellotriangle.out b/cli/tests/webgpu_hellotriangle.out
new file mode 100644
index 000000000..91454dbfc
--- /dev/null
+++ b/cli/tests/webgpu_hellotriangle.out
Binary files differ
diff --git a/cli/tests/webgpu_hellotriangle_shader.wgsl b/cli/tests/webgpu_hellotriangle_shader.wgsl
new file mode 100644
index 000000000..71934415b
--- /dev/null
+++ b/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);
+}
diff --git a/cli/tsc.rs b/cli/tsc.rs
index 6328175b9..8d28b959d 100644
--- a/cli/tsc.rs
+++ b/cli/tsc.rs
@@ -31,6 +31,7 @@ use std::sync::Mutex;
pub static DENO_NS_LIB: &str = include_str!("dts/lib.deno.ns.d.ts");
pub static DENO_WEB_LIB: &str = include_str!(env!("DENO_WEB_LIB_PATH"));
pub static DENO_FETCH_LIB: &str = include_str!(env!("DENO_FETCH_LIB_PATH"));
+pub static DENO_WEBGPU_LIB: &str = include_str!(env!("DENO_WEBGPU_LIB_PATH"));
pub static DENO_WEBSOCKET_LIB: &str =
include_str!(env!("DENO_WEBSOCKET_LIB_PATH"));
pub static SHARED_GLOBALS_LIB: &str =
generated by cgit v1.2.3 (git 2.25.1) at 2025-06-27 13:35:19 +0000