adding vkGetBufferDeviceAddress

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

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",
+            .url = "git+https://github.com/Kbz-8/SPIRV-Interpreter#cba8a4723d9b21c9c58494be8ce477124946f52a",
-            .hash = "SPIRV_Interpreter-0.0.1-ajmpn9_vBwCUxGCNVFXUgZXaDv0OyNz8sqhzuyFD7hTU",
+            .hash = "SPIRV_Interpreter-0.0.1-ajmpnyQ4CABQ7B36QdzikdLDH1E5LMk-ib8O2gAHgwbE",
             .lazy = true,
         },
         //.SPIRV_Interpreter = .{

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 setupWorkgroupBuiltins(data.self, rt, data.local_size, group_count, group_id);
-        try setupSubgroupBuiltins(data.self, rt, group_id, i);
+        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 {
+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 {
-    const spv_module = &self.state.pipeline.?.stages.getPtrAssertContains(.compute).module.module;
+    rt.writeBuiltIn(self.device.device_allocator.allocator(), std.mem.asBytes(&local_size), .WorkgroupSize) catch |err| switch (err) {
-    const workgroup_size = @Vector(3, u32){
+        SpvRuntimeError.NotFound => {},
-        spv_module.reflection_infos.local_size_x,
+        else => return err,
-        spv_module.reflection_infos.local_size_y,
+    };
-        spv_module.reflection_infos.local_size_z,
+    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 {
+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 spv_module = &self.state.pipeline.?.stages.getPtrAssertContains(.compute).module.module;
+    const local_base = local_size * group_id;
-    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;
     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]);
+    local_invocation[2] = @divTrunc(idx, local_size[0] * local_size[1]);
-    idx -= local_invocation[2] * workgroup_size[0] * workgroup_size[1];
+    idx -= local_invocation[2] * local_size[0] * local_size[1];
-    local_invocation[1] = @divTrunc(idx, workgroup_size[0]);
+    local_invocation[1] = @divTrunc(idx, local_size[0]);
-    idx -= local_invocation[1] * workgroup_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(self.device.device_allocator.allocator(), std.mem.asBytes(&local_invocation), .LocalInvocationId) catch |err| switch (err) {
-    rt.writeBuiltIn(std.mem.asBytes(&local_invocation_index_u32), .LocalInvocationIndex) catch {};
+        SpvRuntimeError.NotFound => {},
-    rt.writeBuiltIn(std.mem.asBytes(&global_invocation_index), .GlobalInvocationId) catch {};
+        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,
+    };
 }

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(
+                        writeDescriptorSet(
-                            map,
+                            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(
+                        writeDescriptorSet(
-                            std.mem.asBytes(&addr),
+                            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(
+                        writeDescriptorSet(
-                            std.mem.asBytes(&addr),
+                            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(
+                        writeDescriptorSet(
-                            std.mem.asBytes(&addr),
+                            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 {
+                    var data: SampledImageDescriptor = .{
-                        image: usize,
-                        sampler: usize,
-                    };
-
-                    var data: SampledImage = .{
                         .image = 0,
                         .sampler = 0,
                     };
@@ -173,8 +231,9 @@ pub fn writeDescriptorSets(state: *PipelineState, rt: *spv.Runtime) !void {
                         data.sampler = addr;
                     }
 
-                    rt.writeDescriptorSet(
+                    writeDescriptorSet(
-                        std.mem.asBytes(&data),
+                        rt,
+                        .{ .sampled_image = data },
                         @as(u32, @intCast(set_index)),
                         @as(u32, @intCast(binding_index)),
                         @as(u32, @intCast(descriptor_index)),

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];

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(allocator, std.mem.asBytes(&vertex_index_u32), .VertexIndex);
-    try rt.writeBuiltIn(std.mem.asBytes(&instance_index_u32), .InstanceIndex);
+    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;
     }

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));
+}

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();