11const std = @import ("std.zig" );
22
3- /// Will make `ptr` contain the location of the current invocation within the
4- /// global workgroup. Each component is equal to the index of the local workgroup
5- /// multiplied by the size of the local workgroup plus `localInvocationId`.
6- /// `ptr` must be a reference to variable or struct field.
7- pub fn globalInvocationId (comptime ptr : * addrspace (.input ) @Vector (3 , u32 )) void {
8- asm volatile (
9- \\OpDecorate %ptr BuiltIn GlobalInvocationId
10- :
11- : [ptr ] "" (ptr ),
12- );
13- }
14-
15- /// Will make that variable contain the location of the current cluster
16- /// culling, task, mesh, or compute shader invocation within the local
17- /// workgroup. Each component ranges from zero through to the size of the
18- /// workgroup in that dimension minus one.
19- /// `ptr` must be a reference to variable or struct field.
20- pub fn localInvocationId (comptime ptr : * addrspace (.input ) @Vector (3 , u32 )) void {
21- asm volatile (
22- \\OpDecorate %ptr BuiltIn LocalInvocationId
23- :
24- : [ptr ] "" (ptr ),
25- );
26- }
27-
28- /// Output vertex position from a `Vertex` entrypoint
29- /// `ptr` must be a reference to variable or struct field.
30- pub fn position (comptime ptr : * addrspace (.output ) @Vector (4 , f32 )) void {
31- asm volatile (
32- \\OpDecorate %ptr BuiltIn Position
33- :
34- : [ptr ] "" (ptr ),
35- );
36- }
37-
38- /// Will make `ptr` contain the index of the vertex that is
39- /// being processed by the current vertex shader invocation.
40- /// `ptr` must be a reference to variable or struct field.
41- pub fn vertexIndex (comptime ptr : * addrspace (.input ) u32 ) void {
42- asm volatile (
43- \\OpDecorate %ptr BuiltIn VertexIndex
44- :
45- : [ptr ] "" (ptr ),
46- );
47- }
48-
49- /// Will make `ptr` contain the index of the instance that is
50- /// being processed by the current vertex shader invocation.
51- /// `ptr` must be a reference to variable or struct field.
52- pub fn instanceIndex (comptime ptr : * addrspace (.input ) u32 ) void {
53- asm volatile (
54- \\OpDecorate %ptr BuiltIn InstanceIndex
55- :
56- : [ptr ] "" (ptr ),
57- );
58- }
59-
60- /// Output fragment depth from a `Fragment` entrypoint
61- /// `ptr` must be a reference to variable or struct field.
62- pub fn fragmentCoord (comptime ptr : * addrspace (.input ) @Vector (4 , f32 )) void {
63- asm volatile (
64- \\OpDecorate %ptr BuiltIn FragCoord
65- :
66- : [ptr ] "" (ptr ),
67- );
68- }
69-
70- /// Output fragment depth from a `Fragment` entrypoint
71- /// `ptr` must be a reference to variable or struct field.
72- pub fn fragmentDepth (comptime ptr : * addrspace (.output ) f32 ) void {
73- asm volatile (
74- \\OpDecorate %ptr BuiltIn FragDepth
75- :
76- : [ptr ] "" (ptr ),
77- );
78- }
3+ pub const position_in = @extern (* addrspace (.input ) @Vector (4 , f32 ), .{ .name = "position" });
4+ pub const position_out = @extern (* addrspace (.output ) @Vector (4 , f32 ), .{ .name = "position" });
5+ pub const point_size_in = @extern (* addrspace (.input ) f32 , .{ .name = "point_size" });
6+ pub const point_size_out = @extern (* addrspace (.output ) f32 , .{ .name = "point_size" });
7+ pub extern const invocation_id : u32 addrspace (.input );
8+ pub extern const frag_coord : @Vector (4 , f32 ) addrspace (.input );
9+ pub 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;
14+ pub extern var frag_depth : f32 addrspace (.output );
15+ pub extern const num_workgroups : @Vector (3 , u32 ) addrspace (.input );
16+ pub extern const workgroup_size : @Vector (3 , u32 ) addrspace (.input );
17+ pub extern const workgroup_id : @Vector (3 , u32 ) addrspace (.input );
18+ pub extern const local_invocation_id : @Vector (3 , u32 ) addrspace (.input );
19+ pub extern const global_invocation_id : @Vector (3 , u32 ) addrspace (.input );
20+ pub extern const vertex_index : u32 addrspace (.input );
21+ pub extern const instance_index : u32 addrspace (.input );
7922
8023/// Forms the main linkage for `input` and `output` address spaces.
8124/// `ptr` must be a reference to variable or struct field.
@@ -101,74 +44,85 @@ pub fn binding(comptime ptr: anytype, comptime set: u32, comptime bind: u32) voi
10144 );
10245}
10346
104- pub const Origin = enum (u32 ) {
105- /// Increase toward the right and downward
106- upper_left = 7 ,
107- /// Increase toward the right and upward
108- lower_left = 8 ,
109- };
110-
111- /// The coordinates appear to originate in the specified `origin`.
112- /// Only valid with the `Fragment` calling convention.
113- pub fn fragmentOrigin (comptime entry_point : anytype , comptime origin : Origin ) void {
114- asm volatile (
115- \\OpExecutionMode %entry_point $origin
116- :
117- : [entry_point ] "" (entry_point ),
118- [origin ] "c" (@intFromEnum (origin )),
119- );
120- }
121-
122- pub const DepthMode = enum (u32 ) {
123- /// Declares that this entry point dynamically writes the
124- /// `fragmentDepth` built in-decorated variable.
125- replacing = 12 ,
47+ pub 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 ,
12655 /// Indicates that per-fragment tests may assume that
127- /// any `fragmentDepth ` built in-decorated value written by the shader is
56+ /// any `frag_depth ` built in-decorated value written by the shader is
12857 /// greater-than-or-equal to the fragment’s interpolated depth value
129- greater = 14 ,
58+ depth_greater ,
13059 /// Indicates that per-fragment tests may assume that
131- /// any `fragmentDepth ` built in-decorated value written by the shader is
60+ /// any `frag_depth ` built in-decorated value written by the shader is
13261 /// less-than-or-equal to the fragment’s interpolated depth value
133- less = 15 ,
62+ depth_less ,
13463 /// Indicates that per-fragment tests may assume that
135- /// any `fragmentDepth ` built in-decorated value written by the shader is
64+ /// any `frag_depth ` built in-decorated value written by the shader is
13665 /// the same as the fragment’s interpolated depth value
137- unchanged = 16 ,
138- };
66+ depth_unchanged ,
67+ /// Indicates the workgroup size in the x, y, and z dimensions.
68+ local_size : LocalSize ,
13969
140- /// Only valid with the `Fragment` calling convention.
141- pub fn depthMode ( comptime entry_point : anytype , comptime mode : DepthMode ) void {
142- asm volatile (
143- \\OpExecutionMode %entry_point $mode
144- :
145- : [ entry_point ] "" ( entry_point ) ,
146- [ mode ] "c" ( mode ) ,
147- );
148- }
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+ };
14979
150- /// Indicates the workgroup size in the `x`, `y`, and `z` dimensions.
151- /// Only valid with the `GLCompute` or `Kernel` calling conventions.
152- pub fn workgroupSize (comptime entry_point : anytype , comptime size : @Vector (3 , u32 )) void {
153- asm volatile (
154- \\OpExecutionMode %entry_point LocalSize %x %y %z
155- :
156- : [entry_point ] "" (entry_point ),
157- [x ] "c" (size [0 ]),
158- [y ] "c" (size [1 ]),
159- [z ] "c" (size [2 ]),
160- );
161- }
80+ pub const LocalSize = struct { x : u32 , y : u32 , z : u32 };
81+ };
16282
163- /// A hint to the client, which indicates the workgroup size in the `x`, `y`, and `z` dimensions.
164- /// Only valid with the `GLCompute` or `Kernel` calling conventions.
165- pub fn workgroupSizeHint (comptime entry_point : anytype , comptime size : @Vector (3 , u32 )) void {
166- asm volatile (
167- \\OpExecutionMode %entry_point LocalSizeHint %x %y %z
168- :
169- : [entry_point ] "" (entry_point ),
170- [x ] "c" (size [0 ]),
171- [y ] "c" (size [1 ]),
172- [z ] "c" (size [2 ]),
173- );
83+ /// Declare the mode entry point executes in.
84+ pub 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+ }
174128}
0 commit comments