1const builtin = @import("builtin");
2const std = @import("std.zig");
3
4pub const position_in = @extern(*addrspace(.input) @Vector(4, f32), .{ .name = "position" });
5pub const position_out = @extern(*addrspace(.output) @Vector(4, f32), .{ .name = "position" });
6pub const point_size_in = @extern(*addrspace(.input) f32, .{ .name = "point_size" });
7pub const point_size_out = @extern(*addrspace(.output) f32, .{ .name = "point_size" });
8pub extern const invocation_id: u32 addrspace(.input);
9pub extern const frag_coord: @Vector(4, f32) addrspace(.input);
10pub extern const point_coord: @Vector(2, f32) addrspace(.input);
11// TODO: direct/indirect values
12// pub extern const front_facing: bool addrspace(.input);
13// TODO: runtime array
14// pub extern const sample_mask;
15pub extern var frag_depth: f32 addrspace(.output);
16pub extern const num_workgroups: @Vector(3, u32) addrspace(.input);
17pub extern const workgroup_size: @Vector(3, u32) addrspace(.input);
18pub extern const workgroup_id: @Vector(3, u32) addrspace(.input);
19pub extern const local_invocation_id: @Vector(3, u32) addrspace(.input);
20pub extern const global_invocation_id: @Vector(3, u32) addrspace(.input);
21pub extern const vertex_index: u32 addrspace(.input);
22pub extern const instance_index: u32 addrspace(.input);
23
24pub const Scope = enum(u32) {
25 cross_device = 0,
26 device = 1,
27 workgroup = 2,
28 subgroup = 3,
29 invocation = 4,
30 queue_family = 5,
31 shader_call_khr = 6,
32};
33
34pub const MemorySemantics = packed struct(u32) {
35 _reserved_bit_0: bool = false,
36 acquire: bool = false,
37 release: bool = false,
38 acquire_release: bool = false,
39 sequentially_consistent: bool = false,
40 _reserved_bit_5: bool = false,
41 uniform_memory: bool = false,
42 subgroup_memory: bool = false,
43 workgroup_memory: bool = false,
44 cross_workgroup_memory: bool = false,
45 atomic_counter_memory: bool = false,
46 image_memory: bool = false,
47 output_memory: bool = false,
48 make_available: bool = false,
49 make_visible: bool = false,
50 @"volatile": bool = false,
51 _reserved: u16 = 0,
52
53 pub const none: MemorySemantics = .{};
54};
55
56pub fn controlBarrier(
57 comptime execution: Scope,
58 comptime memory: Scope,
59 comptime semantics: MemorySemantics,
60) void {
61 asm volatile (
62 \\OpControlBarrier %exec %mem %sem
63 :
64 : [exec] "" (@as(u32, @backingInt(execution))),
65 [mem] "" (@as(u32, @backingInt(memory))),
66 [sem] "" (@as(u32, @bitCast(semantics))),
67 );
68}
69
70pub fn memoryBarrier(comptime memory: Scope, comptime semantics: MemorySemantics) void {
71 asm volatile (
72 \\OpMemoryBarrier %mem %sem
73 :
74 : [mem] "" (@as(u32, @backingInt(memory))),
75 [sem] "" (@as(u32, @bitCast(semantics))),
76 );
77}
78
79pub fn workgroupBarrier() void {
80 controlBarrier(
81 .workgroup,
82 .workgroup,
83 .{ .acquire_release = true, .workgroup_memory = true },
84 );
85}
86
87pub fn specConst(T: type, comptime default_value: T, comptime spec_id: u32) T {
88 switch (@typeInfo(T)) {
89 .bool => {
90 const op = if (default_value) "OpSpecConstantTrue" else "OpSpecConstantFalse";
91 return asm ("%ret = " ++ op ++ " %ty\n" ++
92 "OpDecorate %ret SpecId $spec_id"
93 : [ret] "" (-> T),
94 : [ty] "t" (T),
95 [spec_id] "c" (spec_id),
96 );
97 },
98 .int, .float => return asm (
99 \\%ret = OpSpecConstant %ty $default_value
100 \\ OpDecorate %ret SpecId $spec_id
101 : [ret] "" (-> T),
102 : [ty] "t" (T),
103 [default_value] "c" (default_value),
104 [spec_id] "c" (spec_id),
105 ),
106 .vector => return asm (
107 \\%ret = OpSpecConstantComposite %ty %default_value %spec_id
108 : [ret] "" (-> T),
109 : [ty] "t" (T),
110 [default_value] "c" (default_value),
111 [spec_id] "c" (spec_id),
112 ),
113 else => @compileError("Invalid spec-constant type"),
114 }
115}
116
117/// Get the type that specifies a coordinate for a SPIR-V image or sampled image.
118fn ImageCoordinate(Image: type, Element: type) type {
119 const image_info = switch (@typeInfo(Image)) {
120 .spirv => |spirv| switch (spirv) {
121 .sampled_image => |sampled_image| @typeInfo(sampled_image).spirv.image,
122 .image => |image| image,
123 else => @compileError("Expected SPIR-V image or sampled image type, found '" ++ @typeName(Image) ++ "'"),
124 },
125 else => @compileError("Expected SPIR-V image or sampled image type, found '" ++ @typeName(Image) ++ "'"),
126 };
127 const dim = switch (image_info.dim) {
128 .@"1d" => 1 + @as(u8, @intFromBool(image_info.arrayed)),
129 .@"2d" => 2 + @as(u8, @intFromBool(image_info.arrayed)),
130 .@"3d", .cube => 3 + @as(u8, @intFromBool(image_info.arrayed)),
131 };
132 if (dim == 1) return Element else return @Vector(dim, Element);
133}
134
135/// The type of the components that result from sampling or reading from the given SPIR-V image or sampled image type.
136fn ImageSampledType(Image: type) type {
137 const image_info = switch (@typeInfo(Image)) {
138 .spirv => |spirv| switch (spirv) {
139 .sampled_image => |sampled_image| @typeInfo(sampled_image).spirv.image,
140 .image => |image| image,
141 else => @compileError("Expected SPIR-V image or sampled image type, found '" ++ @typeName(Image) ++ "'"),
142 },
143 else => @compileError("Expected SPIR-V image or sampled image type, found '" ++ @typeName(Image) ++ "'"),
144 };
145 return switch (image_info.usage) {
146 inline else => |usage| usage,
147 };
148}
149
150/// The type of `sampled_image` must be a pointer to a SPIR-V sampled image.
151pub fn imageSampleImplicitLod(
152 sampled_image: anytype,
153 coordinate: ImageCoordinate(std.meta.Child(@TypeOf(sampled_image)), f32),
154) @Vector(4, ImageSampledType(std.meta.Child(@TypeOf(sampled_image)))) {
155 const SampledImage = switch (@typeInfo(@TypeOf(sampled_image))) {
156 .pointer => |pointer| pointer.child,
157 else => @compileError("Expected a pointer to SPIR-V sampled image type, found '" ++ @typeName(@TypeOf(sampled_image)) ++ "'"),
158 };
159 const Result = @Vector(4, ImageSampledType(SampledImage));
160
161 const image_info = switch (@typeInfo(SampledImage)) {
162 .spirv => |spirv| switch (spirv) {
163 .sampled_image => |sampled_image_info| @typeInfo(sampled_image_info).spirv.image,
164 else => @compileError("Expected SPIR-V sampled image type, found '" ++ @typeName(SampledImage) ++ "'"),
165 },
166 else => @compileError("Expected SPIR-V sampled image type, found '" ++ @typeName(SampledImage) ++ "'"),
167 };
168
169 if (image_info.multisampled)
170 @compileError("Can not implicitly sample a sampled image that was multisampled");
171
172 // TOOD: If buffer dim is added, throw a compile error if the dimension is a buffer.
173
174 return asm volatile (
175 \\%loaded_sampler = OpLoad %SampledImage %sampled_image
176 \\%ret = OpImageSampleImplicitLod %Result %loaded_sampler %coordinate
177 : [ret] "" (-> Result),
178 : [SampledImage] "t" (SampledImage),
179 [sampled_image] "" (sampled_image),
180 [Result] "t" (Result),
181 [coordinate] "" (coordinate),
182 );
183}
184
185/// Query the dimensions of `image`, with no level of detail.
186pub fn imageQuerySize(
187 image: anytype,
188) ImageCoordinate(std.meta.Child(@TypeOf(image)), u32) {
189 const Image = switch (@typeInfo(@TypeOf(image))) {
190 .pointer => |pointer| pointer.child,
191 else => @compileError("Expected a pointer to SPIR-V image type, found '" ++ @typeName(@TypeOf(image)) ++ "'"),
192 };
193
194 const image_info = switch (@typeInfo(Image)) {
195 .spirv => |spirv| switch (spirv) {
196 .image => |info| info,
197 else => @compileError("Expected SPIR-V image type, found '" ++ @typeName(Image) ++ "'"),
198 },
199 else => @compileError("Expected SPIR-V image type, found '" ++ @typeName(Image) ++ "'"),
200 };
201
202 // TODO: Remove this check if dimension is not 1d, 2d, 3d, or cube (in case buffer is added).
203 if (!image_info.multisampled and image_info.usage != .unknown and image_info.usage != .storage)
204 @compileError("SPIR-V image must be either be multisampled or have an unknown or storage usage");
205
206 const Result = ImageCoordinate(std.meta.Child(@TypeOf(image)), u32);
207
208 return asm volatile (
209 \\%loaded_image = OpLoad %Image %image
210 \\%ret = OpImageQuerySize %Result %loaded_image
211 : [ret] "" (-> Result),
212 : [Image] "t" (Image),
213 [image] "" (image),
214 [Result] "t" (Result),
215 );
216}
217
218/// Write a texel to an image without a sampler.
219/// The type of `image` must be a pointer to a SPIR-V image.
220pub fn imageWrite(
221 image: anytype,
222 T: type,
223 coordinate: ImageCoordinate(std.meta.Child(@TypeOf(image)), T),
224 texel: @Vector(4, ImageSampledType(std.meta.Child(@TypeOf(image)))),
225) void {
226 switch (T) {
227 u32, i32 => {},
228 f32 => if (builtin.target.os.tag != .opencl) {
229 @compileError("Floating point image coordinates only supported by OpenCL");
230 },
231 else => @compileError("Expected one of u32, i32 and f32 types. Found '" ++ @typeName(T) ++ "'"),
232 }
233
234 const Image = switch (@typeInfo(@TypeOf(image))) {
235 .pointer => |pointer| pointer.child,
236 else => @compileError("Expected a pointer to SPIR-V image type, found '" ++ @typeName(@TypeOf(image)) ++ "'"),
237 };
238
239 const image_info = switch (@typeInfo(Image)) {
240 .spirv => |spirv| switch (spirv) {
241 .image => |info| info,
242 else => @compileError("Expected SPIR-V image type, found '" ++ @typeName(Image) ++ "'"),
243 },
244 else => @compileError("Expected SPIR-V image type, found '" ++ @typeName(Image) ++ "'"),
245 };
246
247 switch (image_info.usage) {
248 .unknown, .storage => {},
249 else => @compileError("SPIR-V image must have unknown or storage usage"),
250 }
251
252 // TODO: If SubpassData dim is added, throw a compiler error if the image is arrayed and has the SubpassData dim.
253
254 return asm volatile (
255 \\%loaded_image = OpLoad %Image %image
256 \\ OpImageWrite %loaded_image %coordinate %texel
257 :
258 : [Image] "t" (Image),
259 [image] "" (image),
260 [coordinate] "" (coordinate),
261 [texel] "" (texel),
262 );
263}
264
265pub fn setMeshOutputs(vertex_count: u32, primitive_count: u32) void {
266 asm volatile (
267 \\OpSetMeshOutputsEXT %vertex_count %primitive_count
268 :
269 : [vertex_count] "" (vertex_count),
270 [primitive_count] "" (primitive_count),
271 );
272}