mirror of
https://codeberg.org/ziglang/zig.git
synced 2026-06-21 16:11:50 -04:00
Ali Cheraghi
c9ca79fb48
Exposes OpControlBarrier and OpMemoryBarrier so compute kernels do not need to hand-roll inline asm for workgroup synchronisation. Scope and MemorySemantics mirror the SPIRV-Headers unified1 grammar bit-for-bit; workgroupBarrier() matches the semantics of GLSL barrier(). Co-authored-by: Quint Daenen <quint@daenen.email>
85 lines
2.8 KiB
Zig
85 lines
2.8 KiB
Zig
const std = @import("std.zig");
|
|
|
|
pub const position_in = @extern(*addrspace(.input) @Vector(4, f32), .{ .name = "position" });
|
|
pub const position_out = @extern(*addrspace(.output) @Vector(4, f32), .{ .name = "position" });
|
|
pub const point_size_in = @extern(*addrspace(.input) f32, .{ .name = "point_size" });
|
|
pub const point_size_out = @extern(*addrspace(.output) f32, .{ .name = "point_size" });
|
|
pub extern const invocation_id: u32 addrspace(.input);
|
|
pub extern const frag_coord: @Vector(4, f32) addrspace(.input);
|
|
pub extern const point_coord: @Vector(2, f32) addrspace(.input);
|
|
// TODO: direct/indirect values
|
|
// pub extern const front_facing: bool addrspace(.input);
|
|
// TODO: runtime array
|
|
// pub extern const sample_mask;
|
|
pub extern var frag_depth: f32 addrspace(.output);
|
|
pub extern const num_workgroups: @Vector(3, u32) addrspace(.input);
|
|
pub extern const workgroup_size: @Vector(3, u32) addrspace(.input);
|
|
pub extern const workgroup_id: @Vector(3, u32) addrspace(.input);
|
|
pub extern const local_invocation_id: @Vector(3, u32) addrspace(.input);
|
|
pub extern const global_invocation_id: @Vector(3, u32) addrspace(.input);
|
|
pub extern const vertex_index: u32 addrspace(.input);
|
|
pub extern const instance_index: u32 addrspace(.input);
|
|
|
|
pub const Scope = enum(u32) {
|
|
cross_device = 0,
|
|
device = 1,
|
|
workgroup = 2,
|
|
subgroup = 3,
|
|
invocation = 4,
|
|
queue_family = 5,
|
|
shader_call_khr = 6,
|
|
};
|
|
|
|
pub const MemorySemantics = packed struct(u32) {
|
|
_reserved_bit_0: bool = false,
|
|
acquire: bool = false,
|
|
release: bool = false,
|
|
acquire_release: bool = false,
|
|
sequentially_consistent: bool = false,
|
|
_reserved_bit_5: bool = false,
|
|
uniform_memory: bool = false,
|
|
subgroup_memory: bool = false,
|
|
workgroup_memory: bool = false,
|
|
cross_workgroup_memory: bool = false,
|
|
atomic_counter_memory: bool = false,
|
|
image_memory: bool = false,
|
|
output_memory: bool = false,
|
|
make_available: bool = false,
|
|
make_visible: bool = false,
|
|
@"volatile": bool = false,
|
|
_reserved: u16 = 0,
|
|
|
|
pub const none: MemorySemantics = .{};
|
|
};
|
|
|
|
pub fn controlBarrier(
|
|
comptime execution: Scope,
|
|
comptime memory: Scope,
|
|
comptime semantics: MemorySemantics,
|
|
) void {
|
|
asm volatile (
|
|
\\OpControlBarrier %exec %mem %sem
|
|
:
|
|
: [exec] "" (@as(u32, @intFromEnum(execution))),
|
|
[mem] "" (@as(u32, @intFromEnum(memory))),
|
|
[sem] "" (@as(u32, @bitCast(semantics))),
|
|
);
|
|
}
|
|
|
|
pub fn memoryBarrier(comptime memory: Scope, comptime semantics: MemorySemantics) void {
|
|
asm volatile (
|
|
\\OpMemoryBarrier %mem %sem
|
|
:
|
|
: [mem] "" (@as(u32, @intFromEnum(memory))),
|
|
[sem] "" (@as(u32, @bitCast(semantics))),
|
|
);
|
|
}
|
|
|
|
pub fn workgroupBarrier() void {
|
|
controlBarrier(
|
|
.workgroup,
|
|
.workgroup,
|
|
.{ .acquire_release = true, .workgroup_memory = true },
|
|
);
|
|
}
|