removing spirv reflection code from driver

build.zig

@@ -88,6 +88,7 @@ pub fn build(b: *std.Build) !void {
             .root_source_file = b.path(impl.root_source_file),
             .target = target,
             .optimize = optimize,
+            //.error_tracing = true,
             .imports = &.{
                 .{ .name = "base", .module = base_mod },
                 .{ .name = "vulkan", .module = vulkan },
@@ -133,13 +134,10 @@ pub fn build(b: *std.Build) !void {
         const test_step = b.step(b.fmt("test-{s}", .{impl.name}), b.fmt("Run libvulkan_{s} tests", .{impl.name}));
         test_step.dependOn(&run_tests.step);
 
-        (try addCTS(b, target, &impl, lib, .normal)).dependOn(&lib_install.step);
+        inline for (std.enums.values(RunningMode)) |mode| {
-        (try addCTS(b, target, &impl, lib, .gdb)).dependOn(&lib_install.step);
+            (try addCTS(b, target, &impl, lib, mode)).dependOn(&lib_install.step);
-        (try addCTS(b, target, &impl, lib, .valgrind)).dependOn(&lib_install.step);
+            (try addMultithreadedCTS(b, target, &impl, lib, mode)).dependOn(&lib_install.step);
-
+        }
-        (try addMultithreadedCTS(b, target, &impl, lib, .normal)).dependOn(&lib_install.step);
-        (try addMultithreadedCTS(b, target, &impl, lib, .gdb)).dependOn(&lib_install.step);
-        (try addMultithreadedCTS(b, target, &impl, lib, .valgrind)).dependOn(&lib_install.step);
 
         const impl_autodoc_test = b.addObject(.{
             .name = "lib",
@@ -212,7 +210,7 @@ fn addCTS(b: *std.Build, target: std.Build.ResolvedTarget, impl: *const Implemen
         switch (if (target.query.os_tag) |tag| tag else builtin.target.os.tag) {
             .linux => "linux.x86_64",
             .windows => "windows.exe",
-            else => unreachable,
+            else => return error.NoCTSForPlatform,
         },
     }));
 
@@ -289,7 +287,7 @@ fn addMultithreadedCTS(b: *std.Build, target: std.Build.ResolvedTarget, impl: *c
         switch (if (target.query.os_tag) |tag| tag else builtin.target.os.tag) {
             .linux => "linux.x86_64",
             .windows => "windows.exe",
-            else => unreachable,
+            else => return error.NoCTSForPlatform,
         },
     }));
 

build.zig.zon

@@ -26,8 +26,8 @@
             .hash = "N-V-__8AAF9uOh0I4P_99za7N822J3JwsDaqONrFVrcEQo59",
         },
         .SPIRV_Interpreter = .{
-            .url = "git+https://git.kbz8.me/kbz_8/SPIRV-Interpreter#11be3ea41dd5b48d4547732c4ea156381def3802",
+            .url = "git+https://git.kbz8.me/kbz_8/SPIRV-Interpreter#c868a4481b6be4f7fb07c90cffcb359e4e21a7fb",
-            .hash = "SPIRV_Interpreter-0.0.1-ajmpnwN2BgC6Wg-dhBksmlzTkV6anEx_dFsZRmsXR7ZQ",
+            .hash = "SPIRV_Interpreter-0.0.1-ajmpn4TyBgAgPySLGtH-ewEI0b53c3nRIjCGTugLyunr",
             .lazy = true,
         },
         //.SPIRV_Interpreter = .{

src/soft/device/ComputeDispatcher.zig

