Compare commits
19
Commits
d42c521a96
..
0.1.0
| Author | SHA1 | Date | |
|---|---|---|---|
|
|
6ec53bb909 | ||
|
|
6b933465b0 | ||
|
|
3d9cb4df04 | ||
|
|
c62bead6e9 | ||
|
|
9dbd668de9 | ||
|
|
50bf436fa5 | ||
|
|
df4eb36bc8 | ||
|
|
c7afe28cc6 | ||
|
|
d9cb4c4672 | ||
|
|
3a09e11f54 | ||
|
|
e608681ce1 | ||
|
|
d2b0735540 | ||
|
|
8e137c8f33 | ||
|
|
97d5f9001f | ||
|
|
d5e7f60926 | ||
|
|
2673aef0fd | ||
|
|
7ce6b9cd1d | ||
|
|
f5daf66784 | ||
|
|
c3166e552b |
@@ -0,0 +1,83 @@
|
|||||||
|
# Minimal Zig WebGPU Compute Library
|
||||||
|
|
||||||
|
This is a minimal, self-contained Zig library designed to simplify running compute shaders using WebGPU. It abstracts away much of the boilerplate required for GPU device initialization, memory management, and pipeline execution.
|
||||||
|
|
||||||
|
## Core Modules
|
||||||
|
|
||||||
|
The library exports five primary components:
|
||||||
|
|
||||||
|
* **`GpuDevice`**: Initializes the WebGPU instance, adapter, device, and queue. It is configured to prioritize high performance and automatically requests the `ShaderF16` feature if the adapter supports it. By default, it enforces a 2 GB VRAM limit.
|
||||||
|
* **`GpuArena` / `GpuAllocator`**: A memory management layer that tracks allocated VRAM bytes to prevent exceeding the device budget. The arena automatically destroys and releases all tracked WebGPU buffers when deinitialized.
|
||||||
|
* **`GpuBuffer`**: Wraps native WebGPU buffers. It automatically aligns buffer sizes forward to a multiple of 4 bytes. It provides a `.load()` method for CPU-to-GPU data transfers (handling both aligned and unaligned lengths smoothly) and a `.read()` method that utilizes a staging buffer to map GPU data back to the CPU.
|
||||||
|
* **`GpuProcess`**: Compiles WGSL source code into a compute pipeline. When running a process, it automatically splits the work into manageable chunks (up to 1 GB at a time) and dispatches workgroups of size 256.
|
||||||
|
|
||||||
|
## Quick Start Example
|
||||||
|
|
||||||
|
Below is a complete, self-contained example demonstrating how to initialize the GPU, load data, run a compute shader, and read the results back to the CPU:
|
||||||
|
|
||||||
|
```zig
|
||||||
|
|
||||||
|
const std = @import("std");
|
||||||
|
const gpu = @import("gpu");
|
||||||
|
const GpuDevice = gpu.GpuDevice;
|
||||||
|
const GpuArena = gpu.GpuArena;
|
||||||
|
const GpuBuffer = gpu.GpuBuffer;
|
||||||
|
const GpuProcess = gpu.GpuProcess;
|
||||||
|
|
||||||
|
pub fn main(init: std.process.Init) !void {
|
||||||
|
const allocator = init.gpa;
|
||||||
|
|
||||||
|
// 1. Open GPU Device
|
||||||
|
const device = try GpuDevice.init(.{});
|
||||||
|
defer device.deinit();
|
||||||
|
|
||||||
|
// 2. Create a GPU Arena to manage VRAM
|
||||||
|
var grena = GpuArena.init(allocator, device);
|
||||||
|
defer grena.deinit();
|
||||||
|
const gloc = grena.gpuAllocator();
|
||||||
|
|
||||||
|
// 3. Load the WGSL compute pipeline
|
||||||
|
const add_process = try GpuProcess.init(device, @embedFile("shaders/add.wgsl"));
|
||||||
|
defer add_process.deinit();
|
||||||
|
|
||||||
|
// 4. Setup CPU data
|
||||||
|
const len: usize = 16;
|
||||||
|
const data_a = try allocator.alloc(f16, len);
|
||||||
|
defer allocator.free(data_a);
|
||||||
|
const data_b = try allocator.alloc(f16, len);
|
||||||
|
defer allocator.free(data_b);
|
||||||
|
|
||||||
|
for (0..len) |i| {
|
||||||
|
data_a[i] = @floatFromInt(i);
|
||||||
|
data_b[i] = @floatFromInt(len - 1 - i);
|
||||||
|
}
|
||||||
|
|
||||||
|
// 5. Initialize raw GPU Buffers
|
||||||
|
// We pass the EnumSet inline using `.initMany` since the Enum itself isn't exported
|
||||||
|
const byte_size = len * @sizeOf(f16);
|
||||||
|
const buf_a = try GpuBuffer.init(gloc, byte_size, .initMany(&.{ .Storage, .CopyDst, .CopySrc }));
|
||||||
|
const buf_b = try GpuBuffer.init(gloc, byte_size, .initMany(&.{ .Storage, .CopyDst, .CopySrc }));
|
||||||
|
const buf_out = try GpuBuffer.init(gloc, byte_size, .initMany(&.{ .Storage, .CopyDst, .CopySrc }));
|
||||||
|
|
||||||
|
// Note: The buffers are safely tied to the GpuArena which will automatically
|
||||||
|
// release them at the end. You can also manually call buf_x.deinit() if desired.
|
||||||
|
|
||||||
|
// 6. Transfer data from CPU slices to GPU Buffers
|
||||||
|
try buf_a.load(f16, data_a);
|
||||||
|
try buf_b.load(f16, data_b);
|
||||||
|
|
||||||
|
// 7. Dispatch the Compute Process
|
||||||
|
// We pass the data type (f16) to allow GpuProcess to calculate chunks correctly
|
||||||
|
try add_process.run(gloc, f16, buf_a, buf_b, buf_out);
|
||||||
|
|
||||||
|
// 8. Map and copy the resulting buffer back to the CPU
|
||||||
|
const out = try buf_out.read(allocator, f16);
|
||||||
|
defer allocator.free(out);
|
||||||
|
|
||||||
|
std.debug.print("Result: {any}\n", .{out});
|
||||||
|
}
|
||||||
|
```
|
||||||
|
|
||||||
|
## Dependencies
|
||||||
|
|
||||||
|
* **`wgpu.h`**: The library relies on the WebGPU C API headers to bind to the native system graphics.
|
||||||
@@ -1,94 +1,60 @@
|
|||||||
// build.zig
|
|
||||||
// zig build run
|
|
||||||
//
|
|
||||||
// Expects wgpu-native pre-built in libs/wgpu-native/:
|
|
||||||
// include/wgpu.h
|
|
||||||
// lib/libwgpu_native.a (or .so / .dylib / .dll)
|
|
||||||
//
|
|
||||||
// Download release: https://github.com/gfx-rs/wgpu-native/releases
|
|
||||||
// Pick the archive matching your OS/arch, e.g.:
|
|
||||||
// wgpu-linux-x86_64-release.zip → libwgpu_native.a + wgpu.h
|
|
||||||
// wgpu-macos-aarch64-release.zip
|
|
||||||
// wgpu-windows-x86_64-msvc-release.zip
|
|
||||||
|
|
||||||
const std = @import("std");
|
const std = @import("std");
|
||||||
|
|
||||||
pub fn build(b: *std.Build) void {
|
pub fn build(b: *std.Build) !void {
|
||||||
const target = b.standardTargetOptions(.{});
|
const target = b.standardTargetOptions(.{});
|
||||||
const optimize = b.standardOptimizeOption(.{});
|
const optimize = b.standardOptimizeOption(.{});
|
||||||
|
|
||||||
// 1. Define the module so other projects can import it
|
// Define the module so other projects can import it
|
||||||
_ = b.addModule("zig-wgpu", .{
|
const mod = b.addModule("zig-wgpu", .{
|
||||||
.root_source_file = b.path("src/lib.zig"),
|
.root_source_file = b.path("src/lib.zig"),
|
||||||
|
.target = target,
|
||||||
});
|
});
|
||||||
|
|
||||||
const exe = b.addExecutable(.{
|
mod.addIncludePath(b.path("libs/wgpu-native/include"));
|
||||||
.root_module = b.createModule(.{
|
mod.addLibraryPath(b.path("libs/wgpu-native/lib"));
|
||||||
.root_source_file = b.path("src/bench.zig"),
|
mod.addObjectFile(b.path("libs/wgpu-native/lib/libwgpu_native.a"));
|
||||||
.link_libc = true,
|
|
||||||
.target = target,
|
|
||||||
.optimize = optimize,
|
|
||||||
}),
|
|
||||||
.name = "bench",
|
|
||||||
});
|
|
||||||
|
|
||||||
// wgpu-native headers + pre-built static library
|
|
||||||
exe.root_module.addIncludePath(b.path("libs/wgpu-native/include"));
|
|
||||||
exe.root_module.addLibraryPath(b.path("libs/wgpu-native/lib"));
|
|
||||||
exe.root_module.addObjectFile(b.path("libs/wgpu-native/lib/libwgpu_native.a"));
|
|
||||||
|
|
||||||
// Platform-specific system frameworks needed by wgpu-native
|
// Platform-specific system frameworks needed by wgpu-native
|
||||||
const t = target.result;
|
const t = target.result;
|
||||||
if (t.os.tag == .macos) {
|
if (t.os.tag == .macos) {
|
||||||
exe.root_module.linkFramework("Metal", .{});
|
mod.linkFramework("Metal", .{});
|
||||||
exe.root_module.linkFramework("QuartzCore", .{});
|
mod.linkFramework("QuartzCore", .{});
|
||||||
exe.root_module.linkFramework("Foundation", .{});
|
mod.linkFramework("Foundation", .{});
|
||||||
exe.root_module.linkFramework("CoreGraphics", .{});
|
mod.linkFramework("CoreGraphics", .{});
|
||||||
} else if (t.os.tag == .windows) {
|
} else if (t.os.tag == .windows) {
|
||||||
exe.root_module.linkSystemLibrary("d3d12", .{});
|
mod.linkSystemLibrary("d3d12", .{});
|
||||||
exe.root_module.linkSystemLibrary("dxgi", .{});
|
mod.linkSystemLibrary("dxgi", .{});
|
||||||
exe.root_module.linkSystemLibrary("user32", .{});
|
mod.linkSystemLibrary("user32", .{});
|
||||||
} else {
|
} else {
|
||||||
exe.root_module.linkSystemLibrary("vulkan", .{});
|
mod.linkSystemLibrary("vulkan", .{});
|
||||||
exe.root_module.linkSystemLibrary("gcc_s", .{});
|
mod.linkSystemLibrary("gcc_s", .{});
|
||||||
}
|
}
|
||||||
|
|
||||||
b.installArtifact(exe);
|
var threaded: std.Io.Threaded = .init_single_threaded;
|
||||||
|
const io = threaded.io();
|
||||||
|
|
||||||
const run = b.addRunArtifact(exe);
|
var buf: [1024]u8 = undefined;
|
||||||
run.step.dependOn(b.getInstallStep());
|
const exemples = try std.Io.Dir.cwd().openDir(io, "examples", .{ .access_sub_paths = false, .iterate = true });
|
||||||
b.step("bench", "Benchmark a simple add vector.").dependOn(&run.step);
|
var iter = exemples.iterate();
|
||||||
|
while (try iter.next(io)) |entry| {
|
||||||
|
if (entry.kind != .file) continue;
|
||||||
|
if (!std.mem.eql(u8, entry.name[entry.name.len - 4 ..], ".zig")) continue;
|
||||||
|
|
||||||
const exe_examp = b.addExecutable(.{
|
const exe = b.addExecutable(.{
|
||||||
.root_module = b.createModule(.{
|
.name = entry.name[0 .. entry.name.len - 4],
|
||||||
.root_source_file = b.path("src/example.zig"),
|
.root_module = b.createModule(.{
|
||||||
.link_libc = true,
|
.root_source_file = b.path(try std.fmt.bufPrint(&buf, "examples/{s}", .{entry.name})),
|
||||||
.target = target,
|
.target = target,
|
||||||
.optimize = optimize,
|
.optimize = optimize,
|
||||||
}),
|
.imports = &.{},
|
||||||
.name = "bench",
|
}),
|
||||||
});
|
});
|
||||||
|
exe.root_module.addImport("gpu", mod);
|
||||||
|
|
||||||
// wgpu-native headers + pre-built static library
|
b.installArtifact(exe);
|
||||||
exe_examp.root_module.addIncludePath(b.path("libs/wgpu-native/include"));
|
|
||||||
exe_examp.root_module.addLibraryPath(b.path("libs/wgpu-native/lib"));
|
|
||||||
exe_examp.root_module.addObjectFile(b.path("libs/wgpu-native/lib/libwgpu_native.a"));
|
|
||||||
|
|
||||||
if (t.os.tag == .macos) {
|
const run_step = b.step(entry.name[0 .. entry.name.len - 4], try std.fmt.bufPrint(&buf, "Run {s} demo", .{entry.name}));
|
||||||
exe_examp.root_module.linkFramework("Metal", .{});
|
const run_cmd = b.addRunArtifact(exe);
|
||||||
exe_examp.root_module.linkFramework("QuartzCore", .{});
|
run_step.dependOn(&run_cmd.step);
|
||||||
exe_examp.root_module.linkFramework("Foundation", .{});
|
|
||||||
exe_examp.root_module.linkFramework("CoreGraphics", .{});
|
|
||||||
} else if (t.os.tag == .windows) {
|
|
||||||
exe_examp.root_module.linkSystemLibrary("d3d12", .{});
|
|
||||||
exe_examp.root_module.linkSystemLibrary("dxgi", .{});
|
|
||||||
exe_examp.root_module.linkSystemLibrary("user32", .{});
|
|
||||||
} else {
|
|
||||||
exe_examp.root_module.linkSystemLibrary("vulkan", .{});
|
|
||||||
exe_examp.root_module.linkSystemLibrary("gcc_s", .{});
|
|
||||||
}
|
}
|
||||||
|
|
||||||
const examp = b.addRunArtifact(exe_examp);
|
|
||||||
run.step.dependOn(b.getInstallStep());
|
|
||||||
b.step("example", "Run basic example.").dependOn(&examp.step);
|
|
||||||
}
|
}
|
||||||
|
|||||||
+3
-33
@@ -1,43 +1,13 @@
|
|||||||
.{
|
.{
|
||||||
// This is the default name used by packages depending on this one. For
|
|
||||||
// example, when a user runs `zig fetch --save <url>`, this field is used
|
|
||||||
// as the key in the `dependencies` table. Although the user can choose a
|
|
||||||
// different name, most users will stick with this provided value.
|
|
||||||
//
|
|
||||||
// It is redundant to include "zig" in this name because it is already
|
|
||||||
// within the Zig package namespace.
|
|
||||||
.name = .zig_wgpu,
|
.name = .zig_wgpu,
|
||||||
// This is a [Semantic Version](https://semver.org/).
|
.version = "0.1.0",
|
||||||
// In a future version of Zig it will be used for package deduplication.
|
.fingerprint = 0x5d0e853acbc0c2c6,
|
||||||
.version = "0.0.0",
|
|
||||||
// Together with name, this represents a globally unique package
|
|
||||||
// identifier. This field is generated by the Zig toolchain when the
|
|
||||||
// package is first created, and then *never changes*. This allows
|
|
||||||
// unambiguous detection of one package being an updated version of
|
|
||||||
// another.
|
|
||||||
//
|
|
||||||
// When forking a Zig project, this id should be regenerated (delete the
|
|
||||||
// field and run `zig build`) if the upstream project is still maintained.
|
|
||||||
// Otherwise, the fork is *hostile*, attempting to take control over the
|
|
||||||
// original project's identity. Thus it is recommended to leave the comment
|
|
||||||
// on the following line intact, so that it shows up in code reviews that
|
|
||||||
// modify the field.
|
|
||||||
.fingerprint = 0x5d0e853acbc0c2c6, // Changing this has security and trust implications.
|
|
||||||
// Tracks the earliest Zig version that the package considers to be a
|
|
||||||
// supported use case.
|
|
||||||
.minimum_zig_version = "0.16.0",
|
.minimum_zig_version = "0.16.0",
|
||||||
// This field is optional.
|
|
||||||
// Each dependency must either provide a `url` and `hash`, or a `path`.
|
|
||||||
// `zig build --fetch` can be used to fetch all dependencies of a package, recursively.
|
|
||||||
// Once all dependencies are fetched, `zig build` no longer requires
|
|
||||||
// internet connectivity.
|
|
||||||
.dependencies = .{},
|
.dependencies = .{},
|
||||||
.paths = .{
|
.paths = .{
|
||||||
"build.zig",
|
"build.zig",
|
||||||
"build.zig.zon",
|
"build.zig.zon",
|
||||||
"src",
|
"src",
|
||||||
// For example...
|
"libs",
|
||||||
//"LICENSE",
|
|
||||||
//"README.md",
|
|
||||||
},
|
},
|
||||||
}
|
}
|
||||||
|
|||||||
@@ -0,0 +1,59 @@
|
|||||||
|
const std = @import("std");
|
||||||
|
const gpu = @import("gpu");
|
||||||
|
const GpuDevice = gpu.GpuDevice;
|
||||||
|
const GpuArena = gpu.GpuArena;
|
||||||
|
const GpuBuffer = gpu.GpuBuffer;
|
||||||
|
const GpuProcess = gpu.GpuProcess;
|
||||||
|
|
||||||
|
pub fn main(init: std.process.Init) !void {
|
||||||
|
const allocator = init.gpa;
|
||||||
|
|
||||||
|
// 1. Open GPU Device
|
||||||
|
const device = try GpuDevice.init(.{});
|
||||||
|
defer device.deinit();
|
||||||
|
|
||||||
|
// 2. Create a GPU Arena to manage VRAM
|
||||||
|
var grena = GpuArena.init(allocator, device);
|
||||||
|
defer grena.deinit();
|
||||||
|
const gloc = grena.gpuAllocator();
|
||||||
|
|
||||||
|
// 3. Load the WGSL compute pipeline
|
||||||
|
const add_process = try GpuProcess.init(device, @embedFile("shaders/add.wgsl"));
|
||||||
|
defer add_process.deinit();
|
||||||
|
|
||||||
|
// 4. Setup CPU data
|
||||||
|
const len: usize = 16;
|
||||||
|
const data_a = try allocator.alloc(f16, len);
|
||||||
|
defer allocator.free(data_a);
|
||||||
|
const data_b = try allocator.alloc(f16, len);
|
||||||
|
defer allocator.free(data_b);
|
||||||
|
|
||||||
|
for (0..len) |i| {
|
||||||
|
data_a[i] = @floatFromInt(i);
|
||||||
|
data_b[i] = @floatFromInt(len - 1 - i);
|
||||||
|
}
|
||||||
|
|
||||||
|
// 5. Initialize raw GPU Buffers
|
||||||
|
// We pass the EnumSet inline using `.initMany` since the Enum itself isn't exported
|
||||||
|
const byte_size = len * @sizeOf(f16);
|
||||||
|
const buf_a = try GpuBuffer.init(gloc, byte_size, .initMany(&.{ .Storage, .CopyDst, .CopySrc }));
|
||||||
|
const buf_b = try GpuBuffer.init(gloc, byte_size, .initMany(&.{ .Storage, .CopyDst, .CopySrc }));
|
||||||
|
const buf_out = try GpuBuffer.init(gloc, byte_size, .initMany(&.{ .Storage, .CopyDst, .CopySrc }));
|
||||||
|
|
||||||
|
// Note: The buffers are safely tied to the GpuArena which will automatically
|
||||||
|
// release them at the end. You can also manually call buf_x.deinit() if desired.
|
||||||
|
|
||||||
|
// 6. Transfer data from CPU slices to GPU Buffers
|
||||||
|
try buf_a.load(f16, data_a);
|
||||||
|
try buf_b.load(f16, data_b);
|
||||||
|
|
||||||
|
// 7. Dispatch the Compute Process
|
||||||
|
// We pass the data type (f16) to allow GpuProcess to calculate chunks correctly
|
||||||
|
try add_process.run(gloc, f16, buf_a, buf_b, buf_out);
|
||||||
|
|
||||||
|
// 8. Map and copy the resulting buffer back to the CPU
|
||||||
|
const out = try buf_out.read(allocator, f16);
|
||||||
|
defer allocator.free(out);
|
||||||
|
|
||||||
|
std.debug.print("Result: {any}\n", .{out});
|
||||||
|
}
|
||||||
@@ -1,37 +1,75 @@
|
|||||||
const std = @import("std");
|
const std = @import("std");
|
||||||
const GpuDevice = @import("GpuDevice.zig");
|
const gpu = @import("gpu");
|
||||||
const GpuAllocator = @import("GpuAllocator.zig");
|
const GpuDevice = gpu.GpuDevice;
|
||||||
const GpuPipeline = @import("GpuPipeline.zig");
|
const GpuArena = gpu.GpuArena;
|
||||||
const Vec = @import("Vec.zig");
|
const GpuAllocator = gpu.GpuAllocator;
|
||||||
|
const GpuBuffer = gpu.GpuBuffer;
|
||||||
|
const GpuProcess = gpu.GpuProcess;
|
||||||
|
|
||||||
const c = @import("utils.zig").c;
|
/// Minimal implementation of a f16 Vector
|
||||||
|
const Vec = struct {
|
||||||
|
buf: GpuBuffer,
|
||||||
|
len: usize,
|
||||||
|
|
||||||
|
// Changed: gloc is passed by value (const)
|
||||||
|
pub fn initZero(gloc: GpuAllocator, len: usize) !Vec {
|
||||||
|
return .{
|
||||||
|
.buf = try GpuBuffer.init(
|
||||||
|
gloc,
|
||||||
|
len * @sizeOf(f16),
|
||||||
|
.initMany(&.{ .Storage, .CopyDst, .CopySrc }),
|
||||||
|
),
|
||||||
|
.len = len,
|
||||||
|
};
|
||||||
|
}
|
||||||
|
|
||||||
|
// Changed: gloc is passed by value
|
||||||
|
pub fn initLoad(gloc: GpuAllocator, data: []const f16) !Vec {
|
||||||
|
var self = try initZero(gloc, data.len);
|
||||||
|
try self.load(data); // Direct access via the interface copy
|
||||||
|
return self;
|
||||||
|
}
|
||||||
|
|
||||||
|
pub fn deinit(self: Vec) void {
|
||||||
|
self.buf.deinit();
|
||||||
|
}
|
||||||
|
|
||||||
|
/// CPU to GPU.
|
||||||
|
pub fn load(self: Vec, data: []const f16) !void {
|
||||||
|
try self.buf.load(f16, data);
|
||||||
|
}
|
||||||
|
|
||||||
|
// Changed: gloc is passed by value instead of *GpuAllocator
|
||||||
|
pub fn run(self: Vec, gloc: GpuAllocator, other: Vec, process: GpuProcess) !Vec {
|
||||||
|
std.debug.assert(self.len == other.len);
|
||||||
|
|
||||||
|
const result = try Vec.initZero(gloc, self.len);
|
||||||
|
errdefer result.deinit();
|
||||||
|
|
||||||
|
try process.run(gloc, f16, self.buf, other.buf, result.buf);
|
||||||
|
return result;
|
||||||
|
}
|
||||||
|
|
||||||
|
// Changed: gloc is passed by value instead of *GpuAllocator
|
||||||
|
pub fn read(self: Vec, alloc: std.mem.Allocator) ![]f16 {
|
||||||
|
return self.buf.read(alloc, f16);
|
||||||
|
}
|
||||||
|
};
|
||||||
|
|
||||||
pub fn main(init: std.process.Init) !void {
|
pub fn main(init: std.process.Init) !void {
|
||||||
const device = try GpuDevice.init(.{ .vram_bytes_limit = 4 * 1024 * 1024 * 1024 });
|
const device = try GpuDevice.init(.{ .vram_bytes_limit = 4 * 1024 * 1024 * 1024 });
|
||||||
defer device.deinit();
|
defer device.deinit();
|
||||||
|
|
||||||
var gloc = try GpuAllocator.init(init.gpa, device);
|
var grena = GpuArena.init(init.gpa, device);
|
||||||
defer gloc.deinit();
|
defer grena.deinit();
|
||||||
|
|
||||||
const add_pip = try GpuPipeline.init(device, @embedFile("shaders/add.wgsl"));
|
const gloc = grena.gpuAllocator();
|
||||||
|
|
||||||
|
const add_pip = try GpuProcess.init(device, @embedFile("shaders/add.wgsl"));
|
||||||
defer add_pip.deinit();
|
defer add_pip.deinit();
|
||||||
|
|
||||||
const allocator = init.gpa;
|
const allocator = init.gpa;
|
||||||
|
|
||||||
// --- WARM-UP PHASE ---
|
|
||||||
{
|
|
||||||
var warmup_a = [_]f16{1.0};
|
|
||||||
var warmup_b = [_]f16{1.0};
|
|
||||||
const wa = try Vec.initLoad(&gloc, &warmup_a);
|
|
||||||
defer wa.deinit();
|
|
||||||
const wb = try Vec.initLoad(&gloc, &warmup_b);
|
|
||||||
defer wb.deinit();
|
|
||||||
const wsum = try wa.run(&gloc, wb, add_pip);
|
|
||||||
defer wsum.deinit();
|
|
||||||
const wout = try wsum.read(&gloc, allocator);
|
|
||||||
defer allocator.free(wout);
|
|
||||||
}
|
|
||||||
|
|
||||||
const sizes = [_]usize{
|
const sizes = [_]usize{
|
||||||
1,
|
1,
|
||||||
256,
|
256,
|
||||||
@@ -41,9 +79,9 @@ pub fn main(init: std.process.Init) !void {
|
|||||||
4 * 4 * 4 * 1024,
|
4 * 4 * 4 * 1024,
|
||||||
4 * 4 * 4 * 4 * 1024,
|
4 * 4 * 4 * 4 * 1024,
|
||||||
1024 * 1024,
|
1024 * 1024,
|
||||||
// 4 * 1024 * 1024,
|
4 * 1024 * 1024,
|
||||||
// 4 * 4 * 1024 * 1024,
|
4 * 4 * 1024 * 1024,
|
||||||
// 4 * 4 * 4 * 1024 * 1024,
|
4 * 4 * 4 * 1024 * 1024,
|
||||||
// 4 * 4 * 4 * 4 * 1024 * 1024,
|
// 4 * 4 * 4 * 4 * 1024 * 1024,
|
||||||
// 4 * 4 * 4 * 4 * 4 * 1024 * 1024,
|
// 4 * 4 * 4 * 4 * 4 * 1024 * 1024,
|
||||||
};
|
};
|
||||||
@@ -78,9 +116,9 @@ pub fn main(init: std.process.Init) !void {
|
|||||||
// --- 1. GPU ALLOCATION PHASE ---
|
// --- 1. GPU ALLOCATION PHASE ---
|
||||||
const alloc_start = std.Io.Clock.awake.now(init.io);
|
const alloc_start = std.Io.Clock.awake.now(init.io);
|
||||||
|
|
||||||
const a = try Vec.initLoad(&gloc, data_a);
|
const a = try Vec.initLoad(gloc, data_a);
|
||||||
defer a.deinit();
|
defer a.deinit();
|
||||||
const b = try Vec.initLoad(&gloc, data_b);
|
const b = try Vec.initLoad(gloc, data_b);
|
||||||
defer b.deinit();
|
defer b.deinit();
|
||||||
|
|
||||||
const alloc_duration = alloc_start.durationTo(std.Io.Clock.awake.now(init.io));
|
const alloc_duration = alloc_start.durationTo(std.Io.Clock.awake.now(init.io));
|
||||||
@@ -90,15 +128,15 @@ pub fn main(init: std.process.Init) !void {
|
|||||||
// --- 2. COMPUTE PHASE ---
|
// --- 2. COMPUTE PHASE ---
|
||||||
const compute_start = std.Io.Clock.awake.now(init.io);
|
const compute_start = std.Io.Clock.awake.now(init.io);
|
||||||
|
|
||||||
const sum = try a.run(&gloc, b, add_pip);
|
const sum = try a.run(gloc, b, add_pip);
|
||||||
defer sum.deinit();
|
defer sum.deinit();
|
||||||
|
|
||||||
// All 3 buffers (a, b, sum) are currently resident in VRAM here.
|
// All 3 buffers (a, b, sum) are currently resident in VRAM here.
|
||||||
// Querying now catches the true peak allocation step.
|
// Querying now catches the true peak allocation step.
|
||||||
if (gloc.allocated_vram_bytes > peak_vram_bytes)
|
if (grena.allocated_vram_bytes > peak_vram_bytes)
|
||||||
peak_vram_bytes = gloc.allocated_vram_bytes;
|
peak_vram_bytes = grena.allocated_vram_bytes;
|
||||||
|
|
||||||
_ = c.wgpuDevicePoll(device.device, 1, null);
|
device.poll();
|
||||||
|
|
||||||
const compute_duration = compute_start.durationTo(std.Io.Clock.awake.now(init.io));
|
const compute_duration = compute_start.durationTo(std.Io.Clock.awake.now(init.io));
|
||||||
const compute_ns = @as(u64, @intCast(compute_duration.toNanoseconds()));
|
const compute_ns = @as(u64, @intCast(compute_duration.toNanoseconds()));
|
||||||
@@ -107,7 +145,7 @@ pub fn main(init: std.process.Init) !void {
|
|||||||
// --- 3. TRANSFER PHASE (Device -> Host) ---
|
// --- 3. TRANSFER PHASE (Device -> Host) ---
|
||||||
const transfer_start = std.Io.Clock.awake.now(init.io);
|
const transfer_start = std.Io.Clock.awake.now(init.io);
|
||||||
|
|
||||||
const out = try sum.read(&gloc, allocator);
|
const out = try sum.read(allocator);
|
||||||
defer allocator.free(out);
|
defer allocator.free(out);
|
||||||
|
|
||||||
const transfer_duration = transfer_start.durationTo(std.Io.Clock.awake.now(init.io));
|
const transfer_duration = transfer_start.durationTo(std.Io.Clock.awake.now(init.io));
|
||||||
@@ -3,11 +3,7 @@ enable f16;
|
|||||||
@group(0) @binding(0) var<storage, read> A: array<f16>;
|
@group(0) @binding(0) var<storage, read> A: array<f16>;
|
||||||
@group(0) @binding(1) var<storage, read> B: array<f16>;
|
@group(0) @binding(1) var<storage, read> B: array<f16>;
|
||||||
@group(0) @binding(2) var<storage, read_write> C: array<f16>;
|
@group(0) @binding(2) var<storage, read_write> C: array<f16>;
|
||||||
|
@group(0) @binding(3) var<uniform> size: u32;
|
||||||
struct TensorInfo {
|
|
||||||
size: u32,
|
|
||||||
};
|
|
||||||
@group(0) @binding(3) var<uniform> info: TensorInfo;
|
|
||||||
|
|
||||||
@compute @workgroup_size(256)
|
@compute @workgroup_size(256)
|
||||||
fn main(
|
fn main(
|
||||||
@@ -21,7 +17,7 @@ fn main(
|
|||||||
var index = global_id.x;
|
var index = global_id.x;
|
||||||
|
|
||||||
// 3. Stride through the tensor elements
|
// 3. Stride through the tensor elements
|
||||||
while (index < info.size) {
|
while (index < size) {
|
||||||
C[index] = A[index] + B[index];
|
C[index] = A[index] + B[index];
|
||||||
index += total_threads; // Jump forward by the total thread count
|
index += total_threads; // Jump forward by the total thread count
|
||||||
}
|
}
|
||||||
+15
-49
@@ -1,58 +1,24 @@
|
|||||||
|
// GpuAllocator.zig
|
||||||
const std = @import("std");
|
const std = @import("std");
|
||||||
const GpuDevice = @import("GpuDevice.zig");
|
const GpuDevice = @import("GpuDevice.zig");
|
||||||
const GpuBuffer = @import("GpuBuffer.zig");
|
|
||||||
const c = @import("utils.zig").c;
|
const c = @import("utils.zig").c;
|
||||||
|
|
||||||
|
const GpuAllocator = @This();
|
||||||
|
|
||||||
|
/// The function definitions our underlying implementations must satisfy
|
||||||
|
pub const VTable = struct {
|
||||||
|
alloc: *const fn (ctx: *anyopaque, bytes: u64, usage: c.WGPUBufferUsage) anyerror!c.WGPUBuffer,
|
||||||
|
free: *const fn (ctx: *anyopaque, buf_raw: c.WGPUBuffer, size: u64) void,
|
||||||
|
};
|
||||||
|
|
||||||
device: GpuDevice,
|
device: GpuDevice,
|
||||||
tracked_buffers: std.AutoHashMap(c.WGPUBuffer, void),
|
ptr: *anyopaque,
|
||||||
allocated_vram_bytes: u64 = 0,
|
vtable: *const VTable,
|
||||||
|
|
||||||
pub fn init(cpu_allocator: std.mem.Allocator, device: GpuDevice) !@This() {
|
pub fn allocBuffer(self: GpuAllocator, bytes: u64, usage: c.WGPUBufferUsage) !c.WGPUBuffer {
|
||||||
return .{
|
return self.vtable.alloc(self.ptr, bytes, usage);
|
||||||
.device = device,
|
|
||||||
.tracked_buffers = .init(cpu_allocator),
|
|
||||||
};
|
|
||||||
}
|
}
|
||||||
|
|
||||||
pub fn deinit(self: *@This()) void {
|
pub fn freeBuffer(self: GpuAllocator, buf_raw: c.WGPUBuffer, size: u64) void {
|
||||||
var it = self.tracked_buffers.keyIterator();
|
self.vtable.free(self.ptr, buf_raw, size);
|
||||||
while (it.next()) |buf_ptr| {
|
|
||||||
const buf = buf_ptr.*;
|
|
||||||
c.wgpuBufferDestroy(buf);
|
|
||||||
c.wgpuBufferRelease(buf);
|
|
||||||
}
|
|
||||||
self.tracked_buffers.deinit();
|
|
||||||
}
|
|
||||||
|
|
||||||
pub fn registerBuffer(
|
|
||||||
self: *@This(),
|
|
||||||
bytes: u64,
|
|
||||||
usage: c.WGPUBufferUsage,
|
|
||||||
) !c.WGPUBuffer {
|
|
||||||
if (bytes > self.device.limits.maxBufferSize)
|
|
||||||
return error.SingleBufferExceedsLimit;
|
|
||||||
|
|
||||||
if (bytes + self.allocated_vram_bytes > self.device.config.vram_bytes_limit)
|
|
||||||
return error.ExceedsVramBudget;
|
|
||||||
|
|
||||||
const buf = c.wgpuDeviceCreateBuffer(self.device.device, &.{
|
|
||||||
.usage = usage,
|
|
||||||
.size = bytes,
|
|
||||||
}) orelse return error.BufferAlloc;
|
|
||||||
errdefer {
|
|
||||||
c.wgpuBufferDestroy(buf);
|
|
||||||
c.wgpuBufferRelease(buf);
|
|
||||||
}
|
|
||||||
|
|
||||||
try self.tracked_buffers.put(buf, {});
|
|
||||||
self.allocated_vram_bytes += bytes;
|
|
||||||
return buf;
|
|
||||||
}
|
|
||||||
|
|
||||||
pub fn unregisterAndDestroyBuffer(self: *@This(), buf: GpuBuffer) void {
|
|
||||||
if (self.tracked_buffers.remove(buf.raw)) {
|
|
||||||
c.wgpuBufferDestroy(buf.raw);
|
|
||||||
c.wgpuBufferRelease(buf.raw);
|
|
||||||
self.allocated_vram_bytes -= buf.size;
|
|
||||||
}
|
|
||||||
}
|
}
|
||||||
|
|||||||
@@ -0,0 +1,72 @@
|
|||||||
|
// GpuArena.zig
|
||||||
|
const std = @import("std");
|
||||||
|
const GpuDevice = @import("GpuDevice.zig");
|
||||||
|
const GpuAllocator = @import("GpuAllocator.zig");
|
||||||
|
const c = @import("utils.zig").c;
|
||||||
|
|
||||||
|
const GpuArena = @This();
|
||||||
|
|
||||||
|
device: GpuDevice,
|
||||||
|
tracked_buffers: std.AutoHashMap(c.WGPUBuffer, void),
|
||||||
|
allocated_vram_bytes: u64 = 0,
|
||||||
|
|
||||||
|
pub fn init(cpu_allocator: std.mem.Allocator, device: GpuDevice) GpuArena {
|
||||||
|
return .{
|
||||||
|
.device = device,
|
||||||
|
.tracked_buffers = .init(cpu_allocator),
|
||||||
|
};
|
||||||
|
}
|
||||||
|
|
||||||
|
pub fn deinit(self: *GpuArena) void {
|
||||||
|
var it = self.tracked_buffers.keyIterator();
|
||||||
|
while (it.next()) |buf_ptr| {
|
||||||
|
c.wgpuBufferDestroy(buf_ptr.*);
|
||||||
|
c.wgpuBufferRelease(buf_ptr.*);
|
||||||
|
}
|
||||||
|
self.tracked_buffers.deinit();
|
||||||
|
}
|
||||||
|
|
||||||
|
/// Returns the type-erased immutable interface wrapper
|
||||||
|
pub fn gpuAllocator(self: *GpuArena) GpuAllocator {
|
||||||
|
return .{
|
||||||
|
.device = self.device,
|
||||||
|
.ptr = self,
|
||||||
|
.vtable = &.{
|
||||||
|
.alloc = alloc,
|
||||||
|
.free = free,
|
||||||
|
},
|
||||||
|
};
|
||||||
|
}
|
||||||
|
|
||||||
|
fn alloc(ctx: *anyopaque, bytes: u64, usage: c.WGPUBufferUsage) anyerror!c.WGPUBuffer {
|
||||||
|
const self: *GpuArena = @ptrCast(@alignCast(ctx));
|
||||||
|
|
||||||
|
if (bytes > self.device.limits.maxBufferSize)
|
||||||
|
return error.SingleBufferExceedsLimit;
|
||||||
|
|
||||||
|
if (bytes + self.allocated_vram_bytes > self.device.config.vram_bytes_limit)
|
||||||
|
return error.ExceedsVramBudget;
|
||||||
|
|
||||||
|
const buf = c.wgpuDeviceCreateBuffer(self.device.device, &.{
|
||||||
|
.usage = usage,
|
||||||
|
.size = bytes,
|
||||||
|
}) orelse return error.BufferAlloc;
|
||||||
|
errdefer {
|
||||||
|
c.wgpuBufferDestroy(buf);
|
||||||
|
c.wgpuBufferRelease(buf);
|
||||||
|
}
|
||||||
|
|
||||||
|
try self.tracked_buffers.put(buf, {});
|
||||||
|
self.allocated_vram_bytes += bytes;
|
||||||
|
return buf;
|
||||||
|
}
|
||||||
|
|
||||||
|
fn free(ctx: *anyopaque, buf_raw: c.WGPUBuffer, size: u64) void {
|
||||||
|
const self: *GpuArena = @ptrCast(@alignCast(ctx));
|
||||||
|
|
||||||
|
if (self.tracked_buffers.remove(buf_raw)) {
|
||||||
|
c.wgpuBufferDestroy(buf_raw);
|
||||||
|
c.wgpuBufferRelease(buf_raw);
|
||||||
|
self.allocated_vram_bytes -= size;
|
||||||
|
}
|
||||||
|
}
|
||||||
+105
-17
@@ -3,29 +3,50 @@ const c = @import("utils.zig").c;
|
|||||||
const GpuAllocator = @import("GpuAllocator.zig");
|
const GpuAllocator = @import("GpuAllocator.zig");
|
||||||
|
|
||||||
raw: c.WGPUBuffer,
|
raw: c.WGPUBuffer,
|
||||||
size: u64,
|
size: u64, // Now tracks the 4-byte aligned size directly
|
||||||
usage: c.WGPUBufferUsage,
|
usage: c.WGPUBufferUsage,
|
||||||
gloc: *GpuAllocator,
|
gloc: GpuAllocator,
|
||||||
|
|
||||||
|
const BufferUsage = enum(u64) {
|
||||||
|
None = 0x0000000000000000,
|
||||||
|
MapRead = 0x0000000000000001,
|
||||||
|
MapWrite = 0x0000000000000002,
|
||||||
|
CopySrc = 0x0000000000000004,
|
||||||
|
CopyDst = 0x0000000000000008,
|
||||||
|
Index = 0x0000000000000010,
|
||||||
|
Vertex = 0x0000000000000020,
|
||||||
|
Uniform = 0x0000000000000040,
|
||||||
|
Storage = 0x0000000000000080,
|
||||||
|
Indirect = 0x0000000000000100,
|
||||||
|
QueryResolve = 0x0000000000000200,
|
||||||
|
};
|
||||||
|
|
||||||
/// Allocates the underlying WebGPU handle and registers it to the parent GpuAllocator
|
/// Allocates the underlying WebGPU handle and registers it to the parent GpuAllocator
|
||||||
pub fn init(gloc: *GpuAllocator, T: type, len: usize, usage: c.WGPUBufferUsage) !@This() {
|
pub fn init(gloc: GpuAllocator, size: u64, usage: std.EnumSet(BufferUsage)) !@This() {
|
||||||
switch (@typeInfo(T)) {
|
var use: u64 = 0;
|
||||||
.int, .float => {},
|
var iter = usage.iterator();
|
||||||
else => @compileError("GpuBuffer can only use int and float type"),
|
while (iter.next()) |flag| use |= @intFromEnum(flag);
|
||||||
}
|
|
||||||
const bytes = @sizeOf(T) * len;
|
// Automatically align the buffer size forward to a multiple of 4 bytes under the hood
|
||||||
const raw_handle = try gloc.registerBuffer(bytes, usage);
|
const aligned_size = std.mem.alignForward(u64, size, 4);
|
||||||
|
|
||||||
|
const raw_handle = try gloc.allocBuffer(aligned_size, use);
|
||||||
return .{
|
return .{
|
||||||
.raw = raw_handle,
|
.raw = raw_handle,
|
||||||
.size = bytes,
|
.size = aligned_size, // Expose the aligned size to the rest of the application
|
||||||
.usage = usage,
|
.usage = use,
|
||||||
.gloc = gloc,
|
.gloc = gloc,
|
||||||
};
|
};
|
||||||
}
|
}
|
||||||
|
|
||||||
/// Unregisters from the parent GpuAllocator and cleanly destroys GPU resources
|
/// Unregisters from the parent GpuAllocator and cleanly destroys GPU resources
|
||||||
pub fn deinit(self: @This()) void {
|
pub fn deinit(self: @This()) void {
|
||||||
self.gloc.unregisterAndDestroyBuffer(self);
|
self.gloc.freeBuffer(self.raw, self.size);
|
||||||
|
}
|
||||||
|
|
||||||
|
/// Native getConstMappedRange wrapper
|
||||||
|
pub fn getConstMappedRange(self: @This(), offset: u64, size: u64) ?*const anyopaque {
|
||||||
|
return c.wgpuBufferGetConstMappedRange(self.raw, offset, size);
|
||||||
}
|
}
|
||||||
|
|
||||||
/// Native mapAsync wrapper
|
/// Native mapAsync wrapper
|
||||||
@@ -39,12 +60,79 @@ pub fn mapAsync(
|
|||||||
_ = c.wgpuBufferMapAsync(self.raw, mode, offset, size, callback_info);
|
_ = c.wgpuBufferMapAsync(self.raw, mode, offset, size, callback_info);
|
||||||
}
|
}
|
||||||
|
|
||||||
/// Native getConstMappedRange wrapper
|
|
||||||
pub fn getConstMappedRange(self: @This(), offset: u64, size: u64) ?*const anyopaque {
|
|
||||||
return c.wgpuBufferGetConstMappedRange(self.raw, offset, size);
|
|
||||||
}
|
|
||||||
|
|
||||||
/// Native unmap wrapper
|
/// Native unmap wrapper
|
||||||
pub fn unmap(self: @This()) void {
|
pub fn unmap(self: @This()) void {
|
||||||
c.wgpuBufferUnmap(self.raw);
|
c.wgpuBufferUnmap(self.raw);
|
||||||
}
|
}
|
||||||
|
|
||||||
|
/// CPU to GPU.
|
||||||
|
pub fn load(
|
||||||
|
self: @This(),
|
||||||
|
T: type,
|
||||||
|
data: []const T,
|
||||||
|
) !void {
|
||||||
|
const bytes = data.len * @sizeOf(T);
|
||||||
|
|
||||||
|
if (bytes == self.size) {
|
||||||
|
// Aligned path: direct download
|
||||||
|
c.wgpuQueueWriteBuffer(self.gloc.device.queue, self.raw, 0, data.ptr, self.size);
|
||||||
|
} else {
|
||||||
|
// Unaligned path: Split the write into an aligned chunk and a padded remainder
|
||||||
|
// to support arbitrary lengths without any allocations or large stack arrays.
|
||||||
|
const aligned_part = (bytes / 4) * 4;
|
||||||
|
if (aligned_part > 0) {
|
||||||
|
c.wgpuQueueWriteBuffer(self.gloc.device.queue, self.raw, 0, data.ptr, aligned_part);
|
||||||
|
}
|
||||||
|
|
||||||
|
var remainder_buf: [4]u8 = .{ 0, 0, 0, 0 };
|
||||||
|
const data_bytes = std.mem.sliceAsBytes(data);
|
||||||
|
@memcpy(remainder_buf[0 .. bytes - aligned_part], data_bytes[aligned_part..bytes]);
|
||||||
|
|
||||||
|
c.wgpuQueueWriteBuffer(self.gloc.device.queue, self.raw, aligned_part, &remainder_buf, 4);
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
pub fn read(self: @This(), alloc: std.mem.Allocator, T: type) ![]T {
|
||||||
|
const out = try alloc.alloc(T, @divExact(self.size, @sizeOf(T)));
|
||||||
|
|
||||||
|
const staging = try init(
|
||||||
|
self.gloc,
|
||||||
|
self.size,
|
||||||
|
.initMany(&.{ .MapRead, .CopyDst }),
|
||||||
|
);
|
||||||
|
defer staging.deinit();
|
||||||
|
|
||||||
|
const enc = c.wgpuDeviceCreateCommandEncoder(self.gloc.device.device, null) orelse return error.Encoder;
|
||||||
|
c.wgpuCommandEncoderCopyBufferToBuffer(enc, self.raw, 0, staging.raw, 0, self.size);
|
||||||
|
const cmd = c.wgpuCommandEncoderFinish(enc, null);
|
||||||
|
defer c.wgpuCommandEncoderRelease(enc);
|
||||||
|
defer c.wgpuCommandBufferRelease(cmd);
|
||||||
|
c.wgpuQueueSubmit(self.gloc.device.queue, 1, &cmd);
|
||||||
|
|
||||||
|
var mapped = false;
|
||||||
|
staging.mapAsync(
|
||||||
|
c.WGPUMapMode_Read,
|
||||||
|
0,
|
||||||
|
self.size,
|
||||||
|
.{ .callback = onMapped, .userdata1 = &mapped },
|
||||||
|
);
|
||||||
|
while (!mapped) self.gloc.device.poll();
|
||||||
|
|
||||||
|
const ptr: [*]const T = @ptrCast(@alignCast(
|
||||||
|
staging.getConstMappedRange(0, self.size),
|
||||||
|
));
|
||||||
|
@memcpy(out[0..out.len], ptr[0..out.len]);
|
||||||
|
staging.unmap();
|
||||||
|
|
||||||
|
return out;
|
||||||
|
}
|
||||||
|
|
||||||
|
fn onMapped(
|
||||||
|
status: c.WGPUMapAsyncStatus,
|
||||||
|
_: c.WGPUStringView,
|
||||||
|
userdata1: ?*anyopaque,
|
||||||
|
_: ?*anyopaque,
|
||||||
|
) callconv(.c) void {
|
||||||
|
const flag: *bool = @ptrCast(@alignCast(userdata1.?));
|
||||||
|
flag.* = (status == c.WGPUMapAsyncStatus_Success);
|
||||||
|
}
|
||||||
|
|||||||
+9
-2
@@ -10,6 +10,12 @@ const Ctx = struct {
|
|||||||
const GpuDeviceConfig = struct {
|
const GpuDeviceConfig = struct {
|
||||||
/// VRAM limit. Default 2 GB
|
/// VRAM limit. Default 2 GB
|
||||||
vram_bytes_limit: u64 = 2 * 1024 * 1024 * 1024,
|
vram_bytes_limit: u64 = 2 * 1024 * 1024 * 1024,
|
||||||
|
power_preference: enum(c_uint) {
|
||||||
|
Undefined = 0x00000000,
|
||||||
|
LowPower = 0x00000001,
|
||||||
|
HighPerformance = 0x00000002,
|
||||||
|
Force32 = 0x7FFFFFFF,
|
||||||
|
} = .HighPerformance,
|
||||||
};
|
};
|
||||||
|
|
||||||
instance: c.WGPUInstance,
|
instance: c.WGPUInstance,
|
||||||
@@ -29,7 +35,7 @@ pub fn init(config: GpuDeviceConfig) !@This() {
|
|||||||
var ctx = Ctx{};
|
var ctx = Ctx{};
|
||||||
_ = c.wgpuInstanceRequestAdapter(
|
_ = c.wgpuInstanceRequestAdapter(
|
||||||
instance,
|
instance,
|
||||||
&.{ .powerPreference = c.WGPUPowerPreference_HighPerformance },
|
&.{ .powerPreference = @intFromEnum(config.power_preference) },
|
||||||
.{ .callback = onAdapter, .userdata1 = &ctx },
|
.{ .callback = onAdapter, .userdata1 = &ctx },
|
||||||
);
|
);
|
||||||
c.wgpuInstanceProcessEvents(instance);
|
c.wgpuInstanceProcessEvents(instance);
|
||||||
@@ -87,7 +93,8 @@ pub fn deinit(self: @This()) void {
|
|||||||
c.wgpuInstanceRelease(self.instance);
|
c.wgpuInstanceRelease(self.instance);
|
||||||
}
|
}
|
||||||
|
|
||||||
pub fn poll(self: *@This()) void {
|
/// Wait for thing to be done
|
||||||
|
pub fn poll(self: @This()) void {
|
||||||
_ = c.wgpuDevicePoll(self.device, 1, null);
|
_ = c.wgpuDevicePoll(self.device, 1, null);
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|||||||
@@ -1,25 +0,0 @@
|
|||||||
const std = @import("std");
|
|
||||||
const GpuDevice = @import("GpuDevice.zig");
|
|
||||||
const c = @import("utils.zig").c;
|
|
||||||
const sv = @import("utils.zig").sv;
|
|
||||||
|
|
||||||
raw: c.WGPUComputePipeline,
|
|
||||||
|
|
||||||
pub fn init(device: GpuDevice, wgsl: []const u8) !@This() {
|
|
||||||
var wgsl_src = c.WGPUShaderSourceWGSL{
|
|
||||||
.chain = .{ .sType = c.WGPUSType_ShaderSourceWGSL },
|
|
||||||
.code = sv(wgsl),
|
|
||||||
};
|
|
||||||
const shader = c.wgpuDeviceCreateShaderModule(device.device, &.{
|
|
||||||
.nextInChain = @ptrCast(&wgsl_src),
|
|
||||||
}) orelse return error.Shader;
|
|
||||||
defer c.wgpuShaderModuleRelease(shader);
|
|
||||||
|
|
||||||
return .{ .raw = c.wgpuDeviceCreateComputePipeline(device.device, &.{
|
|
||||||
.compute = .{ .module = shader, .entryPoint = sv("main") },
|
|
||||||
}) orelse return error.Pipeline };
|
|
||||||
}
|
|
||||||
|
|
||||||
pub fn deinit(self: @This()) void {
|
|
||||||
c.wgpuComputePipelineRelease(self.raw);
|
|
||||||
}
|
|
||||||
@@ -0,0 +1,122 @@
|
|||||||
|
/// GpuProcess is just a pipeline with 2 inpout and 1 output
|
||||||
|
/// for now, to see if I make it a bit more generic
|
||||||
|
///
|
||||||
|
const std = @import("std");
|
||||||
|
const c = @import("utils.zig").c;
|
||||||
|
const sv = @import("utils.zig").sv;
|
||||||
|
const GpuAllocator = @import("GpuAllocator.zig");
|
||||||
|
const GpuBuffer = @import("GpuBuffer.zig");
|
||||||
|
const GpuDevice = @import("GpuDevice.zig");
|
||||||
|
|
||||||
|
pip: c.WGPUComputePipeline,
|
||||||
|
|
||||||
|
pub fn init(device: GpuDevice, wgsl: []const u8) !@This() {
|
||||||
|
var wgsl_src = c.WGPUShaderSourceWGSL{
|
||||||
|
.chain = .{ .sType = c.WGPUSType_ShaderSourceWGSL },
|
||||||
|
.code = sv(wgsl),
|
||||||
|
};
|
||||||
|
const shader = c.wgpuDeviceCreateShaderModule(device.device, &.{
|
||||||
|
.nextInChain = @ptrCast(&wgsl_src),
|
||||||
|
}) orelse return error.Shader;
|
||||||
|
defer c.wgpuShaderModuleRelease(shader);
|
||||||
|
|
||||||
|
return .{ .pip = c.wgpuDeviceCreateComputePipeline(device.device, &.{
|
||||||
|
.compute = .{ .module = shader, .entryPoint = sv("main") },
|
||||||
|
}) orelse return error.Pipeline };
|
||||||
|
}
|
||||||
|
|
||||||
|
pub fn deinit(self: @This()) void {
|
||||||
|
c.wgpuComputePipelineRelease(self.pip);
|
||||||
|
}
|
||||||
|
|
||||||
|
fn onMapped(
|
||||||
|
status: c.WGPUMapAsyncStatus,
|
||||||
|
_: c.WGPUStringView,
|
||||||
|
userdata1: ?*anyopaque,
|
||||||
|
_: ?*anyopaque,
|
||||||
|
) callconv(.c) void {
|
||||||
|
const flag: *bool = @ptrCast(@alignCast(userdata1.?));
|
||||||
|
flag.* = (status == c.WGPUMapAsyncStatus_Success);
|
||||||
|
}
|
||||||
|
|
||||||
|
// Changed: gloc is passed by value instead of *GpuAllocator
|
||||||
|
pub fn run(
|
||||||
|
self: @This(),
|
||||||
|
gloc: GpuAllocator,
|
||||||
|
T: type,
|
||||||
|
buf_a: GpuBuffer,
|
||||||
|
buf_b: GpuBuffer,
|
||||||
|
buf_out: GpuBuffer,
|
||||||
|
) !void {
|
||||||
|
const max_chunk_bytes: u64 = 1024 * 1024 * 1024; // 1 GB
|
||||||
|
|
||||||
|
const bytes = buf_a.size;
|
||||||
|
var offset: u64 = 0;
|
||||||
|
while (offset < bytes) {
|
||||||
|
const current_chunk_bytes = @min(max_chunk_bytes, bytes - offset);
|
||||||
|
const current_chunk_elements: u32 = @intCast(current_chunk_bytes / @sizeOf(T));
|
||||||
|
|
||||||
|
const info_buf = try GpuBuffer.init(
|
||||||
|
gloc,
|
||||||
|
@sizeOf(u32),
|
||||||
|
.initMany(&.{ .Uniform, .CopyDst }),
|
||||||
|
);
|
||||||
|
defer info_buf.deinit();
|
||||||
|
|
||||||
|
c.wgpuQueueWriteBuffer(gloc.device.queue, info_buf.raw, 0, ¤t_chunk_elements, @sizeOf(u32));
|
||||||
|
|
||||||
|
const entries = [_]c.WGPUBindGroupEntry{
|
||||||
|
.{ .binding = 0, .buffer = buf_a.raw, .offset = offset, .size = current_chunk_bytes },
|
||||||
|
.{ .binding = 1, .buffer = buf_b.raw, .offset = offset, .size = current_chunk_bytes },
|
||||||
|
.{ .binding = 2, .buffer = buf_out.raw, .offset = offset, .size = current_chunk_bytes },
|
||||||
|
.{ .binding = 3, .buffer = info_buf.raw, .offset = 0, .size = @sizeOf(u32) },
|
||||||
|
};
|
||||||
|
|
||||||
|
try submitPass(gloc, self.pip, &entries, current_chunk_elements);
|
||||||
|
|
||||||
|
offset += current_chunk_bytes;
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
// Changed: gloc is passed by value instead of *GpuAllocator
|
||||||
|
fn submitPass(
|
||||||
|
gloc: GpuAllocator,
|
||||||
|
pipeline: c.WGPUComputePipeline,
|
||||||
|
entries: []const c.WGPUBindGroupEntry,
|
||||||
|
n: usize,
|
||||||
|
) !void {
|
||||||
|
const bgl = c.wgpuComputePipelineGetBindGroupLayout(pipeline, 0);
|
||||||
|
defer c.wgpuBindGroupLayoutRelease(bgl);
|
||||||
|
|
||||||
|
const bg = c.wgpuDeviceCreateBindGroup(gloc.device.device, &.{
|
||||||
|
.layout = bgl,
|
||||||
|
.entries = entries.ptr,
|
||||||
|
.entryCount = entries.len,
|
||||||
|
}) orelse return error.BindGroup;
|
||||||
|
defer c.wgpuBindGroupRelease(bg);
|
||||||
|
|
||||||
|
const enc = c.wgpuDeviceCreateCommandEncoder(gloc.device.device, null) orelse
|
||||||
|
return error.Encoder;
|
||||||
|
const pass = c.wgpuCommandEncoderBeginComputePass(enc, null);
|
||||||
|
c.wgpuComputePassEncoderSetPipeline(pass, pipeline);
|
||||||
|
c.wgpuComputePassEncoderSetBindGroup(pass, 0, bg, 0, null);
|
||||||
|
|
||||||
|
const WORKGROUP_SIZE = 256;
|
||||||
|
const MAX_WORKGROUPS = 65535;
|
||||||
|
|
||||||
|
const desired_workgroups = ceilDiv(n, WORKGROUP_SIZE);
|
||||||
|
const dispatch_count = @min(desired_workgroups, MAX_WORKGROUPS);
|
||||||
|
|
||||||
|
c.wgpuComputePassEncoderDispatchWorkgroups(pass, @intCast(dispatch_count), 1, 1);
|
||||||
|
c.wgpuComputePassEncoderEnd(pass);
|
||||||
|
c.wgpuComputePassEncoderRelease(pass);
|
||||||
|
|
||||||
|
const cmd = c.wgpuCommandEncoderFinish(enc, null);
|
||||||
|
defer c.wgpuCommandEncoderRelease(enc);
|
||||||
|
defer c.wgpuCommandBufferRelease(cmd);
|
||||||
|
c.wgpuQueueSubmit(gloc.device.queue, 1, &cmd);
|
||||||
|
}
|
||||||
|
|
||||||
|
fn ceilDiv(n: usize, d: usize) usize {
|
||||||
|
return (n + d - 1) / d;
|
||||||
|
}
|
||||||
-195
@@ -1,195 +0,0 @@
|
|||||||
/// Dummy
|
|
||||||
const std = @import("std");
|
|
||||||
const c = @import("utils.zig").c;
|
|
||||||
const GpuAllocator = @import("GpuAllocator.zig");
|
|
||||||
const GpuBuffer = @import("GpuBuffer.zig");
|
|
||||||
const GpuDevice = @import("GpuDevice.zig");
|
|
||||||
const GpuPipeline = @import("GpuPipeline.zig");
|
|
||||||
|
|
||||||
const Vec = @This();
|
|
||||||
|
|
||||||
buf: GpuBuffer,
|
|
||||||
len: usize,
|
|
||||||
|
|
||||||
pub fn initZero(gloc: *GpuAllocator, len: usize) !Vec {
|
|
||||||
return .{
|
|
||||||
.buf = try GpuBuffer.init(
|
|
||||||
gloc,
|
|
||||||
f16,
|
|
||||||
len,
|
|
||||||
c.WGPUBufferUsage_Storage | c.WGPUBufferUsage_CopyDst | c.WGPUBufferUsage_CopySrc,
|
|
||||||
),
|
|
||||||
.len = len,
|
|
||||||
};
|
|
||||||
}
|
|
||||||
|
|
||||||
pub fn initLoad(gloc: *GpuAllocator, data: []const f16) !Vec {
|
|
||||||
var self = try initZero(gloc, data.len);
|
|
||||||
try self.load(gloc.device, data);
|
|
||||||
return self;
|
|
||||||
}
|
|
||||||
|
|
||||||
pub fn deinit(self: Vec) void {
|
|
||||||
self.buf.deinit();
|
|
||||||
}
|
|
||||||
|
|
||||||
/// CPU to GPU.
|
|
||||||
pub fn load(
|
|
||||||
self: Vec,
|
|
||||||
device: GpuDevice,
|
|
||||||
data: []const f16,
|
|
||||||
) !void {
|
|
||||||
std.debug.assert(data.len == self.len);
|
|
||||||
const bytes = self.byteSize();
|
|
||||||
c.wgpuQueueWriteBuffer(device.queue, self.buf.raw, 0, data.ptr, bytes);
|
|
||||||
}
|
|
||||||
|
|
||||||
pub fn byteSize(self: Vec) u64 {
|
|
||||||
return @as(u64, self.len) * @sizeOf(f16);
|
|
||||||
}
|
|
||||||
|
|
||||||
pub fn run(self: Vec, gloc: *GpuAllocator, other: Vec, pip: GpuPipeline) !Vec {
|
|
||||||
std.debug.assert(self.len == other.len);
|
|
||||||
|
|
||||||
const result = try Vec.initZero(gloc, self.len);
|
|
||||||
errdefer result.deinit();
|
|
||||||
|
|
||||||
try dispatch2in1out(gloc, pip.raw, self.buf, other.buf, result.buf, self.byteSize());
|
|
||||||
|
|
||||||
return result;
|
|
||||||
}
|
|
||||||
|
|
||||||
/// GPU to CPU.
|
|
||||||
pub fn read(self: Vec, gloc: *GpuAllocator, alloc: std.mem.Allocator) ![]f16 {
|
|
||||||
const out = try alloc.alloc(f16, self.len);
|
|
||||||
const bytes = self.byteSize();
|
|
||||||
|
|
||||||
const staging = try GpuBuffer.init(
|
|
||||||
gloc,
|
|
||||||
f16,
|
|
||||||
self.len,
|
|
||||||
c.WGPUBufferUsage_MapRead | c.WGPUBufferUsage_CopyDst,
|
|
||||||
);
|
|
||||||
defer staging.deinit();
|
|
||||||
|
|
||||||
const enc = c.wgpuDeviceCreateCommandEncoder(gloc.device.device, null) orelse return error.Encoder;
|
|
||||||
c.wgpuCommandEncoderCopyBufferToBuffer(enc, self.buf.raw, 0, staging.raw, 0, bytes);
|
|
||||||
const cmd = c.wgpuCommandEncoderFinish(enc, null);
|
|
||||||
defer c.wgpuCommandEncoderRelease(enc);
|
|
||||||
defer c.wgpuCommandBufferRelease(cmd);
|
|
||||||
c.wgpuQueueSubmit(gloc.device.queue, 1, &cmd);
|
|
||||||
|
|
||||||
var mapped = false;
|
|
||||||
staging.mapAsync(
|
|
||||||
c.WGPUMapMode_Read,
|
|
||||||
0,
|
|
||||||
bytes,
|
|
||||||
.{ .callback = onMapped, .userdata1 = &mapped },
|
|
||||||
);
|
|
||||||
while (!mapped) gloc.device.poll();
|
|
||||||
|
|
||||||
const ptr: [*]const f16 = @ptrCast(@alignCast(
|
|
||||||
staging.getConstMappedRange(0, bytes),
|
|
||||||
));
|
|
||||||
@memcpy(out[0..self.len], ptr[0..self.len]);
|
|
||||||
staging.unmap();
|
|
||||||
|
|
||||||
return out;
|
|
||||||
}
|
|
||||||
|
|
||||||
fn onMapped(
|
|
||||||
status: c.WGPUMapAsyncStatus,
|
|
||||||
_: c.WGPUStringView,
|
|
||||||
userdata1: ?*anyopaque,
|
|
||||||
_: ?*anyopaque,
|
|
||||||
) callconv(.c) void {
|
|
||||||
const flag: *bool = @ptrCast(@alignCast(userdata1.?));
|
|
||||||
flag.* = (status == c.WGPUMapAsyncStatus_Success);
|
|
||||||
}
|
|
||||||
|
|
||||||
/// Encode + submit a 2-input, 1-output compute pass (used by add).
|
|
||||||
fn dispatch2in1out(
|
|
||||||
gloc: *GpuAllocator,
|
|
||||||
pipeline: c.WGPUComputePipeline,
|
|
||||||
buf_a: GpuBuffer,
|
|
||||||
buf_b: GpuBuffer,
|
|
||||||
buf_out: GpuBuffer,
|
|
||||||
bytes: u64,
|
|
||||||
) !void {
|
|
||||||
const max_chunk_bytes: u64 = 1024 * 1024 * 1024; // 1 GB
|
|
||||||
|
|
||||||
var offset: u64 = 0;
|
|
||||||
while (offset < bytes) {
|
|
||||||
// Calculate bounds for the current chunk
|
|
||||||
const current_chunk_bytes = @min(max_chunk_bytes, bytes - offset);
|
|
||||||
const current_chunk_elements: u32 = @intCast(current_chunk_bytes / @sizeOf(f16));
|
|
||||||
|
|
||||||
// Create uniform buffer for this specific chunk's size
|
|
||||||
const info_buf = try GpuBuffer.init(
|
|
||||||
gloc,
|
|
||||||
u32,
|
|
||||||
1,
|
|
||||||
c.WGPUBufferUsage_Uniform | c.WGPUBufferUsage_CopyDst,
|
|
||||||
);
|
|
||||||
defer info_buf.deinit();
|
|
||||||
|
|
||||||
// Write the number of elements *in this chunk* to the uniform buffer
|
|
||||||
c.wgpuQueueWriteBuffer(gloc.device.queue, info_buf.raw, 0, ¤t_chunk_elements, @sizeOf(u32));
|
|
||||||
|
|
||||||
// Bind only the sub-slice for this chunk using `.offset` and `.size`
|
|
||||||
const entries = [_]c.WGPUBindGroupEntry{
|
|
||||||
.{ .binding = 0, .buffer = buf_a.raw, .offset = offset, .size = current_chunk_bytes },
|
|
||||||
.{ .binding = 1, .buffer = buf_b.raw, .offset = offset, .size = current_chunk_bytes },
|
|
||||||
.{ .binding = 2, .buffer = buf_out.raw, .offset = offset, .size = current_chunk_bytes },
|
|
||||||
.{ .binding = 3, .buffer = info_buf.raw, .offset = 0, .size = @sizeOf(u32) },
|
|
||||||
};
|
|
||||||
|
|
||||||
// Submit the pass for this specific chunk
|
|
||||||
try submitPass(gloc, pipeline, &entries, current_chunk_elements);
|
|
||||||
|
|
||||||
offset += current_chunk_bytes;
|
|
||||||
}
|
|
||||||
}
|
|
||||||
|
|
||||||
/// Create bind group, encode pass, submit.
|
|
||||||
fn submitPass(
|
|
||||||
gloc: *GpuAllocator,
|
|
||||||
pipeline: c.WGPUComputePipeline,
|
|
||||||
entries: []const c.WGPUBindGroupEntry,
|
|
||||||
n: usize,
|
|
||||||
) !void {
|
|
||||||
const bgl = c.wgpuComputePipelineGetBindGroupLayout(pipeline, 0);
|
|
||||||
defer c.wgpuBindGroupLayoutRelease(bgl);
|
|
||||||
|
|
||||||
const bg = c.wgpuDeviceCreateBindGroup(gloc.device.device, &.{
|
|
||||||
.layout = bgl,
|
|
||||||
.entries = entries.ptr,
|
|
||||||
.entryCount = entries.len,
|
|
||||||
}) orelse return error.BindGroup;
|
|
||||||
defer c.wgpuBindGroupRelease(bg);
|
|
||||||
|
|
||||||
const enc = c.wgpuDeviceCreateCommandEncoder(gloc.device.device, null) orelse
|
|
||||||
return error.Encoder;
|
|
||||||
const pass = c.wgpuCommandEncoderBeginComputePass(enc, null);
|
|
||||||
c.wgpuComputePassEncoderSetPipeline(pass, pipeline);
|
|
||||||
c.wgpuComputePassEncoderSetBindGroup(pass, 0, bg, 0, null);
|
|
||||||
|
|
||||||
const WORKGROUP_SIZE = 256;
|
|
||||||
const MAX_WORKGROUPS = 65535;
|
|
||||||
|
|
||||||
const desired_workgroups = ceilDiv(n, WORKGROUP_SIZE);
|
|
||||||
const dispatch_count = @min(desired_workgroups, MAX_WORKGROUPS);
|
|
||||||
|
|
||||||
c.wgpuComputePassEncoderDispatchWorkgroups(pass, @intCast(dispatch_count), 1, 1);
|
|
||||||
c.wgpuComputePassEncoderEnd(pass);
|
|
||||||
c.wgpuComputePassEncoderRelease(pass);
|
|
||||||
|
|
||||||
const cmd = c.wgpuCommandEncoderFinish(enc, null);
|
|
||||||
defer c.wgpuCommandEncoderRelease(enc);
|
|
||||||
defer c.wgpuCommandBufferRelease(cmd);
|
|
||||||
c.wgpuQueueSubmit(gloc.device.queue, 1, &cmd);
|
|
||||||
}
|
|
||||||
|
|
||||||
fn ceilDiv(n: usize, d: usize) usize {
|
|
||||||
return (n + d - 1) / d;
|
|
||||||
}
|
|
||||||
@@ -1,43 +0,0 @@
|
|||||||
const std = @import("std");
|
|
||||||
const GpuDevice = @import("GpuDevice.zig");
|
|
||||||
const GpuAllocator = @import("GpuAllocator.zig");
|
|
||||||
const GpuPipeline = @import("GpuPipeline.zig");
|
|
||||||
const Vec = @import("Vec.zig");
|
|
||||||
|
|
||||||
const c = @import("utils.zig").c;
|
|
||||||
|
|
||||||
pub fn main(init: std.process.Init) !void {
|
|
||||||
const allocator = init.gpa;
|
|
||||||
|
|
||||||
const device = try GpuDevice.init(.{ .vram_bytes_limit = 4 * 1024 * 1024 * 1024 });
|
|
||||||
defer device.deinit();
|
|
||||||
|
|
||||||
var gloc = try GpuAllocator.init(allocator, device);
|
|
||||||
defer gloc.deinit();
|
|
||||||
|
|
||||||
const add_pip = try GpuPipeline.init(device, @embedFile("shaders/add.wgsl"));
|
|
||||||
defer add_pip.deinit();
|
|
||||||
|
|
||||||
const data_a = try allocator.alloc(f16, 1024);
|
|
||||||
defer allocator.free(data_a);
|
|
||||||
const data_b = try allocator.alloc(f16, 1024);
|
|
||||||
defer allocator.free(data_b);
|
|
||||||
|
|
||||||
for (0..1024) |i| {
|
|
||||||
data_a[i] = @floatFromInt(i);
|
|
||||||
data_b[i] = @floatFromInt(1024 - 1 - i);
|
|
||||||
}
|
|
||||||
|
|
||||||
const a = try Vec.initLoad(&gloc, data_a);
|
|
||||||
defer a.deinit();
|
|
||||||
const b = try Vec.initLoad(&gloc, data_b);
|
|
||||||
defer b.deinit();
|
|
||||||
|
|
||||||
const sum = try a.run(&gloc, b, add_pip);
|
|
||||||
defer sum.deinit();
|
|
||||||
|
|
||||||
const out = try sum.read(&gloc, allocator);
|
|
||||||
defer allocator.free(out);
|
|
||||||
|
|
||||||
std.debug.print("{any}\n", .{out});
|
|
||||||
}
|
|
||||||
+2
-1
@@ -1,4 +1,5 @@
|
|||||||
pub const GpuAllocator = @import("GpuAllocator.zig");
|
pub const GpuAllocator = @import("GpuAllocator.zig");
|
||||||
|
pub const GpuArena = @import("GpuArena.zig");
|
||||||
pub const GpuBuffer = @import("GpuBuffer.zig");
|
pub const GpuBuffer = @import("GpuBuffer.zig");
|
||||||
pub const GpuDevice = @import("GpuDevice.zig");
|
pub const GpuDevice = @import("GpuDevice.zig");
|
||||||
pub const GpuPipeline = @import("GpuPipeline.zig");
|
pub const GpuProcess = @import("GpuProcess.zig");
|
||||||
|
|||||||
Reference in New Issue
Block a user