Renamed ProcessDef to COmputeDef

Now GpuProcess isnt limited to 2 in, 1 out but to anything.

.gitignore

@@ -1,2 +1,3 @@
 .zig-cache
 zig-out
+examples/mnist

README.md

@@ -1,6 +1,7 @@
 # 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.
+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
 
@@ -9,20 +10,19 @@ 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.
+* **`GpuCompute`**: Compiles WGSL source code into a compute pipeline. When running, 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;
+const GpuCompute = gpu.GpuCompute;
 
 pub fn main(init: std.process.Init) !void {
     const allocator = init.gpa;
@@ -37,8 +37,16 @@ pub fn main(init: std.process.Init) !void {
     const gloc = grena.gpuAllocator();
 
     // 3. Load the WGSL compute pipeline
-    const add_process = try GpuProcess.init(device, @embedFile("shaders/add.wgsl"));
+    const add_cp = try GpuCompute.init(
-    defer add_process.deinit();
+        device,
+        @embedFile("shaders/add.wgsl"),
+        .{ .bindings = &.{
+            .{ .element_size = @sizeOf(f16) },
+            .{ .element_size = @sizeOf(f16) },
+            .{ .element_size = @sizeOf(f16) },
+        } },
+    );
+    defer add_cp.deinit();
 
     // 4. Setup CPU data
     const len: usize = 16;
@@ -53,7 +61,6 @@ pub fn main(init: std.process.Init) !void {
     }
 
     // 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 }));
@@ -66,9 +73,8 @@ pub fn main(init: std.process.Init) !void {
     try buf_a.load(f16, data_a);
     try buf_b.load(f16, data_b);
 
-    // 7. Dispatch the Compute Process
+    // 7. Dispatch the Compute
-    // We pass the data type (f16) to allow GpuProcess to calculate chunks correctly
+    try add_cp.run(gloc, .{ buf_a, buf_b, buf_out });
-    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);
@@ -81,3 +87,55 @@ pub fn main(init: std.process.Init) !void {
 ## Dependencies
 
 * **`wgpu.h`**: The library relies on the WebGPU C API headers to bind to the native system graphics.
+
+## System Requirements
+
+Because this library binds to native system graphics APIs via `wgpu-native`,
+you must ensure the appropriate development headers and libraries are available on your system before compiling.
+
+### Linux (Vulkan)
+You need the Vulkan development headers to compile the project, and a Vulkan-compatible driver to run it. 
+
+Depending on your distribution, install the following packages:
+* **Ubuntu / Debian:**
+  ```bash
+  sudo apt update
+  sudo apt install libvulkan-dev mesa-vulkan-drivers
+  ```
+* **Fedora / RHEL:**
+  ```bash
+  sudo dnf install vulkan-devel mesa-vulkan-drivers
+  ```
+* **Arch Linux:**
+  ```bash
+  sudo pacman -S vulkan-headers vulkan-icd-loader
+  ```
+
+### macOS (Metal)
+No extra installation is required.
+The build script automatically links against the standard Apple frameworks provided by the Xcode Command Line Tools
+(`Metal`, `QuartzCore`, `Foundation`, `CoreGraphics`).
+
+### Windows (DirectX 12)
+No extra installation is required.
+The build script automatically links against the standard Windows SDK libraries (`d3d12`, `dxgi`, `user32`).
+Ensure you have the MSVC build tools installed.
+
+---
+
+## Adding to your project
+
+To use this module in your own Zig project, add it to your `build.zig.zon`:
+```bash
+zig fetch --save git+https://git.bouvais.lu/adrien/zig-wgpu
+```
+
+Then, in your `build.zig`, import and add the module to your executable:
+```zig
+const zig_wgpu = b.dependency("zig-wgpu", .{
+    .target = target,
+    .optimize = optimize,
+});
+
+exe.root_module.addImport("gpu", zig_wgpu.module("zig-wgpu"));
+```

build.zig

@@ -4,7 +4,6 @@ pub fn build(b: *std.Build) !void {
     const target = b.standardTargetOptions(.{});
     const optimize = b.standardOptimizeOption(.{});
 
-    // Define the module so other projects can import it
     const mod = b.addModule("zig-wgpu", .{
         .root_source_file = b.path("src/lib.zig"),
         .target = target,

examples/add.zig

@@ -3,7 +3,7 @@ const gpu = @import("gpu");
 const GpuDevice = gpu.GpuDevice;
 const GpuArena = gpu.GpuArena;
 const GpuBuffer = gpu.GpuBuffer;
-const GpuProcess = gpu.GpuProcess;
+const GpuCompute = gpu.GpuCompute;
 
 pub fn main(init: std.process.Init) !void {
     const allocator = init.gpa;
@@ -18,8 +18,16 @@ pub fn main(init: std.process.Init) !void {
     const gloc = grena.gpuAllocator();
 
     // 3. Load the WGSL compute pipeline
-    const add_process = try GpuProcess.init(device, @embedFile("shaders/add.wgsl"));
+    const add_cp = try GpuCompute.init(
-    defer add_process.deinit();
+        device,
+        @embedFile("shaders/add.wgsl"),
+        .{ .bindings = &.{
+            .{ .element_size = @sizeOf(f16) },
+            .{ .element_size = @sizeOf(f16) },
+            .{ .element_size = @sizeOf(f16) },
+        } },
+    );
+    defer add_cp.deinit();
 
     // 4. Setup CPU data
     const len: usize = 16;
@@ -34,7 +42,6 @@ pub fn main(init: std.process.Init) !void {
     }
 
     // 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 }));
@@ -47,9 +54,8 @@ pub fn main(init: std.process.Init) !void {
     try buf_a.load(f16, data_a);
     try buf_b.load(f16, data_b);
 
-    // 7. Dispatch the Compute Process
+    // 7. Dispatch the Compute
-    // We pass the data type (f16) to allow GpuProcess to calculate chunks correctly
+    try add_cp.run(gloc, .{ buf_a, buf_b, buf_out });
-    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);

examples/bench.zig

@@ -4,7 +4,7 @@ const GpuDevice = gpu.GpuDevice;
 const GpuArena = gpu.GpuArena;
 const GpuAllocator = gpu.GpuAllocator;
 const GpuBuffer = gpu.GpuBuffer;
-const GpuProcess = gpu.GpuProcess;
+const GpuCompute = gpu.GpuCompute;
 
 /// Minimal implementation of a f16 Vector
 const Vec = struct {
@@ -40,13 +40,13 @@ const Vec = struct {
     }
 
     // Changed: gloc is passed by value instead of *GpuAllocator
-    pub fn run(self: Vec, gloc: GpuAllocator, other: Vec, process: GpuProcess) !Vec {
+    pub fn run(self: Vec, gloc: GpuAllocator, other: Vec, process: GpuCompute) !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);
+        try process.run(gloc, .{ self.buf, other.buf, result.buf });
         return result;
     }
 
@@ -65,7 +65,11 @@ pub fn main(init: std.process.Init) !void {
 
     const gloc = grena.gpuAllocator();
 
-    const add_pip = try GpuProcess.init(device, @embedFile("shaders/add.wgsl"));
+    const add_pip = try GpuCompute.init(device, @embedFile("shaders/add.wgsl"), .{ .bindings = &.{
+        .{ .element_size = @sizeOf(f16) },
+        .{ .element_size = @sizeOf(f16) },
+        .{ .element_size = @sizeOf(f16) },
+    } });
     defer add_pip.deinit();
 
     const allocator = init.gpa;

examples/digit.zig

@@ -0,0 +1,77 @@
+// I am using this mnist reduced dataset https://www.kaggle.com/datasets/mohamedgamal07/reduced-mnist
+
+const std = @import("std");
+const gpu = @import("gpu");
+const GpuDevice = gpu.GpuDevice;
+const GpuArena = gpu.GpuArena;
+const GpuBuffer = gpu.GpuBuffer;
+const GpuProcess = gpu.GpuProcess;
+
+const BATCHSIZE = 10;
+const EPOCH = 10;
+
+pub fn main(init: std.process.Init) !void {
+    const allocator = init.gpa;
+    const io = init.io;
+
+    // 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();
+
+    var train_dir = try std.Io.Dir.cwd().openDir(io, "mnist/train", .{});
+
+    var images: [BATCHSIZE * 28 * 28]f16 = undefined;
+    for (EPOCH) |epoch| {
+        // Load random images from train dir
+        train_dir.openDir(io, "0", .{});
+        for (BATCHSIZE) |i| {
+            const file = try train_dir.openFile(io, "0.jpg", .{});
+            images[28 * 28 * i .. 28 * 28 * (i + 1)] = file.read
+        }
+    }
+
+    // 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});
+}

src/GpuAllocator.zig

@@ -1,11 +1,6 @@
-// GpuAllocator.zig
-const std = @import("std");
 const GpuDevice = @import("GpuDevice.zig");
 const c = @import("utils.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,
@@ -15,10 +10,10 @@ device: GpuDevice,
 ptr: *anyopaque,
 vtable: *const VTable,
 
-pub fn allocBuffer(self: GpuAllocator, bytes: u64, usage: c.WGPUBufferUsage) !c.WGPUBuffer {
+pub fn allocBuffer(self: @This(), bytes: u64, usage: c.WGPUBufferUsage) !c.WGPUBuffer {
     return self.vtable.alloc(self.ptr, bytes, usage);
 }
 
-pub fn freeBuffer(self: GpuAllocator, buf_raw: c.WGPUBuffer, size: u64) void {
+pub fn freeBuffer(self: @This(), buf_raw: c.WGPUBuffer, size: u64) void {
     self.vtable.free(self.ptr, buf_raw, size);
 }

src/GpuArena.zig

@@ -1,23 +1,20 @@
-// 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 {
+pub fn init(cpu_allocator: std.mem.Allocator, device: GpuDevice) @This() {
     return .{
         .device = device,
         .tracked_buffers = .init(cpu_allocator),
     };
 }
 
-pub fn deinit(self: *GpuArena) void {
+pub fn deinit(self: *@This()) void {
     var it = self.tracked_buffers.keyIterator();
     while (it.next()) |buf_ptr| {
         c.wgpuBufferDestroy(buf_ptr.*);
@@ -27,7 +24,7 @@ pub fn deinit(self: *GpuArena) void {
 }
 
 /// Returns the type-erased immutable interface wrapper
-pub fn gpuAllocator(self: *GpuArena) GpuAllocator {
+pub fn gpuAllocator(self: *@This()) GpuAllocator {
     return .{
         .device = self.device,
         .ptr = self,
@@ -39,7 +36,7 @@ pub fn gpuAllocator(self: *GpuArena) GpuAllocator {
 }
 
 fn alloc(ctx: *anyopaque, bytes: u64, usage: c.WGPUBufferUsage) anyerror!c.WGPUBuffer {
-    const self: *GpuArena = @ptrCast(@alignCast(ctx));
+    const self: *@This() = @ptrCast(@alignCast(ctx));
 
     if (bytes > self.device.limits.maxBufferSize)
         return error.SingleBufferExceedsLimit;
@@ -62,7 +59,7 @@ fn alloc(ctx: *anyopaque, bytes: u64, usage: c.WGPUBufferUsage) anyerror!c.WGPUB
 }
 
 fn free(ctx: *anyopaque, buf_raw: c.WGPUBuffer, size: u64) void {
-    const self: *GpuArena = @ptrCast(@alignCast(ctx));
+    const self: *@This() = @ptrCast(@alignCast(ctx));
 
     if (self.tracked_buffers.remove(buf_raw)) {
         c.wgpuBufferDestroy(buf_raw);

src/GpuBuffer.zig

@@ -3,7 +3,7 @@ const c = @import("utils.zig").c;
 const GpuAllocator = @import("GpuAllocator.zig");
 
 raw: c.WGPUBuffer,
-size: u64, // Now tracks the 4-byte aligned size directly
+size: u64,
 usage: c.WGPUBufferUsage,
 gloc: GpuAllocator,
 
@@ -33,7 +33,7 @@ pub fn init(gloc: GpuAllocator, size: u64, usage: std.EnumSet(BufferUsage)) !@Th
     const raw_handle = try gloc.allocBuffer(aligned_size, use);
     return .{
         .raw = raw_handle,
-        .size = aligned_size, // Expose the aligned size to the rest of the application
+        .size = aligned_size,
         .usage = use,
         .gloc = gloc,
     };

src/GpuCompute.zig

@@ -0,0 +1,171 @@
+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");
+
+pub const Binding = struct {
+    /// Element size in bytes for this binding. E.g. @sizeOf(f32).
+    /// If 0, no element-based size validation is performed for this buffer.
+    element_size: u32 = 0,
+};
+
+pub const ComputeDef = struct {
+    bindings: []const Binding,
+    workgroup_size: u32 = 256,
+    max_workgroups: u32 = 65535,
+    /// If true, automatically adds a Uniform Buffer containing `elements_count` as a `u32`
+    /// to the next available binding slot.
+    append_info_buffer: bool = true,
+};
+
+pip: c.WGPUComputePipeline,
+def: ComputeDef,
+
+pub fn init(device: GpuDevice, wgsl: []const u8, def: ComputeDef) !@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);
+
+    const pip = c.wgpuDeviceCreateComputePipeline(device.device, &.{
+        .compute = .{ .module = shader, .entryPoint = sv("main") },
+    }) orelse return error.Pipeline;
+
+    return .{
+        .pip = pip,
+        .def = def,
+    };
+}
+
+pub fn deinit(self: @This()) void {
+    c.wgpuComputePipelineRelease(self.pip);
+}
+
+/// Execute the compute pass with arbitrary buffer bindings via a tuple.
+/// Example: `try proc.run(gloc, .{ buf_a, buf_b, buf_out });`
+pub fn run(
+    self: @This(),
+    gloc: GpuAllocator,
+    args: anytype,
+) !void {
+    const type_info = @typeInfo(@TypeOf(args));
+    if (type_info != .@"struct" or !type_info.@"struct".is_tuple)
+        @compileError("Expected a tuple of GpuBuffers for args. E.g. .{ buf_a, buf_b }");
+
+    const fields = type_info.@"struct".fields;
+    if (fields.len != self.def.bindings.len)
+        return error.InvalidArgumentCount;
+
+    var elements_count: u32 = 0;
+
+    // Infer elements_count from the first arg with a defined element_size
+    inline for (fields, 0..) |field, i| {
+        if (elements_count == 0) {
+            const buf = @field(args, field.name);
+            const el_size = self.def.bindings[i].element_size;
+            if (el_size > 0) {
+                elements_count = @intCast(buf.size / el_size);
+            }
+        }
+    }
+
+    // Validate runtime buffer sizes before dispatching
+    inline for (fields, 0..) |field, i| {
+        const buf = @field(args, field.name);
+        const el_size = self.def.bindings[i].element_size;
+        if (el_size > 0) {
+            const expected_min_bytes = @as(u64, elements_count) * el_size;
+            if (buf.size < expected_min_bytes)
+                return error.BufferTooSmall;
+        }
+    }
+
+    var entries_buf: [32]c.WGPUBindGroupEntry = undefined;
+    var entry_count: usize = 0;
+
+    // Unpack tuple into WebGPU BindGroupEntries
+    inline for (fields, 0..) |field, i| {
+        const buf = @field(args, field.name);
+        if (@TypeOf(buf) != GpuBuffer) {
+            @compileError("All arguments in the tuple must be of type GpuBuffer");
+        }
+        entries_buf[entry_count] = .{
+            .binding = @intCast(i),
+            .buffer = buf.raw,
+            .offset = 0,
+            .size = buf.size, // Size exposes the fully allocated length
+        };
+        entry_count += 1;
+    }
+
+    // Optional uniform dispatch buffer appended at the end
+    var info_buf: ?GpuBuffer = null;
+    defer if (info_buf) |b| b.deinit();
+
+    if (self.def.append_info_buffer) {
+        info_buf = try GpuBuffer.init(
+            gloc,
+            @sizeOf(u32),
+            .initMany(&.{ .Uniform, .CopyDst }),
+        );
+        c.wgpuQueueWriteBuffer(gloc.device.queue, info_buf.?.raw, 0, &elements_count, @sizeOf(u32));
+
+        entries_buf[entry_count] = .{
+            .binding = @intCast(entry_count),
+            .buffer = info_buf.?.raw,
+            .offset = 0,
+            .size = @sizeOf(u32),
+        };
+        entry_count += 1;
+    }
+
+    const entries = entries_buf[0..entry_count];
+    try submitPass(gloc, self.pip, entries, elements_count, self.def.workgroup_size, self.def.max_workgroups);
+}
+
+fn submitPass(
+    gloc: GpuAllocator,
+    pipeline: c.WGPUComputePipeline,
+    entries: []const c.WGPUBindGroupEntry,
+    n: usize,
+    workgroup_size: u32,
+    max_workgroups: u32,
+) !void {
+    if (n == 0) return;
+
+    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 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/GpuProcess.zig

@@ -1,122 +0,0 @@
-/// 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/lib.zig

@@ -2,4 +2,4 @@ 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 GpuProcess = @import("GpuProcess.zig");
+pub const GpuCompute = @import("GpuCompute.zig");