Moved add and bench into an examples directory so src is just the library

Otherwise a GpuBuffer of a single f16 is too small (only 2) and a eneven
f16 wouldn't work either

README.md

@@ -0,0 +1,83 @@
+# Minimal Zig WebGPU Compute Library
+
+This is a minimal, self-contained Zig library designed to simplify running compute shaders using WebGPU. It abstracts away much of the boilerplate required for GPU device initialization, memory management, and pipeline execution.
+
+## Core Modules
+
+The library exports five primary components:
+
+* **`GpuDevice`**: Initializes the WebGPU instance, adapter, device, and queue. It is configured to prioritize high performance and automatically requests the `ShaderF16` feature if the adapter supports it. By default, it enforces a 2 GB VRAM limit.
+* **`GpuArena` / `GpuAllocator`**: A memory management layer that tracks allocated VRAM bytes to prevent exceeding the device budget. The arena automatically destroys and releases all tracked WebGPU buffers when deinitialized.
+* **`GpuBuffer`**: Wraps native WebGPU buffers. It automatically aligns buffer sizes forward to a multiple of 4 bytes. It provides a `.load()` method for CPU-to-GPU data transfers (handling both aligned and unaligned lengths smoothly) and a `.read()` method that utilizes a staging buffer to map GPU data back to the CPU.
+* **`GpuProcess`**: Compiles WGSL source code into a compute pipeline. When running a process, it automatically splits the work into manageable chunks (up to 1 GB at a time) and dispatches workgroups of size 256.
+
+## Quick Start Example
+
+Below is a complete, self-contained example demonstrating how to initialize the GPU, load data, run a compute shader, and read the results back to the CPU:
+
+```zig
+
+const std = @import("std");
+const gpu = @import("gpu");
+const GpuDevice = gpu.GpuDevice;
+const GpuArena = gpu.GpuArena;
+const GpuBuffer = gpu.GpuBuffer;
+const GpuProcess = gpu.GpuProcess;
+
+pub fn main(init: std.process.Init) !void {
+    const allocator = init.gpa;
+
+    // 1. Open GPU Device
+    const device = try GpuDevice.init(.{});
+    defer device.deinit();
+
+    // 2. Create a GPU Arena to manage VRAM
+    var grena = GpuArena.init(allocator, device);
+    defer grena.deinit();
+    const gloc = grena.gpuAllocator();
+
+    // 3. Load the WGSL compute pipeline
+    const add_process = try GpuProcess.init(device, @embedFile("shaders/add.wgsl"));
+    defer add_process.deinit();
+
+    // 4. Setup CPU data
+    const len: usize = 16;
+    const data_a = try allocator.alloc(f16, len);
+    defer allocator.free(data_a);
+    const data_b = try allocator.alloc(f16, len);
+    defer allocator.free(data_b);
+
+    for (0..len) |i| {
+        data_a[i] = @floatFromInt(i);
+        data_b[i] = @floatFromInt(len - 1 - i);
+    }
+
+    // 5. Initialize raw GPU Buffers
+    // We pass the EnumSet inline using `.initMany` since the Enum itself isn't exported
+    const byte_size = len * @sizeOf(f16);
+    const buf_a = try GpuBuffer.init(gloc, byte_size, .initMany(&.{ .Storage, .CopyDst, .CopySrc }));
+    const buf_b = try GpuBuffer.init(gloc, byte_size, .initMany(&.{ .Storage, .CopyDst, .CopySrc }));
+    const buf_out = try GpuBuffer.init(gloc, byte_size, .initMany(&.{ .Storage, .CopyDst, .CopySrc }));
+
+    // Note: The buffers are safely tied to the GpuArena which will automatically
+    // release them at the end. You can also manually call buf_x.deinit() if desired.
+
+    // 6. Transfer data from CPU slices to GPU Buffers
+    try buf_a.load(f16, data_a);
+    try buf_b.load(f16, data_b);
+
+    // 7. Dispatch the Compute Process
+    // We pass the data type (f16) to allow GpuProcess to calculate chunks correctly
+    try add_process.run(gloc, f16, buf_a, buf_b, buf_out);
+
+    // 8. Map and copy the resulting buffer back to the CPU
+    const out = try buf_out.read(allocator, f16);
+    defer allocator.free(out);
+
+    std.debug.print("Result: {any}\n", .{out});
+}
+```
+
+## Dependencies
+
+* **`wgpu.h`**: The library relies on the WebGPU C API headers to bind to the native system graphics.

build.zig

@@ -1,61 +1,60 @@
-// build.zig
-// zig build run
-//
-// Expects wgpu-native pre-built in libs/wgpu-native/:
-//   include/wgpu.h
-//   lib/libwgpu_native.a   (or .so / .dylib / .dll)
-//
-// Download release: https://github.com/gfx-rs/wgpu-native/releases
-// Pick the archive matching your OS/arch, e.g.:
-//   wgpu-linux-x86_64-release.zip     → libwgpu_native.a  + wgpu.h
-//   wgpu-macos-aarch64-release.zip
-//   wgpu-windows-x86_64-msvc-release.zip
-
 const std = @import("std");
 
