19 Commits
Author SHA1 Message Date
adrien 6ec53bb909 Moved add and bench into an examples directory so src is just the library 2026-05-18 22:57:13 +02:00
adrien 6b933465b0 Better example 2026-05-18 17:00:44 +02:00
adrien 3d9cb4df04 Removed TensorInfo from example shader 2026-05-18 16:45:55 +02:00
adrien c62bead6e9 Better example (no Vec) 2026-05-18 16:38:30 +02:00
adrien 9dbd668de9 . 2026-05-18 16:33:26 +02:00
adrien 50bf436fa5 Update README 2026-05-18 16:33:07 +02:00
adrien df4eb36bc8 Added README 2026-05-18 16:30:27 +02:00
adrien c7afe28cc6 GpuBuffer align with size 4 as required by WebGPU
Otherwise a GpuBuffer of a single f16 is too small (only 2) and a eneven
f16 wouldn't work either
2026-05-18 16:24:43 +02:00
adrien d9cb4c4672 Better example comments 2026-05-18 16:18:58 +02:00
adrien 3a09e11f54 Deleted Vec.zig 2026-05-18 16:05:50 +02:00
adrien e608681ce1 Added preference power in DeviceCOnf 2026-05-18 16:05:42 +02:00
adrien d2b0735540 Moved all f32 out of GpuXxxx to all be in Vec 2026-05-18 15:35:38 +02:00
adrien 8e137c8f33 Vec no longer need import c 2026-05-18 15:28:28 +02:00
adrien 97d5f9001f Removed GpuPipeline for a GpuProcess 2026-05-18 15:28:02 +02:00
adrien d5e7f60926 Moved read stuff to GpuBuffer directly 2026-05-18 14:27:11 +02:00
adrien 2673aef0fd Moved c.gpuQueuWriteBuffer to GpuBuffer 2026-05-18 14:16:01 +02:00
adrien 7ce6b9cd1d Better example 2026-05-18 14:11:17 +02:00
adrien f5daf66784 Changed GpuAllocator to be like std.mem.Allocator
Now it is 2 ptr and I created a GpuArena. Point is to be like Zig, a
const allocator and a var arena that track everything.
2026-05-18 14:07:28 +02:00
adrien c3166e552b GpuBuffer init now take a EnumSet of BufferUsage and not a c.WGPUBufferUsage 2026-05-18 10:49:11 +02:00
15 changed files with 580 additions and 475 deletions
+83
View File
@@ -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.
+39 -73
View File
@@ -1,94 +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"); const std = @import("std");
pub fn build(b: *std.Build) void { pub fn build(b: *std.Build) !void {
const target = b.standardTargetOptions(.{}); const target = b.standardTargetOptions(.{});
const optimize = b.standardOptimizeOption(.{}); 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"), .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, .target = target,
.optimize = optimize,
}),
.name = "bench",
}); });
// 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 // Platform-specific system frameworks needed by wgpu-native
const t = target.result; const t = target.result;
if (t.os.tag == .macos) { 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) { } 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 { } 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); 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);
const exe_examp = b.addExecutable(.{
.root_module = b.createModule(.{
.root_source_file = b.path("src/example.zig"),
.link_libc = true,
.target = target,
.optimize = optimize,
}),
.name = "bench",
});
// wgpu-native headers + pre-built static library
exe_examp.root_module.addIncludePath(b.path("libs/wgpu-native/include"));
exe_examp.root_module.addLibraryPath(b.path("libs/wgpu-native/lib"));
exe_examp.root_module.addObjectFile(b.path("libs/wgpu-native/lib/libwgpu_native.a"));
if (t.os.tag == .macos) {
exe_examp.root_module.linkFramework("Metal", .{});
exe_examp.root_module.linkFramework("QuartzCore", .{});
exe_examp.root_module.linkFramework("Foundation", .{});
exe_examp.root_module.linkFramework("CoreGraphics", .{});
} else if (t.os.tag == .windows) {
exe_examp.root_module.linkSystemLibrary("d3d12", .{});
exe_examp.root_module.linkSystemLibrary("dxgi", .{});
exe_examp.root_module.linkSystemLibrary("user32", .{});
} else {
exe_examp.root_module.linkSystemLibrary("vulkan", .{});
exe_examp.root_module.linkSystemLibrary("gcc_s", .{});
} }
const examp = b.addRunArtifact(exe_examp);
run.step.dependOn(b.getInstallStep());
b.step("example", "Run basic example.").dependOn(&examp.step);
} }
+3 -33
View File
@@ -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, .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", .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 = .{}, .dependencies = .{},
.paths = .{ .paths = .{
"build.zig", "build.zig",
"build.zig.zon", "build.zig.zon",
"src", "src",
// For example... "libs",
//"LICENSE",
//"README.md",
}, },
} }
+59
View File
@@ -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});
}
+70 -32
View File
@@ -1,37 +1,75 @@
const std = @import("std"); 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("utils.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 { pub fn main(init: std.process.Init) !void {
const device = try GpuDevice.init(.{ .vram_bytes_limit = 4 * 1024 * 1024 * 1024 }); const device = try GpuDevice.init(.{ .vram_bytes_limit = 4 * 1024 * 1024 * 1024 });
defer device.deinit(); 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(); defer add_pip.deinit();
const allocator = init.gpa; const allocator = init.gpa;
// --- WARM-UP PHASE ---
{
var warmup_a = [_]f16{1.0};
var warmup_b = [_]f16{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{ const sizes = [_]usize{
1, 1,
256, 256,
@@ -41,9 +79,9 @@ pub fn main(init: std.process.Init) !void {
4 * 4 * 4 * 1024, 4 * 4 * 4 * 1024,
4 * 4 * 4 * 4 * 1024, 4 * 4 * 4 * 4 * 1024,
1024 * 1024, 1024 * 1024,
// 4 * 1024 * 1024, 4 * 1024 * 1024,
// 4 * 4 * 1024 * 1024, 4 * 4 * 1024 * 1024,
// 4 * 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,
}; };
@@ -78,9 +116,9 @@ pub fn main(init: std.process.Init) !void {
// --- 1. GPU ALLOCATION PHASE --- // --- 1. GPU ALLOCATION PHASE ---
const alloc_start = std.Io.Clock.awake.now(init.io); 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(); defer a.deinit();
const b = try Vec.initLoad(&gloc, data_b); const b = try Vec.initLoad(gloc, data_b);
defer b.deinit(); defer b.deinit();
const alloc_duration = alloc_start.durationTo(std.Io.Clock.awake.now(init.io)); const alloc_duration = alloc_start.durationTo(std.Io.Clock.awake.now(init.io));
@@ -90,15 +128,15 @@ pub fn main(init: std.process.Init) !void {
// --- 2. COMPUTE PHASE --- // --- 2. COMPUTE PHASE ---
const compute_start = std.Io.Clock.awake.now(init.io); 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(); defer sum.deinit();
// All 3 buffers (a, b, sum) are currently resident in VRAM here. // All 3 buffers (a, b, sum) are currently resident in VRAM here.
// Querying now catches the true peak allocation step. // 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_duration = compute_start.durationTo(std.Io.Clock.awake.now(init.io));
const compute_ns = @as(u64, @intCast(compute_duration.toNanoseconds())); const compute_ns = @as(u64, @intCast(compute_duration.toNanoseconds()));
@@ -107,7 +145,7 @@ pub fn main(init: std.process.Init) !void {
// --- 3. TRANSFER PHASE (Device -> Host) --- // --- 3. TRANSFER PHASE (Device -> Host) ---
const transfer_start = std.Io.Clock.awake.now(init.io); 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); defer allocator.free(out);
const transfer_duration = transfer_start.durationTo(std.Io.Clock.awake.now(init.io)); const transfer_duration = transfer_start.durationTo(std.Io.Clock.awake.now(init.io));
@@ -3,11 +3,7 @@ enable f16;
@group(0) @binding(0) var<storage, read> A: array<f16>; @group(0) @binding(0) var<storage, read> A: array<f16>;
@group(0) @binding(1) var<storage, read> B: 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(2) var<storage, read_write> C: array<f16>;
@group(0) @binding(3) var<uniform> size: u32;
struct TensorInfo {
size: u32,
};
@group(0) @binding(3) var<uniform> info: TensorInfo;
@compute @workgroup_size(256) @compute @workgroup_size(256)
fn main( fn main(
@@ -21,7 +17,7 @@ fn main(
var index = global_id.x; var index = global_id.x;
// 3. Stride through the tensor elements // 3. Stride through the tensor elements
while (index < info.size) { while (index < size) {
C[index] = A[index] + B[index]; C[index] = A[index] + B[index];
index += total_threads; // Jump forward by the total thread count index += total_threads; // Jump forward by the total thread count
} }
+15 -49
View File
@@ -1,58 +1,24 @@
// GpuAllocator.zig
const std = @import("std"); const std = @import("std");
const GpuDevice = @import("GpuDevice.zig"); const GpuDevice = @import("GpuDevice.zig");
const GpuBuffer = @import("GpuBuffer.zig");
const c = @import("utils.zig").c; const c = @import("utils.zig").c;
device: GpuDevice, const GpuAllocator = @This();
tracked_buffers: std.AutoHashMap(c.WGPUBuffer, void),
allocated_vram_bytes: u64 = 0,
pub fn init(cpu_allocator: std.mem.Allocator, device: GpuDevice) !@This() { /// The function definitions our underlying implementations must satisfy
return .{ pub const VTable = struct {
.device = device, alloc: *const fn (ctx: *anyopaque, bytes: u64, usage: c.WGPUBufferUsage) anyerror!c.WGPUBuffer,
.tracked_buffers = .init(cpu_allocator), free: *const fn (ctx: *anyopaque, buf_raw: c.WGPUBuffer, size: u64) void,
}; };
device: GpuDevice,
ptr: *anyopaque,
vtable: *const VTable,
pub fn allocBuffer(self: GpuAllocator, bytes: u64, usage: c.WGPUBufferUsage) !c.WGPUBuffer {
return self.vtable.alloc(self.ptr, bytes, usage);
} }
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;
errdefer {
c.wgpuBufferDestroy(buf);
c.wgpuBufferRelease(buf);
}
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;
}
} }
+72
View File
@@ -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;
}
}
+105 -17
View File
@@ -3,29 +3,50 @@ const c = @import("utils.zig").c;
const GpuAllocator = @import("GpuAllocator.zig"); const GpuAllocator = @import("GpuAllocator.zig");
raw: c.WGPUBuffer, raw: c.WGPUBuffer,
size: u64, size: u64, // Now tracks the 4-byte aligned size directly
usage: c.WGPUBufferUsage, 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 /// Allocates the underlying WebGPU handle and registers it to the parent GpuAllocator
pub fn init(gloc: *GpuAllocator, T: type, len: usize, usage: c.WGPUBufferUsage) !@This() { pub fn init(gloc: GpuAllocator, size: u64, usage: std.EnumSet(BufferUsage)) !@This() {
switch (@typeInfo(T)) { var use: u64 = 0;
.int, .float => {}, var iter = usage.iterator();
else => @compileError("GpuBuffer can only use int and float type"), while (iter.next()) |flag| use |= @intFromEnum(flag);
}
const bytes = @sizeOf(T) * len; // Automatically align the buffer size forward to a multiple of 4 bytes under the hood
const raw_handle = try gloc.registerBuffer(bytes, usage); const aligned_size = std.mem.alignForward(u64, size, 4);
const raw_handle = try gloc.allocBuffer(aligned_size, use);
return .{ return .{
.raw = raw_handle, .raw = raw_handle,
.size = bytes, .size = aligned_size, // Expose the aligned size to the rest of the application
.usage = usage, .usage = use,
.gloc = gloc, .gloc = gloc,
}; };
} }
/// Unregisters from the parent GpuAllocator and cleanly destroys GPU resources /// Unregisters from the parent GpuAllocator and cleanly destroys GPU resources
pub fn deinit(self: @This()) void { 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 /// Native mapAsync wrapper
@@ -39,12 +60,79 @@ pub fn mapAsync(
_ = c.wgpuBufferMapAsync(self.raw, mode, offset, size, callback_info); _ = 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 /// Native unmap wrapper
pub fn unmap(self: @This()) void { pub fn unmap(self: @This()) void {
c.wgpuBufferUnmap(self.raw); 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);
}
+9 -2
View File
@@ -10,6 +10,12 @@ const Ctx = struct {
const GpuDeviceConfig = struct { const GpuDeviceConfig = struct {
/// VRAM limit. Default 2 GB /// VRAM limit. Default 2 GB
vram_bytes_limit: u64 = 2 * 1024 * 1024 * 1024, 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, instance: c.WGPUInstance,
@@ -29,7 +35,7 @@ pub fn init(config: GpuDeviceConfig) !@This() {
var ctx = Ctx{}; var ctx = Ctx{};
_ = c.wgpuInstanceRequestAdapter( _ = c.wgpuInstanceRequestAdapter(
instance, instance,
&.{ .powerPreference = c.WGPUPowerPreference_HighPerformance }, &.{ .powerPreference = @intFromEnum(config.power_preference) },
.{ .callback = onAdapter, .userdata1 = &ctx }, .{ .callback = onAdapter, .userdata1 = &ctx },
); );
c.wgpuInstanceProcessEvents(instance); c.wgpuInstanceProcessEvents(instance);
@@ -87,7 +93,8 @@ pub fn deinit(self: @This()) void {
c.wgpuInstanceRelease(self.instance); 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); _ = c.wgpuDevicePoll(self.device, 1, null);
} }
-25
View File
@@ -1,25 +0,0 @@
const std = @import("std");
const GpuDevice = @import("GpuDevice.zig");
const c = @import("utils.zig").c;
const sv = @import("utils.zig").sv;
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);
}
+122
View File
@@ -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;
}
-195
View File
@@ -1,195 +0,0 @@
/// Dummy
const std = @import("std");
const c = @import("utils.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,
f16,
len,
c.WGPUBufferUsage_Storage | c.WGPUBufferUsage_CopyDst | c.WGPUBufferUsage_CopySrc,
),
.len = len,
};
}
pub fn initLoad(gloc: *GpuAllocator, data: []const f16) !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 f16,
) !void {
std.debug.assert(data.len == self.len);
const bytes = self.byteSize();
c.wgpuQueueWriteBuffer(device.queue, self.buf.raw, 0, data.ptr, bytes);
}
pub fn byteSize(self: Vec) u64 {
return @as(u64, self.len) * @sizeOf(f16);
}
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) ![]f16 {
const out = try alloc.alloc(f16, self.len);
const bytes = self.byteSize();
const staging = try GpuBuffer.init(
gloc,
f16,
self.len,
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 f16 = @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(f16));
// Create uniform buffer for this specific chunk's size
const info_buf = try GpuBuffer.init(
gloc,
u32,
1,
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;
}
-43
View File
@@ -1,43 +0,0 @@
const std = @import("std");
const GpuDevice = @import("GpuDevice.zig");
const GpuAllocator = @import("GpuAllocator.zig");
const GpuPipeline = @import("GpuPipeline.zig");
const Vec = @import("Vec.zig");
const c = @import("utils.zig").c;
pub fn main(init: std.process.Init) !void {
const allocator = init.gpa;
const device = try GpuDevice.init(.{ .vram_bytes_limit = 4 * 1024 * 1024 * 1024 });
defer device.deinit();
var gloc = try GpuAllocator.init(allocator, device);
defer gloc.deinit();
const add_pip = try GpuPipeline.init(device, @embedFile("shaders/add.wgsl"));
defer add_pip.deinit();
const data_a = try allocator.alloc(f16, 1024);
defer allocator.free(data_a);
const data_b = try allocator.alloc(f16, 1024);
defer allocator.free(data_b);
for (0..1024) |i| {
data_a[i] = @floatFromInt(i);
data_b[i] = @floatFromInt(1024 - 1 - i);
}
const a = try Vec.initLoad(&gloc, data_a);
defer a.deinit();
const b = try Vec.initLoad(&gloc, data_b);
defer b.deinit();
const sum = try a.run(&gloc, b, add_pip);
defer sum.deinit();
const out = try sum.read(&gloc, allocator);
defer allocator.free(out);
std.debug.print("{any}\n", .{out});
}
+2 -1
View File
@@ -1,4 +1,5 @@
pub const GpuAllocator = @import("GpuAllocator.zig"); pub const GpuAllocator = @import("GpuAllocator.zig");
pub const GpuArena = @import("GpuArena.zig");
pub const GpuBuffer = @import("GpuBuffer.zig"); pub const GpuBuffer = @import("GpuBuffer.zig");
pub const GpuDevice = @import("GpuDevice.zig"); pub const GpuDevice = @import("GpuDevice.zig");
pub const GpuPipeline = @import("GpuPipeline.zig"); pub const GpuProcess = @import("GpuProcess.zig");