removing unsupported panic

.gitea/workflows/1.4-CTS.yml

@@ -17,7 +17,7 @@ jobs:
     - name: Install system dependencies
       run: |
         apt update
-        apt install -y libgl1 libwayland-egl1 libwayland-cursor0
+        apt install -y libgl1 libwayland-egl1 libwayland-cursor0 clang libwayland-server0
 
     - name: Building
       run: zig build
@@ -40,7 +40,7 @@ jobs:
       shell: bash
       run: |
         source $HOME/.cargo/env
-        cargo binstall dioxus-cli --no-confirm --force --version 0.7.2
+        cargo binstall dioxus-cli --no-confirm --force --version 0.7.3
         echo "$HOME/.cargo/bin" >> $GITHUB_PATH
 
     - name: Install deqp-runner
@@ -55,7 +55,7 @@ jobs:
         which deqp-runner && deqp-runner --version || echo "deqp-runner not found"
 
     - name: Run Vulkan CTS
-      run: zig build cts-soft -- --mustpass-list=master
+      run: zig build cts-soft -- --mustpass-list=master -j4
       continue-on-error: true
 
     - name: Verify tests

.gitea/workflows/Test.yml

@@ -20,7 +20,7 @@ jobs:
     - name: Install system dependencies
       run: |
         apt update
-        apt install -y libgl1 libwayland-egl1 libwayland-cursor0
+        apt install -y libgl1 libwayland-egl1 libwayland-cursor0 clang libwayland-server0
 
     - name: Building
       run: zig build
@@ -46,7 +46,7 @@ jobs:
       shell: bash
       run: |
         source $HOME/.cargo/env
-        cargo binstall dioxus-cli --no-confirm --force --version 0.7.2
+        cargo binstall dioxus-cli --no-confirm --force --version 0.7.3
         echo "$HOME/.cargo/bin" >> $GITHUB_PATH
 
     - name: Install deqp-runner
@@ -61,7 +61,7 @@ jobs:
         which deqp-runner && deqp-runner --version || echo "deqp-runner not found"
 
     - name: Run Vulkan CTS
-      run: zig build cts-soft
+      run: zig build cts-soft -- -j4
       continue-on-error: true
 
     - name: Verify tests

.gitignore

@@ -13,3 +13,4 @@ scripts/__pycache__/
 *.xml
 *.html
 *.pyc
+*.spv

README.md

