Added GpuTexture tied to arena like buffer

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,
@@ -57,4 +56,39 @@ pub fn build(b: *std.Build) !void {
         const run_cmd = b.addRunArtifact(exe);
         run_step.dependOn(&run_cmd.step);
     }
+
+    const exe = b.addExecutable(.{
+        .name = "circle",
+        .root_module = b.createModule(.{
+            .root_source_file = b.path("src/circle.zig"),
+            .target = target,
+            .optimize = optimize,
+            .imports = &.{},
+        }),
+    });
+
+    exe.root_module.addIncludePath(b.path("libs/wgpu-native/include"));
+    exe.root_module.addLibraryPath(b.path("libs/wgpu-native/lib"));
+    exe.root_module.addObjectFile(b.path("libs/wgpu-native/lib/libwgpu_native.a"));
+
+    // Platform-specific system frameworks needed by wgpu-native
+    if (t.os.tag == .macos) {
+        exe.root_module.linkFramework("Metal", .{});
+        exe.root_module.linkFramework("QuartzCore", .{});
+        exe.root_module.linkFramework("Foundation", .{});
+        exe.root_module.linkFramework("CoreGraphics", .{});
+    } else if (t.os.tag == .windows) {
+        exe.root_module.linkSystemLibrary("d3d12", .{});
+        exe.root_module.linkSystemLibrary("dxgi", .{});
+        exe.root_module.linkSystemLibrary("user32", .{});
+    } else {
+        exe.root_module.linkSystemLibrary("vulkan", .{});
+        exe.root_module.linkSystemLibrary("gcc_s", .{});
+    }
+
+    b.installArtifact(exe);
+
+    const run_step = b.step("circle", "Run circle");
+    const run_cmd = b.addRunArtifact(exe);
+    run_step.dependOn(&run_cmd.step);
 }

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,24 +1,29 @@
-// 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,
+    allocBuffer: *const fn (ctx: *anyopaque, desc: c.WGPUBufferDescriptor) anyerror!c.WGPUBuffer,
-    free: *const fn (ctx: *anyopaque, buf_raw: c.WGPUBuffer, size: u64) void,
+    freeBuffer: *const fn (ctx: *anyopaque, buf_raw: c.WGPUBuffer) void,
+    allocTexture: *const fn (ctx: *anyopaque, desc: c.WGPUTextureDescriptor) anyerror!c.WGPUTexture,
+    freeTexture: *const fn (ctx: *anyopaque, buf_raw: c.WGPUTexture) void,
 };
 
 device: GpuDevice,
 ptr: *anyopaque,
 vtable: *const VTable,
 
-pub fn allocBuffer(self: GpuAllocator, bytes: u64, usage: c.WGPUBufferUsage) !c.WGPUBuffer {
+pub fn allocBuffer(self: @This(), desc: c.WGPUBufferDescriptor) !c.WGPUBuffer {
-    return self.vtable.alloc(self.ptr, bytes, usage);
+    return self.vtable.allocBuffer(self.ptr, desc);
 }
 
-pub fn freeBuffer(self: GpuAllocator, buf_raw: c.WGPUBuffer, size: u64) void {
+pub fn freeBuffer(self: @This(), buf_raw: c.WGPUBuffer) void {
-    self.vtable.free(self.ptr, buf_raw, size);
+    self.vtable.freeBuffer(self.ptr, buf_raw);
+}
+
+pub fn allocTexture(self: @This(), desc: c.WGPUTextureDescriptor) !c.WGPUTexture {
+    return self.vtable.allocTexture(self.ptr, desc);
+}
+
+pub fn freeTexture(self: @This(), buf_raw: c.WGPUTexture) void {
+    self.vtable.freeTexture(self.ptr, buf_raw);
 }

src/GpuArena.zig

@@ -1,72 +1,107 @@
-// GpuArena.zig
 const std = @import("std");
 const GpuDevice = @import("GpuDevice.zig");
 const GpuAllocator = @import("GpuAllocator.zig");
+const GpuTextureFormat = @import("lib.zig").GpuTextureFormat;
 const c = @import("utils.zig").c;
 
-const GpuArena = @This();
-
 device: GpuDevice,
-tracked_buffers: std.AutoHashMap(c.WGPUBuffer, void),
+tracked_buffers: std.AutoHashMap(c.WGPUBuffer, c.WGPUBufferDescriptor),
+tracked_textures: std.AutoHashMap(c.WGPUTexture, c.WGPUTextureDescriptor),
 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),
+        .tracked_textures = .init(cpu_allocator),
     };
 }
 
