diff --git a/src/intel/FlintPipeline.zig b/src/intel/FlintPipeline.zig
index bee7071..47cb99f 100644
--- a/src/intel/FlintPipeline.zig
+++ b/src/intel/FlintPipeline.zig
@@ -103,7 +103,7 @@ fn compileStage(allocator: std.mem.Allocator, info: *const vk.PipelineShaderStag
         .graphics => if (expected_stage == .compute) return VkError.ValidationFailed,
     }
 
-    const shader_module = base.NonDispatchable(base.ShaderModule).fromHandleObject(info.module) catch |err| return err;
+    const shader_module = try base.NonDispatchable(base.ShaderModule).fromHandleObject(info.module);
     var module = shader_module.instantiateIr(allocator, .{
         .entry_point = std.mem.span(info.p_name),
         .stage = expected_stage,
diff --git a/src/phi/PhiDevice.zig b/src/phi/PhiDevice.zig
index 7c02cf1..078e6e0 100644
--- a/src/phi/PhiDevice.zig
+++ b/src/phi/PhiDevice.zig
@@ -252,7 +252,9 @@ fn launchHostDaemon(instance: *base.Instance, allocator: std.mem.Allocator) VkEr
 
     const local_path = std.fmt.allocPrint(allocator, "/tmp/ape_phi_device_{d}_{d}.host", .{ process_id, thread_id }) catch return VkError.OutOfHostMemory;
     defer allocator.free(local_path);
-    errdefer std.Io.Dir.deleteFileAbsolute(io, local_path) catch {};
+    errdefer std.Io.Dir.deleteFileAbsolute(io, local_path) catch |err| {
+        std.log.scoped(.PhiDevice).warn("Failed to remove Phi host daemon after launch error: {s}", .{@errorName(err)});
+    };
 
     std.Io.Dir.writeFile(.cwd(), io, .{
         .sub_path = local_path,
diff --git a/src/software/SoftCommandBuffer.zig b/src/software/SoftCommandBuffer.zig
index dbf980e..7dd95f3 100644
--- a/src/software/SoftCommandBuffer.zig
+++ b/src/software/SoftCommandBuffer.zig
@@ -132,7 +132,7 @@ pub fn create(device: *base.Device, allocator: std.mem.Allocator, info: *const v
         .command_allocator = undefined,
         .commands = .empty,
     };
-    self.command_allocator = .init(interface.host_allocator.allocator());
+    self.command_allocator = .init(self.interface.host_allocator.allocator());
     return self;
 }
 
diff --git a/src/software/SoftInstance.zig b/src/software/SoftInstance.zig
index 3002029..427f171 100644
--- a/src/software/SoftInstance.zig
+++ b/src/software/SoftInstance.zig
@@ -49,6 +49,7 @@ pub fn create(allocator: std.mem.Allocator, infos: *const vk.InstanceCreateInfo)
         .requestPhysicalDevices = requestPhysicalDevices,
         .releasePhysicalDevices = releasePhysicalDevices,
         .io = io,
+        .enumerate_drm_devices = false,
     };
     return &self.interface;
 }
diff --git a/src/software/SoftPipeline.zig b/src/software/SoftPipeline.zig
index 8a1db99..27dce81 100644
--- a/src/software/SoftPipeline.zig
+++ b/src/software/SoftPipeline.zig
@@ -48,13 +48,14 @@ const Runtime = struct {
     rt: spv.Runtime,
 };
 
-const Shader = struct {
+const SpvShader = struct {
     module: *SoftShaderModule,
     runtimes: []Runtime,
     entry: []const u8,
-    interpreter: ?InterpreterShader,
 };
 
+const Shader = if (base.config.soft_ir_interpreter) InterpreterShader else SpvShader;
+
 const Stages = enum {
     vertex,
     tessellation_control,
@@ -107,7 +108,10 @@ pub fn createCompute(device: *base.Device, allocator: std.mem.Allocator, cache:
     };
 
     self.stages.put(.compute, try createShader(allocator, device_allocator, runtimes_allocator, soft_cache, soft_module, &info.stage, runtimes_count));
-    std.log.scoped(.ComputePipeline).debug("Created {d} runtimes for compute stage", .{runtimes_count});
+    std.log.scoped(.ComputePipeline).debug("Created {d} {s} runtimes for compute stage", .{
+        runtimes_count,
+        if (comptime base.config.soft_ir_interpreter) "IR" else "SPIR-V",
+    });
     return self;
 }
 
@@ -152,7 +156,10 @@ pub fn createGraphics(device: *base.Device, allocator: std.mem.Allocator, cache:
             const soft_module: *SoftShaderModule = @alignCast(@fieldParentPtr("interface", module));
             const shader = try createShader(allocator, device_allocator, runtimes_allocator, soft_cache, soft_module, &stage, runtimes_count);
 
-            std.log.scoped(.GraphicsPipeline).debug("Created {d} runtimes for:", .{runtimes_count});
+            std.log.scoped(.GraphicsPipeline).debug("Created {d} {s} runtimes for:", .{
+                runtimes_count,
+                if (comptime base.config.soft_ir_interpreter) "IR" else "SPIR-V",
+            });
 
             if (stage.stage.contains(.{ .vertex_bit = true })) {
                 std.log.scoped(.GraphicsPipeline).debug(">   Vertex stage", .{});
@@ -183,15 +190,18 @@ pub fn createGraphics(device: *base.Device, allocator: std.mem.Allocator, cache:
 
 pub fn destroy(interface: *Interface, allocator: std.mem.Allocator) void {
     const self: *Self = @alignCast(@fieldParentPtr("interface", interface));
-    const device_allocator = interface.owner.device_allocator.allocator();
 
     var it = self.stages.iterator();
-    while (it.next()) |entry| {
-        if (entry.value.interpreter) |*interpreter|
-            interpreter.deinit();
-        entry.value.module.unref(allocator);
-        for (entry.value.runtimes) |*runtime| {
-            runtime.rt.function_stack.clearAndFree(device_allocator); // Hacky to avoid leaks
+    if (comptime base.config.soft_ir_interpreter) {
+        while (it.next()) |entry|
+            entry.value.deinit();
+    } else {
+        const device_allocator = interface.owner.device_allocator.allocator();
+        while (it.next()) |entry| {
+            entry.value.module.unref(allocator);
+            for (entry.value.runtimes) |*runtime| {
+                runtime.rt.function_stack.clearAndFree(device_allocator); // Hacky to avoid leaks
+            }
         }
     }
     self.runtimes_allocator.deinit();
@@ -207,6 +217,9 @@ fn createShader(
     stage: *const vk.PipelineShaderStageCreateInfo,
     runtimes_count: usize,
 ) VkError!Shader {
+    if (comptime base.config.soft_ir_interpreter)
+        return InterpreterShader.compile(runtimes_allocator, module, stage, runtimes_count);
+
     const entry = std.mem.span(stage.p_name);
     const execution_model = executionModelForStage(stage.stage) orelse return VkError.Unknown;
     const runtimes = runtimes_allocator.alloc(Runtime, runtimes_count) catch return VkError.OutOfDeviceMemory;
@@ -259,15 +272,11 @@ fn createShader(
         }
     }
 
-    var shader: Shader = .{
+    return .{
         .module = module,
         .runtimes = runtimes,
         .entry = runtimes_allocator.dupe(u8, entry) catch return VkError.OutOfDeviceMemory,
-        .interpreter = null,
     };
-    if (comptime base.config.soft_ir_interpreter)
-        shader.interpreter = try InterpreterShader.compile(runtimes_allocator, module, stage, runtimes_count);
-    return shader;
 }
 
 fn initRuntime(allocator: std.mem.Allocator, module: *SoftShaderModule, stage: *const vk.PipelineShaderStageCreateInfo, image_api: spv.Runtime.ImageAPI) VkError!spv.Runtime {
diff --git a/src/software/device/ComputeDispatcher.zig b/src/software/device/ComputeDispatcher.zig
index ed7e8da..9b5f431 100644
--- a/src/software/device/ComputeDispatcher.zig
+++ b/src/software/device/ComputeDispatcher.zig
@@ -71,48 +71,49 @@ pub fn dispatchBase(self: *Self, base_group_x: u32, base_group_y: u32, base_grou
 
     const pipeline = self.state.pipeline orelse return VkError.InvalidPipelineDrv;
     const shader = pipeline.stages.getPtr(.compute) orelse return VkError.InvalidPipelineDrv;
-    if (comptime base.config.soft_ir_interpreter) {
-        if (shader.interpreter) |*interpreter_shader|
-            return ir_compute.dispatch(interpreter_shader, base_group_x, base_group_y, base_group_z, group_count_x, group_count_y, group_count_z);
-    }
-    const spv_module = &shader.module.module;
-    self.batch_size = if (spv_module.reflection_infos.has_atomics) 1 else shader.runtimes.len;
-
-    const allocator = self.device.interface.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);
 
     const io = self.device.interface.io();
     const timer = std.Io.Timestamp.now(io, .real);
     defer if (comptime base.config.logs != .none) {
         const duration = timer.untilNow(io, .real);
         const ms: f32 = @floatFromInt(duration.toMicroseconds());
-        std.log.scoped(.ComputeDispatcher).debug("Compute dispatch took {}ms", .{ms / 1000});
+        std.log.scoped(.ComputeDispatcher).debug("Compute dispatch took {}ms using {s} interpreter", .{ ms / 1000, if (comptime base.config.soft_ir_interpreter) "IR" else "SPIR-V" });
     };
 
-    var wg: std.Io.Group = .init;
-    for (0..@min(self.batch_size, group_count)) |batch_id| {
-        const run_data: RunData = .{
-            .self = self,
-            .batch_id = batch_id,
-            .group_count = group_count,
-            .base_group_x = @as(usize, @intCast(base_group_x)),
-            .base_group_y = @as(usize, @intCast(base_group_y)),
-            .base_group_z = @as(usize, @intCast(base_group_z)),
-            .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,
-            .local_size = local_size,
-            .pipeline = pipeline,
-        };
+    if (comptime base.config.soft_ir_interpreter) {
+        return ir_compute.dispatch(shader, base_group_x, base_group_y, base_group_z, group_count_x, group_count_y, group_count_z);
+    } else {
+        const spv_module = &shader.module.module;
+        self.batch_size = if (spv_module.reflection_infos.has_atomics) 1 else shader.runtimes.len;
 
-        wg.async(self.device.interface.io(), runWrapper, .{run_data});
+        const allocator = self.device.interface.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);
+
+        var wg: std.Io.Group = .init;
+        for (0..@min(self.batch_size, group_count)) |batch_id| {
+            const run_data: RunData = .{
+                .self = self,
+                .batch_id = batch_id,
+                .group_count = group_count,
+                .base_group_x = @as(usize, @intCast(base_group_x)),
+                .base_group_y = @as(usize, @intCast(base_group_y)),
+                .base_group_z = @as(usize, @intCast(base_group_z)),
+                .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,
+                .local_size = local_size,
+                .pipeline = pipeline,
+            };
+
+            wg.async(self.device.interface.io(), runWrapper, .{run_data});
+        }
+        wg.await(self.device.interface.io()) catch return VkError.DeviceLost;
     }
-    wg.await(self.device.interface.io()) catch return VkError.DeviceLost;
 }
 
 fn runWrapper(data: RunData) void {
@@ -138,17 +139,9 @@ inline fn run(data: RunData) !void {
 
     var barrier_runtimes: []spv.Runtime = &.{};
     var barrier_statuses: []spv.Runtime.EntryPointStatus = &.{};
-    if (uses_control_barrier) {
-        barrier_runtimes = try allocator.alloc(spv.Runtime, data.invocations_per_workgroup);
-        barrier_statuses = try allocator.alloc(spv.Runtime.EntryPointStatus, data.invocations_per_workgroup);
-        for (barrier_runtimes) |*barrier_rt| {
-            barrier_rt.* = try spv.Runtime.init(allocator, rt.mod, rt.image_api);
-            try barrier_rt.copySpecializationConstantsFrom(allocator, rt);
-        }
-    }
-
+    var initialized_barrier_runtimes: usize = 0;
     defer {
-        for (barrier_runtimes) |*barrier_rt| {
+        for (barrier_runtimes[0..initialized_barrier_runtimes]) |*barrier_rt| {
             barrier_rt.resetInvocation(allocator);
             barrier_rt.deinit(allocator);
         }
@@ -156,6 +149,19 @@ inline fn run(data: RunData) !void {
         allocator.free(barrier_statuses);
     }
 
+    if (uses_control_barrier) {
+        barrier_runtimes = try allocator.alloc(spv.Runtime, data.invocations_per_workgroup);
+        barrier_statuses = try allocator.alloc(spv.Runtime.EntryPointStatus, data.invocations_per_workgroup);
+        for (barrier_runtimes) |*barrier_rt| {
+            barrier_rt.* = try spv.Runtime.init(allocator, rt.mod, rt.image_api);
+            initialized_barrier_runtimes += 1;
+            try barrier_rt.copySpecializationConstantsFrom(allocator, rt);
+            try prepareRuntime(data.self, barrier_rt);
+        }
+    } else {
+        try prepareRuntime(data.self, rt);
+    }
+
     var group_index: usize = data.batch_id;
     while (group_index < data.group_count) : (group_index += data.self.batch_size) {
         var modulo: usize = group_index;
@@ -188,10 +194,6 @@ inline fn run(data: RunData) !void {
         defer rt.destroyWorkgroupMemory(allocator, workgroup_memory);
 
         rt.resetInvocation(allocator);
-        if (rt.specialization_constants.count() != 0)
-            try rt.applySpecializationInvocationLayout(allocator);
-        try ExecutionDevice.writeDescriptorSets(data.self.state, rt);
-        try rt.populatePushConstants(data.self.state.push_constant_blob[0..]);
         try rt.bindWorkgroupMemory(workgroup_memory);
         try setupWorkgroupBuiltins(data.self, rt, data.local_size, group_count_vec, group_id_vec);
 
@@ -228,6 +230,16 @@ inline fn run(data: RunData) !void {
     }
 }
 
+fn prepareRuntime(self: *Self, rt: *spv.Runtime) !void {
+    const allocator = self.device.interface.device_allocator.allocator();
+
+    rt.resetInvocation(allocator);
+    if (rt.specialization_constants.count() != 0)
+        try rt.applySpecializationInvocationLayout(allocator);
+    try ExecutionDevice.writeDescriptorSets(self.state, rt);
+    try rt.populatePushConstants(self.state.push_constant_blob[0..]);
+}
+
 fn runBarrierWorkgroup(
     data: RunData,
     runtimes: []spv.Runtime,
@@ -242,10 +254,6 @@ fn runBarrierWorkgroup(
     defer runtimes[0].destroyWorkgroupMemory(allocator, workgroup_memory);
     for (runtimes, 0..) |*rt, i| {
         rt.resetInvocation(allocator);
-        if (rt.specialization_constants.count() != 0)
-            try rt.applySpecializationInvocationLayout(allocator);
-        try ExecutionDevice.writeDescriptorSets(data.self.state, rt);
-        try rt.populatePushConstants(data.self.state.push_constant_blob[0..]);
         try rt.bindWorkgroupMemory(workgroup_memory);
         try setupWorkgroupBuiltins(data.self, rt, data.local_size, group_count, group_id);
         try setupSubgroupBuiltins(data.self, rt, data.local_size, group_id, i);
diff --git a/src/software/device/Renderer.zig b/src/software/device/Renderer.zig
index dcfd143..e14161b 100644
--- a/src/software/device/Renderer.zig
+++ b/src/software/device/Renderer.zig
@@ -246,14 +246,16 @@ fn drawCall(self: *Self, bounded_allocator: *BoundedAllocator, vertex_count: usi
     };
 
     const pipeline = self.state.pipeline orelse return VkError.InvalidPipelineDrv;
-    const vertex_shader = pipeline.stages.getPtrAssertContains(.vertex);
-    for (vertex_shader.runtimes[0..]) |*runtime| {
-        ExecutionDevice.writeDescriptorSets(self.state, &runtime.rt) catch return VkError.Unknown;
-    }
-    if (pipeline.stages.getPtr(.fragment)) |fragment_shader| {
-        for (fragment_shader.runtimes[0..]) |*runtime| {
+    if (comptime !base.config.soft_ir_interpreter) {
+        const vertex_shader = pipeline.stages.getPtrAssertContains(.vertex);
+        for (vertex_shader.runtimes) |*runtime| {
             ExecutionDevice.writeDescriptorSets(self.state, &runtime.rt) catch return VkError.Unknown;
         }
+        if (pipeline.stages.getPtr(.fragment)) |fragment_shader| {
+            for (fragment_shader.runtimes) |*runtime| {
+                ExecutionDevice.writeDescriptorSets(self.state, &runtime.rt) catch return VkError.Unknown;
+            }
+        }
     }
 
     self.vertexShaderStage(allocator, &draw_call, vertex_count, instance_count, first_vertex, first_instance, indices, primitive_restart) catch |err| {
diff --git a/src/software/device/fragment.zig b/src/software/device/fragment.zig
index cf75ef8..35f6641 100644
--- a/src/software/device/fragment.zig
+++ b/src/software/device/fragment.zig
@@ -36,6 +36,9 @@ pub fn shaderInvocation(
     inputs: [spv.SPIRV_MAX_OUTPUT_LOCATIONS]VertexInterpolationLocation,
     derivative_inputs: ?DerivativeInputs,
 ) SpvRuntimeError!InvocationResult {
+    if (comptime base.config.soft_ir_interpreter)
+        return SpvRuntimeError.InvalidSpirV;
+
     var fragment_inputs = inputs;
     errdefer freeOwnedInputs(allocator, fragment_inputs);
 
diff --git a/src/software/device/rasterizer.zig b/src/software/device/rasterizer.zig
index 4509812..d2c4820 100644
--- a/src/software/device/rasterizer.zig
+++ b/src/software/device/rasterizer.zig
@@ -447,7 +447,7 @@ fn rasterizeTransformedPoint(
     const point_min_x = vertex.position[0] - (point_size / 2.0);
     const point_min_y = vertex.position[1] - (point_size / 2.0);
     const pipeline = draw_call.renderer.state.pipeline orelse return;
-    const has_fragment_shader = pipeline.stages.getPtr(.fragment) != null;
+    const has_fragment_shader = if (comptime base.config.soft_ir_interpreter) false else pipeline.stages.getPtr(.fragment) != null;
 
     var py = min_y;
     while (py <= max_y) : (py += 1) {
diff --git a/src/software/device/rasterizer/bresenham.zig b/src/software/device/rasterizer/bresenham.zig
index 939f68d..3d1489d 100644
--- a/src/software/device/rasterizer/bresenham.zig
+++ b/src/software/device/rasterizer/bresenham.zig
@@ -134,7 +134,12 @@ fn drawLineBresenham(
 
     const pipeline = draw_call.renderer.state.pipeline orelse return;
     const fragment_stage = pipeline.stages.getPtr(.fragment);
-    const runtimes_count = if (fragment_stage) |stage| stage.runtimes.len else 1;
+    const runtimes_count = if (comptime base.config.soft_ir_interpreter)
+        1
+    else if (fragment_stage) |stage|
+        stage.runtimes.len
+    else
+        1;
     if (runtimes_count == 0)
         return;
 
@@ -170,7 +175,7 @@ fn drawLineBresenham(
             .color_attachment_access = color_attachment_access,
             .depth_attachment_access = depth_attachment_access,
             .stencil_attachment_access = stencil_attachment_access,
-            .has_fragment_shader = fragment_stage != null,
+            .has_fragment_shader = if (comptime base.config.soft_ir_interpreter) false else fragment_stage != null,
         };
 
         draw_call.rasterizer_wait_group.async(io, runWrapper, .{run_data});
@@ -191,7 +196,7 @@ fn drawLineDiamond(
 ) VkError!void {
     const pipeline = draw_call.renderer.state.pipeline orelse return;
     const fragment_stage = pipeline.stages.getPtr(.fragment);
-    const has_fragment_shader = fragment_stage != null;
+    const has_fragment_shader = if (comptime base.config.soft_ir_interpreter) false else fragment_stage != null;
     const batch_id: usize = 0;
 
     const min_x: i32 = @intFromFloat(@floor(@min(v0.position[0], v1.position[0]) - 1.0));
diff --git a/src/software/device/rasterizer/edge_function.zig b/src/software/device/rasterizer/edge_function.zig
index c8c8f46..d279578 100644
--- a/src/software/device/rasterizer/edge_function.zig
+++ b/src/software/device/rasterizer/edge_function.zig
@@ -89,24 +89,37 @@ pub fn drawTriangle(
     }
 
     const fragment_stage = pipeline.stages.getPtr(.fragment);
-    const fragment_uses_derivatives = if (fragment_stage) |stage|
+    const fragment_uses_derivatives = if (comptime base.config.soft_ir_interpreter)
+        false
+    else if (fragment_stage) |stage|
         stage.module.module.reflection_infos.needs_derivatives
     else
         false;
-    const early_fragment_tests = if (fragment_stage) |stage|
+    const early_fragment_tests = if (comptime base.config.soft_ir_interpreter)
+        false
+    else if (fragment_stage) |stage|
         stage.module.module.reflection_infos.early_fragment_tests
     else
         false;
-    const fragment_uses_sample_id = if (fragment_stage) |stage|
+    const fragment_uses_sample_id = if (comptime base.config.soft_ir_interpreter)
+        false
+    else if (fragment_stage) |stage|
         stage.module.module.builtins.get(.SampleId) != null
     else
         false;
-    const fragment_uses_centroid = if (fragment_stage) |stage|
+    const fragment_uses_centroid = if (comptime base.config.soft_ir_interpreter)
+        false
+    else if (fragment_stage) |stage|
         fragmentStageUsesInputDecoration(stage, .Centroid)
     else
         false;
 
-    const runtimes_count = if (fragment_stage) |stage| stage.runtimes.len else 1;
+    const runtimes_count = if (comptime base.config.soft_ir_interpreter)
+        1
+    else if (fragment_stage) |stage|
+        stage.runtimes.len
+    else
+        1;
     if (runtimes_count == 0)
         return;
     const sample_count = pipeline_data.multisample.rasterization_samples.toInt();
@@ -173,7 +186,7 @@ pub fn drawTriangle(
                 .depth_attachment_access = depth_attachment_access,
                 .stencil_attachment_access = stencil_attachment_access,
                 .front_face = front_face,
-                .has_fragment_shader = fragment_stage != null,
+                .has_fragment_shader = if (comptime base.config.soft_ir_interpreter) false else fragment_stage != null,
                 .early_fragment_tests = early_fragment_tests,
                 .fragment_uses_derivatives = fragment_uses_derivatives,
                 .fragment_uses_sample_id = fragment_uses_sample_id,
diff --git a/src/software/device/vertex_dispatcher.zig b/src/software/device/vertex_dispatcher.zig
index 73c1824..9eab481 100644
--- a/src/software/device/vertex_dispatcher.zig
+++ b/src/software/device/vertex_dispatcher.zig
@@ -43,26 +43,20 @@ pub fn runWrapper(data: RunData) void {
 inline fn run(data: RunData) !void {
     const shader = data.pipeline.stages.getPtrAssertContains(.vertex);
     if (comptime base.config.soft_ir_interpreter) {
-        // Interpolation decorations are not represented in the common IR yet,
-        // so fragment-linked graphics pipelines retain the SPIR-V path.
-        if (data.pipeline.stages.getPtr(.fragment) == null) {
-            if (shader.interpreter) |*interpreter_shader| {
-                return ir_vertex.run(
-                    data.allocator,
-                    data.pipeline,
-                    interpreter_shader,
-                    data.batch_id,
-                    data.batch_size,
-                    data.vertex_count,
-                    data.first_vertex,
-                    data.first_instance,
-                    data.indices,
-                    data.primitive_restart,
-                    data.instance_index,
-                    data.draw_call,
-                );
-            }
-        }
+        return ir_vertex.run(
+            data.allocator,
+            data.pipeline,
+            shader,
+            data.batch_id,
+            data.batch_size,
+            data.vertex_count,
+            data.first_vertex,
+            data.first_instance,
+            data.indices,
+            data.primitive_restart,
+            data.instance_index,
+            data.draw_call,
+        );
     }
     const runtime = &shader.runtimes[data.batch_id];
     const mutex = &runtime.mutex;
diff --git a/src/software/interpreter/Program.zig b/src/software/interpreter/Program.zig
index b331eef..7924cf9 100644
--- a/src/software/interpreter/Program.zig
+++ b/src/software/interpreter/Program.zig
@@ -47,24 +47,25 @@ interfaces: []const ?InterfaceBinding,
 pub fn compile(backing_allocator: std.mem.Allocator, module: *const module_ir.Module) !Self {
     try ir.validator.validate(module);
 
-    var result: Self = undefined;
-    result.arena = std.heap.ArenaAllocator.init(backing_allocator);
-    errdefer result.arena.deinit();
+    var arena = std.heap.ArenaAllocator.init(backing_allocator);
+    errdefer arena.deinit();
 
-    var lowerer = try Lowerer.init(result.arena.allocator(), module);
+    var lowerer = try Lowerer.init(arena.allocator(), module);
     try lowerer.lower();
 
-    result.stage = module.stage;
-    result.entry_pc = lowerer.entry_pc;
-    result.register_count = lowerer.register_count;
-    result.scratch_count = lowerer.scratch_count;
-    result.code = lowerer.code.items;
-    result.edges = lowerer.edges.items;
-    result.copies = lowerer.copies.items;
-    result.branches = lowerer.branches.items;
-    result.initializers = lowerer.initializers.items;
-    result.interfaces = lowerer.interfaces;
-    return result;
+    return .{
+        .arena = arena,
+        .stage = module.stage,
+        .entry_pc = lowerer.entry_pc,
+        .register_count = lowerer.register_count,
+        .scratch_count = lowerer.scratch_count,
+        .code = lowerer.code.items,
+        .edges = lowerer.edges.items,
+        .copies = lowerer.copies.items,
+        .branches = lowerer.branches.items,
+        .initializers = lowerer.initializers.items,
+        .interfaces = lowerer.interfaces,
+    };
 }
 
 pub fn deinit(self: *Self) void {
@@ -158,41 +159,32 @@ const Lowerer = struct {
 
     fn allocate(self: *Lowerer, type_id: ids.TypeId) !bc.Span {
         const ty = self.module.types.get(type_id) orelse return CompileError.UnsupportedType;
-        var kind: bc.ValueKind = undefined;
         var components: u8 = 1;
-
-        switch (ty.*) {
-            .boolean => kind = .boolean,
-            .integer => |integer| {
-                if (integer.bits != 32)
-                    return CompileError.UnsupportedType;
-
-                kind = if (integer.signedness == .signed) .signed_integer else .unsigned_integer;
-            },
-            .floating => |floating| {
-                if (floating.bits != 32)
-                    return CompileError.UnsupportedType;
-
-                kind = .floating;
-            },
-            .vector => |vector| {
+        const kind: bc.ValueKind = switch (ty.*) {
+            .boolean => .boolean,
+            .integer => |integer| if (integer.bits == 32)
+                if (integer.signedness == .signed) .signed_integer else .unsigned_integer
+            else
+                return CompileError.UnsupportedType,
+            .floating => |floating| if (floating.bits == 32)
+                .floating
+            else
+                return CompileError.UnsupportedType,
+            .vector => |vector| blk: {
                 const element = self.module.types.get(vector.element_type) orelse return CompileError.UnsupportedType;
                 components = vector.length;
-
-                kind = switch (element.*) {
+                break :blk switch (element.*) {
                     .boolean => .boolean,
-                    .integer => |integer| blk: {
-                        if (integer.bits != 32)
-                            return CompileError.UnsupportedType;
-
-                        break :blk if (integer.signedness == .signed) .signed_integer else .unsigned_integer;
-                    },
+                    .integer => |integer| if (integer.bits == 32)
+                        if (integer.signedness == .signed) .signed_integer else .unsigned_integer
+                    else
+                        return CompileError.UnsupportedType,
                     .floating => |floating| if (floating.bits == 32) .floating else return CompileError.UnsupportedType,
                     else => return CompileError.UnsupportedType,
                 };
             },
             else => return CompileError.UnsupportedType,
-        }
+        };
         const end = std.math.add(usize, self.register_count, components) catch return CompileError.TooManyRegisters;
 
         if (end > @as(usize, std.math.maxInt(bc.Register)) + 1)
diff --git a/src/software/interpreter/Shader.zig b/src/software/interpreter/Shader.zig
index 6895bb5..33b1c0f 100644
--- a/src/software/interpreter/Shader.zig
+++ b/src/software/interpreter/Shader.zig
@@ -21,17 +21,15 @@ program: Program,
 runtimes: []RuntimeSlot,
 workgroup_size: ?[3]u32,
 
-/// Compiles a stage when the current interpreter can execute its complete
-/// interface. `null` deliberately selects the existing SPIR-V runtime.
-pub fn compile(
-    allocator: std.mem.Allocator,
-    module: *SoftShaderModule,
-    stage: *const vk.PipelineShaderStageCreateInfo,
-    runtime_count: usize,
-) VkError!?Self {
-    const expected_stage = commonStage(stage.stage) orelse return null;
-    if (expected_stage == .fragment)
-        return null;
+pub fn compile(allocator: std.mem.Allocator, module: *SoftShaderModule, stage: *const vk.PipelineShaderStageCreateInfo, runtime_count: usize) VkError!Self {
+    const expected_stage = commonStage(stage.stage) orelse {
+        std.log.scoped(.IrInterpreter).err("unsupported shader stage", .{});
+        return VkError.ValidationFailed;
+    };
+    if (expected_stage == .fragment) {
+        std.log.scoped(.IrInterpreter).err("fragment shaders are not supported", .{});
+        return VkError.ValidationFailed;
+    }
 
     const specializations = try specializationValues(allocator, stage.p_specialization_info);
     defer if (specializations.len != 0) allocator.free(specializations);
@@ -43,25 +41,24 @@ pub fn compile(
     }) catch |err| {
         if (err == error.OutOfMemory)
             return VkError.OutOfDeviceMemory;
-        std.log.scoped(.SoftIrInterpreter).debug("IR translation fallback: {s}", .{@errorName(err)});
-        return null;
+        std.log.scoped(.IrInterpreter).err("IR translation failed: {s}", .{@errorName(err)});
+        return VkError.ValidationFailed;
     };
     defer module_ir.deinit();
 
     var program = Program.compile(allocator, &module_ir) catch |err| {
         if (err == error.OutOfMemory)
             return VkError.OutOfDeviceMemory;
-        std.log.scoped(.SoftIrInterpreter).debug("bytecode lowering fallback: {s}", .{@errorName(err)});
-        return null;
+        std.log.scoped(.IrInterpreter).err("bytecode lowering failed: {s}", .{@errorName(err)});
+        return VkError.ValidationFailed;
     };
     errdefer program.deinit();
 
     if (!hasCompatibleInterface(&program, expected_stage) or
         (expected_stage == .compute and module_ir.execution_modes.workgroup_size == null))
     {
-        std.log.scoped(.SoftIrInterpreter).debug("stage interface or execution modes require the SPIR-V runtime", .{});
-        program.deinit();
-        return null;
+        std.log.scoped(.IrInterpreter).err("unsupported stage interface or execution modes", .{});
+        return VkError.ValidationFailed;
     }
 
     const runtimes = allocator.alloc(RuntimeSlot, runtime_count) catch return VkError.OutOfDeviceMemory;
@@ -76,7 +73,7 @@ pub fn compile(
         initialized += 1;
     }
 
-    std.log.scoped(.SoftIrInterpreter).debug("compiled {s} stage to {d} bytecode instructions", .{
+    std.log.scoped(.IrInterpreter).debug("compiled {s} stage to {d} bytecode instructions", .{
         @tagName(expected_stage),
         program.code.len,
     });
diff --git a/src/software/interpreter/vertex.zig b/src/software/interpreter/vertex.zig
index b11b845..7046433 100644
--- a/src/software/interpreter/vertex.zig
+++ b/src/software/interpreter/vertex.zig
@@ -3,7 +3,6 @@ const vk = @import("vulkan");
 const base = @import("base");
 const shader_ir = @import("shader_ir");
 
-const bc = @import("bytecode.zig");
 const Shader = @import("Shader.zig");
 const SoftPipeline = @import("../SoftPipeline.zig");
 const Renderer = @import("../device/Renderer.zig");
@@ -58,14 +57,7 @@ pub fn run(
     }
 }
 
-fn populateInputs(
-    runtime: anytype,
-    program: *const @import("Program.zig"),
-    pipeline: *SoftPipeline,
-    draw_call: *Renderer.DrawCall,
-    vertex_index: u32,
-    instance_index: u32,
-) VkError!void {
+fn populateInputs(runtime: anytype, program: *const @import("Program.zig"), pipeline: *SoftPipeline, draw_call: *Renderer.DrawCall, vertex_index: u32, instance_index: u32) VkError!void {
     for (program.interfaces, 0..) |optional_binding, index| {
         const binding = optional_binding orelse continue;
         if (binding.direction != .input)
@@ -167,7 +159,3 @@ fn findBinding(bindings: []const vk.VertexInputBindingDescription, binding: u32)
         if (description.binding == binding) return description;
     return null;
 }
-
-comptime {
-    _ = bc;
-}
diff --git a/src/vulkan/Instance.zig b/src/vulkan/Instance.zig
index 208765b..c54bfa4 100644
--- a/src/vulkan/Instance.zig
+++ b/src/vulkan/Instance.zig
@@ -34,6 +34,7 @@ pub const VTable = struct {
     releasePhysicalDevices: *const fn (*Self, std.mem.Allocator) VkError!void,
     requestPhysicalDevices: *const fn (*Self, std.mem.Allocator, []lib.drm.Card) VkError!void,
     io: *const fn (*Self) std.Io,
+    enumerate_drm_devices: bool = true,
 };
 
 pub const DispatchTable = struct {
@@ -154,10 +155,13 @@ pub fn releasePhysicalDevices(self: *Self, allocator: std.mem.Allocator) VkError
 }
 
 pub fn requestPhysicalDevices(self: *Self, allocator: std.mem.Allocator) VkError!void {
-    const devices = try drm.enumerateDrmPhysicalDevices(allocator, self);
-    defer allocator.free(devices);
-
-    try self.vtable.requestPhysicalDevices(self, allocator, devices);
+    if (self.vtable.enumerate_drm_devices) {
+        const devices = try drm.enumerateDrmPhysicalDevices(allocator, self);
+        defer allocator.free(devices);
+        try self.vtable.requestPhysicalDevices(self, allocator, devices);
+    } else {
+        try self.vtable.requestPhysicalDevices(self, allocator, &.{});
+    }
 
     if (self.physical_devices.items.len == 0) {
         std.log.scoped(.vkCreateInstance).err("No VkPhysicalDevice found", .{});