From 33ca4ecb8d7cf3bf645296be42b597de7b2c61a9 Mon Sep 17 00:00:00 2001 From: Kbz-8 Date: Wed, 10 Jun 2026 19:18:32 +0200 Subject: [PATCH] removing spirv reflection code from driver --- build.zig | 16 ++-- build.zig.zon | 4 +- src/soft/device/ComputeDispatcher.zig | 44 +++-------- src/soft/device/fragment.zig | 78 ++++++++++++-------- src/soft/device/rasterizer.zig | 1 + src/soft/device/rasterizer/bresenham.zig | 1 + src/soft/device/rasterizer/edge_function.zig | 4 +- 7 files changed, 73 insertions(+), 75 deletions(-) diff --git a/build.zig b/build.zig index 5896a0a..5b26201 100644 --- a/build.zig +++ b/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); - (try addCTS(b, target, &impl, lib, .gdb)).dependOn(&lib_install.step); - (try addCTS(b, target, &impl, lib, .valgrind)).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); + inline for (std.enums.values(RunningMode)) |mode| { + (try addCTS(b, target, &impl, lib, mode)).dependOn(&lib_install.step); + (try addMultithreadedCTS(b, target, &impl, lib, mode)).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, }, })); diff --git a/build.zig.zon b/build.zig.zon index a22c7f1..5f1377e 100644 --- a/build.zig.zon +++ b/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", - .hash = "SPIRV_Interpreter-0.0.1-ajmpnwN2BgC6Wg-dhBksmlzTkV6anEx_dFsZRmsXR7ZQ", + .url = "git+https://git.kbz8.me/kbz_8/SPIRV-Interpreter#c868a4481b6be4f7fb07c90cffcb359e4e21a7fb", + .hash = "SPIRV_Interpreter-0.0.1-ajmpn4TyBgAgPySLGtH-ewEI0b53c3nRIjCGTugLyunr", .lazy = true, }, //.SPIRV_Interpreter = .{ diff --git a/src/soft/device/ComputeDispatcher.zig b/src/soft/device/ComputeDispatcher.zig index abf6354..6b39dbe 100644 --- a/src/soft/device/ComputeDispatcher.zig +++ b/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( - self: *Self, - rt: *spv.Runtime, - group_count: @Vector(3, u32), - group_id: @Vector(3, u32), -) spv.Runtime.RuntimeError!void { +fn setupWorkgroupBuiltins(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.local_size_y, - spv_module.local_size_z, + spv_module.reflection_infos.local_size_x, + spv_module.reflection_infos.local_size_y, + 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( - self: *Self, - rt: *spv.Runtime, - group_id: @Vector(3, u32), - local_invocation_index: usize, -) spv.Runtime.RuntimeError!void { +fn setupSubgroupBuiltins(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.local_size_y, - spv_module.local_size_z, + spv_module.reflection_infos.local_size_x, + spv_module.reflection_infos.local_size_y, + spv_module.reflection_infos.local_size_z, }; const local_base = workgroup_size * group_id; var local_invocation = @Vector(3, u32){ 0, 0, 0 }; diff --git a/src/soft/device/fragment.zig b/src/soft/device/fragment.zig index 3da7fc8..a67532b 100644 --- a/src/soft/device/fragment.zig +++ b/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); - const offset = component * @sizeOf(f32); - try rt.readOutput(outputs[location][offset .. offset + memory_size], result_word); + try readFragmentOutput(allocator, rt, &outputs, location, component, 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| { diff --git a/src/soft/device/rasterizer.zig b/src/soft/device/rasterizer.zig index 27fcc6f..b3b6679 100644 --- a/src/soft/device/rasterizer.zig +++ b/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| { diff --git a/src/soft/device/rasterizer/bresenham.zig b/src/soft/device/rasterizer/bresenham.zig index 1c18d83..9e5be77 100644 --- a/src/soft/device/rasterizer/bresenham.zig +++ b/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| { diff --git a/src/soft/device/rasterizer/edge_function.zig b/src/soft/device/rasterizer/edge_function.zig index fc23e40..b7f2d10 100644 --- a/src/soft/device/rasterizer/edge_function.zig +++ b/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| {