@@ -22,7 +22,7 @@ zig build
 ```
 
 Then ensure thy Vulkan loader is pointed toward the ICD manifest.
-The precise ritual varies by system — consult the tomes of your operating system, or wander the web’s endless mausoleum of documentation.
+The precise ritual varies by system — consult the tomes of your operating system, or wander the web's endless mausoleum of documentation.
 
 Use at your own risk. If thy machine shudders, weeps, or attempts to flee — know that it was warned.
 \

build.zig

@@ -32,6 +32,7 @@ pub fn build(b: *std.Build) !void {
 
     const zdt = b.dependency("zdt", .{}).module("zdt");
     const zigrc = b.dependency("zigrc", .{}).module("zigrc");
+    //const spv_tools = b.dependency("SPIRV_Tools", .{}).module("zigrc");
     const vulkan_headers = b.dependency("vulkan_headers", .{});
     const vulkan_utility_libraries = b.dependency("vulkan_utility_libraries", .{});
 
@@ -89,6 +90,8 @@ pub fn build(b: *std.Build) !void {
 
         lib.step.dependOn(&icd_file.step);
         const lib_install = b.addInstallArtifact(lib, .{});
+        const install_step = b.step(impl.name, b.fmt("Build libvulkan_{s}", .{impl.name}));
+        install_step.dependOn(&lib_install.step);
 
         const lib_tests = b.addTest(.{ .root_module = lib_mod });
 
@@ -131,6 +134,9 @@ fn customSoft(b: *std.Build, lib: *std.Build.Step.Compile) !void {
     lib.addSystemIncludePath(cpuinfo.path("include"));
     lib.linkLibrary(cpuinfo.artifact("cpuinfo"));
 
+    const interface = b.lazyDependency("interface", .{}) orelse return error.UnresolvedDependency;
+    lib.root_module.addImport("interface", interface.module("interface"));
+
     const spv = b.dependency("SPIRV_Interpreter", .{
         .@"no-example" = true,
         .@"no-test" = true,
@@ -254,16 +260,18 @@ fn addMultithreadedCTS(b: *std.Build, target: std.Build.ResolvedTarget, impl: *c
         },
     }));
 
-    const mustpass_override = blk: {
+    var mustpass_override: ?[]const u8 = null;
+    var jobs_count: ?usize = null;
+
     if (b.args) |args| {
         for (args) |arg| {
             if (std.mem.startsWith(u8, arg, "--mustpass-list")) {
-                    break :blk arg["--mustpass-list=".len..];
+                mustpass_override = arg["--mustpass-list=".len..];
+            } else if (std.mem.startsWith(u8, arg, "-j")) {
+                jobs_count = try std.fmt.parseInt(usize, arg["-j".len..], 10);
             }
         }
     }
-        break :blk null;
-    };
 
     const mustpass_path = try cts.path(
         if (mustpass_override) |override|
@@ -286,6 +294,9 @@ fn addMultithreadedCTS(b: *std.Build, target: std.Build.ResolvedTarget, impl: *c
     run.addArg(b.fmt("{s}{s}", .{ cache_path, mustpass_path }));
     run.addArg("--output");
     run.addArg("./cts");
+    if (jobs_count) |count| {
+        run.addArg(b.fmt("-j{d}", .{count}));
+    }
     run.addArg("--");
     run.addArg(b.fmt("--deqp-archive-dir={s}{s}", .{ cache_path, try cts.path("").getPath3(b, null).toString(b.allocator) }));
     run.addArg(b.fmt("--deqp-vk-library-path={s}", .{b.getInstallPath(.lib, impl_lib.out_lib_filename)}));

build.zig.zon

@@ -33,11 +33,6 @@
             .url = "https://github.com/Aandreba/zigrc/archive/refs/tags/1.1.0.tar.gz",
             .hash = "zigrc-1.0.0-lENlWzvQAACulrbkL9PVhWjFsWSkYhi7AmfSbCM-2Xlh",
         },
-        .SPIRV_Interpreter = .{
-            .url = "git+https://git.kbz8.me/kbz_8/SPIRV-Interpreter#e21d26d9975b96a222b70648ceeea9e473e9657f",
-            .hash = "SPIRV_Interpreter-0.0.1-ajmpn2tAAwCBI0oWa3VKlYX3MEM0OxN4iXQ-PwO6_Vhx",
-            //.path = "../SPIRV-Interpreter",
-        },
         .cpuinfo = .{
             .url = "git+https://github.com/Kbz-8/cpuinfo#4883954cfcec3f6c9ca9c4aaddfc26107e08726f",
             .hash = "cpuinfo-0.0.1-RLgIQTLRMgF4dLo8AJ-HvnpFsJe6jmXCJjMWWjil6RF1",
@@ -58,6 +53,15 @@
             .hash = "N-V-__8AABQ7TgCnPlp8MP4YA8znrjd6E-ZjpF1rvrS8J_2I",
             .lazy = true,
         },
+        .interface = .{
+            .url = "git+https://github.com/nilslice/zig-interface#8c0fe8fa9fd0702eee43f50cb75dce1cc5a7e1f4",
+            .hash = "interface-0.0.2-GFlWJ1mcAQARS-V4xJ7qDt5_cutxOHSEz6H9yiK-Sw0A",
+            .lazy = true,
+        },
+        .SPIRV_Interpreter = .{
+            .url = "git+https://git.kbz8.me/kbz_8/SPIRV-Interpreter#e09a41754ffa115ba1668f82698140b2b727b7fd",
+            .hash = "SPIRV_Interpreter-0.0.1-ajmpn5eJAwCe-SNkaCifjF0UUqE6cww18VpqO0Qip_mp",
+        },
     },
 
     .paths = .{

src/soft/SoftBuffer.zig

@@ -46,3 +46,35 @@ pub fn getMemoryRequirements(interface: *Interface, requirements: *vk.MemoryRequ
         requirements.alignment = @max(requirements.alignment, lib.MIN_UNIFORM_BUFFER_ALIGNMENT);
     }
 }
+
+pub fn copyBuffer(self: *const Self, dst: *Self, regions: []const vk.BufferCopy) VkError!void {
+    for (regions) |region| {
+        const src_memory = if (self.interface.memory) |memory| memory else return VkError.InvalidDeviceMemoryDrv;
+        const dst_memory = if (dst.interface.memory) |memory| memory else return VkError.InvalidDeviceMemoryDrv;
+
+        const src_map: []u8 = @as([*]u8, @ptrCast(try src_memory.map(self.interface.offset + region.src_offset, region.size)))[0..region.size];
+        const dst_map: []u8 = @as([*]u8, @ptrCast(try dst_memory.map(dst.interface.offset + region.dst_offset, region.size)))[0..region.size];
+
+        @memcpy(dst_map, src_map);
+
+        src_memory.unmap();
+        dst_memory.unmap();
+    }
+}
+
+pub fn fillBuffer(self: *Self, offset: vk.DeviceSize, size: vk.DeviceSize, data: u32) VkError!void {
+    const memory = if (self.interface.memory) |memory| memory else return VkError.InvalidDeviceMemoryDrv;
+    var memory_map: []u32 = @as([*]u32, @ptrCast(@alignCast(try memory.map(offset, size))))[0..size];
+
+    var bytes = if (size == vk.WHOLE_SIZE) memory.size - offset else size;
+
+    var i: usize = 0;
+    while (bytes >= 4) : ({
+        bytes -= 4;
+        i += 1;
+    }) {
+        memory_map[i] = data;
+    }
+
+    memory.unmap();
+}

src/soft/SoftCommandBuffer.zig

@@ -1,15 +1,32 @@
 const std = @import("std");
 const vk = @import("vulkan");
 const base = @import("base");
+const lib = @import("lib.zig");
+
+const InterfaceFactory = @import("interface").Interface;
 
 const VkError = base.VkError;
 const Device = base.Device;
 
+const SoftBuffer = @import("SoftBuffer.zig");
+const SoftImage = @import("SoftImage.zig");
+const SoftPipeline = @import("SoftPipeline.zig");
+const SoftDescriptorSet = @import("SoftDescriptorSet.zig");
+
+const ExecutionDevice = @import("device/Device.zig");
+
 const Self = @This();
 pub const Interface = base.CommandBuffer;
 
+const Command = InterfaceFactory(.{
+    .execute = fn (*ExecutionDevice) VkError!void,
+}, null);
+
 interface: Interface,
 
+command_allocator: std.heap.ArenaAllocator,
+commands: std.ArrayList(Command),
+
 pub fn create(device: *base.Device, allocator: std.mem.Allocator, info: *const vk.CommandBufferAllocateInfo) VkError!*Self {
     const self = allocator.create(Self) catch return VkError.OutOfHostMemory;
     errdefer allocator.destroy(self);
@@ -22,10 +39,13 @@ pub fn create(device: *base.Device, allocator: std.mem.Allocator, info: *const v
 
     interface.dispatch_table = &.{
         .begin = begin,
+        .bindDescriptorSets = bindDescriptorSets,
+        .bindPipeline = bindPipeline,
         .clearColorImage = clearColorImage,
         .copyBuffer = copyBuffer,
         .copyImage = copyImage,
         .copyImageToBuffer = copyImageToBuffer,
+        .dispatch = dispatch,
         .end = end,
         .fillBuffer = fillBuffer,
         .reset = reset,
@@ -36,7 +56,10 @@ pub fn create(device: *base.Device, allocator: std.mem.Allocator, info: *const v
 
     self.* = .{
         .interface = interface,
+        .command_allocator = undefined,
+        .commands = .empty,
     };
+    self.command_allocator = .init(self.interface.host_allocator.allocator());
     return self;
 }
 
@@ -45,10 +68,16 @@ pub fn destroy(interface: *Interface, allocator: std.mem.Allocator) void {
     allocator.destroy(self);
 }
 
-pub fn begin(interface: *Interface, info: *const vk.CommandBufferBeginInfo) VkError!void {
+pub fn execute(self: *Self, device: *ExecutionDevice) VkError!void {
-    // No-op
+    self.interface.submit() catch return;
-    _ = interface;
+    for (self.commands.items) |command| {
-    _ = info;
+        try command.vtable.execute(command.ptr, device);
+    }
+}
+
+pub fn begin(interface: *Interface, _: *const vk.CommandBufferBeginInfo) VkError!void {
+    const self: *Self = @alignCast(@fieldParentPtr("interface", interface));
+    self.command_allocator.deinit();
 }
 
 pub fn end(interface: *Interface) VkError!void {
@@ -56,57 +85,236 @@ pub fn end(interface: *Interface) VkError!void {
     _ = interface;
 }
 
-pub fn reset(interface: *Interface, flags: vk.CommandBufferResetFlags) VkError!void {
+pub fn reset(interface: *Interface, _: vk.CommandBufferResetFlags) VkError!void {
-    // No-op
+    const self: *Self = @alignCast(@fieldParentPtr("interface", interface));
-    _ = interface;
+    const allocator = self.command_allocator.allocator();
-    _ = flags;
+    self.commands.clearAndFree(allocator);
+    if (!self.command_allocator.reset(.{ .retain_with_limit = 16_384 }))
+        return VkError.OutOfHostMemory;
 }
 
 // Commands ====================================================================================================
 
-pub fn clearColorImage(interface: *Interface, image: *base.Image, layout: vk.ImageLayout, color: *const vk.ClearColorValue, range: vk.ImageSubresourceRange) VkError!void {
+pub fn bindDescriptorSets(interface: *Interface, bind_point: vk.PipelineBindPoint, first_set: u32, sets: [base.VULKAN_MAX_DESCRIPTOR_SETS]?*base.DescriptorSet, dynamic_offsets: []const u32) VkError!void {
-    // No-op
+    const self: *Self = @alignCast(@fieldParentPtr("interface", interface));
-    _ = interface;
+    const allocator = self.command_allocator.allocator();
-    _ = image;
+
-    _ = layout;
+    const CommandImpl = struct {
-    _ = color;
+        const Impl = @This();
-    _ = range;
+
+        bind_point: vk.PipelineBindPoint,
+        first_set: u32,
+        sets: [base.VULKAN_MAX_DESCRIPTOR_SETS]?*base.DescriptorSet,
+        dynamic_offsets: []const u32,
+
+        pub fn execute(impl: *const Impl, device: *ExecutionDevice) VkError!void {
+            for (impl.first_set.., impl.sets[0..]) |i, set| {
+                if (set == null)
+                    break;
+                device.pipeline_states[@intCast(@intFromEnum(impl.bind_point))].sets[i] = @alignCast(@fieldParentPtr("interface", set.?));
+            }
+        }
+    };
+
+    const cmd = allocator.create(CommandImpl) catch return VkError.OutOfHostMemory;
+    errdefer allocator.destroy(cmd);
+    cmd.* = .{
+        .bind_point = bind_point,
+        .first_set = first_set,
+        .sets = sets,
+        .dynamic_offsets = dynamic_offsets,
+    };
+    self.commands.append(allocator, Command.from(cmd)) catch return VkError.OutOfHostMemory;
 }
 
-pub fn fillBuffer(interface: *Interface, buffer: *base.Buffer, offset: vk.DeviceSize, size: vk.DeviceSize, data: u32) VkError!void {
+pub fn bindPipeline(interface: *Interface, bind_point: vk.PipelineBindPoint, pipeline: *base.Pipeline) VkError!void {
-    // No-op
+    const self: *Self = @alignCast(@fieldParentPtr("interface", interface));
-    _ = interface;
+    const allocator = self.command_allocator.allocator();
-    _ = buffer;
+
-    _ = offset;
+    const CommandImpl = struct {
-    _ = size;
+        const Impl = @This();
-    _ = data;
+
+        bind_point: vk.PipelineBindPoint,
+        pipeline: *SoftPipeline,
+
+        pub fn execute(impl: *const Impl, device: *ExecutionDevice) VkError!void {
+            device.pipeline_states[@intCast(@intFromEnum(impl.bind_point))].pipeline = impl.pipeline;
+        }
+    };
+
+    const cmd = allocator.create(CommandImpl) catch return VkError.OutOfHostMemory;
+    errdefer allocator.destroy(cmd);
+    cmd.* = .{
+        .bind_point = bind_point,
+        .pipeline = @alignCast(@fieldParentPtr("interface", pipeline)),
+    };
+    self.commands.append(allocator, Command.from(cmd)) catch return VkError.OutOfHostMemory;
+}
+
+pub fn clearColorImage(interface: *Interface, image: *base.Image, layout: vk.ImageLayout, color: *const vk.ClearColorValue, range: vk.ImageSubresourceRange) VkError!void {
+    const self: *Self = @alignCast(@fieldParentPtr("interface", interface));
+    const allocator = self.command_allocator.allocator();
+
+    const CommandImpl = struct {
+        const Impl = @This();
+
+        image: *SoftImage,
+        layout: vk.ImageLayout,
+        clear_color: vk.ClearColorValue,
+        range: vk.ImageSubresourceRange,
+
+        pub fn execute(impl: *const Impl, _: *ExecutionDevice) VkError!void {
+            impl.image.clearRange(impl.clear_color, impl.range);
+        }
+    };
+
+    const cmd = allocator.create(CommandImpl) catch return VkError.OutOfHostMemory;
+    errdefer allocator.destroy(cmd);
+    cmd.* = .{
+        .image = @alignCast(@fieldParentPtr("interface", image)),
+        .layout = layout,
+        .clear_color = color.*,
+        .range = range,
+    };
+    self.commands.append(allocator, Command.from(cmd)) catch return VkError.OutOfHostMemory;
 }
 
 pub fn copyBuffer(interface: *Interface, src: *base.Buffer, dst: *base.Buffer, regions: []const vk.BufferCopy) VkError!void {
-    // No-op
+    const self: *Self = @alignCast(@fieldParentPtr("interface", interface));
-    _ = interface;
+    const allocator = self.command_allocator.allocator();
-    _ = src;
+
-    _ = dst;
+    const CommandImpl = struct {
-    _ = regions;
+        const Impl = @This();
+
+        src: *const SoftBuffer,
+        dst: *SoftBuffer,
+        regions: []const vk.BufferCopy,
+
+        pub fn execute(impl: *const Impl, _: *ExecutionDevice) VkError!void {
+            try impl.src.copyBuffer(impl.dst, impl.regions);
+        }
+    };
+
+    const cmd = allocator.create(CommandImpl) catch return VkError.OutOfHostMemory;
+    errdefer allocator.destroy(cmd);
+    cmd.* = .{
+        .src = @alignCast(@fieldParentPtr("interface", src)),
+        .dst = @alignCast(@fieldParentPtr("interface", dst)),
+        .regions = allocator.dupe(vk.BufferCopy, regions) catch return VkError.OutOfHostMemory, // Will be freed on cmdbuf reset or destroy
+    };
+    self.commands.append(allocator, Command.from(cmd)) catch return VkError.OutOfHostMemory;
 }
 
 pub fn copyImage(interface: *Interface, src: *base.Image, src_layout: vk.ImageLayout, dst: *base.Image, dst_layout: vk.ImageLayout, regions: []const vk.ImageCopy) VkError!void {
-    // No-op
+    const self: *Self = @alignCast(@fieldParentPtr("interface", interface));
-    _ = interface;
+    const allocator = self.command_allocator.allocator();
-    _ = src;
+
-    _ = src_layout;
+    const CommandImpl = struct {
-    _ = dst;
+        const Impl = @This();
-    _ = dst_layout;
+
-    _ = regions;
+        src: *const SoftImage,
+        src_layout: vk.ImageLayout,
+        dst: *SoftImage,
+        dst_layout: vk.ImageLayout,
+        regions: []const vk.ImageCopy,
+
+        pub fn execute(impl: *const Impl, _: *ExecutionDevice) VkError!void {
+            try impl.src.copyImage(impl.src_layout, impl.dst, impl.dst_layout, impl.regions);
+        }
+    };
+
+    const cmd = allocator.create(CommandImpl) catch return VkError.OutOfHostMemory;
+    errdefer allocator.destroy(cmd);
+    cmd.* = .{
+        .src = @alignCast(@fieldParentPtr("interface", src)),
+        .src_layout = src_layout,
+        .dst = @alignCast(@fieldParentPtr("interface", dst)),
+        .dst_layout = dst_layout,
+        .regions = allocator.dupe(vk.ImageCopy, regions) catch return VkError.OutOfHostMemory, // Will be freed on cmdbuf reset or destroy
+    };
+    self.commands.append(allocator, Command.from(cmd)) catch return VkError.OutOfHostMemory;
 }
 
 pub fn copyImageToBuffer(interface: *Interface, src: *base.Image, src_layout: vk.ImageLayout, dst: *base.Buffer, regions: []const vk.BufferImageCopy) VkError!void {
-    // No-op
+    const self: *Self = @alignCast(@fieldParentPtr("interface", interface));
-    _ = interface;
+    const allocator = self.command_allocator.allocator();
-    _ = src;
+
-    _ = src_layout;
+    const CommandImpl = struct {
-    _ = dst;
+        const Impl = @This();
-    _ = regions;
+
+        src: *const SoftImage,
+        src_layout: vk.ImageLayout,
+        dst: *SoftBuffer,
+        regions: []const vk.BufferImageCopy,
+
+        pub fn execute(impl: *const Impl, _: *ExecutionDevice) VkError!void {
+            try impl.src.copyImageToBuffer(impl.src_layout, impl.dst, impl.regions);
+        }
+    };
+
+    const cmd = allocator.create(CommandImpl) catch return VkError.OutOfHostMemory;
+    errdefer allocator.destroy(cmd);
+    cmd.* = .{
+        .src = @alignCast(@fieldParentPtr("interface", src)),
+        .src_layout = src_layout,
+        .dst = @alignCast(@fieldParentPtr("interface", dst)),
+        .regions = allocator.dupe(vk.BufferImageCopy, regions) catch return VkError.OutOfHostMemory, // Will be freed on cmdbuf reset or destroy
+    };
+    self.commands.append(allocator, Command.from(cmd)) catch return VkError.OutOfHostMemory;
+}
+
+pub fn dispatch(interface: *Interface, group_count_x: u32, group_count_y: u32, group_count_z: u32) VkError!void {
+    const self: *Self = @alignCast(@fieldParentPtr("interface", interface));
+    const allocator = self.command_allocator.allocator();
+
+    const CommandImpl = struct {
+        const Impl = @This();
+
+        group_count_x: u32,
+        group_count_y: u32,
+        group_count_z: u32,
+
+        pub fn execute(impl: *const Impl, device: *ExecutionDevice) VkError!void {
+            try device.compute_routines.dispatch(impl.group_count_x, impl.group_count_y, impl.group_count_z);
+        }
+    };
+
+    const cmd = allocator.create(CommandImpl) catch return VkError.OutOfHostMemory;
+    errdefer allocator.destroy(cmd);
+    cmd.* = .{
+        .group_count_x = group_count_x,
+        .group_count_y = group_count_y,
+        .group_count_z = group_count_z,
+    };
+    self.commands.append(allocator, Command.from(cmd)) catch return VkError.OutOfHostMemory;
+}
+
+pub fn fillBuffer(interface: *Interface, buffer: *base.Buffer, offset: vk.DeviceSize, size: vk.DeviceSize, data: u32) VkError!void {
+    const self: *Self = @alignCast(@fieldParentPtr("interface", interface));
+    const allocator = self.command_allocator.allocator();
+
+    const CommandImpl = struct {
+        const Impl = @This();
+
+        buffer: *SoftBuffer,
+        offset: vk.DeviceSize,
+        size: vk.DeviceSize,
+        data: u32,
+
+        pub fn execute(impl: *const Impl, _: *ExecutionDevice) VkError!void {
+            try impl.buffer.fillBuffer(impl.offset, impl.size, impl.data);
+        }
+    };
+
+    const cmd = allocator.create(CommandImpl) catch return VkError.OutOfHostMemory;
+    errdefer allocator.destroy(cmd);
+    cmd.* = .{
+        .buffer = @alignCast(@fieldParentPtr("interface", buffer)),
+        .offset = offset,
+        .size = size,
+        .data = data,
+    };
+    self.commands.append(allocator, Command.from(cmd)) catch return VkError.OutOfHostMemory;
 }
 
 pub fn resetEvent(interface: *Interface, event: *base.Event, stage: vk.PipelineStageFlags) VkError!void {

src/soft/SoftDescriptorSet.zig

@@ -4,12 +4,28 @@ const base = @import("base");
 
 const VkError = base.VkError;
 const Device = base.Device;
+const Buffer = base.Buffer;
+
+const SoftBuffer = @import("SoftBuffer.zig");
+
+const NonDispatchable = base.NonDispatchable;
 
 const Self = @This();
 pub const Interface = base.DescriptorSet;
 
+const Descriptor = union(enum) {
+    buffer: struct {
+        object: ?*SoftBuffer,
+        offset: vk.DeviceSize,
+        size: vk.DeviceSize,
+    },
+    image: struct {},
+};
+
 interface: Interface,
 
+descriptors: []Descriptor,
+
 pub fn create(device: *base.Device, allocator: std.mem.Allocator, layout: *base.DescriptorSetLayout) VkError!*Self {
     const self = allocator.create(Self) catch return VkError.OutOfHostMemory;
     errdefer allocator.destroy(self);
@@ -17,16 +33,56 @@ pub fn create(device: *base.Device, allocator: std.mem.Allocator, layout: *base.
     var interface = try Interface.init(device, allocator, layout);
 
     interface.vtable = &.{
+        .copy = copy,
         .destroy = destroy,
+        .write = write,
     };
 
+    const descriptors = allocator.alloc(Descriptor, layout.bindings.len) catch return VkError.OutOfHostMemory;
+    errdefer allocator.free(descriptors);
+
     self.* = .{
         .interface = interface,
+        .descriptors = descriptors,
     };
     return self;
 }
 
+pub fn copy(interface: *Interface, copy_data: vk.CopyDescriptorSet) VkError!void {
+    const self: *Self = @alignCast(@fieldParentPtr("interface", interface));
+    _ = self;
+    _ = copy_data;
+}
+
 pub fn destroy(interface: *Interface, allocator: std.mem.Allocator) void {
     const self: *Self = @alignCast(@fieldParentPtr("interface", interface));
+    allocator.free(self.descriptors);
     allocator.destroy(self);
 }
+
+pub fn write(interface: *Interface, write_data: vk.WriteDescriptorSet) VkError!void {
+    const self: *Self = @alignCast(@fieldParentPtr("interface", interface));
+
+    switch (write_data.descriptor_type) {
+        .storage_buffer, .storage_buffer_dynamic => {
+            for (write_data.p_buffer_info, 0..write_data.descriptor_count) |buffer_info, i| {
+                const desc = &self.descriptors[write_data.dst_binding + i];
+                desc.* = .{
+                    .buffer = .{
+                        .object = null,
+                        .offset = buffer_info.offset,
+                        .size = buffer_info.range,
+                    },
+                };
+                if (buffer_info.buffer != .null_handle) {
+                    const buffer = try NonDispatchable(Buffer).fromHandleObject(buffer_info.buffer);
+                    desc.buffer.object = @as(*SoftBuffer, @alignCast(@fieldParentPtr("interface", buffer)));
+                    if (desc.buffer.size == vk.WHOLE_SIZE) {
+                        desc.buffer.size = if (buffer.memory) |memory| memory.size - desc.buffer.offset else return VkError.InvalidDeviceMemoryDrv;
+                    }
+                }
+            }
+        },
+        else => base.unsupported("descriptor type {s} for writting", .{@tagName(write_data.descriptor_type)}),
+    }
+}

src/soft/SoftImage.zig

@@ -7,6 +7,7 @@ const lib = @import("lib.zig");
 const VkError = base.VkError;
 const Device = base.Device;
 
+const SoftBuffer = @import("SoftBuffer.zig");
 const SoftDevice = @import("SoftDevice.zig");
 
 const Self = @This();
@@ -57,3 +58,49 @@ pub fn clearRange(self: *Self, color: vk.ClearColorValue, range: vk.ImageSubreso
         .r32g32b32a32_sfloat;
     self.clear(.{ .color = color }, clear_format, self.interface.format, range, null);
 }
+
+pub fn copyImage(self: *const Self, self_layout: vk.ImageLayout, dst: *Self, dst_layout: vk.ImageLayout, regions: []const vk.ImageCopy) VkError!void {
+    _ = self;
+    _ = self_layout;
+    _ = dst;
+    _ = dst_layout;
+    _ = regions;
+    std.log.scoped(.commandExecutor).warn("FIXME: implement image to image copy", .{});
+}
+
+pub fn copyImageToBuffer(self: *const Self, self_layout: vk.ImageLayout, dst: *SoftBuffer, regions: []const vk.BufferImageCopy) VkError!void {
+    _ = self_layout;
+    for (regions) |region| {
+        const src_memory = if (self.interface.memory) |memory| memory else return VkError.InvalidDeviceMemoryDrv;
+        const dst_memory = if (dst.interface.memory) |memory| memory else return VkError.InvalidDeviceMemoryDrv;
+
+        const pixel_size: u32 = @intCast(self.interface.getPixelSize());
+        const image_row_pitch: u32 = self.interface.extent.width * pixel_size;
+        const image_size: u32 = @intCast(self.interface.getTotalSize());
+
+        const buffer_row_length: u32 = if (region.buffer_row_length != 0) region.buffer_row_length else region.image_extent.width;
+        const buffer_row_pitch: u32 = buffer_row_length * pixel_size;
+        const buffer_size: u32 = buffer_row_pitch * region.image_extent.height * region.image_extent.depth;
+
+        const src_map: []u8 = @as([*]u8, @ptrCast(try src_memory.map(0, image_size)))[0..image_size];
+        const dst_map: []u8 = @as([*]u8, @ptrCast(try dst_memory.map(region.buffer_offset, buffer_size)))[0..buffer_size];
+
+        const row_size = region.image_extent.width * pixel_size;
+        for (0..self.interface.extent.depth) |z| {
+            for (0..self.interface.extent.height) |y| {
+                const z_as_u32: u32 = @intCast(z);
+                const y_as_u32: u32 = @intCast(y);
+
+                const src_offset = ((@as(u32, @intCast(region.image_offset.z)) + z_as_u32) * self.interface.extent.height + @as(u32, @intCast(region.image_offset.y)) + y_as_u32) * image_row_pitch + @as(u32, @intCast(region.image_offset.x)) * pixel_size;
+                const dst_offset = (z_as_u32 * buffer_row_length * region.image_extent.height + y_as_u32 * buffer_row_length) * pixel_size;
+
+                const src_slice = src_map[src_offset..(src_offset + row_size)];
+                const dst_slice = dst_map[dst_offset..(dst_offset + row_size)];
+                @memcpy(dst_slice, src_slice);
+            }
+        }
+
+        src_memory.unmap();
+        dst_memory.unmap();
+    }
+}

src/soft/SoftPhysicalDevice.zig

@@ -52,7 +52,7 @@ pub fn create(allocator: std.mem.Allocator, instance: *const base.Instance) VkEr
         .max_sampler_allocation_count = 4096,
         .buffer_image_granularity = 131072,
         .sparse_address_space_size = 0,
-        .max_bound_descriptor_sets = 4,
+        .max_bound_descriptor_sets = base.VULKAN_MAX_DESCRIPTOR_SETS,
         .max_per_stage_descriptor_samplers = 16,
         .max_per_stage_descriptor_uniform_buffers = 12,
         .max_per_stage_descriptor_storage_buffers = 4,

src/soft/SoftPipeline.zig

@@ -6,10 +6,32 @@ const spv = @import("spv");
 const VkError = base.VkError;
 const Device = base.Device;
 
+const NonDispatchable = base.NonDispatchable;
+const ShaderModule = base.ShaderModule;
+
+const SoftDevice = @import("SoftDevice.zig");
+const SoftShaderModule = @import("SoftShaderModule.zig");
+
 const Self = @This();
 pub const Interface = base.Pipeline;
 
+const Shader = struct {
+    module: *SoftShaderModule,
+    runtimes: []spv.Runtime,
+    entry: []const u8,
+};
+
+const Stages = enum {
+    vertex,
+    tessellation_control,
+    tessellation_evaluation,
+    geometry,
+    fragment,
+    compute,
+};
+
 interface: Interface,
+stages: std.EnumMap(Stages, Shader),
 
 pub fn createCompute(device: *base.Device, allocator: std.mem.Allocator, cache: ?*base.PipelineCache, info: *const vk.ComputePipelineCreateInfo) VkError!*Self {
     const self = allocator.create(Self) catch return VkError.OutOfHostMemory;
@@ -21,8 +43,35 @@ pub fn createCompute(device: *base.Device, allocator: std.mem.Allocator, cache:
         .destroy = destroy,
     };
 
+    const soft_device: *SoftDevice = @alignCast(@fieldParentPtr("interface", device));
+    const module = try NonDispatchable(ShaderModule).fromHandleObject(info.stage.module);
+    const soft_module: *SoftShaderModule = @alignCast(@fieldParentPtr("interface", module));
+
+    const device_allocator = soft_device.device_allocator.allocator();
+
     self.* = .{
         .interface = interface,
+        .stages = std.EnumMap(Stages, Shader).init(.{
+            .compute = .{
+                .module = blk: {
+                    soft_module.ref();
+                    break :blk soft_module;
+                },
+                .runtimes = blk: {
+                    const runtimes = device_allocator.alloc(spv.Runtime, soft_device.workers.getIdCount()) catch return VkError.OutOfHostMemory;
+                    errdefer device_allocator.free(runtimes);
+
+                    for (runtimes) |*runtime| {
+                        runtime.* = spv.Runtime.init(device_allocator, &soft_module.module) catch |err| {
+                            std.log.scoped(.SpvRuntimeInit).err("SPIR-V Runtime failed to initialize, {s}", .{@errorName(err)});
+                            return VkError.Unknown;
+                        };
+                    }
+                    break :blk runtimes;
+                },
+                .entry = allocator.dupe(u8, std.mem.span(info.stage.p_name)) catch return VkError.OutOfHostMemory,
+            },
+        }),
     };
     return self;
 }
@@ -37,13 +86,38 @@ pub fn createGraphics(device: *base.Device, allocator: std.mem.Allocator, cache:
         .destroy = destroy,
     };
 
+    const soft_device: *SoftDevice = @alignCast(@fieldParentPtr("interface", device));
+
+    const runtimes = allocator.alloc(spv.Runtime, soft_device.workers.getIdCount()) catch return VkError.OutOfHostMemory;
+    errdefer allocator.free(runtimes);
+
+    //for (runtimes) |*runtime| {
+    //    runtime.* = spv.Runtime.init() catch |err| {
+    //        std.log.scoped(.SpvRuntimeInit).err("SPIR-V Runtime failed to initialize, {s}", .{@errorName(err)});
+    //        return VkError.Unknown;
+    //    };
+    //}
+
     self.* = .{
         .interface = interface,
+        .stages = std.enums.EnumMap(Stages, Shader).init(.{}),
     };
     return self;
 }
 
 pub fn destroy(interface: *Interface, allocator: std.mem.Allocator) void {
     const self: *Self = @alignCast(@fieldParentPtr("interface", interface));
+    const soft_device: *SoftDevice = @alignCast(@fieldParentPtr("interface", interface.owner));
+    const device_allocator = soft_device.device_allocator.allocator();
+
+    var it = self.stages.iterator();
+    while (it.next()) |stage| {
+        stage.value.module.unref(allocator);
+        for (stage.value.runtimes) |*runtime| {
+            runtime.deinit(device_allocator);
+        }
+        device_allocator.free(stage.value.runtimes);
+        allocator.free(stage.value.entry);
+    }
     allocator.destroy(self);
 }

src/soft/SoftQueue.zig

@@ -4,11 +4,12 @@ const base = @import("base");
 
 const RefCounter = base.RefCounter;
 
-const Device = @import("device/Device.zig");
+const ExecutionDevice = @import("device/Device.zig");
 const Dispatchable = base.Dispatchable;
 
 const CommandBuffer = base.CommandBuffer;
 const SoftDevice = @import("SoftDevice.zig");
+const SoftCommandBuffer = @import("SoftCommandBuffer.zig");
 
 const VkError = base.VkError;
 
@@ -61,7 +62,7 @@ pub fn submit(interface: *Interface, infos: []Interface.SubmitInfo, p_fence: ?*b
     defer self.lock.unlockShared();
 
     for (infos) |info| {
-        // Cloning info to keep them alive until commands dispatch end
+        // Cloning info to keep them alive until command execution ends
         const cloned_info: Interface.SubmitInfo = .{
             .command_buffers = info.command_buffers.clone(allocator) catch return VkError.OutOfDeviceMemory,
         };
@@ -97,14 +98,13 @@ fn taskRunner(self: *Self, info: Interface.SubmitInfo, p_fence: ?*base.Fence, ru
         command_buffers.deinit(soft_device.device_allocator.allocator());
     }
 
-    var device = Device.init();
+    var execution_device: ExecutionDevice = .init;
-    defer device.deinit();
+    execution_device.setup(soft_device);
+    defer execution_device.deinit();
 
-    loop: for (info.command_buffers.items) |command_buffer| {
+    for (info.command_buffers.items) |command_buffer| {
-        command_buffer.submit() catch continue :loop;
+        const soft_command_buffer: *SoftCommandBuffer = @alignCast(@fieldParentPtr("interface", command_buffer));
-        for (command_buffer.commands.items) |command| {
+        soft_command_buffer.execute(&execution_device) catch |err| base.errors.errorLoggerContext(err, "the software execution device");
-            device.dispatch(&command) catch |err| base.errors.errorLoggerContext(err, "the software command dispatcher");
-        }
     }
 
     if (p_fence) |fence| {

src/soft/SoftShaderModule.zig

@@ -13,6 +13,10 @@ pub const Interface = base.ShaderModule;
 interface: Interface,
 module: spv.Module,
 
+/// Pipelines need SPIR-V module reference so shader module may not
+/// be destroy on call to `vkDestroyShaderModule`
+ref_count: std.atomic.Value(usize),
+
 pub fn create(device: *base.Device, allocator: std.mem.Allocator, info: *const vk.ShaderModuleCreateInfo) VkError!*Self {
     const self = allocator.create(Self) catch return VkError.OutOfHostMemory;
     errdefer allocator.destroy(self);
@@ -27,22 +31,33 @@ pub fn create(device: *base.Device, allocator: std.mem.Allocator, info: *const v
 
     self.* = .{
         .interface = interface,
-        .module = spv.Module.init(
+        .module = spv.Module.init(allocator, code, .{
-            allocator,
-            code,
-            .{
             .use_simd_vectors_specializations = !std.process.hasEnvVarConstant(lib.NO_SHADER_SIMD_ENV_NAME),
-            },
+        }) catch |err| switch (err) {
-        ) catch |err| switch (err) {
             spv.Module.ModuleError.OutOfMemory => return VkError.OutOfHostMemory,
             else => return VkError.ValidationFailed,
         },
+        .ref_count = std.atomic.Value(usize).init(1),
     };
     return self;
 }
 
 pub fn destroy(interface: *Interface, allocator: std.mem.Allocator) void {
     const self: *Self = @alignCast(@fieldParentPtr("interface", interface));
+    self.unref(allocator);
+}
+
+pub inline fn drop(self: *Self, allocator: std.mem.Allocator) void {
     self.module.deinit(allocator);
     allocator.destroy(self);
 }
+
+pub inline fn ref(self: *Self) void {
+    _ = self.ref_count.fetchAdd(1, .monotonic);
+}
+
+pub inline fn unref(self: *Self, allocator: std.mem.Allocator) void {
+    if (self.ref_count.fetchSub(1, .release) == 1) {
+        self.drop(allocator);
+    }
+}

src/soft/device/ComputeRoutines.zig

@@ -0,0 +1,218 @@
+const std = @import("std");
+const vk = @import("vulkan");
+const base = @import("base");
+const spv = @import("spv");
+
+const PipelineState = @import("PipelineState.zig");
+
+const SoftDevice = @import("../SoftDevice.zig");
+const SoftPipeline = @import("../SoftPipeline.zig");
+
+const VkError = base.VkError;
+const SpvRuntimeError = spv.Runtime.RuntimeError;
+
+const Self = @This();
+
+const RunData = struct {
+    self: *Self,
+    batch_id: usize,
+    group_count: usize,
+    group_count_x: usize,
+    group_count_y: usize,
+    group_count_z: usize,
+    invocations_per_workgroup: usize,
+    pipeline: *SoftPipeline,
+};
+
+device: *SoftDevice,
+state: *PipelineState,
+batch_size: usize,
+
+pub fn init(device: *SoftDevice, state: *PipelineState) Self {
+    return .{
+        .device = device,
+        .state = state,
+        .batch_size = 0,
+    };
+}
+
+pub fn destroy(self: *Self) void {
+    _ = self;
+}
+
+pub fn dispatch(self: *Self, group_count_x: u32, group_count_y: u32, group_count_z: u32) VkError!void {
+    const group_count: usize = @intCast(group_count_x * group_count_y * group_count_z);
+
+    const pipeline = self.state.pipeline orelse return VkError.InvalidPipelineDrv;
+    const shader = pipeline.stages.getPtr(.compute) orelse return VkError.InvalidPipelineDrv;
+    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;
+
+    //var wg: std.Thread.WaitGroup = .{};
+    for (0..@min(self.batch_size, group_count)) |batch_id| {
+        //self.device.workers.spawnWg(&wg, runWrapper, .{
+        //    RunData{
+        //        .self = self,
+        //        .batch_id = batch_id,
+        //        .group_count = group_count,
+        //        .group_count_x = @as(usize, @intCast(group_count_x)),
+        //        .group_count_y = @as(usize, @intCast(group_count_y)),
+        //        .group_count_z = @as(usize, @intCast(group_count_z)),
+        //        .invocations_per_workgroup = invocations_per_workgroup,
+        //        .pipeline = pipeline,
+        //    },
+        //});
+        runWrapper(
+            RunData{
+                .self = self,
+                .batch_id = batch_id,
+                .group_count = group_count,
+                .group_count_x = @as(usize, @intCast(group_count_x)),
+                .group_count_y = @as(usize, @intCast(group_count_y)),
+                .group_count_z = @as(usize, @intCast(group_count_z)),
+                .invocations_per_workgroup = invocations_per_workgroup,
+                .pipeline = pipeline,
+            },
+        );
+    }
+    //self.device.workers.waitAndWork(&wg);
+}
+
+fn runWrapper(data: RunData) void {
+    @call(.always_inline, run, .{data}) catch |err| {
+        std.log.scoped(.@"SPIR-V runtime").err("SPIR-V runtime catched a '{s}'", .{@errorName(err)});
+        if (@errorReturnTrace()) |trace| {
+            std.debug.dumpStackTrace(trace.*);
+        }
+    };
+}
+
+inline fn run(data: RunData) !void {
+    const allocator = data.self.device.device_allocator.allocator();
+
+    const shader = data.pipeline.stages.getPtrAssertContains(.compute);
+    const rt = &shader.runtimes[data.batch_id];
+
+    const entry = try rt.getEntryPointByName(shader.entry);
+
+    try data.self.syncDescriptorSets(allocator, rt, true);
+
+    var group_index: usize = data.batch_id;
+    while (group_index < data.group_count) : (group_index += data.self.batch_size) {
+        var modulo: usize = group_index;
+
+        const group_z = @divTrunc(modulo, data.group_count_x * data.group_count_y);
+
+        modulo -= group_z * data.group_count_x * data.group_count_y;
+        const group_y = @divTrunc(modulo, data.group_count_x);
+
+        modulo -= group_y * data.group_count_x;
+        const group_x = modulo;
+
+        try setupWorkgroupBuiltins(data.self, rt, .{
+            @as(u32, @intCast(data.group_count_x)),
+            @as(u32, @intCast(data.group_count_y)),
+            @as(u32, @intCast(data.group_count_z)),
+        }, .{
+            @as(u32, @intCast(group_x)),
+            @as(u32, @intCast(group_y)),
+            @as(u32, @intCast(group_z)),
+        });
+
+        for (0..data.invocations_per_workgroup) |i| {
+            try setupSubgroupBuiltins(data.self, rt, .{
+                @as(u32, @intCast(group_x)),
+                @as(u32, @intCast(group_y)),
+                @as(u32, @intCast(group_z)),
+            }, i);
+
+            rt.callEntryPoint(allocator, entry) catch |err| switch (err) {
+                // Some errors can be ignored
+                SpvRuntimeError.OutOfBounds,
+                SpvRuntimeError.Killed,
+                => {},
+                else => return err,
+            };
+        }
+
+        try data.self.syncDescriptorSets(allocator, rt, false);
+    }
+}
+
+fn syncDescriptorSets(self: *Self, allocator: std.mem.Allocator, rt: *spv.Runtime, write: bool) !void {
+    sets: for (self.state.sets[0..], 0..) |set, set_index| {
+        if (set == null)
+            continue :sets;
+
+        bindings: for (set.?.descriptors[0..], 0..) |binding, binding_index| {
+            switch (binding) {
+                .buffer => |buffer_data| if (buffer_data.object) |buffer| {
+                    const memory = if (buffer.interface.memory) |memory| memory else continue :bindings;
+                    const map: []u8 = @as([*]u8, @ptrCast(try memory.map(buffer_data.offset, buffer_data.size)))[0..buffer_data.size];
+                    if (write) {
+                        try rt.writeDescriptorSet(
+                            allocator,
+                            map,
+                            @as(u32, @intCast(set_index)),
+                            @as(u32, @intCast(binding_index)),
+                        );
+                    } else {
+                        try rt.readDescriptorSet(
+                            map,
+                            @as(u32, @intCast(set_index)),
+                            @as(u32, @intCast(binding_index)),
+                        );
+                    }
+                },
+                else => {},
+            }
+        }
+    }
+}
+
+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,
+    };
+
+    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.local_size_x,
+        spv_module.local_size_y,
+        spv_module.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]);
+    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[0] = idx;
+
+    const global_invocation_index = local_base + local_invocation;
+
+    rt.writeBuiltIn(std.mem.asBytes(&global_invocation_index), .GlobalInvocationId) catch {};
+}

src/soft/device/Device.zig

@@ -2,107 +2,38 @@ const std = @import("std");
 const vk = @import("vulkan");
 const base = @import("base");
 
-const SoftImage = @import("../SoftImage.zig");
+const SoftDescriptorSet = @import("../SoftDescriptorSet.zig");
+const SoftDevice = @import("../SoftDevice.zig");
+const SoftPipeline = @import("../SoftPipeline.zig");
+
+const ComputeRoutines = @import("ComputeRoutines.zig");
+const PipelineState = @import("PipelineState.zig");
 
-const cmd = base.commands;
 const VkError = base.VkError;
 
 const Self = @This();
 
-pub fn init() Self {
+compute_routines: ComputeRoutines,
-    return .{};
+
+/// .graphics = 0
+/// .compute = 1
+pipeline_states: [2]PipelineState,
+
+pub const init: Self = .{
+    .compute_routines = undefined,
+    .pipeline_states = undefined,
+};
+
+pub fn setup(self: *Self, device: *SoftDevice) void {
+    for (self.pipeline_states[0..]) |*state| {
+        state.* = .{
+            .pipeline = null,
+            .sets = [_]?*SoftDescriptorSet{null} ** base.VULKAN_MAX_DESCRIPTOR_SETS,
+        };
+    }
+    self.compute_routines = .init(device, &self.pipeline_states[@intFromEnum(vk.PipelineBindPoint.compute)]);
 }
 
 pub fn deinit(self: *Self) void {
-    _ = self;
+    self.compute_routines.destroy();
-}
-
-pub fn dispatch(self: *Self, command: *const cmd.Command) VkError!void {
-    _ = self;
-    switch (command.*) {
-        .ClearColorImage => |data| try clearColorImage(&data),
-        .CopyBuffer => |data| try copyBuffer(&data),
-        .CopyImage => |data| try copyImage(&data),
-        .CopyImageToBuffer => |data| try copyImageToBuffer(&data),
-        .FillBuffer => |data| try fillBuffer(&data),
-        else => {},
-    }
-}
-
-fn clearColorImage(data: *const cmd.CommandClearColorImage) VkError!void {
-    const soft_image: *SoftImage = @alignCast(@fieldParentPtr("interface", data.image));
-    soft_image.clearRange(data.clear_color, data.range);
-}
-
-fn copyBuffer(data: *const cmd.CommandCopyBuffer) VkError!void {
-    for (data.regions) |region| {
-        const src_memory = if (data.src.memory) |memory| memory else return VkError.ValidationFailed;
-        const dst_memory = if (data.dst.memory) |memory| memory else return VkError.ValidationFailed;
-
-        const src_map: []u8 = @as([*]u8, @ptrCast(try src_memory.map(region.src_offset, region.size)))[0..region.size];
-        const dst_map: []u8 = @as([*]u8, @ptrCast(try dst_memory.map(region.dst_offset, region.size)))[0..region.size];
-
-        @memcpy(dst_map, src_map);
-
-        src_memory.unmap();
-        dst_memory.unmap();
-    }
-}
-
-fn copyImage(data: *const cmd.CommandCopyImage) VkError!void {
-    _ = data;
-    std.log.scoped(.commandExecutor).warn("FIXME: implement image to image copy", .{});
-}
-
-fn copyImageToBuffer(data: *const cmd.CommandCopyImageToBuffer) VkError!void {
-    for (data.regions) |region| {
-        const src_memory = if (data.src.memory) |memory| memory else return VkError.ValidationFailed;
-        const dst_memory = if (data.dst.memory) |memory| memory else return VkError.ValidationFailed;
-
-        const pixel_size: u32 = @intCast(data.src.getPixelSize());
-        const image_row_pitch: u32 = data.src.extent.width * pixel_size;
-        const image_size: u32 = @intCast(data.src.getTotalSize());
-
-        const buffer_row_length: u32 = if (region.buffer_row_length != 0) region.buffer_row_length else region.image_extent.width;
-        const buffer_row_pitch: u32 = buffer_row_length * pixel_size;
-        const buffer_size: u32 = buffer_row_pitch * region.image_extent.height * region.image_extent.depth;
-
-        const src_map: []u8 = @as([*]u8, @ptrCast(try src_memory.map(0, image_size)))[0..image_size];
-        const dst_map: []u8 = @as([*]u8, @ptrCast(try dst_memory.map(region.buffer_offset, buffer_size)))[0..buffer_size];
-
-        const row_size = region.image_extent.width * pixel_size;
-        for (0..data.src.extent.depth) |z| {
-            for (0..data.src.extent.height) |y| {
-                const z_as_u32: u32 = @intCast(z);
-                const y_as_u32: u32 = @intCast(y);
-
-                const src_offset = ((@as(u32, @intCast(region.image_offset.z)) + z_as_u32) * data.src.extent.height + @as(u32, @intCast(region.image_offset.y)) + y_as_u32) * image_row_pitch + @as(u32, @intCast(region.image_offset.x)) * pixel_size;
-                const dst_offset = (z_as_u32 * buffer_row_length * region.image_extent.height + y_as_u32 * buffer_row_length) * pixel_size;
-
-                const src_slice = src_map[src_offset..(src_offset + row_size)];
-                const dst_slice = dst_map[dst_offset..(dst_offset + row_size)];
-                @memcpy(dst_slice, src_slice);
-            }
-        }
-
-        src_memory.unmap();
-        dst_memory.unmap();
-    }
-}
-
-fn fillBuffer(data: *const cmd.CommandFillBuffer) VkError!void {
-    const memory = if (data.buffer.memory) |memory| memory else return VkError.ValidationFailed;
-    var memory_map: []u32 = @as([*]u32, @ptrCast(@alignCast(try memory.map(data.offset, data.size))))[0..data.size];
-
-    var bytes = if (data.size == vk.WHOLE_SIZE) memory.size - data.offset else data.size;
-
-    var i: usize = 0;
-    while (bytes >= 4) : ({
-        bytes -= 4;
-        i += 1;
-    }) {
-        memory_map[i] = data.data;
-    }
-
-    memory.unmap();
 }

src/soft/device/PipelineState.zig

@@ -0,0 +1,9 @@
+const std = @import("std");
+const vk = @import("vulkan");
+const base = @import("base");
+
+const SoftDescriptorSet = @import("../SoftDescriptorSet.zig");
+const SoftPipeline = @import("../SoftPipeline.zig");
+
+pipeline: ?*SoftPipeline,
+sets: [base.VULKAN_MAX_DESCRIPTOR_SETS]?*SoftDescriptorSet,

src/vulkan/Buffer.zig

@@ -40,7 +40,7 @@ pub inline fn destroy(self: *Self, allocator: std.mem.Allocator) void {
 }
 
 pub inline fn bindMemory(self: *Self, memory: *DeviceMemory, offset: vk.DeviceSize) VkError!void {
-    if (offset >= self.size or !self.allowed_memory_types.isSet(memory.memory_type_index)) {
+    if (offset > memory.size or !self.allowed_memory_types.isSet(memory.memory_type_index)) {
         return VkError.ValidationFailed;
     }
     self.memory = memory;

src/vulkan/CommandBuffer.zig

@@ -1,7 +1,6 @@
 const std = @import("std");
 const vk = @import("vulkan");
-
+const lib = @import("lib.zig");
-const cmd = @import("commands.zig");
 
 const NonDispatchable = @import("NonDispatchable.zig").NonDispatchable;
 const VkError = @import("error_set.zig").VkError;
@@ -13,8 +12,8 @@ const Buffer = @import("Buffer.zig");
 const CommandPool = @import("CommandPool.zig");
 const Event = @import("Event.zig");
 const Image = @import("Image.zig");
-
+const Pipeline = @import("Pipeline.zig");
-const COMMAND_BUFFER_BASE_CAPACITY = 256;
+const DescriptorSet = @import("DescriptorSet.zig");
 
 const State = enum {
     Initial,
@@ -32,18 +31,20 @@ pool: *CommandPool,
 state: State,
 begin_info: ?vk.CommandBufferBeginInfo,
 host_allocator: VulkanAllocator,
-commands: std.ArrayList(cmd.Command),
 state_mutex: std.Thread.Mutex,
 
 vtable: *const VTable,
 dispatch_table: *const DispatchTable,
 
 pub const DispatchTable = struct {
+    bindDescriptorSets: *const fn (*Self, vk.PipelineBindPoint, u32, [lib.VULKAN_MAX_DESCRIPTOR_SETS]?*DescriptorSet, []const u32) VkError!void,
+    bindPipeline: *const fn (*Self, vk.PipelineBindPoint, *Pipeline) VkError!void,
     begin: *const fn (*Self, *const vk.CommandBufferBeginInfo) VkError!void,
     clearColorImage: *const fn (*Self, *Image, vk.ImageLayout, *const vk.ClearColorValue, vk.ImageSubresourceRange) VkError!void,
     copyBuffer: *const fn (*Self, *Buffer, *Buffer, []const vk.BufferCopy) VkError!void,
     copyImage: *const fn (*Self, *Image, vk.ImageLayout, *Image, vk.ImageLayout, []const vk.ImageCopy) VkError!void,
     copyImageToBuffer: *const fn (*Self, *Image, vk.ImageLayout, *Buffer, []const vk.BufferImageCopy) VkError!void,
+    dispatch: *const fn (*Self, u32, u32, u32) VkError!void,
     end: *const fn (*Self) VkError!void,
     fillBuffer: *const fn (*Self, *Buffer, vk.DeviceSize, vk.DeviceSize, u32) VkError!void,
     reset: *const fn (*Self, vk.CommandBufferResetFlags) VkError!void,
@@ -63,7 +64,6 @@ pub fn init(device: *Device, allocator: std.mem.Allocator, info: *const vk.Comma
         .state = .Initial,
         .begin_info = null,
         .host_allocator = VulkanAllocator.from(allocator).cloneWithScope(.object),
-        .commands = std.ArrayList(cmd.Command).initCapacity(allocator, COMMAND_BUFFER_BASE_CAPACITY) catch return VkError.OutOfHostMemory,
         .state_mutex = .{},
         .vtable = undefined,
         .dispatch_table = undefined,
@@ -80,8 +80,6 @@ inline fn transitionState(self: *Self, target: State, from_allowed: []const Stat
 }
 
 pub inline fn destroy(self: *Self, allocator: std.mem.Allocator) void {
-    self.cleanCommandList();
-    self.commands.deinit(allocator);
     self.vtable.destroy(self, allocator);
 }
 
@@ -104,7 +102,6 @@ pub inline fn reset(self: *Self, flags: vk.CommandBufferResetFlags) VkError!void
     if (!self.pool.flags.reset_command_buffer_bit) {
         return VkError.ValidationFailed;
     }
-    defer self.cleanCommandList();
 
     self.transitionState(.Initial, &.{ .Initial, .Recording, .Executable, .Invalid }) catch return VkError.ValidationFailed;
     try self.dispatch_table.reset(self, flags);
@@ -119,74 +116,43 @@ pub inline fn submit(self: *Self) VkError!void {
     self.transitionState(.Pending, &.{ .Pending, .Executable }) catch return VkError.ValidationFailed;
 }
 
-fn cleanCommandList(self: *Self) void {
-    const allocator = self.host_allocator.allocator();
-    for (self.commands.items) |command| {
-        switch (command) {
-            .CopyBuffer => |data| allocator.free(data.regions),
-            .CopyImage => |data| allocator.free(data.regions),
-            .CopyImageToBuffer => |data| allocator.free(data.regions),
-            else => {},
-        }
-    }
-}
-
 // Commands ====================================================================================================
 
+pub inline fn bindDescriptorSets(self: *Self, bind_point: vk.PipelineBindPoint, first_set: u32, sets: []const vk.DescriptorSet, dynamic_offsets: []const u32) VkError!void {
+    var inner_sets = [_]?*DescriptorSet{null} ** lib.VULKAN_MAX_DESCRIPTOR_SETS;
+    for (sets, inner_sets[0..sets.len]) |set, *inner_set| {
+        inner_set.* = try NonDispatchable(DescriptorSet).fromHandleObject(set);
+    }
+    try self.dispatch_table.bindDescriptorSets(self, bind_point, first_set, inner_sets, dynamic_offsets);
+}
+
+pub inline fn bindPipeline(self: *Self, bind_point: vk.PipelineBindPoint, pipeline: *Pipeline) VkError!void {
+    try self.dispatch_table.bindPipeline(self, bind_point, pipeline);
+}
+
 pub inline fn clearColorImage(self: *Self, image: *Image, layout: vk.ImageLayout, color: *const vk.ClearColorValue, ranges: []const vk.ImageSubresourceRange) VkError!void {
-    const allocator = self.host_allocator.allocator();
     for (ranges) |range| {
-        self.commands.append(allocator, .{ .ClearColorImage = .{
-            .image = image,
-            .layout = layout,
-            .clear_color = color.*,
-            .range = range,
-        } }) catch return VkError.OutOfHostMemory;
         try self.dispatch_table.clearColorImage(self, image, layout, color, range);
     }
 }
 
 pub inline fn copyBuffer(self: *Self, src: *Buffer, dst: *Buffer, regions: []const vk.BufferCopy) VkError!void {
-    const allocator = self.host_allocator.allocator();
-    self.commands.append(allocator, .{ .CopyBuffer = .{
-        .src = src,
-        .dst = dst,
-        .regions = allocator.dupe(vk.BufferCopy, regions) catch return VkError.OutOfHostMemory,
-    } }) catch return VkError.OutOfHostMemory;
     try self.dispatch_table.copyBuffer(self, src, dst, regions);
 }
 
 pub inline fn copyImage(self: *Self, src: *Image, src_layout: vk.ImageLayout, dst: *Image, dst_layout: vk.ImageLayout, regions: []const vk.ImageCopy) VkError!void {
-    const allocator = self.host_allocator.allocator();
-    self.commands.append(allocator, .{ .CopyImage = .{
-        .src = src,
-        .src_layout = src_layout,
-        .dst = dst,
-        .dst_layout = dst_layout,
-        .regions = allocator.dupe(vk.ImageCopy, regions) catch return VkError.OutOfHostMemory,
-    } }) catch return VkError.OutOfHostMemory;
     try self.dispatch_table.copyImage(self, src, src_layout, dst, dst_layout, regions);
 }
 
 pub inline fn copyImageToBuffer(self: *Self, src: *Image, src_layout: vk.ImageLayout, dst: *Buffer, regions: []const vk.BufferImageCopy) VkError!void {
-    const allocator = self.host_allocator.allocator();
-    self.commands.append(allocator, .{ .CopyImageToBuffer = .{
-        .src = src,
-        .src_layout = src_layout,
-        .dst = dst,
-        .regions = allocator.dupe(vk.BufferImageCopy, regions) catch return VkError.OutOfHostMemory,
-    } }) catch return VkError.OutOfHostMemory;
     try self.dispatch_table.copyImageToBuffer(self, src, src_layout, dst, regions);
 }
 
+pub inline fn dispatch(self: *Self, group_count_x: u32, group_count_y: u32, group_count_z: u32) VkError!void {
+    try self.dispatch_table.dispatch(self, group_count_x, group_count_y, group_count_z);
+}
+
 pub inline fn fillBuffer(self: *Self, buffer: *Buffer, offset: vk.DeviceSize, size: vk.DeviceSize, data: u32) VkError!void {
-    const allocator = self.host_allocator.allocator();
-    self.commands.append(allocator, .{ .FillBuffer = .{
-        .buffer = buffer,
-        .offset = offset,
-        .size = if (size == vk.WHOLE_SIZE) buffer.size else size,
-        .data = data,
-    } }) catch return VkError.OutOfHostMemory;
     try self.dispatch_table.fillBuffer(self, buffer, offset, size, data);
 }
 

src/vulkan/DescriptorSet.zig

@@ -18,11 +18,14 @@ layout: *DescriptorSetLayout,
 vtable: *const VTable,
 
 pub const VTable = struct {
+    copy: *const fn (*Self, vk.CopyDescriptorSet) VkError!void,
     destroy: *const fn (*Self, std.mem.Allocator) void,
+    write: *const fn (*Self, vk.WriteDescriptorSet) VkError!void,
 };
 
 pub fn init(device: *Device, allocator: std.mem.Allocator, layout: *DescriptorSetLayout) VkError!Self {
     _ = allocator;
+    layout.ref();
     return .{
         .owner = device,
         .layout = layout,
@@ -30,7 +33,15 @@ pub fn init(device: *Device, allocator: std.mem.Allocator, layout: *DescriptorSe
     };
 }
 
+pub inline fn copy(self: *Self, copy_data: vk.CopyDescriptorSet) VkError!void {
+    try self.vtable.copy(self, copy_data);
+}
+
 pub inline fn destroy(self: *Self, allocator: std.mem.Allocator) void {
-    allocator.free(self.layouts);
+    self.layout.unref(allocator);
     self.vtable.destroy(self, allocator);
 }
+
+pub inline fn write(self: *Self, write_data: vk.WriteDescriptorSet) VkError!void {
+    try self.vtable.write(self, write_data);
+}

src/vulkan/DescriptorSetLayout.zig

@@ -1,14 +1,53 @@
 const std = @import("std");
 const vk = @import("vulkan");
 
+const VulkanAllocator = @import("VulkanAllocator.zig");
+
 const VkError = @import("error_set.zig").VkError;
 const Device = @import("Device.zig");
+const Sampler = @import("Sampler.zig");
 
 const Self = @This();
 pub const ObjectType: vk.ObjectType = .descriptor_set_layout;
 
+const BindingLayout = struct {
+    descriptor_type: vk.DescriptorType,
+    dynamic_index: usize,
+    array_size: usize,
+
+    /// This slice points to an array located after the binding layouts array
+    immutable_samplers: []*const Sampler,
+
+    driver_data: *anyopaque,
+};
+
 owner: *Device,
-bindings: ?[]const vk.DescriptorSetLayoutBinding,
+
+/// Memory containing actual binding layouts array and immutable samplers array
+heap: []u8,
+
+bindings: []BindingLayout,
+
+dynamic_offset_count: usize,
+dynamic_descriptor_count: usize,
+
+/// Shader stages affected by this descriptor set
+stages: vk.ShaderStageFlags,
+
+/// Mesa's common Vulkan runtime states:
+///
+/// It's often necessary to store a pointer to the descriptor set layout in
+/// the descriptor so that any entrypoint which has access to a descriptor
+/// set also has the layout. While layouts are often passed into various
+/// entrypoints, they're notably missing from vkUpdateDescriptorSets(). In
+/// order to implement descriptor writes, you either need to stash a pointer
+/// to the descriptor set layout in the descriptor set or you need to copy
+/// all of the relevant information.  Storing a pointer is a lot cheaper.
+///
+/// Because descriptor set layout lifetimes and descriptor set lifetimes are
+/// not guaranteed to coincide, we have to reference count if we're going to
+/// do this.
+ref_count: std.atomic.Value(usize),
 
 vtable: *const VTable,
 
@@ -17,18 +56,97 @@ pub const VTable = struct {
 };
 
 pub fn init(device: *Device, allocator: std.mem.Allocator, info: *const vk.DescriptorSetLayoutCreateInfo) VkError!Self {
-    const bindings = if (info.p_bindings) |bindings|
+    const command_allocator = VulkanAllocator.from(allocator).cloneWithScope(.command).allocator();
-        allocator.dupe(vk.DescriptorSetLayoutBinding, bindings[0..info.binding_count]) catch return VkError.OutOfHostMemory
+
-    else
+    var binding_count: usize = 0;
-        null;
+    var immutable_samplers_count: usize = 0;
+
+    if (info.p_bindings) |binding_infos| {
+        for (binding_infos, 0..info.binding_count) |binding, _| {
+            binding_count = @max(binding_count, binding.binding + 1);
+            if (bindingHasImmutableSamplers(binding)) {
+                immutable_samplers_count += binding.descriptor_count;
+            }
+        }
+    }
+
+    const size = (binding_count * @sizeOf(BindingLayout)) + (immutable_samplers_count * @sizeOf(*Sampler));
+
+    // Clean way to put the immutable samplers array right after the binding layouts one
+    const heap = allocator.alloc(u8, size) catch return VkError.OutOfHostMemory;
+    errdefer allocator.free(heap);
+
+    var local_heap = std.heap.FixedBufferAllocator.init(heap);
+    const local_allocator = local_heap.allocator();
+
+    const bindings = local_allocator.alloc(BindingLayout, binding_count) catch return VkError.OutOfHostMemory;
+    const immutable_samplers = local_allocator.alloc(*const Sampler, immutable_samplers_count) catch return VkError.OutOfHostMemory;
+
+    var stages: vk.ShaderStageFlags = .{};
+
+    if (info.p_bindings) |binding_infos| {
+        const sorted_bindings = command_allocator.dupe(vk.DescriptorSetLayoutBinding, binding_infos[0..info.binding_count]) catch return VkError.OutOfHostMemory;
+        defer command_allocator.free(sorted_bindings);
+        std.mem.sort(vk.DescriptorSetLayoutBinding, sorted_bindings, .{}, sortBindings);
+
+        for (sorted_bindings) |binding_info| {
+            const binding_index = binding_info.binding;
+
+            const descriptor_count = switch (binding_info.descriptor_type) {
+                .inline_uniform_block => 1,
+                else => binding_info.descriptor_count,
+            };
+
+            bindings[binding_index] = .{
+                .descriptor_type = binding_info.descriptor_type,
+                .array_size = descriptor_count,
+                .dynamic_index = 0,
+                .immutable_samplers = immutable_samplers[0..],
+                .driver_data = undefined,
+            };
+
+            stages = stages.merge(binding_info.stage_flags);
+        }
+    }
 
     return .{
         .owner = device,
+        .heap = heap,
         .bindings = bindings,
+        .dynamic_offset_count = 0,
+        .dynamic_descriptor_count = 0,
+        .stages = stages,
+        .ref_count = std.atomic.Value(usize).init(1),
         .vtable = undefined,
     };
 }
 
+fn sortBindings(_: @TypeOf(.{}), lhs: vk.DescriptorSetLayoutBinding, rhs: vk.DescriptorSetLayoutBinding) bool {
+    return lhs.binding < rhs.binding;
+}
+
+inline fn bindingHasImmutableSamplers(binding: vk.DescriptorSetLayoutBinding) bool {
+    return switch (binding.descriptor_type) {
+        .sampler, .combined_image_sampler => binding.p_immutable_samplers != null,
+        else => false,
+    };
+}
+
 pub inline fn destroy(self: *Self, allocator: std.mem.Allocator) void {
+    self.unref(allocator);
+}
+
+pub inline fn drop(self: *Self, allocator: std.mem.Allocator) void {
+    allocator.free(self.heap);
     self.vtable.destroy(self, allocator);
 }
+
+pub inline fn ref(self: *Self) void {
+    _ = self.ref_count.fetchAdd(1, .monotonic);
+}
+
+pub inline fn unref(self: *Self, allocator: std.mem.Allocator) void {
+    if (self.ref_count.fetchSub(1, .release) == 1) {
+        self.drop(allocator);
+    }
+}

src/vulkan/Dispatchable.zig

@@ -47,11 +47,11 @@ pub fn Dispatchable(comptime T: type) type {
         pub inline fn fromHandle(vk_handle: anytype) VkError!*Self {
             const handle = @intFromEnum(vk_handle);
             if (handle == 0) {
-                return VkError.ValidationFailed;
+                return VkError.InvalidHandleDrv;
             }
             const dispatchable: *Self = @ptrFromInt(handle);
             if (dispatchable.object_type != T.ObjectType) {
-                return VkError.ValidationFailed;
+                return VkError.InvalidHandleDrv;
             }
             return dispatchable;
         }

src/vulkan/Image.zig

@@ -79,11 +79,11 @@ pub inline fn getClearFormat(self: *Self) vk.Format {
         .r32g32b32a32_sfloat;
 }
 
-pub inline fn getPixelSize(self: *Self) usize {
+pub inline fn getPixelSize(self: *const Self) usize {
     return lib.vku.vkuFormatTexelBlockSize(@intCast(@intFromEnum(self.format)));
 }
 
-pub inline fn getTotalSize(self: *Self) usize {
+pub inline fn getTotalSize(self: *const Self) usize {
     const pixel_size = self.getPixelSize();
     return self.extent.width * self.extent.height * self.extent.depth * pixel_size;
 }
@@ -92,7 +92,7 @@ pub inline fn getFormatPixelSize(format: vk.Format) usize {
     return lib.vku.vkuFormatTexelBlockSize(@intCast(@intFromEnum(format)));
 }
 
-pub inline fn getFormatTotalSize(self: *Self, format: vk.Format) usize {
+pub inline fn getFormatTotalSize(self: *const Self, format: vk.Format) usize {
     const pixel_size = self.getFormatPixelSize(format);
     return self.extent.width * self.extent.height * self.extent.depth * pixel_size;
 }

src/vulkan/NonDispatchable.zig

@@ -42,11 +42,11 @@ pub fn NonDispatchable(comptime T: type) type {
         pub inline fn fromHandle(vk_handle: anytype) VkError!*Self {
             const handle = @intFromEnum(vk_handle);
             if (handle == 0) {
-                return VkError.ValidationFailed;
+                return VkError.InvalidHandleDrv;
             }
             const non_dispatchable: *Self = @ptrFromInt(handle);
             if (non_dispatchable.object_type != T.ObjectType) {
-                return VkError.ValidationFailed;
+                return VkError.InvalidHandleDrv;
             }
             return non_dispatchable;
         }

src/vulkan/PipelineLayout.zig

@@ -1,7 +1,9 @@
 const std = @import("std");
 const vk = @import("vulkan");
+const lib = @import("lib.zig");
 
-const NonDispatchable = @import("NonDispatchable.zig");
+const NonDispatchable = @import("NonDispatchable.zig").NonDispatchable;
+const DescriptorSetLayout = @import("DescriptorSetLayout.zig");
 
 const VkError = @import("error_set.zig").VkError;
 
@@ -12,6 +14,30 @@ pub const ObjectType: vk.ObjectType = .pipeline_layout;
 
 owner: *Device,
 
+set_count: usize,
+
+set_layouts: [lib.VULKAN_MAX_DESCRIPTOR_SETS]?*DescriptorSetLayout,
+
+dynamic_descriptor_offsets: [lib.VULKAN_MAX_DESCRIPTOR_SETS]usize,
+
+push_ranges_count: usize,
+push_ranges: [lib.VULKAN_MAX_PUSH_CONSTANT_RANGES]vk.PushConstantRange,
+
+/// Mesa's common Vulkan runtime states:
+///
+/// It's often necessary to store a pointer to the descriptor set layout in
+/// the descriptor so that any entrypoint which has access to a descriptor
+/// set also has the layout. While layouts are often passed into various
+/// entrypoints, they're notably missing from vkUpdateDescriptorSets(). In
+/// order to implement descriptor writes, you either need to stash a pointer
+/// to the descriptor set layout in the descriptor set or you need to copy
+/// all of the relevant information.  Storing a pointer is a lot cheaper.
+///
+/// Because descriptor set layout lifetimes and descriptor set lifetimes are
+/// not guaranteed to coincide, we have to reference count if we're going to
+/// do this.
+ref_count: std.atomic.Value(usize),
+
 vtable: *const VTable,
 
 pub const VTable = struct {
@@ -20,13 +46,52 @@ pub const VTable = struct {
 
 pub fn init(device: *Device, allocator: std.mem.Allocator, info: *const vk.PipelineLayoutCreateInfo) VkError!Self {
     _ = allocator;
-    _ = info;
+    var self: Self = .{
-    return .{
         .owner = device,
+        .set_count = info.set_layout_count,
+        .set_layouts = [_]?*DescriptorSetLayout{null} ** lib.VULKAN_MAX_DESCRIPTOR_SETS,
+        .dynamic_descriptor_offsets = [_]usize{0} ** lib.VULKAN_MAX_DESCRIPTOR_SETS,
+        .push_ranges_count = info.push_constant_range_count,
+        .push_ranges = undefined,
+        .ref_count = std.atomic.Value(usize).init(1),
         .vtable = undefined,
     };
+
+    if (info.p_set_layouts) |set_layouts| {
+        var dynamic_descriptor_count: usize = 0;
+        for (set_layouts, 0..info.set_layout_count) |handle, i| {
+            self.dynamic_descriptor_offsets[i] = dynamic_descriptor_count;
+            if (handle != .null_handle) {
+                const layout = try NonDispatchable(DescriptorSetLayout).fromHandleObject(handle);
+                self.set_layouts[i] = layout;
+                layout.ref();
+                dynamic_descriptor_count += layout.dynamic_descriptor_count;
+            }
+        }
+    }
+
+    return self;
 }
 
 pub inline fn destroy(self: *Self, allocator: std.mem.Allocator) void {
+    self.unref(allocator);
+}
+
+pub inline fn drop(self: *Self, allocator: std.mem.Allocator) void {
+    for (self.set_layouts) |set_layout| {
+        if (set_layout) |layout| {
+            layout.unref(allocator);
+        }
+    }
     self.vtable.destroy(self, allocator);
 }
+
+pub inline fn ref(self: *Self) void {
+    _ = self.ref_count.fetchAdd(1, .monotonic);
+}
+
+pub inline fn unref(self: *Self, allocator: std.mem.Allocator) void {
+    if (self.ref_count.fetchSub(1, .release) == 1) {
+        self.drop(allocator);
+    }
+}

src/vulkan/ShaderModule.zig

@@ -20,12 +20,8 @@ pub const VTable = struct {
 };
 
 pub fn init(device: *Device, allocator: std.mem.Allocator, info: *const vk.ShaderModuleCreateInfo) VkError!Self {
-    if (std.process.hasEnvVarConstant(lib.DRIVER_LOG_SPIRV_ENV_NAME)) {
+    _ = allocator;
-        logShaderModule(allocator, info) catch |e| {
+    _ = info;
-            std.log.scoped(.ShaderModule).err("Failed to disassemble SPIR-V module to readable text: {s}", .{@errorName(e)});
-        };
-    }
-
     return .{
         .owner = device,
         .vtable = undefined,
@@ -35,35 +31,3 @@ pub fn init(device: *Device, allocator: std.mem.Allocator, info: *const vk.Shade
 pub inline fn destroy(self: *Self, allocator: std.mem.Allocator) void {
     self.vtable.destroy(self, allocator);
 }
-
-fn logShaderModule(allocator: std.mem.Allocator, info: *const vk.ShaderModuleCreateInfo) !void {
-    std.log.scoped(.ShaderModule).info("Logging SPIR-V module", .{});
-
-    var process = std.process.Child.init(&[_][]const u8{ "spirv-dis", "--no-color", "/home/kbz8/Documents/Code/Zig/SPIRV-Interpreter/example/shader.spv" }, allocator);
-
-    process.stdout_behavior = .Pipe;
-    process.stderr_behavior = .Pipe;
-    process.stdin_behavior = .Pipe;
-
-    var stdout: std.ArrayList(u8) = .empty;
-    defer stdout.deinit(allocator);
-    var stderr: std.ArrayList(u8) = .empty;
-    defer stderr.deinit(allocator);
-
-    try process.spawn();
-    errdefer {
-        _ = process.kill() catch {};
-    }
-
-    if (process.stdin) |stdin| {
-        _ = try stdin.write(@ptrCast(@alignCast(info.p_code[0..@divExact(info.code_size, 4)])));
-    }
-    try process.collectOutput(allocator, &stdout, &stderr, 1024 * 1024);
-    _ = try process.wait();
-
-    if (stderr.items.len != 0) {
-        std.log.scoped(.ShaderModule).err("Failed to disassemble SPIR-V module to readable text.\nError:\n{s}", .{stderr.items});
-    } else if (stdout.items.len != 0) {
-        std.log.scoped(.ShaderModule).info("{s}\n{d}", .{ stdout.items, stdout.items.len });
-    }
-}

src/vulkan/commands.zig

@@ -1,97 +0,0 @@
-const std = @import("std");
-const vk = @import("vulkan");
-
-const Buffer = @import("Buffer.zig");
-const Image = @import("Image.zig");
-
-pub const CommandType = enum {
-    BindPipeline,
-    BindVertexBuffer,
-    ClearColorImage,
-    CopyBuffer,
-    CopyImage,
-    CopyImageToBuffer,
-    Draw,
-    DrawIndexed,
-    DrawIndexedIndirect,
-    DrawIndirect,
-    FillBuffer,
-};
-
-pub const CommandBindPipeline = struct {
-    bind_point: vk.PipelineBindPoint,
-};
-pub const CommandBindVertexBuffer = struct {
-    buffers: []*const Buffer,
-    offsets: []vk.DeviceSize,
-    first_binding: u32,
-};
-pub const CommandClearColorImage = struct {
-    image: *Image,
-    layout: vk.ImageLayout,
-    clear_color: vk.ClearColorValue,
-    range: vk.ImageSubresourceRange,
-};
-pub const CommandCopyBuffer = struct {
-    src: *Buffer,
-    dst: *Buffer,
-    regions: []const vk.BufferCopy,
-};
-pub const CommandCopyImage = struct {
-    src: *Image,
-    src_layout: vk.ImageLayout,
-    dst: *Image,
-    dst_layout: vk.ImageLayout,
-    regions: []const vk.ImageCopy,
-};
-pub const CommandCopyImageToBuffer = struct {
-    src: *Image,
-    src_layout: vk.ImageLayout,
-    dst: *Buffer,
-    regions: []const vk.BufferImageCopy,
-};
-pub const CommandDraw = struct {
-    vertex_count: u32,
-    instance_count: u32,
-    first_vertex: u32,
-    first_instance: u32,
-};
-pub const CommandDrawIndexed = struct {
-    index_count: u32,
-    instance_count: u32,
-    first_index: u32,
-    vertex_offset: i32,
-    first_instance: u32,
-};
-pub const CommandDrawIndexedIndirect = struct {
-    buffer: *Buffer,
-    offset: vk.DeviceSize,
-    count: u32,
-    stride: u32,
-};
-pub const CommandDrawIndirect = struct {
-    buffer: *Buffer,
-    offset: vk.DeviceSize,
-    count: u32,
-    stride: u32,
-};
-pub const CommandFillBuffer = struct {
-    buffer: *Buffer,
-    offset: vk.DeviceSize,
-    size: vk.DeviceSize,
-    data: u32,
-};
-
-pub const Command = union(CommandType) {
-    BindPipeline: CommandBindPipeline,
-    BindVertexBuffer: CommandBindVertexBuffer,
-    ClearColorImage: CommandClearColorImage,
-    CopyBuffer: CommandCopyBuffer,
-    CopyImage: CommandCopyImage,
-    CopyImageToBuffer: CommandCopyImageToBuffer,
-    Draw: CommandDraw,
-    DrawIndexed: CommandDrawIndexed,
-    DrawIndexedIndirect: CommandDrawIndexedIndirect,
-    DrawIndirect: CommandDrawIndirect,
-    FillBuffer: CommandFillBuffer,
-};

src/vulkan/error_set.zig

@@ -50,6 +50,10 @@ pub const VkError = error{
     IncompatibleShaderBinaryExt,
     PipelineBinaryMissingKhr,
     NotEnoughSpaceKhr,
+    // ====== Internal errors
+    InvalidHandleDrv,
+    InvalidPipelineDrv,
+    InvalidDeviceMemoryDrv,
 };
 
 pub inline fn errorLogger(err: VkError) void {
@@ -80,7 +84,6 @@ pub inline fn toVkResult(err: VkError) vk.Result {
         VkError.TooManyObjects => .error_too_many_objects,
         VkError.FormatNotSupported => .error_format_not_supported,
         VkError.FragmentedPool => .error_fragmented_pool,
-        VkError.Unknown => .error_unknown,
         VkError.ValidationFailed => .error_validation_failed,
         VkError.OutOfPoolMemory => .error_out_of_pool_memory,
         VkError.InvalidExternalHandle => .error_invalid_external_handle,
@@ -111,5 +114,7 @@ pub inline fn toVkResult(err: VkError) vk.Result {
         VkError.IncompatibleShaderBinaryExt => .incompatible_shader_binary_ext,
         VkError.PipelineBinaryMissingKhr => .pipeline_binary_missing_khr,
         VkError.NotEnoughSpaceKhr => .error_not_enough_space_khr,
+        VkError.InvalidHandleDrv => .error_validation_failed,
+        else => .error_unknown,
     };
 }

src/vulkan/lib.zig

@@ -7,7 +7,6 @@ pub const vku = @cImport({
     @cInclude("vulkan/utility/vk_format_utils.h");
 });
 
-pub const commands = @import("commands.zig");
 pub const errors = @import("error_set.zig");
 pub const lib_vulkan = @import("lib_vulkan.zig");
 pub const logger = @import("logger/logger.zig");
@@ -50,13 +49,35 @@ pub const VULKAN_VENDOR_ID = @typeInfo(vk.VendorId).@"enum".fields[@typeInfo(vk.
 pub const DRIVER_DEBUG_ALLOCATOR_ENV_NAME = "STROLL_DEBUG_ALLOCATOR";
 pub const DRIVER_LOGS_ENV_NAME = "STROLL_LOGS_LEVEL";
 
-pub const DRIVER_LOG_SPIRV_ENV_NAME = "STROLL_LOG_SPIRV";
-
 /// Default driver name
 pub const DRIVER_NAME = "Unnamed Stroll Driver";
 /// Default Vulkan version
 pub const VULKAN_VERSION = vk.makeApiVersion(0, 1, 0, 0);
 
+/// Maximum number of descriptor sets per pipeline
+pub const VULKAN_MAX_DESCRIPTOR_SETS = 4;
+
+/// The number of push constant ranges is effectively bounded
+/// by the number of possible shader stages.  Not the number of stages that can
+/// be compiled together (a pipeline layout can be used in multiple pipelnes
+/// wth different sets of shaders) but the total number of stage bits supported
+/// by the implementation. Currently, those are
+/// - VK_SHADER_STAGE_VERTEX_BIT
+/// - VK_SHADER_STAGE_TESSELLATION_CONTROL_BIT
+/// - VK_SHADER_STAGE_TESSELLATION_EVALUATION_BIT
+/// - VK_SHADER_STAGE_GEOMETRY_BIT
+/// - VK_SHADER_STAGE_FRAGMENT_BIT
+/// - VK_SHADER_STAGE_COMPUTE_BIT
+/// - VK_SHADER_STAGE_RAYGEN_BIT_KHR
+/// - VK_SHADER_STAGE_ANY_HIT_BIT_KHR
+/// - VK_SHADER_STAGE_CLOSEST_HIT_BIT_KHR
+/// - VK_SHADER_STAGE_MISS_BIT_KHR
+/// - VK_SHADER_STAGE_INTERSECTION_BIT_KHR
+/// - VK_SHADER_STAGE_CALLABLE_BIT_KHR
+/// - VK_SHADER_STAGE_TASK_BIT_EXT
+/// - VK_SHADER_STAGE_MESH_BIT_EXT
+pub const VULKAN_MAX_PUSH_CONSTANT_RANGES = 14;
+
 pub const std_options: std.Options = .{
     .log_level = .debug,
     .logFn = logger.log,
@@ -83,13 +104,9 @@ pub inline fn getLogVerboseLevel() LogVerboseLevel {
         .Standard;
 }
 
-pub fn unsupported(comptime fmt: []const u8, args: anytype) void {
+pub inline fn unsupported(comptime fmt: []const u8, args: anytype) void {
-    if (builtin.mode == std.builtin.OptimizeMode.Debug) {
-        std.debug.panic("UNSUPPORTED " ++ fmt, args);
-    } else {
     std.log.scoped(.UNSUPPORTED).warn(fmt, args);
 }
-}
 
 comptime {
     _ = lib_vulkan;

src/vulkan/lib_vulkan.zig

@@ -741,6 +741,7 @@ pub export fn strollCreateComputePipelines(p_device: vk.Device, p_cache: vk.Pipe
     defer entryPointEndLogTrace();
 
     const allocator = VulkanAllocator.init(callbacks, .object).allocator();
+
     const device = Dispatchable(Device).fromHandleObject(p_device) catch |err| return toVkResult(err);
     const cache = if (p_cache == .null_handle) null else NonDispatchable(PipelineCache).fromHandleObject(p_cache) catch |err| return toVkResult(err);
 
@@ -772,6 +773,7 @@ pub export fn strollCreateComputePipelines(p_device: vk.Device, p_cache: vk.Pipe
             global_res = local_res;
         }
     }
+
     return global_res;
 }
 
@@ -798,7 +800,8 @@ pub export fn strollCreateDescriptorSetLayout(p_device: vk.Device, info: *const
         return .error_validation_failed;
     }
 
-    const allocator = VulkanAllocator.init(callbacks, .object).allocator();
+    // Device scoped because we're reference counting and layout may not be destroyed when vkDestroyDescriptorSetLayout is called
+    const allocator = VulkanAllocator.init(callbacks, .device).allocator();
     const device = Dispatchable(Device).fromHandleObject(p_device) catch |err| return toVkResult(err);
     const layout = device.createDescriptorSetLayout(allocator, info) catch |err| return toVkResult(err);
     p_layout.* = (NonDispatchable(DescriptorSetLayout).wrap(allocator, layout) catch |err| return toVkResult(err)).toVkHandle(vk.DescriptorSetLayout);
@@ -940,7 +943,8 @@ pub export fn strollCreatePipelineLayout(p_device: vk.Device, info: *const vk.Pi
         return .error_validation_failed;
     }
 
-    const allocator = VulkanAllocator.init(callbacks, .object).allocator();
+    // Device scoped because we're reference counting and layout may not be destroyed when vkDestroyPipelineLayout is called
+    const allocator = VulkanAllocator.init(callbacks, .device).allocator();
     const device = Dispatchable(Device).fromHandleObject(p_device) catch |err| return toVkResult(err);
     const layout = device.createPipelineLayout(allocator, info) catch |err| return toVkResult(err);
     p_layout.* = (NonDispatchable(PipelineLayout).wrap(allocator, layout) catch |err| return toVkResult(err)).toVkHandle(vk.PipelineLayout);
@@ -1425,7 +1429,7 @@ pub export fn strollGetPipelineCacheData(p_device: vk.Device, p_cache: vk.Pipeli
     _ = size;
     _ = data;
 
-    return .error_unknown;
+    return .success;
 }
 
 pub export fn strollGetQueryPoolResults(
@@ -1586,15 +1590,17 @@ pub export fn strollUpdateDescriptorSets(p_device: vk.Device, write_count: u32,
     entryPointBeginLogTrace(.vkUpdateDescriptorSets);
     defer entryPointEndLogTrace();
 
-    const device = Dispatchable(Device).fromHandleObject(p_device) catch |err| return errorLogger(err);
+    Dispatchable(Device).checkHandleValidity(p_device) catch |err| return errorLogger(err);
 
-    notImplementedWarning();
+    for (writes, 0..write_count) |write, _| {
+        const set = NonDispatchable(DescriptorSet).fromHandleObject(write.dst_set) catch |err| return errorLogger(err);
+        set.write(write) catch |err| return errorLogger(err);
+    }
 
-    _ = device;
+    for (copies, 0..copy_count) |copy, _| {
-    _ = write_count;
+        const set = NonDispatchable(DescriptorSet).fromHandleObject(copy.dst_set) catch |err| return errorLogger(err);
-    _ = writes;
+        set.copy(copy) catch |err| return errorLogger(err);
-    _ = copy_count;
+    }
-    _ = copies;
 }
 
 pub export fn strollWaitForFences(p_device: vk.Device, count: u32, p_fences: [*]const vk.Fence, waitForAll: vk.Bool32, timeout: u64) callconv(vk.vulkan_call_conv) vk.Result {
@@ -1668,17 +1674,9 @@ pub export fn strollCmdBindDescriptorSets(
     defer entryPointEndLogTrace();
 
     const cmd = Dispatchable(CommandBuffer).fromHandleObject(p_cmd) catch |err| return errorLogger(err);
+    cmd.bindDescriptorSets(bind_point, first, sets[0..count], dynamic_offsets[0..dynamic_offset_count]) catch |err| return errorLogger(err);
 
-    notImplementedWarning();
-
-    _ = cmd;
-    _ = bind_point;
     _ = layout;
-    _ = first;
-    _ = count;
-    _ = sets;
-    _ = dynamic_offsets;
-    _ = dynamic_offset_count;
 }
 
 pub export fn strollCmdBindIndexBuffer(p_cmd: vk.CommandBuffer, p_buffer: vk.Buffer, offset: vk.DeviceSize, index_type: vk.IndexType) callconv(vk.vulkan_call_conv) void {
@@ -1701,12 +1699,8 @@ pub export fn strollCmdBindPipeline(p_cmd: vk.CommandBuffer, bind_point: vk.Pipe
     defer entryPointEndLogTrace();
 
     const cmd = Dispatchable(CommandBuffer).fromHandleObject(p_cmd) catch |err| return errorLogger(err);
-
+    const pipeline = NonDispatchable(Pipeline).fromHandleObject(p_pipeline) catch |err| return errorLogger(err);
-    notImplementedWarning();
+    cmd.bindPipeline(bind_point, pipeline) catch |err| return errorLogger(err);
-
-    _ = cmd;
-    _ = bind_point;
-    _ = p_pipeline;
 }
 
 pub export fn strollCmdBindVertexBuffers(p_cmd: vk.CommandBuffer, first: u32, count: u32, p_buffers: [*]const vk.Buffer, offsets: [*]const vk.DeviceSize) callconv(vk.vulkan_call_conv) void {
@@ -1866,13 +1860,7 @@ pub export fn strollCmdDispatch(p_cmd: vk.CommandBuffer, group_count_x: u32, gro
     defer entryPointEndLogTrace();
 
     const cmd = Dispatchable(CommandBuffer).fromHandleObject(p_cmd) catch |err| return errorLogger(err);
-
+    cmd.dispatch(group_count_x, group_count_y, group_count_z) catch |err| return errorLogger(err);
-    notImplementedWarning();
-
-    _ = cmd;
-    _ = group_count_x;
-    _ = group_count_y;
-    _ = group_count_z;
 }
 
 pub export fn strollCmdDispatchIndirect(p_cmd: vk.CommandBuffer, p_buffer: vk.Buffer, offset: vk.DeviceSize) callconv(vk.vulkan_call_conv) void {

src/vulkan/logger/logger.zig

@@ -66,14 +66,21 @@ pub fn log(comptime level: std.log.Level, comptime scope: @Type(.enum_literal),
         .warn, .err => stderr_file,
     };
 
-    var buffer = std.mem.zeroes([512]u8);
-    var out_config = std.Io.tty.Config.detect(file);
-    var writer = std.Io.Writer.fixed(&buffer);
-
     var timezone = zdt.Timezone.tzLocal(std.heap.page_allocator) catch zdt.Timezone.UTC;
     defer timezone.deinit();
     const now = zdt.Datetime.now(.{ .tz = &timezone }) catch zdt.Datetime{};
 
+    var fmt_buffer = std.mem.zeroes([4096]u8);
+    var fmt_writer = std.Io.Writer.fixed(&fmt_buffer);
+    fmt_writer.print(format ++ "\n", args) catch {};
+    fmt_writer.flush() catch return;
+
+    var last_pos: usize = 0;
+    while (std.mem.indexOfScalarPos(u8, &fmt_buffer, last_pos, '\n')) |pos| {
+        var buffer = std.mem.zeroes([512]u8);
+        var out_config = std.Io.tty.Config.detect(file);
+        var writer = std.Io.Writer.fixed(&buffer);
+
         out_config.setColor(&writer, .magenta) catch {};
         writer.print("[StrollDriver ", .{}) catch {};
         if (!builtin.is_test) {
@@ -102,7 +109,8 @@ pub fn log(comptime level: std.log.Level, comptime scope: @Type(.enum_literal),
                 writer.print(">   ", .{}) catch {};
             }
         }
-    writer.print(format ++ "\n", args) catch {};
+
+        writer.print("{s}\n", .{fmt_buffer[last_pos..pos]}) catch {};
         writer.flush() catch return;
 
         if (level == .debug and lib.getLogVerboseLevel() == .Standard) {
@@ -127,4 +135,6 @@ pub fn log(comptime level: std.log.Level, comptime scope: @Type(.enum_literal),
             .info, .debug => _ = stdout_file.write(&buffer) catch {},
             .warn, .err => _ = stderr_file.write(&buffer) catch {},
         }
+        last_pos = pos + 1;
+    }
 }