1const std = @import("std.zig");
  2
  3pub const position_in = @extern(*addrspace(.input) @Vector(4, f32), .{ .name = "position" });
  4pub const position_out = @extern(*addrspace(.output) @Vector(4, f32), .{ .name = "position" });
  5pub const point_size_in = @extern(*addrspace(.input) f32, .{ .name = "point_size" });
  6pub const point_size_out = @extern(*addrspace(.output) f32, .{ .name = "point_size" });
  7pub extern const invocation_id: u32 addrspace(.input);
  8pub extern const frag_coord: @Vector(4, f32) addrspace(.input);
  9pub extern const point_coord: @Vector(2, f32) addrspace(.input);
 10// TODO: direct/indirect values
 11// pub extern const front_facing: bool addrspace(.input);
 12// TODO: runtime array
 13// pub extern const sample_mask;
 14pub extern var frag_depth: f32 addrspace(.output);
 15pub extern const num_workgroups: @Vector(3, u32) addrspace(.input);
 16pub extern const workgroup_size: @Vector(3, u32) addrspace(.input);
 17pub extern const workgroup_id: @Vector(3, u32) addrspace(.input);
 18pub extern const local_invocation_id: @Vector(3, u32) addrspace(.input);
 19pub extern const global_invocation_id: @Vector(3, u32) addrspace(.input);
 20pub extern const vertex_index: u32 addrspace(.input);
 21pub extern const instance_index: u32 addrspace(.input);
 22
 23/// Forms the main linkage for `input` and `output` address spaces.
 24/// `ptr` must be a reference to variable or struct field.
 25pub fn location(comptime ptr: anytype, comptime loc: u32) void {
 26    asm volatile (
 27        \\OpDecorate %ptr Location $loc
 28        :
 29        : [ptr] "" (ptr),
 30          [loc] "c" (loc),
 31    );
 32}
 33
 34/// Forms the main linkage for `input` and `output` address spaces.
 35/// `ptr` must be a reference to variable or struct field.
 36pub fn binding(comptime ptr: anytype, comptime set: u32, comptime bind: u32) void {
 37    asm volatile (
 38        \\OpDecorate %ptr DescriptorSet $set
 39        \\OpDecorate %ptr Binding $bind
 40        :
 41        : [ptr] "" (ptr),
 42          [set] "c" (set),
 43          [bind] "c" (bind),
 44    );
 45}
 46
 47pub const ExecutionMode = union(Tag) {
 48    /// Sets origin of the framebuffer to the upper-left corner
 49    origin_upper_left,
 50    /// Sets origin of the framebuffer to the lower-left corner
 51    origin_lower_left,
 52    /// Indicates that the fragment shader writes to `frag_depth`,
 53    /// replacing the fixed-function depth value.
 54    depth_replacing,
 55    /// Indicates that per-fragment tests may assume that
 56    /// any `frag_depth` built in-decorated value written by the shader is
 57    /// greater-than-or-equal to the fragment’s interpolated depth value
 58    depth_greater,
 59    /// Indicates that per-fragment tests may assume that
 60    /// any `frag_depth` built in-decorated value written by the shader is
 61    /// less-than-or-equal to the fragment’s interpolated depth value
 62    depth_less,
 63    /// Indicates that per-fragment tests may assume that
 64    /// any `frag_depth` built in-decorated value written by the shader is
 65    /// the same as the fragment’s interpolated depth value
 66    depth_unchanged,
 67    /// Indicates the workgroup size in the x, y, and z dimensions.
 68    local_size: LocalSize,
 69
 70    pub const Tag = enum(u32) {
 71        origin_upper_left = 7,
 72        origin_lower_left = 8,
 73        depth_replacing = 12,
 74        depth_greater = 14,
 75        depth_less = 15,
 76        depth_unchanged = 16,
 77        local_size = 17,
 78    };
 79
 80    pub const LocalSize = struct { x: u32, y: u32, z: u32 };
 81};
 82
 83/// Declare the mode entry point executes in.
 84pub fn executionMode(comptime entry_point: anytype, comptime mode: ExecutionMode) void {
 85    const cc = @typeInfo(@TypeOf(entry_point)).@"fn".calling_convention;
 86    switch (mode) {
 87        .origin_upper_left,
 88        .origin_lower_left,
 89        .depth_replacing,
 90        .depth_greater,
 91        .depth_less,
 92        .depth_unchanged,
 93        => {
 94            if (cc != .spirv_fragment) {
 95                @compileError(
 96                    \\invalid execution mode '
 97                ++ @tagName(mode) ++
 98                    \\' for function with '
 99                ++ @tagName(cc) ++
100                    \\' calling convention
101                );
102            }
103            asm volatile (
104                \\OpExecutionMode %entry_point $mode
105                :
106                : [entry_point] "" (entry_point),
107                  [mode] "c" (@intFromEnum(mode)),
108            );
109        },
110        .local_size => |size| {
111            if (cc != .spirv_kernel) {
112                @compileError(
113                    \\invalid execution mode 'local_size' for function with '
114                ++ @tagName(cc) ++
115                    \\' calling convention
116                );
117            }
118            asm volatile (
119                \\OpExecutionMode %entry_point LocalSize $x $y $z
120                :
121                : [entry_point] "" (entry_point),
122                  [x] "c" (size.x),
123                  [y] "c" (size.y),
124                  [z] "c" (size.z),
125            );
126        },
127    }
128}