From 17a51c401a6ec494310f4a1f8b8cc07057bcc656 Mon Sep 17 00:00:00 2001 From: Leo Kettmeir Date: Tue, 18 Feb 2025 09:29:45 +0100 Subject: [PATCH] feat(jupyter): make GPUTexture and GPUBuffer displayable (#28117) --- cli/js/40_jupyter.js | 27 +++- cli/ops/jupyter.rs | 223 ++++++++++++++++++++++++++++++++ cli/tools/test/mod.rs | 5 + ext/canvas/lib.rs | 3 +- ext/web/lib.rs | 2 +- ext/webgpu/device.rs | 2 + ext/webgpu/error.rs | 7 ++ ext/webgpu/lib.rs | 6 +- ext/webgpu/surface.rs | 2 + ext/webgpu/texture.rs | 4 +- runtime/js/99_main.js | 2 + tests/unit/jupyter_test.ts | 252 ++++++++++++++++++++++++++++++++++++- tests/unit/ops_test.ts | 2 +- 13 files changed, 528 insertions(+), 9 deletions(-) diff --git a/cli/js/40_jupyter.js b/cli/js/40_jupyter.js index f392af1d43..f11b0463f3 100644 --- a/cli/js/40_jupyter.js +++ b/cli/js/40_jupyter.js @@ -299,6 +299,19 @@ async function format(obj) { "text/html": obj.outerHTML, }; } + if (obj instanceof GPUTexture) { + return { "image/png": core.ops.op_jupyter_create_png_from_texture(obj) }; + } + if (obj instanceof GPUBuffer) { + return { + "text/plain": Deno[Deno.internal].inspectArgs([ + "%o", + core.ops.op_jupyter_get_buffer(obj), + ], { + colors: !Deno.noColor, + }), + }; + } return { "text/plain": Deno[Deno.internal].inspectArgs(["%o", obj], { colors: !Deno.noColor, @@ -387,8 +400,14 @@ function image(obj) { return makeDisplayable({ "image/png": core.ops.op_base64_encode(obj) }); } + if (obj instanceof GPUTexture) { + return makeDisplayable({ + "image/png": core.ops.op_jupyter_create_png_from_texture(obj), + }); + } + throw new TypeError( - "Object is not a valid image or a path to an image. `Deno.jupyter.image` supports displaying JPG or PNG images.", + "Object is not a valid image or a path to an image. `Deno.jupyter.image` supports displaying JPG or PNG images, or GPUTextures.", ); } @@ -452,6 +471,12 @@ function enableJupyter() { evalue: err.message, traceback: stack.split("\n"), }); + } else if (err instanceof GPUError) { + await broadcast("error", { + ename: err.constructor.name, + evalue: err.message, + traceback: [], + }); } else if (typeof err == "string") { await broadcast("error", { ename: "Error", diff --git a/cli/ops/jupyter.rs b/cli/ops/jupyter.rs index 3160f991bf..999ab11e0d 100644 --- a/cli/ops/jupyter.rs +++ b/cli/ops/jupyter.rs @@ -13,6 +13,7 @@ use deno_core::op2; use deno_core::parking_lot::Mutex; use deno_core::serde_json; use deno_core::OpState; +use deno_error::JsErrorBox; use jupyter_runtime::InputRequest; use jupyter_runtime::JupyterMessage; use jupyter_runtime::JupyterMessageContent; @@ -26,6 +27,8 @@ deno_core::extension!(deno_jupyter, ops = [ op_jupyter_broadcast, op_jupyter_input, + op_jupyter_create_png_from_texture, + op_jupyter_get_buffer, ], options = { sender: mpsc::UnboundedSender, @@ -39,6 +42,21 @@ deno_core::extension!(deno_jupyter, }, ); +deno_core::extension!(deno_jupyter_for_test, + ops = [ + op_jupyter_broadcast, + op_jupyter_input, + op_jupyter_create_png_from_texture, + op_jupyter_get_buffer, + ], + options = { + sender: mpsc::UnboundedSender, + }, + state = |state, options| { + state.put(options.sender); + }, +); + #[op2] #[string] pub fn op_jupyter_input( @@ -166,3 +184,208 @@ pub fn op_print(state: &mut OpState, #[string] msg: &str, is_err: bool) { log::error!("Failed to send stdout message: {}", err); } } + +#[op2] +#[string] +pub fn op_jupyter_create_png_from_texture( + #[cppgc] texture: &deno_runtime::deno_webgpu::texture::GPUTexture, +) -> Result { + use deno_runtime::deno_canvas::image::ExtendedColorType; + use deno_runtime::deno_canvas::image::ImageEncoder; + use deno_runtime::deno_webgpu::error::GPUError; + use deno_runtime::deno_webgpu::*; + use texture::GPUTextureFormat; + + // We only support the 8 bit per pixel formats with 4 channels + // as such a pixel has 4 bytes + const BYTES_PER_PIXEL: u32 = 4; + + let unpadded_bytes_per_row = texture.size.width * BYTES_PER_PIXEL; + let padded_bytes_per_row_padding = (wgpu_types::COPY_BYTES_PER_ROW_ALIGNMENT + - (unpadded_bytes_per_row % wgpu_types::COPY_BYTES_PER_ROW_ALIGNMENT)) + % wgpu_types::COPY_BYTES_PER_ROW_ALIGNMENT; + let padded_bytes_per_row = + unpadded_bytes_per_row + padded_bytes_per_row_padding; + + let (buffer, maybe_err) = texture.instance.device_create_buffer( + texture.device_id, + &wgpu_types::BufferDescriptor { + label: None, + size: (padded_bytes_per_row * texture.size.height) as _, + usage: wgpu_types::BufferUsages::MAP_READ + | wgpu_types::BufferUsages::COPY_DST, + mapped_at_creation: false, + }, + None, + ); + if let Some(maybe_err) = maybe_err { + return Err(JsErrorBox::from_err::(maybe_err.into())); + } + + let (command_encoder, maybe_err) = + texture.instance.device_create_command_encoder( + texture.device_id, + &wgpu_types::CommandEncoderDescriptor { label: None }, + None, + ); + if let Some(maybe_err) = maybe_err { + return Err(JsErrorBox::from_err::(maybe_err.into())); + } + + texture + .instance + .command_encoder_copy_texture_to_buffer( + command_encoder, + &wgpu_types::TexelCopyTextureInfo { + texture: texture.id, + mip_level: 0, + origin: Default::default(), + aspect: Default::default(), + }, + &wgpu_types::TexelCopyBufferInfo { + buffer, + layout: wgpu_types::TexelCopyBufferLayout { + offset: 0, + bytes_per_row: Some(padded_bytes_per_row), + rows_per_image: None, + }, + }, + &texture.size, + ) + .map_err(|e| JsErrorBox::from_err::(e.into()))?; + + let (command_buffer, maybe_err) = texture.instance.command_encoder_finish( + command_encoder, + &wgpu_types::CommandBufferDescriptor { label: None }, + ); + if let Some(maybe_err) = maybe_err { + return Err(JsErrorBox::from_err::(maybe_err.into())); + } + + let maybe_err = texture + .instance + .queue_submit(texture.queue_id, &[command_buffer]) + .err(); + if let Some((_, maybe_err)) = maybe_err { + return Err(JsErrorBox::from_err::(maybe_err.into())); + } + + let index = texture + .instance + .buffer_map_async( + buffer, + 0, + None, + wgpu_core::resource::BufferMapOperation { + host: wgpu_core::device::HostMap::Read, + callback: None, + }, + ) + .map_err(|e| JsErrorBox::from_err::(e.into()))?; + + texture + .instance + .device_poll( + texture.device_id, + wgpu_types::Maintain::WaitForSubmissionIndex(index), + ) + .map_err(|e| JsErrorBox::from_err::(e.into()))?; + + let (slice_pointer, range_size) = texture + .instance + .buffer_get_mapped_range(buffer, 0, None) + .map_err(|e| JsErrorBox::from_err::(e.into()))?; + + let data = { + // SAFETY: creating a slice from pointer and length provided by wgpu and + // then dropping it before unmapping + let slice = unsafe { + std::slice::from_raw_parts(slice_pointer.as_ptr(), range_size as usize) + }; + + let mut unpadded = + Vec::with_capacity((unpadded_bytes_per_row * texture.size.height) as _); + + for i in 0..texture.size.height { + unpadded.extend_from_slice( + &slice[((i * padded_bytes_per_row) as usize) + ..(((i + 1) * padded_bytes_per_row) as usize)] + [..(unpadded_bytes_per_row as usize)], + ); + } + + unpadded + }; + + let color_type = match texture.format { + GPUTextureFormat::Rgba8unorm => ExtendedColorType::Rgba8, + GPUTextureFormat::Rgba8unormSrgb => ExtendedColorType::Rgba8, + GPUTextureFormat::Rgba8snorm => ExtendedColorType::Rgba8, + GPUTextureFormat::Rgba8uint => ExtendedColorType::Rgba8, + GPUTextureFormat::Rgba8sint => ExtendedColorType::Rgba8, + GPUTextureFormat::Bgra8unorm => ExtendedColorType::Bgra8, + GPUTextureFormat::Bgra8unormSrgb => ExtendedColorType::Bgra8, + _ => { + return Err(JsErrorBox::type_error(format!( + "Unsupported texture format '{}'", + texture.format.as_str() + ))) + } + }; + + let mut out: Vec = vec![]; + + let img = + deno_runtime::deno_canvas::image::codecs::png::PngEncoder::new(&mut out); + img + .write_image(&data, texture.size.width, texture.size.height, color_type) + .map_err(|e| JsErrorBox::type_error(e.to_string()))?; + + texture + .instance + .buffer_unmap(buffer) + .map_err(|e| JsErrorBox::from_err::(e.into()))?; + texture.instance.buffer_drop(buffer); + + Ok(deno_runtime::deno_web::forgiving_base64_encode(&out)) +} + +#[op2] +#[serde] +pub fn op_jupyter_get_buffer( + #[cppgc] buffer: &deno_runtime::deno_webgpu::buffer::GPUBuffer, +) -> Result, deno_runtime::deno_webgpu::error::GPUError> { + use deno_runtime::deno_webgpu::*; + let index = buffer.instance.buffer_map_async( + buffer.id, + 0, + None, + wgpu_core::resource::BufferMapOperation { + host: wgpu_core::device::HostMap::Read, + callback: None, + }, + )?; + + buffer.instance.device_poll( + buffer.device, + wgpu_types::Maintain::WaitForSubmissionIndex(index), + )?; + + let (slice_pointer, range_size) = buffer + .instance + .buffer_get_mapped_range(buffer.id, 0, None)?; + + let data = { + // SAFETY: creating a slice from pointer and length provided by wgpu and + // then dropping it before unmapping + let slice = unsafe { + std::slice::from_raw_parts(slice_pointer.as_ptr(), range_size as usize) + }; + + slice.to_vec() + }; + + buffer.instance.buffer_unmap(buffer.id)?; + + Ok(data) +} diff --git a/cli/tools/test/mod.rs b/cli/tools/test/mod.rs index 697f99aa17..c3b9d5edc0 100644 --- a/cli/tools/test/mod.rs +++ b/cli/tools/test/mod.rs @@ -67,6 +67,7 @@ use rand::SeedableRng; use regex::Regex; use serde::Deserialize; use tokio::signal; +use tokio::sync::mpsc::UnboundedSender; use crate::args::CliOptions; use crate::args::Flags; @@ -615,6 +616,7 @@ async fn configure_main_worker( permissions_container: PermissionsContainer, worker_sender: TestEventWorkerSender, options: &TestSpecifierOptions, + sender: UnboundedSender, ) -> Result< (Option>, MainWorker), CreateCustomWorkerError, @@ -627,6 +629,7 @@ async fn configure_main_worker( vec![ ops::testing::deno_test::init_ops(worker_sender.sender), ops::lint::deno_lint_ext_for_test::init_ops(), + ops::jupyter::deno_jupyter_for_test::init_ops(sender), ], Stdio { stdin: StdioPipe::inherit(), @@ -672,12 +675,14 @@ pub async fn test_specifier( if fail_fast_tracker.should_stop() { return Ok(()); } + let jupyter_channel = tokio::sync::mpsc::unbounded_channel(); let (coverage_collector, mut worker) = configure_main_worker( worker_factory, &specifier, permissions_container, worker_sender, &options, + jupyter_channel.0, ) .await?; diff --git a/ext/canvas/lib.rs b/ext/canvas/lib.rs index 83aee3cb81..fed86d28ba 100644 --- a/ext/canvas/lib.rs +++ b/ext/canvas/lib.rs @@ -4,6 +4,7 @@ use std::path::PathBuf; mod image_ops; mod op_create_image_bitmap; +pub use image; use image::ColorType; use op_create_image_bitmap::op_create_image_bitmap; @@ -11,7 +12,7 @@ use op_create_image_bitmap::op_create_image_bitmap; pub enum CanvasError { /// Image formats that is 32-bit depth are not supported currently due to the following reasons: /// - e.g. OpenEXR, it's not covered by the spec. - /// - JPEG XL supported by WebKit, but it cannot be called a standard today. + /// - JPEG XL supported by WebKit, but it cannot be called a standard today. /// https://github.com/whatwg/mimesniff/issues/143 /// #[class(type)] diff --git a/ext/web/lib.rs b/ext/web/lib.rs index 9b6d215f83..b802f4791a 100644 --- a/ext/web/lib.rs +++ b/ext/web/lib.rs @@ -193,7 +193,7 @@ fn op_base64_btoa(#[serde] s: ByteString) -> String { /// See #[inline] -fn forgiving_base64_encode(s: &[u8]) -> String { +pub fn forgiving_base64_encode(s: &[u8]) -> String { base64_simd::STANDARD.encode_to_string(s) } diff --git a/ext/webgpu/device.rs b/ext/webgpu/device.rs index e6e27d4251..a18b09a2ae 100644 --- a/ext/webgpu/device.rs +++ b/ext/webgpu/device.rs @@ -209,6 +209,8 @@ impl GPUDevice { instance: self.instance.clone(), error_handler: self.error_handler.clone(), id, + device_id: self.id, + queue_id: self.queue, label: descriptor.label, size: wgpu_descriptor.size, mip_level_count: wgpu_descriptor.mip_level_count, diff --git a/ext/webgpu/error.rs b/ext/webgpu/error.rs index 056301e876..c9dd7fa5e6 100644 --- a/ext/webgpu/error.rs +++ b/ext/webgpu/error.rs @@ -20,6 +20,7 @@ use wgpu_core::command::RenderPassError; use wgpu_core::device::queue::QueueSubmitError; use wgpu_core::device::queue::QueueWriteError; use wgpu_core::device::DeviceError; +use wgpu_core::device::WaitIdleError; use wgpu_core::pipeline::CreateComputePipelineError; use wgpu_core::pipeline::CreateRenderPipelineError; use wgpu_core::pipeline::CreateShaderModuleError; @@ -354,3 +355,9 @@ impl From for GPUError { GPUError::Validation(fmt_err(&err)) } } + +impl From for GPUError { + fn from(err: WaitIdleError) -> Self { + GPUError::Validation(fmt_err(&err)) + } +} diff --git a/ext/webgpu/lib.rs b/ext/webgpu/lib.rs index af609e7a19..ebf690172f 100644 --- a/ext/webgpu/lib.rs +++ b/ext/webgpu/lib.rs @@ -18,14 +18,14 @@ use wgpu_types::PowerPreference; mod adapter; mod bind_group; mod bind_group_layout; -mod buffer; +pub mod buffer; mod byow; mod command_buffer; mod command_encoder; mod compute_pass; mod compute_pipeline; mod device; -mod error; +pub mod error; mod pipeline_layout; mod query_set; mod queue; @@ -35,7 +35,7 @@ mod render_pipeline; mod sampler; mod shader; mod surface; -mod texture; +pub mod texture; mod webidl; pub const UNSTABLE_FEATURE_NAME: &str = "webgpu"; diff --git a/ext/webgpu/surface.rs b/ext/webgpu/surface.rs index a37f03513f..7f3204004d 100644 --- a/ext/webgpu/surface.rs +++ b/ext/webgpu/surface.rs @@ -132,6 +132,8 @@ impl GPUCanvasContext { instance: config.device.instance.clone(), error_handler: config.device.error_handler.clone(), id, + device_id: config.device.id, + queue_id: config.device.queue, label: "".to_string(), size: wgpu_types::Extent3d { width: *self.width.borrow(), diff --git a/ext/webgpu/texture.rs b/ext/webgpu/texture.rs index 28faf1fe34..f7c75e945c 100644 --- a/ext/webgpu/texture.rs +++ b/ext/webgpu/texture.rs @@ -42,6 +42,8 @@ pub struct GPUTexture { pub error_handler: super::error::ErrorHandler, pub id: wgpu_core::id::TextureId, + pub device_id: wgpu_core::id::DeviceId, + pub queue_id: wgpu_core::id::QueueId, pub label: String, @@ -286,7 +288,7 @@ impl From for TextureDimension { #[derive(WebIDL, Clone)] #[webidl(enum)] -pub(crate) enum GPUTextureFormat { +pub enum GPUTextureFormat { #[webidl(rename = "r8unorm")] R8unorm, #[webidl(rename = "r8snorm")] diff --git a/runtime/js/99_main.js b/runtime/js/99_main.js index 47c6b419a5..58cce3b94a 100644 --- a/runtime/js/99_main.js +++ b/runtime/js/99_main.js @@ -661,6 +661,8 @@ const NOT_IMPORTED_OPS = [ // Related to `Deno.jupyter` API "op_jupyter_broadcast", "op_jupyter_input", + "op_jupyter_create_png_from_texture", + "op_jupyter_get_buffer", // Used in jupyter API "op_base64_encode", diff --git a/tests/unit/jupyter_test.ts b/tests/unit/jupyter_test.ts index e29bb2b300..fd615fb68a 100644 --- a/tests/unit/jupyter_test.ts +++ b/tests/unit/jupyter_test.ts @@ -1,6 +1,6 @@ // Copyright 2018-2025 the Deno authors. MIT license. -import { assertEquals, assertThrows } from "./test_util.ts"; +import { assert, assertEquals, assertThrows } from "./test_util.ts"; // @ts-expect-error TypeScript (as of 3.7) does not support indexing namespaces by symbol const format = Deno[Deno.internal].jupyter.formatInner; @@ -77,3 +77,253 @@ Deno.test( await assertFormattedAs(example, { "application/json": { x: 3 } }); }, ); + +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 x86 CI because these do not have virtual GPUs. +const isCIWithoutGPU = (Deno.build.os === "linux" || + (Deno.build.os === "darwin" && Deno.build.arch === "x86_64")) && isCI; +// Skip these tests in WSL because it doesn't have good GPU support. +const isWsl = await checkIsWsl(); +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; + } + } +} + +Deno.test( + { + ignore: isWsl || isCIWithoutGPU, + name: "display GPUTexture", + }, + async () => { + const dimensions = { + width: 200, + height: 200, + }; + + const adapter = await navigator.gpu.requestAdapter(); + const device = await adapter?.requestDevice(); + + assert(device); + + const shaderCode = ` +@vertex +fn vs_main(@builtin(vertex_index) in_vertex_index: u32) -> @builtin(position) vec4 { + let x = f32(i32(in_vertex_index) - 1); + let y = f32(i32(in_vertex_index & 1u) * 2 - 1); + return vec4(x, y, 0.0, 1.0); +} + +@fragment +fn fs_main() -> @location(0) vec4 { + return vec4(1.0, 0.0, 0.0, 1.0); +} +`; + + 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 texture = device.createTexture({ + label: "Capture", + size: { + width: dimensions.width, + height: dimensions.height, + }, + format: "rgba8unorm-srgb", + usage: GPUTextureUsage.RENDER_ATTACHMENT | GPUTextureUsage.COPY_SRC, + }); + + const encoder = device.createCommandEncoder(); + const renderPass = encoder.beginRenderPass({ + colorAttachments: [ + { + view: texture.createView(), + storeOp: "store", + loadOp: "clear", + clearValue: [0, 1, 0, 1], + }, + ], + }); + renderPass.setPipeline(renderPipeline); + renderPass.draw(3, 1); + renderPass.end(); + + device.queue.submit([encoder.finish()]); + + await assertFormattedAs(texture, { + "image/png": + "iVBORw0KGgoAAAANSUhEUgAAAMgAAADICAYAAACtWK6eAAAHoklEQVR4Ae3gAZAkSZIkSRKLqpm7R0REZmZmVlVVVVV3d3d3d/fMzMzMzMzMzMzMzMzMzMzMzMzMzMzMzMzMzMzMzMzMzMzMzMzMzMzMzMzMzMzMzMzMzMzMzMzMzMzMzMzMzMzMzMzMzMzMzMzMzMzMzMzMzMzMzMzMdHd3d3dXV1VVVVVmZkZGRIS7m5kKz0xmV3d1d3dPz8zMzMxMYsWYq6666vmhctX/GBaXyVz1PwOVq6666gUhuOp/BItnsbjqfwYqV1111QtCcNV/O4vnYXHVfz8qV1111QtCcNV/K4sXyOKq/15UrrrqqheE4Kr/Nhb/Iour/vtQueqqq14Qgqv+W1i8yCyu+u9B5aqrrnpBCK76L2fxr2Zx1X89KlddddULQnDVfymLfzOLq/5rUbnqqqteEIKr/stY/LtZXPVfh8pVV131ghBc9V/C4j+MxVX/NahcddVVLwjBVf/pLP7DWVz1n4/KVVdd9YIQXPWfyuI/jcVV/7moXHXVVS8IwVX/aSz+01lc9Z+HylVXXfWCEFz1n8Liv4zFVf85qFx11VUvCMFV/+Es/stZXPUfj8pVV131ghBc9R/K4r+NxVX/sahcddVVLwjBVf9hLP7bWVz1H4fKVVdd9YIQXPUfwuJ/DIur/mNQueqqq14Qgqv+3Sz+x7G46t+PylVXXfWCEFz172LxP5bFVf8+VK666qoXhOCqfzOL//Esrvq3o3LVVVe9IARX/ZtY/K9hcdW/DZWrrrrqBSG46l/N4n8di6v+9ahcddVVLwjBVf8qFv9rWVz1r0PlqquuekEIrnqRWfyvZ3HVi47KVVdd9YIQXPUisfg/w+KqFw2Vq6666gUhuOpfZPF/jsVV/zIqV1111QtCcNULZfF/lsVVLxyVq6666gUhuOoFsvg/z+KqF4zKVVdd9YIQXPV8Wfy/YXHV80flqquuekEIrnoeFv/vWFz1vKhcddVVLwjBVc/B4v8ti6ueE5WrrrrqBSG46lks/t+zuOrZqFx11VUvCMFVl1lc9UwWV11B5aqrrnpBCK7C4qrnYnEVULnqqqteEIL/5yyuegEs/r+jctVVV70gBP+PWVz1L7D4/4zKVVdd9YIQ/D9lcdWLyOL/KypXXXXVC0Lw/5DFVf9KFv8fUbnqqqteEIL/Zyyu+jey+P+GylVXXfWCEPw/YnHVv5PF/ydUrrrqqheE4P8Ji6v+g1j8f0HlqquuekEI/h+wuOo/mMX/B1SuuuqqF4Tg/ziLq/6TWPxfR+Wqq656QQj+D7O46j+Zxf9lVK666qoXhOD/KIur/otY/F9F5aqrrnpBCP4Psrjqv5jF/0VUrrrqqheE4P8Yi6v+m1j8X0PlqquuekEI/g+xuOq/mcX/JVSuuuqqF4Tg/wiLq/6HsPi/gspVV131ghD8H2Bx1f8wFv8XULnqqqteEIL/5Syu+h/K4n87KlddddULQvC/mMVV/8NZ/G9G5aqrrnpBCP6XsrjqfwmL/62oXHXVVS8Iwf9CFlf9L2PxvxGVq6666gUh+F/G4qr/pSz+t6Fy1VVXvSAE/4tYXPW/nMX/JlSuuuqqF4TgfwmLq/6PsPjfgspVV131ghD8L2Bx1f8xFv8bULnqqqteEIL/4Syu+j/K4n86KlddddULQvA/mMVV/8dZ/E9G5aqrrnpBCP6Hsrjq/wmL/6moXHXVVS8Iwf9AFlf9P2PxPxGVq6666gUh+B/G4qr/pyz+p6Fy1VVXvSAE/4NYXPX/nMX/JFSuuuqqF4TgfwiLq666zOJ/CipXXXXVC0LwP4DFVVc9B4v/CahcddVVLwjBfzOLq656viz+u1G56qqrXhCC/0YWV131Qln8d6Jy1VVXvSAE/00srrrqRWLx34XKVVdd9YIQ/DewuOqqfxWL/w5UrrrqqheE4L+YxVVX/ZtY/FejctVVV70gBP+FLK666t/F4r8SlauuuuoFIfgvYnHVVf8hLP6rULnqqqteEIL/AhZXXfUfyuK/ApWrrrrqBSH4T2Zx1VX/KSz+s1G56qqrXhCC/0QWV131n8riPxOVq6666gUh+E9icdVV/yUs/rNQueqqq14Qgv8EFldd9V/K4j8DlauuuuoFIfgPZnHVVf8tLP6jUbnqqqteEIL/QBZXXfXfyuI/EpWrrrrqBSH4D2Jx1VX/I1j8R6Fy1VVXvSAE/wEsrrrqfxSL/whUrrrqqheE4N/J4qqr/key+PeictVVV70gBP8OFldd9T+axb8HlauuuuoFIfg3srjqqv8VLP6tqFx11VUvCMG/gcVVV/2vYvFvQeWqq656QQj+lSyuuup/JYt/LSpXXXXVC0Lwr2Bx1VX/q1n8a1C56qqrXhCCF5HFVVf9n2DxoqJy1VVXvSAELwKLq676P8XiRUHlqquuekEI/gUWV131f5LFv4TKVVdd9YIQvBAWV131f5rFC0PlqquuekEIXgCLq676f8HiBaFy1VVXvSAEz4fFVVf9v2Lx/FC56qqrXhCC52Jx1VX/L1k8NypXXXXVC0LwABZXXfX/msUDUbnqqqteEIJnsrjqqqsAi/tRueqqq14QAsDiqquuegALACpXXXXVC4IM5qqrrnp++EdXhsxWnWgkVAAAAABJRU5ErkJggg==", + }); + }, +); + +Deno.test( + { + ignore: isWsl || isCIWithoutGPU, + name: "display GPUBuffer", + }, + async () => { + // Get some numbers from the command line, or use the default 1, 4, 3, 295. + let numbers: Uint32Array; + if (Deno.args.length > 0) { + numbers = new Uint32Array(Deno.args.map((a) => parseInt(a))); + } else { + numbers = new Uint32Array([1, 4, 3, 295]); + } + + const adapter = await navigator.gpu.requestAdapter(); + const device = await adapter?.requestDevice(); + assert(device); + + const shaderCode = ` +@group(0) +@binding(0) +var v_indices: array; // 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) { + v_indices[global_id.x] = collatz_iterations(v_indices[global_id.x]); +} +`; + + const shaderModule = device.createShaderModule({ + code: shaderCode, + }); + + const stagingBuffer = device.createBuffer({ + size: numbers.byteLength, + usage: GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST, + }); + + const contents = new Uint8Array(numbers.buffer); + + const alignMask = 4 - 1; + const paddedSize = Math.max( + (contents.byteLength + alignMask) & ~alignMask, + 4, + ); + + const storageBuffer = device.createBuffer({ + label: "Storage Buffer", + usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST | + GPUBufferUsage.COPY_SRC, + mappedAtCreation: true, + size: paddedSize, + }); + const data = new Uint8Array(storageBuffer.getMappedRange()); + data.set(contents); + 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, + numbers.byteLength, + ); + + device.queue.submit([encoder.finish()]); + + await assertFormattedAs(stagingBuffer, { + "text/plain": + "[\n \x1b[33m0\x1b[39m, \x1b[33m0\x1b[39m, \x1b[33m0\x1b[39m, \x1b[33m0\x1b[39m, \x1b[33m2\x1b[39m, \x1b[33m0\x1b[39m,\n \x1b[33m0\x1b[39m, \x1b[33m0\x1b[39m, \x1b[33m7\x1b[39m, \x1b[33m0\x1b[39m, \x1b[33m0\x1b[39m, \x1b[33m0\x1b[39m,\n \x1b[33m55\x1b[39m, \x1b[33m0\x1b[39m, \x1b[33m0\x1b[39m, \x1b[33m0\x1b[39m\n]", + }); + }, +); diff --git a/tests/unit/ops_test.ts b/tests/unit/ops_test.ts index 9998ad6d80..c9f721786e 100644 --- a/tests/unit/ops_test.ts +++ b/tests/unit/ops_test.ts @@ -1,6 +1,6 @@ // Copyright 2018-2025 the Deno authors. MIT license. -const EXPECTED_OP_COUNT = 14; +const EXPECTED_OP_COUNT = 18; Deno.test(function checkExposedOps() { // @ts-ignore TS doesn't allow to index with symbol