feat: nonblocking buffer readback

CHANGELOG.md

@@ -2,6 +2,18 @@
 
 _Disclaimer: this changelog is updated using generative AI, but is still verified manually._
 
+## Unreleased
+
+### Added
+- **Non-blocking GPU→CPU readback.** A new `GpuReadback` type stages a buffer copy and returns immediately: `request`/`request_copy` start the transfer and `try_take` polls for completion (`is_idle` reports whether a readback is in flight). `GpuTimestamps` gained matching `request_read`/`try_take`/`is_idle` for non-blocking timestamp readback. Both are driven by a new `Backend::poll` maintenance method (a non-blocking counterpart to `synchronize`) implemented across the WebGPU, CUDA, Metal, and CPU backends. The `metal` feature now pulls in `block` for command-buffer completion handlers.
+- `WebGpu::from_device`: build a WebGPU backend on top of an already-created wgpu instance/adapter/device/queue instead of creating its own.
+- `from_wgpu` constructors on `GpuBufferSlice`, `GpuBufferSliceMut`, and `WebGpuBufferSlice` to wrap a foreign `wgpu::Buffer` (owned by another library sharing the device) as a khal buffer slice.
+- `force_cpu_coroutines` option for `#[spirv_bindgen]`: forces the CPU backend to dispatch a kernel through the coroutine scheduler so `barrier_wait()` synchronizes, even for kernels with no `#[spirv(workgroup)]` parameter. Replaces the old dummy-workgroup-parameter hack.
+- The CUDA backend's `load_module_bytes` now accepts a pre-linked CUBIN (detected via the ELF magic) in addition to PTX text, for modules referencing symbols the driver JIT cannot resolve on its own (e.g. libdevice `__nv_*` math).
+
+### Fixed
+- WebGPU entry-point lookup on wasm now mirrors naga's identifier sanitization, which appends a trailing `_` to names ending in a digit, so kernels like `reduce_add_f32` resolve correctly.
+
 ## v0.2.0
 
 ### Added

crates/khal-derive/src/spirv_bindgen/cpu.rs