-pub fn deinit(self: *GpuArena) void {
+pub fn deinit(self: *@This()) void {
-    var it = self.tracked_buffers.keyIterator();
+    var it_buffer = self.tracked_buffers.keyIterator();
-    while (it.next()) |buf_ptr| {
+    while (it_buffer.next()) |buf_ptr| {
         c.wgpuBufferDestroy(buf_ptr.*);
         c.wgpuBufferRelease(buf_ptr.*);
     }
     self.tracked_buffers.deinit();
+
+    var it_texture = self.tracked_textures.keyIterator();
+    while (it_texture.next()) |tex_ptr|
+        c.wgpuTextureRelease(tex_ptr.*);
+    self.tracked_textures.deinit();
 }
 
 /// 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,
         .vtable = &.{
-            .alloc = alloc,
+            .allocBuffer = allocBuffer,
-            .free = free,
+            .freeBuffer = freeBuffer,
+            .allocTexture = allocTexture,
+            .freeTexture = freeTexture,
         },
     };
 }
 
-fn alloc(ctx: *anyopaque, bytes: u64, usage: c.WGPUBufferUsage) anyerror!c.WGPUBuffer {
+fn allocBuffer(ctx: *anyopaque, desc: c.WGPUBufferDescriptor) anyerror!c.WGPUBuffer {
-    const self: *GpuArena = @ptrCast(@alignCast(ctx));
+    const self: *@This() = @ptrCast(@alignCast(ctx));
 
-    if (bytes > self.device.limits.maxBufferSize)
+    if (desc.size > self.device.limits.maxBufferSize)
         return error.SingleBufferExceedsLimit;
 
-    if (bytes + self.allocated_vram_bytes > self.device.config.vram_bytes_limit)
+    if (desc.size + self.allocated_vram_bytes > self.device.config.vram_bytes_limit)
         return error.ExceedsVramBudget;
 
-    const buf = c.wgpuDeviceCreateBuffer(self.device.device, &.{
+    const buf = c.wgpuDeviceCreateBuffer(self.device.device, &desc) orelse return error.BufferAlloc;
-        .usage = usage,
-        .size = bytes,
-    }) orelse return error.BufferAlloc;
     errdefer {
         c.wgpuBufferDestroy(buf);
         c.wgpuBufferRelease(buf);
     }
 
-    try self.tracked_buffers.put(buf, {});
+    try self.tracked_buffers.put(buf, desc);
-    self.allocated_vram_bytes += bytes;
+    self.allocated_vram_bytes += desc.size;
     return buf;
 }
 
-fn free(ctx: *anyopaque, buf_raw: c.WGPUBuffer, size: u64) void {
+fn freeBuffer(ctx: *anyopaque, buf_raw: c.WGPUBuffer) void {
-    const self: *GpuArena = @ptrCast(@alignCast(ctx));
+    const self: *@This() = @ptrCast(@alignCast(ctx));
 
-    if (self.tracked_buffers.remove(buf_raw)) {
+    if (self.tracked_buffers.fetchRemove(buf_raw)) |kv| {
         c.wgpuBufferDestroy(buf_raw);
         c.wgpuBufferRelease(buf_raw);
-        self.allocated_vram_bytes -= size;
+        self.allocated_vram_bytes -= kv.value.size;
+    }
+}
+
+fn allocTexture(ctx: *anyopaque, desc: c.WGPUTextureDescriptor) anyerror!c.WGPUTexture {
+    const self: *@This() = @ptrCast(@alignCast(ctx));
+
+    const format: GpuTextureFormat = @enumFromInt(desc.format);
+    const bytes_size = desc.size.width * desc.size.height * format.bytesPerPixel();
+    if (bytes_size > self.device.limits.maxBufferSize)
+        return error.SingleBufferExceedsLimit;
+
+    if (bytes_size + self.allocated_vram_bytes > self.device.config.vram_bytes_limit)
+        return error.ExceedsVramBudget;
+
+    const texture = c.wgpuDeviceCreateTexture(self.device.device, &desc) orelse return error.Texture;
+
+    try self.tracked_textures.put(texture, desc);
+    self.allocated_vram_bytes += bytes_size;
+    return texture;
+}
+
+fn freeTexture(ctx: *anyopaque, texture_raw: c.WGPUTexture) void {
+    const self: *@This() = @ptrCast(@alignCast(ctx));
+
+    if (self.tracked_textures.fetchRemove(texture_raw)) |kv| {
+        c.wgpuTextureRelease(texture_raw);
+
+        const desc = kv.value;
+        const format: GpuTextureFormat = @enumFromInt(desc.format);
+        const bytes_size = desc.size.width * desc.size.height * format.bytesPerPixel();
+        self.allocated_vram_bytes -= bytes_size;
     }
 }

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,
 
@@ -30,10 +30,10 @@ pub fn init(gloc: GpuAllocator, size: u64, usage: std.EnumSet(BufferUsage)) !@Th
     // 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);
+    const raw_handle = try gloc.allocBuffer(.{ .size = aligned_size, .usage = 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,
     };
@@ -41,7 +41,7 @@ pub fn init(gloc: GpuAllocator, size: u64, usage: std.EnumSet(BufferUsage)) !@Th
 
 /// Unregisters from the parent GpuAllocator and cleanly destroys GPU resources
 pub fn deinit(self: @This()) void {
-    self.gloc.freeBuffer(self.raw, self.size);
+    self.gloc.freeBuffer(self.raw);
 }
 
 /// Native getConstMappedRange wrapper

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/GpuRender.zig

@@ -0,0 +1,167 @@
+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");
+
+pub const Binding = struct {
+    element_size: u32 = 0,
+};
+
+pub const RenderDef = struct {
+    bindings: []const Binding = &.{},
+    /// The surface texture format we are rendering to (e.g., BGRA8Unorm)
+    texture_format: c.WGPUTextureFormat,
+    /// The names of the entry points inside your WGSL code
+    vertex_entry: []const u8 = "vs_main",
+    fragment_entry: []const u8 = "fs_main",
+    /// Primitive topology, default to triangle list
+    topology: c.WGPUPrimitiveTopology = c.WGPUPrimitiveTopology_TriangleList,
+};
+
+pip: c.WGPURenderPipeline,
+def: RenderDef,
+
+pub fn init(device: GpuDevice, wgsl: []const u8, def: RenderDef) !@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);
+
+    // 1. Setup the Color Target State (where the fragment shader outputs)
+    const blend = c.WGPUBlendState{
+        .color = .{ .operation = c.WGPUBlendOperation_Add, .srcFactor = c.WGPUBlendFactor_SrcAlpha, .dstFactor = c.WGPUBlendFactor_OneMinusSrcAlpha },
+        .alpha = .{ .operation = c.WGPUBlendOperation_Add, .srcFactor = c.WGPUBlendFactor_One, .dstFactor = c.WGPUBlendFactor_Zero },
+    };
+
+    const color_target = c.WGPUColorTargetState{
+        .format = def.texture_format,
+        .blend = &blend,
+        .writeMask = c.WGPUColorWriteMask_All,
+    };
+
+    // 2. Setup the Fragment State
+    const fragment_state = c.WGPUFragmentState{
+        .module = shader,
+        .entryPoint = sv(def.fragment_entry),
+        .targetCount = 1,
+        .targets = &color_target,
+    };
+
+    // 3. Compile the Complete Render Pipeline
+    const pip = c.wgpuDeviceCreateRenderPipeline(device.device, &.{
+        .vertex = .{
+            .module = shader,
+            .entryPoint = sv(def.vertex_entry),
+        },
+        .primitive = .{
+            .topology = def.topology,
+            .stripIndexFormat = c.WGPUIndexFormat_Undefined,
+            .frontFace = c.WGPUFrontFace_CCW,
+            .cullMode = c.WGPUCullMode_None,
+        },
+        .multisample = .{
+            .count = 1,
+            .mask = 0xFFFFFFFF,
+            .alphaToCoverageEnabled = 0,
+        },
+        .fragment = &fragment_state,
+    }) orelse return error.Pipeline;
+
+    return .{
+        .pip = pip,
+        .def = def,
+    };
+}
+
+pub fn deinit(self: @This()) void {
+    c.wgpuRenderPipelineRelease(self.pip);
+}
+
+/// Execute the render pass targeting a specific frame texture view.
+/// Passes bind groups via a tuple exactly like your original compute setup.
+pub fn draw(
+    self: @This(),
+    gloc: GpuAllocator,
+    target_view: c.WGPUTextureView,
+    vertex_count: u32,
+    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. .{ uniform_buf }");
+
+    const fields = type_info.@"struct".fields;
+    if (fields.len != self.def.bindings.len)
+        return error.InvalidArgumentCount;
+
+    var entries_buf: [32]c.WGPUBindGroupEntry = undefined;
+
+    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[i] = .{
+            .binding = @intCast(i),
+            .buffer = buf.raw,
+            .offset = 0,
+            .size = buf.size,
+        };
+    }
+
+    const entries = entries_buf[0..fields.len];
+
+    // Create Render Bind Group from layout
+    const bgl = c.wgpuRenderPipelineGetBindGroupLayout(self.pip, 0);
+    defer c.wgpuBindGroupLayoutRelease(bgl);
+
+    const bg = c.wgpuDeviceCreateBindGroup(gloc.device.device, &.{
+        .layout = bgl,
+        .entries = entries.ptr,
+        .entryCount = @intCast(entries.len),
+    }) orelse return error.BindGroup;
+    defer c.wgpuBindGroupRelease(bg);
+
+    // Encode Render Command
+    const enc = c.wgpuDeviceCreateCommandEncoder(gloc.device.device, null) orelse return error.Encoder;
+    defer c.wgpuCommandEncoderRelease(enc);
+
+    const color_attachment = c.WGPURenderPassColorAttachment{
+        .view = target_view,
+        .resolveTarget = null,
+        .loadOp = c.WGPULoadOp_Clear,
+        .storeOp = c.WGPUStoreOp_Store,
+        .clearValue = .{ .r = 0.1, .g = 0.1, .b = 0.1, .a = 1.0 },
+        .depthSlice = c.WGPU_DEPTH_SLICE_UNDEFINED,
+    };
+
+    const pass_desc = c.WGPURenderPassDescriptor{
+        .colorAttachmentCount = 1,
+        .colorAttachments = &color_attachment,
+        .depthStencilAttachment = null,
+    };
+
+    const pass = c.wgpuCommandEncoderBeginRenderPass(enc, &pass_desc);
+    c.wgpuRenderPassEncoderSetPipeline(pass, self.pip);
+
+    if (fields.len > 0) {
+        c.wgpuRenderPassEncoderSetBindGroup(pass, 0, bg, 0, null);
+    }
+
+    // Draw! (Instead of Compute Dispatch)
+    c.wgpuRenderPassEncoderDraw(pass, vertex_count, 1, 0, 0);
+
+    c.wgpuRenderPassEncoderEnd(pass);
+    c.wgpuRenderPassEncoderRelease(pass);
+
+    const cmd = c.wgpuCommandEncoderFinish(enc, null);
+    defer c.wgpuCommandBufferRelease(cmd);
+
+    c.wgpuQueueSubmit(gloc.device.queue, 1, &cmd);
+}

src/GpuTexture.zig

@@ -0,0 +1,137 @@
+const std = @import("std");
+const c = @import("utils.zig").c;
+const GpuAllocator = @import("GpuAllocator.zig");
+const GpuTextureFormat = @import("lib.zig").GpuTextureFormat;
+
+const TextureUsage = enum(u64) {
+    None = 0x0000000000000000,
+    CopySrc = 0x0000000000000001,
+    CopyDst = 0x0000000000000002,
+    TextureBinding = 0x0000000000000004,
+    StorageBinding = 0x0000000000000008,
+    RenderAttachment = 0x0000000000000010,
+    TransientAttachment = 0x0000000000000020,
+};
+
+raw: c.WGPUTexture,
+size: c.WGPUExtent3D,
+usage: c.WGPUTextureUsage,
+format: GpuTextureFormat,
+gloc: GpuAllocator,
+
+/// Allocates the underlying WebGPU handle and registers it to the parent GpuAllocator
+pub fn init(gloc: GpuAllocator, format: GpuTextureFormat, size: c.WGPUExtent3D, usage: std.EnumSet(TextureUsage)) !@This() {
+    var use: u64 = 0;
+    var iter = usage.iterator();
+    while (iter.next()) |flag| use |= @intFromEnum(flag);
+
+    const desc = c.WGPUTextureDescriptor{
+        .usage = use,
+        .dimension = c.WGPUTextureDimension_2D,
+        .size = size,
+        .format = @intCast(@intFromEnum(format)),
+        .mipLevelCount = 1,
+        .sampleCount = 1,
+    };
+    const raw = try gloc.allocTexture(desc);
+
+    return .{ .gloc = gloc, .raw = raw, .size = size, .format = format, .usage = use };
+}
+
+/// Unregisters from the parent GpuAllocator and cleanly destroys GPU resources
+pub fn deinit(self: @This()) void {
+    self.gloc.freeTexture(self.raw);
+}
+
+/// Native getConstMappedRange wrapper
+pub fn getConstMappedRange(self: @This(), offset: u64, size: u64) ?*const anyopaque {
+    return c.wgpuBufferGetConstMappedRange(self.raw, offset, size);
+}
+
+/// Native mapAsync wrapper
+pub fn mapAsync(
+    self: @This(),
+    mode: c.WGPUMapMode,
+    offset: u64,
+    size: u64,
+    callback_info: c.WGPUBufferMapCallbackInfo,
+) void {
+    _ = c.wgpuBufferMapAsync(self.raw, mode, offset, size, callback_info);
+}
+
+/// 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/circle.zig

@@ -0,0 +1,113 @@
+const std = @import("std");
+const gpu = @import("lib.zig");
+const c = @import("utils.zig").c;
+const sv = @import("utils.zig").sv;
+const GpuDevice = gpu.GpuDevice;
+const GpuArena = gpu.GpuArena;
+const GpuBuffer = gpu.GpuBuffer;
+const GpuRender = gpu.GpuRender;
+const GpuTexture = gpu.GpuTexture;
+
+pub fn main(init: std.process.Init) !void {
+    const allocator = init.gpa;
+
+    // 1. Open the raw headless GPU Device you shared
+    const device = try GpuDevice.init(.{});
+    defer device.deinit();
+
+    var grena = GpuArena.init(allocator, device);
+    defer grena.deinit();
+    const gloc = grena.gpuAllocator();
+
+    const width: u32 = 512;
+    const height: u32 = 512;
+    const render_format = c.WGPUTextureFormat_RGBA8Unorm;
+
+    // 2. Load our Render Pipeline (Procedural Triangle Strip)
+    const circle_rp = try GpuRender.init(
+        device,
+        @embedFile("shaders/circle.wgsl"),
+        .{
+            .bindings = &.{},
+            .texture_format = render_format,
+            .topology = c.WGPUPrimitiveTopology_TriangleStrip,
+        },
+    );
+    defer circle_rp.deinit();
+
+    // 3. Create the offscreen VRAM texture to render into
+    const texture = try GpuTexture.init(
+        gloc,
+        .RGBA8Unorm,
+        .{ .width = width, .height = height, .depthOrArrayLayers = 1 },
+        .initMany(&.{ .RenderAttachment, .CopySrc }),
+    );
+    defer texture.deinit();
+
+    const target_view = c.wgpuTextureCreateView(texture.raw, null) orelse return error.View;
+    defer c.wgpuTextureViewRelease(target_view);
+
+    // 4. Create a staging buffer to pull pixels from VRAM to CPU
+    // 4 bytes per pixel (RGBA8)
+    const row_bytes = width * 4;
+    const buffer_bytes = row_bytes * height;
+
+    // Create a regular GpuBuffer set up to receive texture copy transfers
+    const cpu_staging_buf = try GpuBuffer.init(gloc, buffer_bytes, .initMany(&.{ .CopyDst, .CopySrc }));
+
+    // 5. Draw the Circle Frame into the texture view!
+    try circle_rp.draw(gloc, target_view, 4, .{});
+
+    // 6. Copy the texture data into our CPU staging buffer
+    const enc = c.wgpuDeviceCreateCommandEncoder(device.device, null) orelse return error.Encoder;
+    defer c.wgpuCommandEncoderRelease(enc);
+
+    const src_copy = c.WGPUTexelCopyTextureInfo{
+        .texture = texture.raw,
+        .mipLevel = 0,
+        .origin = .{ .x = 0, .y = 0, .z = 0 },
+        .aspect = c.WGPUTextureAspect_All,
+    };
+    const dst_copy = c.WGPUTexelCopyBufferInfo{
+        .buffer = cpu_staging_buf.raw,
+        .layout = .{
+            .offset = 0,
+            .bytesPerRow = row_bytes,
+            .rowsPerImage = height,
+        },
+    };
+    const copy_size = c.WGPUExtent3D{ .width = width, .height = height, .depthOrArrayLayers = 1 };
+
+    c.wgpuCommandEncoderCopyTextureToBuffer(enc, &src_copy, &dst_copy, &copy_size);
+
+    const cmd = c.wgpuCommandEncoderFinish(enc, null);
+    defer c.wgpuCommandBufferRelease(cmd);
+    c.wgpuQueueSubmit(device.queue, 1, &cmd);
+
+    // 7. Map and read the raw image bytes back to CPU
+    // (This uses whatever slice-reading helpers your `GpuBuffer` wrapper provides)
+    const pixels = try cpu_staging_buf.read(allocator, u8);
+    defer allocator.free(pixels);
+
+    // Now you have the raw binary image data! Let's output a simple Netpbm PPM image file
+    // so you can actually open and look at your rendered circle.
+    try savePpm(init.io, "circle.ppm", width, height, pixels);
+    std.debug.print("Successfully rendered circle to circle.ppm!\n", .{});
+}
+
+fn savePpm(io: std.Io, filename: []const u8, w: u32, h: u32, rgba_pixels: []const u8) !void {
+    const file = try std.Io.Dir.cwd().createFile(io, filename, .{});
+    defer file.close(io);
+
+    var buf: [255]u8 = undefined;
+    var writer = file.writer(io, &buf);
+
+    // PPM Header: P6 format means raw RGB bytes
+    try writer.interface.print("P6\n{d} {d}\n255\n", .{ w, h });
+
+    // Strip Alpha channel when writing out to standard RGB PPM format
+    var i: usize = 0;
+    while (i < rgba_pixels.len) : (i += 4) {
+        try writer.interface.writeAll(rgba_pixels[i .. i + 3]);
+    }
+}

src/lib.zig

@@ -2,4 +2,175 @@ 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");
+pub const GpuRender = @import("GpuRender.zig");
+pub const GpuTexture = @import("GpuTexture.zig");
+
+pub const GpuTextureFormat = enum(c_int) {
+    Undefined = 0,
+    R8Unorm = 1,
+    R8Snorm = 2,
+    R8Uint = 3,
+    R8Sint = 4,
+    R16Unorm = 5,
+    R16Snorm = 6,
+    R16Uint = 7,
+    R16Sint = 8,
+    R16Float = 9,
+    RG8Unorm = 10,
+    RG8Snorm = 11,
+    RG8Uint = 12,
+    RG8Sint = 13,
+    R32Float = 14,
+    R32Uint = 15,
+    R32Sint = 16,
+    RG16Unorm = 17,
+    RG16Snorm = 18,
+    RG16Uint = 19,
+    RG16Sint = 20,
+    RG16Float = 21,
+    RGBA8Unorm = 22,
+    RGBA8UnormSrgb = 23,
+    RGBA8Snorm = 24,
+    RGBA8Uint = 25,
+    RGBA8Sint = 26,
+    BGRA8Unorm = 27,
+    BGRA8UnormSrgb = 28,
+    RGB10A2Uint = 29,
+    RGB10A2Unorm = 30,
+    RG11B10Ufloat = 31,
+    RGB9E5Ufloat = 32,
+    RG32Float = 33,
+    RG32Uint = 34,
+    RG32Sint = 35,
+    RGBA16Unorm = 36,
+    RGBA16Snorm = 37,
+    RGBA16Uint = 38,
+    RGBA16Sint = 39,
+    RGBA16Float = 40,
+    RGBA32Float = 41,
+    RGBA32Uint = 42,
+    RGBA32Sint = 43,
+    Stencil8 = 44,
+    Depth16Unorm = 45,
+    Depth24Plus = 46,
+    Depth24PlusStencil8 = 47,
+    Depth32Float = 48,
+    Depth32FloatStencil8 = 49,
+    BC1RGBAUnorm = 50,
+    BC1RGBAUnormSrgb = 51,
+    BC2RGBAUnorm = 52,
+    BC2RGBAUnormSrgb = 53,
+    BC3RGBAUnorm = 54,
+    BC3RGBAUnormSrgb = 55,
+    BC4RUnorm = 56,
+    BC4RSnorm = 57,
+    BC5RGUnorm = 58,
+    BC5RGSnorm = 59,
+    BC6HRGBUfloat = 60,
+    BC6HRGBFloat = 61,
+    BC7RGBAUnorm = 62,
+    BC7RGBAUnormSrgb = 63,
+    ETC2RGB8Unorm = 64,
+    ETC2RGB8UnormSrgb = 65,
+    ETC2RGB8A1Unorm = 66,
+    ETC2RGB8A1UnormSrgb = 67,
+    ETC2RGBA8Unorm = 68,
+    ETC2RGBA8UnormSrgb = 69,
+    EACR11Unorm = 70,
+    EACR11Snorm = 71,
+    EACRG11Unorm = 72,
+    EACRG11Snorm = 73,
+    ASTC4x4Unorm = 74,
+    ASTC4x4UnormSrgb = 75,
+    ASTC5x4Unorm = 76,
+    ASTC5x4UnormSrgb = 77,
+    ASTC5x5Unorm = 78,
+    ASTC5x5UnormSrgb = 79,
+    ASTC6x5Unorm = 80,
+    ASTC6x5UnormSrgb = 81,
+    ASTC6x6Unorm = 82,
+    ASTC6x6UnormSrgb = 83,
+    ASTC8x5Unorm = 84,
+    ASTC8x5UnormSrgb = 85,
+    ASTC8x6Unorm = 86,
+    ASTC8x6UnormSrgb = 87,
+    ASTC8x8Unorm = 88,
+    ASTC8x8UnormSrgb = 89,
+    ASTC10x5Unorm = 90,
+    ASTC10x5UnormSrgb = 91,
+    ASTC10x6Unorm = 92,
+    ASTC10x6UnormSrgb = 93,
+    ASTC10x8Unorm = 94,
+    ASTC10x8UnormSrgb = 95,
+    ASTC10x10Unorm = 96,
+    ASTC10x10UnormSrgb = 97,
+    ASTC12x10Unorm = 98,
+    ASTC12x10UnormSrgb = 99,
+    ASTC12x12Unorm = 100,
+    ASTC12x12UnormSrgb = 101,
+    Force32 = 2147483647,
+
+    pub fn bytesPerPixel(format: GpuTextureFormat) u32 {
+        return switch (format) {
+            // 8-bit formats (1 byte)
+            .R8Unorm, .R8Snorm, .R8Uint, .R8Sint, .Stencil8 => 1,
+
+            // 16-bit formats (2 bytes)
+            .R16Unorm,
+            .R16Snorm,
+            .R16Uint,
+            .R16Sint,
+            .R16Float,
+            .RG8Unorm,
+            .RG8Snorm,
+            .RG8Uint,
+            .RG8Sint,
+            .Depth16Unorm,
+            => 2,
+
+            // 32-bit formats (4 bytes)
+            .R32Float,
+            .R32Uint,
+            .R32Sint,
+            .RG16Unorm,
+            .RG16Snorm,
+            .RG16Uint,
+            .RG16Sint,
+            .RG16Float,
+            .RGBA8Unorm,
+            .RGBA8UnormSrgb,
+            .RGBA8Snorm,
+            .RGBA8Uint,
+            .RGBA8Sint,
+            .BGRA8Unorm,
+            .BGRA8UnormSrgb,
+            .RGB10A2Uint,
+            .RGB10A2Unorm,
+            .RG11B10Ufloat,
+            .RGB9E5Ufloat,
+            .Depth24Plus,
+            .Depth32Float,
+            => 4,
+
+            // 64-bit formats (8 bytes)
+            .RG32Float,
+            .RG32Uint,
+            .RG32Sint,
+            .RGBA16Unorm,
+            .RGBA16Snorm,
+            .RGBA16Uint,
+            .RGBA16Sint,
+            .RGBA16Float,
+            .Depth24PlusStencil8, // 24-bit depth + 8-bit stencil layout padded to 4+4 or 1+3
+            .Depth32FloatStencil8, // 32-bit float depth + 8-bit stencil (padded to 8 bytes)
+            => 8,
+
+            // 128-bit formats (16 bytes)
+            .RGBA32Float, .RGBA32Uint, .RGBA32Sint => 16,
+
+            // Block Compressed Formats (Handled separately)
+            else => 0,
+        };
+    }
+};

src/shaders/add.wgsl

@@ -0,0 +1,24 @@
+enable f16;
+
+@group(0) @binding(0) var<storage, read> A: array<f16>;
+@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> size: u32; 
+
+@compute @workgroup_size(256)
+fn main(
+    @builtin(global_invocation_id) global_id : vec3<u32>,
+    @builtin(num_workgroups) num_workgroups: vec3<u32>
+) {
+    // 1. Calculate the total number of threads across the entire grid
+    let total_threads = num_workgroups.x * 256u; 
+    
+    // 2. Start at this thread's unique global ID
+    var index = global_id.x;
+    
+    // 3. Stride through the tensor elements
+    while (index < size) {
+        C[index] = A[index] + B[index];
+        index += total_threads; // Jump forward by the total thread count
+    }
+}

src/shaders/circle.wgsl

@@ -0,0 +1,39 @@
+struct VertexOutput {
+    @builtin(position) position: vec4f,
+    @location(0) uv: vec2f,
+};
+
+@vertex
+fn vs_main(@builtin(vertex_index) vertex_index: u32) -> VertexOutput {
+    var output: VertexOutput;
+    // Hardcoded fullscreen quad layout using 4 vertices (Triangle Strip)
+    // Indexes: 0: Top-Left, 1: Bottom-Left, 2: Top-Right, 3: Bottom-Right
+    var pos = array<vec2f, 4>(
+        vec2f(-1.0,  1.0),
+        vec2f(-1.0, -1.0),
+        vec2f( 1.0,  1.0),
+        vec2f( 1.0, -1.0)
+    );
+    
+    output.position = vec4f(pos[vertex_index], 0.0, 1.0);
+    output.uv = pos[vertex_index]; // Ranges cleanly from -1.0 to 1.0
+    return output;
+}
+
+@fragment
+fn fs_main(input: VertexOutput) -> @location(0) vec4f {
+    // Distance from the center (0,0)
+    let distance = length(input.uv);
+    let radius = 0.5;
+    
+    // Smooth out pixel edges (anti-aliasing)
+    let edge_softness = 0.005;
+    let alpha = 1.0 - smoothstep(radius - edge_softness, radius + edge_softness, distance);
+    
+    if (alpha <= 0.0) {
+        discard; 
+    }
+    
+    // Draw a sharp/smooth red circle
+    return vec4f(1.0, 0.3, 0.3, alpha);
+}