Compare commits
| Author | SHA1 | Date | |
|---|---|---|---|
|
|
9c329a5a06 |
@@ -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.
|
||||
* **`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.
|
||||
* **`GpuRender` / `GpuTexture` / `GpuTextureView`**: Components used to initialize render pipelines, set up render attachments (textures), and bind render targets for offscreen drawing.
|
||||
|
||||
@@ -36,20 +36,17 @@ pub fn main(init: std.process.Init) !void {
|
||||
// 2. Create a GPU Arena to manage VRAM
|
||||
var grena = GpuArenaAllocator.init(allocator, device.gpuAllocator());
|
||||
defer grena.deinit();
|
||||
const glloc = grena.gpuAllocator();
|
||||
const gloc = grena.gpuAllocator();
|
||||
|
||||
// 3. Load the WGSL compute pipeline
|
||||
const add_cp = try GpuCompute.init(
|
||||
glloc,
|
||||
gloc,
|
||||
@embedFile("shaders/add.wgsl"),
|
||||
.{
|
||||
.label = "add",
|
||||
.bindings = &.{
|
||||
.{ .element_size = @sizeOf(f16) },
|
||||
.{ .element_size = @sizeOf(f16) },
|
||||
.{ .element_size = @sizeOf(f16) },
|
||||
},
|
||||
},
|
||||
.{ .bindings = &.{
|
||||
.{ .element_size = @sizeOf(f16) },
|
||||
.{ .element_size = @sizeOf(f16) },
|
||||
.{ .element_size = @sizeOf(f16) },
|
||||
} },
|
||||
);
|
||||
|
||||
// 4. Setup CPU data
|
||||
@@ -66,20 +63,19 @@ pub fn main(init: std.process.Init) !void {
|
||||
|
||||
// 5. Initialize raw GPU Buffers
|
||||
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_b = try GpuBuffer.init(glloc, .{ .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_a = try GpuBuffer.init(gloc, byte_size, .initMany(&.{ .Storage, .CopyDst, .CopySrc }));
|
||||
const buf_b = try GpuBuffer.init(gloc, byte_size, .initMany(&.{ .Storage, .CopyDst, .CopySrc }));
|
||||
const buf_out = try GpuBuffer.init(gloc, byte_size, .initMany(&.{ .Storage, .CopyDst, .CopySrc }));
|
||||
|
||||
// Note: 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.
|
||||
// This will also release pipelines, textures, ect. Everything using a GpuAllocator to init.
|
||||
// Note: Buffers, pipelines, and other objects initialized with 'gloc'
|
||||
// are safely tied to the GpuArenaAllocator and will automatically release.
|
||||
|
||||
// 6. Transfer data from CPU slices to GPU Buffers
|
||||
try buf_a.load(f16, data_a);
|
||||
try buf_b.load(f16, data_b);
|
||||
|
||||
// 7. Dispatch the Compute
|
||||
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
|
||||
const out = try buf_out.read(allocator, f16);
|
||||
@@ -100,7 +96,6 @@ and pull the frame pixels back to the CPU to write a standard image file:
|
||||
const std = @import("std");
|
||||
const gpu = @import("gpu");
|
||||
const GpuDevice = gpu.GpuDevice;
|
||||
const GpuArenaAllocator = gpu.GpuArenaAllocator;
|
||||
const GpuBuffer = gpu.GpuBuffer;
|
||||
const GpuRender = gpu.GpuRender;
|
||||
const GpuTexture = gpu.GpuTexture;
|
||||
@@ -116,21 +111,19 @@ pub fn main(init: std.process.Init) !void {
|
||||
const device = try GpuDevice.init(.{});
|
||||
defer device.deinit();
|
||||
|
||||
// 2. Init VRAM Arena
|
||||
var grena = GpuArenaAllocator.init(allocator, device.gpuAllocator());
|
||||
defer grena.deinit();
|
||||
const glloc = grena.gpuAllocator();
|
||||
// 2. Get base device GPU Allocator
|
||||
const gloc = device.gpuAllocator();
|
||||
|
||||
// 3. Load Render Pipeline
|
||||
const circle_rp = try GpuRender.init(
|
||||
glloc,
|
||||
gloc,
|
||||
@embedFile("shaders/circle.wgsl"),
|
||||
.{ .bindings = &.{}, .texture_format = .RGBA8Unorm, .topology = .TriangleStrip },
|
||||
);
|
||||
defer circle_rp.deinit();
|
||||
|
||||
// 4. Create VRAM texture to render into
|
||||
const texture = try GpuTexture.init(glloc, .{
|
||||
const texture = try GpuTexture.init(gloc, .{
|
||||
.format = .RGBA8Unorm,
|
||||
.size = .{ .width = width, .height = height, .depthOrArrayLayers = 1 },
|
||||
.usage = .initMany(&.{ .RenderAttachment, .CopySrc }),
|
||||
@@ -138,22 +131,21 @@ pub fn main(init: std.process.Init) !void {
|
||||
defer texture.deinit();
|
||||
|
||||
// 5. Create a view from texture
|
||||
const view = try GpuTextureView.init(glloc, texture, .{});
|
||||
const view = try GpuTextureView.init(gloc, texture, .{});
|
||||
defer view.deinit();
|
||||
|
||||
// 6. Run the rendering pipeline
|
||||
try circle_rp.draw(glloc, view, 4, .{});
|
||||
try circle_rp.draw(gloc, view, 4, .{});
|
||||
|
||||
// 7. Load Texture into GpuBuffer
|
||||
const cpu_staging_cpu = try texture.buffCopy(glloc);
|
||||
defer cpu_staging_cpu.deinit();
|
||||
// 7. Copy Texture into a readable GPU staging buffer
|
||||
const cpu_staging_buf = try texture.buffCopy(gloc);
|
||||
defer cpu_staging_buf.deinit();
|
||||
|
||||
// 8. Read GpuBuffer to CPU
|
||||
// This need to be free manually because CPU memory
|
||||
const pixels = try cpu_staging_cpu.read(allocator, u8);
|
||||
// 8. Read GpuBuffer to CPU memory
|
||||
const pixels = try cpu_staging_buf.read(allocator, u8);
|
||||
defer allocator.free(pixels);
|
||||
|
||||
// 9. Write a simple ppm image
|
||||
// 9. Write out to a simple PPM image
|
||||
try savePpm(init.io, "circle.ppm", width, height, pixels);
|
||||
}
|
||||
|
||||
@@ -223,7 +215,7 @@ No extra installation required. Automatically links against `d3d12`, `dxgi`, and
|
||||
Add it to your `build.zig.zon`:
|
||||
|
||||
```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`:
|
||||
|
||||
+1
-1
@@ -1,6 +1,6 @@
|
||||
.{
|
||||
.name = .zig_wgpu,
|
||||
.version = "0.2.3",
|
||||
.version = "0.2.0",
|
||||
.fingerprint = 0x5d0e853acbc0c2c6,
|
||||
.minimum_zig_version = "0.16.0",
|
||||
.dependencies = .{},
|
||||
|
||||
+20
-21
@@ -6,27 +6,26 @@ const GpuAllocator = gpu.GpuAllocator;
|
||||
const GpuBuffer = gpu.GpuBuffer;
|
||||
const GpuCompute = gpu.GpuCompute;
|
||||
|
||||
pub const std_options = std.Options{ .log_level = .info };
|
||||
|
||||
/// Minimal implementation of a f16 Vector
|
||||
const Vec = struct {
|
||||
buf: GpuBuffer,
|
||||
len: usize,
|
||||
|
||||
// Changed: glloc is passed by value (const)
|
||||
pub fn initZero(glloc: GpuAllocator, len: usize) !Vec {
|
||||
// Changed: gloc is passed by value (const)
|
||||
pub fn initZero(gloc: GpuAllocator, len: usize) !Vec {
|
||||
return .{
|
||||
.buf = try GpuBuffer.init(glloc, .{
|
||||
.size = len * @sizeOf(f16),
|
||||
.usage = .initMany(&.{ .Storage, .CopyDst, .CopySrc }),
|
||||
}),
|
||||
.buf = try GpuBuffer.init(
|
||||
gloc,
|
||||
len * @sizeOf(f16),
|
||||
.initMany(&.{ .Storage, .CopyDst, .CopySrc }),
|
||||
),
|
||||
.len = len,
|
||||
};
|
||||
}
|
||||
|
||||
// Changed: glloc is passed by value
|
||||
pub fn initLoad(glloc: GpuAllocator, data: []const f16) !Vec {
|
||||
var self = try initZero(glloc, data.len);
|
||||
// Changed: gloc is passed by value
|
||||
pub fn initLoad(gloc: GpuAllocator, data: []const f16) !Vec {
|
||||
var self = try initZero(gloc, data.len);
|
||||
try self.load(data); // Direct access via the interface copy
|
||||
return self;
|
||||
}
|
||||
@@ -40,18 +39,18 @@ const Vec = struct {
|
||||
try self.buf.load(f16, data);
|
||||
}
|
||||
|
||||
// Changed: glloc is passed by value instead of *GpuAllocator
|
||||
pub fn run(self: Vec, glloc: GpuAllocator, other: Vec, process: GpuCompute) !Vec {
|
||||
// Changed: gloc is passed by value instead of *GpuAllocator
|
||||
pub fn run(self: Vec, gloc: GpuAllocator, other: Vec, process: GpuCompute) !Vec {
|
||||
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();
|
||||
|
||||
try process.run(glloc, .{ self.buf, other.buf, result.buf });
|
||||
try process.run(gloc, .{ self.buf, other.buf, result.buf });
|
||||
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 {
|
||||
return self.buf.read(alloc, f16);
|
||||
}
|
||||
@@ -63,9 +62,9 @@ pub fn main(init: std.process.Init) !void {
|
||||
|
||||
var grena = GpuArenaAllocator.init(init.gpa, device.gpuAllocator());
|
||||
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) },
|
||||
@@ -120,9 +119,9 @@ pub fn main(init: std.process.Init) !void {
|
||||
// --- 1. GPU ALLOCATION PHASE ---
|
||||
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();
|
||||
const b = try Vec.initLoad(glloc, data_b);
|
||||
const b = try Vec.initLoad(gloc, data_b);
|
||||
defer b.deinit();
|
||||
|
||||
const alloc_duration = alloc_start.durationTo(std.Io.Clock.awake.now(init.io));
|
||||
@@ -132,7 +131,7 @@ pub fn main(init: std.process.Init) !void {
|
||||
// --- 2. COMPUTE PHASE ---
|
||||
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();
|
||||
|
||||
// All 3 buffers (a, b, sum) are currently resident in VRAM here.
|
||||
|
||||
+6
-8
@@ -18,20 +18,18 @@ pub fn main(init: std.process.Init) !void {
|
||||
defer device.deinit();
|
||||
|
||||
// 2. Init VRAM Arena
|
||||
var grena = GpuArenaAllocator.init(allocator, device.gpuAllocator());
|
||||
defer grena.deinit();
|
||||
const glloc = grena.gpuAllocator();
|
||||
const gloc = device.gpuAllocator();
|
||||
|
||||
// 3. Load Render Pipeline
|
||||
const circle_rp = try GpuRender.init(
|
||||
glloc,
|
||||
gloc,
|
||||
@embedFile("shaders/circle.wgsl"),
|
||||
.{ .bindings = &.{}, .texture_format = .RGBA8Unorm, .topology = .TriangleStrip },
|
||||
);
|
||||
defer circle_rp.deinit();
|
||||
|
||||
// 4. Create VRAM texture to render into
|
||||
const texture = try GpuTexture.init(glloc, .{
|
||||
const texture = try GpuTexture.init(gloc, .{
|
||||
.format = .RGBA8Unorm,
|
||||
.size = .{ .width = width, .height = height, .depthOrArrayLayers = 1 },
|
||||
.usage = .initMany(&.{ .RenderAttachment, .CopySrc }),
|
||||
@@ -39,14 +37,14 @@ pub fn main(init: std.process.Init) !void {
|
||||
defer texture.deinit();
|
||||
|
||||
// 5. Create a view from texture
|
||||
const view = try GpuTextureView.init(glloc, texture, .{});
|
||||
const view = try GpuTextureView.init(gloc, texture, .{});
|
||||
defer view.deinit();
|
||||
|
||||
// 6. Run the rendering pipeline
|
||||
try circle_rp.draw(glloc, view, 4, .{});
|
||||
try circle_rp.draw(gloc, view, 4, .{});
|
||||
|
||||
// 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();
|
||||
|
||||
// 8. Read GpuBuffer to CPU
|
||||
|
||||
+11
-14
@@ -15,20 +15,17 @@ pub fn main(init: std.process.Init) !void {
|
||||
// 2. Create a GPU Arena to manage VRAM
|
||||
var grena = GpuArenaAllocator.init(allocator, device.gpuAllocator());
|
||||
defer grena.deinit();
|
||||
const glloc = grena.gpuAllocator();
|
||||
const gloc = grena.gpuAllocator();
|
||||
|
||||
// 3. Load the WGSL compute pipeline
|
||||
const add_cp = try GpuCompute.init(
|
||||
glloc,
|
||||
gloc,
|
||||
@embedFile("shaders/add.wgsl"),
|
||||
.{
|
||||
.label = "add",
|
||||
.bindings = &.{
|
||||
.{ .element_size = @sizeOf(f16) },
|
||||
.{ .element_size = @sizeOf(f16) },
|
||||
.{ .element_size = @sizeOf(f16) },
|
||||
},
|
||||
},
|
||||
.{ .bindings = &.{
|
||||
.{ .element_size = @sizeOf(f16) },
|
||||
.{ .element_size = @sizeOf(f16) },
|
||||
.{ .element_size = @sizeOf(f16) },
|
||||
} },
|
||||
);
|
||||
|
||||
// 4. Setup CPU data
|
||||
@@ -45,9 +42,9 @@ pub fn main(init: std.process.Init) !void {
|
||||
|
||||
// 5. Initialize raw GPU Buffers
|
||||
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_b = try GpuBuffer.init(glloc, .{ .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_a = try GpuBuffer.init(gloc, byte_size, .initMany(&.{ .Storage, .CopyDst, .CopySrc }));
|
||||
const buf_b = try GpuBuffer.init(gloc, byte_size, .initMany(&.{ .Storage, .CopyDst, .CopySrc }));
|
||||
const buf_out = try GpuBuffer.init(gloc, byte_size, .initMany(&.{ .Storage, .CopyDst, .CopySrc }));
|
||||
|
||||
// Note: 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.
|
||||
@@ -58,7 +55,7 @@ pub fn main(init: std.process.Init) !void {
|
||||
try buf_b.load(f16, data_b);
|
||||
|
||||
// 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
|
||||
const out = try buf_out.read(allocator, f16);
|
||||
|
||||
@@ -3,7 +3,6 @@ const GpuDevice = @import("GpuDevice.zig");
|
||||
const GpuAllocator = @import("GpuAllocator.zig");
|
||||
const GpuTextureFormat = @import("lib.zig").GpuTextureFormat;
|
||||
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
|
||||
tracked_buffers: std.AutoHashMap(c.WGPUBuffer, c.WGPUBufferDescriptor),
|
||||
@@ -25,40 +24,30 @@ pub fn init(cpu_allocator: std.mem.Allocator, child_allocator: GpuAllocator) @Th
|
||||
}
|
||||
|
||||
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();
|
||||
while (it_buffer.next()) |buf_ptr|
|
||||
freeBuffer(self, buf_ptr.*);
|
||||
self.child_allocator.freeBuffer(buf_ptr.*);
|
||||
self.tracked_buffers.deinit();
|
||||
|
||||
var it_tex = self.tracked_textures.keyIterator();
|
||||
while (it_tex.next()) |buf_ptr|
|
||||
freeTexture(self, buf_ptr.*);
|
||||
self.child_allocator.freeTexture(buf_ptr.*);
|
||||
self.tracked_textures.deinit();
|
||||
|
||||
var it_view = self.tracked_views.keyIterator();
|
||||
while (it_view.next()) |buf_ptr|
|
||||
freeTextureView(self, buf_ptr.*);
|
||||
self.child_allocator.freeTextureView(buf_ptr.*);
|
||||
self.tracked_views.deinit();
|
||||
|
||||
var it_render = self.tracked_renders.keyIterator();
|
||||
while (it_render.next()) |buf_ptr|
|
||||
freeRenderPipeline(self, buf_ptr.*);
|
||||
self.child_allocator.freeRenderPipeline(buf_ptr.*);
|
||||
self.tracked_renders.deinit();
|
||||
|
||||
var it_compute = self.tracked_computes.keyIterator();
|
||||
while (it_compute.next()) |buf_ptr|
|
||||
freeComputePipeline(self, buf_ptr.*);
|
||||
self.child_allocator.freeComputePipeline(buf_ptr.*);
|
||||
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
|
||||
@@ -90,14 +79,6 @@ fn allocBuffer(ctx: *anyopaque, desc: c.WGPUBufferDescriptor) anyerror!c.WGPUBuf
|
||||
const raw = try self.child_allocator.allocBuffer(desc);
|
||||
self.tracked_buffers.putAssumeCapacity(raw, desc);
|
||||
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;
|
||||
}
|
||||
|
||||
@@ -106,12 +87,6 @@ fn freeBuffer(ctx: *anyopaque, raw: c.WGPUBuffer) void {
|
||||
if (self.tracked_buffers.fetchRemove(raw)) |kv| {
|
||||
self.child_allocator.freeBuffer(raw);
|
||||
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,
|
||||
});
|
||||
}
|
||||
}
|
||||
|
||||
@@ -122,21 +97,13 @@ fn allocTexture(ctx: *anyopaque, desc: c.WGPUTextureDescriptor) anyerror!c.WGPUT
|
||||
const format: GpuTextureFormat = @enumFromInt(desc.format);
|
||||
const bytes_size = desc.size.width * desc.size.height * format.bytesPerPixel();
|
||||
|
||||
if (bytes_size + self.allocated_vram_bytes > self.child_allocator.device.def.vram_bytes_limit)
|
||||
if (bytes_size + self.allocated_vram_bytes > self.child_allocator.device.config.vram_bytes_limit)
|
||||
return error.ExceedsVramBudget;
|
||||
|
||||
const raw = try self.child_allocator.allocTexture(desc);
|
||||
|
||||
self.tracked_textures.putAssumeCapacity(raw, desc);
|
||||
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;
|
||||
}
|
||||
|
||||
@@ -150,12 +117,6 @@ fn freeTexture(ctx: *anyopaque, raw: c.WGPUTexture) void {
|
||||
const format: GpuTextureFormat = @enumFromInt(desc.format);
|
||||
const bytes_size = desc.size.width * desc.size.height * format.bytesPerPixel();
|
||||
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,
|
||||
});
|
||||
}
|
||||
}
|
||||
|
||||
@@ -164,17 +125,13 @@ fn allocTextureView(ctx: *anyopaque, texture: c.WGPUTexture, desc: c.WGPUTexture
|
||||
try self.tracked_views.ensureTotalCapacity(self.tracked_views.count() + 1);
|
||||
const raw = try self.child_allocator.allocTextureView(texture, desc);
|
||||
self.tracked_views.putAssumeCapacity(raw, desc);
|
||||
std.log.debug("Allocated Texture View '{s}'", .{viewStr(desc.label)});
|
||||
return raw;
|
||||
}
|
||||
|
||||
fn freeTextureView(ctx: *anyopaque, raw: c.WGPUTextureView) void {
|
||||
const self: *@This() = @ptrCast(@alignCast(ctx));
|
||||
if (self.tracked_views.fetchRemove(raw)) |kv| {
|
||||
if (self.tracked_views.remove(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 {
|
||||
@@ -182,17 +139,13 @@ fn allocRenderPipeline(ctx: *anyopaque, desc: c.WGPURenderPipelineDescriptor) an
|
||||
try self.tracked_renders.ensureTotalCapacity(self.tracked_renders.count() + 1);
|
||||
const raw = try self.child_allocator.allocRenderPipeline(desc);
|
||||
self.tracked_renders.putAssumeCapacity(raw, desc);
|
||||
std.log.debug("Allocated Render Pipeline '{s}'", .{viewStr(desc.label)});
|
||||
return raw;
|
||||
}
|
||||
|
||||
fn freeRenderPipeline(ctx: *anyopaque, raw: c.WGPURenderPipeline) void {
|
||||
const self: *@This() = @ptrCast(@alignCast(ctx));
|
||||
if (self.tracked_renders.fetchRemove(raw)) |kv| {
|
||||
if (self.tracked_renders.remove(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 {
|
||||
@@ -200,15 +153,11 @@ fn allocComputePipeline(ctx: *anyopaque, desc: c.WGPUComputePipelineDescriptor)
|
||||
try self.tracked_computes.ensureTotalCapacity(self.tracked_computes.count() + 1);
|
||||
const raw = try self.child_allocator.allocComputePipeline(desc);
|
||||
self.tracked_computes.putAssumeCapacity(raw, desc);
|
||||
std.log.debug("Allocated Compute Pipeline '{s}'", .{viewStr(desc.label)});
|
||||
return raw;
|
||||
}
|
||||
|
||||
fn freeComputePipeline(ctx: *anyopaque, raw: c.WGPUComputePipeline) void {
|
||||
const self: *@This() = @ptrCast(@alignCast(ctx));
|
||||
if (self.tracked_computes.fetchRemove(raw)) |kv| {
|
||||
if (self.tracked_computes.remove(raw))
|
||||
self.child_allocator.freeComputePipeline(raw);
|
||||
const desc = kv.value;
|
||||
std.log.debug("Freed Compute Pipeline '{s}'", .{viewStr(desc.label)});
|
||||
}
|
||||
}
|
||||
|
||||
+18
-29
@@ -1,12 +1,11 @@
|
||||
const std = @import("std");
|
||||
const c = @import("utils.zig").c;
|
||||
const GpuAllocator = @import("GpuAllocator.zig");
|
||||
const svOpt = @import("utils.zig").svOpt;
|
||||
|
||||
raw: c.WGPUBuffer,
|
||||
size: u64,
|
||||
usage: c.WGPUBufferUsage,
|
||||
glloc: GpuAllocator,
|
||||
gloc: GpuAllocator,
|
||||
|
||||
const BufferUsage = enum(u64) {
|
||||
None = 0x0000000000000000,
|
||||
@@ -22,35 +21,25 @@ const BufferUsage = enum(u64) {
|
||||
QueryResolve = 0x0000000000000200,
|
||||
};
|
||||
|
||||
const GpuBufferDef = struct {
|
||||
label: ?[]const u8 = null,
|
||||
size: u64,
|
||||
usage: std.EnumSet(BufferUsage),
|
||||
};
|
||||
|
||||
pub fn init(glloc: GpuAllocator, def: GpuBufferDef) !@This() {
|
||||
pub fn init(gloc: GpuAllocator, size: u64, usage: std.EnumSet(BufferUsage)) !@This() {
|
||||
var use: u64 = 0;
|
||||
var iter = def.usage.iterator();
|
||||
var iter = usage.iterator();
|
||||
while (iter.next()) |flag| use |= @intFromEnum(flag);
|
||||
|
||||
// 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, size, 4);
|
||||
|
||||
const raw_handle = try glloc.allocBuffer(.{
|
||||
.size = aligned_size,
|
||||
.usage = use,
|
||||
.label = svOpt(def.label),
|
||||
});
|
||||
const raw_handle = try gloc.allocBuffer(.{ .size = aligned_size, .usage = use });
|
||||
return .{
|
||||
.raw = raw_handle,
|
||||
.size = aligned_size,
|
||||
.usage = use,
|
||||
.glloc = glloc,
|
||||
.gloc = gloc,
|
||||
};
|
||||
}
|
||||
|
||||
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 {
|
||||
@@ -81,20 +70,20 @@ pub fn load(
|
||||
|
||||
if (bytes == self.size) {
|
||||
// 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.size);
|
||||
} else {
|
||||
// Unaligned path: Split the write into an aligned chunk and a padded remainder
|
||||
// to support arbitrary lengths without any allocations or large stack arrays.
|
||||
const aligned_part = (bytes / 4) * 4;
|
||||
if (aligned_part > 0) {
|
||||
c.wgpuQueueWriteBuffer(self.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 };
|
||||
const data_bytes = std.mem.sliceAsBytes(data);
|
||||
@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);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -102,19 +91,19 @@ pub fn load(
|
||||
pub fn read(self: @This(), alloc: std.mem.Allocator, T: type) ![]T {
|
||||
const out = try alloc.alloc(T, @divExact(self.size, @sizeOf(T)));
|
||||
|
||||
const staging = try init(self.glloc, .{
|
||||
.size = self.size,
|
||||
.usage = .initMany(&.{ .MapRead, .CopyDst }),
|
||||
.label = "staging_read_buffer",
|
||||
});
|
||||
const staging = try init(
|
||||
self.gloc,
|
||||
self.size,
|
||||
.initMany(&.{ .MapRead, .CopyDst }),
|
||||
);
|
||||
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;
|
||||
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);
|
||||
c.wgpuQueueSubmit(self.gloc.device.queue, 1, &cmd);
|
||||
|
||||
var mapped = false;
|
||||
staging.mapAsync(
|
||||
@@ -123,7 +112,7 @@ pub fn read(self: @This(), alloc: std.mem.Allocator, T: type) ![]T {
|
||||
self.size,
|
||||
.{ .callback = onMapped, .userdata1 = &mapped },
|
||||
);
|
||||
while (!mapped) self.glloc.device.poll();
|
||||
while (!mapped) self.gloc.device.poll();
|
||||
|
||||
const ptr: [*]const T = @ptrCast(@alignCast(
|
||||
staging.getConstMappedRange(0, self.size),
|
||||
|
||||
+19
-24
@@ -1,6 +1,5 @@
|
||||
const c = @import("utils.zig").c;
|
||||
const sv = @import("utils.zig").sv;
|
||||
const svOpt = @import("utils.zig").svOpt;
|
||||
const GpuAllocator = @import("GpuAllocator.zig");
|
||||
const GpuBuffer = @import("GpuBuffer.zig");
|
||||
const GpuDevice = @import("GpuDevice.zig");
|
||||
@@ -12,7 +11,6 @@ pub const Binding = struct {
|
||||
};
|
||||
|
||||
pub const ComputeDef = struct {
|
||||
label: ?[]const u8 = null,
|
||||
bindings: []const Binding,
|
||||
workgroup_size: u32 = 256,
|
||||
max_workgroups: u32 = 65535,
|
||||
@@ -22,40 +20,37 @@ pub const ComputeDef = struct {
|
||||
};
|
||||
|
||||
pip: c.WGPUComputePipeline,
|
||||
glloc: GpuAllocator,
|
||||
gloc: GpuAllocator,
|
||||
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{
|
||||
.chain = .{ .sType = c.WGPUSType_ShaderSourceWGSL },
|
||||
.code = sv(wgsl),
|
||||
};
|
||||
const shader = c.wgpuDeviceCreateShaderModule(glloc.device.device, &.{
|
||||
const shader = c.wgpuDeviceCreateShaderModule(gloc.device.device, &.{
|
||||
.nextInChain = @ptrCast(&wgsl_src),
|
||||
}) orelse return error.Shader;
|
||||
defer c.wgpuShaderModuleRelease(shader);
|
||||
|
||||
const pip = try glloc.allocComputePipeline(.{
|
||||
.label = svOpt(def.label),
|
||||
.compute = .{ .module = shader, .entryPoint = sv("main") },
|
||||
});
|
||||
const pip = try gloc.allocComputePipeline(.{ .compute = .{ .module = shader, .entryPoint = sv("main") } });
|
||||
|
||||
return .{
|
||||
.glloc = glloc,
|
||||
.gloc = gloc,
|
||||
.pip = pip,
|
||||
.def = def,
|
||||
};
|
||||
}
|
||||
|
||||
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.
|
||||
/// 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(
|
||||
self: @This(),
|
||||
glloc: GpuAllocator,
|
||||
gloc: GpuAllocator,
|
||||
args: anytype,
|
||||
) !void {
|
||||
const type_info = @typeInfo(@TypeOf(args));
|
||||
@@ -113,12 +108,12 @@ pub fn run(
|
||||
defer if (info_buf) |b| b.deinit();
|
||||
|
||||
if (self.def.append_info_buffer) {
|
||||
info_buf = try GpuBuffer.init(glloc, .{
|
||||
.size = @sizeOf(u32),
|
||||
.usage = .initMany(&.{ .Uniform, .CopyDst }),
|
||||
.label = "compute_info_buffer",
|
||||
});
|
||||
c.wgpuQueueWriteBuffer(glloc.device.queue, info_buf.?.raw, 0, &elements_count, @sizeOf(u32));
|
||||
info_buf = try GpuBuffer.init(
|
||||
gloc,
|
||||
@sizeOf(u32),
|
||||
.initMany(&.{ .Uniform, .CopyDst }),
|
||||
);
|
||||
c.wgpuQueueWriteBuffer(gloc.device.queue, info_buf.?.raw, 0, &elements_count, @sizeOf(u32));
|
||||
|
||||
entries_buf[entry_count] = .{
|
||||
.binding = @intCast(entry_count),
|
||||
@@ -130,11 +125,11 @@ pub fn run(
|
||||
}
|
||||
|
||||
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(
|
||||
glloc: GpuAllocator,
|
||||
gloc: GpuAllocator,
|
||||
pipeline: c.WGPUComputePipeline,
|
||||
entries: []const c.WGPUBindGroupEntry,
|
||||
n: usize,
|
||||
@@ -146,14 +141,14 @@ fn submitPass(
|
||||
const bgl = c.wgpuComputePipelineGetBindGroupLayout(pipeline, 0);
|
||||
defer c.wgpuBindGroupLayoutRelease(bgl);
|
||||
|
||||
const bg = c.wgpuDeviceCreateBindGroup(glloc.device.device, &.{
|
||||
const bg = c.wgpuDeviceCreateBindGroup(gloc.device.device, &.{
|
||||
.layout = bgl,
|
||||
.entries = entries.ptr,
|
||||
.entryCount = entries.len,
|
||||
}) orelse return error.BindGroup;
|
||||
defer c.wgpuBindGroupRelease(bg);
|
||||
|
||||
const enc = c.wgpuDeviceCreateCommandEncoder(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);
|
||||
c.wgpuComputePassEncoderSetPipeline(pass, pipeline);
|
||||
c.wgpuComputePassEncoderSetBindGroup(pass, 0, bg, 0, null);
|
||||
@@ -168,7 +163,7 @@ fn submitPass(
|
||||
const cmd = c.wgpuCommandEncoderFinish(enc, null);
|
||||
defer c.wgpuCommandEncoderRelease(enc);
|
||||
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 {
|
||||
|
||||
+11
-27
@@ -1,17 +1,20 @@
|
||||
const std = @import("std");
|
||||
const c = @import("utils.zig").c;
|
||||
const sv = @import("utils.zig").sv;
|
||||
const svOpt = @import("utils.zig").svOpt;
|
||||
const GpuAllocator = @import("GpuAllocator.zig");
|
||||
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 {
|
||||
adapter: c.WGPUAdapter = null,
|
||||
device: c.WGPUDevice = null,
|
||||
};
|
||||
|
||||
const GpuDeviceDef = struct {
|
||||
label: ?[]const u8 = null,
|
||||
const GpuDeviceConfig = struct {
|
||||
/// VRAM limit. Default 2 GB
|
||||
vram_bytes_limit: u64 = 2 * 1024 * 1024 * 1024,
|
||||
power_preference: enum(c_uint) {
|
||||
@@ -28,9 +31,9 @@ device: c.WGPUDevice,
|
||||
queue: c.WGPUQueue,
|
||||
limits: c.WGPULimits,
|
||||
|
||||
def: GpuDeviceDef,
|
||||
config: GpuDeviceConfig,
|
||||
|
||||
pub fn init(def: GpuDeviceDef) !@This() {
|
||||
pub fn init(config: GpuDeviceConfig) !@This() {
|
||||
const instance = c.wgpuCreateInstance(
|
||||
&std.mem.zeroes(c.WGPUInstanceDescriptor),
|
||||
) orelse return error.NoInstance;
|
||||
@@ -39,25 +42,13 @@ pub fn init(def: GpuDeviceDef) !@This() {
|
||||
var ctx = Ctx{};
|
||||
_ = c.wgpuInstanceRequestAdapter(
|
||||
instance,
|
||||
&.{ .powerPreference = @intFromEnum(def.power_preference) },
|
||||
&.{ .powerPreference = @intFromEnum(config.power_preference) },
|
||||
.{ .callback = onAdapter, .userdata1 = &ctx },
|
||||
);
|
||||
c.wgpuInstanceProcessEvents(instance);
|
||||
const adapter = ctx.adapter orelse return error.NoAdapter;
|
||||
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);
|
||||
c.wgpuAdapterGetFeatures(adapter, &supported_features);
|
||||
|
||||
@@ -65,11 +56,6 @@ pub fn init(def: GpuDeviceDef) !@This() {
|
||||
supported_limits.nextInChain = null;
|
||||
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;
|
||||
for (0..supported_features.featureCount) |i| {
|
||||
if (supported_features.features[i] == c.WGPUFeatureName_ShaderF16) {
|
||||
@@ -77,8 +63,6 @@ pub fn init(def: GpuDeviceDef) !@This() {
|
||||
break;
|
||||
}
|
||||
}
|
||||
std.log.info(" Shader F16 : {}", .{has_f16});
|
||||
std.log.info("=================================", .{});
|
||||
|
||||
var feature_buf = [_]c.WGPUFeatureName{c.WGPUFeatureName_ShaderF16};
|
||||
const required_features: []const c.WGPUFeatureName =
|
||||
@@ -86,7 +70,7 @@ pub fn init(def: GpuDeviceDef) !@This() {
|
||||
|
||||
const device_descriptor = c.WGPUDeviceDescriptor{
|
||||
.nextInChain = null,
|
||||
.label = svOpt(def.label),
|
||||
.label = sv("TensorCompilerDevice"),
|
||||
.requiredFeatureCount = required_features.len,
|
||||
.requiredFeatures = if (required_features.len > 0) required_features.ptr else null,
|
||||
.requiredLimits = &supported_limits,
|
||||
@@ -105,7 +89,7 @@ pub fn init(def: GpuDeviceDef) !@This() {
|
||||
.device = device,
|
||||
.queue = c.wgpuDeviceGetQueue(device),
|
||||
.limits = supported_limits,
|
||||
.def = def,
|
||||
.config = config,
|
||||
};
|
||||
}
|
||||
|
||||
|
||||
+10
-13
@@ -1,7 +1,6 @@
|
||||
const std = @import("std");
|
||||
const c = @import("utils.zig").c;
|
||||
const sv = @import("utils.zig").sv;
|
||||
const svOpt = @import("utils.zig").svOpt;
|
||||
const GpuAllocator = @import("GpuAllocator.zig");
|
||||
const GpuBuffer = @import("GpuBuffer.zig");
|
||||
const GpuDevice = @import("GpuDevice.zig");
|
||||
@@ -13,7 +12,6 @@ pub const Binding = struct {
|
||||
};
|
||||
|
||||
pub const GpuRenderDef = struct {
|
||||
label: ?[]const u8 = null,
|
||||
bindings: []const Binding = &.{},
|
||||
/// The surface texture format we are rendering to (e.g., BGRA8Unorm)
|
||||
texture_format: GpuTextureFormat,
|
||||
@@ -34,16 +32,16 @@ const GpuPrimitiveTopology = enum(c_uint) {
|
||||
Force32 = 0x7FFFFFFF,
|
||||
};
|
||||
|
||||
glloc: GpuAllocator,
|
||||
gloc: GpuAllocator,
|
||||
pip: c.WGPURenderPipeline,
|
||||
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{
|
||||
.chain = .{ .sType = c.WGPUSType_ShaderSourceWGSL },
|
||||
.code = sv(wgsl),
|
||||
};
|
||||
const shader = c.wgpuDeviceCreateShaderModule(glloc.device.device, &.{
|
||||
const shader = c.wgpuDeviceCreateShaderModule(gloc.device.device, &.{
|
||||
.nextInChain = @ptrCast(&wgsl_src),
|
||||
}) orelse return error.Shader;
|
||||
defer c.wgpuShaderModuleRelease(shader);
|
||||
@@ -69,8 +67,7 @@ pub fn init(glloc: GpuAllocator, wgsl: []const u8, def: GpuRenderDef) !@This() {
|
||||
};
|
||||
|
||||
// 3. Compile the Complete Render Pipeline
|
||||
const pip = try glloc.allocRenderPipeline(.{
|
||||
.label = svOpt(def.label),
|
||||
const pip = try gloc.allocRenderPipeline(.{
|
||||
.vertex = .{
|
||||
.module = shader,
|
||||
.entryPoint = sv(def.vertex_entry),
|
||||
@@ -90,21 +87,21 @@ pub fn init(glloc: GpuAllocator, wgsl: []const u8, def: GpuRenderDef) !@This() {
|
||||
});
|
||||
|
||||
return .{
|
||||
.glloc = glloc,
|
||||
.gloc = gloc,
|
||||
.pip = pip,
|
||||
.def = def,
|
||||
};
|
||||
}
|
||||
|
||||
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.
|
||||
/// Passes bind groups via a tuple exactly like your original compute setup.
|
||||
pub fn draw(
|
||||
self: @This(),
|
||||
glloc: GpuAllocator,
|
||||
gloc: GpuAllocator,
|
||||
target_view: GpuTextureView,
|
||||
vertex_count: u32,
|
||||
args: anytype,
|
||||
@@ -138,7 +135,7 @@ pub fn draw(
|
||||
const bgl = c.wgpuRenderPipelineGetBindGroupLayout(self.pip, 0);
|
||||
defer c.wgpuBindGroupLayoutRelease(bgl);
|
||||
|
||||
const bg = c.wgpuDeviceCreateBindGroup(glloc.device.device, &.{
|
||||
const bg = c.wgpuDeviceCreateBindGroup(gloc.device.device, &.{
|
||||
.layout = bgl,
|
||||
.entries = entries.ptr,
|
||||
.entryCount = @intCast(entries.len),
|
||||
@@ -146,7 +143,7 @@ pub fn draw(
|
||||
defer c.wgpuBindGroupRelease(bg);
|
||||
|
||||
// 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);
|
||||
|
||||
const color_attachment = c.WGPURenderPassColorAttachment{
|
||||
@@ -180,5 +177,5 @@ pub fn draw(
|
||||
const cmd = c.wgpuCommandEncoderFinish(enc, null);
|
||||
defer c.wgpuCommandBufferRelease(cmd);
|
||||
|
||||
c.wgpuQueueSubmit(glloc.device.queue, 1, &cmd);
|
||||
c.wgpuQueueSubmit(gloc.device.queue, 1, &cmd);
|
||||
}
|
||||
|
||||
+35
-56
@@ -1,29 +1,26 @@
|
||||
const std = @import("std");
|
||||
const c = @import("utils.zig").c;
|
||||
const svOpt = @import("utils.zig").svOpt;
|
||||
const GpuAllocator = @import("GpuAllocator.zig");
|
||||
const GpuBuffer = @import("GpuBuffer.zig");
|
||||
const GpuTextureFormat = @import("lib.zig").GpuTextureFormat;
|
||||
const GpuTextureUsage = @import("lib.zig").GpuTextureUsage;
|
||||
|
||||
pub const GpuTextureDef = struct {
|
||||
label: ?[]const u8 = null,
|
||||
size: c.WGPUExtent3D,
|
||||
usage: std.EnumSet(GpuTextureUsage),
|
||||
format: GpuTextureFormat,
|
||||
};
|
||||
|
||||
raw: c.WGPUTexture,
|
||||
glloc: GpuAllocator,
|
||||
gloc: GpuAllocator,
|
||||
def: GpuTextureDef,
|
||||
|
||||
pub fn init(glloc: GpuAllocator, def: GpuTextureDef) !@This() {
|
||||
pub fn init(gloc: GpuAllocator, def: GpuTextureDef) !@This() {
|
||||
var use: u64 = 0;
|
||||
var iter = def.usage.iterator();
|
||||
while (iter.next()) |flag| use |= @intFromEnum(flag);
|
||||
|
||||
const desc = c.WGPUTextureDescriptor{
|
||||
.label = svOpt(def.label),
|
||||
.usage = use,
|
||||
.dimension = c.WGPUTextureDimension_2D,
|
||||
.size = def.size,
|
||||
@@ -31,13 +28,13 @@ pub fn init(glloc: GpuAllocator, def: GpuTextureDef) !@This() {
|
||||
.mipLevelCount = 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 {
|
||||
self.glloc.freeTexture(self.raw);
|
||||
self.gloc.freeTexture(self.raw);
|
||||
}
|
||||
|
||||
pub fn getConstMappedRange(self: @This(), offset: u64, size: u64) ?*const anyopaque {
|
||||
@@ -53,14 +50,10 @@ pub fn bytesSizeRow(self: @This()) u32 {
|
||||
}
|
||||
|
||||
/// Return a GpuBuffer containing a copy of the texture.
|
||||
pub fn buffCopy(self: @This(), glloc: GpuAllocator) !GpuBuffer {
|
||||
const buf = try GpuBuffer.init(glloc, .{
|
||||
.size = self.bytesSize(),
|
||||
.usage = .initMany(&.{ .CopyDst, .CopySrc }),
|
||||
.label = "texture_copy_buffer",
|
||||
});
|
||||
pub fn buffCopy(self: @This(), gloc: GpuAllocator) !GpuBuffer {
|
||||
const buf = try GpuBuffer.init(gloc, self.bytesSize(), .initMany(&.{ .CopyDst, .CopySrc }));
|
||||
|
||||
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);
|
||||
|
||||
const src_copy = c.WGPUTexelCopyTextureInfo{
|
||||
@@ -82,7 +75,7 @@ pub fn buffCopy(self: @This(), glloc: GpuAllocator) !GpuBuffer {
|
||||
|
||||
const cmd = c.wgpuCommandEncoderFinish(enc, null);
|
||||
defer c.wgpuCommandBufferRelease(cmd);
|
||||
c.wgpuQueueSubmit(glloc.device.queue, 1, &cmd);
|
||||
c.wgpuQueueSubmit(gloc.device.queue, 1, &cmd);
|
||||
|
||||
return buf;
|
||||
}
|
||||
@@ -109,56 +102,42 @@ pub fn load(
|
||||
) !void {
|
||||
const bytes = data.len * @sizeOf(T);
|
||||
|
||||
c.wgpuQueueWriteTexture(
|
||||
self.glloc.device.queue,
|
||||
&.{
|
||||
.texture = self.raw,
|
||||
.mipLevel = 0,
|
||||
.origin = .{ .x = 0, .y = 0, .z = 0 },
|
||||
.aspect = c.WGPUTextureAspect_All,
|
||||
},
|
||||
data.ptr,
|
||||
bytes,
|
||||
&.{
|
||||
.offset = 0,
|
||||
.bytesPerRow = self.bytesSizeRow(),
|
||||
.rowsPerImage = self.def.size.height,
|
||||
},
|
||||
&self.def.size,
|
||||
);
|
||||
if (bytes == self.size) {
|
||||
// Aligned path: direct download
|
||||
c.wgpuQueueWriteBuffer(self.gloc.device.queue, self.raw, 0, data.ptr, self.size);
|
||||
} else {
|
||||
// Unaligned path: Split the write into an aligned chunk and a padded remainder
|
||||
// to support arbitrary lengths without any allocations or large stack arrays.
|
||||
const aligned_part = (bytes / 4) * 4;
|
||||
if (aligned_part > 0) {
|
||||
c.wgpuQueueWriteBuffer(self.gloc.device.queue, self.raw, 0, data.ptr, aligned_part);
|
||||
}
|
||||
|
||||
var remainder_buf: [4]u8 = .{ 0, 0, 0, 0 };
|
||||
const data_bytes = std.mem.sliceAsBytes(data);
|
||||
@memcpy(remainder_buf[0 .. bytes - aligned_part], data_bytes[aligned_part..bytes]);
|
||||
|
||||
c.wgpuQueueWriteBuffer(self.gloc.device.queue, self.raw, aligned_part, &remainder_buf, 4);
|
||||
}
|
||||
}
|
||||
|
||||
// GPU to CPU
|
||||
pub fn read(self: @This(), alloc: std.mem.Allocator, T: type) ![]T {
|
||||
const out = try alloc.alloc(T, @divExact(self.size, @sizeOf(T)));
|
||||
|
||||
const staging = try init(self.glloc, .{
|
||||
.size = self.size,
|
||||
.usage = .initMany(&.{ .MapRead, .CopyDst }),
|
||||
.label = "texture_read_staging",
|
||||
});
|
||||
const staging = try init(
|
||||
self.gloc,
|
||||
self.size,
|
||||
.initMany(&.{ .MapRead, .CopyDst }),
|
||||
);
|
||||
defer staging.deinit();
|
||||
|
||||
const enc = c.wgpuDeviceCreateCommandEncoder(self.glloc.device.device, null) orelse return error.Encoder;
|
||||
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 enc = c.wgpuDeviceCreateCommandEncoder(self.gloc.device.device, null) orelse return error.Encoder;
|
||||
c.wgpuCommandEncoderCopyBufferToBuffer(enc, self.raw, 0, staging.raw, 0, self.size);
|
||||
const cmd = c.wgpuCommandEncoderFinish(enc, null);
|
||||
defer c.wgpuCommandEncoderRelease(enc);
|
||||
defer c.wgpuCommandBufferRelease(cmd);
|
||||
c.wgpuQueueSubmit(self.glloc.device.queue, 1, &cmd);
|
||||
c.wgpuQueueSubmit(self.gloc.device.queue, 1, &cmd);
|
||||
|
||||
var mapped = false;
|
||||
staging.mapAsync(
|
||||
@@ -167,7 +146,7 @@ pub fn read(self: @This(), alloc: std.mem.Allocator, T: type) ![]T {
|
||||
self.size,
|
||||
.{ .callback = onMapped, .userdata1 = &mapped },
|
||||
);
|
||||
while (!mapped) self.glloc.device.poll();
|
||||
while (!mapped) self.gloc.device.poll();
|
||||
|
||||
const ptr: [*]const T = @ptrCast(@alignCast(
|
||||
staging.getConstMappedRange(0, self.size),
|
||||
|
||||
@@ -1,35 +1,32 @@
|
||||
const std = @import("std");
|
||||
const c = @import("utils.zig").c;
|
||||
const svOpt = @import("utils.zig").svOpt;
|
||||
const GpuAllocator = @import("GpuAllocator.zig");
|
||||
const GpuTexture = @import("lib.zig").GpuTexture;
|
||||
const GpuTextureFormat = @import("lib.zig").GpuTextureFormat;
|
||||
const GpuTextureUsage = @import("lib.zig").GpuTextureUsage;
|
||||
|
||||
pub const GpuViewDef = struct {
|
||||
label: ?[]const u8 = null,
|
||||
usage: std.EnumSet(GpuTextureUsage) = .empty,
|
||||
format: GpuTextureFormat = .Undefined,
|
||||
};
|
||||
|
||||
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 iter = def.usage.iterator();
|
||||
while (iter.next()) |flag| use |= @intFromEnum(flag);
|
||||
|
||||
const raw = try glloc.allocTextureView(texture.raw, .{
|
||||
.label = svOpt(def.label),
|
||||
const raw = try gloc.allocTextureView(texture.raw, .{
|
||||
.format = @intFromEnum(def.format),
|
||||
.usage = use,
|
||||
.mipLevelCount = 1,
|
||||
.arrayLayerCount = 1,
|
||||
});
|
||||
return .{ .glloc = glloc, .raw = raw };
|
||||
return .{ .gloc = gloc, .raw = raw };
|
||||
}
|
||||
|
||||
pub fn deinit(self: @This()) void {
|
||||
self.glloc.freeTextureView(self.raw);
|
||||
self.gloc.freeTextureView(self.raw);
|
||||
}
|
||||
|
||||
@@ -3,17 +3,3 @@ pub const c = @cImport(@cInclude("wgpu.h"));
|
||||
pub fn sv(s: []const u8) c.WGPUStringView {
|
||||
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";
|
||||
}
|
||||
|
||||
Reference in New Issue
Block a user