@@ -55,7 +55,7 @@ pub fn dispatch(self: *Self, group_count_x: u32, group_count_y: u32, group_count
     const spv_module = &shader.module.module;
     self.batch_size = shader.runtimes.len;
 
-    const invocations_per_workgroup = spv_module.local_size_x * spv_module.local_size_y * spv_module.local_size_z;
+    const invocations_per_workgroup = spv_module.reflection_infos.local_size_x * spv_module.reflection_infos.local_size_y * spv_module.reflection_infos.local_size_z;
 
     self.invocation_index.store(0, .monotonic);
 
@@ -107,7 +107,7 @@ inline fn run(data: RunData) !void {
     const rt = &shader.runtimes[data.batch_id].rt;
 
     const entry = try rt.getEntryPointByName(shader.entry);
-    const uses_control_barrier = hasControlBarrier(rt.mod.code);
+    const uses_control_barrier = rt.mod.reflection_infos.has_control_barriers;
 
     var barrier_runtimes: []spv.Runtime = &.{};
     var barrier_statuses: []spv.Runtime.EntryPointStatus = &.{};
@@ -235,20 +235,6 @@ fn runBarrierWorkgroup(
     }
 }
 
-/// TODO: Move this in the SPIR-V Interpreter
-fn hasControlBarrier(code: []const spv.SpvWord) bool {
-    var i: usize = 5;
-    while (i < code.len) {
-        const opcode_data = code[i];
-        const word_count = (opcode_data & (~spv.spv.SpvOpCodeMask)) >> spv.spv.SpvWordCountShift;
-        const opcode: spv.spv.SpvOp = @enumFromInt(opcode_data & spv.spv.SpvOpCodeMask);
-        if (opcode == .ControlBarrier)
-            return true;
-        i += @max(word_count, 1);
-    }
-    return false;
-}
-
 inline fn dumpResultsTable(allocator: std.mem.Allocator, io: std.Io, rt: *spv.Runtime, is_early: bool) !void {
     @branchHint(.cold);
     const file = try std.Io.Dir.cwd().createFile(
@@ -262,17 +248,12 @@ inline fn dumpResultsTable(allocator: std.mem.Allocator, io: std.Io, rt: *spv.Ru
     try rt.dumpResultsTable(allocator, &writer.interface);
 }
 
-fn setupWorkgroupBuiltins(
+fn setupWorkgroupBuiltins(self: *Self, rt: *spv.Runtime, group_count: @Vector(3, u32), group_id: @Vector(3, u32)) spv.Runtime.RuntimeError!void {
-    self: *Self,
-    rt: *spv.Runtime,
-    group_count: @Vector(3, u32),
-    group_id: @Vector(3, u32),
-) spv.Runtime.RuntimeError!void {
     const spv_module = &self.state.pipeline.?.stages.getPtrAssertContains(.compute).module.module;
     const workgroup_size = @Vector(3, u32){
-        spv_module.local_size_x,
+        spv_module.reflection_infos.local_size_x,
-        spv_module.local_size_y,
+        spv_module.reflection_infos.local_size_y,
-        spv_module.local_size_z,
+        spv_module.reflection_infos.local_size_z,
     };
 
     rt.writeBuiltIn(std.mem.asBytes(&workgroup_size), .WorkgroupSize) catch {};
@@ -280,17 +261,12 @@ fn setupWorkgroupBuiltins(
     rt.writeBuiltIn(std.mem.asBytes(&group_id), .WorkgroupId) catch {};
 }
 
-fn setupSubgroupBuiltins(
+fn setupSubgroupBuiltins(self: *Self, rt: *spv.Runtime, group_id: @Vector(3, u32), local_invocation_index: usize) spv.Runtime.RuntimeError!void {
-    self: *Self,
-    rt: *spv.Runtime,
-    group_id: @Vector(3, u32),
-    local_invocation_index: usize,
-) spv.Runtime.RuntimeError!void {
     const spv_module = &self.state.pipeline.?.stages.getPtrAssertContains(.compute).module.module;
     const workgroup_size = @Vector(3, u32){
-        spv_module.local_size_x,
+        spv_module.reflection_infos.local_size_x,
-        spv_module.local_size_y,
+        spv_module.reflection_infos.local_size_y,
-        spv_module.local_size_z,
+        spv_module.reflection_infos.local_size_z,
     };
     const local_base = workgroup_size * group_id;
     var local_invocation = @Vector(3, u32){ 0, 0, 0 };

src/soft/device/fragment.zig

@@ -18,36 +18,12 @@ pub const DerivativeInputs = struct {
     dy: [spv.SPIRV_MAX_OUTPUT_LOCATIONS]VertexInterpolationLocation,
 };
 
-pub fn shaderUsesDerivatives(code: []const spv.SpvWord) bool {
-    var i: usize = 5;
-    while (i < code.len) {
-        const opcode_data = code[i];
-        const word_count = (opcode_data & (~spv.spv.SpvOpCodeMask)) >> spv.spv.SpvWordCountShift;
-        if (word_count == 0)
-            return false;
-
-        const opcode = opcode_data & spv.spv.SpvOpCodeMask;
-        switch (opcode) {
-            @intFromEnum(spv.spv.SpvOp.DPdx),
-            @intFromEnum(spv.spv.SpvOp.DPdy),
-            @intFromEnum(spv.spv.SpvOp.DPdxFine),
-            @intFromEnum(spv.spv.SpvOp.DPdyFine),
-            @intFromEnum(spv.spv.SpvOp.DPdxCoarse),
-            @intFromEnum(spv.spv.SpvOp.DPdyCoarse),
-            => return true,
-            else => {},
-        }
-
-        i += word_count;
-    }
-    return false;
-}
-
 pub fn shaderInvocation(
     allocator: std.mem.Allocator,
     draw_call: *Renderer.DrawCall,
     batch_id: usize,
     position: zm.F32x4,
+    front_face: bool,
     inputs: [spv.SPIRV_MAX_OUTPUT_LOCATIONS]VertexInterpolationLocation,
     derivative_inputs: ?DerivativeInputs,
 ) SpvRuntimeError![spv.SPIRV_MAX_OUTPUT_LOCATIONS][@sizeOf(zm.F32x4)]u8 {
@@ -77,6 +53,10 @@ pub fn shaderInvocation(
         SpvRuntimeError.NotFound => {},
         else => return err,
     };
+    rt.writeBuiltIn(std.mem.asBytes(&front_face), .FrontFacing) catch |err| switch (err) {
+        SpvRuntimeError.NotFound => {},
+        else => return err,
+    };
 
     const entry = try rt.getEntryPointByName(shader.entry);
 
@@ -143,9 +123,7 @@ pub fn shaderInvocation(
                     SpvRuntimeError.NotFound => continue,
                     else => return err,
                 };
-                const memory_size = try rt.getResultMemorySize(result_word);
+                try readFragmentOutput(allocator, rt, &outputs, location, component, result_word);
-                const offset = component * @sizeOf(f32);
-                try rt.readOutput(outputs[location][offset .. offset + memory_size], result_word);
             }
             continue;
         }
@@ -154,7 +132,7 @@ pub fn shaderInvocation(
             SpvRuntimeError.NotFound => continue,
             else => return err,
         };
-        try rt.readOutput(&outputs[location], result_word);
+        try readFragmentOutput(allocator, rt, &outputs, location, 0, result_word);
     }
 
     try rt.flushDescriptorSets(allocator);
@@ -167,6 +145,48 @@ pub fn shaderInvocation(
     return outputs;
 }
 
+fn readFragmentOutput(
+    allocator: std.mem.Allocator,
+    rt: anytype,
+    outputs: *[spv.SPIRV_MAX_OUTPUT_LOCATIONS][@sizeOf(zm.F32x4)]u8,
+    location: usize,
+    component: usize,
+    result_word: spv.SpvWord,
+) SpvRuntimeError!void {
+    const memory_size = try rt.getResultMemorySize(result_word);
+    const offset = component * @sizeOf(f32);
+    const location_size = @sizeOf(zm.F32x4);
+    const direct_capacity = location_size - offset;
+
+    if (memory_size <= direct_capacity) {
+        try rt.readOutput(outputs[location][offset .. offset + memory_size], result_word);
+        return;
+    }
+
+    const output = allocator.alloc(u8, memory_size + INTERFACE_BLOB_PADDING) catch return SpvRuntimeError.OutOfMemory;
+    defer allocator.free(output);
+    @memset(output, 0);
+
+    try rt.readOutput(output, result_word);
+
+    var source_offset: usize = 0;
+    var target_location = location;
+    var target_offset = offset;
+    while (source_offset < memory_size and target_location < outputs.len) {
+        const copy_size = @min(memory_size - source_offset, location_size - target_offset);
+        @memcpy(
+            outputs[target_location][target_offset .. target_offset + copy_size],
+            output[source_offset .. source_offset + copy_size],
+        );
+        source_offset += copy_size;
+        target_location += 1;
+        target_offset = offset;
+    }
+
+    if (source_offset != memory_size)
+        return SpvRuntimeError.OutOfBounds;
+}
+
 fn freeOwnedInputs(allocator: std.mem.Allocator, inputs: [spv.SPIRV_MAX_OUTPUT_LOCATIONS]VertexInterpolationLocation) void {
     for (inputs) |location| {
         for (location) |input| {

src/soft/device/rasterizer.zig

@@ -311,6 +311,7 @@ fn clipTransformAndRasterizePoint(
                     draw_call,
                     0,
                     zm.f32x4(@floatFromInt(px), @floatFromInt(py), transformed.position[2], 1.0),
+                    true,
                     try common.interpolateVertexOutputs(allocator, &transformed, &transformed, &transformed, 1.0, 0.0, 0.0),
                     null,
                 ) catch |err| {

src/soft/device/rasterizer/bresenham.zig

@@ -159,6 +159,7 @@ inline fn run(data: RunData) !void {
                 data.draw_call,
                 data.batch_id,
                 zm.f32x4(@floatFromInt(pixel_x), @floatFromInt(pixel_y), z, 1.0),
+                true,
                 try common.interpolateLineOutputs(data.allocator, data.start_vertex, data.end_vertex, t),
                 null,
             ) catch |err| {

src/soft/device/rasterizer/edge_function.zig

@@ -58,9 +58,10 @@ pub fn drawTriangle(
     const pipeline = draw_call.renderer.state.pipeline orelse return;
     const fragment_stage = pipeline.stages.getPtr(.fragment);
     const fragment_uses_derivatives = if (fragment_stage) |stage|
-        fragment.shaderUsesDerivatives(stage.module.module.code)
+        stage.module.module.reflection_infos.needs_derivatives
     else
         false;
+
     const runtimes_count = if (fragment_stage) |stage| stage.runtimes.len else 1;
     if (runtimes_count == 0)
         return;
@@ -218,6 +219,7 @@ inline fn run(data: RunData) !void {
                     data.draw_call,
                     data.batch_id,
                     zm.f32x4(@floatFromInt(x), @floatFromInt(y), z, 1.0),
+                    data.front_face,
                     inputs,
                     derivative_inputs,
                 ) catch |err| {