feat(jupyter): make GPUTexture and GPUBuffer displayable (#28117)

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",

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<StreamContent>,
@@ -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<StreamContent>,
+  },
+  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<String, JsErrorBox> {
+  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::<GPUError>(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::<GPUError>(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::<GPUError>(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::<GPUError>(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::<GPUError>(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::<GPUError>(e.into()))?;
+
+  texture
+    .instance
+    .device_poll(
+      texture.device_id,
+      wgpu_types::Maintain::WaitForSubmissionIndex(index),
+    )
+    .map_err(|e| JsErrorBox::from_err::<GPUError>(e.into()))?;
+
+  let (slice_pointer, range_size) = texture
+    .instance
+    .buffer_get_mapped_range(buffer, 0, None)
+    .map_err(|e| JsErrorBox::from_err::<GPUError>(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<u8> = 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::<GPUError>(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<Vec<u8>, 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)
+}

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<jupyter_runtime::messaging::content::StreamContent>,
 ) -> Result<
   (Option<Box<dyn CoverageCollector>>, 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?;
 

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)]

ext/web/lib.rs

@@ -193,7 +193,7 @@ fn op_base64_btoa(#[serde] s: ByteString) -> String {
 
 /// See <https://infra.spec.whatwg.org/#forgiving-base64>
 #[inline]
-fn forgiving_base64_encode(s: &[u8]) -> String {
+pub fn forgiving_base64_encode(s: &[u8]) -> String {
   base64_simd::STANDARD.encode_to_string(s)
 }
 

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,

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<ConfigureSurfaceError> for GPUError {
     GPUError::Validation(fmt_err(&err))
   }
 }
+
+impl From<WaitIdleError> for GPUError {
+  fn from(err: WaitIdleError) -> Self {
+    GPUError::Validation(fmt_err(&err))
+  }
+}

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";

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(),

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<GPUTextureDimension> for TextureDimension {
 
 #[derive(WebIDL, Clone)]
 #[webidl(enum)]
-pub(crate) enum GPUTextureFormat {
+pub enum GPUTextureFormat {
   #[webidl(rename = "r8unorm")]
   R8unorm,
   #[webidl(rename = "r8snorm")]

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",
 

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<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);
+}
+`;
+
+    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<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]);
+}
+`;
+
+    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]",
+    });
+  },
+);

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