-pub fn build(b: *std.Build) void {
+pub fn build(b: *std.Build) !void {
     const target = b.standardTargetOptions(.{});
     const optimize = b.standardOptimizeOption(.{});
 
-    // 1. Define the module so other projects can import it
+    // Define the module so other projects can import it
-    _ = b.addModule("zig-wgpu", .{
+    const mod = b.addModule("zig-wgpu", .{
         .root_source_file = b.path("src/lib.zig"),
-    });
-
-    const exe = b.addExecutable(.{
-        .root_module = b.createModule(.{
-            .root_source_file = b.path("src/bench.zig"),
-            .link_libc = true,
         .target = target,
-            .optimize = optimize,
-        }),
-        .name = "gpu_matrix_add",
     });
 
-    // wgpu-native headers + pre-built static library
+    mod.addIncludePath(b.path("libs/wgpu-native/include"));
-    exe.root_module.addIncludePath(b.path("libs/wgpu-native/include"));
+    mod.addLibraryPath(b.path("libs/wgpu-native/lib"));
-    exe.root_module.addLibraryPath(b.path("libs/wgpu-native/lib"));
+    mod.addObjectFile(b.path("libs/wgpu-native/lib/libwgpu_native.a"));
-    exe.root_module.addObjectFile(b.path("libs/wgpu-native/lib/libwgpu_native.a"));
 
     // Platform-specific system frameworks needed by wgpu-native
     const t = target.result;
     if (t.os.tag == .macos) {
-        exe.root_module.linkFramework("Metal", .{});
+        mod.linkFramework("Metal", .{});
-        exe.root_module.linkFramework("QuartzCore", .{});
+        mod.linkFramework("QuartzCore", .{});
-        exe.root_module.linkFramework("Foundation", .{});
+        mod.linkFramework("Foundation", .{});
-        exe.root_module.linkFramework("CoreGraphics", .{});
+        mod.linkFramework("CoreGraphics", .{});
     } else if (t.os.tag == .windows) {
-        exe.root_module.linkSystemLibrary("d3d12", .{});
+        mod.linkSystemLibrary("d3d12", .{});
-        exe.root_module.linkSystemLibrary("dxgi", .{});
+        mod.linkSystemLibrary("dxgi", .{});
-        exe.root_module.linkSystemLibrary("user32", .{});
+        mod.linkSystemLibrary("user32", .{});
     } else {
-        exe.root_module.linkSystemLibrary("vulkan", .{});
+        mod.linkSystemLibrary("vulkan", .{});
-        exe.root_module.linkSystemLibrary("gcc_s", .{});
+        mod.linkSystemLibrary("gcc_s", .{});
     }
 
+    var threaded: std.Io.Threaded = .init_single_threaded;
+    const io = threaded.io();
+
+    var buf: [1024]u8 = undefined;
+    const exemples = try std.Io.Dir.cwd().openDir(io, "examples", .{ .access_sub_paths = false, .iterate = true });
+    var iter = exemples.iterate();
+    while (try iter.next(io)) |entry| {
+        if (entry.kind != .file) continue;
+        if (!std.mem.eql(u8, entry.name[entry.name.len - 4 ..], ".zig")) continue;
+
+        const exe = b.addExecutable(.{
+            .name = entry.name[0 .. entry.name.len - 4],
+            .root_module = b.createModule(.{
+                .root_source_file = b.path(try std.fmt.bufPrint(&buf, "examples/{s}", .{entry.name})),
+                .target = target,
+                .optimize = optimize,
+                .imports = &.{},
+            }),
+        });
+        exe.root_module.addImport("gpu", mod);
+
         b.installArtifact(exe);
 
-    const run = b.addRunArtifact(exe);
+        const run_step = b.step(entry.name[0 .. entry.name.len - 4], try std.fmt.bufPrint(&buf, "Run {s} demo", .{entry.name}));
-    run.step.dependOn(b.getInstallStep());
+        const run_cmd = b.addRunArtifact(exe);
-    b.step("bench", "Benchmark a simple add vector").dependOn(&run.step);
+        run_step.dependOn(&run_cmd.step);
+    }
 }

build.zig.zon

@@ -1,43 +1,13 @@
 .{
-    // This is the default name used by packages depending on this one. For
-    // example, when a user runs `zig fetch --save <url>`, this field is used
-    // as the key in the `dependencies` table. Although the user can choose a
-    // different name, most users will stick with this provided value.
-    //
-    // It is redundant to include "zig" in this name because it is already
-    // within the Zig package namespace.
     .name = .zig_wgpu,
-    // This is a [Semantic Version](https://semver.org/).
+    .version = "0.1.0",
-    // In a future version of Zig it will be used for package deduplication.
+    .fingerprint = 0x5d0e853acbc0c2c6,
-    .version = "0.0.0",
-    // Together with name, this represents a globally unique package
-    // identifier. This field is generated by the Zig toolchain when the
-    // package is first created, and then *never changes*. This allows
-    // unambiguous detection of one package being an updated version of
-    // another.
-    //
-    // When forking a Zig project, this id should be regenerated (delete the
-    // field and run `zig build`) if the upstream project is still maintained.
-    // Otherwise, the fork is *hostile*, attempting to take control over the
-    // original project's identity. Thus it is recommended to leave the comment
-    // on the following line intact, so that it shows up in code reviews that
-    // modify the field.
-    .fingerprint = 0x5d0e853acbc0c2c6, // Changing this has security and trust implications.
-    // Tracks the earliest Zig version that the package considers to be a
-    // supported use case.
     .minimum_zig_version = "0.16.0",
-    // This field is optional.
-    // Each dependency must either provide a `url` and `hash`, or a `path`.
-    // `zig build --fetch` can be used to fetch all dependencies of a package, recursively.
-    // Once all dependencies are fetched, `zig build` no longer requires
-    // internet connectivity.
     .dependencies = .{},
     .paths = .{
         "build.zig",
         "build.zig.zon",
         "src",
-        // For example...
+        "libs",
-        //"LICENSE",
-        //"README.md",
     },
 }

examples/add.zig

@@ -0,0 +1,59 @@
+const std = @import("std");
+const gpu = @import("gpu");
+const GpuDevice = gpu.GpuDevice;
+const GpuArena = gpu.GpuArena;
+const GpuBuffer = gpu.GpuBuffer;
+const GpuProcess = gpu.GpuProcess;
+
+pub fn main(init: std.process.Init) !void {
+    const allocator = init.gpa;
+
+    // 1. Open GPU Device
+    const device = try GpuDevice.init(.{});
+    defer device.deinit();
+
+    // 2. Create a GPU Arena to manage VRAM
+    var grena = GpuArena.init(allocator, device);
+    defer grena.deinit();
+    const gloc = grena.gpuAllocator();
+
+    // 3. Load the WGSL compute pipeline
+    const add_process = try GpuProcess.init(device, @embedFile("shaders/add.wgsl"));
+    defer add_process.deinit();
+
+    // 4. Setup CPU data
+    const len: usize = 16;
+    const data_a = try allocator.alloc(f16, len);
+    defer allocator.free(data_a);
+    const data_b = try allocator.alloc(f16, len);
+    defer allocator.free(data_b);
+
+    for (0..len) |i| {
+        data_a[i] = @floatFromInt(i);
+        data_b[i] = @floatFromInt(len - 1 - i);
+    }
+
+    // 5. Initialize raw GPU Buffers
+    // We pass the EnumSet inline using `.initMany` since the Enum itself isn't exported
+    const byte_size = len * @sizeOf(f16);
+    const buf_a = try GpuBuffer.init(gloc, byte_size, .initMany(&.{ .Storage, .CopyDst, .CopySrc }));
+    const buf_b = try GpuBuffer.init(gloc, byte_size, .initMany(&.{ .Storage, .CopyDst, .CopySrc }));
+    const buf_out = try GpuBuffer.init(gloc, byte_size, .initMany(&.{ .Storage, .CopyDst, .CopySrc }));
+
+    // Note: The buffers are safely tied to the GpuArena which will automatically
+    // release them at the end. You can also manually call buf_x.deinit() if desired.
+
+    // 6. Transfer data from CPU slices to GPU Buffers
+    try buf_a.load(f16, data_a);
+    try buf_b.load(f16, data_b);
+
+    // 7. Dispatch the Compute Process
+    // We pass the data type (f16) to allow GpuProcess to calculate chunks correctly
+    try add_process.run(gloc, f16, buf_a, buf_b, buf_out);
+
+    // 8. Map and copy the resulting buffer back to the CPU
+    const out = try buf_out.read(allocator, f16);
+    defer allocator.free(out);
+
+    std.debug.print("Result: {any}\n", .{out});
+}

examples/bench.zig

@@ -1,37 +1,75 @@
 const std = @import("std");
-const GpuDevice = @import("GpuDevice.zig");
+const gpu = @import("gpu");
-const GpuAllocator = @import("GpuAllocator.zig");
+const GpuDevice = gpu.GpuDevice;
-const GpuPipeline = @import("GpuPipeline.zig");
+const GpuArena = gpu.GpuArena;
-const Vec = @import("Vec.zig");
+const GpuAllocator = gpu.GpuAllocator;
+const GpuBuffer = gpu.GpuBuffer;
+const GpuProcess = gpu.GpuProcess;
 
-const c = @import("c.zig").c;
+/// Minimal implementation of a f16 Vector
+const Vec = struct {
+    buf: GpuBuffer,
+    len: usize,
+
+    // Changed: gloc is passed by value (const)
+    pub fn initZero(gloc: GpuAllocator, len: usize) !Vec {
+        return .{
+            .buf = try GpuBuffer.init(
+                gloc,
+                len * @sizeOf(f16),
+                .initMany(&.{ .Storage, .CopyDst, .CopySrc }),
+            ),
+            .len = len,
+        };
+    }
+
+    // Changed: gloc is passed by value
+    pub fn initLoad(gloc: GpuAllocator, data: []const f16) !Vec {
+        var self = try initZero(gloc, data.len);
+        try self.load(data); // Direct access via the interface copy
+        return self;
+    }
+
+    pub fn deinit(self: Vec) void {
+        self.buf.deinit();
+    }
+
+    /// CPU to GPU.
+    pub fn load(self: Vec, data: []const f16) !void {
+        try self.buf.load(f16, data);
+    }
+
+    // Changed: gloc is passed by value instead of *GpuAllocator
+    pub fn run(self: Vec, gloc: GpuAllocator, other: Vec, process: GpuProcess) !Vec {
+        std.debug.assert(self.len == other.len);
+
+        const result = try Vec.initZero(gloc, self.len);
+        errdefer result.deinit();
+
+        try process.run(gloc, f16, self.buf, other.buf, result.buf);
+        return result;
+    }
+
+    // Changed: gloc is passed by value instead of *GpuAllocator
+    pub fn read(self: Vec, alloc: std.mem.Allocator) ![]f16 {
+        return self.buf.read(alloc, f16);
+    }
+};
 
 pub fn main(init: std.process.Init) !void {
-    const device = try GpuDevice.init();
+    const device = try GpuDevice.init(.{ .vram_bytes_limit = 4 * 1024 * 1024 * 1024 });
     defer device.deinit();
 
-    var gloc = try GpuAllocator.init(init.gpa, device);
+    var grena = GpuArena.init(init.gpa, device);
-    defer gloc.deinit();
+    defer grena.deinit();
 
-    const add_pip = try GpuPipeline.init(device, @embedFile("shaders/add.wgsl"));
+    const gloc = grena.gpuAllocator();
+
+    const add_pip = try GpuProcess.init(device, @embedFile("shaders/add.wgsl"));
     defer add_pip.deinit();
 
     const allocator = init.gpa;
 
-    // --- WARM-UP PHASE ---
-    {
-        var warmup_a = [_]f32{1.0};
-        var warmup_b = [_]f32{1.0};
-        const wa = try Vec.initLoad(&gloc, &warmup_a);
-        defer wa.deinit();
-        const wb = try Vec.initLoad(&gloc, &warmup_b);
-        defer wb.deinit();
-        const wsum = try wa.run(&gloc, wb, add_pip);
-        defer wsum.deinit();
-        const wout = try wsum.read(&gloc, allocator);
-        defer allocator.free(wout);
-    }
-
     const sizes = [_]usize{
         1,
         256,
@@ -44,8 +82,8 @@ pub fn main(init: std.process.Init) !void {
         4 * 1024 * 1024,
         4 * 4 * 1024 * 1024,
         4 * 4 * 4 * 1024 * 1024,
-        4 * 4 * 4 * 4 * 1024 * 1024,
+        // 4 * 4 * 4 * 4 * 1024 * 1024,
-        4 * 4 * 4 * 4 * 4 * 1024 * 1024,
+        // 4 * 4 * 4 * 4 * 4 * 1024 * 1024,
     };
 
     const iterations = 10;
@@ -56,9 +94,9 @@ pub fn main(init: std.process.Init) !void {
 
     for (sizes) |size| {
         // --- Phase 1: Host Init/Alloc (Outside the iteration loop for pure host prep) ---
-        var data_a = try allocator.alloc(f32, size);
+        const data_a = try allocator.alloc(f16, size);
         defer allocator.free(data_a);
-        var data_b = try allocator.alloc(f32, size);
+        const data_b = try allocator.alloc(f16, size);
         defer allocator.free(data_b);
 
         for (0..size) |i| {
@@ -72,15 +110,15 @@ pub fn main(init: std.process.Init) !void {
         var min_compute_ns: u64 = std.math.maxInt(u64);
 
         // Track peak VRAM usage observed during the iterations
-        var peak_vram_bytes: usize = 0;
+        var peak_vram_bytes: u64 = 0;
 
         for (0..iterations) |_| {
             // --- 1. GPU ALLOCATION PHASE ---
             const alloc_start = std.Io.Clock.awake.now(init.io);
 
-            const a = try Vec.initLoad(&gloc, data_a);
+            const a = try Vec.initLoad(gloc, data_a);
             defer a.deinit();
-            const b = try Vec.initLoad(&gloc, data_b);
+            const b = try Vec.initLoad(gloc, data_b);
             defer b.deinit();
 
             const alloc_duration = alloc_start.durationTo(std.Io.Clock.awake.now(init.io));
@@ -90,16 +128,15 @@ pub fn main(init: std.process.Init) !void {
             // --- 2. COMPUTE PHASE ---
             const compute_start = std.Io.Clock.awake.now(init.io);
 
-            const sum = try a.run(&gloc, b, add_pip);
+            const sum = try a.run(gloc, b, add_pip);
             defer sum.deinit();
 
             // All 3 buffers (a, b, sum) are currently resident in VRAM here.
             // Querying now catches the true peak allocation step.
-            if (gloc.allocated_vram_bytes > peak_vram_bytes) {
+            if (grena.allocated_vram_bytes > peak_vram_bytes)
-                peak_vram_bytes = gloc.allocated_vram_bytes;
+                peak_vram_bytes = grena.allocated_vram_bytes;
-            }
 
-            _ = c.wgpuDevicePoll(device.device, 1, null);
+            device.poll();
 
             const compute_duration = compute_start.durationTo(std.Io.Clock.awake.now(init.io));
             const compute_ns = @as(u64, @intCast(compute_duration.toNanoseconds()));
@@ -108,7 +145,7 @@ pub fn main(init: std.process.Init) !void {
             // --- 3. TRANSFER PHASE (Device -> Host) ---
             const transfer_start = std.Io.Clock.awake.now(init.io);
 
-            const out = try sum.read(&gloc, allocator);
+            const out = try sum.read(allocator);
             defer allocator.free(out);
 
             const transfer_duration = transfer_start.durationTo(std.Io.Clock.awake.now(init.io));
@@ -118,7 +155,7 @@ pub fn main(init: std.process.Init) !void {
 
         // --- Metrics Calculations ---
         const f_size = @as(f64, @floatFromInt(size));
-        const element_bytes = f_size * @as(f64, @floatFromInt(@sizeOf(f32)));
+        const element_bytes = f_size * @as(f64, @floatFromInt(@sizeOf(f16)));
         const mb = element_bytes / (1024.0 * 1024.0);
 
         // Individual Phase Timings (ms)

examples/shaders/add.wgsl

@@ -1,11 +1,9 @@
-@group(0) @binding(0) var<storage, read> A: array<f32>;
+enable f16;
-@group(0) @binding(1) var<storage, read> B: array<f32>;
-@group(0) @binding(2) var<storage, read_write> C: array<f32>;
 
-struct TensorInfo {
+@group(0) @binding(0) var<storage, read> A: array<f16>;
-    size: u32,
+@group(0) @binding(1) var<storage, read> B: array<f16>;
-};
+@group(0) @binding(2) var<storage, read_write> C: array<f16>;
-@group(0) @binding(3) var<uniform> info: TensorInfo; 
+@group(0) @binding(3) var<uniform> size: u32; 
 
 @compute @workgroup_size(256)
 fn main(
@@ -19,7 +17,7 @@ fn main(
     var index = global_id.x;
     
     // 3. Stride through the tensor elements
-    while (index < info.size) {
+    while (index < size) {
         C[index] = A[index] + B[index];
         index += total_threads; // Jump forward by the total thread count
     }

src/GpuAllocator.zig

@@ -1,55 +1,24 @@
+// GpuAllocator.zig
 const std = @import("std");
 const GpuDevice = @import("GpuDevice.zig");
-const GpuBuffer = @import("GpuBuffer.zig");
+const c = @import("utils.zig").c;
-const c = @import("c.zig").c;
+
+const GpuAllocator = @This();
+
+/// The function definitions our underlying implementations must satisfy
+pub const VTable = struct {
+    alloc: *const fn (ctx: *anyopaque, bytes: u64, usage: c.WGPUBufferUsage) anyerror!c.WGPUBuffer,
+    free: *const fn (ctx: *anyopaque, buf_raw: c.WGPUBuffer, size: u64) void,
+};
 
 device: GpuDevice,
-tracked_buffers: std.AutoHashMap(c.WGPUBuffer, void),
+ptr: *anyopaque,
-allocated_vram_bytes: u64 = 0,
+vtable: *const VTable,
 
-pub fn init(cpu_allocator: std.mem.Allocator, device: GpuDevice) !@This() {
+pub fn allocBuffer(self: GpuAllocator, bytes: u64, usage: c.WGPUBufferUsage) !c.WGPUBuffer {
-    return .{
+    return self.vtable.alloc(self.ptr, bytes, usage);
-        .device = device,
-        .tracked_buffers = .init(cpu_allocator),
-    };
 }
 
-pub fn deinit(self: *@This()) void {
+pub fn freeBuffer(self: GpuAllocator, buf_raw: c.WGPUBuffer, size: u64) void {
-    var it = self.tracked_buffers.keyIterator();
+    self.vtable.free(self.ptr, buf_raw, size);
-    while (it.next()) |buf_ptr| {
-        const buf = buf_ptr.*;
-        c.wgpuBufferDestroy(buf);
-        c.wgpuBufferRelease(buf);
-    }
-    self.tracked_buffers.deinit();
-}
-
-pub fn registerBuffer(
-    self: *@This(),
-    bytes: u64,
-    usage: c.WGPUBufferUsage,
-) !c.WGPUBuffer {
-    if (bytes > self.device.limits.maxBufferSize)
-        return error.SingleBufferExceedsLimit;
-
-    if (bytes + self.allocated_vram_bytes > self.device.config.vram_bytes_limit)
-        return error.ExceedsVramBudget;
-
-    const buf = c.wgpuDeviceCreateBuffer(self.device.device, &.{
-        .usage = usage,
-        .size = bytes,
-    }) orelse return error.BufferAlloc;
-
-    try self.tracked_buffers.put(buf, {});
-    self.allocated_vram_bytes += bytes;
-    return buf;
-}
-
-pub fn unregisterAndDestroyBuffer(self: *@This(), buf: GpuBuffer) void {
-    if (self.tracked_buffers.remove(buf.raw)) {
-        c.wgpuBufferDestroy(buf.raw);
-        c.wgpuBufferRelease(buf.raw);
-        self.allocated_vram_bytes -= buf.size;
-        self.device.poll();
-    }
 }

src/GpuArena.zig

@@ -0,0 +1,72 @@
+// GpuArena.zig
+const std = @import("std");
+const GpuDevice = @import("GpuDevice.zig");
+const GpuAllocator = @import("GpuAllocator.zig");
+const c = @import("utils.zig").c;
+
+const GpuArena = @This();
+
+device: GpuDevice,
+tracked_buffers: std.AutoHashMap(c.WGPUBuffer, void),
+allocated_vram_bytes: u64 = 0,
+
+pub fn init(cpu_allocator: std.mem.Allocator, device: GpuDevice) GpuArena {
+    return .{
+        .device = device,
+        .tracked_buffers = .init(cpu_allocator),
+    };
+}
+
+pub fn deinit(self: *GpuArena) void {
+    var it = self.tracked_buffers.keyIterator();
+    while (it.next()) |buf_ptr| {
+        c.wgpuBufferDestroy(buf_ptr.*);
+        c.wgpuBufferRelease(buf_ptr.*);
+    }
+    self.tracked_buffers.deinit();
+}
+
+/// Returns the type-erased immutable interface wrapper
+pub fn gpuAllocator(self: *GpuArena) GpuAllocator {
+    return .{
+        .device = self.device,
+        .ptr = self,
+        .vtable = &.{
+            .alloc = alloc,
+            .free = free,
+        },
+    };
+}
+
+fn alloc(ctx: *anyopaque, bytes: u64, usage: c.WGPUBufferUsage) anyerror!c.WGPUBuffer {
+    const self: *GpuArena = @ptrCast(@alignCast(ctx));
+
+    if (bytes > self.device.limits.maxBufferSize)
+        return error.SingleBufferExceedsLimit;
+
+    if (bytes + self.allocated_vram_bytes > self.device.config.vram_bytes_limit)
+        return error.ExceedsVramBudget;
+
+    const buf = c.wgpuDeviceCreateBuffer(self.device.device, &.{
+        .usage = usage,
+        .size = bytes,
+    }) orelse return error.BufferAlloc;
+    errdefer {
+        c.wgpuBufferDestroy(buf);
+        c.wgpuBufferRelease(buf);
+    }
+
+    try self.tracked_buffers.put(buf, {});
+    self.allocated_vram_bytes += bytes;
+    return buf;
+}
+
+fn free(ctx: *anyopaque, buf_raw: c.WGPUBuffer, size: u64) void {
+    const self: *GpuArena = @ptrCast(@alignCast(ctx));
+
+    if (self.tracked_buffers.remove(buf_raw)) {
+        c.wgpuBufferDestroy(buf_raw);
+        c.wgpuBufferRelease(buf_raw);
+        self.allocated_vram_bytes -= size;
+    }
+}

src/GpuBuffer.zig

@@ -1,26 +1,52 @@
 const std = @import("std");
-const c = @import("c.zig").c;
+const c = @import("utils.zig").c;
 const GpuAllocator = @import("GpuAllocator.zig");
 
 raw: c.WGPUBuffer,
-size: u64,
+size: u64, // Now tracks the 4-byte aligned size directly
 usage: c.WGPUBufferUsage,
-gloc: *GpuAllocator,
+gloc: GpuAllocator,
+
+const BufferUsage = enum(u64) {
+    None = 0x0000000000000000,
+    MapRead = 0x0000000000000001,
+    MapWrite = 0x0000000000000002,
+    CopySrc = 0x0000000000000004,
+    CopyDst = 0x0000000000000008,
+    Index = 0x0000000000000010,
+    Vertex = 0x0000000000000020,
+    Uniform = 0x0000000000000040,
+    Storage = 0x0000000000000080,
+    Indirect = 0x0000000000000100,
+    QueryResolve = 0x0000000000000200,
+};
 
 /// Allocates the underlying WebGPU handle and registers it to the parent GpuAllocator
-pub fn init(gloc: *GpuAllocator, bytes: u64, usage: c.WGPUBufferUsage) !@This() {
+pub fn init(gloc: GpuAllocator, size: u64, usage: std.EnumSet(BufferUsage)) !@This() {
-    const raw_handle = try gloc.registerBuffer(bytes, usage);
+    var use: u64 = 0;
+    var iter = usage.iterator();
+    while (iter.next()) |flag| use |= @intFromEnum(flag);
+
+    // Automatically align the buffer size forward to a multiple of 4 bytes under the hood
+    const aligned_size = std.mem.alignForward(u64, size, 4);
+
+    const raw_handle = try gloc.allocBuffer(aligned_size, use);
     return .{
         .raw = raw_handle,
-        .size = bytes,
+        .size = aligned_size, // Expose the aligned size to the rest of the application
-        .usage = usage,
+        .usage = use,
         .gloc = gloc,
     };
 }
 
 /// Unregisters from the parent GpuAllocator and cleanly destroys GPU resources
 pub fn deinit(self: @This()) void {
-    self.gloc.unregisterAndDestroyBuffer(self);
+    self.gloc.freeBuffer(self.raw, self.size);
+}
+
+/// Native getConstMappedRange wrapper
+pub fn getConstMappedRange(self: @This(), offset: u64, size: u64) ?*const anyopaque {
+    return c.wgpuBufferGetConstMappedRange(self.raw, offset, size);
 }
 
 /// Native mapAsync wrapper
@@ -34,12 +60,79 @@ pub fn mapAsync(
     _ = c.wgpuBufferMapAsync(self.raw, mode, offset, size, callback_info);
 }
 
-/// Native getConstMappedRange wrapper
-pub fn getConstMappedRange(self: @This(), offset: u64, size: u64) ?*const anyopaque {
-    return c.wgpuBufferGetConstMappedRange(self.raw, offset, size);
-}
-
 /// Native unmap wrapper
 pub fn unmap(self: @This()) void {
     c.wgpuBufferUnmap(self.raw);
 }
+
+/// CPU to GPU.
+pub fn load(
+    self: @This(),
+    T: type,
+    data: []const T,
+) !void {
+    const bytes = data.len * @sizeOf(T);
+
+    if (bytes == self.size) {
+        // Aligned path: direct download
+        c.wgpuQueueWriteBuffer(self.gloc.device.queue, self.raw, 0, data.ptr, self.size);
+    } else {
+        // Unaligned path: Split the write into an aligned chunk and a padded remainder
+        // to support arbitrary lengths without any allocations or large stack arrays.
+        const aligned_part = (bytes / 4) * 4;
+        if (aligned_part > 0) {
+            c.wgpuQueueWriteBuffer(self.gloc.device.queue, self.raw, 0, data.ptr, aligned_part);
+        }
+
+        var remainder_buf: [4]u8 = .{ 0, 0, 0, 0 };
+        const data_bytes = std.mem.sliceAsBytes(data);
+        @memcpy(remainder_buf[0 .. bytes - aligned_part], data_bytes[aligned_part..bytes]);
+
+        c.wgpuQueueWriteBuffer(self.gloc.device.queue, self.raw, aligned_part, &remainder_buf, 4);
+    }
+}
+
+pub fn read(self: @This(), alloc: std.mem.Allocator, T: type) ![]T {
+    const out = try alloc.alloc(T, @divExact(self.size, @sizeOf(T)));
+
+    const staging = try init(
+        self.gloc,
+        self.size,
+        .initMany(&.{ .MapRead, .CopyDst }),
+    );
+    defer staging.deinit();
+
+    const enc = c.wgpuDeviceCreateCommandEncoder(self.gloc.device.device, null) orelse return error.Encoder;
+    c.wgpuCommandEncoderCopyBufferToBuffer(enc, self.raw, 0, staging.raw, 0, self.size);
+    const cmd = c.wgpuCommandEncoderFinish(enc, null);
+    defer c.wgpuCommandEncoderRelease(enc);
+    defer c.wgpuCommandBufferRelease(cmd);
+    c.wgpuQueueSubmit(self.gloc.device.queue, 1, &cmd);
+
+    var mapped = false;
+    staging.mapAsync(
+        c.WGPUMapMode_Read,
+        0,
+        self.size,
+        .{ .callback = onMapped, .userdata1 = &mapped },
+    );
+    while (!mapped) self.gloc.device.poll();
+
+    const ptr: [*]const T = @ptrCast(@alignCast(
+        staging.getConstMappedRange(0, self.size),
+    ));
+    @memcpy(out[0..out.len], ptr[0..out.len]);
+    staging.unmap();
+
+    return out;
+}
+
+fn onMapped(
+    status: c.WGPUMapAsyncStatus,
+    _: c.WGPUStringView,
+    userdata1: ?*anyopaque,
+    _: ?*anyopaque,
+) callconv(.c) void {
+    const flag: *bool = @ptrCast(@alignCast(userdata1.?));
+    flag.* = (status == c.WGPUMapAsyncStatus_Success);
+}

src/GpuDevice.zig

@@ -1,22 +1,32 @@
 const std = @import("std");
-const c = @import("c.zig").c;
+const c = @import("utils.zig").c;
+const sv = @import("utils.zig").sv;
 
 const Ctx = struct {
     adapter: c.WGPUAdapter = null,
     device: c.WGPUDevice = null,
 };
 
+const GpuDeviceConfig = struct {
+    /// VRAM limit. Default 2 GB
+    vram_bytes_limit: u64 = 2 * 1024 * 1024 * 1024,
+    power_preference: enum(c_uint) {
+        Undefined = 0x00000000,
+        LowPower = 0x00000001,
+        HighPerformance = 0x00000002,
+        Force32 = 0x7FFFFFFF,
+    } = .HighPerformance,
+};
+
 instance: c.WGPUInstance,
 adapter: c.WGPUAdapter,
 device: c.WGPUDevice,
 queue: c.WGPUQueue,
 limits: c.WGPULimits,
 
-config: struct {
+config: GpuDeviceConfig,
-    vram_bytes_limit: u64 = 10 * 1024 * 1024 * 1024, // 10 GB
-} = .{},
 
-pub fn init() !@This() {
+pub fn init(config: GpuDeviceConfig) !@This() {
     const instance = c.wgpuCreateInstance(
         &std.mem.zeroes(c.WGPUInstanceDescriptor),
     ) orelse return error.NoInstance;
@@ -25,28 +35,39 @@ pub fn init() !@This() {
     var ctx = Ctx{};
     _ = c.wgpuInstanceRequestAdapter(
         instance,
-        &.{ .powerPreference = c.WGPUPowerPreference_HighPerformance },
+        &.{ .powerPreference = @intFromEnum(config.power_preference) },
         .{ .callback = onAdapter, .userdata1 = &ctx },
     );
     c.wgpuInstanceProcessEvents(instance);
     const adapter = ctx.adapter orelse return error.NoAdapter;
     errdefer c.wgpuAdapterRelease(adapter);
 
-    // --- QUERY HARDWARE LIMITS ---
+    var supported_features = std.mem.zeroes(c.WGPUSupportedFeatures);
+    c.wgpuAdapterGetFeatures(adapter, &supported_features);
+
     var supported_limits = std.mem.zeroes(c.WGPULimits);
     supported_limits.nextInChain = null;
-
-    // Fetch what your physical graphic card can actually handle
     if (c.wgpuAdapterGetLimits(adapter, &supported_limits) != 1) return error.FailedToGetAdapterLimits;
 
+    var has_f16 = false;
+    for (0..supported_features.featureCount) |i| {
+        if (supported_features.features[i] == c.WGPUFeatureName_ShaderF16) {
+            has_f16 = true;
+            break;
+        }
+    }
+
+    var feature_buf = [_]c.WGPUFeatureName{c.WGPUFeatureName_ShaderF16};
+    const required_features: []const c.WGPUFeatureName =
+        if (has_f16) feature_buf[0..1] else &.{};
+
     const device_descriptor = c.WGPUDeviceDescriptor{
         .nextInChain = null,
         .label = sv("TensorCompilerDevice"),
-        .requiredFeatureCount = 0,
+        .requiredFeatureCount = required_features.len,
-        .requiredFeatures = null,
+        .requiredFeatures = if (required_features.len > 0) required_features.ptr else null,
         .requiredLimits = &supported_limits,
     };
-
     _ = c.wgpuAdapterRequestDevice(
         adapter,
         &device_descriptor,
@@ -61,6 +82,7 @@ pub fn init() !@This() {
         .device = device,
         .queue = c.wgpuDeviceGetQueue(device),
         .limits = supported_limits,
+        .config = config,
     };
 }
 
@@ -71,7 +93,8 @@ pub fn deinit(self: @This()) void {
     c.wgpuInstanceRelease(self.instance);
 }
 
-pub fn poll(self: *@This()) void {
+/// Wait for thing to be done
+pub fn poll(self: @This()) void {
     _ = c.wgpuDevicePoll(self.device, 1, null);
 }
 
@@ -104,7 +127,3 @@ fn onDevice(
     const ctx: *Ctx = @ptrCast(@alignCast(userdata1.?));
     ctx.device = device;
 }
-
-fn sv(s: []const u8) c.WGPUStringView {
-    return .{ .data = s.ptr, .length = s.len };
-}

src/GpuPipeline.zig

@@ -1,28 +0,0 @@
-const std = @import("std");
-const GpuDevice = @import("GpuDevice.zig");
-const c = @import("c.zig").c;
-
-raw: c.WGPUComputePipeline,
-
-pub fn init(device: GpuDevice, wgsl: []const u8) !@This() {
-    var wgsl_src = c.WGPUShaderSourceWGSL{
-        .chain = .{ .sType = c.WGPUSType_ShaderSourceWGSL },
-        .code = sv(wgsl),
-    };
-    const shader = c.wgpuDeviceCreateShaderModule(device.device, &.{
-        .nextInChain = @ptrCast(&wgsl_src),
-    }) orelse return error.Shader;
-    defer c.wgpuShaderModuleRelease(shader);
-
-    return .{ .raw = c.wgpuDeviceCreateComputePipeline(device.device, &.{
-        .compute = .{ .module = shader, .entryPoint = sv("main") },
-    }) orelse return error.Pipeline };
-}
-
-pub fn deinit(self: @This()) void {
-    c.wgpuComputePipelineRelease(self.raw);
-}
-
-fn sv(s: []const u8) c.WGPUStringView {
-    return .{ .data = s.ptr, .length = s.len };
-}

src/GpuProcess.zig

@@ -0,0 +1,122 @@
+/// GpuProcess is just a pipeline with 2 inpout and 1 output
+/// for now, to see if I make it a bit more generic
+///
+const std = @import("std");
+const c = @import("utils.zig").c;
+const sv = @import("utils.zig").sv;
+const GpuAllocator = @import("GpuAllocator.zig");
+const GpuBuffer = @import("GpuBuffer.zig");
+const GpuDevice = @import("GpuDevice.zig");
+
+pip: c.WGPUComputePipeline,
+
+pub fn init(device: GpuDevice, wgsl: []const u8) !@This() {
+    var wgsl_src = c.WGPUShaderSourceWGSL{
+        .chain = .{ .sType = c.WGPUSType_ShaderSourceWGSL },
+        .code = sv(wgsl),
+    };
+    const shader = c.wgpuDeviceCreateShaderModule(device.device, &.{
+        .nextInChain = @ptrCast(&wgsl_src),
+    }) orelse return error.Shader;
+    defer c.wgpuShaderModuleRelease(shader);
+
+    return .{ .pip = c.wgpuDeviceCreateComputePipeline(device.device, &.{
+        .compute = .{ .module = shader, .entryPoint = sv("main") },
+    }) orelse return error.Pipeline };
+}
+
+pub fn deinit(self: @This()) void {
+    c.wgpuComputePipelineRelease(self.pip);
+}
+
+fn onMapped(
+    status: c.WGPUMapAsyncStatus,
+    _: c.WGPUStringView,
+    userdata1: ?*anyopaque,
+    _: ?*anyopaque,
+) callconv(.c) void {
+    const flag: *bool = @ptrCast(@alignCast(userdata1.?));
+    flag.* = (status == c.WGPUMapAsyncStatus_Success);
+}
+
+// Changed: gloc is passed by value instead of *GpuAllocator
+pub fn run(
+    self: @This(),
+    gloc: GpuAllocator,
+    T: type,
+    buf_a: GpuBuffer,
+    buf_b: GpuBuffer,
+    buf_out: GpuBuffer,
+) !void {
+    const max_chunk_bytes: u64 = 1024 * 1024 * 1024; // 1 GB
+
+    const bytes = buf_a.size;
+    var offset: u64 = 0;
+    while (offset < bytes) {
+        const current_chunk_bytes = @min(max_chunk_bytes, bytes - offset);
+        const current_chunk_elements: u32 = @intCast(current_chunk_bytes / @sizeOf(T));
+
+        const info_buf = try GpuBuffer.init(
+            gloc,
+            @sizeOf(u32),
+            .initMany(&.{ .Uniform, .CopyDst }),
+        );
+        defer info_buf.deinit();
+
+        c.wgpuQueueWriteBuffer(gloc.device.queue, info_buf.raw, 0, &current_chunk_elements, @sizeOf(u32));
+
+        const entries = [_]c.WGPUBindGroupEntry{
+            .{ .binding = 0, .buffer = buf_a.raw, .offset = offset, .size = current_chunk_bytes },
+            .{ .binding = 1, .buffer = buf_b.raw, .offset = offset, .size = current_chunk_bytes },
+            .{ .binding = 2, .buffer = buf_out.raw, .offset = offset, .size = current_chunk_bytes },
+            .{ .binding = 3, .buffer = info_buf.raw, .offset = 0, .size = @sizeOf(u32) },
+        };
+
+        try submitPass(gloc, self.pip, &entries, current_chunk_elements);
+
+        offset += current_chunk_bytes;
+    }
+}
+
+// Changed: gloc is passed by value instead of *GpuAllocator
+fn submitPass(
+    gloc: GpuAllocator,
+    pipeline: c.WGPUComputePipeline,
+    entries: []const c.WGPUBindGroupEntry,
+    n: usize,
+) !void {
+    const bgl = c.wgpuComputePipelineGetBindGroupLayout(pipeline, 0);
+    defer c.wgpuBindGroupLayoutRelease(bgl);
+
+    const bg = c.wgpuDeviceCreateBindGroup(gloc.device.device, &.{
+        .layout = bgl,
+        .entries = entries.ptr,
+        .entryCount = entries.len,
+    }) orelse return error.BindGroup;
+    defer c.wgpuBindGroupRelease(bg);
+
+    const enc = c.wgpuDeviceCreateCommandEncoder(gloc.device.device, null) orelse
+        return error.Encoder;
+    const pass = c.wgpuCommandEncoderBeginComputePass(enc, null);
+    c.wgpuComputePassEncoderSetPipeline(pass, pipeline);
+    c.wgpuComputePassEncoderSetBindGroup(pass, 0, bg, 0, null);
+
+    const WORKGROUP_SIZE = 256;
+    const MAX_WORKGROUPS = 65535;
+
+    const desired_workgroups = ceilDiv(n, WORKGROUP_SIZE);
+    const dispatch_count = @min(desired_workgroups, MAX_WORKGROUPS);
+
+    c.wgpuComputePassEncoderDispatchWorkgroups(pass, @intCast(dispatch_count), 1, 1);
+    c.wgpuComputePassEncoderEnd(pass);
+    c.wgpuComputePassEncoderRelease(pass);
+
+    const cmd = c.wgpuCommandEncoderFinish(enc, null);
+    defer c.wgpuCommandEncoderRelease(enc);
+    defer c.wgpuCommandBufferRelease(cmd);
+    c.wgpuQueueSubmit(gloc.device.queue, 1, &cmd);
+}
+
+fn ceilDiv(n: usize, d: usize) usize {
+    return (n + d - 1) / d;
+}

src/Vec.zig

@@ -1,192 +0,0 @@
-/// Dummy
-const std = @import("std");
-const c = @import("c.zig").c;
-const GpuAllocator = @import("GpuAllocator.zig");
-const GpuBuffer = @import("GpuBuffer.zig");
-const GpuDevice = @import("GpuDevice.zig");
-const GpuPipeline = @import("GpuPipeline.zig");
-
-const Vec = @This();
-
-buf: GpuBuffer,
-len: usize,
-
-pub fn initZero(gloc: *GpuAllocator, len: usize) !Vec {
-    return .{
-        .buf = try GpuBuffer.init(
-            gloc,
-            len * @sizeOf(f32),
-            c.WGPUBufferUsage_Storage | c.WGPUBufferUsage_CopyDst | c.WGPUBufferUsage_CopySrc,
-        ),
-        .len = len,
-    };
-}
-
-pub fn initLoad(gloc: *GpuAllocator, data: []const f32) !Vec {
-    var self = try initZero(gloc, data.len);
-    try self.load(gloc.device, data);
-    return self;
-}
-
-pub fn deinit(self: Vec) void {
-    self.buf.deinit();
-}
-
-/// CPU to GPU.
-pub fn load(
-    self: Vec,
-    device: GpuDevice,
-    data: []const f32,
-) !void {
-    std.debug.assert(data.len == self.len);
-    const bytes = data.len * @sizeOf(f32);
-    c.wgpuQueueWriteBuffer(device.queue, self.buf.raw, 0, data.ptr, bytes);
-}
-
-pub fn byteSize(self: Vec) u64 {
-    return @as(u64, self.len) * @sizeOf(f32);
-}
-
-pub fn run(self: Vec, gloc: *GpuAllocator, other: Vec, pip: GpuPipeline) !Vec {
-    std.debug.assert(self.len == other.len);
-
-    const result = try Vec.initZero(gloc, self.len);
-    errdefer result.deinit();
-
-    try dispatch2in1out(gloc, pip.raw, self.buf, other.buf, result.buf, self.byteSize());
-
-    return result;
-}
-
-/// GPU to CPU.
-pub fn read(self: Vec, gloc: *GpuAllocator, alloc: std.mem.Allocator) ![]f32 {
-    const out = try alloc.alloc(f32, self.len);
-    const bytes = self.byteSize();
-
-    const staging = try GpuBuffer.init(
-        gloc,
-        bytes,
-        c.WGPUBufferUsage_MapRead | c.WGPUBufferUsage_CopyDst,
-    );
-    defer staging.deinit();
-
-    const enc = c.wgpuDeviceCreateCommandEncoder(gloc.device.device, null) orelse return error.Encoder;
-    c.wgpuCommandEncoderCopyBufferToBuffer(enc, self.buf.raw, 0, staging.raw, 0, bytes);
-    const cmd = c.wgpuCommandEncoderFinish(enc, null);
-    defer c.wgpuCommandEncoderRelease(enc);
-    defer c.wgpuCommandBufferRelease(cmd);
-    c.wgpuQueueSubmit(gloc.device.queue, 1, &cmd);
-
-    var mapped = false;
-    staging.mapAsync(
-        c.WGPUMapMode_Read,
-        0,
-        bytes,
-        .{ .callback = onMapped, .userdata1 = &mapped },
-    );
-    while (!mapped) gloc.device.poll();
-
-    const ptr: [*]const f32 = @ptrCast(@alignCast(
-        staging.getConstMappedRange(0, bytes),
-    ));
-    @memcpy(out[0..self.len], ptr[0..self.len]);
-    staging.unmap();
-
-    return out;
-}
-
-fn onMapped(
-    status: c.WGPUMapAsyncStatus,
-    _: c.WGPUStringView,
-    userdata1: ?*anyopaque,
-    _: ?*anyopaque,
-) callconv(.c) void {
-    const flag: *bool = @ptrCast(@alignCast(userdata1.?));
-    flag.* = (status == c.WGPUMapAsyncStatus_Success);
-}
-
-/// Encode + submit a 2-input, 1-output compute pass (used by add).
-fn dispatch2in1out(
-    gloc: *GpuAllocator,
-    pipeline: c.WGPUComputePipeline,
-    buf_a: GpuBuffer,
-    buf_b: GpuBuffer,
-    buf_out: GpuBuffer,
-    bytes: u64,
-) !void {
-    const max_chunk_bytes: u64 = 1024 * 1024 * 1024; // 1 GB
-
-    var offset: u64 = 0;
-    while (offset < bytes) {
-        // Calculate bounds for the current chunk
-        const current_chunk_bytes = @min(max_chunk_bytes, bytes - offset);
-        const current_chunk_elements: u32 = @intCast(current_chunk_bytes / @sizeOf(f32));
-
-        // Create uniform buffer for this specific chunk's size
-        const info_buf = try GpuBuffer.init(
-            gloc,
-            @sizeOf(u32),
-            c.WGPUBufferUsage_Uniform | c.WGPUBufferUsage_CopyDst,
-        );
-        defer info_buf.deinit();
-
-        // Write the number of elements *in this chunk* to the uniform buffer
-        c.wgpuQueueWriteBuffer(gloc.device.queue, info_buf.raw, 0, &current_chunk_elements, @sizeOf(u32));
-
-        // Bind only the sub-slice for this chunk using `.offset` and `.size`
-        const entries = [_]c.WGPUBindGroupEntry{
-            .{ .binding = 0, .buffer = buf_a.raw, .offset = offset, .size = current_chunk_bytes },
-            .{ .binding = 1, .buffer = buf_b.raw, .offset = offset, .size = current_chunk_bytes },
-            .{ .binding = 2, .buffer = buf_out.raw, .offset = offset, .size = current_chunk_bytes },
-            .{ .binding = 3, .buffer = info_buf.raw, .offset = 0, .size = @sizeOf(u32) },
-        };
-
-        // Submit the pass for this specific chunk
-        try submitPass(gloc, pipeline, &entries, current_chunk_elements);
-
-        offset += current_chunk_bytes;
-    }
-}
-
-/// Create bind group, encode pass, submit.
-fn submitPass(
-    gloc: *GpuAllocator,
-    pipeline: c.WGPUComputePipeline,
-    entries: []const c.WGPUBindGroupEntry,
-    n: usize,
-) !void {
-    const bgl = c.wgpuComputePipelineGetBindGroupLayout(pipeline, 0);
-    defer c.wgpuBindGroupLayoutRelease(bgl);
-
-    const bg = c.wgpuDeviceCreateBindGroup(gloc.device.device, &.{
-        .layout = bgl,
-        .entries = entries.ptr,
-        .entryCount = entries.len,
-    }) orelse return error.BindGroup;
-    defer c.wgpuBindGroupRelease(bg);
-
-    const enc = c.wgpuDeviceCreateCommandEncoder(gloc.device.device, null) orelse
-        return error.Encoder;
-    const pass = c.wgpuCommandEncoderBeginComputePass(enc, null);
-    c.wgpuComputePassEncoderSetPipeline(pass, pipeline);
-    c.wgpuComputePassEncoderSetBindGroup(pass, 0, bg, 0, null);
-
-    const WORKGROUP_SIZE = 256;
-    const MAX_WORKGROUPS = 65535;
-
-    const desired_workgroups = ceilDiv(n, WORKGROUP_SIZE);
-    const dispatch_count = @min(desired_workgroups, MAX_WORKGROUPS);
-
-    c.wgpuComputePassEncoderDispatchWorkgroups(pass, @intCast(dispatch_count), 1, 1);
-    c.wgpuComputePassEncoderEnd(pass);
-    c.wgpuComputePassEncoderRelease(pass);
-
-    const cmd = c.wgpuCommandEncoderFinish(enc, null);
-    defer c.wgpuCommandEncoderRelease(enc);
-    defer c.wgpuCommandBufferRelease(cmd);
-    c.wgpuQueueSubmit(gloc.device.queue, 1, &cmd);
-}
-
-fn ceilDiv(n: usize, d: usize) usize {
-    return (n + d - 1) / d;
-}

src/c.zig

@@ -1 +0,0 @@
-pub const c = @cImport(@cInclude("wgpu.h"));

src/lib.zig

@@ -1,4 +1,5 @@
 pub const GpuAllocator = @import("GpuAllocator.zig");
+pub const GpuArena = @import("GpuArena.zig");
 pub const GpuBuffer = @import("GpuBuffer.zig");
 pub const GpuDevice = @import("GpuDevice.zig");
-pub const GpuPipeline = @import("GpuPipeline.zig");
+pub const GpuProcess = @import("GpuProcess.zig");

src/utils.zig

@@ -0,0 +1,5 @@
+pub const c = @cImport(@cInclude("wgpu.h"));
+
+pub fn sv(s: []const u8) c.WGPUStringView {
+    return .{ .data = s.ptr, .length = s.len };
+}