4 Commits
11 changed files with 155 additions and 132 deletions
+14 -14
View File
@@ -8,7 +8,7 @@ The library exports the following 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. It provides the base `GpuAllocator` for raw VRAM allocations. * **`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. It provides the base `GpuAllocator` for raw VRAM allocations.
* **`GpuArenaAllocator`**: A memory management layer that wraps a base allocator to track and automatically destroy all allocated WebGPU buffers, textures, views, and pipelines when deinitialized. * **`GpuArenaAllocator`**: A memory management layer that wraps a base allocator to track and automatically destroy all allocated WebGPU buffers, textures, views, and pipelines when deinitialized.
* **`GpuBuffer`**: Wraps native WebGPU buffers. It provides a `.load()` method for CPU-to-GPU data transfers and a `.read()` method to map GPU data back to the CPU. * **`GpuBuffer`**: Wraps native WebGPU buffers. It provides a `.load()` method for CPU-to-GPU data transfers and a `.read()` method that utilizes a staging buffer to map GPU data back to the CPU.
* **`GpuCompute`**: Compiles WGSL source code into a compute pipeline and dispatches compute workgroups. * **`GpuCompute`**: Compiles WGSL source code into a compute pipeline and dispatches compute workgroups.
* **`GpuRender` / `GpuTexture` / `GpuTextureView`**: Components used to initialize render pipelines, set up render attachments (textures), and bind render targets for offscreen drawing. * **`GpuRender` / `GpuTexture` / `GpuTextureView`**: Components used to initialize render pipelines, set up render attachments (textures), and bind render targets for offscreen drawing.
@@ -36,11 +36,11 @@ pub fn main(init: std.process.Init) !void {
// 2. Create a GPU Arena to manage VRAM // 2. Create a GPU Arena to manage VRAM
var grena = GpuArenaAllocator.init(allocator, device.gpuAllocator()); var grena = GpuArenaAllocator.init(allocator, device.gpuAllocator());
defer grena.deinit(); defer grena.deinit();
const glloc = grena.gpuAllocator(); const gloc = grena.gpuAllocator();
// 3. Load the WGSL compute pipeline // 3. Load the WGSL compute pipeline
const add_cp = try GpuCompute.init( const add_cp = try GpuCompute.init(
glloc, gloc,
@embedFile("shaders/add.wgsl"), @embedFile("shaders/add.wgsl"),
.{ .{
.label = "add", .label = "add",
@@ -66,9 +66,9 @@ pub fn main(init: std.process.Init) !void {
// 5. Initialize raw GPU Buffers // 5. Initialize raw GPU Buffers
const byte_size = len * @sizeOf(f16); const byte_size = len * @sizeOf(f16);
const buf_a = try GpuBuffer.init(glloc, .{ .label = "a", .size = byte_size, .usage = .initMany(&.{ .Storage, .CopyDst, .CopySrc }) }); const buf_a = try GpuBuffer.init(gloc, .{ .label = "a", .size = byte_size, .usage = .initMany(&.{ .Storage, .CopyDst, .CopySrc }) });
const buf_b = try GpuBuffer.init(glloc, .{ .label = "b", .size = byte_size, .usage = .initMany(&.{ .Storage, .CopyDst, .CopySrc }) }); const buf_b = try GpuBuffer.init(gloc, .{ .label = "b", .size = byte_size, .usage = .initMany(&.{ .Storage, .CopyDst, .CopySrc }) });
const buf_out = try GpuBuffer.init(glloc, .{ .label = "out", .size = byte_size, .usage = .initMany(&.{ .Storage, .CopyDst, .CopySrc }) }); const buf_out = try GpuBuffer.init(gloc, .{ .label = "out", .size = byte_size, .usage = .initMany(&.{ .Storage, .CopyDst, .CopySrc }) });
// Note: Buffers are safely tied to the GpuArenaAllocator which will automatically // Note: Buffers are safely tied to the GpuArenaAllocator which will automatically
// release them at the end. You can also manually call buf_x.deinit() if desired. // release them at the end. You can also manually call buf_x.deinit() if desired.
@@ -79,7 +79,7 @@ pub fn main(init: std.process.Init) !void {
try buf_b.load(f16, data_b); try buf_b.load(f16, data_b);
// 7. Dispatch the Compute // 7. Dispatch the Compute
try add_cp.run(glloc, .{ buf_a, buf_b, buf_out }); try add_cp.run(gloc, .{ buf_a, buf_b, buf_out });
// 8. Map and copy the resulting buffer back to the CPU // 8. Map and copy the resulting buffer back to the CPU
const out = try buf_out.read(allocator, f16); const out = try buf_out.read(allocator, f16);
@@ -119,18 +119,18 @@ pub fn main(init: std.process.Init) !void {
// 2. Init VRAM Arena // 2. Init VRAM Arena
var grena = GpuArenaAllocator.init(allocator, device.gpuAllocator()); var grena = GpuArenaAllocator.init(allocator, device.gpuAllocator());
defer grena.deinit(); defer grena.deinit();
const glloc = grena.gpuAllocator(); const gloc = grena.gpuAllocator();
// 3. Load Render Pipeline // 3. Load Render Pipeline
const circle_rp = try GpuRender.init( const circle_rp = try GpuRender.init(
glloc, gloc,
@embedFile("shaders/circle.wgsl"), @embedFile("shaders/circle.wgsl"),
.{ .bindings = &.{}, .texture_format = .RGBA8Unorm, .topology = .TriangleStrip }, .{ .bindings = &.{}, .texture_format = .RGBA8Unorm, .topology = .TriangleStrip },
); );
defer circle_rp.deinit(); defer circle_rp.deinit();
// 4. Create VRAM texture to render into // 4. Create VRAM texture to render into
const texture = try GpuTexture.init(glloc, .{ const texture = try GpuTexture.init(gloc, .{
.format = .RGBA8Unorm, .format = .RGBA8Unorm,
.size = .{ .width = width, .height = height, .depthOrArrayLayers = 1 }, .size = .{ .width = width, .height = height, .depthOrArrayLayers = 1 },
.usage = .initMany(&.{ .RenderAttachment, .CopySrc }), .usage = .initMany(&.{ .RenderAttachment, .CopySrc }),
@@ -138,14 +138,14 @@ pub fn main(init: std.process.Init) !void {
defer texture.deinit(); defer texture.deinit();
// 5. Create a view from texture // 5. Create a view from texture
const view = try GpuTextureView.init(glloc, texture, .{}); const view = try GpuTextureView.init(gloc, texture, .{});
defer view.deinit(); defer view.deinit();
// 6. Run the rendering pipeline // 6. Run the rendering pipeline
try circle_rp.draw(glloc, view, 4, .{}); try circle_rp.draw(gloc, view, 4, .{});
// 7. Load Texture into GpuBuffer // 7. Load Texture into GpuBuffer
const cpu_staging_cpu = try texture.buffCopy(glloc); const cpu_staging_cpu = try texture.buffCopy(gloc);
defer cpu_staging_cpu.deinit(); defer cpu_staging_cpu.deinit();
// 8. Read GpuBuffer to CPU // 8. Read GpuBuffer to CPU
@@ -223,7 +223,7 @@ No extra installation required. Automatically links against `d3d12`, `dxgi`, and
Add it to your `build.zig.zon`: Add it to your `build.zig.zon`:
```bash ```bash
zig fetch --save git+https://git.bouvais.lu/adrien/zig-wgpu#ref=0.2.1 zig fetch --save git+[https://git.bouvais.lu/adrien/zig-wgpu#ref=0.2.0](https://git.bouvais.lu/adrien/zig-wgpu)
``` ```
Then, expose it in your `build.zig`: Then, expose it in your `build.zig`:
+1 -1
View File
@@ -1,6 +1,6 @@
.{ .{
.name = .zig_wgpu, .name = .zig_wgpu,
.version = "0.2.3", .version = "0.2.0",
.fingerprint = 0x5d0e853acbc0c2c6, .fingerprint = 0x5d0e853acbc0c2c6,
.minimum_zig_version = "0.16.0", .minimum_zig_version = "0.16.0",
.dependencies = .{}, .dependencies = .{},
+16 -16
View File
@@ -13,10 +13,10 @@ const Vec = struct {
buf: GpuBuffer, buf: GpuBuffer,
len: usize, len: usize,
// Changed: glloc is passed by value (const) // Changed: gloc is passed by value (const)
pub fn initZero(glloc: GpuAllocator, len: usize) !Vec { pub fn initZero(gloc: GpuAllocator, len: usize) !Vec {
return .{ return .{
.buf = try GpuBuffer.init(glloc, .{ .buf = try GpuBuffer.init(gloc, .{
.size = len * @sizeOf(f16), .size = len * @sizeOf(f16),
.usage = .initMany(&.{ .Storage, .CopyDst, .CopySrc }), .usage = .initMany(&.{ .Storage, .CopyDst, .CopySrc }),
}), }),
@@ -24,9 +24,9 @@ const Vec = struct {
}; };
} }
// Changed: glloc is passed by value // Changed: gloc is passed by value
pub fn initLoad(glloc: GpuAllocator, data: []const f16) !Vec { pub fn initLoad(gloc: GpuAllocator, data: []const f16) !Vec {
var self = try initZero(glloc, data.len); var self = try initZero(gloc, data.len);
try self.load(data); // Direct access via the interface copy try self.load(data); // Direct access via the interface copy
return self; return self;
} }
@@ -40,18 +40,18 @@ const Vec = struct {
try self.buf.load(f16, data); try self.buf.load(f16, data);
} }
// Changed: glloc is passed by value instead of *GpuAllocator // Changed: gloc is passed by value instead of *GpuAllocator
pub fn run(self: Vec, glloc: GpuAllocator, other: Vec, process: GpuCompute) !Vec { pub fn run(self: Vec, gloc: GpuAllocator, other: Vec, process: GpuCompute) !Vec {
std.debug.assert(self.len == other.len); std.debug.assert(self.len == other.len);
const result = try Vec.initZero(glloc, self.len); const result = try Vec.initZero(gloc, self.len);
errdefer result.deinit(); errdefer result.deinit();
try process.run(glloc, .{ self.buf, other.buf, result.buf }); try process.run(gloc, .{ self.buf, other.buf, result.buf });
return result; return result;
} }
// Changed: glloc is passed by value instead of *GpuAllocator // Changed: gloc is passed by value instead of *GpuAllocator
pub fn read(self: Vec, alloc: std.mem.Allocator) ![]f16 { pub fn read(self: Vec, alloc: std.mem.Allocator) ![]f16 {
return self.buf.read(alloc, f16); return self.buf.read(alloc, f16);
} }
@@ -63,9 +63,9 @@ pub fn main(init: std.process.Init) !void {
var grena = GpuArenaAllocator.init(init.gpa, device.gpuAllocator()); var grena = GpuArenaAllocator.init(init.gpa, device.gpuAllocator());
defer grena.deinit(); defer grena.deinit();
const glloc = grena.gpuAllocator(); const gloc = grena.gpuAllocator();
const add_pip = try GpuCompute.init(glloc, @embedFile("shaders/add.wgsl"), .{ .bindings = &.{ const add_pip = try GpuCompute.init(gloc, @embedFile("shaders/add.wgsl"), .{ .bindings = &.{
.{ .element_size = @sizeOf(f16) }, .{ .element_size = @sizeOf(f16) },
.{ .element_size = @sizeOf(f16) }, .{ .element_size = @sizeOf(f16) },
.{ .element_size = @sizeOf(f16) }, .{ .element_size = @sizeOf(f16) },
@@ -120,9 +120,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(glloc, data_a); const a = try Vec.initLoad(gloc, data_a);
defer a.deinit(); defer a.deinit();
const b = try Vec.initLoad(glloc, 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));
@@ -132,7 +132,7 @@ 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(glloc, 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.
+6 -6
View File
@@ -20,18 +20,18 @@ pub fn main(init: std.process.Init) !void {
// 2. Init VRAM Arena // 2. Init VRAM Arena
var grena = GpuArenaAllocator.init(allocator, device.gpuAllocator()); var grena = GpuArenaAllocator.init(allocator, device.gpuAllocator());
defer grena.deinit(); defer grena.deinit();
const glloc = grena.gpuAllocator(); const gloc = grena.gpuAllocator();
// 3. Load Render Pipeline // 3. Load Render Pipeline
const circle_rp = try GpuRender.init( const circle_rp = try GpuRender.init(
glloc, gloc,
@embedFile("shaders/circle.wgsl"), @embedFile("shaders/circle.wgsl"),
.{ .bindings = &.{}, .texture_format = .RGBA8Unorm, .topology = .TriangleStrip }, .{ .bindings = &.{}, .texture_format = .RGBA8Unorm, .topology = .TriangleStrip },
); );
defer circle_rp.deinit(); defer circle_rp.deinit();
// 4. Create VRAM texture to render into // 4. Create VRAM texture to render into
const texture = try GpuTexture.init(glloc, .{ const texture = try GpuTexture.init(gloc, .{
.format = .RGBA8Unorm, .format = .RGBA8Unorm,
.size = .{ .width = width, .height = height, .depthOrArrayLayers = 1 }, .size = .{ .width = width, .height = height, .depthOrArrayLayers = 1 },
.usage = .initMany(&.{ .RenderAttachment, .CopySrc }), .usage = .initMany(&.{ .RenderAttachment, .CopySrc }),
@@ -39,14 +39,14 @@ pub fn main(init: std.process.Init) !void {
defer texture.deinit(); defer texture.deinit();
// 5. Create a view from texture // 5. Create a view from texture
const view = try GpuTextureView.init(glloc, texture, .{}); const view = try GpuTextureView.init(gloc, texture, .{});
defer view.deinit(); defer view.deinit();
// 6. Run the rendering pipeline // 6. Run the rendering pipeline
try circle_rp.draw(glloc, view, 4, .{}); try circle_rp.draw(gloc, view, 4, .{});
// 7. Load Texture into GpuBuffer // 7. Load Texture into GpuBuffer
const cpu_staging_cpu = try texture.buffCopy(glloc); const cpu_staging_cpu = try texture.buffCopy(gloc);
defer cpu_staging_cpu.deinit(); defer cpu_staging_cpu.deinit();
// 8. Read GpuBuffer to CPU // 8. Read GpuBuffer to CPU
+16 -9
View File
@@ -15,11 +15,11 @@ pub fn main(init: std.process.Init) !void {
// 2. Create a GPU Arena to manage VRAM // 2. Create a GPU Arena to manage VRAM
var grena = GpuArenaAllocator.init(allocator, device.gpuAllocator()); var grena = GpuArenaAllocator.init(allocator, device.gpuAllocator());
defer grena.deinit(); defer grena.deinit();
const glloc = grena.gpuAllocator(); const gloc = grena.gpuAllocator();
// 3. Load the WGSL compute pipeline // 3. Load the WGSL compute pipeline
const add_cp = try GpuCompute.init( const add_cp = try GpuCompute.init(
glloc, gloc,
@embedFile("shaders/add.wgsl"), @embedFile("shaders/add.wgsl"),
.{ .{
.label = "add", .label = "add",
@@ -32,7 +32,7 @@ pub fn main(init: std.process.Init) !void {
); );
// 4. Setup CPU data // 4. Setup CPU data
const len: usize = 16; const len: usize = 1024;
const data_a = try allocator.alloc(f16, len); const data_a = try allocator.alloc(f16, len);
defer allocator.free(data_a); defer allocator.free(data_a);
const data_b = try allocator.alloc(f16, len); const data_b = try allocator.alloc(f16, len);
@@ -45,9 +45,9 @@ pub fn main(init: std.process.Init) !void {
// 5. Initialize raw GPU Buffers // 5. Initialize raw GPU Buffers
const byte_size = len * @sizeOf(f16); const byte_size = len * @sizeOf(f16);
const buf_a = try GpuBuffer.init(glloc, .{ .label = "a", .size = byte_size, .usage = .initMany(&.{ .Storage, .CopyDst, .CopySrc }) }); const buf_a = try GpuBuffer.init(gloc, .{ .label = "a", .size = byte_size, .usage = .initMany(&.{ .Storage, .CopyDst, .CopySrc }) });
const buf_b = try GpuBuffer.init(glloc, .{ .label = "b", .size = byte_size, .usage = .initMany(&.{ .Storage, .CopyDst, .CopySrc }) }); const buf_b = try GpuBuffer.init(gloc, .{ .label = "b", .size = byte_size, .usage = .initMany(&.{ .Storage, .CopyDst, .CopySrc }) });
const buf_out = try GpuBuffer.init(glloc, .{ .label = "out", .size = byte_size, .usage = .initMany(&.{ .Storage, .CopyDst, .CopySrc }) }); const buf_out = try GpuBuffer.init(gloc, .{ .label = "out", .size = byte_size, .usage = .initMany(&.{ .Storage, .CopyDst, .CopySrc }) });
// Note: Buffers are safely tied to the GpuArenaAllocator which will automatically // Note: Buffers are safely tied to the GpuArenaAllocator which will automatically
// release them at the end. You can also manually call buf_x.deinit() if desired. // release them at the end. You can also manually call buf_x.deinit() if desired.
@@ -58,11 +58,18 @@ pub fn main(init: std.process.Init) !void {
try buf_b.load(f16, data_b); try buf_b.load(f16, data_b);
// 7. Dispatch the Compute // 7. Dispatch the Compute
try add_cp.run(glloc, .{ buf_a, buf_b, buf_out }); try add_cp.run(gloc, .{ buf_a, buf_b, buf_out });
// 8. Map and copy the resulting buffer back to the CPU // 8. Map and copy the resulting buffer back to the CPU
const out = try buf_out.read(allocator, f16); const staging = try GpuBuffer.init(gloc, .{
.size = byte_size,
.usage = .initMany(&.{ .MapRead, .CopyDst }),
});
defer staging.deinit();
try buf_out.copy(staging);
const out = try staging.read(allocator, f16);
defer allocator.free(out); defer allocator.free(out);
std.debug.print("Result: {any}\n", .{out}); std.debug.print("Result: {any}\n", .{out[0..@min(6, len)]});
} }
+8
View File
@@ -87,6 +87,14 @@ pub fn gpuAllocator(self: *@This()) GpuAllocator {
fn allocBuffer(ctx: *anyopaque, desc: c.WGPUBufferDescriptor) anyerror!c.WGPUBuffer { fn allocBuffer(ctx: *anyopaque, desc: c.WGPUBufferDescriptor) anyerror!c.WGPUBuffer {
const self: *@This() = @ptrCast(@alignCast(ctx)); const self: *@This() = @ptrCast(@alignCast(ctx));
try self.tracked_buffers.ensureTotalCapacity(self.tracked_buffers.count() + 1); try self.tracked_buffers.ensureTotalCapacity(self.tracked_buffers.count() + 1);
const bytes = desc.size;
if (bytes > self.child_allocator.device.limits.maxBufferSize)
return error.SingleBufferExceedsLimit;
if (bytes + self.allocated_vram_bytes > self.child_allocator.device.def.vram_bytes_limit)
return error.ExceedsVramBudget;
const raw = try self.child_allocator.allocBuffer(desc); const raw = try self.child_allocator.allocBuffer(desc);
self.tracked_buffers.putAssumeCapacity(raw, desc); self.tracked_buffers.putAssumeCapacity(raw, desc);
self.allocated_vram_bytes += desc.size; self.allocated_vram_bytes += desc.size;
+47 -39
View File
@@ -4,11 +4,10 @@ const GpuAllocator = @import("GpuAllocator.zig");
const svOpt = @import("utils.zig").svOpt; const svOpt = @import("utils.zig").svOpt;
raw: c.WGPUBuffer, raw: c.WGPUBuffer,
size: u64, gloc: GpuAllocator,
usage: c.WGPUBufferUsage, def: GpuBufferDef,
glloc: GpuAllocator,
const BufferUsage = enum(u64) { pub const GpuBufferUsage = enum(u64) {
None = 0x0000000000000000, None = 0x0000000000000000,
MapRead = 0x0000000000000001, MapRead = 0x0000000000000001,
MapWrite = 0x0000000000000002, MapWrite = 0x0000000000000002,
@@ -20,37 +19,40 @@ const BufferUsage = enum(u64) {
Storage = 0x0000000000000080, Storage = 0x0000000000000080,
Indirect = 0x0000000000000100, Indirect = 0x0000000000000100,
QueryResolve = 0x0000000000000200, QueryResolve = 0x0000000000000200,
fn enumSetToWGPUBufferUsage(set: std.EnumSet(GpuBufferUsage)) c.WGPUBufferUsage {
var use: u64 = 0;
var iter = set.iterator();
while (iter.next()) |flag| use |= @intFromEnum(flag);
return use;
}
}; };
const GpuBufferDef = struct { pub const GpuBufferDef = struct {
label: ?[]const u8 = null, label: ?[]const u8 = null,
size: u64, size: u64,
usage: std.EnumSet(BufferUsage), usage: std.EnumSet(GpuBufferUsage),
}; };
pub fn init(glloc: GpuAllocator, def: GpuBufferDef) !@This() { pub fn init(gloc: GpuAllocator, def: GpuBufferDef) !@This() {
var use: u64 = 0;
var iter = def.usage.iterator();
while (iter.next()) |flag| use |= @intFromEnum(flag);
// Automatically align the buffer size forward to a multiple of 4 bytes under the hood // Automatically align the buffer size forward to a multiple of 4 bytes under the hood
const aligned_size = std.mem.alignForward(u64, def.size, 4); const aligned_size = std.mem.alignForward(u64, def.size, 4);
const raw_handle = try glloc.allocBuffer(.{ const raw_handle = try gloc.allocBuffer(.{
.size = aligned_size, .size = aligned_size,
.usage = use, .usage = GpuBufferUsage.enumSetToWGPUBufferUsage(def.usage),
.label = svOpt(def.label), .label = svOpt(def.label),
}); });
return .{ return .{
.raw = raw_handle, .raw = raw_handle,
.size = aligned_size, .def = def,
.usage = use, .gloc = gloc,
.glloc = glloc,
}; };
} }
pub fn deinit(self: @This()) void { pub fn deinit(self: @This()) void {
self.glloc.freeBuffer(self.raw); self.gloc.freeBuffer(self.raw);
} }
pub fn getConstMappedRange(self: @This(), offset: u64, size: u64) ?*const anyopaque { pub fn getConstMappedRange(self: @This(), offset: u64, size: u64) ?*const anyopaque {
@@ -79,57 +81,46 @@ pub fn load(
) !void { ) !void {
const bytes = data.len * @sizeOf(T); const bytes = data.len * @sizeOf(T);
if (bytes == self.size) { if (bytes == self.def.size) {
// Aligned path: direct download // Aligned path: direct download
c.wgpuQueueWriteBuffer(self.glloc.device.queue, self.raw, 0, data.ptr, self.size); c.wgpuQueueWriteBuffer(self.gloc.device.queue, self.raw, 0, data.ptr, self.def.size);
} else { } else {
// Unaligned path: Split the write into an aligned chunk and a padded remainder // Unaligned path: Split the write into an aligned chunk and a padded remainder
// to support arbitrary lengths without any allocations or large stack arrays. // to support arbitrary lengths without any allocations or large stack arrays.
const aligned_part = (bytes / 4) * 4; const aligned_part = (bytes / 4) * 4;
if (aligned_part > 0) { if (aligned_part > 0) {
c.wgpuQueueWriteBuffer(self.glloc.device.queue, self.raw, 0, data.ptr, aligned_part); c.wgpuQueueWriteBuffer(self.gloc.device.queue, self.raw, 0, data.ptr, aligned_part);
} }
var remainder_buf: [4]u8 = .{ 0, 0, 0, 0 }; var remainder_buf: [4]u8 = .{ 0, 0, 0, 0 };
const data_bytes = std.mem.sliceAsBytes(data); const data_bytes = std.mem.sliceAsBytes(data);
@memcpy(remainder_buf[0 .. bytes - aligned_part], data_bytes[aligned_part..bytes]); @memcpy(remainder_buf[0 .. bytes - aligned_part], data_bytes[aligned_part..bytes]);
c.wgpuQueueWriteBuffer(self.glloc.device.queue, self.raw, aligned_part, &remainder_buf, 4); c.wgpuQueueWriteBuffer(self.gloc.device.queue, self.raw, aligned_part, &remainder_buf, 4);
} }
} }
/// GPU to CPU /// GPU to CPU
/// Buffer must have MapRead usage or returns error.BufferNotMappable.
pub fn read(self: @This(), alloc: std.mem.Allocator, T: type) ![]T { pub fn read(self: @This(), alloc: std.mem.Allocator, T: type) ![]T {
const out = try alloc.alloc(T, @divExact(self.size, @sizeOf(T))); if (!self.def.usage.contains(.MapRead)) return error.BufferNotMappable;
const staging = try init(self.glloc, .{ const out = try alloc.alloc(T, @divExact(self.def.size, @sizeOf(T)));
.size = self.size,
.usage = .initMany(&.{ .MapRead, .CopyDst }),
.label = "staging_read_buffer",
});
defer staging.deinit();
const enc = c.wgpuDeviceCreateCommandEncoder(self.glloc.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.glloc.device.queue, 1, &cmd);
var mapped = false; var mapped = false;
staging.mapAsync( self.mapAsync(
c.WGPUMapMode_Read, c.WGPUMapMode_Read,
0, 0,
self.size, self.def.size,
.{ .callback = onMapped, .userdata1 = &mapped }, .{ .callback = onMapped, .userdata1 = &mapped },
); );
while (!mapped) self.glloc.device.poll(); while (!mapped) self.gloc.device.poll();
const ptr: [*]const T = @ptrCast(@alignCast( const ptr: [*]const T = @ptrCast(@alignCast(
staging.getConstMappedRange(0, self.size), self.getConstMappedRange(0, self.def.size),
)); ));
@memcpy(out[0..out.len], ptr[0..out.len]); @memcpy(out[0..out.len], ptr[0..out.len]);
staging.unmap(); self.unmap();
return out; return out;
} }
@@ -143,3 +134,20 @@ fn onMapped(
const flag: *bool = @ptrCast(@alignCast(userdata1.?)); const flag: *bool = @ptrCast(@alignCast(userdata1.?));
flag.* = (status == c.WGPUMapAsyncStatus_Success); flag.* = (status == c.WGPUMapAsyncStatus_Success);
} }
/// GPU to GPU. Both buffers must be same size, src needs CopySrc, dst needs CopyDst.
pub fn copy(src: @This(), dst: @This()) !void {
if (src.def.size != dst.def.size) return error.SizeMismatch;
const copy_src: u64 = @intFromEnum(GpuBufferUsage.CopySrc);
const copy_dst: u64 = @intFromEnum(GpuBufferUsage.CopyDst);
if (@as(u64, GpuBufferUsage.enumSetToWGPUBufferUsage(src.def.usage)) & copy_src == 0) return error.SrcNotCopyable;
if (@as(u64, GpuBufferUsage.enumSetToWGPUBufferUsage(dst.def.usage)) & copy_dst == 0) return error.DstNotWritable;
const enc = c.wgpuDeviceCreateCommandEncoder(src.gloc.device.device, null) orelse return error.Encoder;
c.wgpuCommandEncoderCopyBufferToBuffer(enc, src.raw, 0, dst.raw, 0, src.def.size);
const cmd = c.wgpuCommandEncoderFinish(enc, null);
defer c.wgpuCommandEncoderRelease(enc);
defer c.wgpuCommandBufferRelease(cmd);
c.wgpuQueueSubmit(src.gloc.device.queue, 1, &cmd);
}
+18 -18
View File
@@ -22,40 +22,40 @@ pub const ComputeDef = struct {
}; };
pip: c.WGPUComputePipeline, pip: c.WGPUComputePipeline,
glloc: GpuAllocator, gloc: GpuAllocator,
def: ComputeDef, def: ComputeDef,
pub fn init(glloc: GpuAllocator, wgsl: []const u8, def: ComputeDef) !@This() { pub fn init(gloc: GpuAllocator, wgsl: []const u8, def: ComputeDef) !@This() {
var wgsl_src = c.WGPUShaderSourceWGSL{ var wgsl_src = c.WGPUShaderSourceWGSL{
.chain = .{ .sType = c.WGPUSType_ShaderSourceWGSL }, .chain = .{ .sType = c.WGPUSType_ShaderSourceWGSL },
.code = sv(wgsl), .code = sv(wgsl),
}; };
const shader = c.wgpuDeviceCreateShaderModule(glloc.device.device, &.{ const shader = c.wgpuDeviceCreateShaderModule(gloc.device.device, &.{
.nextInChain = @ptrCast(&wgsl_src), .nextInChain = @ptrCast(&wgsl_src),
}) orelse return error.Shader; }) orelse return error.Shader;
defer c.wgpuShaderModuleRelease(shader); defer c.wgpuShaderModuleRelease(shader);
const pip = try glloc.allocComputePipeline(.{ const pip = try gloc.allocComputePipeline(.{
.label = svOpt(def.label), .label = svOpt(def.label),
.compute = .{ .module = shader, .entryPoint = sv("main") }, .compute = .{ .module = shader, .entryPoint = sv("main") },
}); });
return .{ return .{
.glloc = glloc, .gloc = gloc,
.pip = pip, .pip = pip,
.def = def, .def = def,
}; };
} }
pub fn deinit(self: @This()) void { pub fn deinit(self: @This()) void {
self.glloc.freeComputePipeline(self.pip); self.gloc.freeComputePipeline(self.pip);
} }
/// Execute the compute pass with arbitrary buffer bindings via a tuple. /// Execute the compute pass with arbitrary buffer bindings via a tuple.
/// Example: `try proc.run(glloc, .{ buf_a, buf_b, buf_out });` /// Example: `try proc.run(gloc, .{ buf_a, buf_b, buf_out });`
pub fn run( pub fn run(
self: @This(), self: @This(),
glloc: GpuAllocator, gloc: GpuAllocator,
args: anytype, args: anytype,
) !void { ) !void {
const type_info = @typeInfo(@TypeOf(args)); const type_info = @typeInfo(@TypeOf(args));
@@ -74,7 +74,7 @@ pub fn run(
const buf = @field(args, field.name); const buf = @field(args, field.name);
const el_size = self.def.bindings[i].element_size; const el_size = self.def.bindings[i].element_size;
if (el_size > 0) { if (el_size > 0) {
elements_count = @intCast(buf.size / el_size); elements_count = @intCast(buf.def.size / el_size);
} }
} }
} }
@@ -85,7 +85,7 @@ pub fn run(
const el_size = self.def.bindings[i].element_size; const el_size = self.def.bindings[i].element_size;
if (el_size > 0) { if (el_size > 0) {
const expected_min_bytes = @as(u64, elements_count) * el_size; const expected_min_bytes = @as(u64, elements_count) * el_size;
if (buf.size < expected_min_bytes) if (buf.def.size < expected_min_bytes)
return error.BufferTooSmall; return error.BufferTooSmall;
} }
} }
@@ -103,7 +103,7 @@ pub fn run(
.binding = @intCast(i), .binding = @intCast(i),
.buffer = buf.raw, .buffer = buf.raw,
.offset = 0, .offset = 0,
.size = buf.size, // Size exposes the fully allocated length .size = buf.def.size, // Size exposes the fully allocated length
}; };
entry_count += 1; entry_count += 1;
} }
@@ -113,12 +113,12 @@ pub fn run(
defer if (info_buf) |b| b.deinit(); defer if (info_buf) |b| b.deinit();
if (self.def.append_info_buffer) { if (self.def.append_info_buffer) {
info_buf = try GpuBuffer.init(glloc, .{ info_buf = try GpuBuffer.init(gloc, .{
.size = @sizeOf(u32), .size = @sizeOf(u32),
.usage = .initMany(&.{ .Uniform, .CopyDst }), .usage = .initMany(&.{ .Uniform, .CopyDst }),
.label = "compute_info_buffer", .label = "compute_info_buffer",
}); });
c.wgpuQueueWriteBuffer(glloc.device.queue, info_buf.?.raw, 0, &elements_count, @sizeOf(u32)); c.wgpuQueueWriteBuffer(gloc.device.queue, info_buf.?.raw, 0, &elements_count, @sizeOf(u32));
entries_buf[entry_count] = .{ entries_buf[entry_count] = .{
.binding = @intCast(entry_count), .binding = @intCast(entry_count),
@@ -130,11 +130,11 @@ pub fn run(
} }
const entries = entries_buf[0..entry_count]; const entries = entries_buf[0..entry_count];
try submitPass(glloc, self.pip, entries, elements_count, self.def.workgroup_size, self.def.max_workgroups); try submitPass(gloc, self.pip, entries, elements_count, self.def.workgroup_size, self.def.max_workgroups);
} }
fn submitPass( fn submitPass(
glloc: GpuAllocator, gloc: GpuAllocator,
pipeline: c.WGPUComputePipeline, pipeline: c.WGPUComputePipeline,
entries: []const c.WGPUBindGroupEntry, entries: []const c.WGPUBindGroupEntry,
n: usize, n: usize,
@@ -146,14 +146,14 @@ fn submitPass(
const bgl = c.wgpuComputePipelineGetBindGroupLayout(pipeline, 0); const bgl = c.wgpuComputePipelineGetBindGroupLayout(pipeline, 0);
defer c.wgpuBindGroupLayoutRelease(bgl); defer c.wgpuBindGroupLayoutRelease(bgl);
const bg = c.wgpuDeviceCreateBindGroup(glloc.device.device, &.{ const bg = c.wgpuDeviceCreateBindGroup(gloc.device.device, &.{
.layout = bgl, .layout = bgl,
.entries = entries.ptr, .entries = entries.ptr,
.entryCount = entries.len, .entryCount = entries.len,
}) orelse return error.BindGroup; }) orelse return error.BindGroup;
defer c.wgpuBindGroupRelease(bg); defer c.wgpuBindGroupRelease(bg);
const enc = c.wgpuDeviceCreateCommandEncoder(glloc.device.device, null) orelse return error.Encoder; const enc = c.wgpuDeviceCreateCommandEncoder(gloc.device.device, null) orelse return error.Encoder;
const pass = c.wgpuCommandEncoderBeginComputePass(enc, null); const pass = c.wgpuCommandEncoderBeginComputePass(enc, null);
c.wgpuComputePassEncoderSetPipeline(pass, pipeline); c.wgpuComputePassEncoderSetPipeline(pass, pipeline);
c.wgpuComputePassEncoderSetBindGroup(pass, 0, bg, 0, null); c.wgpuComputePassEncoderSetBindGroup(pass, 0, bg, 0, null);
@@ -168,7 +168,7 @@ fn submitPass(
const cmd = c.wgpuCommandEncoderFinish(enc, null); const cmd = c.wgpuCommandEncoderFinish(enc, null);
defer c.wgpuCommandEncoderRelease(enc); defer c.wgpuCommandEncoderRelease(enc);
defer c.wgpuCommandBufferRelease(cmd); defer c.wgpuCommandBufferRelease(cmd);
c.wgpuQueueSubmit(glloc.device.queue, 1, &cmd); c.wgpuQueueSubmit(gloc.device.queue, 1, &cmd);
} }
fn ceilDiv(n: usize, d: usize) usize { fn ceilDiv(n: usize, d: usize) usize {
+10 -10
View File
@@ -34,16 +34,16 @@ const GpuPrimitiveTopology = enum(c_uint) {
Force32 = 0x7FFFFFFF, Force32 = 0x7FFFFFFF,
}; };
glloc: GpuAllocator, gloc: GpuAllocator,
pip: c.WGPURenderPipeline, pip: c.WGPURenderPipeline,
def: GpuRenderDef, def: GpuRenderDef,
pub fn init(glloc: GpuAllocator, wgsl: []const u8, def: GpuRenderDef) !@This() { pub fn init(gloc: GpuAllocator, wgsl: []const u8, def: GpuRenderDef) !@This() {
var wgsl_src = c.WGPUShaderSourceWGSL{ var wgsl_src = c.WGPUShaderSourceWGSL{
.chain = .{ .sType = c.WGPUSType_ShaderSourceWGSL }, .chain = .{ .sType = c.WGPUSType_ShaderSourceWGSL },
.code = sv(wgsl), .code = sv(wgsl),
}; };
const shader = c.wgpuDeviceCreateShaderModule(glloc.device.device, &.{ const shader = c.wgpuDeviceCreateShaderModule(gloc.device.device, &.{
.nextInChain = @ptrCast(&wgsl_src), .nextInChain = @ptrCast(&wgsl_src),
}) orelse return error.Shader; }) orelse return error.Shader;
defer c.wgpuShaderModuleRelease(shader); defer c.wgpuShaderModuleRelease(shader);
@@ -69,7 +69,7 @@ pub fn init(glloc: GpuAllocator, wgsl: []const u8, def: GpuRenderDef) !@This() {
}; };
// 3. Compile the Complete Render Pipeline // 3. Compile the Complete Render Pipeline
const pip = try glloc.allocRenderPipeline(.{ const pip = try gloc.allocRenderPipeline(.{
.label = svOpt(def.label), .label = svOpt(def.label),
.vertex = .{ .vertex = .{
.module = shader, .module = shader,
@@ -90,21 +90,21 @@ pub fn init(glloc: GpuAllocator, wgsl: []const u8, def: GpuRenderDef) !@This() {
}); });
return .{ return .{
.glloc = glloc, .gloc = gloc,
.pip = pip, .pip = pip,
.def = def, .def = def,
}; };
} }
pub fn deinit(self: @This()) void { pub fn deinit(self: @This()) void {
self.glloc.freeRenderPipeline(self.pip); self.gloc.freeRenderPipeline(self.pip);
} }
/// Execute the render pass targeting a specific frame texture view. /// Execute the render pass targeting a specific frame texture view.
/// Passes bind groups via a tuple exactly like your original compute setup. /// Passes bind groups via a tuple exactly like your original compute setup.
pub fn draw( pub fn draw(
self: @This(), self: @This(),
glloc: GpuAllocator, gloc: GpuAllocator,
target_view: GpuTextureView, target_view: GpuTextureView,
vertex_count: u32, vertex_count: u32,
args: anytype, args: anytype,
@@ -138,7 +138,7 @@ pub fn draw(
const bgl = c.wgpuRenderPipelineGetBindGroupLayout(self.pip, 0); const bgl = c.wgpuRenderPipelineGetBindGroupLayout(self.pip, 0);
defer c.wgpuBindGroupLayoutRelease(bgl); defer c.wgpuBindGroupLayoutRelease(bgl);
const bg = c.wgpuDeviceCreateBindGroup(glloc.device.device, &.{ const bg = c.wgpuDeviceCreateBindGroup(gloc.device.device, &.{
.layout = bgl, .layout = bgl,
.entries = entries.ptr, .entries = entries.ptr,
.entryCount = @intCast(entries.len), .entryCount = @intCast(entries.len),
@@ -146,7 +146,7 @@ pub fn draw(
defer c.wgpuBindGroupRelease(bg); defer c.wgpuBindGroupRelease(bg);
// Encode Render Command // Encode Render Command
const enc = c.wgpuDeviceCreateCommandEncoder(glloc.device.device, null) orelse return error.Encoder; const enc = c.wgpuDeviceCreateCommandEncoder(gloc.device.device, null) orelse return error.Encoder;
defer c.wgpuCommandEncoderRelease(enc); defer c.wgpuCommandEncoderRelease(enc);
const color_attachment = c.WGPURenderPassColorAttachment{ const color_attachment = c.WGPURenderPassColorAttachment{
@@ -180,5 +180,5 @@ pub fn draw(
const cmd = c.wgpuCommandEncoderFinish(enc, null); const cmd = c.wgpuCommandEncoderFinish(enc, null);
defer c.wgpuCommandBufferRelease(cmd); defer c.wgpuCommandBufferRelease(cmd);
c.wgpuQueueSubmit(glloc.device.queue, 1, &cmd); c.wgpuQueueSubmit(gloc.device.queue, 1, &cmd);
} }
+14 -14
View File
@@ -14,10 +14,10 @@ pub const GpuTextureDef = struct {
}; };
raw: c.WGPUTexture, raw: c.WGPUTexture,
glloc: GpuAllocator, gloc: GpuAllocator,
def: GpuTextureDef, def: GpuTextureDef,
pub fn init(glloc: GpuAllocator, def: GpuTextureDef) !@This() { pub fn init(gloc: GpuAllocator, def: GpuTextureDef) !@This() {
var use: u64 = 0; var use: u64 = 0;
var iter = def.usage.iterator(); var iter = def.usage.iterator();
while (iter.next()) |flag| use |= @intFromEnum(flag); while (iter.next()) |flag| use |= @intFromEnum(flag);
@@ -31,13 +31,13 @@ pub fn init(glloc: GpuAllocator, def: GpuTextureDef) !@This() {
.mipLevelCount = 1, .mipLevelCount = 1,
.sampleCount = 1, .sampleCount = 1,
}; };
const raw = try glloc.allocTexture(desc); const raw = try gloc.allocTexture(desc);
return .{ .glloc = glloc, .raw = raw, .def = def }; return .{ .gloc = gloc, .raw = raw, .def = def };
} }
pub fn deinit(self: @This()) void { pub fn deinit(self: @This()) void {
self.glloc.freeTexture(self.raw); self.gloc.freeTexture(self.raw);
} }
pub fn getConstMappedRange(self: @This(), offset: u64, size: u64) ?*const anyopaque { pub fn getConstMappedRange(self: @This(), offset: u64, size: u64) ?*const anyopaque {
@@ -53,14 +53,14 @@ pub fn bytesSizeRow(self: @This()) u32 {
} }
/// Return a GpuBuffer containing a copy of the texture. /// Return a GpuBuffer containing a copy of the texture.
pub fn buffCopy(self: @This(), glloc: GpuAllocator) !GpuBuffer { pub fn buffCopy(self: @This(), gloc: GpuAllocator) !GpuBuffer {
const buf = try GpuBuffer.init(glloc, .{ const buf = try GpuBuffer.init(gloc, .{
.size = self.bytesSize(), .size = self.bytesSize(),
.usage = .initMany(&.{ .CopyDst, .CopySrc }), .usage = .initMany(&.{ .CopyDst, .CopySrc }),
.label = "texture_copy_buffer", .label = "texture_copy_buffer",
}); });
const enc = c.wgpuDeviceCreateCommandEncoder(glloc.device.device, null) orelse return error.Encoder; const enc = c.wgpuDeviceCreateCommandEncoder(gloc.device.device, null) orelse return error.Encoder;
defer c.wgpuCommandEncoderRelease(enc); defer c.wgpuCommandEncoderRelease(enc);
const src_copy = c.WGPUTexelCopyTextureInfo{ const src_copy = c.WGPUTexelCopyTextureInfo{
@@ -82,7 +82,7 @@ pub fn buffCopy(self: @This(), glloc: GpuAllocator) !GpuBuffer {
const cmd = c.wgpuCommandEncoderFinish(enc, null); const cmd = c.wgpuCommandEncoderFinish(enc, null);
defer c.wgpuCommandBufferRelease(cmd); defer c.wgpuCommandBufferRelease(cmd);
c.wgpuQueueSubmit(glloc.device.queue, 1, &cmd); c.wgpuQueueSubmit(gloc.device.queue, 1, &cmd);
return buf; return buf;
} }
@@ -110,7 +110,7 @@ pub fn load(
const bytes = data.len * @sizeOf(T); const bytes = data.len * @sizeOf(T);
c.wgpuQueueWriteTexture( c.wgpuQueueWriteTexture(
self.glloc.device.queue, self.gloc.device.queue,
&.{ &.{
.texture = self.raw, .texture = self.raw,
.mipLevel = 0, .mipLevel = 0,
@@ -132,14 +132,14 @@ pub fn load(
pub fn read(self: @This(), alloc: std.mem.Allocator, T: type) ![]T { pub fn read(self: @This(), alloc: std.mem.Allocator, T: type) ![]T {
const out = try alloc.alloc(T, @divExact(self.size, @sizeOf(T))); const out = try alloc.alloc(T, @divExact(self.size, @sizeOf(T)));
const staging = try init(self.glloc, .{ const staging = try init(self.gloc, .{
.size = self.size, .size = self.size,
.usage = .initMany(&.{ .MapRead, .CopyDst }), .usage = .initMany(&.{ .MapRead, .CopyDst }),
.label = "texture_read_staging", .label = "texture_read_staging",
}); });
defer staging.deinit(); defer staging.deinit();
const enc = c.wgpuDeviceCreateCommandEncoder(self.glloc.device.device, null) orelse return error.Encoder; const enc = c.wgpuDeviceCreateCommandEncoder(self.gloc.device.device, null) orelse return error.Encoder;
const src_copy = c.WGPUTexelCopyTextureInfo{ const src_copy = c.WGPUTexelCopyTextureInfo{
.texture = self.raw, .texture = self.raw,
.mipLevel = 0, .mipLevel = 0,
@@ -158,7 +158,7 @@ pub fn read(self: @This(), alloc: std.mem.Allocator, T: type) ![]T {
const cmd = c.wgpuCommandEncoderFinish(enc, null); const cmd = c.wgpuCommandEncoderFinish(enc, null);
defer c.wgpuCommandEncoderRelease(enc); defer c.wgpuCommandEncoderRelease(enc);
defer c.wgpuCommandBufferRelease(cmd); defer c.wgpuCommandBufferRelease(cmd);
c.wgpuQueueSubmit(self.glloc.device.queue, 1, &cmd); c.wgpuQueueSubmit(self.gloc.device.queue, 1, &cmd);
var mapped = false; var mapped = false;
staging.mapAsync( staging.mapAsync(
@@ -167,7 +167,7 @@ pub fn read(self: @This(), alloc: std.mem.Allocator, T: type) ![]T {
self.size, self.size,
.{ .callback = onMapped, .userdata1 = &mapped }, .{ .callback = onMapped, .userdata1 = &mapped },
); );
while (!mapped) self.glloc.device.poll(); while (!mapped) self.gloc.device.poll();
const ptr: [*]const T = @ptrCast(@alignCast( const ptr: [*]const T = @ptrCast(@alignCast(
staging.getConstMappedRange(0, self.size), staging.getConstMappedRange(0, self.size),
+5 -5
View File
@@ -13,23 +13,23 @@ pub const GpuViewDef = struct {
}; };
raw: c.WGPUTextureView, raw: c.WGPUTextureView,
glloc: GpuAllocator, gloc: GpuAllocator,
pub fn init(glloc: GpuAllocator, texture: GpuTexture, def: GpuViewDef) !@This() { pub fn init(gloc: GpuAllocator, texture: GpuTexture, def: GpuViewDef) !@This() {
var use: u64 = 0; var use: u64 = 0;
var iter = def.usage.iterator(); var iter = def.usage.iterator();
while (iter.next()) |flag| use |= @intFromEnum(flag); while (iter.next()) |flag| use |= @intFromEnum(flag);
const raw = try glloc.allocTextureView(texture.raw, .{ const raw = try gloc.allocTextureView(texture.raw, .{
.label = svOpt(def.label), .label = svOpt(def.label),
.format = @intFromEnum(def.format), .format = @intFromEnum(def.format),
.usage = use, .usage = use,
.mipLevelCount = 1, .mipLevelCount = 1,
.arrayLayerCount = 1, .arrayLayerCount = 1,
}); });
return .{ .glloc = glloc, .raw = raw }; return .{ .gloc = gloc, .raw = raw };
} }
pub fn deinit(self: @This()) void { pub fn deinit(self: @This()) void {
self.glloc.freeTextureView(self.raw); self.gloc.freeTextureView(self.raw);
} }