authorgravatar for alichraghi@proton.meAli Chraghi <alichraghi@proton.me> 2025-05-11 15:45:44+03:30
committergravatar for alichraghi@proton.meAli Chraghi <alichraghi@proton.me> 2025-05-21 13:01:21+03:30
log8fa54eb7987bdb8138c625f03aa9fb91239dba48
tree4565cd41420208f0da774c967dc0af0074a504cc
parent9209f4b16acc4453f89a06caabf54691f6253f62

spirv: error when execution mode is set more than once

5 files changed, 143 insertions(+), 166 deletions(-)

lib/std/gpu.zig+91-137
...@@ -1,81 +1,24 @@...@@ -1,81 +1,24 @@
1const std = @import("std.zig");1const std = @import("std.zig");
22
3/// Will make `ptr` contain the location of the current invocation within the3pub const position_in = @extern(*addrspace(.input) @Vector(4, f32), .{ .name = "position" });
4/// global workgroup. Each component is equal to the index of the local workgroup4pub const position_out = @extern(*addrspace(.output) @Vector(4, f32), .{ .name = "position" });
5/// multiplied by the size of the local workgroup plus `localInvocationId`.5pub const point_size_in = @extern(*addrspace(.input) f32, .{ .name = "point_size" });
6/// `ptr` must be a reference to variable or struct field.6pub const point_size_out = @extern(*addrspace(.output) f32, .{ .name = "point_size" });
7pub fn globalInvocationId(comptime ptr: *addrspace(.input) @Vector(3, u32)) void {7pub extern const invocation_id: u32 addrspace(.input);
8 asm volatile (8pub extern const frag_coord: @Vector(4, f32) addrspace(.input);
9 \\OpDecorate %ptr BuiltIn GlobalInvocationId9pub extern const point_coord: @Vector(2, f32) addrspace(.input);
10 :10// TODO: direct/indirect values
11 : [ptr] "" (ptr),11// pub extern const front_facing: bool addrspace(.input);
12 );12// TODO: runtime array
13}13// pub extern const sample_mask;
1414pub extern var frag_depth: f32 addrspace(.output);
15/// Will make that variable contain the location of the current cluster15pub extern const num_workgroups: @Vector(3, u32) addrspace(.input);
16/// culling, task, mesh, or compute shader invocation within the local16pub extern const workgroup_size: @Vector(3, u32) addrspace(.input);
17/// workgroup. Each component ranges from zero through to the size of the17pub extern const workgroup_id: @Vector(3, u32) addrspace(.input);
18/// workgroup in that dimension minus one.18pub extern const local_invocation_id: @Vector(3, u32) addrspace(.input);
19/// `ptr` must be a reference to variable or struct field.19pub extern const global_invocation_id: @Vector(3, u32) addrspace(.input);
20pub fn localInvocationId(comptime ptr: *addrspace(.input) @Vector(3, u32)) void {20pub extern const vertex_index: u32 addrspace(.input);
21 asm volatile (21pub extern const instance_index: u32 addrspace(.input);
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.
30pub 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.
41pub 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.
52pub 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.
62pub 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.
72pub fn fragmentDepth(comptime ptr: *addrspace(.output) f32) void {
73 asm volatile (
74 \\OpDecorate %ptr BuiltIn FragDepth
75 :
76 : [ptr] "" (ptr),
77 );
78}
7922
80/// Forms the main linkage for `input` and `output` address spaces.23/// Forms the main linkage for `input` and `output` address spaces.
81/// `ptr` must be a reference to variable or struct field.24/// `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...@@ -101,74 +44,85 @@ pub fn binding(comptime ptr: anytype, comptime set: u32, comptime bind: u32) voi
101 );44 );
102}45}
10346
104pub const Origin = enum(u32) {47pub const ExecutionMode = union(Tag) {
105 /// Increase toward the right and downward48 /// Sets origin of the framebuffer to the upper-left corner
106 upper_left = 7,49 origin_upper_left,
107 /// Increase toward the right and upward50 /// Sets origin of the framebuffer to the lower-left corner
108 lower_left = 8,51 origin_lower_left,
109};52 /// Indicates that the fragment shader writes to `frag_depth`,
11053 /// replacing the fixed-function depth value.
111/// The coordinates appear to originate in the specified `origin`.54 depth_replacing,
112/// Only valid with the `Fragment` calling convention.
113pub 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
122pub const DepthMode = enum(u32) {
123 /// Declares that this entry point dynamically writes the
124 /// `fragmentDepth` built in-decorated variable.
125 replacing = 12,
126 /// Indicates that per-fragment tests may assume that55 /// Indicates that per-fragment tests may assume that
127 /// any `fragmentDepth` built in-decorated value written by the shader is56 /// any `frag_depth` built in-decorated value written by the shader is
128 /// greater-than-or-equal to the fragment’s interpolated depth value57 /// greater-than-or-equal to the fragment’s interpolated depth value
129 greater = 14,58 depth_greater,
130 /// Indicates that per-fragment tests may assume that59 /// Indicates that per-fragment tests may assume that
131 /// any `fragmentDepth` built in-decorated value written by the shader is60 /// any `frag_depth` built in-decorated value written by the shader is
132 /// less-than-or-equal to the fragment’s interpolated depth value61 /// less-than-or-equal to the fragment’s interpolated depth value
133 less = 15,62 depth_less,
134 /// Indicates that per-fragment tests may assume that63 /// Indicates that per-fragment tests may assume that
135 /// any `fragmentDepth` built in-decorated value written by the shader is64 /// any `frag_depth` built in-decorated value written by the shader is
136 /// the same as the fragment’s interpolated depth value65 /// the same as the fragment’s interpolated depth value
137 unchanged = 16,66 depth_unchanged,
138};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.70 pub const Tag = enum(u32) {
141pub fn depthMode(comptime entry_point: anytype, comptime mode: DepthMode) void {71 origin_upper_left = 7,
142 asm volatile (72 origin_lower_left = 8,
143 \\OpExecutionMode %entry_point $mode73 depth_replacing = 12,
144 :74 depth_greater = 14,
145 : [entry_point] "" (entry_point),75 depth_less = 15,
146 [mode] "c" (mode),76 depth_unchanged = 16,
147 );77 local_size = 17,
148}78 };
14979
150/// Indicates the workgroup size in the `x`, `y`, and `z` dimensions.80 pub const LocalSize = struct { x: u32, y: u32, z: u32 };
151/// Only valid with the `GLCompute` or `Kernel` calling conventions.81};
152pub 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}
16282
163/// A hint to the client, which indicates the workgroup size in the `x`, `y`, and `z` dimensions.83/// Declare the mode entry point executes in.
164/// Only valid with the `GLCompute` or `Kernel` calling conventions.84pub fn executionMode(comptime entry_point: anytype, comptime mode: ExecutionMode) void {
165pub fn workgroupSizeHint(comptime entry_point: anytype, comptime size: @Vector(3, u32)) void {85 const cc = @typeInfo(@TypeOf(entry_point)).@"fn".calling_convention;
166 asm volatile (86 switch (mode) {
167 \\OpExecutionMode %entry_point LocalSizeHint %x %y %z87 .origin_upper_left,
168 :88 .origin_lower_left,
169 : [entry_point] "" (entry_point),89 .depth_replacing,
170 [x] "c" (size[0]),90 .depth_greater,
171 [y] "c" (size[1]),91 .depth_less,
172 [z] "c" (size[2]),92 .depth_unchanged,
173 );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 }
174}128}
src/codegen/spirv.zig+1-7
...@@ -2870,7 +2870,7 @@ const NavGen = struct {...@@ -2870,7 +2870,7 @@ const NavGen = struct {
2870 };2870 };
28712871
2872 try self.spv.declareDeclDeps(spv_decl_index, decl_deps.items);2872 try self.spv.declareDeclDeps(spv_decl_index, decl_deps.items);
2873 try self.spv.declareEntryPoint(spv_decl_index, test_name, execution_mode);2873 try self.spv.declareEntryPoint(spv_decl_index, test_name, execution_mode, null);
2874 }2874 }
28752875
2876 fn genNav(self: *NavGen, do_codegen: bool) !void {2876 fn genNav(self: *NavGen, do_codegen: bool) !void {
...@@ -2976,10 +2976,6 @@ const NavGen = struct {...@@ -2976,10 +2976,6 @@ const NavGen = struct {
2976 try self.spv.decorate(result_id, .{ .BuiltIn = .{ .built_in = .Position } });2976 try self.spv.decorate(result_id, .{ .BuiltIn = .{ .built_in = .Position } });
2977 } else if (nav.fqn.eqlSlice("point_size", ip)) {2977 } else if (nav.fqn.eqlSlice("point_size", ip)) {
2978 try self.spv.decorate(result_id, .{ .BuiltIn = .{ .built_in = .PointSize } });2978 try self.spv.decorate(result_id, .{ .BuiltIn = .{ .built_in = .PointSize } });
2979 } else if (nav.fqn.eqlSlice("vertex_id", ip)) {
2980 try self.spv.decorate(result_id, .{ .BuiltIn = .{ .built_in = .VertexId } });
2981 } else if (nav.fqn.eqlSlice("instance_id", ip)) {
2982 try self.spv.decorate(result_id, .{ .BuiltIn = .{ .built_in = .InstanceId } });
2983 } else if (nav.fqn.eqlSlice("invocation_id", ip)) {2979 } else if (nav.fqn.eqlSlice("invocation_id", ip)) {
2984 try self.spv.decorate(result_id, .{ .BuiltIn = .{ .built_in = .InvocationId } });2980 try self.spv.decorate(result_id, .{ .BuiltIn = .{ .built_in = .InvocationId } });
2985 } else if (nav.fqn.eqlSlice("frag_coord", ip)) {2981 } else if (nav.fqn.eqlSlice("frag_coord", ip)) {
...@@ -2990,8 +2986,6 @@ const NavGen = struct {...@@ -2990,8 +2986,6 @@ const NavGen = struct {
2990 try self.spv.decorate(result_id, .{ .BuiltIn = .{ .built_in = .FrontFacing } });2986 try self.spv.decorate(result_id, .{ .BuiltIn = .{ .built_in = .FrontFacing } });
2991 } else if (nav.fqn.eqlSlice("sample_mask", ip)) {2987 } else if (nav.fqn.eqlSlice("sample_mask", ip)) {
2992 try self.spv.decorate(result_id, .{ .BuiltIn = .{ .built_in = .SampleMask } });2988 try self.spv.decorate(result_id, .{ .BuiltIn = .{ .built_in = .SampleMask } });
2993 } else if (nav.fqn.eqlSlice("sample_mask", ip)) {
2994 try self.spv.decorate(result_id, .{ .BuiltIn = .{ .built_in = .SampleMask } });
2995 } else if (nav.fqn.eqlSlice("frag_depth", ip)) {2989 } else if (nav.fqn.eqlSlice("frag_depth", ip)) {
2996 try self.spv.decorate(result_id, .{ .BuiltIn = .{ .built_in = .FragDepth } });2990 try self.spv.decorate(result_id, .{ .BuiltIn = .{ .built_in = .FragDepth } });
2997 } else if (nav.fqn.eqlSlice("num_workgroups", ip)) {2991 } else if (nav.fqn.eqlSlice("num_workgroups", ip)) {
src/codegen/spirv/Assembler.zig+18-4
...@@ -296,12 +296,26 @@ fn processInstruction(self: *Assembler) !void {...@@ -296,12 +296,26 @@ fn processInstruction(self: *Assembler) !void {
296 };296 };
297 break :blk .{ .value = try self.spv.importInstructionSet(set_tag) };297 break :blk .{ .value = try self.spv.importInstructionSet(set_tag) };
298 },298 },
299 .OpExecutionMode, .OpExecutionModeId => {
300 assert(try self.processGenericInstruction() == null);
301 const entry_point_id = try self.resolveRefId(self.inst.operands.items[0].ref_id);
302 const exec_mode: spec.ExecutionMode = @enumFromInt(self.inst.operands.items[1].value);
303 const gop = try self.spv.entry_points.getOrPut(self.gpa, entry_point_id);
304 if (!gop.found_existing) {
305 gop.value_ptr.* = .{};
306 } else if (gop.value_ptr.exec_mode != null) {
307 return self.fail(
308 self.currentToken().start,
309 "cannot set execution mode more than once to any entry point",
310 .{},
311 );
312 }
313 gop.value_ptr.exec_mode = exec_mode;
314 return;
315 },
299 else => switch (self.inst.opcode.class()) {316 else => switch (self.inst.opcode.class()) {
300 .TypeDeclaration => try self.processTypeInstruction(),317 .TypeDeclaration => try self.processTypeInstruction(),
301 else => if (try self.processGenericInstruction()) |result|318 else => (try self.processGenericInstruction()) orelse return,
302 result
303 else
304 return,
305 },319 },
306 };320 };
307321
src/codegen/spirv/Module.zig+30-16
...@@ -92,11 +92,12 @@ pub const Decl = struct {...@@ -92,11 +92,12 @@ pub const Decl = struct {
92/// This models a kernel entry point.92/// This models a kernel entry point.
93pub const EntryPoint = struct {93pub const EntryPoint = struct {
94 /// The declaration that should be exported.94 /// The declaration that should be exported.
95 decl_index: Decl.Index,95 decl_index: ?Decl.Index = null,
96 /// The name of the kernel to be exported.96 /// The name of the kernel to be exported.
97 name: []const u8,97 name: ?[]const u8 = null,
98 /// Calling Convention98 /// Calling Convention
99 execution_model: spec.ExecutionModel,99 exec_model: ?spec.ExecutionModel = null,
100 exec_mode: ?spec.ExecutionMode = null,
100};101};
101102
102/// A general-purpose allocator which may be used to allocate resources for this module103/// A general-purpose allocator which may be used to allocate resources for this module
...@@ -184,7 +185,7 @@ decls: std.ArrayListUnmanaged(Decl) = .empty,...@@ -184,7 +185,7 @@ decls: std.ArrayListUnmanaged(Decl) = .empty,
184decl_deps: std.ArrayListUnmanaged(Decl.Index) = .empty,185decl_deps: std.ArrayListUnmanaged(Decl.Index) = .empty,
185186
186/// The list of entry points that should be exported from this module.187/// The list of entry points that should be exported from this module.
187entry_points: std.ArrayListUnmanaged(EntryPoint) = .empty,188entry_points: std.AutoArrayHashMapUnmanaged(IdRef, EntryPoint) = .empty,
188189
189pub fn init(gpa: Allocator, target: std.Target) Module {190pub fn init(gpa: Allocator, target: std.Target) Module {
190 const version_minor: u8 = blk: {191 const version_minor: u8 = blk: {
...@@ -304,19 +305,30 @@ fn entryPoints(self: *Module) !Section {...@@ -304,19 +305,30 @@ fn entryPoints(self: *Module) !Section {
304 var seen = try std.DynamicBitSetUnmanaged.initEmpty(self.gpa, self.decls.items.len);305 var seen = try std.DynamicBitSetUnmanaged.initEmpty(self.gpa, self.decls.items.len);
305 defer seen.deinit(self.gpa);306 defer seen.deinit(self.gpa);
306307
307 for (self.entry_points.items) |entry_point| {308 for (self.entry_points.keys(), self.entry_points.values()) |entry_point_id, entry_point| {
308 interface.items.len = 0;309 interface.items.len = 0;
309 seen.setRangeValue(.{ .start = 0, .end = self.decls.items.len }, false);310 seen.setRangeValue(.{ .start = 0, .end = self.decls.items.len }, false);
310311
311 try self.addEntryPointDeps(entry_point.decl_index, &seen, &interface);312 try self.addEntryPointDeps(entry_point.decl_index.?, &seen, &interface);
312
313 const entry_point_id = self.declPtr(entry_point.decl_index).result_id;
314 try entry_points.emit(self.gpa, .OpEntryPoint, .{313 try entry_points.emit(self.gpa, .OpEntryPoint, .{
315 .execution_model = entry_point.execution_model,314 .execution_model = entry_point.exec_model.?,
316 .entry_point = entry_point_id,315 .entry_point = entry_point_id,
317 .name = entry_point.name,316 .name = entry_point.name.?,
318 .interface = interface.items,317 .interface = interface.items,
319 });318 });
319
320 if (entry_point.exec_mode == null and entry_point.exec_model == .Fragment) {
321 switch (self.target.os.tag) {
322 .vulkan, .opengl => |tag| {
323 try self.sections.execution_modes.emit(self.gpa, .OpExecutionMode, .{
324 .entry_point = entry_point_id,
325 .mode = if (tag == .vulkan) .OriginUpperLeft else .OriginLowerLeft,
326 });
327 },
328 .opencl => {},
329 else => unreachable,
330 }
331 }
320 }332 }
321333
322 return entry_points;334 return entry_points;
...@@ -749,13 +761,15 @@ pub fn declareEntryPoint(...@@ -749,13 +761,15 @@ pub fn declareEntryPoint(
749 self: *Module,761 self: *Module,
750 decl_index: Decl.Index,762 decl_index: Decl.Index,
751 name: []const u8,763 name: []const u8,
752 execution_model: spec.ExecutionModel,764 exec_model: spec.ExecutionModel,
765 exec_mode: ?spec.ExecutionMode,
753) !void {766) !void {
754 try self.entry_points.append(self.gpa, .{767 const gop = try self.entry_points.getOrPut(self.gpa, self.declPtr(decl_index).result_id);
755 .decl_index = decl_index,768 gop.value_ptr.decl_index = decl_index;
756 .name = try self.arena.allocator().dupe(u8, name),769 gop.value_ptr.name = try self.arena.allocator().dupe(u8, name);
757 .execution_model = execution_model,770 gop.value_ptr.exec_model = exec_model;
758 });771 // Might've been set by assembler
772 if (!gop.found_existing) gop.value_ptr.exec_mode = exec_mode;
759}773}
760774
761pub fn debugName(self: *Module, target: IdResult, name: []const u8) !void {775pub fn debugName(self: *Module, target: IdResult, name: []const u8) !void {
src/link/SpirV.zig+3-2
...@@ -162,7 +162,7 @@ pub fn updateExports(...@@ -162,7 +162,7 @@ pub fn updateExports(
162 if (ip.isFunctionType(nav_ty)) {162 if (ip.isFunctionType(nav_ty)) {
163 const spv_decl_index = try self.object.resolveNav(zcu, nav_index);163 const spv_decl_index = try self.object.resolveNav(zcu, nav_index);
164 const cc = Type.fromInterned(nav_ty).fnCallingConvention(zcu);164 const cc = Type.fromInterned(nav_ty).fnCallingConvention(zcu);
165 const execution_model: spec.ExecutionModel = switch (target.os.tag) {165 const exec_model: spec.ExecutionModel = switch (target.os.tag) {
166 .vulkan, .opengl => switch (cc) {166 .vulkan, .opengl => switch (cc) {
167 .spirv_vertex => .Vertex,167 .spirv_vertex => .Vertex,
168 .spirv_fragment => .Fragment,168 .spirv_fragment => .Fragment,
...@@ -185,7 +185,8 @@ pub fn updateExports(...@@ -185,7 +185,8 @@ pub fn updateExports(
185 try self.object.spv.declareEntryPoint(185 try self.object.spv.declareEntryPoint(
186 spv_decl_index,186 spv_decl_index,
187 exp.opts.name.toSlice(ip),187 exp.opts.name.toSlice(ip),
188 execution_model,188 exec_model,
189 null,
189 );190 );
190 }191 }
191 }192 }