From b13a256eb6f1e106cd30267ebdb9ede32cf6bad1 Mon Sep 17 00:00:00 2001 From: Kbz-8 Date: Fri, 26 Jun 2026 23:28:32 +0200 Subject: [PATCH] adding vkGetBufferDeviceAddress --- README.md | 3 + build.zig.zon | 6 +- src/software/device/ComputeDispatcher.zig | 82 ++++++++++++-------- src/software/device/Device.zig | 91 +++++++++++++++++++---- src/software/device/fragment.zig | 8 +- src/software/device/vertex_dispatcher.zig | 12 +-- src/vulkan/Buffer.zig | 6 ++ src/vulkan/lib_vulkan.zig | 31 ++++++++ 8 files changed, 180 insertions(+), 59 deletions(-) diff --git a/README.md b/README.md index e4e933d..7b669d9 100644 --- a/README.md +++ b/README.md @@ -162,6 +162,9 @@ vkFlushMappedMemoryRanges | ✅ Implemented vkFreeCommandBuffers | ✅ Implemented vkFreeDescriptorSets | ✅ Implemented vkFreeMemory | ✅ Implemented +vkGetBufferDeviceAddress | ✅ Implemented +vkGetBufferDeviceAddressEXT | ✅ Implemented +vkGetBufferDeviceAddressKHR | ✅ Implemented vkGetBufferMemoryRequirements | ✅ Implemented vkGetDeviceGroupPeerMemoryFeaturesKHR | ✅ Implemented vkGetDeviceGroupPresentCapabilitiesKHR | ✅ Implemented diff --git a/build.zig.zon b/build.zig.zon index 17355f3..50b8a03 100644 --- a/build.zig.zon +++ b/build.zig.zon @@ -26,14 +26,14 @@ .hash = "N-V-__8AAF9uOh0I4P_99za7N822J3JwsDaqONrFVrcEQo59", }, .drm = .{ - .url = "git+https://git.kbz8.me/kbz_8/zig-drm.git#409f58daa8f5174b2fcb8897f1c30f0b0729b611", + .url = "git+https://github.com/Kbz-8/zig-drm#409f58daa8f5174b2fcb8897f1c30f0b0729b611", .hash = "drm-0.0.1-uWWar5YwAQCiTNOcqtYqt_yL19B0_Q-YLjPNkRysFSnR", }, // Soft dependencies .SPIRV_Interpreter = .{ - .url = "git+https://github.com/Kbz-8/SPIRV-Interpreter#ae0914647bbc48683dd10108b5cae9acc54fe759", - .hash = "SPIRV_Interpreter-0.0.1-ajmpn9_vBwCUxGCNVFXUgZXaDv0OyNz8sqhzuyFD7hTU", + .url = "git+https://github.com/Kbz-8/SPIRV-Interpreter#cba8a4723d9b21c9c58494be8ce477124946f52a", + .hash = "SPIRV_Interpreter-0.0.1-ajmpnyQ4CABQ7B36QdzikdLDH1E5LMk-ib8O2gAHgwbE", .lazy = true, }, //.SPIRV_Interpreter = .{ diff --git a/src/software/device/ComputeDispatcher.zig b/src/software/device/ComputeDispatcher.zig index 2656531..3eba655 100644 --- a/src/software/device/ComputeDispatcher.zig +++ b/src/software/device/ComputeDispatcher.zig @@ -26,6 +26,7 @@ const RunData = struct { group_count_y: usize, group_count_z: usize, invocations_per_workgroup: usize, + local_size: @Vector(3, u32), pipeline: *SoftPipeline, }; @@ -53,6 +54,18 @@ pub fn dispatch(self: *Self, group_count_x: u32, group_count_y: u32, group_count try self.dispatchBase(0, 0, 0, group_count_x, group_count_y, group_count_z); } +fn getLocalSize(rt: *spv.Runtime, allocator: std.mem.Allocator, spv_module: *const spv.Module) VkError!@Vector(3, u32) { + if (rt.getWorkgroupSize(allocator) catch return VkError.ValidationFailed) |workgroup_size| { + return workgroup_size; + } + + return .{ + spv_module.reflection_infos.local_size_x, + spv_module.reflection_infos.local_size_y, + spv_module.reflection_infos.local_size_z, + }; +} + pub fn dispatchBase(self: *Self, base_group_x: u32, base_group_y: u32, base_group_z: u32, group_count_x: u32, group_count_y: u32, group_count_z: u32) VkError!void { const group_count_xy = std.math.mul(usize, group_count_x, group_count_y) catch return VkError.ValidationFailed; const group_count = std.math.mul(usize, group_count_xy, group_count_z) catch return VkError.ValidationFailed; @@ -62,7 +75,10 @@ pub fn dispatchBase(self: *Self, base_group_x: u32, base_group_y: u32, base_grou const spv_module = &shader.module.module; self.batch_size = shader.runtimes.len; - 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; + const allocator = self.device.device_allocator.allocator(); + const local_size = try getLocalSize(&shader.runtimes[0].rt, allocator, spv_module); + const local_size_xy = std.math.mul(usize, local_size[0], local_size[1]) catch return VkError.ValidationFailed; + const invocations_per_workgroup = std.math.mul(usize, local_size_xy, local_size[2]) catch return VkError.ValidationFailed; self.invocation_index.store(0, .monotonic); @@ -87,6 +103,7 @@ pub fn dispatchBase(self: *Self, base_group_x: u32, base_group_y: u32, base_grou .group_count_y = @as(usize, @intCast(group_count_y)), .group_count_z = @as(usize, @intCast(group_count_z)), .invocations_per_workgroup = invocations_per_workgroup, + .local_size = local_size, .pipeline = pipeline, }; @@ -170,11 +187,11 @@ inline fn run(data: RunData) !void { for (0..data.invocations_per_workgroup) |i| { rt.resetInvocation(allocator); - try setupWorkgroupBuiltins(data.self, rt, group_count_vec, group_id_vec); + try setupWorkgroupBuiltins(data.self, rt, data.local_size, group_count_vec, group_id_vec); const invocation_index = data.self.invocation_index.fetchAdd(1, .monotonic); - try setupSubgroupBuiltins(data.self, rt, .{ + try setupSubgroupBuiltins(data.self, rt, data.local_size, .{ @as(u32, @intCast(data.base_group_x + group_x)), @as(u32, @intCast(data.base_group_y + group_y)), @as(u32, @intCast(data.base_group_z + group_z)), @@ -216,8 +233,8 @@ fn runBarrierWorkgroup( rt.resetInvocation(allocator); try ExecutionDevice.writeDescriptorSets(data.self.state, rt); try rt.populatePushConstants(data.self.state.push_constant_blob[0..]); - try setupWorkgroupBuiltins(data.self, rt, group_count, group_id); - try setupSubgroupBuiltins(data.self, rt, group_id, i); + try setupWorkgroupBuiltins(data.self, rt, data.local_size, group_count, group_id); + try setupSubgroupBuiltins(data.self, rt, data.local_size, group_id, i); statuses[i] = try rt.beginEntryPoint(allocator, entry); try rt.flushDescriptorSets(allocator); } @@ -255,40 +272,45 @@ fn dumpResultsTable(allocator: std.mem.Allocator, io: std.Io, rt: *spv.Runtime, 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 { - const spv_module = &self.state.pipeline.?.stages.getPtrAssertContains(.compute).module.module; - const workgroup_size = @Vector(3, u32){ - spv_module.reflection_infos.local_size_x, - spv_module.reflection_infos.local_size_y, - spv_module.reflection_infos.local_size_z, +fn setupWorkgroupBuiltins(self: *Self, rt: *spv.Runtime, local_size: @Vector(3, u32), group_count: @Vector(3, u32), group_id: @Vector(3, u32)) spv.Runtime.RuntimeError!void { + rt.writeBuiltIn(self.device.device_allocator.allocator(), std.mem.asBytes(&local_size), .WorkgroupSize) catch |err| switch (err) { + SpvRuntimeError.NotFound => {}, + else => return err, + }; + rt.writeBuiltIn(self.device.device_allocator.allocator(), std.mem.asBytes(&group_count), .NumWorkgroups) catch |err| switch (err) { + SpvRuntimeError.NotFound => {}, + else => return err, + }; + rt.writeBuiltIn(self.device.device_allocator.allocator(), std.mem.asBytes(&group_id), .WorkgroupId) catch |err| switch (err) { + SpvRuntimeError.NotFound => {}, + else => return err, }; - - rt.writeBuiltIn(std.mem.asBytes(&workgroup_size), .WorkgroupSize) catch {}; - rt.writeBuiltIn(std.mem.asBytes(&group_count), .NumWorkgroups) catch {}; - 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 { - const spv_module = &self.state.pipeline.?.stages.getPtrAssertContains(.compute).module.module; - const workgroup_size = @Vector(3, u32){ - 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; +fn setupSubgroupBuiltins(self: *Self, rt: *spv.Runtime, local_size: @Vector(3, u32), group_id: @Vector(3, u32), local_invocation_index: usize) spv.Runtime.RuntimeError!void { + const local_base = local_size * group_id; var local_invocation = @Vector(3, u32){ 0, 0, 0 }; var idx: u32 = @intCast(local_invocation_index); - local_invocation[2] = @divTrunc(idx, workgroup_size[0] * workgroup_size[1]); - idx -= local_invocation[2] * workgroup_size[0] * workgroup_size[1]; - local_invocation[1] = @divTrunc(idx, workgroup_size[0]); - idx -= local_invocation[1] * workgroup_size[0]; + local_invocation[2] = @divTrunc(idx, local_size[0] * local_size[1]); + idx -= local_invocation[2] * local_size[0] * local_size[1]; + local_invocation[1] = @divTrunc(idx, local_size[0]); + idx -= local_invocation[1] * local_size[0]; local_invocation[0] = idx; const global_invocation_index = local_base + local_invocation; const local_invocation_index_u32: u32 = @intCast(local_invocation_index); - rt.writeBuiltIn(std.mem.asBytes(&local_invocation), .LocalInvocationId) catch {}; - rt.writeBuiltIn(std.mem.asBytes(&local_invocation_index_u32), .LocalInvocationIndex) catch {}; - rt.writeBuiltIn(std.mem.asBytes(&global_invocation_index), .GlobalInvocationId) catch {}; + rt.writeBuiltIn(self.device.device_allocator.allocator(), std.mem.asBytes(&local_invocation), .LocalInvocationId) catch |err| switch (err) { + SpvRuntimeError.NotFound => {}, + else => return err, + }; + rt.writeBuiltIn(self.device.device_allocator.allocator(), std.mem.asBytes(&local_invocation_index_u32), .LocalInvocationIndex) catch |err| switch (err) { + SpvRuntimeError.NotFound => {}, + else => return err, + }; + rt.writeBuiltIn(self.device.device_allocator.allocator(), std.mem.asBytes(&global_invocation_index), .GlobalInvocationId) catch |err| switch (err) { + SpvRuntimeError.NotFound => {}, + else => return err, + }; } diff --git a/src/software/device/Device.zig b/src/software/device/Device.zig index bcedd9b..641e9f0 100644 --- a/src/software/device/Device.zig +++ b/src/software/device/Device.zig @@ -76,6 +76,65 @@ pub fn deinit(self: *Self, allocator: std.mem.Allocator) void { self.active_occlusion_queries.deinit(allocator); } +const SampledImageDescriptor = struct { + image: usize, + sampler: usize, +}; + +const DescriptorPayload = union(enum) { + raw: []const u8, + sampled_image: SampledImageDescriptor, +}; + +fn writeDescriptorValue(value: anytype, payload: DescriptorPayload, descriptor_index: u32) spv.Runtime.RuntimeError!void { + const dst = switch (value.*) { + .Array => |arr| blk: { + if (descriptor_index >= arr.values.len) + return spv.Runtime.RuntimeError.NotFound; + break :blk &arr.values[descriptor_index]; + }, + else => blk: { + if (descriptor_index != 0) + return spv.Runtime.RuntimeError.NotFound; + break :blk value; + }, + }; + + switch (payload) { + .raw => |data| _ = try dst.write(data), + .sampled_image => |data| switch (dst.*) { + .Image => _ = try dst.write(std.mem.asBytes(&data.image)), + .Sampler => _ = try dst.write(std.mem.asBytes(&data.sampler)), + .SampledImage => _ = try dst.write(std.mem.asBytes(&data)), + else => return spv.Runtime.RuntimeError.InvalidValueType, + }, + } +} + +fn writeDescriptorSet(rt: *spv.Runtime, payload: DescriptorPayload, set: u32, binding: u32, descriptor_index: u32) spv.Runtime.RuntimeError!void { + var found = false; + for (rt.mod.bindings.items) |entry| { + if (entry.set != set or entry.binding != binding) + continue; + + const variable = switch (rt.results[entry.result].variant orelse continue) { + .Variable => |*variable| variable, + else => continue, + }; + + writeDescriptorValue(&variable.value, payload, descriptor_index) catch |err| switch (err) { + spv.Runtime.RuntimeError.InvalidValueType, + spv.Runtime.RuntimeError.OutOfBounds, + => continue, + else => return err, + }; + found = true; + } + + if (!found) + return spv.Runtime.RuntimeError.NotFound; +} + pub fn writeDescriptorSets(state: *PipelineState, rt: *spv.Runtime) !void { sets: for (state.sets[0..], 0..) |set, set_index| { if (set == null) @@ -95,8 +154,9 @@ pub fn writeDescriptorSets(state: *PipelineState, rt: *spv.Runtime) !void { }; const map = buffer.mapAsSliceWithAddedOffset(u8, buffer_data.offset + dynamic_offset, buffer_data.size) catch continue :bindings; - rt.writeDescriptorSet( - map, + writeDescriptorSet( + rt, + .{ .raw = map }, @as(u32, @intCast(set_index)), @as(u32, @intCast(binding_index)), @as(u32, @intCast(descriptor_index)), @@ -110,8 +170,9 @@ pub fn writeDescriptorSets(state: *PipelineState, rt: *spv.Runtime) !void { .image => |image_data_array| for (image_data_array, 0..) |image_data, descriptor_index| { if (image_data.object) |image_view| { const addr: usize = @intFromPtr(image_view); - rt.writeDescriptorSet( - std.mem.asBytes(&addr), + writeDescriptorSet( + rt, + .{ .raw = std.mem.asBytes(&addr) }, @as(u32, @intCast(set_index)), @as(u32, @intCast(binding_index)), @as(u32, @intCast(descriptor_index)), @@ -125,8 +186,9 @@ pub fn writeDescriptorSets(state: *PipelineState, rt: *spv.Runtime) !void { .sampler => |sampler_data_array| for (sampler_data_array, 0..) |sampler_data, descriptor_index| { if (sampler_data.object) |sampler| { const addr: usize = @intFromPtr(sampler); - rt.writeDescriptorSet( - std.mem.asBytes(&addr), + writeDescriptorSet( + rt, + .{ .raw = std.mem.asBytes(&addr) }, @as(u32, @intCast(set_index)), @as(u32, @intCast(binding_index)), @as(u32, @intCast(descriptor_index)), @@ -140,8 +202,9 @@ pub fn writeDescriptorSets(state: *PipelineState, rt: *spv.Runtime) !void { .texel_buffer => |texel_data_array| for (texel_data_array, 0..) |texel_data, descriptor_index| { if (texel_data.object) |buffer_view| { const addr: usize = @intFromPtr(buffer_view); - rt.writeDescriptorSet( - std.mem.asBytes(&addr), + writeDescriptorSet( + rt, + .{ .raw = std.mem.asBytes(&addr) }, @as(u32, @intCast(set_index)), @as(u32, @intCast(binding_index)), @as(u32, @intCast(descriptor_index)), @@ -153,12 +216,7 @@ pub fn writeDescriptorSets(state: *PipelineState, rt: *spv.Runtime) !void { }, .texture => |texture_data_array| for (texture_data_array, 0..) |texture_data, descriptor_index| { - const SampledImage = packed struct { - image: usize, - sampler: usize, - }; - - var data: SampledImage = .{ + var data: SampledImageDescriptor = .{ .image = 0, .sampler = 0, }; @@ -173,8 +231,9 @@ pub fn writeDescriptorSets(state: *PipelineState, rt: *spv.Runtime) !void { data.sampler = addr; } - rt.writeDescriptorSet( - std.mem.asBytes(&data), + writeDescriptorSet( + rt, + .{ .sampled_image = data }, @as(u32, @intCast(set_index)), @as(u32, @intCast(binding_index)), @as(u32, @intCast(descriptor_index)), diff --git a/src/software/device/fragment.zig b/src/software/device/fragment.zig index d8e4c40..aba8f73 100644 --- a/src/software/device/fragment.zig +++ b/src/software/device/fragment.zig @@ -57,17 +57,17 @@ pub fn shaderInvocation( rt.resetInvocation(allocator); try rt.populatePushConstants(draw_call.renderer.state.push_constant_blob[0..]); - rt.writeBuiltIn(std.mem.asBytes(&position), .FragCoord) catch |err| switch (err) { + rt.writeBuiltIn(allocator, std.mem.asBytes(&position), .FragCoord) catch |err| switch (err) { SpvRuntimeError.NotFound => {}, else => return err, }; if (point_coord) |coord| { - rt.writeBuiltIn(std.mem.asBytes(&coord), .PointCoord) catch |err| switch (err) { + rt.writeBuiltIn(allocator, std.mem.asBytes(&coord), .PointCoord) catch |err| switch (err) { SpvRuntimeError.NotFound => {}, else => return err, }; } - rt.writeBuiltIn(std.mem.asBytes(&front_face), .FrontFacing) catch |err| switch (err) { + rt.writeBuiltIn(allocator, std.mem.asBytes(&front_face), .FrontFacing) catch |err| switch (err) { SpvRuntimeError.NotFound => {}, else => return err, }; @@ -104,7 +104,7 @@ pub fn shaderInvocation( } if (input.blob.len != 0) { - try rt.writeInput(input.blob, result_word); + try rt.writeInput(allocator, input.blob, result_word); if (derivatives) |derivative| { const dx = derivative.dx[location][component]; const dy = derivative.dy[location][component]; diff --git a/src/software/device/vertex_dispatcher.zig b/src/software/device/vertex_dispatcher.zig index c8f2a68..24c415e 100644 --- a/src/software/device/vertex_dispatcher.zig +++ b/src/software/device/vertex_dispatcher.zig @@ -68,7 +68,7 @@ inline fn run(data: RunData) !void { const vertex_index: usize = vertex_index_u32; const instance_index = data.first_instance + data.instance_index; - setupBuiltins(rt, vertex_index_u32, instance_index) catch |err| switch (err) { + setupBuiltins(rt, data.allocator, vertex_index_u32, instance_index) catch |err| switch (err) { SpvRuntimeError.NotFound => {}, else => return err, }; @@ -212,11 +212,11 @@ fn isConstantZero(rt: *spv.Runtime, result_word: spv.SpvWord) bool { } } -fn setupBuiltins(rt: *spv.Runtime, vertex_index_u32: u32, instance_index: usize) !void { +fn setupBuiltins(rt: *spv.Runtime, allocator: std.mem.Allocator, vertex_index_u32: u32, instance_index: usize) !void { const instance_index_u32: u32 = @intCast(instance_index); - try rt.writeBuiltIn(std.mem.asBytes(&vertex_index_u32), .VertexIndex); - try rt.writeBuiltIn(std.mem.asBytes(&instance_index_u32), .InstanceIndex); + try rt.writeBuiltIn(allocator, std.mem.asBytes(&vertex_index_u32), .VertexIndex); + try rt.writeBuiltIn(allocator, std.mem.asBytes(&instance_index_u32), .InstanceIndex); } fn writeVertexInput(rt: *spv.Runtime, allocator: std.mem.Allocator, raw_input: []const u8, format: vk.Format, location: u32) !void { @@ -243,7 +243,7 @@ fn writeVertexInput(rt: *spv.Runtime, allocator: std.mem.Allocator, raw_input: [ const raw_offset = component * @sizeOf(f32); if (raw_offset + input_memory_size <= expanded_slice.len) { - try rt.writeInput(expanded_slice[raw_offset .. raw_offset + input_memory_size], result_word); + try rt.writeInput(allocator, expanded_slice[raw_offset .. raw_offset + input_memory_size], result_word); continue; } @@ -266,7 +266,7 @@ fn writeVertexInput(rt: *spv.Runtime, allocator: std.mem.Allocator, raw_input: [ } } - try rt.writeInput(input, result_word); + try rt.writeInput(allocator, input, result_word); } return; } diff --git a/src/vulkan/Buffer.zig b/src/vulkan/Buffer.zig index 4fd7d07..d2e4d1d 100644 --- a/src/vulkan/Buffer.zig +++ b/src/vulkan/Buffer.zig @@ -52,3 +52,9 @@ pub inline fn getMemoryRequirements(self: *Self, requirements: *vk.MemoryRequire requirements.memory_type_bits = self.allowed_memory_types.mask; self.vtable.getMemoryRequirements(self, requirements); } + +pub inline fn getDeviceAddress(self: *Self) VkError!vk.DeviceAddress { + const memory = self.memory orelse return 0; + const map = try memory.map(self.offset, self.size); + return @intCast(@intFromPtr(map.ptr)); +} diff --git a/src/vulkan/lib_vulkan.zig b/src/vulkan/lib_vulkan.zig index 5ee295e..a699f67 100644 --- a/src/vulkan/lib_vulkan.zig +++ b/src/vulkan/lib_vulkan.zig @@ -239,6 +239,9 @@ const device_pfn_map = block: { functionMapEntryPoint("vkFreeDescriptorSets"), functionMapEntryPoint("vkFreeMemory"), functionMapEntryPoint("vkGetBufferMemoryRequirements"), + functionMapEntryPoint("vkGetBufferDeviceAddress"), + functionMapEntryPoint("vkGetBufferDeviceAddressEXT"), + functionMapEntryPoint("vkGetBufferDeviceAddressKHR"), functionMapEntryPoint("vkGetDeviceMemoryCommitment"), functionMapEntryPoint("vkGetDeviceGroupPeerMemoryFeatures"), functionMapEntryPoint("vkGetDeviceGroupPeerMemoryFeaturesKHR"), @@ -1398,6 +1401,34 @@ pub export fn apeGetBufferMemoryRequirements(p_device: vk.Device, p_buffer: vk.B buffer.getMemoryRequirements(requirements); } +pub export fn apeGetBufferDeviceAddress(p_device: vk.Device, info: *const vk.BufferDeviceAddressInfo) callconv(vk.vulkan_call_conv) vk.DeviceAddress { + entryPointBeginLogTrace(.vkGetBufferDeviceAddress); + defer entryPointEndLogTrace(); + + Dispatchable(Device).checkHandleValidity(p_device) catch |err| { + errorLogger(err); + return 0; + }; + + const buffer = NonDispatchable(Buffer).fromHandleObject(info.buffer) catch |err| { + errorLogger(err); + return 0; + }; + + return buffer.getDeviceAddress() catch |err| { + errorLogger(err); + return 0; + }; +} + +pub export fn apeGetBufferDeviceAddressEXT(p_device: vk.Device, info: *const vk.BufferDeviceAddressInfo) callconv(vk.vulkan_call_conv) vk.DeviceAddress { + return @call(.always_inline, apeGetBufferDeviceAddress, .{ p_device, info }); +} + +pub export fn apeGetBufferDeviceAddressKHR(p_device: vk.Device, info: *const vk.BufferDeviceAddressInfo) callconv(vk.vulkan_call_conv) vk.DeviceAddress { + return @call(.always_inline, apeGetBufferDeviceAddress, .{ p_device, info }); +} + pub export fn apeGetDeviceMemoryCommitment(p_device: vk.Device, p_memory: vk.DeviceMemory, committed_memory: *vk.DeviceSize) callconv(vk.vulkan_call_conv) void { entryPointBeginLogTrace(.vkGetDeviceMemoryCommitment); defer entryPointEndLogTrace();