@@ -222,6 +222,7 @@ pub(super) fn generate_cpu_dispatch_block(
     original_params: &[OriginalParam],
     workgroup_size: [u32; 3],
     func_ident: &syn::Ident,
+    force_coroutines: bool,
 ) -> proc_macro2::TokenStream {
     let wg_x = workgroup_size[0];
     let wg_y = workgroup_size[1];
@@ -250,7 +251,10 @@ pub(super) fn generate_cpu_dispatch_block(
     if cfg_variant_sets.is_empty() {
         // No cfg-gated params: single variant with all params.
         let all_params: Vec<&OriginalParam> = original_params.iter().collect();
-        let body = generate_variant_body(&all_params, workgroup_size);
+        let mut body = generate_variant_body(&all_params, workgroup_size);
+        // `force_cpu_coroutines`: force the coroutine path even with no workgroup param,
+        // so `barrier_wait()` actually synchronizes.
+        body.has_workgroup |= force_coroutines;
         let dispatch_loop = gen_dispatch_loop(&body, workgroup_size, func_ident);
 
         quote! {
@@ -284,7 +288,8 @@ pub(super) fn generate_cpu_dispatch_block(
                     })
                     .collect();
 
-                let body = generate_variant_body(&active_params, workgroup_size);
+                let mut body = generate_variant_body(&active_params, workgroup_size);
+                body.has_workgroup |= force_coroutines;
                 let dispatch_loop = gen_dispatch_loop(&body, workgroup_size, func_ident);
 
                 quote! {

crates/khal-derive/src/spirv_bindgen/mod.rs

@@ -344,7 +344,17 @@ pub(crate) fn spirv_bindgen(attr: TokenStream, item: TokenStream) -> TokenStream
     //   #[spirv_bindgen(MyName)]
     //   #[spirv_bindgen(spirv_passthrough)]
     //   #[spirv_bindgen(MyName, spirv_passthrough)]
+    //   #[spirv_bindgen(force_cpu_coroutines)]            // force CPU coroutine dispatch
+    //   #[spirv_bindgen(MyName, force_cpu_coroutines)]
     let mut spirv_passthrough = false;
+    // When set, the CPU backend always dispatches this kernel through the
+    // coroutine scheduler (the path that makes `barrier_wait()` actually
+    // synchronize), even when the kernel declares no `#[spirv(workgroup)]`
+    // parameter. Use for kernels that rely on barriers but exchange data through
+    // storage buffers rather than shared memory — otherwise their threads run
+    // sequentially with no-op barriers on the CPU backend. Replaces the old
+    // "dummy `#[spirv(workgroup)] _cpu_marker` parameter" hack.
+    let mut force_cpu_coroutines = false;
     let struct_name: syn::Ident = if attr.is_empty() {
         // No explicit name provided - derive from function name (snake_case to PascalCase)
         let func_name = func.sig.ident.to_string();
@@ -358,6 +368,8 @@ pub(crate) fn spirv_bindgen(attr: TokenStream, item: TokenStream) -> TokenStream
         for ident in &args {
             if ident == "spirv_passthrough" {
                 spirv_passthrough = true;
+            } else if ident == "force_cpu_coroutines" {
+                force_cpu_coroutines = true;
             } else {
                 if name.is_some() {
                     return syn::Error::new_spanned(
@@ -618,8 +630,12 @@ pub(crate) fn spirv_bindgen(attr: TokenStream, item: TokenStream) -> TokenStream
 
     // ── CPU dispatch code generation ─────────────────────────────────────
     let func_ident = &func.sig.ident;
-    let cpu_dispatch_block =
-        cpu::generate_cpu_dispatch_block(&original_params, workgroup_size, func_ident);
+    let cpu_dispatch_block = cpu::generate_cpu_dispatch_block(
+        &original_params,
+        workgroup_size,
+        func_ident,
+        force_cpu_coroutines,
+    );
 
     // ── CUDA (nvptx64) kernel entry point generation ─────────────────────
     let func_name_str = func.sig.ident.to_string();
@@ -743,8 +759,20 @@ pub(crate) fn spirv_bindgen(attr: TokenStream, item: TokenStream) -> TokenStream
                     format!("{}::{}", module, entry_point)
                 };
 
+                // On wasm/WebGPU, wgpu transpiles the SPIR-V to WGSL via naga and matches
+                // the requested entry point against the generated WGSL function name. naga's
+                // Namer sanitizes identifiers: `::` -> `_`, and additionally appends a
+                // trailing `_` when the name ends in a digit (it reserves that suffix for its
+                // own numeric disambiguation). Mirror both rules so the lookup matches, e.g.
+                // `linalg::reduce::reduce_add_f32` -> `linalg_reduce_reduce_add_f32_`.
                 #[cfg(target_arch = "wasm32")]
-                let full_entry = full_entry.replace("::", "_");
+                let full_entry = {
+                    let mut full_entry = full_entry.replace("::", "_");
+                    if matches!(full_entry.bytes().last(), Some(b'0'..=b'9')) {
+                        full_entry.push('_');
+                    }
+                    full_entry
+                };
 
                 Self::from_bytes(backend, file.contents(), &full_entry)
             }

crates/khal/Cargo.toml

@@ -13,7 +13,7 @@ derive = ["khal-derive"]
 webgpu = ["khal-derive?/webgpu"]
 cpu = ["khal-derive?/cpu"]
 cuda = ["dep:cudarc", "khal-derive?/cuda"]
-metal = ["dep:metal", "dep:naga", "khal-derive?/webgpu"]
+metal = ["dep:metal", "dep:naga", "dep:block", "khal-derive?/webgpu"]
 push_constants = []
 subgroup_ops = []
 
@@ -40,6 +40,9 @@ paste = "1"
 [target.'cfg(target_os = "macos")'.dependencies]
 metal = { version = "0.32", optional = true }
 naga = { version = "29", optional = true, features = ["spv-in", "msl-out"] }
+# Objective-C block support for Metal command-buffer completion handlers
+# (non-blocking readback completion). Matches metal-rs's own `block` dep.
+block = { version = "0.1.6", optional = true }
 
 [lints.rust]
 # The `objc` crate's `msg_send!` macro (used by the Metal backend) expands to