WIP: add cts_runner and deno_webgpu crate (#1859)

* WIP: add cts_runner and deno_webgpu crate * add test * remove Cargo.lock * review comment * simplify * fix bugs * improve cts_runner to work with crowlKats/webgpu-examples * fix * remove build.rs cts_runner binaries are now not portable anymore. Also startup will now print a bunch of cargo:rerun-if-changed=. This will be fixed in deno_core. * remove d.ts * add original deno license file
2024-11-25 00:03:29 +00:00 · 2021-09-03 19:23:35 +02:00 · 2021-09-03 19:23:35 +02:00 · d5ba0b439d
commit d5ba0b439d
parent 663f64c571
31 changed files with 13340 additions and 1991 deletions
--- a/.gitignore
+++ b/.gitignore
@ -1,6 +1,7 @@
 # Generated by Cargo
 # will have compiled files and executables
 /target/
+Cargo.lock

 # These are backup files generated by rustfmt
 **/*.rs.bk
--- a/Cargo.lock
+++ b/Cargo.lock
--- a/Cargo.toml
+++ b/Cargo.toml
@ -1,6 +1,8 @@
 [workspace]
 resolver = "2"
 members = [
+    "cts_runner",
+    "deno_webgpu",
    "dummy",
    "player",
    "wgpu",
--- a/cts_runner/Cargo.toml
+++ b/cts_runner/Cargo.toml
@ -0,0 +1,31 @@
+[package]
+name = "cts_runner"
+version = "0.1.0"
+authors = [
+	"Luca Casonato <hello@lcas.dev>",
+]
+edition = "2018"
+description = "CTS runner for wgpu"
+license = "MIT OR Apache-2.0"
+publish = false
+
+[dependencies]
+deno_console = { git = "https://github.com/denoland/deno", rev = "ca75752e5a9499a0a997809f02b18c2ba1ecd58d" }
+deno_core = { git = "https://github.com/denoland/deno", rev = "ca75752e5a9499a0a997809f02b18c2ba1ecd58d" }
+deno_timers = { git = "https://github.com/denoland/deno", rev = "ca75752e5a9499a0a997809f02b18c2ba1ecd58d" }
+deno_url = { git = "https://github.com/denoland/deno", rev = "ca75752e5a9499a0a997809f02b18c2ba1ecd58d" }
+deno_web = { git = "https://github.com/denoland/deno", rev = "ca75752e5a9499a0a997809f02b18c2ba1ecd58d" }
+deno_webidl = { git = "https://github.com/denoland/deno", rev = "ca75752e5a9499a0a997809f02b18c2ba1ecd58d" }
+deno_webgpu = { path = "../deno_webgpu" }
+tokio = { version = "1.10.0", features = ["full"] }
+termcolor = "1.1.2"
+
+[build-dependencies]
+deno_console = { git = "https://github.com/denoland/deno", rev = "ca75752e5a9499a0a997809f02b18c2ba1ecd58d" }
+deno_core = { git = "https://github.com/denoland/deno", rev = "ca75752e5a9499a0a997809f02b18c2ba1ecd58d" }
+deno_timers = { git = "https://github.com/denoland/deno", rev = "ca75752e5a9499a0a997809f02b18c2ba1ecd58d" }
+deno_url = { git = "https://github.com/denoland/deno", rev = "ca75752e5a9499a0a997809f02b18c2ba1ecd58d" }
+deno_web = { git = "https://github.com/denoland/deno", rev = "ca75752e5a9499a0a997809f02b18c2ba1ecd58d" }
+deno_webidl = { git = "https://github.com/denoland/deno", rev = "ca75752e5a9499a0a997809f02b18c2ba1ecd58d" }
+deno_webgpu = { path = "../deno_webgpu" }
+tokio = { version = "1.10.0", features = ["full"] }
--- a/cts_runner/examples/hello-compute.js
+++ b/cts_runner/examples/hello-compute.js
@ -0,0 +1,123 @@
+const adapter = await navigator.gpu.requestAdapter();
+
+const numbers = [1, 4, 3, 295];
+
+const device = await adapter.requestDevice();
+
+const shaderCode = `[[block]]
+struct PrimeIndices {
+    data: [[stride(4)]] array<u32>;
+}; // this is used as both input and output for convenience
+[[group(0), binding(0)]]
+var<storage, read_write> v_indices: 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 {
+            // Overflow? (i.e. 3*n + 1 > 0xffffffffu?)
+            if (n >= 1431655765u) {   // 0x55555555u
+                return 4294967295u;   // 0xffffffffu
+            }
+            n = 3u * n + 1u;
+        }
+        i = i + 1u;
+    }
+    return i;
+}
+[[stage(compute), workgroup_size(1)]]
+fn main([[builtin(global_invocation_id)]] global_id: vec3<u32>) {
+    v_indices.data[global_id.x] = collatz_iterations(v_indices.data[global_id.x]);
+}`;
+
+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 computePipeline = device.createComputePipeline({
+  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.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();
+
+function isTypedArrayEqual(a, b) {
+  if (a.byteLength !== b.byteLength) return false;
+  return a.every((val, i) => val === b[i]);
+}
+
+const actual = new Uint32Array(data);
+const expected = new Uint32Array([0, 2, 7, 55]);
+
+console.error("actual", actual);
+console.error("expected", expected);
+
+if (!isTypedArrayEqual(actual, expected)) {
+  throw new TypeError("Actual does not equal expected!");
+}
+
+stagingBuffer.unmap();
+
+device.destroy();
--- a/cts_runner/src/bootstrap.js
+++ b/cts_runner/src/bootstrap.js
@ -0,0 +1,232 @@
+// Adapted from https://github.com/denoland/deno/blob/6abf126c2a7a451cded8c6b5e6ddf1b69c84055d/runtime/js/99_main.js
+
+// Removes the `__proto__` for security reasons.  This intentionally makes
+// Deno non compliant with ECMA-262 Annex B.2.2.1
+//
+delete Object.prototype.__proto__;
+
+((window) => {
+  const core = Deno.core;
+  const {
+    Error,
+    ObjectDefineProperty,
+    ObjectDefineProperties,
+    Symbol,
+    StringPrototypeReplace,
+  } = window.__bootstrap.primordials;
+  const webidl = window.__bootstrap.webidl;
+  const eventTarget = window.__bootstrap.eventTarget;
+  const globalInterfaces = window.__bootstrap.globalInterfaces;
+  const { Console } = window.__bootstrap.console;
+  const timers = window.__bootstrap.timers;
+  const base64 = window.__bootstrap.base64;
+  const encoding = window.__bootstrap.encoding;
+  const url = window.__bootstrap.url;
+  const domException = window.__bootstrap.domException;
+  const performance = window.__bootstrap.performance;
+  const webgpu = window.__bootstrap.webgpu;
+
+  const util = {
+    immutableDefine(o, p, value) {
+      ObjectDefineProperty(o, p, {
+        value,
+        configurable: false,
+        writable: false,
+      });
+    },
+    writable(value) {
+      return {
+        value,
+        writable: true,
+        enumerable: true,
+        configurable: true,
+      };
+    },
+    nonEnumerable(value) {
+      return {
+        value,
+        writable: true,
+        configurable: true,
+      };
+    },
+    readOnly(value) {
+      return {
+        value,
+        enumerable: true,
+      };
+    },
+  };
+
+  class Navigator {
+    constructor() {
+      webidl.illegalConstructor();
+    }
+
+    [Symbol.for("Deno.customInspect")](inspect) {
+      return `${this.constructor.name} ${inspect({})}`;
+    }
+  }
+
+  const navigator = webidl.createBranded(Navigator);
+
+  ObjectDefineProperties(Navigator.prototype, {
+    gpu: {
+      configurable: true,
+      enumerable: true,
+      get() {
+        webidl.assertBranded(this, Navigator);
+        return webgpu.gpu;
+      },
+    },
+  });
+
+  const windowOrWorkerGlobalScope = {
+    CloseEvent: util.nonEnumerable(CloseEvent),
+    CustomEvent: util.nonEnumerable(CustomEvent),
+    DOMException: util.nonEnumerable(domException.DOMException),
+    ErrorEvent: util.nonEnumerable(ErrorEvent),
+    Event: util.nonEnumerable(Event),
+    EventTarget: util.nonEnumerable(EventTarget),
+    Navigator: util.nonEnumerable(Navigator),
+    navigator: {
+      configurable: true,
+      enumerable: true,
+      get: () => navigator,
+    },
+    Performance: util.nonEnumerable(performance.Performance),
+    PerformanceEntry: util.nonEnumerable(performance.PerformanceEntry),
+    PerformanceMark: util.nonEnumerable(performance.PerformanceMark),
+    PerformanceMeasure: util.nonEnumerable(performance.PerformanceMeasure),
+    TextDecoder: util.nonEnumerable(encoding.TextDecoder),
+    TextEncoder: util.nonEnumerable(encoding.TextEncoder),
+    URL: util.nonEnumerable(url.URL),
+    URLSearchParams: util.nonEnumerable(url.URLSearchParams),
+    atob: util.writable(base64.atob),
+    btoa: util.writable(base64.btoa),
+    console: util.writable(new Console(core.print)),
+    setInterval: util.writable(timers.setInterval),
+    setTimeout: util.writable(timers.setTimeout),
+    clearInterval: util.writable(timers.clearInterval),
+    clearTimeout: util.writable(timers.clearTimeout),
+    performance: util.writable(performance.performance),
+
+    GPU: util.nonEnumerable(webgpu.GPU),
+    GPUAdapter: util.nonEnumerable(webgpu.GPUAdapter),
+    GPUAdapterLimits: util.nonEnumerable(webgpu.GPUAdapterLimits),
+    GPUSupportedFeatures: util.nonEnumerable(webgpu.GPUSupportedFeatures),
+    GPUDevice: util.nonEnumerable(webgpu.GPUDevice),
+    GPUQueue: util.nonEnumerable(webgpu.GPUQueue),
+    GPUBuffer: util.nonEnumerable(webgpu.GPUBuffer),
+    GPUBufferUsage: util.nonEnumerable(webgpu.GPUBufferUsage),
+    GPUMapMode: util.nonEnumerable(webgpu.GPUMapMode),
+    GPUTexture: util.nonEnumerable(webgpu.GPUTexture),
+    GPUTextureUsage: util.nonEnumerable(webgpu.GPUTextureUsage),
+    GPUTextureView: util.nonEnumerable(webgpu.GPUTextureView),
+    GPUSampler: util.nonEnumerable(webgpu.GPUSampler),
+    GPUBindGroupLayout: util.nonEnumerable(webgpu.GPUBindGroupLayout),
+    GPUPipelineLayout: util.nonEnumerable(webgpu.GPUPipelineLayout),
+    GPUBindGroup: util.nonEnumerable(webgpu.GPUBindGroup),
+    GPUShaderModule: util.nonEnumerable(webgpu.GPUShaderModule),
+    GPUShaderStage: util.nonEnumerable(webgpu.GPUShaderStage),
+    GPUComputePipeline: util.nonEnumerable(webgpu.GPUComputePipeline),
+    GPURenderPipeline: util.nonEnumerable(webgpu.GPURenderPipeline),
+    GPUColorWrite: util.nonEnumerable(webgpu.GPUColorWrite),
+    GPUCommandEncoder: util.nonEnumerable(webgpu.GPUCommandEncoder),
+    GPURenderPassEncoder: util.nonEnumerable(webgpu.GPURenderPassEncoder),
+    GPUComputePassEncoder: util.nonEnumerable(webgpu.GPUComputePassEncoder),
+    GPUCommandBuffer: util.nonEnumerable(webgpu.GPUCommandBuffer),
+    GPURenderBundleEncoder: util.nonEnumerable(webgpu.GPURenderBundleEncoder),
+    GPURenderBundle: util.nonEnumerable(webgpu.GPURenderBundle),
+    GPUQuerySet: util.nonEnumerable(webgpu.GPUQuerySet),
+    GPUOutOfMemoryError: util.nonEnumerable(webgpu.GPUOutOfMemoryError),
+    GPUValidationError: util.nonEnumerable(webgpu.GPUValidationError),
+  };
+
+  windowOrWorkerGlobalScope.console.enumerable = false;
+
+  const mainRuntimeGlobalProperties = {
+    Window: globalInterfaces.windowConstructorDescriptor,
+    window: util.readOnly(globalThis),
+    self: util.readOnly(globalThis),
+  };
+
+  // Taken from deno/runtime/js/06_util.js
+  function pathFromURL(pathOrUrl) {
+    if (pathOrUrl instanceof URL) {
+      if (pathOrUrl.protocol != "file:") {
+        throw new TypeError("Must be a file URL.");
+      }
+      if (pathOrUrl.hostname !== "") {
+        throw new TypeError("Host must be empty.");
+      }
+      return decodeURIComponent(
+        StringPrototypeReplace(
+          pathOrUrl.pathname,
+          /%(?![0-9A-Fa-f]{2})/g,
+          "%25",
+        ),
+      );
+    }
+    return pathOrUrl;
+  }
+
+  const denoNs = {
+    exit(code) {
+      core.opSync("op_exit", code);
+    },
+    readFileSync(path) {
+      return core.opSync("op_read_file_sync", pathFromURL(path));
+    },
+    readTextFileSync(path) {
+      const buf = core.opSync("op_read_file_sync", pathFromURL(path));
+      const decoder = new TextDecoder();
+      return decoder.decode(buf);
+    },
+    writeFileSync(path, buf) {
+      return core.opSync("op_write_file_sync", pathFromURL(path), buf);
+    },
+  };
+
+  function registerErrors() {
+    core.registerErrorBuilder(
+      "DOMExceptionOperationError",
+      function DOMExceptionOperationError(msg) {
+        return new DOMException(msg, "OperationError");
+      },
+    );
+  }
+
+  let hasBootstrapped = false;
+
+  function bootstrapRuntime({ args, cwd }) {
+    core.setMacrotaskCallback(timers.handleTimerMacrotask);
+    if (hasBootstrapped) {
+      throw new Error("Runtime has already been bootstrapped.");
+    }
+    delete globalThis.__bootstrap;
+    delete globalThis.bootstrap;
+    hasBootstrapped = true;
+
+    registerErrors();
+
+    Object.defineProperties(globalThis, windowOrWorkerGlobalScope);
+    Object.defineProperties(globalThis, mainRuntimeGlobalProperties);
+    Object.setPrototypeOf(globalThis, Window.prototype);
+    eventTarget.setEventTargetData(globalThis);
+
+    denoNs.args = args;
+    denoNs.cwd = () => cwd;
+    util.immutableDefine(globalThis, "Deno", denoNs);
+    Object.freeze(globalThis.Deno);
+
+    core.ops();
+    Error.prepareStackTrace = core.createPrepareStackTrace();
+  }
+
+  ObjectDefineProperties(globalThis, {
+    bootstrap: {
+      value: bootstrapRuntime,
+      configurable: true,
+    },
+  });
+})(globalThis);
--- a/cts_runner/src/main.rs
+++ b/cts_runner/src/main.rs
@ -0,0 +1,149 @@
+use std::fmt;
+use std::io::Read;
+use std::io::Write;
+use std::rc::Rc;
+
+use deno_core::error::anyhow;
+use deno_core::error::AnyError;
+use deno_core::located_script_name;
+use deno_core::resolve_url_or_path;
+use deno_core::serde_json;
+use deno_core::serde_json::json;
+use deno_core::JsRuntime;
+use deno_core::OpState;
+use deno_core::RuntimeOptions;
+use deno_core::ZeroCopyBuf;
+use deno_web::BlobStore;
+use termcolor::Ansi;
+use termcolor::Color::Red;
+use termcolor::ColorSpec;
+use termcolor::WriteColor;
+
+#[tokio::main(flavor = "current_thread")]
+async fn main() {
+    unwrap_or_exit(run().await)
+}
+
+async fn run() -> Result<(), AnyError> {
+    let args = std::env::args().collect::<Vec<_>>();
+    let url = args
+        .get(1)
+        .ok_or_else(|| anyhow!("missing specifier in first command line argument"))?;
+    let specifier = resolve_url_or_path(url)?;
+
+    let options = RuntimeOptions {
+        module_loader: Some(Rc::new(deno_core::FsModuleLoader)),
+        get_error_class_fn: Some(&get_error_class_name),
+        extensions: vec![
+            deno_webidl::init(),
+            deno_console::init(),
+            deno_url::init(),
+            deno_web::init(BlobStore::default(), None),
+            deno_timers::init::<deno_timers::NoTimersPermission>(),
+            deno_webgpu::init(true),
+            extension(),
+        ],
+        ..Default::default()
+    };
+    let mut isolate = JsRuntime::new(options);
+    let args: Vec<String> = std::env::args().skip(2).collect();
+    let cfg = json!({"args": args, "cwd": std::env::current_dir().unwrap().to_string_lossy() });
+    let bootstrap_script = format!("globalThis.bootstrap({})", serde_json::to_string(&cfg)?);
+    isolate.execute_script(&located_script_name!(), &bootstrap_script)?;
+
+    isolate
+        .op_state()
+        .borrow_mut()
+        .put(deno_timers::NoTimersPermission);
+
+    let mod_id = isolate.load_module(&specifier, None).await?;
+    let mod_rx = isolate.mod_evaluate(mod_id);
+
+    let rx = tokio::spawn(async move {
+        match mod_rx.await {
+            Ok(err @ Err(_)) => err,
+            _ => Ok(()),
+        }
+    });
+
+    isolate.run_event_loop(false).await?;
+    rx.await.unwrap()?;
+
+    Ok(())
+}
+
+fn extension() -> deno_core::Extension {
+    deno_core::Extension::builder()
+        .ops(vec![
+            ("op_exit", deno_core::op_sync(op_exit)),
+            ("op_read_file_sync", deno_core::op_sync(op_read_file_sync)),
+            ("op_write_file_sync", deno_core::op_sync(op_write_file_sync)),
+        ])
+        .js(deno_core::include_js_files!(
+          prefix "deno:cts_runner",
+          "src/bootstrap.js",
+        ))
+        .build()
+}
+
+fn op_exit(_state: &mut OpState, code: i32, _: ()) -> Result<(), AnyError> {
+    std::process::exit(code)
+}
+
+fn op_read_file_sync(_state: &mut OpState, path: String, _: ()) -> Result<ZeroCopyBuf, AnyError> {
+    let path = std::path::Path::new(&path);
+    let mut file = std::fs::File::open(path)?;
+    let mut buf = Vec::new();
+    file.read_to_end(&mut buf)?;
+    Ok(ZeroCopyBuf::from(buf))
+}
+
+fn op_write_file_sync(
+    _state: &mut OpState,
+    path: String,
+    buf: ZeroCopyBuf,
+) -> Result<(), AnyError> {
+    let path = std::path::Path::new(&path);
+    let mut file = std::fs::File::create(path)?;
+    file.write_all(&buf)?;
+    Ok(())
+}
+
+fn get_error_class_name(e: &AnyError) -> &'static str {
+    deno_core::error::get_custom_error_class(e)
+        .or_else(|| deno_webgpu::error::get_error_class_name(e))
+        .unwrap_or_else(|| {
+            panic!(
+                "Error '{}' contains boxed error of unsupported type:{}",
+                e,
+                e.chain()
+                    .map(|e| format!("\n  {:?}", e))
+                    .collect::<String>()
+            );
+        })
+}
+
+fn unwrap_or_exit<T>(result: Result<T, AnyError>) -> T {
+    match result {
+        Ok(value) => value,
+        Err(error) => {
+            eprintln!("{}: {:?}", red_bold("error"), error);
+            std::process::exit(1);
+        }
+    }
+}
+
+fn style<S: AsRef<str>>(s: S, colorspec: ColorSpec) -> impl fmt::Display {
+    let mut v = Vec::new();
+    let mut ansi_writer = Ansi::new(&mut v);
+    ansi_writer.set_color(&colorspec).unwrap();
+    ansi_writer.write_all(s.as_ref().as_bytes()).unwrap();
+    ansi_writer.reset().unwrap();
+    String::from_utf8_lossy(&v).into_owned()
+}
+
+fn red_bold<S: AsRef<str>>(s: S) -> impl fmt::Display {
+    let mut style_spec = ColorSpec::new();
+    style_spec.set_fg(Some(Red)).set_bold(true);
+    style(s, style_spec)
+}
--- a/cts_runner/tests/integration.rs
+++ b/cts_runner/tests/integration.rs
@ -0,0 +1,27 @@
+use std::path::PathBuf;
+
+pub fn target_dir() -> PathBuf {
+    let current_exe = std::env::current_exe().unwrap();
+    let target_dir = current_exe.parent().unwrap().parent().unwrap();
+    target_dir.into()
+}
+
+pub fn cts_runner_exe_path() -> PathBuf {
+    // Something like /Users/lucacasonato/src/wgpu/target/debug/cts_runner
+    let mut p = target_dir().join("cts_runner");
+    if cfg!(windows) {
+        p.set_extension("exe");
+    }
+    p
+}
+
+#[test]
+fn hello_compute_example() {
+    let output = std::process::Command::new(cts_runner_exe_path())
+        .arg("examples/hello-compute.js")
+        .spawn()
+        .unwrap()
+        .wait_with_output()
+        .unwrap();
+    assert!(output.status.success())
+}
--- a/deno_webgpu/01_webgpu.js
+++ b/deno_webgpu/01_webgpu.js
--- a/deno_webgpu/02_idl_types.js
+++ b/deno_webgpu/02_idl_types.js
--- a/deno_webgpu/Cargo.toml
+++ b/deno_webgpu/Cargo.toml
@ -0,0 +1,21 @@
+# Copyright 2018-2021 the Deno authors. All rights reserved. MIT license.
+
+[package]
+name = "deno_webgpu"
+version = "0.17.0"
+authors = ["the Deno authors"]
+edition = "2018"
+license = "MIT"
+readme = "README.md"
+repository = "https://github.com/denoland/deno"
+description = "WebGPU implementation for Deno"
+
+[lib]
+path = "lib.rs"
+
+[dependencies]
+deno_core = { git = "https://github.com/denoland/deno", rev = "ca75752e5a9499a0a997809f02b18c2ba1ecd58d" }
+serde = { version = "1.0", features = ["derive"] }
+tokio = { version = "1.10", features = ["full"] }
+wgpu-core = { path = "../wgpu-core", features = ["trace", "replay", "serde"] }
+wgpu-types = { path = "../wgpu-types", features = ["trace", "replay", "serde"] }
--- a/deno_webgpu/LICENSE.md
+++ b/deno_webgpu/LICENSE.md
@ -0,0 +1,20 @@
+MIT License
+
+Copyright 2018-2021 the Deno authors
+
+Permission is hereby granted, free of charge, to any person obtaining a copy of
+this software and associated documentation files (the "Software"), to deal in
+the Software without restriction, including without limitation the rights to
+use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies of
+the Software, and to permit persons to whom the Software is furnished to do so,
+subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in all
+copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, FITNESS
+FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR
+COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER
+IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN
+CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE.
--- a/deno_webgpu/README.md
+++ b/deno_webgpu/README.md
@ -0,0 +1,35 @@
+# deno_webgpu
+
+This op crate implements the WebGPU API as defined in
+https://gpuweb.github.io/gpuweb/ in Deno. The implementation targets the spec
+draft as of February 22, 2021. The spec is still very much in flux. This op
+crate tries to stay up to date with the spec, but is constrained by the features
+implemented in our GPU backend library [wgpu](https://github.com/gfx-rs/wgpu).
+
+The spec is still very bare bones, and is still missing many details. As the
+spec becomes more concrete, we will implement to follow the spec more closely.
+
+In addition, setting the `DENO_WEBGPU_TRACE` environmental variable will output
+a
+[wgpu trace](https://github.com/gfx-rs/wgpu/wiki/Debugging-wgpu-Applications#tracing-infrastructure)
+to the specified directory.
+
+For testing this op crate will make use of the WebGPU conformance tests suite,
+running through our WPT runner. This will be used to validate implementation
+conformance.
+
+GitHub CI doesn't run with GPUs, so testing relies on software like DX WARP &
+Vulkan lavapipe. Currently only using DX WARP works, so tests are only run on
+Windows.
+
+## Links
+
+Specification: https://gpuweb.github.io/gpuweb/
+
+Design documents: https://github.com/gpuweb/gpuweb/tree/main/design
+
+Conformance tests suite: https://github.com/gpuweb/cts
+
+WebGPU examples for Deno: https://github.com/crowlKats/webgpu-examples
+
+wgpu-users matrix channel: https://matrix.to/#/#wgpu-users:matrix.org
--- a/deno_webgpu/binding.rs
+++ b/deno_webgpu/binding.rs
@ -0,0 +1,354 @@
+// Copyright 2018-2021 the Deno authors. All rights reserved. MIT license.
+
+use deno_core::error::AnyError;
+use deno_core::ResourceId;
+use deno_core::{OpState, Resource};
+use serde::Deserialize;
+use std::borrow::Cow;
+use std::convert::{TryFrom, TryInto};
+
+use super::error::WebGpuResult;
+
+pub(crate) struct WebGpuBindGroupLayout(pub(crate) wgpu_core::id::BindGroupLayoutId);
+impl Resource for WebGpuBindGroupLayout {
+    fn name(&self) -> Cow<str> {
+        "webGPUBindGroupLayout".into()
+    }
+}
+
+pub(crate) struct WebGpuBindGroup(pub(crate) wgpu_core::id::BindGroupId);
+impl Resource for WebGpuBindGroup {
+    fn name(&self) -> Cow<str> {
+        "webGPUBindGroup".into()
+    }
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+struct GpuBufferBindingLayout {
+    r#type: GpuBufferBindingType,
+    has_dynamic_offset: bool,
+    min_binding_size: u64,
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "kebab-case")]
+enum GpuBufferBindingType {
+    Uniform,
+    Storage,
+    ReadOnlyStorage,
+}
+
+impl From<GpuBufferBindingType> for wgpu_types::BufferBindingType {
+    fn from(binding_type: GpuBufferBindingType) -> Self {
+        match binding_type {
+            GpuBufferBindingType::Uniform => wgpu_types::BufferBindingType::Uniform,
+            GpuBufferBindingType::Storage => {
+                wgpu_types::BufferBindingType::Storage { read_only: false }
+            }
+            GpuBufferBindingType::ReadOnlyStorage => {
+                wgpu_types::BufferBindingType::Storage { read_only: true }
+            }
+        }
+    }
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+struct GpuSamplerBindingLayout {
+    r#type: GpuSamplerBindingType,
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "kebab-case")]
+enum GpuSamplerBindingType {
+    Filtering,
+    NonFiltering,
+    Comparison,
+}
+
+impl From<GpuSamplerBindingType> for wgpu_types::BindingType {
+    fn from(binding_type: GpuSamplerBindingType) -> Self {
+        match binding_type {
+            GpuSamplerBindingType::Filtering => wgpu_types::BindingType::Sampler {
+                filtering: true,
+                comparison: false,
+            },
+            GpuSamplerBindingType::NonFiltering => wgpu_types::BindingType::Sampler {
+                filtering: false,
+                comparison: false,
+            },
+            GpuSamplerBindingType::Comparison => wgpu_types::BindingType::Sampler {
+                filtering: true,
+                comparison: true,
+            },
+        }
+    }
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+struct GpuTextureBindingLayout {
+    sample_type: GpuTextureSampleType,
+    view_dimension: wgpu_types::TextureViewDimension,
+    multisampled: bool,
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "kebab-case")]
+enum GpuTextureSampleType {
+    Float,
+    UnfilterableFloat,
+    Depth,
+    Sint,
+    Uint,
+}
+
+impl From<GpuTextureSampleType> for wgpu_types::TextureSampleType {
+    fn from(sample_type: GpuTextureSampleType) -> Self {
+        match sample_type {
+            GpuTextureSampleType::Float => {
+                wgpu_types::TextureSampleType::Float { filterable: true }
+            }
+            GpuTextureSampleType::UnfilterableFloat => {
+                wgpu_types::TextureSampleType::Float { filterable: false }
+            }
+            GpuTextureSampleType::Depth => wgpu_types::TextureSampleType::Depth,
+            GpuTextureSampleType::Sint => wgpu_types::TextureSampleType::Sint,
+            GpuTextureSampleType::Uint => wgpu_types::TextureSampleType::Uint,
+        }
+    }
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+struct GpuStorageTextureBindingLayout {
+    access: GpuStorageTextureAccess,
+    format: wgpu_types::TextureFormat,
+    view_dimension: wgpu_types::TextureViewDimension,
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "kebab-case")]
+enum GpuStorageTextureAccess {
+    WriteOnly,
+}
+
+impl From<GpuStorageTextureAccess> for wgpu_types::StorageTextureAccess {
+    fn from(access: GpuStorageTextureAccess) -> Self {
+        match access {
+            GpuStorageTextureAccess::WriteOnly => wgpu_types::StorageTextureAccess::WriteOnly,
+        }
+    }
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+struct GpuBindGroupLayoutEntry {
+    binding: u32,
+    visibility: u32,
+    #[serde(flatten)]
+    binding_type: GpuBindingType,
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+enum GpuBindingType {
+    Buffer(GpuBufferBindingLayout),
+    Sampler(GpuSamplerBindingLayout),
+    Texture(GpuTextureBindingLayout),
+    StorageTexture(GpuStorageTextureBindingLayout),
+}
+
+impl TryFrom<GpuBindingType> for wgpu_types::BindingType {
+    type Error = AnyError;
+
+    fn try_from(binding_type: GpuBindingType) -> Result<wgpu_types::BindingType, Self::Error> {
+        let binding_type = match binding_type {
+            GpuBindingType::Buffer(buffer) => wgpu_types::BindingType::Buffer {
+                ty: buffer.r#type.into(),
+                has_dynamic_offset: buffer.has_dynamic_offset,
+                min_binding_size: std::num::NonZeroU64::new(buffer.min_binding_size),
+            },
+            GpuBindingType::Sampler(sampler) => sampler.r#type.into(),
+            GpuBindingType::Texture(texture) => wgpu_types::BindingType::Texture {
+                sample_type: texture.sample_type.into(),
+                view_dimension: texture.view_dimension,
+                multisampled: texture.multisampled,
+            },
+            GpuBindingType::StorageTexture(storage_texture) => {
+                wgpu_types::BindingType::StorageTexture {
+                    access: storage_texture.access.into(),
+                    format: storage_texture.format,
+                    view_dimension: storage_texture.view_dimension,
+                }
+            }
+        };
+        Ok(binding_type)
+    }
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct CreateBindGroupLayoutArgs {
+    device_rid: ResourceId,
+    label: Option<String>,
+    entries: Vec<GpuBindGroupLayoutEntry>,
+}
+
+pub fn op_webgpu_create_bind_group_layout(
+    state: &mut OpState,
+    args: CreateBindGroupLayoutArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let instance = state.borrow::<super::Instance>();
+    let device_resource = state
+        .resource_table
+        .get::<super::WebGpuDevice>(args.device_rid)?;
+    let device = device_resource.0;
+
+    let mut entries = vec![];
+
+    for entry in args.entries {
+        entries.push(wgpu_types::BindGroupLayoutEntry {
+            binding: entry.binding,
+            visibility: wgpu_types::ShaderStages::from_bits(entry.visibility).unwrap(),
+            ty: entry.binding_type.try_into()?,
+            count: None, // native-only
+        });
+    }
+
+    let descriptor = wgpu_core::binding_model::BindGroupLayoutDescriptor {
+        label: args.label.map(Cow::from),
+        entries: Cow::from(entries),
+    };
+
+    gfx_put!(device => instance.device_create_bind_group_layout(
+    device,
+    &descriptor,
+    std::marker::PhantomData
+  ) => state, WebGpuBindGroupLayout)
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct CreatePipelineLayoutArgs {
+    device_rid: ResourceId,
+    label: Option<String>,
+    bind_group_layouts: Vec<u32>,
+}
+
+pub fn op_webgpu_create_pipeline_layout(
+    state: &mut OpState,
+    args: CreatePipelineLayoutArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let instance = state.borrow::<super::Instance>();
+    let device_resource = state
+        .resource_table
+        .get::<super::WebGpuDevice>(args.device_rid)?;
+    let device = device_resource.0;
+
+    let mut bind_group_layouts = vec![];
+
+    for rid in &args.bind_group_layouts {
+        let bind_group_layout = state.resource_table.get::<WebGpuBindGroupLayout>(*rid)?;
+        bind_group_layouts.push(bind_group_layout.0);
+    }
+
+    let descriptor = wgpu_core::binding_model::PipelineLayoutDescriptor {
+        label: args.label.map(Cow::from),
+        bind_group_layouts: Cow::from(bind_group_layouts),
+        push_constant_ranges: Default::default(),
+    };
+
+    gfx_put!(device => instance.device_create_pipeline_layout(
+    device,
+    &descriptor,
+    std::marker::PhantomData
+  ) => state, super::pipeline::WebGpuPipelineLayout)
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+struct GpuBindGroupEntry {
+    binding: u32,
+    kind: String,
+    resource: ResourceId,
+    offset: Option<u64>,
+    size: Option<u64>,
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct CreateBindGroupArgs {
+    device_rid: ResourceId,
+    label: Option<String>,
+    layout: ResourceId,
+    entries: Vec<GpuBindGroupEntry>,
+}
+
+pub fn op_webgpu_create_bind_group(
+    state: &mut OpState,
+    args: CreateBindGroupArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let instance = state.borrow::<super::Instance>();
+    let device_resource = state
+        .resource_table
+        .get::<super::WebGpuDevice>(args.device_rid)?;
+    let device = device_resource.0;
+
+    let mut entries = vec![];
+
+    for entry in &args.entries {
+        let e = wgpu_core::binding_model::BindGroupEntry {
+            binding: entry.binding,
+            resource: match entry.kind.as_str() {
+                "GPUSampler" => {
+                    let sampler_resource = state
+                        .resource_table
+                        .get::<super::sampler::WebGpuSampler>(entry.resource)?;
+                    wgpu_core::binding_model::BindingResource::Sampler(sampler_resource.0)
+                }
+                "GPUTextureView" => {
+                    let texture_view_resource =
+                        state
+                            .resource_table
+                            .get::<super::texture::WebGpuTextureView>(entry.resource)?;
+                    wgpu_core::binding_model::BindingResource::TextureView(texture_view_resource.0)
+                }
+                "GPUBufferBinding" => {
+                    let buffer_resource = state
+                        .resource_table
+                        .get::<super::buffer::WebGpuBuffer>(entry.resource)?;
+                    wgpu_core::binding_model::BindingResource::Buffer(
+                        wgpu_core::binding_model::BufferBinding {
+                            buffer_id: buffer_resource.0,
+                            offset: entry.offset.unwrap_or(0),
+                            size: std::num::NonZeroU64::new(entry.size.unwrap_or(0)),
+                        },
+                    )
+                }
+                _ => unreachable!(),
+            },
+        };
+        entries.push(e);
+    }
+
+    let bind_group_layout = state
+        .resource_table
+        .get::<WebGpuBindGroupLayout>(args.layout)?;
+
+    let descriptor = wgpu_core::binding_model::BindGroupDescriptor {
+        label: args.label.map(Cow::from),
+        layout: bind_group_layout.0,
+        entries: Cow::from(entries),
+    };
+
+    gfx_put!(device => instance.device_create_bind_group(
+    device,
+    &descriptor,
+    std::marker::PhantomData
+  ) => state, WebGpuBindGroup)
+}
--- a/deno_webgpu/buffer.rs
+++ b/deno_webgpu/buffer.rs
@ -0,0 +1,225 @@
+// Copyright 2018-2021 the Deno authors. All rights reserved. MIT license.
+
+use deno_core::error::null_opbuf;
+use deno_core::error::type_error;
+use deno_core::error::AnyError;
+use deno_core::futures::channel::oneshot;
+use deno_core::OpState;
+use deno_core::Resource;
+use deno_core::ResourceId;
+use deno_core::ZeroCopyBuf;
+use serde::Deserialize;
+use std::borrow::Cow;
+use std::cell::RefCell;
+use std::rc::Rc;
+use std::time::Duration;
+
+use super::error::DomExceptionOperationError;
+use super::error::WebGpuResult;
+
+pub(crate) struct WebGpuBuffer(pub(crate) wgpu_core::id::BufferId);
+impl Resource for WebGpuBuffer {
+    fn name(&self) -> Cow<str> {
+        "webGPUBuffer".into()
+    }
+}
+
+struct WebGpuBufferMapped(*mut u8, usize);
+impl Resource for WebGpuBufferMapped {
+    fn name(&self) -> Cow<str> {
+        "webGPUBufferMapped".into()
+    }
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct CreateBufferArgs {
+    device_rid: ResourceId,
+    label: Option<String>,
+    size: u64,
+    usage: u32,
+    mapped_at_creation: bool,
+}
+
+pub fn op_webgpu_create_buffer(
+    state: &mut OpState,
+    args: CreateBufferArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let instance = state.borrow::<super::Instance>();
+    let device_resource = state
+        .resource_table
+        .get::<super::WebGpuDevice>(args.device_rid)?;
+    let device = device_resource.0;
+
+    let descriptor = wgpu_core::resource::BufferDescriptor {
+        label: args.label.map(Cow::from),
+        size: args.size,
+        usage: wgpu_types::BufferUsages::from_bits(args.usage)
+            .ok_or_else(|| type_error("usage is not valid"))?,
+        mapped_at_creation: args.mapped_at_creation,
+    };
+
+    gfx_put!(device => instance.device_create_buffer(
+    device,
+    &descriptor,
+    std::marker::PhantomData
+  ) => state, WebGpuBuffer)
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct BufferGetMapAsyncArgs {
+    buffer_rid: ResourceId,
+    device_rid: ResourceId,
+    mode: u32,
+    offset: u64,
+    size: u64,
+}
+
+pub async fn op_webgpu_buffer_get_map_async(
+    state: Rc<RefCell<OpState>>,
+    args: BufferGetMapAsyncArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let (sender, receiver) = oneshot::channel::<Result<(), AnyError>>();
+
+    let device;
+    {
+        let state_ = state.borrow();
+        let instance = state_.borrow::<super::Instance>();
+        let buffer_resource = state_.resource_table.get::<WebGpuBuffer>(args.buffer_rid)?;
+        let buffer = buffer_resource.0;
+        let device_resource = state_
+            .resource_table
+            .get::<super::WebGpuDevice>(args.device_rid)?;
+        device = device_resource.0;
+
+        let boxed_sender = Box::new(sender);
+        let sender_ptr = Box::into_raw(boxed_sender) as *mut u8;
+
+        extern "C" fn buffer_map_future_wrapper(
+            status: wgpu_core::resource::BufferMapAsyncStatus,
+            user_data: *mut u8,
+        ) {
+            let sender_ptr = user_data as *mut oneshot::Sender<Result<(), AnyError>>;
+            let boxed_sender = unsafe { Box::from_raw(sender_ptr) };
+            boxed_sender
+                .send(match status {
+                    wgpu_core::resource::BufferMapAsyncStatus::Success => Ok(()),
+                    _ => unreachable!(), // TODO
+                })
+                .unwrap();
+        }
+
+        // TODO(lucacasonato): error handling
+        let maybe_err = gfx_select!(buffer => instance.buffer_map_async(
+          buffer,
+          args.offset..(args.offset + args.size),
+          wgpu_core::resource::BufferMapOperation {
+            host: match args.mode {
+              1 => wgpu_core::device::HostMap::Read,
+              2 => wgpu_core::device::HostMap::Write,
+              _ => unreachable!(),
+            },
+            callback: buffer_map_future_wrapper,
+            user_data: sender_ptr,
+          }
+        ))
+        .err();
+
+        if maybe_err.is_some() {
+            return Ok(WebGpuResult::maybe_err(maybe_err));
+        }
+    }
+
+    let done = Rc::new(RefCell::new(false));
+    let done_ = done.clone();
+    let device_poll_fut = async move {
+        while !*done.borrow() {
+            {
+                let state = state.borrow();
+                let instance = state.borrow::<super::Instance>();
+                gfx_select!(device => instance.device_poll(device, false)).unwrap()
+            }
+            tokio::time::sleep(Duration::from_millis(10)).await;
+        }
+        Ok::<(), AnyError>(())
+    };
+
+    let receiver_fut = async move {
+        receiver.await??;
+        let mut done = done_.borrow_mut();
+        *done = true;
+        Ok::<(), AnyError>(())
+    };
+
+    tokio::try_join!(device_poll_fut, receiver_fut)?;
+
+    Ok(WebGpuResult::empty())
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct BufferGetMappedRangeArgs {
+    buffer_rid: ResourceId,
+    offset: u64,
+    size: Option<u64>,
+}
+
+pub fn op_webgpu_buffer_get_mapped_range(
+    state: &mut OpState,
+    args: BufferGetMappedRangeArgs,
+    zero_copy: Option<ZeroCopyBuf>,
+) -> Result<WebGpuResult, AnyError> {
+    let mut zero_copy = zero_copy.ok_or_else(null_opbuf)?;
+    let instance = state.borrow::<super::Instance>();
+    let buffer_resource = state.resource_table.get::<WebGpuBuffer>(args.buffer_rid)?;
+    let buffer = buffer_resource.0;
+
+    let (slice_pointer, range_size) = gfx_select!(buffer => instance.buffer_get_mapped_range(
+      buffer,
+      args.offset,
+      args.size
+    ))
+    .map_err(|e| DomExceptionOperationError::new(&e.to_string()))?;
+
+    let slice = unsafe { std::slice::from_raw_parts_mut(slice_pointer, range_size as usize) };
+    zero_copy.copy_from_slice(slice);
+
+    let rid = state
+        .resource_table
+        .add(WebGpuBufferMapped(slice_pointer, range_size as usize));
+
+    Ok(WebGpuResult::rid(rid))
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct BufferUnmapArgs {
+    buffer_rid: ResourceId,
+    mapped_rid: ResourceId,
+}
+
+pub fn op_webgpu_buffer_unmap(
+    state: &mut OpState,
+    args: BufferUnmapArgs,
+    zero_copy: Option<ZeroCopyBuf>,
+) -> Result<WebGpuResult, AnyError> {
+    let mapped_resource = state
+        .resource_table
+        .take::<WebGpuBufferMapped>(args.mapped_rid)?;
+    let instance = state.borrow::<super::Instance>();
+    let buffer_resource = state.resource_table.get::<WebGpuBuffer>(args.buffer_rid)?;
+    let buffer = buffer_resource.0;
+
+    let slice_pointer = mapped_resource.0;
+    let size = mapped_resource.1;
+
+    if let Some(buffer) = zero_copy {
+        let slice = unsafe { std::slice::from_raw_parts_mut(slice_pointer, size) };
+        slice.copy_from_slice(&buffer);
+    }
+
+    gfx_ok!(buffer => instance.buffer_unmap(buffer))
+}
--- a/deno_webgpu/bundle.rs
+++ b/deno_webgpu/bundle.rs
@ -0,0 +1,442 @@
+// Copyright 2018-2021 the Deno authors. All rights reserved. MIT license.
+
+use deno_core::error::AnyError;
+use deno_core::ResourceId;
+use deno_core::ZeroCopyBuf;
+use deno_core::{OpState, Resource};
+use serde::Deserialize;
+use std::borrow::Cow;
+use std::cell::RefCell;
+use std::rc::Rc;
+
+use super::error::WebGpuResult;
+
+struct WebGpuRenderBundleEncoder(RefCell<wgpu_core::command::RenderBundleEncoder>);
+impl Resource for WebGpuRenderBundleEncoder {
+    fn name(&self) -> Cow<str> {
+        "webGPURenderBundleEncoder".into()
+    }
+}
+
+pub(crate) struct WebGpuRenderBundle(pub(crate) wgpu_core::id::RenderBundleId);
+impl Resource for WebGpuRenderBundle {
+    fn name(&self) -> Cow<str> {
+        "webGPURenderBundle".into()
+    }
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct CreateRenderBundleEncoderArgs {
+    device_rid: ResourceId,
+    label: Option<String>,
+    color_formats: Vec<wgpu_types::TextureFormat>,
+    depth_stencil_format: Option<wgpu_types::TextureFormat>,
+    sample_count: u32,
+    depth_read_only: bool,
+    stencil_read_only: bool,
+}
+
+pub fn op_webgpu_create_render_bundle_encoder(
+    state: &mut OpState,
+    args: CreateRenderBundleEncoderArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let device_resource = state
+        .resource_table
+        .get::<super::WebGpuDevice>(args.device_rid)?;
+    let device = device_resource.0;
+
+    let mut color_formats = vec![];
+
+    for format in args.color_formats {
+        color_formats.push(format);
+    }
+
+    let depth_stencil = if let Some(format) = args.depth_stencil_format {
+        Some(wgpu_types::RenderBundleDepthStencil {
+            format,
+            depth_read_only: args.depth_read_only,
+            stencil_read_only: args.stencil_read_only,
+        })
+    } else {
+        None
+    };
+
+    let descriptor = wgpu_core::command::RenderBundleEncoderDescriptor {
+        label: args.label.map(Cow::from),
+        color_formats: Cow::from(color_formats),
+        sample_count: args.sample_count,
+        depth_stencil,
+    };
+
+    let res = wgpu_core::command::RenderBundleEncoder::new(&descriptor, device, None);
+    let (render_bundle_encoder, maybe_err) = match res {
+        Ok(encoder) => (encoder, None),
+        Err(e) => (
+            wgpu_core::command::RenderBundleEncoder::dummy(device),
+            Some(e),
+        ),
+    };
+
+    let rid = state
+        .resource_table
+        .add(WebGpuRenderBundleEncoder(RefCell::new(
+            render_bundle_encoder,
+        )));
+
+    Ok(WebGpuResult::rid_err(rid, maybe_err))
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct RenderBundleEncoderFinishArgs {
+    render_bundle_encoder_rid: ResourceId,
+    label: Option<String>,
+}
+
+pub fn op_webgpu_render_bundle_encoder_finish(
+    state: &mut OpState,
+    args: RenderBundleEncoderFinishArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let render_bundle_encoder_resource = state
+        .resource_table
+        .take::<WebGpuRenderBundleEncoder>(args.render_bundle_encoder_rid)?;
+    let render_bundle_encoder = Rc::try_unwrap(render_bundle_encoder_resource)
+        .ok()
+        .expect("unwrapping render_bundle_encoder_resource should succeed")
+        .0
+        .into_inner();
+    let instance = state.borrow::<super::Instance>();
+
+    gfx_put!(render_bundle_encoder.parent() => instance.render_bundle_encoder_finish(
+    render_bundle_encoder,
+    &wgpu_core::command::RenderBundleDescriptor {
+      label: args.label.map(Cow::from),
+    },
+    std::marker::PhantomData
+  ) => state, WebGpuRenderBundle)
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct RenderBundleEncoderSetBindGroupArgs {
+    render_bundle_encoder_rid: ResourceId,
+    index: u32,
+    bind_group: ResourceId,
+    dynamic_offsets_data: ZeroCopyBuf,
+    dynamic_offsets_data_start: usize,
+    dynamic_offsets_data_length: usize,
+}
+
+pub fn op_webgpu_render_bundle_encoder_set_bind_group(
+    state: &mut OpState,
+    args: RenderBundleEncoderSetBindGroupArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let bind_group_resource = state
+        .resource_table
+        .get::<super::binding::WebGpuBindGroup>(args.bind_group)?;
+    let render_bundle_encoder_resource = state
+        .resource_table
+        .get::<WebGpuRenderBundleEncoder>(args.render_bundle_encoder_rid)?;
+
+    // Align the data
+    assert!(args.dynamic_offsets_data.len() % std::mem::size_of::<u32>() == 0);
+    // SAFETY: A u8 to u32 cast is safe because we asserted that the length is a
+    // multiple of 4.
+    let (prefix, dynamic_offsets_data, suffix) =
+        unsafe { args.dynamic_offsets_data.align_to::<u32>() };
+    assert!(prefix.is_empty());
+    assert!(suffix.is_empty());
+
+    let start = args.dynamic_offsets_data_start;
+    let len = args.dynamic_offsets_data_length;
+
+    // Assert that length and start are both in bounds
+    assert!(start <= dynamic_offsets_data.len());
+    assert!(len <= dynamic_offsets_data.len() - start);
+
+    let dynamic_offsets_data: &[u32] = &dynamic_offsets_data[start..start + len];
+
+    // SAFETY: the raw pointer and length are of the same slice, and that slice
+    // lives longer than the below function invocation.
+    unsafe {
+        wgpu_core::command::bundle_ffi::wgpu_render_bundle_set_bind_group(
+            &mut render_bundle_encoder_resource.0.borrow_mut(),
+            args.index,
+            bind_group_resource.0,
+            dynamic_offsets_data.as_ptr(),
+            dynamic_offsets_data.len(),
+        );
+    }
+
+    Ok(WebGpuResult::empty())
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct RenderBundleEncoderPushDebugGroupArgs {
+    render_bundle_encoder_rid: ResourceId,
+    group_label: String,
+}
+
+pub fn op_webgpu_render_bundle_encoder_push_debug_group(
+    state: &mut OpState,
+    args: RenderBundleEncoderPushDebugGroupArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let render_bundle_encoder_resource = state
+        .resource_table
+        .get::<WebGpuRenderBundleEncoder>(args.render_bundle_encoder_rid)?;
+
+    let label = std::ffi::CString::new(args.group_label).unwrap();
+    // SAFETY: the string the raw pointer points to lives longer than the below
+    // function invocation.
+    unsafe {
+        wgpu_core::command::bundle_ffi::wgpu_render_bundle_push_debug_group(
+            &mut render_bundle_encoder_resource.0.borrow_mut(),
+            label.as_ptr(),
+        );
+    }
+
+    Ok(WebGpuResult::empty())
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct RenderBundleEncoderPopDebugGroupArgs {
+    render_bundle_encoder_rid: ResourceId,
+}
+
+pub fn op_webgpu_render_bundle_encoder_pop_debug_group(
+    state: &mut OpState,
+    args: RenderBundleEncoderPopDebugGroupArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let render_bundle_encoder_resource = state
+        .resource_table
+        .get::<WebGpuRenderBundleEncoder>(args.render_bundle_encoder_rid)?;
+
+    wgpu_core::command::bundle_ffi::wgpu_render_bundle_pop_debug_group(
+        &mut render_bundle_encoder_resource.0.borrow_mut(),
+    );
+
+    Ok(WebGpuResult::empty())
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct RenderBundleEncoderInsertDebugMarkerArgs {
+    render_bundle_encoder_rid: ResourceId,
+    marker_label: String,
+}
+
+pub fn op_webgpu_render_bundle_encoder_insert_debug_marker(
+    state: &mut OpState,
+    args: RenderBundleEncoderInsertDebugMarkerArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let render_bundle_encoder_resource = state
+        .resource_table
+        .get::<WebGpuRenderBundleEncoder>(args.render_bundle_encoder_rid)?;
+
+    let label = std::ffi::CString::new(args.marker_label).unwrap();
+    // SAFETY: the string the raw pointer points to lives longer than the below
+    // function invocation.
+    unsafe {
+        wgpu_core::command::bundle_ffi::wgpu_render_bundle_insert_debug_marker(
+            &mut render_bundle_encoder_resource.0.borrow_mut(),
+            label.as_ptr(),
+        );
+    }
+
+    Ok(WebGpuResult::empty())
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct RenderBundleEncoderSetPipelineArgs {
+    render_bundle_encoder_rid: ResourceId,
+    pipeline: ResourceId,
+}
+
+pub fn op_webgpu_render_bundle_encoder_set_pipeline(
+    state: &mut OpState,
+    args: RenderBundleEncoderSetPipelineArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let render_pipeline_resource = state
+        .resource_table
+        .get::<super::pipeline::WebGpuRenderPipeline>(args.pipeline)?;
+    let render_bundle_encoder_resource = state
+        .resource_table
+        .get::<WebGpuRenderBundleEncoder>(args.render_bundle_encoder_rid)?;
+
+    wgpu_core::command::bundle_ffi::wgpu_render_bundle_set_pipeline(
+        &mut render_bundle_encoder_resource.0.borrow_mut(),
+        render_pipeline_resource.0,
+    );
+
+    Ok(WebGpuResult::empty())
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct RenderBundleEncoderSetIndexBufferArgs {
+    render_bundle_encoder_rid: ResourceId,
+    buffer: ResourceId,
+    index_format: wgpu_types::IndexFormat,
+    offset: u64,
+    size: u64,
+}
+
+pub fn op_webgpu_render_bundle_encoder_set_index_buffer(
+    state: &mut OpState,
+    args: RenderBundleEncoderSetIndexBufferArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let buffer_resource = state
+        .resource_table
+        .get::<super::buffer::WebGpuBuffer>(args.buffer)?;
+    let render_bundle_encoder_resource = state
+        .resource_table
+        .get::<WebGpuRenderBundleEncoder>(args.render_bundle_encoder_rid)?;
+
+    render_bundle_encoder_resource
+        .0
+        .borrow_mut()
+        .set_index_buffer(
+            buffer_resource.0,
+            args.index_format,
+            args.offset,
+            std::num::NonZeroU64::new(args.size),
+        );
+
+    Ok(WebGpuResult::empty())
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct RenderBundleEncoderSetVertexBufferArgs {
+    render_bundle_encoder_rid: ResourceId,
+    slot: u32,
+    buffer: ResourceId,
+    offset: u64,
+    size: u64,
+}
+
+pub fn op_webgpu_render_bundle_encoder_set_vertex_buffer(
+    state: &mut OpState,
+    args: RenderBundleEncoderSetVertexBufferArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let buffer_resource = state
+        .resource_table
+        .get::<super::buffer::WebGpuBuffer>(args.buffer)?;
+    let render_bundle_encoder_resource = state
+        .resource_table
+        .get::<WebGpuRenderBundleEncoder>(args.render_bundle_encoder_rid)?;
+
+    wgpu_core::command::bundle_ffi::wgpu_render_bundle_set_vertex_buffer(
+        &mut render_bundle_encoder_resource.0.borrow_mut(),
+        args.slot,
+        buffer_resource.0,
+        args.offset,
+        std::num::NonZeroU64::new(args.size),
+    );
+
+    Ok(WebGpuResult::empty())
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct RenderBundleEncoderDrawArgs {
+    render_bundle_encoder_rid: ResourceId,
+    vertex_count: u32,
+    instance_count: u32,
+    first_vertex: u32,
+    first_instance: u32,
+}
+
+pub fn op_webgpu_render_bundle_encoder_draw(
+    state: &mut OpState,
+    args: RenderBundleEncoderDrawArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let render_bundle_encoder_resource = state
+        .resource_table
+        .get::<WebGpuRenderBundleEncoder>(args.render_bundle_encoder_rid)?;
+
+    wgpu_core::command::bundle_ffi::wgpu_render_bundle_draw(
+        &mut render_bundle_encoder_resource.0.borrow_mut(),
+        args.vertex_count,
+        args.instance_count,
+        args.first_vertex,
+        args.first_instance,
+    );
+
+    Ok(WebGpuResult::empty())
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct RenderBundleEncoderDrawIndexedArgs {
+    render_bundle_encoder_rid: ResourceId,
+    index_count: u32,
+    instance_count: u32,
+    first_index: u32,
+    base_vertex: i32,
+    first_instance: u32,
+}
+
+pub fn op_webgpu_render_bundle_encoder_draw_indexed(
+    state: &mut OpState,
+    args: RenderBundleEncoderDrawIndexedArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let render_bundle_encoder_resource = state
+        .resource_table
+        .get::<WebGpuRenderBundleEncoder>(args.render_bundle_encoder_rid)?;
+
+    wgpu_core::command::bundle_ffi::wgpu_render_bundle_draw_indexed(
+        &mut render_bundle_encoder_resource.0.borrow_mut(),
+        args.index_count,
+        args.instance_count,
+        args.first_index,
+        args.base_vertex,
+        args.first_instance,
+    );
+
+    Ok(WebGpuResult::empty())
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct RenderBundleEncoderDrawIndirectArgs {
+    render_bundle_encoder_rid: ResourceId,
+    indirect_buffer: ResourceId,
+    indirect_offset: u64,
+}
+
+pub fn op_webgpu_render_bundle_encoder_draw_indirect(
+    state: &mut OpState,
+    args: RenderBundleEncoderDrawIndirectArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let buffer_resource = state
+        .resource_table
+        .get::<super::buffer::WebGpuBuffer>(args.indirect_buffer)?;
+    let render_bundle_encoder_resource = state
+        .resource_table
+        .get::<WebGpuRenderBundleEncoder>(args.render_bundle_encoder_rid)?;
+
+    wgpu_core::command::bundle_ffi::wgpu_render_bundle_draw_indirect(
+        &mut render_bundle_encoder_resource.0.borrow_mut(),
+        buffer_resource.0,
+        args.indirect_offset,
+    );
+
+    Ok(WebGpuResult::empty())
+}
--- a/deno_webgpu/command_encoder.rs
+++ b/deno_webgpu/command_encoder.rs
@ -0,0 +1,591 @@
+// Copyright 2018-2021 the Deno authors. All rights reserved. MIT license.
+
+use deno_core::error::AnyError;
+use deno_core::ResourceId;
+use deno_core::{OpState, Resource};
+use serde::Deserialize;
+use std::borrow::Cow;
+use std::cell::RefCell;
+use std::num::NonZeroU32;
+
+use super::error::WebGpuResult;
+
+pub(crate) struct WebGpuCommandEncoder(pub(crate) wgpu_core::id::CommandEncoderId);
+impl Resource for WebGpuCommandEncoder {
+    fn name(&self) -> Cow<str> {
+        "webGPUCommandEncoder".into()
+    }
+}
+
+pub(crate) struct WebGpuCommandBuffer(pub(crate) wgpu_core::id::CommandBufferId);
+impl Resource for WebGpuCommandBuffer {
+    fn name(&self) -> Cow<str> {
+        "webGPUCommandBuffer".into()
+    }
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct CreateCommandEncoderArgs {
+    device_rid: ResourceId,
+    label: Option<String>,
+    _measure_execution_time: Option<bool>, // not yet implemented
+}
+
+pub fn op_webgpu_create_command_encoder(
+    state: &mut OpState,
+    args: CreateCommandEncoderArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let instance = state.borrow::<super::Instance>();
+    let device_resource = state
+        .resource_table
+        .get::<super::WebGpuDevice>(args.device_rid)?;
+    let device = device_resource.0;
+
+    let descriptor = wgpu_types::CommandEncoderDescriptor {
+        label: args.label.map(Cow::from),
+    };
+
+    gfx_put!(device => instance.device_create_command_encoder(
+    device,
+    &descriptor,
+    std::marker::PhantomData
+  ) => state, WebGpuCommandEncoder)
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct GpuRenderPassColorAttachment {
+    view: ResourceId,
+    resolve_target: Option<ResourceId>,
+    load_op: GpuLoadOp<wgpu_types::Color>,
+    store_op: wgpu_core::command::StoreOp,
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "kebab-case")]
+enum GpuLoadOp<T> {
+    Load,
+    Clear(T),
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+struct GpuRenderPassDepthStencilAttachment {
+    view: ResourceId,
+    depth_load_op: GpuLoadOp<f32>,
+    depth_store_op: wgpu_core::command::StoreOp,
+    depth_read_only: bool,
+    stencil_load_op: GpuLoadOp<u32>,
+    stencil_store_op: wgpu_core::command::StoreOp,
+    stencil_read_only: bool,
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct CommandEncoderBeginRenderPassArgs {
+    command_encoder_rid: ResourceId,
+    label: Option<String>,
+    color_attachments: Vec<GpuRenderPassColorAttachment>,
+    depth_stencil_attachment: Option<GpuRenderPassDepthStencilAttachment>,
+    _occlusion_query_set: Option<u32>, // not yet implemented
+}
+
+pub fn op_webgpu_command_encoder_begin_render_pass(
+    state: &mut OpState,
+    args: CommandEncoderBeginRenderPassArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let command_encoder_resource = state
+        .resource_table
+        .get::<WebGpuCommandEncoder>(args.command_encoder_rid)?;
+
+    let mut color_attachments = vec![];
+
+    for color_attachment in args.color_attachments {
+        let texture_view_resource = state
+            .resource_table
+            .get::<super::texture::WebGpuTextureView>(color_attachment.view)?;
+
+        let resolve_target = color_attachment
+            .resolve_target
+            .map(|rid| {
+                state
+                    .resource_table
+                    .get::<super::texture::WebGpuTextureView>(rid)
+            })
+            .transpose()?
+            .map(|texture| texture.0);
+
+        let attachment = wgpu_core::command::RenderPassColorAttachment {
+            view: texture_view_resource.0,
+            resolve_target,
+            channel: match color_attachment.load_op {
+                GpuLoadOp::Load => wgpu_core::command::PassChannel {
+                    load_op: wgpu_core::command::LoadOp::Load,
+                    store_op: color_attachment.store_op,
+                    clear_value: Default::default(),
+                    read_only: false,
+                },
+                GpuLoadOp::Clear(color) => wgpu_core::command::PassChannel {
+                    load_op: wgpu_core::command::LoadOp::Clear,
+                    store_op: color_attachment.store_op,
+                    clear_value: color,
+                    read_only: false,
+                },
+            },
+        };
+
+        color_attachments.push(attachment)
+    }
+
+    let mut depth_stencil_attachment = None;
+
+    if let Some(attachment) = args.depth_stencil_attachment {
+        let texture_view_resource = state
+            .resource_table
+            .get::<super::texture::WebGpuTextureView>(attachment.view)?;
+
+        depth_stencil_attachment = Some(wgpu_core::command::RenderPassDepthStencilAttachment {
+            view: texture_view_resource.0,
+            depth: match attachment.depth_load_op {
+                GpuLoadOp::Load => wgpu_core::command::PassChannel {
+                    load_op: wgpu_core::command::LoadOp::Load,
+                    store_op: attachment.depth_store_op,
+                    clear_value: 0.0,
+                    read_only: attachment.depth_read_only,
+                },
+                GpuLoadOp::Clear(value) => wgpu_core::command::PassChannel {
+                    load_op: wgpu_core::command::LoadOp::Clear,
+                    store_op: attachment.depth_store_op,
+                    clear_value: value,
+                    read_only: attachment.depth_read_only,
+                },
+            },
+            stencil: match attachment.stencil_load_op {
+                GpuLoadOp::Load => wgpu_core::command::PassChannel {
+                    load_op: wgpu_core::command::LoadOp::Load,
+                    store_op: attachment.stencil_store_op,
+                    clear_value: 0,
+                    read_only: attachment.stencil_read_only,
+                },
+                GpuLoadOp::Clear(value) => wgpu_core::command::PassChannel {
+                    load_op: wgpu_core::command::LoadOp::Clear,
+                    store_op: attachment.stencil_store_op,
+                    clear_value: value,
+                    read_only: attachment.stencil_read_only,
+                },
+            },
+        });
+    }
+
+    let descriptor = wgpu_core::command::RenderPassDescriptor {
+        label: args.label.map(Cow::from),
+        color_attachments: Cow::from(color_attachments),
+        depth_stencil_attachment: depth_stencil_attachment.as_ref(),
+    };
+
+    let render_pass = wgpu_core::command::RenderPass::new(command_encoder_resource.0, &descriptor);
+
+    let rid = state
+        .resource_table
+        .add(super::render_pass::WebGpuRenderPass(RefCell::new(
+            render_pass,
+        )));
+
+    Ok(WebGpuResult::rid(rid))
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct CommandEncoderBeginComputePassArgs {
+    command_encoder_rid: ResourceId,
+    label: Option<String>,
+}
+
+pub fn op_webgpu_command_encoder_begin_compute_pass(
+    state: &mut OpState,
+    args: CommandEncoderBeginComputePassArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let command_encoder_resource = state
+        .resource_table
+        .get::<WebGpuCommandEncoder>(args.command_encoder_rid)?;
+
+    let descriptor = wgpu_core::command::ComputePassDescriptor {
+        label: args.label.map(Cow::from),
+    };
+
+    let compute_pass =
+        wgpu_core::command::ComputePass::new(command_encoder_resource.0, &descriptor);
+
+    let rid = state
+        .resource_table
+        .add(super::compute_pass::WebGpuComputePass(RefCell::new(
+            compute_pass,
+        )));
+
+    Ok(WebGpuResult::rid(rid))
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct CommandEncoderCopyBufferToBufferArgs {
+    command_encoder_rid: ResourceId,
+    source: ResourceId,
+    source_offset: u64,
+    destination: ResourceId,
+    destination_offset: u64,
+    size: u64,
+}
+
+pub fn op_webgpu_command_encoder_copy_buffer_to_buffer(
+    state: &mut OpState,
+    args: CommandEncoderCopyBufferToBufferArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let instance = state.borrow::<super::Instance>();
+    let command_encoder_resource = state
+        .resource_table
+        .get::<WebGpuCommandEncoder>(args.command_encoder_rid)?;
+    let command_encoder = command_encoder_resource.0;
+    let source_buffer_resource = state
+        .resource_table
+        .get::<super::buffer::WebGpuBuffer>(args.source)?;
+    let source_buffer = source_buffer_resource.0;
+    let destination_buffer_resource = state
+        .resource_table
+        .get::<super::buffer::WebGpuBuffer>(args.destination)?;
+    let destination_buffer = destination_buffer_resource.0;
+
+    gfx_ok!(command_encoder => instance.command_encoder_copy_buffer_to_buffer(
+      command_encoder,
+      source_buffer,
+      args.source_offset,
+      destination_buffer,
+      args.destination_offset,
+      args.size
+    ))
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct GpuImageCopyBuffer {
+    buffer: ResourceId,
+    offset: u64,
+    bytes_per_row: Option<u32>,
+    rows_per_image: Option<u32>,
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct GpuImageCopyTexture {
+    pub texture: ResourceId,
+    pub mip_level: u32,
+    pub origin: wgpu_types::Origin3d,
+    pub aspect: wgpu_types::TextureAspect,
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct CommandEncoderCopyBufferToTextureArgs {
+    command_encoder_rid: ResourceId,
+    source: GpuImageCopyBuffer,
+    destination: GpuImageCopyTexture,
+    copy_size: wgpu_types::Extent3d,
+}
+
+pub fn op_webgpu_command_encoder_copy_buffer_to_texture(
+    state: &mut OpState,
+    args: CommandEncoderCopyBufferToTextureArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let instance = state.borrow::<super::Instance>();
+    let command_encoder_resource = state
+        .resource_table
+        .get::<WebGpuCommandEncoder>(args.command_encoder_rid)?;
+    let command_encoder = command_encoder_resource.0;
+    let source_buffer_resource = state
+        .resource_table
+        .get::<super::buffer::WebGpuBuffer>(args.source.buffer)?;
+    let destination_texture_resource = state
+        .resource_table
+        .get::<super::texture::WebGpuTexture>(args.destination.texture)?;
+
+    let source = wgpu_core::command::ImageCopyBuffer {
+        buffer: source_buffer_resource.0,
+        layout: wgpu_types::ImageDataLayout {
+            offset: args.source.offset,
+            bytes_per_row: NonZeroU32::new(args.source.bytes_per_row.unwrap_or(0)),
+            rows_per_image: NonZeroU32::new(args.source.rows_per_image.unwrap_or(0)),
+        },
+    };
+    let destination = wgpu_core::command::ImageCopyTexture {
+        texture: destination_texture_resource.0,
+        mip_level: args.destination.mip_level,
+        origin: args.destination.origin,
+        aspect: args.destination.aspect,
+    };
+    gfx_ok!(command_encoder => instance.command_encoder_copy_buffer_to_texture(
+      command_encoder,
+      &source,
+      &destination,
+      &args.copy_size
+    ))
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct CommandEncoderCopyTextureToBufferArgs {
+    command_encoder_rid: ResourceId,
+    source: GpuImageCopyTexture,
+    destination: GpuImageCopyBuffer,
+    copy_size: wgpu_types::Extent3d,
+}
+
+pub fn op_webgpu_command_encoder_copy_texture_to_buffer(
+    state: &mut OpState,
+    args: CommandEncoderCopyTextureToBufferArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let instance = state.borrow::<super::Instance>();
+    let command_encoder_resource = state
+        .resource_table
+        .get::<WebGpuCommandEncoder>(args.command_encoder_rid)?;
+    let command_encoder = command_encoder_resource.0;
+    let source_texture_resource = state
+        .resource_table
+        .get::<super::texture::WebGpuTexture>(args.source.texture)?;
+    let destination_buffer_resource = state
+        .resource_table
+        .get::<super::buffer::WebGpuBuffer>(args.destination.buffer)?;
+
+    let source = wgpu_core::command::ImageCopyTexture {
+        texture: source_texture_resource.0,
+        mip_level: args.source.mip_level,
+        origin: args.source.origin,
+        aspect: args.source.aspect,
+    };
+    let destination = wgpu_core::command::ImageCopyBuffer {
+        buffer: destination_buffer_resource.0,
+        layout: wgpu_types::ImageDataLayout {
+            offset: args.destination.offset,
+            bytes_per_row: NonZeroU32::new(args.destination.bytes_per_row.unwrap_or(0)),
+            rows_per_image: NonZeroU32::new(args.destination.rows_per_image.unwrap_or(0)),
+        },
+    };
+    gfx_ok!(command_encoder => instance.command_encoder_copy_texture_to_buffer(
+      command_encoder,
+      &source,
+      &destination,
+      &args.copy_size
+    ))
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct CommandEncoderCopyTextureToTextureArgs {
+    command_encoder_rid: ResourceId,
+    source: GpuImageCopyTexture,
+    destination: GpuImageCopyTexture,
+    copy_size: wgpu_types::Extent3d,
+}
+
+pub fn op_webgpu_command_encoder_copy_texture_to_texture(
+    state: &mut OpState,
+    args: CommandEncoderCopyTextureToTextureArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let instance = state.borrow::<super::Instance>();
+    let command_encoder_resource = state
+        .resource_table
+        .get::<WebGpuCommandEncoder>(args.command_encoder_rid)?;
+    let command_encoder = command_encoder_resource.0;
+    let source_texture_resource = state
+        .resource_table
+        .get::<super::texture::WebGpuTexture>(args.source.texture)?;
+    let destination_texture_resource = state
+        .resource_table
+        .get::<super::texture::WebGpuTexture>(args.destination.texture)?;
+
+    let source = wgpu_core::command::ImageCopyTexture {
+        texture: source_texture_resource.0,
+        mip_level: args.source.mip_level,
+        origin: args.source.origin,
+        aspect: args.source.aspect,
+    };
+    let destination = wgpu_core::command::ImageCopyTexture {
+        texture: destination_texture_resource.0,
+        mip_level: args.destination.mip_level,
+        origin: args.destination.origin,
+        aspect: args.destination.aspect,
+    };
+    gfx_ok!(command_encoder => instance.command_encoder_copy_texture_to_texture(
+      command_encoder,
+      &source,
+      &destination,
+      &args.copy_size
+    ))
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct CommandEncoderPushDebugGroupArgs {
+    command_encoder_rid: ResourceId,
+    group_label: String,
+}
+
+pub fn op_webgpu_command_encoder_push_debug_group(
+    state: &mut OpState,
+    args: CommandEncoderPushDebugGroupArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let instance = state.borrow::<super::Instance>();
+    let command_encoder_resource = state
+        .resource_table
+        .get::<WebGpuCommandEncoder>(args.command_encoder_rid)?;
+    let command_encoder = command_encoder_resource.0;
+
+    gfx_ok!(command_encoder => instance
+    .command_encoder_push_debug_group(command_encoder, &args.group_label))
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct CommandEncoderPopDebugGroupArgs {
+    command_encoder_rid: ResourceId,
+}
+
+pub fn op_webgpu_command_encoder_pop_debug_group(
+    state: &mut OpState,
+    args: CommandEncoderPopDebugGroupArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let instance = state.borrow::<super::Instance>();
+    let command_encoder_resource = state
+        .resource_table
+        .get::<WebGpuCommandEncoder>(args.command_encoder_rid)?;
+    let command_encoder = command_encoder_resource.0;
+
+    gfx_ok!(command_encoder => instance.command_encoder_pop_debug_group(command_encoder))
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct CommandEncoderInsertDebugMarkerArgs {
+    command_encoder_rid: ResourceId,
+    marker_label: String,
+}
+
+pub fn op_webgpu_command_encoder_insert_debug_marker(
+    state: &mut OpState,
+    args: CommandEncoderInsertDebugMarkerArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let instance = state.borrow::<super::Instance>();
+    let command_encoder_resource = state
+        .resource_table
+        .get::<WebGpuCommandEncoder>(args.command_encoder_rid)?;
+    let command_encoder = command_encoder_resource.0;
+
+    gfx_ok!(command_encoder => instance.command_encoder_insert_debug_marker(
+      command_encoder,
+      &args.marker_label
+    ))
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct CommandEncoderWriteTimestampArgs {
+    command_encoder_rid: ResourceId,
+    query_set: ResourceId,
+    query_index: u32,
+}
+
+pub fn op_webgpu_command_encoder_write_timestamp(
+    state: &mut OpState,
+    args: CommandEncoderWriteTimestampArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let instance = state.borrow::<super::Instance>();
+    let command_encoder_resource = state
+        .resource_table
+        .get::<WebGpuCommandEncoder>(args.command_encoder_rid)?;
+    let command_encoder = command_encoder_resource.0;
+    let query_set_resource = state
+        .resource_table
+        .get::<super::WebGpuQuerySet>(args.query_set)?;
+
+    gfx_ok!(command_encoder => instance.command_encoder_write_timestamp(
+      command_encoder,
+      query_set_resource.0,
+      args.query_index
+    ))
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct CommandEncoderResolveQuerySetArgs {
+    command_encoder_rid: ResourceId,
+    query_set: ResourceId,
+    first_query: u32,
+    query_count: u32,
+    destination: ResourceId,
+    destination_offset: u64,
+}
+
+pub fn op_webgpu_command_encoder_resolve_query_set(
+    state: &mut OpState,
+    args: CommandEncoderResolveQuerySetArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let instance = state.borrow::<super::Instance>();
+    let command_encoder_resource = state
+        .resource_table
+        .get::<WebGpuCommandEncoder>(args.command_encoder_rid)?;
+    let command_encoder = command_encoder_resource.0;
+    let query_set_resource = state
+        .resource_table
+        .get::<super::WebGpuQuerySet>(args.query_set)?;
+    let destination_resource = state
+        .resource_table
+        .get::<super::buffer::WebGpuBuffer>(args.destination)?;
+
+    gfx_ok!(command_encoder => instance.command_encoder_resolve_query_set(
+      command_encoder,
+      query_set_resource.0,
+      args.first_query,
+      args.query_count,
+      destination_resource.0,
+      args.destination_offset
+    ))
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct CommandEncoderFinishArgs {
+    command_encoder_rid: ResourceId,
+    label: Option<String>,
+}
+
+pub fn op_webgpu_command_encoder_finish(
+    state: &mut OpState,
+    args: CommandEncoderFinishArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let command_encoder_resource = state
+        .resource_table
+        .take::<WebGpuCommandEncoder>(args.command_encoder_rid)?;
+    let command_encoder = command_encoder_resource.0;
+    let instance = state.borrow::<super::Instance>();
+
+    let descriptor = wgpu_types::CommandBufferDescriptor {
+        label: args.label.map(Cow::from),
+    };
+
+    gfx_put!(command_encoder => instance.command_encoder_finish(
+    command_encoder,
+    &descriptor
+  ) => state, WebGpuCommandBuffer)
+}
--- a/deno_webgpu/compute_pass.rs
+++ b/deno_webgpu/compute_pass.rs
@ -0,0 +1,349 @@
+// Copyright 2018-2021 the Deno authors. All rights reserved. MIT license.
+
+use deno_core::error::AnyError;
+use deno_core::ResourceId;
+use deno_core::ZeroCopyBuf;
+use deno_core::{OpState, Resource};
+use serde::Deserialize;
+use std::borrow::Cow;
+use std::cell::RefCell;
+
+use super::error::WebGpuResult;
+
+pub(crate) struct WebGpuComputePass(pub(crate) RefCell<wgpu_core::command::ComputePass>);
+impl Resource for WebGpuComputePass {
+    fn name(&self) -> Cow<str> {
+        "webGPUComputePass".into()
+    }
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct ComputePassSetPipelineArgs {
+    compute_pass_rid: ResourceId,
+    pipeline: ResourceId,
+}
+
+pub fn op_webgpu_compute_pass_set_pipeline(
+    state: &mut OpState,
+    args: ComputePassSetPipelineArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let compute_pipeline_resource = state
+        .resource_table
+        .get::<super::pipeline::WebGpuComputePipeline>(args.pipeline)?;
+    let compute_pass_resource = state
+        .resource_table
+        .get::<WebGpuComputePass>(args.compute_pass_rid)?;
+
+    wgpu_core::command::compute_ffi::wgpu_compute_pass_set_pipeline(
+        &mut compute_pass_resource.0.borrow_mut(),
+        compute_pipeline_resource.0,
+    );
+
+    Ok(WebGpuResult::empty())
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct ComputePassDispatchArgs {
+    compute_pass_rid: ResourceId,
+    x: u32,
+    y: u32,
+    z: u32,
+}
+
+pub fn op_webgpu_compute_pass_dispatch(
+    state: &mut OpState,
+    args: ComputePassDispatchArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let compute_pass_resource = state
+        .resource_table
+        .get::<WebGpuComputePass>(args.compute_pass_rid)?;
+
+    wgpu_core::command::compute_ffi::wgpu_compute_pass_dispatch(
+        &mut compute_pass_resource.0.borrow_mut(),
+        args.x,
+        args.y,
+        args.z,
+    );
+
+    Ok(WebGpuResult::empty())
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct ComputePassDispatchIndirectArgs {
+    compute_pass_rid: ResourceId,
+    indirect_buffer: ResourceId,
+    indirect_offset: u64,
+}
+
+pub fn op_webgpu_compute_pass_dispatch_indirect(
+    state: &mut OpState,
+    args: ComputePassDispatchIndirectArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let buffer_resource = state
+        .resource_table
+        .get::<super::buffer::WebGpuBuffer>(args.indirect_buffer)?;
+    let compute_pass_resource = state
+        .resource_table
+        .get::<WebGpuComputePass>(args.compute_pass_rid)?;
+
+    wgpu_core::command::compute_ffi::wgpu_compute_pass_dispatch_indirect(
+        &mut compute_pass_resource.0.borrow_mut(),
+        buffer_resource.0,
+        args.indirect_offset,
+    );
+
+    Ok(WebGpuResult::empty())
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct ComputePassBeginPipelineStatisticsQueryArgs {
+    compute_pass_rid: ResourceId,
+    query_set: ResourceId,
+    query_index: u32,
+}
+
+pub fn op_webgpu_compute_pass_begin_pipeline_statistics_query(
+    state: &mut OpState,
+    args: ComputePassBeginPipelineStatisticsQueryArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let compute_pass_resource = state
+        .resource_table
+        .get::<WebGpuComputePass>(args.compute_pass_rid)?;
+    let query_set_resource = state
+        .resource_table
+        .get::<super::WebGpuQuerySet>(args.query_set)?;
+
+    wgpu_core::command::compute_ffi::wgpu_compute_pass_begin_pipeline_statistics_query(
+        &mut compute_pass_resource.0.borrow_mut(),
+        query_set_resource.0,
+        args.query_index,
+    );
+
+    Ok(WebGpuResult::empty())
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct ComputePassEndPipelineStatisticsQueryArgs {
+    compute_pass_rid: ResourceId,
+}
+
+pub fn op_webgpu_compute_pass_end_pipeline_statistics_query(
+    state: &mut OpState,
+    args: ComputePassEndPipelineStatisticsQueryArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let compute_pass_resource = state
+        .resource_table
+        .get::<WebGpuComputePass>(args.compute_pass_rid)?;
+
+    wgpu_core::command::compute_ffi::wgpu_compute_pass_end_pipeline_statistics_query(
+        &mut compute_pass_resource.0.borrow_mut(),
+    );
+
+    Ok(WebGpuResult::empty())
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct ComputePassWriteTimestampArgs {
+    compute_pass_rid: ResourceId,
+    query_set: ResourceId,
+    query_index: u32,
+}
+
+pub fn op_webgpu_compute_pass_write_timestamp(
+    state: &mut OpState,
+    args: ComputePassWriteTimestampArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let compute_pass_resource = state
+        .resource_table
+        .get::<WebGpuComputePass>(args.compute_pass_rid)?;
+    let query_set_resource = state
+        .resource_table
+        .get::<super::WebGpuQuerySet>(args.query_set)?;
+
+    wgpu_core::command::compute_ffi::wgpu_compute_pass_write_timestamp(
+        &mut compute_pass_resource.0.borrow_mut(),
+        query_set_resource.0,
+        args.query_index,
+    );
+
+    Ok(WebGpuResult::empty())
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct ComputePassEndPassArgs {
+    command_encoder_rid: ResourceId,
+    compute_pass_rid: ResourceId,
+}
+
+pub fn op_webgpu_compute_pass_end_pass(
+    state: &mut OpState,
+    args: ComputePassEndPassArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let command_encoder_resource =
+        state
+            .resource_table
+            .get::<super::command_encoder::WebGpuCommandEncoder>(args.command_encoder_rid)?;
+    let command_encoder = command_encoder_resource.0;
+    let compute_pass_resource = state
+        .resource_table
+        .take::<WebGpuComputePass>(args.compute_pass_rid)?;
+    let compute_pass = &compute_pass_resource.0.borrow();
+    let instance = state.borrow::<super::Instance>();
+
+    gfx_ok!(command_encoder => instance.command_encoder_run_compute_pass(
+      command_encoder,
+      compute_pass
+    ))
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct ComputePassSetBindGroupArgs {
+    compute_pass_rid: ResourceId,
+    index: u32,
+    bind_group: ResourceId,
+    dynamic_offsets_data: ZeroCopyBuf,
+    dynamic_offsets_data_start: usize,
+    dynamic_offsets_data_length: usize,
+}
+
+pub fn op_webgpu_compute_pass_set_bind_group(
+    state: &mut OpState,
+    args: ComputePassSetBindGroupArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let bind_group_resource = state
+        .resource_table
+        .get::<super::binding::WebGpuBindGroup>(args.bind_group)?;
+    let compute_pass_resource = state
+        .resource_table
+        .get::<WebGpuComputePass>(args.compute_pass_rid)?;
+
+    // Align the data
+    assert!(args.dynamic_offsets_data_start % std::mem::size_of::<u32>() == 0);
+    // SAFETY: A u8 to u32 cast is safe because we asserted that the length is a
+    // multiple of 4.
+    let (prefix, dynamic_offsets_data, suffix) =
+        unsafe { args.dynamic_offsets_data.align_to::<u32>() };
+    assert!(prefix.is_empty());
+    assert!(suffix.is_empty());
+
+    let start = args.dynamic_offsets_data_start;
+    let len = args.dynamic_offsets_data_length;
+
+    // Assert that length and start are both in bounds
+    assert!(start <= dynamic_offsets_data.len());
+    assert!(len <= dynamic_offsets_data.len() - start);
+
+    let dynamic_offsets_data: &[u32] = &dynamic_offsets_data[start..start + len];
+
+    // SAFETY: the raw pointer and length are of the same slice, and that slice
+    // lives longer than the below function invocation.
+    unsafe {
+        wgpu_core::command::compute_ffi::wgpu_compute_pass_set_bind_group(
+            &mut compute_pass_resource.0.borrow_mut(),
+            args.index,
+            bind_group_resource.0,
+            dynamic_offsets_data.as_ptr(),
+            dynamic_offsets_data.len(),
+        );
+    }
+
+    Ok(WebGpuResult::empty())
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct ComputePassPushDebugGroupArgs {
+    compute_pass_rid: ResourceId,
+    group_label: String,
+}
+
+pub fn op_webgpu_compute_pass_push_debug_group(
+    state: &mut OpState,
+    args: ComputePassPushDebugGroupArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let compute_pass_resource = state
+        .resource_table
+        .get::<WebGpuComputePass>(args.compute_pass_rid)?;
+
+    let label = std::ffi::CString::new(args.group_label).unwrap();
+    // SAFETY: the string the raw pointer points to lives longer than the below
+    // function invocation.
+    unsafe {
+        wgpu_core::command::compute_ffi::wgpu_compute_pass_push_debug_group(
+            &mut compute_pass_resource.0.borrow_mut(),
+            label.as_ptr(),
+            0, // wgpu#975
+        );
+    }
+
+    Ok(WebGpuResult::empty())
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct ComputePassPopDebugGroupArgs {
+    compute_pass_rid: ResourceId,
+}
+
+pub fn op_webgpu_compute_pass_pop_debug_group(
+    state: &mut OpState,
+    args: ComputePassPopDebugGroupArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let compute_pass_resource = state
+        .resource_table
+        .get::<WebGpuComputePass>(args.compute_pass_rid)?;
+
+    wgpu_core::command::compute_ffi::wgpu_compute_pass_pop_debug_group(
+        &mut compute_pass_resource.0.borrow_mut(),
+    );
+
+    Ok(WebGpuResult::empty())
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct ComputePassInsertDebugMarkerArgs {
+    compute_pass_rid: ResourceId,
+    marker_label: String,
+}
+
+pub fn op_webgpu_compute_pass_insert_debug_marker(
+    state: &mut OpState,
+    args: ComputePassInsertDebugMarkerArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let compute_pass_resource = state
+        .resource_table
+        .get::<WebGpuComputePass>(args.compute_pass_rid)?;
+
+    let label = std::ffi::CString::new(args.marker_label).unwrap();
+    // SAFETY: the string the raw pointer points to lives longer than the below
+    // function invocation.
+    unsafe {
+        wgpu_core::command::compute_ffi::wgpu_compute_pass_insert_debug_marker(
+            &mut compute_pass_resource.0.borrow_mut(),
+            label.as_ptr(),
+            0, // wgpu#975
+        );
+    }
+
+    Ok(WebGpuResult::empty())
+}
--- a/deno_webgpu/error.rs
+++ b/deno_webgpu/error.rs
@ -0,0 +1,282 @@
+// Copyright 2018-2021 the Deno authors. All rights reserved. MIT license.
+use deno_core::error::AnyError;
+use deno_core::ResourceId;
+use serde::Serialize;
+use std::convert::From;
+use std::fmt;
+use wgpu_core::binding_model::CreateBindGroupError;
+use wgpu_core::binding_model::CreateBindGroupLayoutError;
+use wgpu_core::binding_model::CreatePipelineLayoutError;
+use wgpu_core::binding_model::GetBindGroupLayoutError;
+use wgpu_core::command::CommandEncoderError;
+use wgpu_core::command::ComputePassError;
+use wgpu_core::command::CopyError;
+use wgpu_core::command::CreateRenderBundleError;
+use wgpu_core::command::QueryError;
+use wgpu_core::command::RenderBundleError;
+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::pipeline::CreateComputePipelineError;
+use wgpu_core::pipeline::CreateRenderPipelineError;
+use wgpu_core::pipeline::CreateShaderModuleError;
+use wgpu_core::resource::BufferAccessError;
+use wgpu_core::resource::CreateBufferError;
+use wgpu_core::resource::CreateQuerySetError;
+use wgpu_core::resource::CreateSamplerError;
+use wgpu_core::resource::CreateTextureError;
+use wgpu_core::resource::CreateTextureViewError;
+
+#[derive(Serialize)]
+pub struct WebGpuResult {
+    pub rid: Option<ResourceId>,
+    pub err: Option<WebGpuError>,
+}
+
+impl WebGpuResult {
+    pub fn rid(rid: ResourceId) -> Self {
+        Self {
+            rid: Some(rid),
+            err: None,
+        }
+    }
+
+    pub fn rid_err<T: Into<WebGpuError>>(rid: ResourceId, err: Option<T>) -> Self {
+        Self {
+            rid: Some(rid),
+            err: err.map(|e| e.into()),
+        }
+    }
+
+    pub fn maybe_err<T: Into<WebGpuError>>(err: Option<T>) -> Self {
+        Self {
+            rid: None,
+            err: err.map(|e| e.into()),
+        }
+    }
+
+    pub fn empty() -> Self {
+        Self {
+            rid: None,
+            err: None,
+        }
+    }
+}
+
+#[derive(Serialize)]
+#[serde(tag = "type", content = "value")]
+#[serde(rename_all = "kebab-case")]
+pub enum WebGpuError {
+    Lost,
+    OutOfMemory,
+    Validation(String),
+}
+
+impl From<CreateBufferError> for WebGpuError {
+    fn from(err: CreateBufferError) -> Self {
+        match err {
+            CreateBufferError::Device(err) => err.into(),
+            CreateBufferError::AccessError(err) => err.into(),
+            err => WebGpuError::Validation(err.to_string()),
+        }
+    }
+}
+
+impl From<DeviceError> for WebGpuError {
+    fn from(err: DeviceError) -> Self {
+        match err {
+            DeviceError::Lost => WebGpuError::Lost,
+            DeviceError::OutOfMemory => WebGpuError::OutOfMemory,
+            DeviceError::Invalid => WebGpuError::Validation(err.to_string()),
+        }
+    }
+}
+
+impl From<BufferAccessError> for WebGpuError {
+    fn from(err: BufferAccessError) -> Self {
+        match err {
+            BufferAccessError::Device(err) => err.into(),
+            err => WebGpuError::Validation(err.to_string()),
+        }
+    }
+}
+
+impl From<CreateBindGroupLayoutError> for WebGpuError {
+    fn from(err: CreateBindGroupLayoutError) -> Self {
+        match err {
+            CreateBindGroupLayoutError::Device(err) => err.into(),
+            err => WebGpuError::Validation(err.to_string()),
+        }
+    }
+}
+
+impl From<CreatePipelineLayoutError> for WebGpuError {
+    fn from(err: CreatePipelineLayoutError) -> Self {
+        match err {
+            CreatePipelineLayoutError::Device(err) => err.into(),
+            err => WebGpuError::Validation(err.to_string()),
+        }
+    }
+}
+
+impl From<CreateBindGroupError> for WebGpuError {
+    fn from(err: CreateBindGroupError) -> Self {
+        match err {
+            CreateBindGroupError::Device(err) => err.into(),
+            err => WebGpuError::Validation(err.to_string()),
+        }
+    }
+}
+
+impl From<RenderBundleError> for WebGpuError {
+    fn from(err: RenderBundleError) -> Self {
+        WebGpuError::Validation(err.to_string())
+    }
+}
+
+impl From<CreateRenderBundleError> for WebGpuError {
+    fn from(err: CreateRenderBundleError) -> Self {
+        WebGpuError::Validation(err.to_string())
+    }
+}
+
+impl From<CopyError> for WebGpuError {
+    fn from(err: CopyError) -> Self {
+        WebGpuError::Validation(err.to_string())
+    }
+}
+
+impl From<CommandEncoderError> for WebGpuError {
+    fn from(err: CommandEncoderError) -> Self {
+        WebGpuError::Validation(err.to_string())
+    }
+}
+
+impl From<QueryError> for WebGpuError {
+    fn from(err: QueryError) -> Self {
+        WebGpuError::Validation(err.to_string())
+    }
+}
+
+impl From<ComputePassError> for WebGpuError {
+    fn from(err: ComputePassError) -> Self {
+        WebGpuError::Validation(err.to_string())
+    }
+}
+
+impl From<CreateComputePipelineError> for WebGpuError {
+    fn from(err: CreateComputePipelineError) -> Self {
+        match err {
+            CreateComputePipelineError::Device(err) => err.into(),
+            err => WebGpuError::Validation(err.to_string()),
+        }
+    }
+}
+
+impl From<GetBindGroupLayoutError> for WebGpuError {
+    fn from(err: GetBindGroupLayoutError) -> Self {
+        WebGpuError::Validation(err.to_string())
+    }
+}
+
+impl From<CreateRenderPipelineError> for WebGpuError {
+    fn from(err: CreateRenderPipelineError) -> Self {
+        match err {
+            CreateRenderPipelineError::Device(err) => err.into(),
+            err => WebGpuError::Validation(err.to_string()),
+        }
+    }
+}
+
+impl From<RenderPassError> for WebGpuError {
+    fn from(err: RenderPassError) -> Self {
+        WebGpuError::Validation(err.to_string())
+    }
+}
+
+impl From<CreateSamplerError> for WebGpuError {
+    fn from(err: CreateSamplerError) -> Self {
+        match err {
+            CreateSamplerError::Device(err) => err.into(),
+            err => WebGpuError::Validation(err.to_string()),
+        }
+    }
+}
+
+impl From<CreateShaderModuleError> for WebGpuError {
+    fn from(err: CreateShaderModuleError) -> Self {
+        match err {
+            CreateShaderModuleError::Device(err) => err.into(),
+            err => WebGpuError::Validation(err.to_string()),
+        }
+    }
+}
+
+impl From<CreateTextureError> for WebGpuError {
+    fn from(err: CreateTextureError) -> Self {
+        match err {
+            CreateTextureError::Device(err) => err.into(),
+            err => WebGpuError::Validation(err.to_string()),
+        }
+    }
+}
+
+impl From<CreateTextureViewError> for WebGpuError {
+    fn from(err: CreateTextureViewError) -> Self {
+        WebGpuError::Validation(err.to_string())
+    }
+}
+
+impl From<CreateQuerySetError> for WebGpuError {
+    fn from(err: CreateQuerySetError) -> Self {
+        match err {
+            CreateQuerySetError::Device(err) => err.into(),
+            err => WebGpuError::Validation(err.to_string()),
+        }
+    }
+}
+
+impl From<QueueSubmitError> for WebGpuError {
+    fn from(err: QueueSubmitError) -> Self {
+        match err {
+            QueueSubmitError::Queue(err) => err.into(),
+            err => WebGpuError::Validation(err.to_string()),
+        }
+    }
+}
+
+impl From<QueueWriteError> for WebGpuError {
+    fn from(err: QueueWriteError) -> Self {
+        match err {
+            QueueWriteError::Queue(err) => err.into(),
+            err => WebGpuError::Validation(err.to_string()),
+        }
+    }
+}
+
+#[derive(Debug)]
+pub struct DomExceptionOperationError {
+    pub msg: String,
+}
+
+impl DomExceptionOperationError {
+    pub fn new(msg: &str) -> Self {
+        DomExceptionOperationError {
+            msg: msg.to_string(),
+        }
+    }
+}
+
+impl fmt::Display for DomExceptionOperationError {
+    fn fmt(&self, f: &mut fmt::Formatter) -> fmt::Result {
+        f.pad(&self.msg)
+    }
+}
+
+impl std::error::Error for DomExceptionOperationError {}
+
+pub fn get_error_class_name(e: &AnyError) -> Option<&'static str> {
+    e.downcast_ref::<DomExceptionOperationError>()
+        .map(|_| "DOMExceptionOperationError")
+}
--- a/deno_webgpu/lib.rs
+++ b/deno_webgpu/lib.rs
@ -0,0 +1,850 @@
+// Copyright 2018-2021 the Deno authors. All rights reserved. MIT license.
+
+use deno_core::error::AnyError;
+use deno_core::include_js_files;
+use deno_core::op_async;
+use deno_core::op_sync;
+use deno_core::Extension;
+use deno_core::OpFn;
+use deno_core::OpState;
+use deno_core::Resource;
+use deno_core::ResourceId;
+use serde::Deserialize;
+use serde::Serialize;
+use std::borrow::Cow;
+use std::cell::RefCell;
+use std::collections::HashSet;
+use std::rc::Rc;
+pub use wgpu_core;
+pub use wgpu_types;
+use wgpu_types::PowerPreference;
+
+use error::DomExceptionOperationError;
+use error::WebGpuResult;
+
+#[macro_use]
+mod macros {
+    macro_rules! gfx_select {
+    ($id:expr => $global:ident.$method:ident( $($param:expr),* )) => {
+      match $id.backend() {
+        #[cfg(not(target_os = "macos"))]
+        wgpu_types::Backend::Vulkan => $global.$method::<wgpu_core::api::Vulkan>( $($param),* ),
+        #[cfg(target_os = "macos")]
+        wgpu_types::Backend::Metal => $global.$method::<wgpu_core::api::Metal>( $($param),* ),
+        #[cfg(windows)]
+        wgpu_types::Backend::Dx12 => $global.$method::<wgpu_core::api::Dx12>( $($param),* ),
+        #[cfg(all(unix, not(target_os = "macos")))]
+        wgpu_types::Backend::Gl => $global.$method::<wgpu_core::api::Gles>( $($param),+ ),
+        other => panic!("Unexpected backend {:?}", other),
+      }
+    };
+  }
+
+    macro_rules! gfx_put {
+    ($id:expr => $global:ident.$method:ident( $($param:expr),* ) => $state:expr, $rc:expr) => {{
+      let (val, maybe_err) = gfx_select!($id => $global.$method($($param),*));
+      let rid = $state.resource_table.add($rc(val));
+      Ok(WebGpuResult::rid_err(rid, maybe_err))
+    }};
+  }
+
+    macro_rules! gfx_ok {
+    ($id:expr => $global:ident.$method:ident( $($param:expr),* )) => {{
+      let maybe_err = gfx_select!($id => $global.$method($($param),*)).err();
+      Ok(WebGpuResult::maybe_err(maybe_err))
+    }};
+  }
+}
+
+pub mod binding;
+pub mod buffer;
+pub mod bundle;
+pub mod command_encoder;
+pub mod compute_pass;
+pub mod error;
+pub mod pipeline;
+pub mod queue;
+pub mod render_pass;
+pub mod sampler;
+pub mod shader;
+pub mod texture;
+
+pub struct Unstable(pub bool);
+
+fn check_unstable(state: &OpState, api_name: &str) {
+    let unstable = state.borrow::<Unstable>();
+    if !unstable.0 {
+        eprintln!(
+            "Unstable API '{}'. The --unstable flag must be provided.",
+            api_name
+        );
+        std::process::exit(70);
+    }
+}
+
+type Instance = wgpu_core::hub::Global<wgpu_core::hub::IdentityManagerFactory>;
+
+struct WebGpuAdapter(wgpu_core::id::AdapterId);
+impl Resource for WebGpuAdapter {
+    fn name(&self) -> Cow<str> {
+        "webGPUAdapter".into()
+    }
+}
+
+struct WebGpuDevice(wgpu_core::id::DeviceId);
+impl Resource for WebGpuDevice {
+    fn name(&self) -> Cow<str> {
+        "webGPUDevice".into()
+    }
+}
+
+struct WebGpuQuerySet(wgpu_core::id::QuerySetId);
+impl Resource for WebGpuQuerySet {
+    fn name(&self) -> Cow<str> {
+        "webGPUQuerySet".into()
+    }
+}
+
+pub fn init(unstable: bool) -> Extension {
+    Extension::builder()
+        .js(include_js_files!(
+          prefix "deno:deno_webgpu",
+          "01_webgpu.js",
+          "02_idl_types.js",
+        ))
+        .ops(declare_webgpu_ops())
+        .state(move |state| {
+            // TODO: check & possibly streamline this
+            // Unstable might be able to be OpMiddleware
+            // let unstable_checker = state.borrow::<super::UnstableChecker>();
+            // let unstable = unstable_checker.unstable;
+            state.put(Unstable(unstable));
+            Ok(())
+        })
+        .build()
+}
+
+fn deserialize_features(features: &wgpu_types::Features) -> Vec<&'static str> {
+    let mut return_features: Vec<&'static str> = vec![];
+
+    if features.contains(wgpu_types::Features::DEPTH_CLAMPING) {
+        return_features.push("depth-clamping");
+    }
+    if features.contains(wgpu_types::Features::PIPELINE_STATISTICS_QUERY) {
+        return_features.push("pipeline-statistics-query");
+    }
+    if features.contains(wgpu_types::Features::TEXTURE_COMPRESSION_BC) {
+        return_features.push("texture-compression-bc");
+    }
+    if features.contains(wgpu_types::Features::TIMESTAMP_QUERY) {
+        return_features.push("timestamp-query");
+    }
+
+    // extended from spec
+    if features.contains(wgpu_types::Features::MAPPABLE_PRIMARY_BUFFERS) {
+        return_features.push("mappable-primary-buffers");
+    }
+    if features.contains(wgpu_types::Features::TEXTURE_BINDING_ARRAY) {
+        return_features.push("texture-binding-array");
+    }
+    if features.contains(wgpu_types::Features::BUFFER_BINDING_ARRAY) {
+        return_features.push("buffer-binding-array");
+    }
+    if features.contains(wgpu_types::Features::STORAGE_RESOURCE_BINDING_ARRAY) {
+        return_features.push("storage-resource-binding-array");
+    }
+    if features.contains(
+        wgpu_types::Features::SAMPLED_TEXTURE_AND_STORAGE_BUFFER_ARRAY_NON_UNIFORM_INDEXING,
+    ) {
+        return_features.push("sampled-texture-and-storage-buffer-array-non-uniform-indexing");
+    }
+    if features.contains(
+        wgpu_types::Features::UNIFORM_BUFFER_AND_STORAGE_TEXTURE_ARRAY_NON_UNIFORM_INDEXING,
+    ) {
+        return_features.push("uniform-buffer-and-storage-buffer-texture-non-uniform-indexing");
+    }
+    if features.contains(wgpu_types::Features::UNSIZED_BINDING_ARRAY) {
+        return_features.push("unsized-binding-array");
+    }
+    if features.contains(wgpu_types::Features::MULTI_DRAW_INDIRECT) {
+        return_features.push("multi-draw-indirect");
+    }
+    if features.contains(wgpu_types::Features::MULTI_DRAW_INDIRECT_COUNT) {
+        return_features.push("multi-draw-indirect-count");
+    }
+    if features.contains(wgpu_types::Features::PUSH_CONSTANTS) {
+        return_features.push("push-constants");
+    }
+    if features.contains(wgpu_types::Features::ADDRESS_MODE_CLAMP_TO_BORDER) {
+        return_features.push("address-mode-clamp-to-border");
+    }
+    if features.contains(wgpu_types::Features::TEXTURE_COMPRESSION_ETC2) {
+        return_features.push("texture-compression-etc2");
+    }
+    if features.contains(wgpu_types::Features::TEXTURE_COMPRESSION_ASTC_LDR) {
+        return_features.push("texture-compression-astc-ldr");
+    }
+    if features.contains(wgpu_types::Features::TEXTURE_ADAPTER_SPECIFIC_FORMAT_FEATURES) {
+        return_features.push("texture-adapter-specific-format-features");
+    }
+    if features.contains(wgpu_types::Features::SHADER_FLOAT64) {
+        return_features.push("shader-float64");
+    }
+    if features.contains(wgpu_types::Features::VERTEX_ATTRIBUTE_64BIT) {
+        return_features.push("vertex-attribute-64bit");
+    }
+    if features.contains(wgpu_types::Features::CONSERVATIVE_RASTERIZATION) {
+        return_features.push("conservative-rasterization");
+    }
+    if features.contains(wgpu_types::Features::VERTEX_WRITABLE_STORAGE) {
+        return_features.push("vertex-writable-storage");
+    }
+    if features.contains(wgpu_types::Features::CLEAR_COMMANDS) {
+        return_features.push("clear-commands");
+    }
+    if features.contains(wgpu_types::Features::SPIRV_SHADER_PASSTHROUGH) {
+        return_features.push("spirv-shader-passthrough");
+    }
+    if features.contains(wgpu_types::Features::SHADER_PRIMITIVE_INDEX) {
+        return_features.push("shader-primitive-index");
+    }
+
+    return_features
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct RequestAdapterArgs {
+    power_preference: Option<wgpu_types::PowerPreference>,
+}
+
+#[derive(Serialize)]
+#[serde(untagged)]
+pub enum GpuAdapterDeviceOrErr {
+    Error { err: String },
+    Features(GpuAdapterDevice),
+}
+
+#[derive(Serialize)]
+#[serde(rename_all = "camelCase")]
+pub struct GpuAdapterDevice {
+    rid: ResourceId,
+    name: Option<String>,
+    limits: wgpu_types::Limits,
+    features: Vec<&'static str>,
+    is_software: bool,
+}
+
+pub async fn op_webgpu_request_adapter(
+    state: Rc<RefCell<OpState>>,
+    args: RequestAdapterArgs,
+    _: (),
+) -> Result<GpuAdapterDeviceOrErr, AnyError> {
+    let mut state = state.borrow_mut();
+    check_unstable(&state, "navigator.gpu.requestAdapter");
+    let instance = if let Some(instance) = state.try_borrow::<Instance>() {
+        instance
+    } else {
+        state.put(wgpu_core::hub::Global::new(
+            "webgpu",
+            wgpu_core::hub::IdentityManagerFactory,
+            wgpu_types::Backends::PRIMARY,
+        ));
+        state.borrow::<Instance>()
+    };
+
+    let descriptor = wgpu_core::instance::RequestAdapterOptions {
+        power_preference: match args.power_preference {
+            Some(power_preference) => power_preference.into(),
+            None => PowerPreference::default(),
+        },
+        // TODO(lucacasonato): respect forceFallbackAdapter
+        compatible_surface: None, // windowless
+    };
+    let res = instance.request_adapter(
+        &descriptor,
+        wgpu_core::instance::AdapterInputs::Mask(wgpu_types::Backends::PRIMARY, |_| {
+            std::marker::PhantomData
+        }),
+    );
+
+    let adapter = match res {
+        Ok(adapter) => adapter,
+        Err(err) => {
+            return Ok(GpuAdapterDeviceOrErr::Error {
+                err: err.to_string(),
+            })
+        }
+    };
+    let name = gfx_select!(adapter => instance.adapter_get_info(adapter))?.name;
+    let adapter_features = gfx_select!(adapter => instance.adapter_features(adapter))?;
+    let features = deserialize_features(&adapter_features);
+    let adapter_limits = gfx_select!(adapter => instance.adapter_limits(adapter))?;
+
+    let rid = state.resource_table.add(WebGpuAdapter(adapter));
+
+    Ok(GpuAdapterDeviceOrErr::Features(GpuAdapterDevice {
+        rid,
+        name: Some(name),
+        features,
+        limits: adapter_limits,
+        is_software: false,
+    }))
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct RequestDeviceArgs {
+    adapter_rid: ResourceId,
+    label: Option<String>,
+    required_features: Option<GpuRequiredFeatures>,
+    required_limits: Option<wgpu_types::Limits>,
+}
+
+#[derive(Deserialize)]
+pub struct GpuRequiredFeatures(HashSet<String>);
+
+impl From<GpuRequiredFeatures> for wgpu_types::Features {
+    fn from(required_features: GpuRequiredFeatures) -> wgpu_types::Features {
+        let mut features: wgpu_types::Features = wgpu_types::Features::empty();
+
+        if required_features.0.contains("depth-clamping") {
+            features.set(wgpu_types::Features::DEPTH_CLAMPING, true);
+        }
+        if required_features.0.contains("pipeline-statistics-query") {
+            features.set(wgpu_types::Features::PIPELINE_STATISTICS_QUERY, true);
+        }
+        if required_features.0.contains("texture-compression-bc") {
+            features.set(wgpu_types::Features::TEXTURE_COMPRESSION_BC, true);
+        }
+        if required_features.0.contains("timestamp-query") {
+            features.set(wgpu_types::Features::TIMESTAMP_QUERY, true);
+        }
+
+        // extended from spec
+        if required_features.0.contains("mappable-primary-buffers") {
+            features.set(wgpu_types::Features::MAPPABLE_PRIMARY_BUFFERS, true);
+        }
+        if required_features.0.contains("texture-binding-array") {
+            features.set(wgpu_types::Features::TEXTURE_BINDING_ARRAY, true);
+        }
+        if required_features.0.contains("buffer-binding-array") {
+            features.set(wgpu_types::Features::BUFFER_BINDING_ARRAY, true);
+        }
+        if required_features
+            .0
+            .contains("storage-resource-binding-array")
+        {
+            features.set(wgpu_types::Features::STORAGE_RESOURCE_BINDING_ARRAY, true);
+        }
+        if required_features
+            .0
+            .contains("sampled-texture-and-storage-buffer-array-non-uniform-indexing")
+        {
+            features.set(
+                wgpu_types::Features::SAMPLED_TEXTURE_AND_STORAGE_BUFFER_ARRAY_NON_UNIFORM_INDEXING,
+                true,
+            );
+        }
+        if required_features
+            .0
+            .contains("uniform-buffer-and-storage-buffer-texture-non-uniform-indexing")
+        {
+            features.set(
+                wgpu_types::Features::UNIFORM_BUFFER_AND_STORAGE_TEXTURE_ARRAY_NON_UNIFORM_INDEXING,
+                true,
+            );
+        }
+        if required_features.0.contains("unsized-binding-array") {
+            features.set(wgpu_types::Features::UNSIZED_BINDING_ARRAY, true);
+        }
+        if required_features.0.contains("multi-draw-indirect") {
+            features.set(wgpu_types::Features::MULTI_DRAW_INDIRECT, true);
+        }
+        if required_features.0.contains("multi-draw-indirect-count") {
+            features.set(wgpu_types::Features::MULTI_DRAW_INDIRECT_COUNT, true);
+        }
+        if required_features.0.contains("push-constants") {
+            features.set(wgpu_types::Features::PUSH_CONSTANTS, true);
+        }
+        if required_features.0.contains("address-mode-clamp-to-border") {
+            features.set(wgpu_types::Features::ADDRESS_MODE_CLAMP_TO_BORDER, true);
+        }
+        if required_features.0.contains("texture-compression-etc2") {
+            features.set(wgpu_types::Features::TEXTURE_COMPRESSION_ETC2, true);
+        }
+        if required_features.0.contains("texture-compression-astc-ldr") {
+            features.set(wgpu_types::Features::TEXTURE_COMPRESSION_ASTC_LDR, true);
+        }
+        if required_features
+            .0
+            .contains("texture-adapter-specific-format-features")
+        {
+            features.set(
+                wgpu_types::Features::TEXTURE_ADAPTER_SPECIFIC_FORMAT_FEATURES,
+                true,
+            );
+        }
+        if required_features.0.contains("shader-float64") {
+            features.set(wgpu_types::Features::SHADER_FLOAT64, true);
+        }
+        if required_features.0.contains("vertex-attribute-64bit") {
+            features.set(wgpu_types::Features::VERTEX_ATTRIBUTE_64BIT, true);
+        }
+        if required_features.0.contains("conservative-rasterization") {
+            features.set(wgpu_types::Features::CONSERVATIVE_RASTERIZATION, true);
+        }
+        if required_features.0.contains("vertex-writable-storage") {
+            features.set(wgpu_types::Features::VERTEX_WRITABLE_STORAGE, true);
+        }
+        if required_features.0.contains("clear-commands") {
+            features.set(wgpu_types::Features::CLEAR_COMMANDS, true);
+        }
+        if required_features.0.contains("spirv-shader-passthrough") {
+            features.set(wgpu_types::Features::SPIRV_SHADER_PASSTHROUGH, true);
+        }
+        if required_features.0.contains("shader-primitive-index") {
+            features.set(wgpu_types::Features::SHADER_PRIMITIVE_INDEX, true);
+        }
+
+        features
+    }
+}
+
+pub async fn op_webgpu_request_device(
+    state: Rc<RefCell<OpState>>,
+    args: RequestDeviceArgs,
+    _: (),
+) -> Result<GpuAdapterDevice, AnyError> {
+    let mut state = state.borrow_mut();
+    let adapter_resource = state
+        .resource_table
+        .get::<WebGpuAdapter>(args.adapter_rid)?;
+    let adapter = adapter_resource.0;
+    let instance = state.borrow::<Instance>();
+
+    let descriptor = wgpu_types::DeviceDescriptor {
+        label: args.label.map(Cow::from),
+        features: args.required_features.map(Into::into).unwrap_or_default(),
+        limits: args.required_limits.map(Into::into).unwrap_or_default(),
+    };
+
+    let (device, maybe_err) = gfx_select!(adapter => instance.adapter_request_device(
+      adapter,
+      &descriptor,
+      std::env::var("DENO_WEBGPU_TRACE").ok().as_ref().map(std::path::Path::new),
+      std::marker::PhantomData
+    ));
+    if let Some(err) = maybe_err {
+        return Err(DomExceptionOperationError::new(&err.to_string()).into());
+    }
+
+    let device_features = gfx_select!(device => instance.device_features(device))?;
+    let features = deserialize_features(&device_features);
+    let limits = gfx_select!(device => instance.device_limits(device))?;
+
+    let rid = state.resource_table.add(WebGpuDevice(device));
+
+    Ok(GpuAdapterDevice {
+        rid,
+        name: None,
+        features,
+        limits,
+        // TODO(lucacasonato): report correctly from wgpu
+        is_software: false,
+    })
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct CreateQuerySetArgs {
+    device_rid: ResourceId,
+    label: Option<String>,
+    #[serde(flatten)]
+    r#type: GpuQueryType,
+    count: u32,
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "kebab-case", tag = "type")]
+enum GpuQueryType {
+    Occlusion,
+    #[serde(rename_all = "camelCase")]
+    PipelineStatistics {
+        pipeline_statistics: HashSet<String>,
+    },
+    Timestamp,
+}
+
+impl From<GpuQueryType> for wgpu_types::QueryType {
+    fn from(query_type: GpuQueryType) -> Self {
+        match query_type {
+            GpuQueryType::Occlusion => wgpu_types::QueryType::Occlusion,
+            GpuQueryType::PipelineStatistics {
+                pipeline_statistics,
+            } => {
+                use wgpu_types::PipelineStatisticsTypes;
+
+                let mut types = PipelineStatisticsTypes::empty();
+
+                if pipeline_statistics.contains("vertex-shader-invocations") {
+                    types.set(PipelineStatisticsTypes::VERTEX_SHADER_INVOCATIONS, true);
+                }
+                if pipeline_statistics.contains("clipper-invocations") {
+                    types.set(PipelineStatisticsTypes::CLIPPER_INVOCATIONS, true);
+                }
+                if pipeline_statistics.contains("clipper-primitives-out") {
+                    types.set(PipelineStatisticsTypes::CLIPPER_PRIMITIVES_OUT, true);
+                }
+                if pipeline_statistics.contains("fragment-shader-invocations") {
+                    types.set(PipelineStatisticsTypes::FRAGMENT_SHADER_INVOCATIONS, true);
+                }
+                if pipeline_statistics.contains("compute-shader-invocations") {
+                    types.set(PipelineStatisticsTypes::COMPUTE_SHADER_INVOCATIONS, true);
+                }
+
+                wgpu_types::QueryType::PipelineStatistics(types)
+            }
+            GpuQueryType::Timestamp => wgpu_types::QueryType::Timestamp,
+        }
+    }
+}
+
+pub fn op_webgpu_create_query_set(
+    state: &mut OpState,
+    args: CreateQuerySetArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let device_resource = state.resource_table.get::<WebGpuDevice>(args.device_rid)?;
+    let device = device_resource.0;
+    let instance = &state.borrow::<Instance>();
+
+    let descriptor = wgpu_types::QuerySetDescriptor {
+        label: args.label.map(Cow::from),
+        ty: args.r#type.into(),
+        count: args.count,
+    };
+
+    gfx_put!(device => instance.device_create_query_set(
+    device,
+    &descriptor,
+    std::marker::PhantomData
+  ) => state, WebGpuQuerySet)
+}
+
+fn declare_webgpu_ops() -> Vec<(&'static str, Box<OpFn>)> {
+    vec![
+        // Request device/adapter
+        (
+            "op_webgpu_request_adapter",
+            op_async(op_webgpu_request_adapter),
+        ),
+        (
+            "op_webgpu_request_device",
+            op_async(op_webgpu_request_device),
+        ),
+        // Query Set
+        (
+            "op_webgpu_create_query_set",
+            op_sync(op_webgpu_create_query_set),
+        ),
+        // buffer
+        (
+            "op_webgpu_create_buffer",
+            op_sync(buffer::op_webgpu_create_buffer),
+        ),
+        (
+            "op_webgpu_buffer_get_mapped_range",
+            op_sync(buffer::op_webgpu_buffer_get_mapped_range),
+        ),
+        (
+            "op_webgpu_buffer_unmap",
+            op_sync(buffer::op_webgpu_buffer_unmap),
+        ),
+        // buffer async
+        (
+            "op_webgpu_buffer_get_map_async",
+            op_async(buffer::op_webgpu_buffer_get_map_async),
+        ),
+        // remaining sync ops
+
+        // texture
+        (
+            "op_webgpu_create_texture",
+            op_sync(texture::op_webgpu_create_texture),
+        ),
+        (
+            "op_webgpu_create_texture_view",
+            op_sync(texture::op_webgpu_create_texture_view),
+        ),
+        // sampler
+        (
+            "op_webgpu_create_sampler",
+            op_sync(sampler::op_webgpu_create_sampler),
+        ),
+        // binding
+        (
+            "op_webgpu_create_bind_group_layout",
+            op_sync(binding::op_webgpu_create_bind_group_layout),
+        ),
+        (
+            "op_webgpu_create_pipeline_layout",
+            op_sync(binding::op_webgpu_create_pipeline_layout),
+        ),
+        (
+            "op_webgpu_create_bind_group",
+            op_sync(binding::op_webgpu_create_bind_group),
+        ),
+        // pipeline
+        (
+            "op_webgpu_create_compute_pipeline",
+            op_sync(pipeline::op_webgpu_create_compute_pipeline),
+        ),
+        (
+            "op_webgpu_compute_pipeline_get_bind_group_layout",
+            op_sync(pipeline::op_webgpu_compute_pipeline_get_bind_group_layout),
+        ),
+        (
+            "op_webgpu_create_render_pipeline",
+            op_sync(pipeline::op_webgpu_create_render_pipeline),
+        ),
+        (
+            "op_webgpu_render_pipeline_get_bind_group_layout",
+            op_sync(pipeline::op_webgpu_render_pipeline_get_bind_group_layout),
+        ),
+        // command_encoder
+        (
+            "op_webgpu_create_command_encoder",
+            op_sync(command_encoder::op_webgpu_create_command_encoder),
+        ),
+        (
+            "op_webgpu_command_encoder_begin_render_pass",
+            op_sync(command_encoder::op_webgpu_command_encoder_begin_render_pass),
+        ),
+        (
+            "op_webgpu_command_encoder_begin_compute_pass",
+            op_sync(command_encoder::op_webgpu_command_encoder_begin_compute_pass),
+        ),
+        (
+            "op_webgpu_command_encoder_copy_buffer_to_buffer",
+            op_sync(command_encoder::op_webgpu_command_encoder_copy_buffer_to_buffer),
+        ),
+        (
+            "op_webgpu_command_encoder_copy_buffer_to_texture",
+            op_sync(command_encoder::op_webgpu_command_encoder_copy_buffer_to_texture),
+        ),
+        (
+            "op_webgpu_command_encoder_copy_texture_to_buffer",
+            op_sync(command_encoder::op_webgpu_command_encoder_copy_texture_to_buffer),
+        ),
+        (
+            "op_webgpu_command_encoder_copy_texture_to_texture",
+            op_sync(command_encoder::op_webgpu_command_encoder_copy_texture_to_texture),
+        ),
+        (
+            "op_webgpu_command_encoder_push_debug_group",
+            op_sync(command_encoder::op_webgpu_command_encoder_push_debug_group),
+        ),
+        (
+            "op_webgpu_command_encoder_pop_debug_group",
+            op_sync(command_encoder::op_webgpu_command_encoder_pop_debug_group),
+        ),
+        (
+            "op_webgpu_command_encoder_insert_debug_marker",
+            op_sync(command_encoder::op_webgpu_command_encoder_insert_debug_marker),
+        ),
+        (
+            "op_webgpu_command_encoder_write_timestamp",
+            op_sync(command_encoder::op_webgpu_command_encoder_write_timestamp),
+        ),
+        (
+            "op_webgpu_command_encoder_resolve_query_set",
+            op_sync(command_encoder::op_webgpu_command_encoder_resolve_query_set),
+        ),
+        (
+            "op_webgpu_command_encoder_finish",
+            op_sync(command_encoder::op_webgpu_command_encoder_finish),
+        ),
+        // render_pass
+        (
+            "op_webgpu_render_pass_set_viewport",
+            op_sync(render_pass::op_webgpu_render_pass_set_viewport),
+        ),
+        (
+            "op_webgpu_render_pass_set_scissor_rect",
+            op_sync(render_pass::op_webgpu_render_pass_set_scissor_rect),
+        ),
+        (
+            "op_webgpu_render_pass_set_blend_constant",
+            op_sync(render_pass::op_webgpu_render_pass_set_blend_constant),
+        ),
+        (
+            "op_webgpu_render_pass_set_stencil_reference",
+            op_sync(render_pass::op_webgpu_render_pass_set_stencil_reference),
+        ),
+        (
+            "op_webgpu_render_pass_begin_pipeline_statistics_query",
+            op_sync(render_pass::op_webgpu_render_pass_begin_pipeline_statistics_query),
+        ),
+        (
+            "op_webgpu_render_pass_end_pipeline_statistics_query",
+            op_sync(render_pass::op_webgpu_render_pass_end_pipeline_statistics_query),
+        ),
+        (
+            "op_webgpu_render_pass_write_timestamp",
+            op_sync(render_pass::op_webgpu_render_pass_write_timestamp),
+        ),
+        (
+            "op_webgpu_render_pass_execute_bundles",
+            op_sync(render_pass::op_webgpu_render_pass_execute_bundles),
+        ),
+        (
+            "op_webgpu_render_pass_end_pass",
+            op_sync(render_pass::op_webgpu_render_pass_end_pass),
+        ),
+        (
+            "op_webgpu_render_pass_set_bind_group",
+            op_sync(render_pass::op_webgpu_render_pass_set_bind_group),
+        ),
+        (
+            "op_webgpu_render_pass_push_debug_group",
+            op_sync(render_pass::op_webgpu_render_pass_push_debug_group),
+        ),
+        (
+            "op_webgpu_render_pass_pop_debug_group",
+            op_sync(render_pass::op_webgpu_render_pass_pop_debug_group),
+        ),
+        (
+            "op_webgpu_render_pass_insert_debug_marker",
+            op_sync(render_pass::op_webgpu_render_pass_insert_debug_marker),
+        ),
+        (
+            "op_webgpu_render_pass_set_pipeline",
+            op_sync(render_pass::op_webgpu_render_pass_set_pipeline),
+        ),
+        (
+            "op_webgpu_render_pass_set_index_buffer",
+            op_sync(render_pass::op_webgpu_render_pass_set_index_buffer),
+        ),
+        (
+            "op_webgpu_render_pass_set_vertex_buffer",
+            op_sync(render_pass::op_webgpu_render_pass_set_vertex_buffer),
+        ),
+        (
+            "op_webgpu_render_pass_draw",
+            op_sync(render_pass::op_webgpu_render_pass_draw),
+        ),
+        (
+            "op_webgpu_render_pass_draw_indexed",
+            op_sync(render_pass::op_webgpu_render_pass_draw_indexed),
+        ),
+        (
+            "op_webgpu_render_pass_draw_indirect",
+            op_sync(render_pass::op_webgpu_render_pass_draw_indirect),
+        ),
+        (
+            "op_webgpu_render_pass_draw_indexed_indirect",
+            op_sync(render_pass::op_webgpu_render_pass_draw_indexed_indirect),
+        ),
+        // compute_pass
+        (
+            "op_webgpu_compute_pass_set_pipeline",
+            op_sync(compute_pass::op_webgpu_compute_pass_set_pipeline),
+        ),
+        (
+            "op_webgpu_compute_pass_dispatch",
+            op_sync(compute_pass::op_webgpu_compute_pass_dispatch),
+        ),
+        (
+            "op_webgpu_compute_pass_dispatch_indirect",
+            op_sync(compute_pass::op_webgpu_compute_pass_dispatch_indirect),
+        ),
+        (
+            "op_webgpu_compute_pass_end_pass",
+            op_sync(compute_pass::op_webgpu_compute_pass_end_pass),
+        ),
+        (
+            "op_webgpu_compute_pass_set_bind_group",
+            op_sync(compute_pass::op_webgpu_compute_pass_set_bind_group),
+        ),
+        (
+            "op_webgpu_compute_pass_push_debug_group",
+            op_sync(compute_pass::op_webgpu_compute_pass_push_debug_group),
+        ),
+        (
+            "op_webgpu_compute_pass_pop_debug_group",
+            op_sync(compute_pass::op_webgpu_compute_pass_pop_debug_group),
+        ),
+        (
+            "op_webgpu_compute_pass_insert_debug_marker",
+            op_sync(compute_pass::op_webgpu_compute_pass_insert_debug_marker),
+        ),
+        // bundle
+        (
+            "op_webgpu_create_render_bundle_encoder",
+            op_sync(bundle::op_webgpu_create_render_bundle_encoder),
+        ),
+        (
+            "op_webgpu_render_bundle_encoder_finish",
+            op_sync(bundle::op_webgpu_render_bundle_encoder_finish),
+        ),
+        (
+            "op_webgpu_render_bundle_encoder_set_bind_group",
+            op_sync(bundle::op_webgpu_render_bundle_encoder_set_bind_group),
+        ),
+        (
+            "op_webgpu_render_bundle_encoder_push_debug_group",
+            op_sync(bundle::op_webgpu_render_bundle_encoder_push_debug_group),
+        ),
+        (
+            "op_webgpu_render_bundle_encoder_pop_debug_group",
+            op_sync(bundle::op_webgpu_render_bundle_encoder_pop_debug_group),
+        ),
+        (
+            "op_webgpu_render_bundle_encoder_insert_debug_marker",
+            op_sync(bundle::op_webgpu_render_bundle_encoder_insert_debug_marker),
+        ),
+        (
+            "op_webgpu_render_bundle_encoder_set_pipeline",
+            op_sync(bundle::op_webgpu_render_bundle_encoder_set_pipeline),
+        ),
+        (
+            "op_webgpu_render_bundle_encoder_set_index_buffer",
+            op_sync(bundle::op_webgpu_render_bundle_encoder_set_index_buffer),
+        ),
+        (
+            "op_webgpu_render_bundle_encoder_set_vertex_buffer",
+            op_sync(bundle::op_webgpu_render_bundle_encoder_set_vertex_buffer),
+        ),
+        (
+            "op_webgpu_render_bundle_encoder_draw",
+            op_sync(bundle::op_webgpu_render_bundle_encoder_draw),
+        ),
+        (
+            "op_webgpu_render_bundle_encoder_draw_indexed",
+            op_sync(bundle::op_webgpu_render_bundle_encoder_draw_indexed),
+        ),
+        (
+            "op_webgpu_render_bundle_encoder_draw_indirect",
+            op_sync(bundle::op_webgpu_render_bundle_encoder_draw_indirect),
+        ),
+        // queue
+        (
+            "op_webgpu_queue_submit",
+            op_sync(queue::op_webgpu_queue_submit),
+        ),
+        (
+            "op_webgpu_write_buffer",
+            op_sync(queue::op_webgpu_write_buffer),
+        ),
+        (
+            "op_webgpu_write_texture",
+            op_sync(queue::op_webgpu_write_texture),
+        ),
+        // shader
+        (
+            "op_webgpu_create_shader_module",
+            op_sync(shader::op_webgpu_create_shader_module),
+        ),
+    ]
+}
--- a/deno_webgpu/pipeline.rs
+++ b/deno_webgpu/pipeline.rs
@ -0,0 +1,417 @@
+// Copyright 2018-2021 the Deno authors. All rights reserved. MIT license.
+
+use deno_core::error::AnyError;
+use deno_core::ResourceId;
+use deno_core::{OpState, Resource};
+use serde::Deserialize;
+use serde::Serialize;
+use std::borrow::Cow;
+use std::convert::{TryFrom, TryInto};
+
+use super::error::WebGpuError;
+use super::error::WebGpuResult;
+
+const MAX_BIND_GROUPS: usize = 8;
+
+pub(crate) struct WebGpuPipelineLayout(pub(crate) wgpu_core::id::PipelineLayoutId);
+impl Resource for WebGpuPipelineLayout {
+    fn name(&self) -> Cow<str> {
+        "webGPUPipelineLayout".into()
+    }
+}
+
+pub(crate) struct WebGpuComputePipeline(pub(crate) wgpu_core::id::ComputePipelineId);
+impl Resource for WebGpuComputePipeline {
+    fn name(&self) -> Cow<str> {
+        "webGPUComputePipeline".into()
+    }
+}
+
+pub(crate) struct WebGpuRenderPipeline(pub(crate) wgpu_core::id::RenderPipelineId);
+impl Resource for WebGpuRenderPipeline {
+    fn name(&self) -> Cow<str> {
+        "webGPURenderPipeline".into()
+    }
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+struct GpuProgrammableStage {
+    module: ResourceId,
+    entry_point: String,
+    // constants: HashMap<String, GPUPipelineConstantValue>
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct CreateComputePipelineArgs {
+    device_rid: ResourceId,
+    label: Option<String>,
+    layout: Option<ResourceId>,
+    compute: GpuProgrammableStage,
+}
+
+pub fn op_webgpu_create_compute_pipeline(
+    state: &mut OpState,
+    args: CreateComputePipelineArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let instance = state.borrow::<super::Instance>();
+    let device_resource = state
+        .resource_table
+        .get::<super::WebGpuDevice>(args.device_rid)?;
+    let device = device_resource.0;
+
+    let pipeline_layout = if let Some(rid) = args.layout {
+        let id = state.resource_table.get::<WebGpuPipelineLayout>(rid)?;
+        Some(id.0)
+    } else {
+        None
+    };
+
+    let compute_shader_module_resource = state
+        .resource_table
+        .get::<super::shader::WebGpuShaderModule>(args.compute.module)?;
+
+    let descriptor = wgpu_core::pipeline::ComputePipelineDescriptor {
+        label: args.label.map(Cow::from),
+        layout: pipeline_layout,
+        stage: wgpu_core::pipeline::ProgrammableStageDescriptor {
+            module: compute_shader_module_resource.0,
+            entry_point: Cow::from(args.compute.entry_point),
+            // TODO(lucacasonato): support args.compute.constants
+        },
+    };
+    let implicit_pipelines = match args.layout {
+        Some(_) => None,
+        None => Some(wgpu_core::device::ImplicitPipelineIds {
+            root_id: std::marker::PhantomData,
+            group_ids: &[std::marker::PhantomData; MAX_BIND_GROUPS],
+        }),
+    };
+
+    let (compute_pipeline, maybe_err) = gfx_select!(device => instance.device_create_compute_pipeline(
+      device,
+      &descriptor,
+      std::marker::PhantomData,
+      implicit_pipelines
+    ));
+
+    let rid = state
+        .resource_table
+        .add(WebGpuComputePipeline(compute_pipeline));
+
+    Ok(WebGpuResult::rid_err(rid, maybe_err))
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct ComputePipelineGetBindGroupLayoutArgs {
+    compute_pipeline_rid: ResourceId,
+    index: u32,
+}
+
+#[derive(Serialize)]
+#[serde(rename_all = "camelCase")]
+pub struct PipelineLayout {
+    rid: ResourceId,
+    label: String,
+    err: Option<WebGpuError>,
+}
+
+pub fn op_webgpu_compute_pipeline_get_bind_group_layout(
+    state: &mut OpState,
+    args: ComputePipelineGetBindGroupLayoutArgs,
+    _: (),
+) -> Result<PipelineLayout, AnyError> {
+    let instance = state.borrow::<super::Instance>();
+    let compute_pipeline_resource = state
+        .resource_table
+        .get::<WebGpuComputePipeline>(args.compute_pipeline_rid)?;
+    let compute_pipeline = compute_pipeline_resource.0;
+
+    let (bind_group_layout, maybe_err) = gfx_select!(compute_pipeline => instance.compute_pipeline_get_bind_group_layout(compute_pipeline, args.index, std::marker::PhantomData));
+
+    let label =
+        gfx_select!(bind_group_layout => instance.bind_group_layout_label(bind_group_layout));
+
+    let rid = state
+        .resource_table
+        .add(super::binding::WebGpuBindGroupLayout(bind_group_layout));
+
+    Ok(PipelineLayout {
+        rid,
+        label,
+        err: maybe_err.map(WebGpuError::from),
+    })
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "kebab-case")]
+pub enum GpuCullMode {
+    None,
+    Front,
+    Back,
+}
+
+impl From<GpuCullMode> for Option<wgpu_types::Face> {
+    fn from(value: GpuCullMode) -> Option<wgpu_types::Face> {
+        match value {
+            GpuCullMode::None => None,
+            GpuCullMode::Front => Some(wgpu_types::Face::Front),
+            GpuCullMode::Back => Some(wgpu_types::Face::Back),
+        }
+    }
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+struct GpuPrimitiveState {
+    topology: wgpu_types::PrimitiveTopology,
+    strip_index_format: Option<wgpu_types::IndexFormat>,
+    front_face: wgpu_types::FrontFace,
+    cull_mode: GpuCullMode,
+    clamp_depth: bool,
+}
+
+impl From<GpuPrimitiveState> for wgpu_types::PrimitiveState {
+    fn from(value: GpuPrimitiveState) -> wgpu_types::PrimitiveState {
+        wgpu_types::PrimitiveState {
+            topology: value.topology,
+            strip_index_format: value.strip_index_format,
+            front_face: value.front_face,
+            cull_mode: value.cull_mode.into(),
+            clamp_depth: value.clamp_depth,
+            polygon_mode: Default::default(), // native-only
+            conservative: false,              // native-only
+        }
+    }
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+struct GpuDepthStencilState {
+    format: wgpu_types::TextureFormat,
+    depth_write_enabled: bool,
+    depth_compare: wgpu_types::CompareFunction,
+    stencil_front: wgpu_types::StencilFaceState,
+    stencil_back: wgpu_types::StencilFaceState,
+    stencil_read_mask: u32,
+    stencil_write_mask: u32,
+    depth_bias: i32,
+    depth_bias_slope_scale: f32,
+    depth_bias_clamp: f32,
+}
+
+impl TryFrom<GpuDepthStencilState> for wgpu_types::DepthStencilState {
+    type Error = AnyError;
+    fn try_from(state: GpuDepthStencilState) -> Result<wgpu_types::DepthStencilState, AnyError> {
+        Ok(wgpu_types::DepthStencilState {
+            format: state.format,
+            depth_write_enabled: state.depth_write_enabled,
+            depth_compare: state.depth_compare.into(),
+            stencil: wgpu_types::StencilState {
+                front: state.stencil_front.into(),
+                back: state.stencil_back.into(),
+                read_mask: state.stencil_read_mask,
+                write_mask: state.stencil_write_mask,
+            },
+            bias: wgpu_types::DepthBiasState {
+                constant: state.depth_bias,
+                slope_scale: state.depth_bias_slope_scale,
+                clamp: state.depth_bias_clamp,
+            },
+        })
+    }
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+struct GpuVertexBufferLayout {
+    array_stride: u64,
+    step_mode: wgpu_types::VertexStepMode,
+    attributes: Vec<wgpu_types::VertexAttribute>,
+}
+
+impl<'a> From<GpuVertexBufferLayout> for wgpu_core::pipeline::VertexBufferLayout<'a> {
+    fn from(layout: GpuVertexBufferLayout) -> wgpu_core::pipeline::VertexBufferLayout<'a> {
+        wgpu_core::pipeline::VertexBufferLayout {
+            array_stride: layout.array_stride,
+            step_mode: layout.step_mode,
+            attributes: Cow::Owned(layout.attributes),
+        }
+    }
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+struct GpuVertexState {
+    module: ResourceId,
+    entry_point: String,
+    buffers: Vec<Option<GpuVertexBufferLayout>>,
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+struct GpuMultisampleState {
+    count: u32,
+    mask: u64,
+    alpha_to_coverage_enabled: bool,
+}
+
+impl From<GpuMultisampleState> for wgpu_types::MultisampleState {
+    fn from(gms: GpuMultisampleState) -> wgpu_types::MultisampleState {
+        wgpu_types::MultisampleState {
+            count: gms.count,
+            mask: gms.mask,
+            alpha_to_coverage_enabled: gms.alpha_to_coverage_enabled,
+        }
+    }
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+struct GpuFragmentState {
+    targets: Vec<wgpu_types::ColorTargetState>,
+    module: u32,
+    entry_point: String,
+    // TODO(lucacasonato): constants
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct CreateRenderPipelineArgs {
+    device_rid: ResourceId,
+    label: Option<String>,
+    layout: Option<ResourceId>,
+    vertex: GpuVertexState,
+    primitive: GpuPrimitiveState,
+    depth_stencil: Option<GpuDepthStencilState>,
+    multisample: wgpu_types::MultisampleState,
+    fragment: Option<GpuFragmentState>,
+}
+
+pub fn op_webgpu_create_render_pipeline(
+    state: &mut OpState,
+    args: CreateRenderPipelineArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let instance = state.borrow::<super::Instance>();
+    let device_resource = state
+        .resource_table
+        .get::<super::WebGpuDevice>(args.device_rid)?;
+    let device = device_resource.0;
+
+    let layout = if let Some(rid) = args.layout {
+        let pipeline_layout_resource = state.resource_table.get::<WebGpuPipelineLayout>(rid)?;
+        Some(pipeline_layout_resource.0)
+    } else {
+        None
+    };
+
+    let vertex_shader_module_resource = state
+        .resource_table
+        .get::<super::shader::WebGpuShaderModule>(args.vertex.module)?;
+
+    let fragment = if let Some(fragment) = args.fragment {
+        let fragment_shader_module_resource =
+            state
+                .resource_table
+                .get::<super::shader::WebGpuShaderModule>(fragment.module)?;
+
+        let mut targets = Vec::with_capacity(fragment.targets.len());
+
+        for target in fragment.targets {
+            targets.push(target.try_into()?);
+        }
+
+        Some(wgpu_core::pipeline::FragmentState {
+            stage: wgpu_core::pipeline::ProgrammableStageDescriptor {
+                module: fragment_shader_module_resource.0,
+                entry_point: Cow::from(fragment.entry_point),
+            },
+            targets: Cow::from(targets),
+        })
+    } else {
+        None
+    };
+
+    let vertex_buffers = args
+        .vertex
+        .buffers
+        .into_iter()
+        .flatten()
+        .map(Into::into)
+        .collect();
+
+    let descriptor = wgpu_core::pipeline::RenderPipelineDescriptor {
+        label: args.label.map(Cow::Owned),
+        layout,
+        vertex: wgpu_core::pipeline::VertexState {
+            stage: wgpu_core::pipeline::ProgrammableStageDescriptor {
+                module: vertex_shader_module_resource.0,
+                entry_point: Cow::Owned(args.vertex.entry_point),
+            },
+            buffers: Cow::Owned(vertex_buffers),
+        },
+        primitive: args.primitive.into(),
+        depth_stencil: args.depth_stencil.map(TryInto::try_into).transpose()?,
+        multisample: args.multisample.into(),
+        fragment,
+    };
+
+    let implicit_pipelines = match args.layout {
+        Some(_) => None,
+        None => Some(wgpu_core::device::ImplicitPipelineIds {
+            root_id: std::marker::PhantomData,
+            group_ids: &[std::marker::PhantomData; MAX_BIND_GROUPS],
+        }),
+    };
+
+    let (render_pipeline, maybe_err) = gfx_select!(device => instance.device_create_render_pipeline(
+      device,
+      &descriptor,
+      std::marker::PhantomData,
+      implicit_pipelines
+    ));
+
+    let rid = state
+        .resource_table
+        .add(WebGpuRenderPipeline(render_pipeline));
+
+    Ok(WebGpuResult::rid_err(rid, maybe_err))
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct RenderPipelineGetBindGroupLayoutArgs {
+    render_pipeline_rid: ResourceId,
+    index: u32,
+}
+
+pub fn op_webgpu_render_pipeline_get_bind_group_layout(
+    state: &mut OpState,
+    args: RenderPipelineGetBindGroupLayoutArgs,
+    _: (),
+) -> Result<PipelineLayout, AnyError> {
+    let instance = state.borrow::<super::Instance>();
+    let render_pipeline_resource = state
+        .resource_table
+        .get::<WebGpuRenderPipeline>(args.render_pipeline_rid)?;
+    let render_pipeline = render_pipeline_resource.0;
+
+    let (bind_group_layout, maybe_err) = gfx_select!(render_pipeline => instance.render_pipeline_get_bind_group_layout(render_pipeline, args.index, std::marker::PhantomData));
+
+    let label =
+        gfx_select!(bind_group_layout => instance.bind_group_layout_label(bind_group_layout));
+
+    let rid = state
+        .resource_table
+        .add(super::binding::WebGpuBindGroupLayout(bind_group_layout));
+
+    Ok(PipelineLayout {
+        rid,
+        label,
+        err: maybe_err.map(WebGpuError::from),
+    })
+}
--- a/deno_webgpu/queue.rs
+++ b/deno_webgpu/queue.rs
@ -0,0 +1,140 @@
+// Copyright 2018-2021 the Deno authors. All rights reserved. MIT license.
+
+use std::num::NonZeroU32;
+
+use deno_core::error::null_opbuf;
+use deno_core::error::AnyError;
+use deno_core::OpState;
+use deno_core::ResourceId;
+use deno_core::ZeroCopyBuf;
+use serde::Deserialize;
+
+use super::error::WebGpuResult;
+
+type WebGpuQueue = super::WebGpuDevice;
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct QueueSubmitArgs {
+    queue_rid: ResourceId,
+    command_buffers: Vec<ResourceId>,
+}
+
+pub fn op_webgpu_queue_submit(
+    state: &mut OpState,
+    args: QueueSubmitArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let instance = state.borrow::<super::Instance>();
+    let queue_resource = state.resource_table.get::<WebGpuQueue>(args.queue_rid)?;
+    let queue = queue_resource.0;
+
+    let mut ids = vec![];
+
+    for rid in args.command_buffers {
+        let buffer_resource = state
+            .resource_table
+            .get::<super::command_encoder::WebGpuCommandBuffer>(rid)?;
+        ids.push(buffer_resource.0);
+    }
+
+    let maybe_err = gfx_select!(queue => instance.queue_submit(queue, &ids)).err();
+
+    Ok(WebGpuResult::maybe_err(maybe_err))
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+struct GpuImageDataLayout {
+    offset: u64,
+    bytes_per_row: Option<u32>,
+    rows_per_image: Option<u32>,
+}
+
+impl From<GpuImageDataLayout> for wgpu_types::ImageDataLayout {
+    fn from(layout: GpuImageDataLayout) -> Self {
+        wgpu_types::ImageDataLayout {
+            offset: layout.offset,
+            bytes_per_row: NonZeroU32::new(layout.bytes_per_row.unwrap_or(0)),
+            rows_per_image: NonZeroU32::new(layout.rows_per_image.unwrap_or(0)),
+        }
+    }
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct QueueWriteBufferArgs {
+    queue_rid: ResourceId,
+    buffer: ResourceId,
+    buffer_offset: u64,
+    data_offset: usize,
+    size: Option<usize>,
+}
+
+pub fn op_webgpu_write_buffer(
+    state: &mut OpState,
+    args: QueueWriteBufferArgs,
+    zero_copy: Option<ZeroCopyBuf>,
+) -> Result<WebGpuResult, AnyError> {
+    let zero_copy = zero_copy.ok_or_else(null_opbuf)?;
+    let instance = state.borrow::<super::Instance>();
+    let buffer_resource = state
+        .resource_table
+        .get::<super::buffer::WebGpuBuffer>(args.buffer)?;
+    let buffer = buffer_resource.0;
+    let queue_resource = state.resource_table.get::<WebGpuQueue>(args.queue_rid)?;
+    let queue = queue_resource.0;
+
+    let data = match args.size {
+        Some(size) => &zero_copy[args.data_offset..(args.data_offset + size)],
+        None => &zero_copy[args.data_offset..],
+    };
+    let maybe_err = gfx_select!(queue => instance.queue_write_buffer(
+      queue,
+      buffer,
+      args.buffer_offset,
+      data
+    ))
+    .err();
+
+    Ok(WebGpuResult::maybe_err(maybe_err))
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct QueueWriteTextureArgs {
+    queue_rid: ResourceId,
+    destination: super::command_encoder::GpuImageCopyTexture,
+    data_layout: GpuImageDataLayout,
+    size: wgpu_types::Extent3d,
+}
+
+pub fn op_webgpu_write_texture(
+    state: &mut OpState,
+    args: QueueWriteTextureArgs,
+    zero_copy: Option<ZeroCopyBuf>,
+) -> Result<WebGpuResult, AnyError> {
+    let zero_copy = zero_copy.ok_or_else(null_opbuf)?;
+    let instance = state.borrow::<super::Instance>();
+    let texture_resource = state
+        .resource_table
+        .get::<super::texture::WebGpuTexture>(args.destination.texture)?;
+    let queue_resource = state.resource_table.get::<WebGpuQueue>(args.queue_rid)?;
+    let queue = queue_resource.0;
+
+    let destination = wgpu_core::command::ImageCopyTexture {
+        texture: texture_resource.0,
+        mip_level: args.destination.mip_level,
+        origin: args.destination.origin,
+        aspect: args.destination.aspect,
+    };
+    let data_layout = args.data_layout.into();
+
+    gfx_ok!(queue => instance.queue_write_texture(
+      queue,
+      &destination,
+      &*zero_copy,
+      &data_layout,
+      &args.size
+    ))
+}
--- a/deno_webgpu/render_pass.rs
+++ b/deno_webgpu/render_pass.rs
@ -0,0 +1,643 @@
+// Copyright 2018-2021 the Deno authors. All rights reserved. MIT license.
+
+use deno_core::error::type_error;
+use deno_core::error::AnyError;
+use deno_core::ResourceId;
+use deno_core::ZeroCopyBuf;
+use deno_core::{OpState, Resource};
+use serde::Deserialize;
+use std::borrow::Cow;
+use std::cell::RefCell;
+
+use super::error::WebGpuResult;
+
+pub(crate) struct WebGpuRenderPass(pub(crate) RefCell<wgpu_core::command::RenderPass>);
+impl Resource for WebGpuRenderPass {
+    fn name(&self) -> Cow<str> {
+        "webGPURenderPass".into()
+    }
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct RenderPassSetViewportArgs {
+    render_pass_rid: ResourceId,
+    x: f32,
+    y: f32,
+    width: f32,
+    height: f32,
+    min_depth: f32,
+    max_depth: f32,
+}
+
+pub fn op_webgpu_render_pass_set_viewport(
+    state: &mut OpState,
+    args: RenderPassSetViewportArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let render_pass_resource = state
+        .resource_table
+        .get::<WebGpuRenderPass>(args.render_pass_rid)?;
+
+    wgpu_core::command::render_ffi::wgpu_render_pass_set_viewport(
+        &mut render_pass_resource.0.borrow_mut(),
+        args.x,
+        args.y,
+        args.width,
+        args.height,
+        args.min_depth,
+        args.max_depth,
+    );
+
+    Ok(WebGpuResult::empty())
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct RenderPassSetScissorRectArgs {
+    render_pass_rid: ResourceId,
+    x: u32,
+    y: u32,
+    width: u32,
+    height: u32,
+}
+
+pub fn op_webgpu_render_pass_set_scissor_rect(
+    state: &mut OpState,
+    args: RenderPassSetScissorRectArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let render_pass_resource = state
+        .resource_table
+        .get::<WebGpuRenderPass>(args.render_pass_rid)?;
+
+    wgpu_core::command::render_ffi::wgpu_render_pass_set_scissor_rect(
+        &mut render_pass_resource.0.borrow_mut(),
+        args.x,
+        args.y,
+        args.width,
+        args.height,
+    );
+
+    Ok(WebGpuResult::empty())
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct RenderPassSetBlendConstantArgs {
+    render_pass_rid: ResourceId,
+    color: wgpu_types::Color,
+}
+
+pub fn op_webgpu_render_pass_set_blend_constant(
+    state: &mut OpState,
+    args: RenderPassSetBlendConstantArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let render_pass_resource = state
+        .resource_table
+        .get::<WebGpuRenderPass>(args.render_pass_rid)?;
+
+    wgpu_core::command::render_ffi::wgpu_render_pass_set_blend_constant(
+        &mut render_pass_resource.0.borrow_mut(),
+        &args.color,
+    );
+
+    Ok(WebGpuResult::empty())
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct RenderPassSetStencilReferenceArgs {
+    render_pass_rid: ResourceId,
+    reference: u32,
+}
+
+pub fn op_webgpu_render_pass_set_stencil_reference(
+    state: &mut OpState,
+    args: RenderPassSetStencilReferenceArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let render_pass_resource = state
+        .resource_table
+        .get::<WebGpuRenderPass>(args.render_pass_rid)?;
+
+    wgpu_core::command::render_ffi::wgpu_render_pass_set_stencil_reference(
+        &mut render_pass_resource.0.borrow_mut(),
+        args.reference,
+    );
+
+    Ok(WebGpuResult::empty())
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct RenderPassBeginPipelineStatisticsQueryArgs {
+    render_pass_rid: ResourceId,
+    query_set: u32,
+    query_index: u32,
+}
+
+pub fn op_webgpu_render_pass_begin_pipeline_statistics_query(
+    state: &mut OpState,
+    args: RenderPassBeginPipelineStatisticsQueryArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let render_pass_resource = state
+        .resource_table
+        .get::<WebGpuRenderPass>(args.render_pass_rid)?;
+    let query_set_resource = state
+        .resource_table
+        .get::<super::WebGpuQuerySet>(args.query_set)?;
+
+    wgpu_core::command::render_ffi::wgpu_render_pass_begin_pipeline_statistics_query(
+        &mut render_pass_resource.0.borrow_mut(),
+        query_set_resource.0,
+        args.query_index,
+    );
+
+    Ok(WebGpuResult::empty())
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct RenderPassEndPipelineStatisticsQueryArgs {
+    render_pass_rid: ResourceId,
+}
+
+pub fn op_webgpu_render_pass_end_pipeline_statistics_query(
+    state: &mut OpState,
+    args: RenderPassEndPipelineStatisticsQueryArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let render_pass_resource = state
+        .resource_table
+        .get::<WebGpuRenderPass>(args.render_pass_rid)?;
+
+    wgpu_core::command::render_ffi::wgpu_render_pass_end_pipeline_statistics_query(
+        &mut render_pass_resource.0.borrow_mut(),
+    );
+
+    Ok(WebGpuResult::empty())
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct RenderPassWriteTimestampArgs {
+    render_pass_rid: ResourceId,
+    query_set: u32,
+    query_index: u32,
+}
+
+pub fn op_webgpu_render_pass_write_timestamp(
+    state: &mut OpState,
+    args: RenderPassWriteTimestampArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let render_pass_resource = state
+        .resource_table
+        .get::<WebGpuRenderPass>(args.render_pass_rid)?;
+    let query_set_resource = state
+        .resource_table
+        .get::<super::WebGpuQuerySet>(args.query_set)?;
+
+    wgpu_core::command::render_ffi::wgpu_render_pass_write_timestamp(
+        &mut render_pass_resource.0.borrow_mut(),
+        query_set_resource.0,
+        args.query_index,
+    );
+
+    Ok(WebGpuResult::empty())
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct RenderPassExecuteBundlesArgs {
+    render_pass_rid: ResourceId,
+    bundles: Vec<u32>,
+}
+
+pub fn op_webgpu_render_pass_execute_bundles(
+    state: &mut OpState,
+    args: RenderPassExecuteBundlesArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let mut render_bundle_ids = vec![];
+
+    for rid in &args.bundles {
+        let render_bundle_resource = state
+            .resource_table
+            .get::<super::bundle::WebGpuRenderBundle>(*rid)?;
+        render_bundle_ids.push(render_bundle_resource.0);
+    }
+
+    let render_pass_resource = state
+        .resource_table
+        .get::<WebGpuRenderPass>(args.render_pass_rid)?;
+
+    // SAFETY: the raw pointer and length are of the same slice, and that slice
+    // lives longer than the below function invocation.
+    unsafe {
+        wgpu_core::command::render_ffi::wgpu_render_pass_execute_bundles(
+            &mut render_pass_resource.0.borrow_mut(),
+            render_bundle_ids.as_ptr(),
+            render_bundle_ids.len(),
+        );
+    }
+
+    Ok(WebGpuResult::empty())
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct RenderPassEndPassArgs {
+    command_encoder_rid: ResourceId,
+    render_pass_rid: ResourceId,
+}
+
+pub fn op_webgpu_render_pass_end_pass(
+    state: &mut OpState,
+    args: RenderPassEndPassArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let command_encoder_resource =
+        state
+            .resource_table
+            .get::<super::command_encoder::WebGpuCommandEncoder>(args.command_encoder_rid)?;
+    let command_encoder = command_encoder_resource.0;
+    let render_pass_resource = state
+        .resource_table
+        .take::<WebGpuRenderPass>(args.render_pass_rid)?;
+    let render_pass = &render_pass_resource.0.borrow();
+    let instance = state.borrow::<super::Instance>();
+
+    gfx_ok!(command_encoder => instance.command_encoder_run_render_pass(command_encoder, render_pass))
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct RenderPassSetBindGroupArgs {
+    render_pass_rid: ResourceId,
+    index: u32,
+    bind_group: u32,
+    dynamic_offsets_data: ZeroCopyBuf,
+    dynamic_offsets_data_start: usize,
+    dynamic_offsets_data_length: usize,
+}
+
+pub fn op_webgpu_render_pass_set_bind_group(
+    state: &mut OpState,
+    args: RenderPassSetBindGroupArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let bind_group_resource = state
+        .resource_table
+        .get::<super::binding::WebGpuBindGroup>(args.bind_group)?;
+    let render_pass_resource = state
+        .resource_table
+        .get::<WebGpuRenderPass>(args.render_pass_rid)?;
+
+    // Align the data
+    assert!(args.dynamic_offsets_data_start % std::mem::size_of::<u32>() == 0);
+    // SAFETY: A u8 to u32 cast is safe because we asserted that the length is a
+    // multiple of 4.
+    let (prefix, dynamic_offsets_data, suffix) =
+        unsafe { args.dynamic_offsets_data.align_to::<u32>() };
+    assert!(prefix.is_empty());
+    assert!(suffix.is_empty());
+
+    let start = args.dynamic_offsets_data_start;
+    let len = args.dynamic_offsets_data_length;
+
+    // Assert that length and start are both in bounds
+    assert!(start <= dynamic_offsets_data.len());
+    assert!(len <= dynamic_offsets_data.len() - start);
+
+    let dynamic_offsets_data: &[u32] = &dynamic_offsets_data[start..start + len];
+
+    // SAFETY: the raw pointer and length are of the same slice, and that slice
+    // lives longer than the below function invocation.
+    unsafe {
+        wgpu_core::command::render_ffi::wgpu_render_pass_set_bind_group(
+            &mut render_pass_resource.0.borrow_mut(),
+            args.index,
+            bind_group_resource.0,
+            dynamic_offsets_data.as_ptr(),
+            dynamic_offsets_data.len(),
+        );
+    }
+
+    Ok(WebGpuResult::empty())
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct RenderPassPushDebugGroupArgs {
+    render_pass_rid: ResourceId,
+    group_label: String,
+}
+
+pub fn op_webgpu_render_pass_push_debug_group(
+    state: &mut OpState,
+    args: RenderPassPushDebugGroupArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let render_pass_resource = state
+        .resource_table
+        .get::<WebGpuRenderPass>(args.render_pass_rid)?;
+
+    let label = std::ffi::CString::new(args.group_label).unwrap();
+    // SAFETY: the string the raw pointer points to lives longer than the below
+    // function invocation.
+    unsafe {
+        wgpu_core::command::render_ffi::wgpu_render_pass_push_debug_group(
+            &mut render_pass_resource.0.borrow_mut(),
+            label.as_ptr(),
+            0, // wgpu#975
+        );
+    }
+
+    Ok(WebGpuResult::empty())
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct RenderPassPopDebugGroupArgs {
+    render_pass_rid: ResourceId,
+}
+
+pub fn op_webgpu_render_pass_pop_debug_group(
+    state: &mut OpState,
+    args: RenderPassPopDebugGroupArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let render_pass_resource = state
+        .resource_table
+        .get::<WebGpuRenderPass>(args.render_pass_rid)?;
+
+    wgpu_core::command::render_ffi::wgpu_render_pass_pop_debug_group(
+        &mut render_pass_resource.0.borrow_mut(),
+    );
+
+    Ok(WebGpuResult::empty())
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct RenderPassInsertDebugMarkerArgs {
+    render_pass_rid: ResourceId,
+    marker_label: String,
+}
+
+pub fn op_webgpu_render_pass_insert_debug_marker(
+    state: &mut OpState,
+    args: RenderPassInsertDebugMarkerArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let render_pass_resource = state
+        .resource_table
+        .get::<WebGpuRenderPass>(args.render_pass_rid)?;
+
+    let label = std::ffi::CString::new(args.marker_label).unwrap();
+    // SAFETY: the string the raw pointer points to lives longer than the below
+    // function invocation.
+    unsafe {
+        wgpu_core::command::render_ffi::wgpu_render_pass_insert_debug_marker(
+            &mut render_pass_resource.0.borrow_mut(),
+            label.as_ptr(),
+            0, // wgpu#975
+        );
+    }
+
+    Ok(WebGpuResult::empty())
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct RenderPassSetPipelineArgs {
+    render_pass_rid: ResourceId,
+    pipeline: u32,
+}
+
+pub fn op_webgpu_render_pass_set_pipeline(
+    state: &mut OpState,
+    args: RenderPassSetPipelineArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let render_pipeline_resource = state
+        .resource_table
+        .get::<super::pipeline::WebGpuRenderPipeline>(args.pipeline)?;
+    let render_pass_resource = state
+        .resource_table
+        .get::<WebGpuRenderPass>(args.render_pass_rid)?;
+
+    wgpu_core::command::render_ffi::wgpu_render_pass_set_pipeline(
+        &mut render_pass_resource.0.borrow_mut(),
+        render_pipeline_resource.0,
+    );
+
+    Ok(WebGpuResult::empty())
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct RenderPassSetIndexBufferArgs {
+    render_pass_rid: ResourceId,
+    buffer: u32,
+    index_format: wgpu_types::IndexFormat,
+    offset: u64,
+    size: Option<u64>,
+}
+
+pub fn op_webgpu_render_pass_set_index_buffer(
+    state: &mut OpState,
+    args: RenderPassSetIndexBufferArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let buffer_resource = state
+        .resource_table
+        .get::<super::buffer::WebGpuBuffer>(args.buffer)?;
+    let render_pass_resource = state
+        .resource_table
+        .get::<WebGpuRenderPass>(args.render_pass_rid)?;
+
+    let size = if let Some(size) = args.size {
+        Some(
+            std::num::NonZeroU64::new(size)
+                .ok_or_else(|| type_error("size must be larger than 0"))?,
+        )
+    } else {
+        None
+    };
+
+    render_pass_resource.0.borrow_mut().set_index_buffer(
+        buffer_resource.0,
+        args.index_format,
+        args.offset,
+        size,
+    );
+
+    Ok(WebGpuResult::empty())
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct RenderPassSetVertexBufferArgs {
+    render_pass_rid: ResourceId,
+    slot: u32,
+    buffer: u32,
+    offset: u64,
+    size: Option<u64>,
+}
+
+pub fn op_webgpu_render_pass_set_vertex_buffer(
+    state: &mut OpState,
+    args: RenderPassSetVertexBufferArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let buffer_resource = state
+        .resource_table
+        .get::<super::buffer::WebGpuBuffer>(args.buffer)?;
+    let render_pass_resource = state
+        .resource_table
+        .get::<WebGpuRenderPass>(args.render_pass_rid)?;
+
+    let size = if let Some(size) = args.size {
+        Some(
+            std::num::NonZeroU64::new(size)
+                .ok_or_else(|| type_error("size must be larger than 0"))?,
+        )
+    } else {
+        None
+    };
+
+    wgpu_core::command::render_ffi::wgpu_render_pass_set_vertex_buffer(
+        &mut render_pass_resource.0.borrow_mut(),
+        args.slot,
+        buffer_resource.0,
+        args.offset,
+        size,
+    );
+
+    Ok(WebGpuResult::empty())
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct RenderPassDrawArgs {
+    render_pass_rid: ResourceId,
+    vertex_count: u32,
+    instance_count: u32,
+    first_vertex: u32,
+    first_instance: u32,
+}
+
+pub fn op_webgpu_render_pass_draw(
+    state: &mut OpState,
+    args: RenderPassDrawArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let render_pass_resource = state
+        .resource_table
+        .get::<WebGpuRenderPass>(args.render_pass_rid)?;
+
+    wgpu_core::command::render_ffi::wgpu_render_pass_draw(
+        &mut render_pass_resource.0.borrow_mut(),
+        args.vertex_count,
+        args.instance_count,
+        args.first_vertex,
+        args.first_instance,
+    );
+
+    Ok(WebGpuResult::empty())
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct RenderPassDrawIndexedArgs {
+    render_pass_rid: ResourceId,
+    index_count: u32,
+    instance_count: u32,
+    first_index: u32,
+    base_vertex: i32,
+    first_instance: u32,
+}
+
+pub fn op_webgpu_render_pass_draw_indexed(
+    state: &mut OpState,
+    args: RenderPassDrawIndexedArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let render_pass_resource = state
+        .resource_table
+        .get::<WebGpuRenderPass>(args.render_pass_rid)?;
+
+    wgpu_core::command::render_ffi::wgpu_render_pass_draw_indexed(
+        &mut render_pass_resource.0.borrow_mut(),
+        args.index_count,
+        args.instance_count,
+        args.first_index,
+        args.base_vertex,
+        args.first_instance,
+    );
+
+    Ok(WebGpuResult::empty())
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct RenderPassDrawIndirectArgs {
+    render_pass_rid: ResourceId,
+    indirect_buffer: u32,
+    indirect_offset: u64,
+}
+
+pub fn op_webgpu_render_pass_draw_indirect(
+    state: &mut OpState,
+    args: RenderPassDrawIndirectArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let buffer_resource = state
+        .resource_table
+        .get::<super::buffer::WebGpuBuffer>(args.indirect_buffer)?;
+    let render_pass_resource = state
+        .resource_table
+        .get::<WebGpuRenderPass>(args.render_pass_rid)?;
+
+    wgpu_core::command::render_ffi::wgpu_render_pass_draw_indirect(
+        &mut render_pass_resource.0.borrow_mut(),
+        buffer_resource.0,
+        args.indirect_offset,
+    );
+
+    Ok(WebGpuResult::empty())
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct RenderPassDrawIndexedIndirectArgs {
+    render_pass_rid: ResourceId,
+    indirect_buffer: u32,
+    indirect_offset: u64,
+}
+
+pub fn op_webgpu_render_pass_draw_indexed_indirect(
+    state: &mut OpState,
+    args: RenderPassDrawIndexedIndirectArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let buffer_resource = state
+        .resource_table
+        .get::<super::buffer::WebGpuBuffer>(args.indirect_buffer)?;
+    let render_pass_resource = state
+        .resource_table
+        .get::<WebGpuRenderPass>(args.render_pass_rid)?;
+
+    wgpu_core::command::render_ffi::wgpu_render_pass_draw_indexed_indirect(
+        &mut render_pass_resource.0.borrow_mut(),
+        buffer_resource.0,
+        args.indirect_offset,
+    );
+
+    Ok(WebGpuResult::empty())
+}
--- a/deno_webgpu/sampler.rs
+++ b/deno_webgpu/sampler.rs
@ -0,0 +1,68 @@
+// Copyright 2018-2021 the Deno authors. All rights reserved. MIT license.
+
+use deno_core::error::AnyError;
+use deno_core::ResourceId;
+use deno_core::{OpState, Resource};
+use serde::Deserialize;
+use std::borrow::Cow;
+
+use super::error::WebGpuResult;
+
+pub(crate) struct WebGpuSampler(pub(crate) wgpu_core::id::SamplerId);
+impl Resource for WebGpuSampler {
+    fn name(&self) -> Cow<str> {
+        "webGPUSampler".into()
+    }
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct CreateSamplerArgs {
+    device_rid: ResourceId,
+    label: Option<String>,
+    address_mode_u: wgpu_types::AddressMode,
+    address_mode_v: wgpu_types::AddressMode,
+    address_mode_w: wgpu_types::AddressMode,
+    mag_filter: wgpu_types::FilterMode,
+    min_filter: wgpu_types::FilterMode,
+    mipmap_filter: wgpu_types::FilterMode,
+    lod_min_clamp: f32,
+    lod_max_clamp: f32,
+    compare: Option<wgpu_types::CompareFunction>,
+    max_anisotropy: u8,
+}
+
+pub fn op_webgpu_create_sampler(
+    state: &mut OpState,
+    args: CreateSamplerArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let instance = state.borrow::<super::Instance>();
+    let device_resource = state
+        .resource_table
+        .get::<super::WebGpuDevice>(args.device_rid)?;
+    let device = device_resource.0;
+
+    let descriptor = wgpu_core::resource::SamplerDescriptor {
+        label: args.label.map(Cow::from),
+        address_modes: [
+            args.address_mode_u,
+            args.address_mode_v,
+            args.address_mode_w,
+        ],
+        mag_filter: args.mag_filter,
+        min_filter: args.min_filter,
+        mipmap_filter: args.mipmap_filter,
+        lod_min_clamp: args.lod_min_clamp,
+        lod_max_clamp: args.lod_max_clamp,
+        compare: args.compare,
+        anisotropy_clamp: std::num::NonZeroU8::new(args.max_anisotropy),
+        border_color: None, // native-only
+    };
+
+    gfx_put!(device => instance.device_create_sampler(
+    device,
+    &descriptor,
+    std::marker::PhantomData
+  ) => state, WebGpuSampler)
+}
--- a/deno_webgpu/shader.rs
+++ b/deno_webgpu/shader.rs
@ -0,0 +1,50 @@
+// Copyright 2018-2021 the Deno authors. All rights reserved. MIT license.
+
+use deno_core::error::AnyError;
+use deno_core::ResourceId;
+use deno_core::{OpState, Resource};
+use serde::Deserialize;
+use std::borrow::Cow;
+
+use super::error::WebGpuResult;
+
+pub(crate) struct WebGpuShaderModule(pub(crate) wgpu_core::id::ShaderModuleId);
+impl Resource for WebGpuShaderModule {
+    fn name(&self) -> Cow<str> {
+        "webGPUShaderModule".into()
+    }
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct CreateShaderModuleArgs {
+    device_rid: ResourceId,
+    label: Option<String>,
+    code: String,
+    _source_map: Option<()>, // not yet implemented
+}
+
+pub fn op_webgpu_create_shader_module(
+    state: &mut OpState,
+    args: CreateShaderModuleArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let instance = state.borrow::<super::Instance>();
+    let device_resource = state
+        .resource_table
+        .get::<super::WebGpuDevice>(args.device_rid)?;
+    let device = device_resource.0;
+
+    let source = wgpu_core::pipeline::ShaderModuleSource::Wgsl(Cow::from(args.code));
+
+    let descriptor = wgpu_core::pipeline::ShaderModuleDescriptor {
+        label: args.label.map(Cow::from),
+    };
+
+    gfx_put!(device => instance.device_create_shader_module(
+    device,
+    &descriptor,
+    source,
+    std::marker::PhantomData
+  ) => state, WebGpuShaderModule)
+}
--- a/deno_webgpu/texture.rs
+++ b/deno_webgpu/texture.rs
@ -0,0 +1,108 @@
+// Copyright 2018-2021 the Deno authors. All rights reserved. MIT license.
+
+use deno_core::error::AnyError;
+use deno_core::ResourceId;
+use deno_core::{OpState, Resource};
+use serde::Deserialize;
+use std::borrow::Cow;
+
+use super::error::WebGpuResult;
+pub(crate) struct WebGpuTexture(pub(crate) wgpu_core::id::TextureId);
+impl Resource for WebGpuTexture {
+    fn name(&self) -> Cow<str> {
+        "webGPUTexture".into()
+    }
+}
+
+pub(crate) struct WebGpuTextureView(pub(crate) wgpu_core::id::TextureViewId);
+impl Resource for WebGpuTextureView {
+    fn name(&self) -> Cow<str> {
+        "webGPUTextureView".into()
+    }
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct CreateTextureArgs {
+    device_rid: ResourceId,
+    label: Option<String>,
+    size: wgpu_types::Extent3d,
+    mip_level_count: u32,
+    sample_count: u32,
+    dimension: wgpu_types::TextureDimension,
+    format: wgpu_types::TextureFormat,
+    usage: u32,
+}
+
+pub fn op_webgpu_create_texture(
+    state: &mut OpState,
+    args: CreateTextureArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let instance = state.borrow::<super::Instance>();
+    let device_resource = state
+        .resource_table
+        .get::<super::WebGpuDevice>(args.device_rid)?;
+    let device = device_resource.0;
+
+    let descriptor = wgpu_core::resource::TextureDescriptor {
+        label: args.label.map(Cow::from),
+        size: args.size,
+        mip_level_count: args.mip_level_count,
+        sample_count: args.sample_count,
+        dimension: args.dimension,
+        format: args.format,
+        usage: wgpu_types::TextureUsages::from_bits_truncate(args.usage),
+    };
+
+    gfx_put!(device => instance.device_create_texture(
+    device,
+    &descriptor,
+    std::marker::PhantomData
+  ) => state, WebGpuTexture)
+}
+
+#[derive(Deserialize)]
+#[serde(rename_all = "camelCase")]
+pub struct CreateTextureViewArgs {
+    texture_rid: ResourceId,
+    label: Option<String>,
+    format: Option<wgpu_types::TextureFormat>,
+    dimension: Option<wgpu_types::TextureViewDimension>,
+    aspect: wgpu_types::TextureAspect,
+    base_mip_level: u32,
+    mip_level_count: Option<u32>,
+    base_array_layer: u32,
+    array_layer_count: Option<u32>,
+}
+
+pub fn op_webgpu_create_texture_view(
+    state: &mut OpState,
+    args: CreateTextureViewArgs,
+    _: (),
+) -> Result<WebGpuResult, AnyError> {
+    let instance = state.borrow::<super::Instance>();
+    let texture_resource = state
+        .resource_table
+        .get::<WebGpuTexture>(args.texture_rid)?;
+    let texture = texture_resource.0;
+
+    let descriptor = wgpu_core::resource::TextureViewDescriptor {
+        label: args.label.map(Cow::from),
+        format: args.format,
+        dimension: args.dimension,
+        range: wgpu_types::ImageSubresourceRange {
+            aspect: args.aspect,
+            base_mip_level: args.base_mip_level,
+            mip_level_count: std::num::NonZeroU32::new(args.mip_level_count.unwrap_or(0)),
+            base_array_layer: args.base_array_layer,
+            array_layer_count: std::num::NonZeroU32::new(args.array_layer_count.unwrap_or(0)),
+        },
+    };
+
+    gfx_put!(texture => instance.texture_create_view(
+    texture,
+    &descriptor,
+    std::marker::PhantomData
+  ) => state, WebGpuTextureView)
+}
--- a/deno_webgpu/webgpu.idl
+++ b/deno_webgpu/webgpu.idl
--- a/wgpu-core/src/command/render.rs
+++ b/wgpu-core/src/command/render.rs
@ -55,6 +55,7 @@ pub enum LoadOp {
 #[derive(Copy, Clone, Debug, Hash, Eq, PartialEq)]
 #[cfg_attr(any(feature = "serial-pass", feature = "trace"), derive(Serialize))]
 #[cfg_attr(any(feature = "serial-pass", feature = "replay"), derive(Deserialize))]
+#[cfg_attr(feature = "serde", serde(rename_all = "kebab-case"))]
 pub enum StoreOp {
    /// Discards the content of the render target. If you don't care about the contents of the target, this can be faster.
    Discard = 0,
--- a/wgpu-core/src/pipeline.rs
+++ b/wgpu-core/src/pipeline.rs
@ -148,6 +148,7 @@ impl<A: hal::Api> Resource for ComputePipeline<A> {
 #[derive(Clone, Debug)]
 #[cfg_attr(feature = "trace", derive(serde::Serialize))]
 #[cfg_attr(feature = "replay", derive(serde::Deserialize))]
+#[cfg_attr(feature = "serde", serde(rename_all = "camelCase"))]
 pub struct VertexBufferLayout<'a> {
    /// The stride, in bytes, between elements of this buffer.
    pub array_stride: wgt::BufferAddress,
--- a/wgpu-types/Cargo.toml
+++ b/wgpu-types/Cargo.toml
@ -12,9 +12,10 @@ license = "MIT OR Apache-2.0"
 [lib]

 [features]
-trace = ["serde"]
-replay = ["serde"]
+trace = ["serde", "bitflags_serde_shim"]
+replay = ["serde", "bitflags_serde_shim"]

 [dependencies]
 bitflags = "1.0"
 serde = { version = "1.0", features = ["serde_derive"], optional = true }
+bitflags_serde_shim = { version = "0.2", optional = true }
--- a/wgpu-types/src/lib.rs
+++ b/wgpu-types/src/lib.rs
@ -93,8 +93,6 @@ impl Default for PowerPreference {
 bitflags::bitflags! {
    /// Represents the backends that wgpu will use.
    #[repr(transparent)]
-    #[cfg_attr(feature = "trace", derive(Serialize))]
-    #[cfg_attr(feature = "replay", derive(Deserialize))]
    pub struct Backends: u32 {
        /// Supported on Windows, Linux/Android, and macOS/iOS via Vulkan Portability (with the Vulkan feature enabled)
        const VULKAN = 1 << Backend::Vulkan as u32;
@ -123,6 +121,9 @@ bitflags::bitflags! {
    }
 }

+#[cfg(feature = "bitflags_serde_shim")]
+bitflags_serde_shim::impl_serde_for_bitflags!(Backends);
+
 impl From<Backend> for Backends {
    fn from(backend: Backend) -> Self {
        Self::from_bits(1 << backend as u32).unwrap()
@ -164,8 +165,6 @@ bitflags::bitflags! {
    /// will panic.
    #[repr(transparent)]
    #[derive(Default)]
-    #[cfg_attr(feature = "trace", derive(Serialize))]
-    #[cfg_attr(feature = "replay", derive(Deserialize))]
    pub struct Features: u64 {
        /// By default, polygon depth is clipped to 0-1 range. Anything outside of that range
        /// is rejected, and respective fragments are not touched.
@ -529,6 +528,9 @@ bitflags::bitflags! {
    }
 }

+#[cfg(feature = "bitflags_serde_shim")]
+bitflags_serde_shim::impl_serde_for_bitflags!(Features);
+
 impl Features {
    /// Mask of all features which are part of the upstream WebGPU standard.
    pub const fn all_webgpu_mask() -> Self {
@ -784,6 +786,9 @@ bitflags::bitflags! {
    }
 }

+#[cfg(feature = "bitflags_serde_shim")]
+bitflags_serde_shim::impl_serde_for_bitflags!(DownlevelFlags);
+
 impl DownlevelFlags {
    /// All flags that indicate if the backend is WebGPU compliant
    pub const fn compliant() -> Self {
@ -877,7 +882,6 @@ bitflags::bitflags! {
    ///
    /// `ShaderStages::VERTEX | ShaderStages::FRAGMENT`
    #[repr(transparent)]
-    #[cfg_attr(feature = "serde", derive(Deserialize, Serialize))]
    pub struct ShaderStages: u32 {
        /// Binding is not visible from any shader stage.
        const NONE = 0;
@ -892,6 +896,9 @@ bitflags::bitflags! {
    }
 }

+#[cfg(feature = "bitflags_serde_shim")]
+bitflags_serde_shim::impl_serde_for_bitflags!(ShaderStages);
+
 /// Dimensions of a particular texture view.
 #[repr(C)]
 #[derive(Copy, Clone, Debug, Hash, Eq, PartialEq)]
@ -1289,8 +1296,6 @@ impl Default for MultisampleState {
 bitflags::bitflags! {
    /// Feature flags for a texture format.
    #[repr(transparent)]
-    #[cfg_attr(feature = "trace", derive(Serialize))]
-    #[cfg_attr(feature = "replay", derive(Deserialize))]
    pub struct TextureFormatFeatureFlags: u32 {
        /// When used as a STORAGE texture, then a texture with this format can be bound with
        /// [`StorageTextureAccess::ReadOnly`] or [`StorageTextureAccess::ReadWrite`].
@ -1301,6 +1306,9 @@ bitflags::bitflags! {
    }
 }

+#[cfg(feature = "bitflags_serde_shim")]
+bitflags_serde_shim::impl_serde_for_bitflags!(TextureFormatFeatureFlags);
+
 /// Features supported by a given texture format
 ///
 /// Features are defined by WebGPU specification unless `Features::TEXTURE_ADAPTER_SPECIFIC_FORMAT_FEATURES` is enabled.
@ -1938,8 +1946,6 @@ impl TextureFormat {
 bitflags::bitflags! {
    /// Color write mask. Disabled color channels will not be written to.
    #[repr(transparent)]
-    #[cfg_attr(feature = "trace", derive(Serialize))]
-    #[cfg_attr(feature = "replay", derive(Deserialize))]
    pub struct ColorWrites: u32 {
        /// Enable red channel writes
        const RED = 1 << 0;
@ -1956,6 +1962,9 @@ bitflags::bitflags! {
    }
 }

+#[cfg(feature = "bitflags_serde_shim")]
+bitflags_serde_shim::impl_serde_for_bitflags!(ColorWrites);
+
 impl Default for ColorWrites {
    fn default() -> Self {
        Self::ALL
@ -2327,8 +2336,6 @@ bitflags::bitflags! {
    /// The usages determine what kind of memory the buffer is allocated from and what
    /// actions the buffer can partake in.
    #[repr(transparent)]
-    #[cfg_attr(feature = "trace", derive(Serialize))]
-    #[cfg_attr(feature = "replay", derive(Deserialize))]
    pub struct BufferUsages: u32 {
        /// Allow a buffer to be mapped for reading using [`Buffer::map_async`] + [`Buffer::get_mapped_range`].
        /// This does not include creating a buffer with [`BufferDescriptor::mapped_at_creation`] set.
@ -2361,6 +2368,9 @@ bitflags::bitflags! {
    }
 }

+#[cfg(feature = "bitflags_serde_shim")]
+bitflags_serde_shim::impl_serde_for_bitflags!(BufferUsages);
+
 /// Describes a [`Buffer`].
 #[repr(C)]
 #[derive(Clone, Debug, PartialEq, Eq, Hash)]
@ -2444,8 +2454,6 @@ bitflags::bitflags! {
    /// The usages determine what kind of memory the texture is allocated from and what
    /// actions the texture can partake in.
    #[repr(transparent)]
-    #[cfg_attr(feature = "trace", derive(Serialize))]
-    #[cfg_attr(feature = "replay", derive(Deserialize))]
    pub struct TextureUsages: u32 {
        /// Allows a texture to be the source in a [`CommandEncoder::copy_texture_to_buffer`] or
        /// [`CommandEncoder::copy_texture_to_texture`] operation.
@ -2462,6 +2470,9 @@ bitflags::bitflags! {
    }
 }

+#[cfg(feature = "bitflags_serde_shim")]
+bitflags_serde_shim::impl_serde_for_bitflags!(TextureUsages);
+
 /// Configures a [`Surface`] for presentation.
 #[repr(C)]
 #[derive(Clone, Debug, PartialEq, Eq, Hash)]
@ -3338,8 +3349,6 @@ bitflags::bitflags! {
    /// the first 8 bytes being the primitive out value, the last 8
    /// bytes being the compute shader invocation count.
    #[repr(transparent)]
-    #[cfg_attr(feature = "trace", derive(Serialize))]
-    #[cfg_attr(feature = "replay", derive(Deserialize))]
    pub struct PipelineStatisticsTypes : u8 {
        /// Amount of times the vertex shader is ran. Accounts for
        /// the vertex cache when doing indexed rendering.
@ -3361,6 +3370,9 @@ bitflags::bitflags! {
    }
 }

+#[cfg(feature = "bitflags_serde_shim")]
+bitflags_serde_shim::impl_serde_for_bitflags!(PipelineStatisticsTypes);
+
 /// Argument buffer layout for draw_indirect commands.
 #[repr(C)]
 #[derive(Clone, Copy, Debug)]