5 Commits
14 changed files with 334 additions and 194 deletions
+32 -24
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 that utilizes a staging buffer 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 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,17 +36,20 @@ 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 gloc = grena.gpuAllocator(); const glloc = 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(
gloc, glloc,
@embedFile("shaders/add.wgsl"), @embedFile("shaders/add.wgsl"),
.{ .bindings = &.{ .{
.label = "add",
.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) },
} }, },
},
); );
// 4. Setup CPU data // 4. Setup CPU data
@@ -63,19 +66,20 @@ 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(gloc, byte_size, .initMany(&.{ .Storage, .CopyDst, .CopySrc })); const buf_a = try GpuBuffer.init(glloc, .{ .label = "a", .size = byte_size, .usage = .initMany(&.{ .Storage, .CopyDst, .CopySrc }) });
const buf_b = try GpuBuffer.init(gloc, byte_size, .initMany(&.{ .Storage, .CopyDst, .CopySrc })); const buf_b = try GpuBuffer.init(glloc, .{ .label = "b", .size = byte_size, .usage = .initMany(&.{ .Storage, .CopyDst, .CopySrc }) });
const buf_out = try GpuBuffer.init(gloc, byte_size, .initMany(&.{ .Storage, .CopyDst, .CopySrc })); const buf_out = try GpuBuffer.init(glloc, .{ .label = "out", .size = byte_size, .usage = .initMany(&.{ .Storage, .CopyDst, .CopySrc }) });
// Note: Buffers, pipelines, and other objects initialized with 'gloc' // Note: Buffers are safely tied to the GpuArenaAllocator which will automatically
// are safely tied to the GpuArenaAllocator and will automatically release. // release them at the end. You can also manually call buf_x.deinit() if desired.
// This will also release pipelines, textures, ect. Everything using a GpuAllocator to init.
// 6. Transfer data from CPU slices to GPU Buffers // 6. Transfer data from CPU slices to GPU Buffers
try buf_a.load(f16, data_a); try buf_a.load(f16, data_a);
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(gloc, .{ buf_a, buf_b, buf_out }); try add_cp.run(glloc, .{ 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);
@@ -96,6 +100,7 @@ and pull the frame pixels back to the CPU to write a standard image file:
const std = @import("std"); const std = @import("std");
const gpu = @import("gpu"); const gpu = @import("gpu");
const GpuDevice = gpu.GpuDevice; const GpuDevice = gpu.GpuDevice;
const GpuArenaAllocator = gpu.GpuArenaAllocator;
const GpuBuffer = gpu.GpuBuffer; const GpuBuffer = gpu.GpuBuffer;
const GpuRender = gpu.GpuRender; const GpuRender = gpu.GpuRender;
const GpuTexture = gpu.GpuTexture; const GpuTexture = gpu.GpuTexture;
@@ -111,19 +116,21 @@ pub fn main(init: std.process.Init) !void {
const device = try GpuDevice.init(.{}); const device = try GpuDevice.init(.{});
defer device.deinit(); defer device.deinit();
// 2. Get base device GPU Allocator // 2. Init VRAM Arena
const gloc = device.gpuAllocator(); var grena = GpuArenaAllocator.init(allocator, device.gpuAllocator());
defer grena.deinit();
const glloc = grena.gpuAllocator();
// 3. Load Render Pipeline // 3. Load Render Pipeline
const circle_rp = try GpuRender.init( const circle_rp = try GpuRender.init(
gloc, glloc,
@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(gloc, .{ const texture = try GpuTexture.init(glloc, .{
.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 }),
@@ -131,21 +138,22 @@ 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(gloc, texture, .{}); const view = try GpuTextureView.init(glloc, texture, .{});
defer view.deinit(); defer view.deinit();
// 6. Run the rendering pipeline // 6. Run the rendering pipeline
try circle_rp.draw(gloc, view, 4, .{}); try circle_rp.draw(glloc, view, 4, .{});
// 7. Copy Texture into a readable GPU staging buffer // 7. Load Texture into GpuBuffer
const cpu_staging_buf = try texture.buffCopy(gloc); const cpu_staging_cpu = try texture.buffCopy(glloc);
defer cpu_staging_buf.deinit(); defer cpu_staging_cpu.deinit();
// 8. Read GpuBuffer to CPU memory // 8. Read GpuBuffer to CPU
const pixels = try cpu_staging_buf.read(allocator, u8); // This need to be free manually because CPU memory
const pixels = try cpu_staging_cpu.read(allocator, u8);
defer allocator.free(pixels); defer allocator.free(pixels);
// 9. Write out to a simple PPM image // 9. Write a simple ppm image
try savePpm(init.io, "circle.ppm", width, height, pixels); try savePpm(init.io, "circle.ppm", width, height, pixels);
} }
@@ -215,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.0](https://git.bouvais.lu/adrien/zig-wgpu) zig fetch --save git+https://git.bouvais.lu/adrien/zig-wgpu#ref=0.2.1
``` ```
Then, expose it in your `build.zig`: Then, expose it in your `build.zig`:
+2
View File
@@ -29,6 +29,7 @@ pub fn build(b: *std.Build) !void {
mod.linkSystemLibrary("gcc_s", .{}); mod.linkSystemLibrary("gcc_s", .{});
} }
if (b.pkg_hash.len == 0) {
var threaded: std.Io.Threaded = .init_single_threaded; var threaded: std.Io.Threaded = .init_single_threaded;
const io = threaded.io(); const io = threaded.io();
@@ -56,4 +57,5 @@ pub fn build(b: *std.Build) !void {
const run_cmd = b.addRunArtifact(exe); const run_cmd = b.addRunArtifact(exe);
run_step.dependOn(&run_cmd.step); run_step.dependOn(&run_cmd.step);
} }
}
} }
+1 -1
View File
@@ -1,6 +1,6 @@
.{ .{
.name = .zig_wgpu, .name = .zig_wgpu,
.version = "0.2.0", .version = "0.2.3",
.fingerprint = 0x5d0e853acbc0c2c6, .fingerprint = 0x5d0e853acbc0c2c6,
.minimum_zig_version = "0.16.0", .minimum_zig_version = "0.16.0",
.dependencies = .{}, .dependencies = .{},
+21 -20
View File
@@ -6,26 +6,27 @@ const GpuAllocator = gpu.GpuAllocator;
const GpuBuffer = gpu.GpuBuffer; const GpuBuffer = gpu.GpuBuffer;
const GpuCompute = gpu.GpuCompute; const GpuCompute = gpu.GpuCompute;
pub const std_options = std.Options{ .log_level = .info };
/// Minimal implementation of a f16 Vector /// Minimal implementation of a f16 Vector
const Vec = struct { const Vec = struct {
buf: GpuBuffer, buf: GpuBuffer,
len: usize, len: usize,
// Changed: gloc is passed by value (const) // Changed: glloc is passed by value (const)
pub fn initZero(gloc: GpuAllocator, len: usize) !Vec { pub fn initZero(glloc: GpuAllocator, len: usize) !Vec {
return .{ return .{
.buf = try GpuBuffer.init( .buf = try GpuBuffer.init(glloc, .{
gloc, .size = len * @sizeOf(f16),
len * @sizeOf(f16), .usage = .initMany(&.{ .Storage, .CopyDst, .CopySrc }),
.initMany(&.{ .Storage, .CopyDst, .CopySrc }), }),
),
.len = len, .len = len,
}; };
} }
// Changed: gloc is passed by value // Changed: glloc is passed by value
pub fn initLoad(gloc: GpuAllocator, data: []const f16) !Vec { pub fn initLoad(glloc: GpuAllocator, data: []const f16) !Vec {
var self = try initZero(gloc, data.len); var self = try initZero(glloc, 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;
} }
@@ -39,18 +40,18 @@ const Vec = struct {
try self.buf.load(f16, data); try self.buf.load(f16, data);
} }
// Changed: gloc is passed by value instead of *GpuAllocator // Changed: glloc is passed by value instead of *GpuAllocator
pub fn run(self: Vec, gloc: GpuAllocator, other: Vec, process: GpuCompute) !Vec { pub fn run(self: Vec, glloc: 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(gloc, self.len); const result = try Vec.initZero(glloc, self.len);
errdefer result.deinit(); errdefer result.deinit();
try process.run(gloc, .{ self.buf, other.buf, result.buf }); try process.run(glloc, .{ self.buf, other.buf, result.buf });
return result; return result;
} }
// Changed: gloc is passed by value instead of *GpuAllocator // Changed: glloc 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);
} }
@@ -62,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 gloc = grena.gpuAllocator(); const glloc = grena.gpuAllocator();
const add_pip = try GpuCompute.init(gloc, @embedFile("shaders/add.wgsl"), .{ .bindings = &.{ const add_pip = try GpuCompute.init(glloc, @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) },
@@ -119,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(gloc, data_a); const a = try Vec.initLoad(glloc, data_a);
defer a.deinit(); defer a.deinit();
const b = try Vec.initLoad(gloc, data_b); const b = try Vec.initLoad(glloc, 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));
@@ -131,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(gloc, b, add_pip); const sum = try a.run(glloc, 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.
+8 -6
View File
@@ -18,18 +18,20 @@ pub fn main(init: std.process.Init) !void {
defer device.deinit(); defer device.deinit();
// 2. Init VRAM Arena // 2. Init VRAM Arena
const gloc = device.gpuAllocator(); var grena = GpuArenaAllocator.init(allocator, device.gpuAllocator());
defer grena.deinit();
const glloc = grena.gpuAllocator();
// 3. Load Render Pipeline // 3. Load Render Pipeline
const circle_rp = try GpuRender.init( const circle_rp = try GpuRender.init(
gloc, glloc,
@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(gloc, .{ const texture = try GpuTexture.init(glloc, .{
.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 }),
@@ -37,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(gloc, texture, .{}); const view = try GpuTextureView.init(glloc, texture, .{});
defer view.deinit(); defer view.deinit();
// 6. Run the rendering pipeline // 6. Run the rendering pipeline
try circle_rp.draw(gloc, view, 4, .{}); try circle_rp.draw(glloc, view, 4, .{});
// 7. Load Texture into GpuBuffer // 7. Load Texture into GpuBuffer
const cpu_staging_cpu = try texture.buffCopy(gloc); const cpu_staging_cpu = try texture.buffCopy(glloc);
defer cpu_staging_cpu.deinit(); defer cpu_staging_cpu.deinit();
// 8. Read GpuBuffer to CPU // 8. Read GpuBuffer to CPU
+11 -8
View File
@@ -15,17 +15,20 @@ 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 gloc = grena.gpuAllocator(); const glloc = 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(
gloc, glloc,
@embedFile("shaders/add.wgsl"), @embedFile("shaders/add.wgsl"),
.{ .bindings = &.{ .{
.label = "add",
.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) },
} }, },
},
); );
// 4. Setup CPU data // 4. Setup CPU data
@@ -42,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(gloc, byte_size, .initMany(&.{ .Storage, .CopyDst, .CopySrc })); const buf_a = try GpuBuffer.init(glloc, .{ .label = "a", .size = byte_size, .usage = .initMany(&.{ .Storage, .CopyDst, .CopySrc }) });
const buf_b = try GpuBuffer.init(gloc, byte_size, .initMany(&.{ .Storage, .CopyDst, .CopySrc })); const buf_b = try GpuBuffer.init(glloc, .{ .label = "b", .size = byte_size, .usage = .initMany(&.{ .Storage, .CopyDst, .CopySrc }) });
const buf_out = try GpuBuffer.init(gloc, byte_size, .initMany(&.{ .Storage, .CopyDst, .CopySrc })); const buf_out = try GpuBuffer.init(glloc, .{ .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.
@@ -55,7 +58,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(gloc, .{ buf_a, buf_b, buf_out }); try add_cp.run(glloc, .{ 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);
+60 -9
View File
@@ -3,6 +3,7 @@ const GpuDevice = @import("GpuDevice.zig");
const GpuAllocator = @import("GpuAllocator.zig"); const GpuAllocator = @import("GpuAllocator.zig");
const GpuTextureFormat = @import("lib.zig").GpuTextureFormat; const GpuTextureFormat = @import("lib.zig").GpuTextureFormat;
const c = @import("utils.zig").c; const c = @import("utils.zig").c;
const viewStr = @import("utils.zig").viewStr;
child_allocator: GpuAllocator, // I use Zig naming child_allocator, but that should be a parent for me. Likely something idk child_allocator: GpuAllocator, // I use Zig naming child_allocator, but that should be a parent for me. Likely something idk
tracked_buffers: std.AutoHashMap(c.WGPUBuffer, c.WGPUBufferDescriptor), tracked_buffers: std.AutoHashMap(c.WGPUBuffer, c.WGPUBufferDescriptor),
@@ -24,30 +25,40 @@ pub fn init(cpu_allocator: std.mem.Allocator, child_allocator: GpuAllocator) @Th
} }
pub fn deinit(self: *@This()) void { pub fn deinit(self: *@This()) void {
std.log.debug("Freeing GpuArenaAllocator (Used VRAM: {}/{} MB)", .{
self.allocated_vram_bytes / 1024 / 1024,
self.child_allocator.device.def.vram_bytes_limit / 1024 / 1024,
});
var it_buffer = self.tracked_buffers.keyIterator(); var it_buffer = self.tracked_buffers.keyIterator();
while (it_buffer.next()) |buf_ptr| while (it_buffer.next()) |buf_ptr|
self.child_allocator.freeBuffer(buf_ptr.*); freeBuffer(self, buf_ptr.*);
self.tracked_buffers.deinit(); self.tracked_buffers.deinit();
var it_tex = self.tracked_textures.keyIterator(); var it_tex = self.tracked_textures.keyIterator();
while (it_tex.next()) |buf_ptr| while (it_tex.next()) |buf_ptr|
self.child_allocator.freeTexture(buf_ptr.*); freeTexture(self, buf_ptr.*);
self.tracked_textures.deinit(); self.tracked_textures.deinit();
var it_view = self.tracked_views.keyIterator(); var it_view = self.tracked_views.keyIterator();
while (it_view.next()) |buf_ptr| while (it_view.next()) |buf_ptr|
self.child_allocator.freeTextureView(buf_ptr.*); freeTextureView(self, buf_ptr.*);
self.tracked_views.deinit(); self.tracked_views.deinit();
var it_render = self.tracked_renders.keyIterator(); var it_render = self.tracked_renders.keyIterator();
while (it_render.next()) |buf_ptr| while (it_render.next()) |buf_ptr|
self.child_allocator.freeRenderPipeline(buf_ptr.*); freeRenderPipeline(self, buf_ptr.*);
self.tracked_renders.deinit(); self.tracked_renders.deinit();
var it_compute = self.tracked_computes.keyIterator(); var it_compute = self.tracked_computes.keyIterator();
while (it_compute.next()) |buf_ptr| while (it_compute.next()) |buf_ptr|
self.child_allocator.freeComputePipeline(buf_ptr.*); freeComputePipeline(self, buf_ptr.*);
self.tracked_computes.deinit(); self.tracked_computes.deinit();
std.log.debug("Freed GpuArenaAllocator (Used VRAM: {}/{} MB)", .{
self.allocated_vram_bytes / 1024 / 1024,
self.child_allocator.device.def.vram_bytes_limit / 1024 / 1024,
});
} }
/// Returns the type-erased immutable interface wrapper /// Returns the type-erased immutable interface wrapper
@@ -79,6 +90,14 @@ fn allocBuffer(ctx: *anyopaque, desc: c.WGPUBufferDescriptor) anyerror!c.WGPUBuf
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;
std.log.debug("Allocated Buffer '{s}': {d} B (Total VRAM: {}/{} MB)", .{
viewStr(desc.label),
desc.size,
self.allocated_vram_bytes / 1024 / 1024,
self.child_allocator.device.def.vram_bytes_limit / 1024 / 1024,
});
return raw; return raw;
} }
@@ -87,6 +106,12 @@ fn freeBuffer(ctx: *anyopaque, raw: c.WGPUBuffer) void {
if (self.tracked_buffers.fetchRemove(raw)) |kv| { if (self.tracked_buffers.fetchRemove(raw)) |kv| {
self.child_allocator.freeBuffer(raw); self.child_allocator.freeBuffer(raw);
self.allocated_vram_bytes -= kv.value.size; self.allocated_vram_bytes -= kv.value.size;
std.log.debug("Freed Buffer '{s}' (Total VRAM: {}/{} MB)", .{
viewStr(kv.value.label),
self.allocated_vram_bytes / 1024 / 1024,
self.child_allocator.device.def.vram_bytes_limit / 1024 / 1024,
});
} }
} }
@@ -97,13 +122,21 @@ fn allocTexture(ctx: *anyopaque, desc: c.WGPUTextureDescriptor) anyerror!c.WGPUT
const format: GpuTextureFormat = @enumFromInt(desc.format); const format: GpuTextureFormat = @enumFromInt(desc.format);
const bytes_size = desc.size.width * desc.size.height * format.bytesPerPixel(); const bytes_size = desc.size.width * desc.size.height * format.bytesPerPixel();
if (bytes_size + self.allocated_vram_bytes > self.child_allocator.device.config.vram_bytes_limit) if (bytes_size + self.allocated_vram_bytes > self.child_allocator.device.def.vram_bytes_limit)
return error.ExceedsVramBudget; return error.ExceedsVramBudget;
const raw = try self.child_allocator.allocTexture(desc); const raw = try self.child_allocator.allocTexture(desc);
self.tracked_textures.putAssumeCapacity(raw, desc); self.tracked_textures.putAssumeCapacity(raw, desc);
self.allocated_vram_bytes += bytes_size; self.allocated_vram_bytes += bytes_size;
std.log.debug("Allocated Texture '{s}': {d} B (Total VRAM: {}/{} MB)", .{
viewStr(desc.label),
bytes_size,
self.allocated_vram_bytes / 1024 / 1024,
self.child_allocator.device.def.vram_bytes_limit / 1024 / 1024,
});
return raw; return raw;
} }
@@ -117,6 +150,12 @@ fn freeTexture(ctx: *anyopaque, raw: c.WGPUTexture) void {
const format: GpuTextureFormat = @enumFromInt(desc.format); const format: GpuTextureFormat = @enumFromInt(desc.format);
const bytes_size = desc.size.width * desc.size.height * format.bytesPerPixel(); const bytes_size = desc.size.width * desc.size.height * format.bytesPerPixel();
self.allocated_vram_bytes -= bytes_size; self.allocated_vram_bytes -= bytes_size;
std.log.debug("Freed Texture '{s}' (Total VRAM: {}/{} MB)", .{
viewStr(desc.label),
self.allocated_vram_bytes / 1024 / 1024,
self.child_allocator.device.def.vram_bytes_limit / 1024 / 1024,
});
} }
} }
@@ -125,13 +164,17 @@ fn allocTextureView(ctx: *anyopaque, texture: c.WGPUTexture, desc: c.WGPUTexture
try self.tracked_views.ensureTotalCapacity(self.tracked_views.count() + 1); try self.tracked_views.ensureTotalCapacity(self.tracked_views.count() + 1);
const raw = try self.child_allocator.allocTextureView(texture, desc); const raw = try self.child_allocator.allocTextureView(texture, desc);
self.tracked_views.putAssumeCapacity(raw, desc); self.tracked_views.putAssumeCapacity(raw, desc);
std.log.debug("Allocated Texture View '{s}'", .{viewStr(desc.label)});
return raw; return raw;
} }
fn freeTextureView(ctx: *anyopaque, raw: c.WGPUTextureView) void { fn freeTextureView(ctx: *anyopaque, raw: c.WGPUTextureView) void {
const self: *@This() = @ptrCast(@alignCast(ctx)); const self: *@This() = @ptrCast(@alignCast(ctx));
if (self.tracked_views.remove(raw)) if (self.tracked_views.fetchRemove(raw)) |kv| {
self.child_allocator.freeTextureView(raw); self.child_allocator.freeTextureView(raw);
const desc = kv.value;
std.log.debug("Freed Texture View '{s}'", .{viewStr(desc.label)});
}
} }
fn allocRenderPipeline(ctx: *anyopaque, desc: c.WGPURenderPipelineDescriptor) anyerror!c.WGPURenderPipeline { fn allocRenderPipeline(ctx: *anyopaque, desc: c.WGPURenderPipelineDescriptor) anyerror!c.WGPURenderPipeline {
@@ -139,13 +182,17 @@ fn allocRenderPipeline(ctx: *anyopaque, desc: c.WGPURenderPipelineDescriptor) an
try self.tracked_renders.ensureTotalCapacity(self.tracked_renders.count() + 1); try self.tracked_renders.ensureTotalCapacity(self.tracked_renders.count() + 1);
const raw = try self.child_allocator.allocRenderPipeline(desc); const raw = try self.child_allocator.allocRenderPipeline(desc);
self.tracked_renders.putAssumeCapacity(raw, desc); self.tracked_renders.putAssumeCapacity(raw, desc);
std.log.debug("Allocated Render Pipeline '{s}'", .{viewStr(desc.label)});
return raw; return raw;
} }
fn freeRenderPipeline(ctx: *anyopaque, raw: c.WGPURenderPipeline) void { fn freeRenderPipeline(ctx: *anyopaque, raw: c.WGPURenderPipeline) void {
const self: *@This() = @ptrCast(@alignCast(ctx)); const self: *@This() = @ptrCast(@alignCast(ctx));
if (self.tracked_renders.remove(raw)) if (self.tracked_renders.fetchRemove(raw)) |kv| {
self.child_allocator.freeRenderPipeline(raw); self.child_allocator.freeRenderPipeline(raw);
const desc = kv.value;
std.log.debug("Freed Render Pipeline '{s}'", .{viewStr(desc.label)});
}
} }
fn allocComputePipeline(ctx: *anyopaque, desc: c.WGPUComputePipelineDescriptor) anyerror!c.WGPUComputePipeline { fn allocComputePipeline(ctx: *anyopaque, desc: c.WGPUComputePipelineDescriptor) anyerror!c.WGPUComputePipeline {
@@ -153,11 +200,15 @@ fn allocComputePipeline(ctx: *anyopaque, desc: c.WGPUComputePipelineDescriptor)
try self.tracked_computes.ensureTotalCapacity(self.tracked_computes.count() + 1); try self.tracked_computes.ensureTotalCapacity(self.tracked_computes.count() + 1);
const raw = try self.child_allocator.allocComputePipeline(desc); const raw = try self.child_allocator.allocComputePipeline(desc);
self.tracked_computes.putAssumeCapacity(raw, desc); self.tracked_computes.putAssumeCapacity(raw, desc);
std.log.debug("Allocated Compute Pipeline '{s}'", .{viewStr(desc.label)});
return raw; return raw;
} }
fn freeComputePipeline(ctx: *anyopaque, raw: c.WGPUComputePipeline) void { fn freeComputePipeline(ctx: *anyopaque, raw: c.WGPUComputePipeline) void {
const self: *@This() = @ptrCast(@alignCast(ctx)); const self: *@This() = @ptrCast(@alignCast(ctx));
if (self.tracked_computes.remove(raw)) if (self.tracked_computes.fetchRemove(raw)) |kv| {
self.child_allocator.freeComputePipeline(raw); self.child_allocator.freeComputePipeline(raw);
const desc = kv.value;
std.log.debug("Freed Compute Pipeline '{s}'", .{viewStr(desc.label)});
}
} }
+29 -18
View File
@@ -1,11 +1,12 @@
const std = @import("std"); const std = @import("std");
const c = @import("utils.zig").c; const c = @import("utils.zig").c;
const GpuAllocator = @import("GpuAllocator.zig"); const GpuAllocator = @import("GpuAllocator.zig");
const svOpt = @import("utils.zig").svOpt;
raw: c.WGPUBuffer, raw: c.WGPUBuffer,
size: u64, size: u64,
usage: c.WGPUBufferUsage, usage: c.WGPUBufferUsage,
gloc: GpuAllocator, glloc: GpuAllocator,
const BufferUsage = enum(u64) { const BufferUsage = enum(u64) {
None = 0x0000000000000000, None = 0x0000000000000000,
@@ -21,25 +22,35 @@ const BufferUsage = enum(u64) {
QueryResolve = 0x0000000000000200, QueryResolve = 0x0000000000000200,
}; };
pub fn init(gloc: GpuAllocator, size: u64, usage: std.EnumSet(BufferUsage)) !@This() { const GpuBufferDef = struct {
label: ?[]const u8 = null,
size: u64,
usage: std.EnumSet(BufferUsage),
};
pub fn init(glloc: GpuAllocator, def: GpuBufferDef) !@This() {
var use: u64 = 0; var use: u64 = 0;
var iter = usage.iterator(); var iter = def.usage.iterator();
while (iter.next()) |flag| use |= @intFromEnum(flag); 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, size, 4); const aligned_size = std.mem.alignForward(u64, def.size, 4);
const raw_handle = try gloc.allocBuffer(.{ .size = aligned_size, .usage = use }); const raw_handle = try glloc.allocBuffer(.{
.size = aligned_size,
.usage = use,
.label = svOpt(def.label),
});
return .{ return .{
.raw = raw_handle, .raw = raw_handle,
.size = aligned_size, .size = aligned_size,
.usage = use, .usage = use,
.gloc = gloc, .glloc = glloc,
}; };
} }
pub fn deinit(self: @This()) void { pub fn deinit(self: @This()) void {
self.gloc.freeBuffer(self.raw); self.glloc.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 {
@@ -70,20 +81,20 @@ pub fn load(
if (bytes == self.size) { if (bytes == self.size) {
// Aligned path: direct download // Aligned path: direct download
c.wgpuQueueWriteBuffer(self.gloc.device.queue, self.raw, 0, data.ptr, self.size); c.wgpuQueueWriteBuffer(self.glloc.device.queue, self.raw, 0, data.ptr, self.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.gloc.device.queue, self.raw, 0, data.ptr, aligned_part); c.wgpuQueueWriteBuffer(self.glloc.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.gloc.device.queue, self.raw, aligned_part, &remainder_buf, 4); c.wgpuQueueWriteBuffer(self.glloc.device.queue, self.raw, aligned_part, &remainder_buf, 4);
} }
} }
@@ -91,19 +102,19 @@ 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( const staging = try init(self.glloc, .{
self.gloc, .size = self.size,
self.size, .usage = .initMany(&.{ .MapRead, .CopyDst }),
.initMany(&.{ .MapRead, .CopyDst }), .label = "staging_read_buffer",
); });
defer staging.deinit(); defer staging.deinit();
const enc = c.wgpuDeviceCreateCommandEncoder(self.gloc.device.device, null) orelse return error.Encoder; const enc = c.wgpuDeviceCreateCommandEncoder(self.glloc.device.device, null) orelse return error.Encoder;
c.wgpuCommandEncoderCopyBufferToBuffer(enc, self.raw, 0, staging.raw, 0, self.size); c.wgpuCommandEncoderCopyBufferToBuffer(enc, self.raw, 0, staging.raw, 0, self.size);
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.gloc.device.queue, 1, &cmd); c.wgpuQueueSubmit(self.glloc.device.queue, 1, &cmd);
var mapped = false; var mapped = false;
staging.mapAsync( staging.mapAsync(
@@ -112,7 +123,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.gloc.device.poll(); while (!mapped) self.glloc.device.poll();
const ptr: [*]const T = @ptrCast(@alignCast( const ptr: [*]const T = @ptrCast(@alignCast(
staging.getConstMappedRange(0, self.size), staging.getConstMappedRange(0, self.size),
+24 -19
View File
@@ -1,5 +1,6 @@
const c = @import("utils.zig").c; const c = @import("utils.zig").c;
const sv = @import("utils.zig").sv; const sv = @import("utils.zig").sv;
const svOpt = @import("utils.zig").svOpt;
const GpuAllocator = @import("GpuAllocator.zig"); const GpuAllocator = @import("GpuAllocator.zig");
const GpuBuffer = @import("GpuBuffer.zig"); const GpuBuffer = @import("GpuBuffer.zig");
const GpuDevice = @import("GpuDevice.zig"); const GpuDevice = @import("GpuDevice.zig");
@@ -11,6 +12,7 @@ pub const Binding = struct {
}; };
pub const ComputeDef = struct { pub const ComputeDef = struct {
label: ?[]const u8 = null,
bindings: []const Binding, bindings: []const Binding,
workgroup_size: u32 = 256, workgroup_size: u32 = 256,
max_workgroups: u32 = 65535, max_workgroups: u32 = 65535,
@@ -20,37 +22,40 @@ pub const ComputeDef = struct {
}; };
pip: c.WGPUComputePipeline, pip: c.WGPUComputePipeline,
gloc: GpuAllocator, glloc: GpuAllocator,
def: ComputeDef, def: ComputeDef,
pub fn init(gloc: GpuAllocator, wgsl: []const u8, def: ComputeDef) !@This() { pub fn init(glloc: 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(gloc.device.device, &.{ const shader = c.wgpuDeviceCreateShaderModule(glloc.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 gloc.allocComputePipeline(.{ .compute = .{ .module = shader, .entryPoint = sv("main") } }); const pip = try glloc.allocComputePipeline(.{
.label = svOpt(def.label),
.compute = .{ .module = shader, .entryPoint = sv("main") },
});
return .{ return .{
.gloc = gloc, .glloc = glloc,
.pip = pip, .pip = pip,
.def = def, .def = def,
}; };
} }
pub fn deinit(self: @This()) void { pub fn deinit(self: @This()) void {
self.gloc.freeComputePipeline(self.pip); self.glloc.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(gloc, .{ buf_a, buf_b, buf_out });` /// Example: `try proc.run(glloc, .{ buf_a, buf_b, buf_out });`
pub fn run( pub fn run(
self: @This(), self: @This(),
gloc: GpuAllocator, glloc: GpuAllocator,
args: anytype, args: anytype,
) !void { ) !void {
const type_info = @typeInfo(@TypeOf(args)); const type_info = @typeInfo(@TypeOf(args));
@@ -108,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( info_buf = try GpuBuffer.init(glloc, .{
gloc, .size = @sizeOf(u32),
@sizeOf(u32), .usage = .initMany(&.{ .Uniform, .CopyDst }),
.initMany(&.{ .Uniform, .CopyDst }), .label = "compute_info_buffer",
); });
c.wgpuQueueWriteBuffer(gloc.device.queue, info_buf.?.raw, 0, &elements_count, @sizeOf(u32)); c.wgpuQueueWriteBuffer(glloc.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),
@@ -125,11 +130,11 @@ pub fn run(
} }
const entries = entries_buf[0..entry_count]; const entries = entries_buf[0..entry_count];
try submitPass(gloc, self.pip, entries, elements_count, self.def.workgroup_size, self.def.max_workgroups); try submitPass(glloc, self.pip, entries, elements_count, self.def.workgroup_size, self.def.max_workgroups);
} }
fn submitPass( fn submitPass(
gloc: GpuAllocator, glloc: GpuAllocator,
pipeline: c.WGPUComputePipeline, pipeline: c.WGPUComputePipeline,
entries: []const c.WGPUBindGroupEntry, entries: []const c.WGPUBindGroupEntry,
n: usize, n: usize,
@@ -141,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(gloc.device.device, &.{ const bg = c.wgpuDeviceCreateBindGroup(glloc.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(gloc.device.device, null) orelse return error.Encoder; const enc = c.wgpuDeviceCreateCommandEncoder(glloc.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);
@@ -163,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(gloc.device.queue, 1, &cmd); c.wgpuQueueSubmit(glloc.device.queue, 1, &cmd);
} }
fn ceilDiv(n: usize, d: usize) usize { fn ceilDiv(n: usize, d: usize) usize {
+27 -11
View File
@@ -1,20 +1,17 @@
const std = @import("std"); const std = @import("std");
const c = @import("utils.zig").c; const c = @import("utils.zig").c;
const sv = @import("utils.zig").sv; const sv = @import("utils.zig").sv;
const svOpt = @import("utils.zig").svOpt;
const GpuAllocator = @import("GpuAllocator.zig"); const GpuAllocator = @import("GpuAllocator.zig");
const GpuTextureFormat = @import("lib.zig").GpuTextureFormat; const GpuTextureFormat = @import("lib.zig").GpuTextureFormat;
// TODO: Make Allocator more zig like
// - GpuDevice can return a GpuAllocator that just allocate and nothing else
// - From this GpuAllocator, can create a GpuArena like std.heap.ArenaAllocator.init(allocator)
// - Rename GpuArenaAllocator too
const Ctx = struct { const Ctx = struct {
adapter: c.WGPUAdapter = null, adapter: c.WGPUAdapter = null,
device: c.WGPUDevice = null, device: c.WGPUDevice = null,
}; };
const GpuDeviceConfig = struct { const GpuDeviceDef = struct {
label: ?[]const u8 = null,
/// 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) { power_preference: enum(c_uint) {
@@ -31,9 +28,9 @@ device: c.WGPUDevice,
queue: c.WGPUQueue, queue: c.WGPUQueue,
limits: c.WGPULimits, limits: c.WGPULimits,
config: GpuDeviceConfig, def: GpuDeviceDef,
pub fn init(config: GpuDeviceConfig) !@This() { pub fn init(def: GpuDeviceDef) !@This() {
const instance = c.wgpuCreateInstance( const instance = c.wgpuCreateInstance(
&std.mem.zeroes(c.WGPUInstanceDescriptor), &std.mem.zeroes(c.WGPUInstanceDescriptor),
) orelse return error.NoInstance; ) orelse return error.NoInstance;
@@ -42,13 +39,25 @@ pub fn init(config: GpuDeviceConfig) !@This() {
var ctx = Ctx{}; var ctx = Ctx{};
_ = c.wgpuInstanceRequestAdapter( _ = c.wgpuInstanceRequestAdapter(
instance, instance,
&.{ .powerPreference = @intFromEnum(config.power_preference) }, &.{ .powerPreference = @intFromEnum(def.power_preference) },
.{ .callback = onAdapter, .userdata1 = &ctx }, .{ .callback = onAdapter, .userdata1 = &ctx },
); );
c.wgpuInstanceProcessEvents(instance); c.wgpuInstanceProcessEvents(instance);
const adapter = ctx.adapter orelse return error.NoAdapter; const adapter = ctx.adapter orelse return error.NoAdapter;
errdefer c.wgpuAdapterRelease(adapter); errdefer c.wgpuAdapterRelease(adapter);
var adapter_info = std.mem.zeroes(c.WGPUAdapterInfo);
_ = c.wgpuAdapterGetInfo(adapter, &adapter_info);
std.log.info("=== WebGPU Device Initialized ===", .{});
if (adapter_info.device.length > 0 and adapter_info.device.data != null) {
std.log.info(" Device Name : {s}", .{adapter_info.device.data[0..adapter_info.device.length]});
}
if (adapter_info.architecture.length > 0 and adapter_info.architecture.data != null) {
std.log.info(" Architecture : {s}", .{adapter_info.architecture.data[0..adapter_info.architecture.length]});
}
std.log.info(" Backend Type : {d}", .{adapter_info.backendType});
var supported_features = std.mem.zeroes(c.WGPUSupportedFeatures); var supported_features = std.mem.zeroes(c.WGPUSupportedFeatures);
c.wgpuAdapterGetFeatures(adapter, &supported_features); c.wgpuAdapterGetFeatures(adapter, &supported_features);
@@ -56,6 +65,11 @@ pub fn init(config: GpuDeviceConfig) !@This() {
supported_limits.nextInChain = null; supported_limits.nextInChain = null;
if (c.wgpuAdapterGetLimits(adapter, &supported_limits) != 1) return error.FailedToGetAdapterLimits; if (c.wgpuAdapterGetLimits(adapter, &supported_limits) != 1) return error.FailedToGetAdapterLimits;
std.log.info(" Max Buf Size : {d} MB", .{supported_limits.maxBufferSize / 1024 / 1024});
std.log.info(" Max Storage : {d} MB", .{supported_limits.maxStorageBufferBindingSize / 1024 / 1024});
std.log.info(" Max Workgroup: X: {d}, Y: {d}, Z: {d}", .{ supported_limits.maxComputeWorkgroupSizeX, supported_limits.maxComputeWorkgroupSizeY, supported_limits.maxComputeWorkgroupSizeZ });
std.log.info(" VRAM Budget : {d} MB", .{def.vram_bytes_limit / 1024 / 1024});
var has_f16 = false; var has_f16 = false;
for (0..supported_features.featureCount) |i| { for (0..supported_features.featureCount) |i| {
if (supported_features.features[i] == c.WGPUFeatureName_ShaderF16) { if (supported_features.features[i] == c.WGPUFeatureName_ShaderF16) {
@@ -63,6 +77,8 @@ pub fn init(config: GpuDeviceConfig) !@This() {
break; break;
} }
} }
std.log.info(" Shader F16 : {}", .{has_f16});
std.log.info("=================================", .{});
var feature_buf = [_]c.WGPUFeatureName{c.WGPUFeatureName_ShaderF16}; var feature_buf = [_]c.WGPUFeatureName{c.WGPUFeatureName_ShaderF16};
const required_features: []const c.WGPUFeatureName = const required_features: []const c.WGPUFeatureName =
@@ -70,7 +86,7 @@ pub fn init(config: GpuDeviceConfig) !@This() {
const device_descriptor = c.WGPUDeviceDescriptor{ const device_descriptor = c.WGPUDeviceDescriptor{
.nextInChain = null, .nextInChain = null,
.label = sv("TensorCompilerDevice"), .label = svOpt(def.label),
.requiredFeatureCount = required_features.len, .requiredFeatureCount = required_features.len,
.requiredFeatures = if (required_features.len > 0) required_features.ptr else null, .requiredFeatures = if (required_features.len > 0) required_features.ptr else null,
.requiredLimits = &supported_limits, .requiredLimits = &supported_limits,
@@ -89,7 +105,7 @@ pub fn init(config: GpuDeviceConfig) !@This() {
.device = device, .device = device,
.queue = c.wgpuDeviceGetQueue(device), .queue = c.wgpuDeviceGetQueue(device),
.limits = supported_limits, .limits = supported_limits,
.config = config, .def = def,
}; };
} }
+13 -10
View File
@@ -1,6 +1,7 @@
const std = @import("std"); const std = @import("std");
const c = @import("utils.zig").c; const c = @import("utils.zig").c;
const sv = @import("utils.zig").sv; const sv = @import("utils.zig").sv;
const svOpt = @import("utils.zig").svOpt;
const GpuAllocator = @import("GpuAllocator.zig"); const GpuAllocator = @import("GpuAllocator.zig");
const GpuBuffer = @import("GpuBuffer.zig"); const GpuBuffer = @import("GpuBuffer.zig");
const GpuDevice = @import("GpuDevice.zig"); const GpuDevice = @import("GpuDevice.zig");
@@ -12,6 +13,7 @@ pub const Binding = struct {
}; };
pub const GpuRenderDef = struct { pub const GpuRenderDef = struct {
label: ?[]const u8 = null,
bindings: []const Binding = &.{}, bindings: []const Binding = &.{},
/// The surface texture format we are rendering to (e.g., BGRA8Unorm) /// The surface texture format we are rendering to (e.g., BGRA8Unorm)
texture_format: GpuTextureFormat, texture_format: GpuTextureFormat,
@@ -32,16 +34,16 @@ const GpuPrimitiveTopology = enum(c_uint) {
Force32 = 0x7FFFFFFF, Force32 = 0x7FFFFFFF,
}; };
gloc: GpuAllocator, glloc: GpuAllocator,
pip: c.WGPURenderPipeline, pip: c.WGPURenderPipeline,
def: GpuRenderDef, def: GpuRenderDef,
pub fn init(gloc: GpuAllocator, wgsl: []const u8, def: GpuRenderDef) !@This() { pub fn init(glloc: 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(gloc.device.device, &.{ const shader = c.wgpuDeviceCreateShaderModule(glloc.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);
@@ -67,7 +69,8 @@ pub fn init(gloc: GpuAllocator, wgsl: []const u8, def: GpuRenderDef) !@This() {
}; };
// 3. Compile the Complete Render Pipeline // 3. Compile the Complete Render Pipeline
const pip = try gloc.allocRenderPipeline(.{ const pip = try glloc.allocRenderPipeline(.{
.label = svOpt(def.label),
.vertex = .{ .vertex = .{
.module = shader, .module = shader,
.entryPoint = sv(def.vertex_entry), .entryPoint = sv(def.vertex_entry),
@@ -87,21 +90,21 @@ pub fn init(gloc: GpuAllocator, wgsl: []const u8, def: GpuRenderDef) !@This() {
}); });
return .{ return .{
.gloc = gloc, .glloc = glloc,
.pip = pip, .pip = pip,
.def = def, .def = def,
}; };
} }
pub fn deinit(self: @This()) void { pub fn deinit(self: @This()) void {
self.gloc.freeRenderPipeline(self.pip); self.glloc.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(),
gloc: GpuAllocator, glloc: GpuAllocator,
target_view: GpuTextureView, target_view: GpuTextureView,
vertex_count: u32, vertex_count: u32,
args: anytype, args: anytype,
@@ -135,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(gloc.device.device, &.{ const bg = c.wgpuDeviceCreateBindGroup(glloc.device.device, &.{
.layout = bgl, .layout = bgl,
.entries = entries.ptr, .entries = entries.ptr,
.entryCount = @intCast(entries.len), .entryCount = @intCast(entries.len),
@@ -143,7 +146,7 @@ pub fn draw(
defer c.wgpuBindGroupRelease(bg); defer c.wgpuBindGroupRelease(bg);
// Encode Render Command // Encode Render Command
const enc = c.wgpuDeviceCreateCommandEncoder(gloc.device.device, null) orelse return error.Encoder; const enc = c.wgpuDeviceCreateCommandEncoder(glloc.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{
@@ -177,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(gloc.device.queue, 1, &cmd); c.wgpuQueueSubmit(glloc.device.queue, 1, &cmd);
} }
+56 -35
View File
@@ -1,26 +1,29 @@
const std = @import("std"); const std = @import("std");
const c = @import("utils.zig").c; const c = @import("utils.zig").c;
const svOpt = @import("utils.zig").svOpt;
const GpuAllocator = @import("GpuAllocator.zig"); const GpuAllocator = @import("GpuAllocator.zig");
const GpuBuffer = @import("GpuBuffer.zig"); const GpuBuffer = @import("GpuBuffer.zig");
const GpuTextureFormat = @import("lib.zig").GpuTextureFormat; const GpuTextureFormat = @import("lib.zig").GpuTextureFormat;
const GpuTextureUsage = @import("lib.zig").GpuTextureUsage; const GpuTextureUsage = @import("lib.zig").GpuTextureUsage;
pub const GpuTextureDef = struct { pub const GpuTextureDef = struct {
label: ?[]const u8 = null,
size: c.WGPUExtent3D, size: c.WGPUExtent3D,
usage: std.EnumSet(GpuTextureUsage), usage: std.EnumSet(GpuTextureUsage),
format: GpuTextureFormat, format: GpuTextureFormat,
}; };
raw: c.WGPUTexture, raw: c.WGPUTexture,
gloc: GpuAllocator, glloc: GpuAllocator,
def: GpuTextureDef, def: GpuTextureDef,
pub fn init(gloc: GpuAllocator, def: GpuTextureDef) !@This() { pub fn init(glloc: 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);
const desc = c.WGPUTextureDescriptor{ const desc = c.WGPUTextureDescriptor{
.label = svOpt(def.label),
.usage = use, .usage = use,
.dimension = c.WGPUTextureDimension_2D, .dimension = c.WGPUTextureDimension_2D,
.size = def.size, .size = def.size,
@@ -28,13 +31,13 @@ pub fn init(gloc: GpuAllocator, def: GpuTextureDef) !@This() {
.mipLevelCount = 1, .mipLevelCount = 1,
.sampleCount = 1, .sampleCount = 1,
}; };
const raw = try gloc.allocTexture(desc); const raw = try glloc.allocTexture(desc);
return .{ .gloc = gloc, .raw = raw, .def = def }; return .{ .glloc = glloc, .raw = raw, .def = def };
} }
pub fn deinit(self: @This()) void { pub fn deinit(self: @This()) void {
self.gloc.freeTexture(self.raw); self.glloc.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 {
@@ -50,10 +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(), gloc: GpuAllocator) !GpuBuffer { pub fn buffCopy(self: @This(), glloc: GpuAllocator) !GpuBuffer {
const buf = try GpuBuffer.init(gloc, self.bytesSize(), .initMany(&.{ .CopyDst, .CopySrc })); const buf = try GpuBuffer.init(glloc, .{
.size = self.bytesSize(),
.usage = .initMany(&.{ .CopyDst, .CopySrc }),
.label = "texture_copy_buffer",
});
const enc = c.wgpuDeviceCreateCommandEncoder(gloc.device.device, null) orelse return error.Encoder; const enc = c.wgpuDeviceCreateCommandEncoder(glloc.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{
@@ -75,7 +82,7 @@ pub fn buffCopy(self: @This(), gloc: GpuAllocator) !GpuBuffer {
const cmd = c.wgpuCommandEncoderFinish(enc, null); const cmd = c.wgpuCommandEncoderFinish(enc, null);
defer c.wgpuCommandBufferRelease(cmd); defer c.wgpuCommandBufferRelease(cmd);
c.wgpuQueueSubmit(gloc.device.queue, 1, &cmd); c.wgpuQueueSubmit(glloc.device.queue, 1, &cmd);
return buf; return buf;
} }
@@ -102,42 +109,56 @@ pub fn load(
) !void { ) !void {
const bytes = data.len * @sizeOf(T); const bytes = data.len * @sizeOf(T);
if (bytes == self.size) { c.wgpuQueueWriteTexture(
// Aligned path: direct download self.glloc.device.queue,
c.wgpuQueueWriteBuffer(self.gloc.device.queue, self.raw, 0, data.ptr, self.size); &.{
} else { .texture = self.raw,
// Unaligned path: Split the write into an aligned chunk and a padded remainder .mipLevel = 0,
// to support arbitrary lengths without any allocations or large stack arrays. .origin = .{ .x = 0, .y = 0, .z = 0 },
const aligned_part = (bytes / 4) * 4; .aspect = c.WGPUTextureAspect_All,
if (aligned_part > 0) { },
c.wgpuQueueWriteBuffer(self.gloc.device.queue, self.raw, 0, data.ptr, aligned_part); data.ptr,
} bytes,
&.{
var remainder_buf: [4]u8 = .{ 0, 0, 0, 0 }; .offset = 0,
const data_bytes = std.mem.sliceAsBytes(data); .bytesPerRow = self.bytesSizeRow(),
@memcpy(remainder_buf[0 .. bytes - aligned_part], data_bytes[aligned_part..bytes]); .rowsPerImage = self.def.size.height,
},
c.wgpuQueueWriteBuffer(self.gloc.device.queue, self.raw, aligned_part, &remainder_buf, 4); &self.def.size,
} );
} }
// GPU to CPU // GPU to CPU
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( const staging = try init(self.glloc, .{
self.gloc, .size = self.size,
self.size, .usage = .initMany(&.{ .MapRead, .CopyDst }),
.initMany(&.{ .MapRead, .CopyDst }), .label = "texture_read_staging",
); });
defer staging.deinit(); defer staging.deinit();
const enc = c.wgpuDeviceCreateCommandEncoder(self.gloc.device.device, null) orelse return error.Encoder; 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 src_copy = c.WGPUTexelCopyTextureInfo{
.texture = self.raw,
.mipLevel = 0,
.origin = .{ .x = 0, .y = 0, .z = 0 },
.aspect = c.WGPUTextureAspect_All,
};
const dst_copy = c.WGPUTexelCopyBufferInfo{
.buffer = staging.raw,
.layout = .{
.offset = 0,
.bytesPerRow = self.bytesSizeRow(),
.rowsPerImage = self.def.size.height,
},
};
c.wgpuCommandEncoderCopyTextureToBuffer(enc, &src_copy, &dst_copy, &self.def.size);
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.gloc.device.queue, 1, &cmd); c.wgpuQueueSubmit(self.glloc.device.queue, 1, &cmd);
var mapped = false; var mapped = false;
staging.mapAsync( staging.mapAsync(
@@ -146,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.gloc.device.poll(); while (!mapped) self.glloc.device.poll();
const ptr: [*]const T = @ptrCast(@alignCast( const ptr: [*]const T = @ptrCast(@alignCast(
staging.getConstMappedRange(0, self.size), staging.getConstMappedRange(0, self.size),
+8 -5
View File
@@ -1,32 +1,35 @@
const std = @import("std"); const std = @import("std");
const c = @import("utils.zig").c; const c = @import("utils.zig").c;
const svOpt = @import("utils.zig").svOpt;
const GpuAllocator = @import("GpuAllocator.zig"); const GpuAllocator = @import("GpuAllocator.zig");
const GpuTexture = @import("lib.zig").GpuTexture; const GpuTexture = @import("lib.zig").GpuTexture;
const GpuTextureFormat = @import("lib.zig").GpuTextureFormat; const GpuTextureFormat = @import("lib.zig").GpuTextureFormat;
const GpuTextureUsage = @import("lib.zig").GpuTextureUsage; const GpuTextureUsage = @import("lib.zig").GpuTextureUsage;
pub const GpuViewDef = struct { pub const GpuViewDef = struct {
label: ?[]const u8 = null,
usage: std.EnumSet(GpuTextureUsage) = .empty, usage: std.EnumSet(GpuTextureUsage) = .empty,
format: GpuTextureFormat = .Undefined, format: GpuTextureFormat = .Undefined,
}; };
raw: c.WGPUTextureView, raw: c.WGPUTextureView,
gloc: GpuAllocator, glloc: GpuAllocator,
pub fn init(gloc: GpuAllocator, texture: GpuTexture, def: GpuViewDef) !@This() { pub fn init(glloc: 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 gloc.allocTextureView(texture.raw, .{ const raw = try glloc.allocTextureView(texture.raw, .{
.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 .{ .gloc = gloc, .raw = raw }; return .{ .glloc = glloc, .raw = raw };
} }
pub fn deinit(self: @This()) void { pub fn deinit(self: @This()) void {
self.gloc.freeTextureView(self.raw); self.glloc.freeTextureView(self.raw);
} }
+14
View File
@@ -3,3 +3,17 @@ pub const c = @cImport(@cInclude("wgpu.h"));
pub fn sv(s: []const u8) c.WGPUStringView { pub fn sv(s: []const u8) c.WGPUStringView {
return .{ .data = s.ptr, .length = s.len }; return .{ .data = s.ptr, .length = s.len };
} }
/// Allows safely passing an optional Zig string to a WebGPU string view.
pub fn svOpt(s: ?[]const u8) c.WGPUStringView {
if (s) |str| return sv(str);
return .{ .data = null, .length = 0 };
}
/// Helper to print a WGPUStringView in your logs.
pub fn viewStr(view: c.WGPUStringView) []const u8 {
if (view.data != null and view.length > 0) {
return view.data[0..view.length];
}
return "unnamed";
}