[IR] adding external resources managements

[Soft] adding descriptor sets management

build.zig

@@ -218,8 +218,10 @@ pub fn build(b: *std.Build) !void {
         test_step.dependOn(&run_tests.step);
 
         inline for (std.enums.values(RunningMode)) |mode| {
-            (try addCTS(b, target, &impl, lib, mode)).dependOn(&lib_install.step);
+            if (addCTS(b, target, &impl, lib, mode) catch null) |step|
-            (try addMultithreadedCTS(b, target, &impl, lib, mode)).dependOn(&lib_install.step);
+                step.dependOn(&lib_install.step);
+            if (addMultithreadedCTS(b, target, &impl, lib, mode) catch null) |step|
+                step.dependOn(&lib_install.step);
         }
 
         const impl_autodoc_test = b.addObject(.{

src/compiler/README.md

@@ -30,7 +30,7 @@ The printer uses these prefixes:
 | Prefix  | Meaning                                                          | Example               |
 | ------- | ---------------------------------------------------------------- | --------------------- |
 | `%id`   | An SSA value, whether constant, parameter, or instruction result | `%3`, `%merged_value` |
-| `@name` | A function or interface declaration                              | `@main`, `@out_color` |
+| `@name` | A function, interface, or resource declaration                   | `@main`, `@out_color` |
 | `.name` | A basic block                                                    | `.entry`, `.merge`    |
 | `#N`    | A constant-store identity used within composite constants        | `#2`                  |
 
@@ -55,6 +55,7 @@ The outer structure has this shape:
 shader <stage> @<entry-point>
 {
     <interface declarations>
+    <resource declarations>
     <constant declarations>
 
     fn @<name>(<parameters>) -> <type>
@@ -67,8 +68,8 @@ shader <stage> @<entry-point>
 }
 ```
 
-Execution modes, resources, source locations, and structured-control metadata
+Execution modes, source locations, and structured-control metadata exist in
-exist in memory, but the printer does not display them yet.
+memory, but the printer does not display them yet.
 
 ## Parsing
 
@@ -113,6 +114,25 @@ The current resource kinds are `uniform_buffer`, `storage_buffer`,
 `sampled_image`, `storage_image`, and `sampler`. A resource handle may also
 carry an optional data type in memory, although the printer omits that type.
 
+## Resources
+
+Resources are declared at module scope. Every `ResourceKind` uses the same
+`set` and `binding` syntax:
+
+```text
+@name: TYPE = storage_buffer[set(N), binding(N)]
+```
+
+The declaration `TYPE` is the storage buffer's block or payload aggregate type;
+it does not constrain the type of each byte-addressed access. Both operations
+accept storage buffers, and byte offsets must have a scalar unsigned integer
+type. Access values may be integer or floating-point scalars or vectors thereof.
+
+```text
+%value: TYPE = load_buffer @name, %offset
+store_buffer @name, %offset, %value
+```
+
 ## Constants
 
 Constants live at module scope and also have ordinary numeric or named `%id` value identities.
@@ -188,8 +208,8 @@ metadata. They are not terminators and do not create graph edges themselves.
 ## Common instruction rules
 
 An instruction belongs to one block, has zero or one result, and may carry a
-source location. Except for `store_interface` and `call`, current operations are
+source location. Except for `store_interface`, `store_buffer`, and `call`, current
-treated as side-effect free by the rewriter. A block's terminator is stored
+operations are treated as side-effect free by the rewriter. A block's terminator is stored
 separately from its ordinary instructions.
 
 Most arithmetic operations are intended for scalars or vectors of their named
@@ -379,6 +399,28 @@ The stored value must equal the interface variable's type. As with
 `load_interface`, an optional unprinted `element_index` is reserved for later
 arrayed-interface work. This operation has side effects.
 
+### `load_buffer`
+
+Reads a numeric scalar or vector at an explicit byte offset. The resource must
+be a storage buffer, the byte offset must be a scalar unsigned integer, and the
+instruction must have a result. The result type is independent of the resource's
+block or payload aggregate type.
+
+```text
+%value: u32 = load_buffer @data, %offset
+```
+
+### `store_buffer`
+
+Writes a numeric scalar or vector to a storage buffer at an explicit byte
+offset. It produces no SSA result, and both the unsigned integer offset and
+stored value are ordinary value uses. The value type is independent of the
+resource's block or payload aggregate type. This operation has side effects.
+
+```text
+store_buffer @data, %offset, %value
+```
+
 ### `call`
 
 Invokes another IR function. Arguments must match the callee's parameters in

src/compiler/ir/Builder.zig

@@ -169,6 +169,16 @@ pub fn addInterfaceVariable(
     });
 }
 
+pub fn addResource(self: *Self, ty: ids.TypeId, kind: type_ir.ResourceKind, set: u32, binding: u32, name: ?[]const u8) !ids.ResourceId {
+    return self.module.resources.add(self.module.allocator(), .{
+        .kind = kind,
+        .set = set,
+        .binding = binding,
+        .type = ty,
+        .name = try self.copyName(name),
+    });
+}
+
 pub fn edge(self: *Self, target: ids.BlockId, arguments: []const ids.ValueId) !module_ir.Edge {
     return .{
         .target = target,

src/compiler/ir/instruction.zig

@@ -6,6 +6,7 @@ pub const ValueId = ids.ValueId;
 pub const BlockId = ids.BlockId;
 pub const FunctionId = ids.FunctionId;
 pub const InterfaceVariableId = ids.InterfaceVariableId;
+pub const ResourceId = ids.ResourceId;
 
 pub const SourceLocation = struct {
     file: ?[]const u8 = null,
@@ -98,6 +99,17 @@ pub const StoreInterface = struct {
     element_index: ?ValueId = null,
 };
 
+pub const LoadBuffer = struct {
+    resource: ResourceId,
+    byte_offset: ValueId,
+};
+
+pub const StoreBuffer = struct {
+    resource: ResourceId,
+    byte_offset: ValueId,
+    value: ValueId,
+};
+
 pub const Call = struct {
     function: FunctionId,
     arguments: []const ValueId,
@@ -113,6 +125,8 @@ pub const Operation = union(enum) {
     composite_extract: CompositeExtract,
     load_interface: LoadInterface,
     store_interface: StoreInterface,
+    load_buffer: LoadBuffer,
+    store_buffer: StoreBuffer,
     call: Call,
 
     pub fn visitValueUses(self: Operation, context: anytype, comptime visitor: anytype) void {
@@ -140,6 +154,11 @@ pub const Operation = union(enum) {
                 if (op.element_index) |index|
                     visitor(context, index);
             },
+            .load_buffer => |op| visitor(context, op.byte_offset),
+            .store_buffer => |op| {
+                visitor(context, op.byte_offset);
+                visitor(context, op.value);
+            },
             .call => |op| {
                 for (op.arguments) |argument|
                     visitor(context, argument);
@@ -176,6 +195,11 @@ pub const Operation = union(enum) {
                 if (op.element_index) |*index|
                     replaceOne(index, old, replacement, &count);
             },
+            .load_buffer => |*op| replaceOne(&op.byte_offset, old, replacement, &count),
+            .store_buffer => |*op| {
+                replaceOne(&op.byte_offset, old, replacement, &count);
+                replaceOne(&op.value, old, replacement, &count);
+            },
             .call => |*op| op.arguments = try replaceSlice(allocator, op.arguments, old, replacement, &count),
         }
         return count;
@@ -183,7 +207,11 @@ pub const Operation = union(enum) {
 
     pub fn hasSideEffects(self: Operation) bool {
         return switch (self) {
-            .store_interface, .call => true,
+            .store_interface,
+            .store_buffer,
+            .call,
+            => true,
+
             else => false,
         };
     }

src/compiler/ir/parser/ast.zig

@@ -2,12 +2,14 @@ const std = @import("std");
 const ids = @import("../id.zig");
 const inst_ir = @import("../instruction.zig");
 const module_ir = @import("../module.zig");
+const type_ir = @import("../type.zig");
 
 pub const ValueRef = []const u8;
 
 pub const ParsedModule = struct {
     entry_point_name: ?[]const u8,
     interfaces: std.ArrayList(ParsedInterface) = .empty,
+    resources: std.ArrayList(ParsedResource) = .empty,
     constants: std.ArrayList(ParsedConstant) = .empty,
     functions: std.ArrayList(ParsedFunction) = .empty,
 };
@@ -19,6 +21,14 @@ pub const ParsedInterface = struct {
     semantic: module_ir.InterfaceSemantic,
 };
 
+pub const ParsedResource = struct {
+    kind: type_ir.ResourceKind,
+    name: []const u8,
+    ty: ids.TypeId,
+    set: u32,
+    binding: u32,
+};
+
 pub const ParsedConstantValue = union(enum) {
     boolean: bool,
     integer_bits: u64,
@@ -89,5 +99,7 @@ pub const ParsedOperation = union(enum) {
     composite_extract: struct { composite: ValueRef, indices: []const u32 },
     load_interface: []const u8,
     store_interface: struct { interface_name: []const u8, value: ValueRef },
+    load_buffer: struct { resource_name: []const u8, byte_offset: ValueRef },
+    store_buffer: struct { resource_name: []const u8, byte_offset: ValueRef, value: ValueRef },
     call: struct { function_name: []const u8, arguments: []const ValueRef },
 };

src/compiler/ir/parser/lower.zig

@@ -22,6 +22,7 @@ pub fn lower(allocator: std.mem.Allocator, module: *module_ir.Module, parsed: *P
     var values: std.StringHashMapUnmanaged(ids.ValueId) = .empty;
     var constants: std.AutoHashMapUnmanaged(u32, ids.ConstantId) = .empty;
     var interfaces: std.StringHashMapUnmanaged(ids.InterfaceVariableId) = .empty;
+    var resources: std.StringHashMapUnmanaged(ids.ResourceId) = .empty;
     var functions: std.StringHashMapUnmanaged(ids.FunctionId) = .empty;
 
     for (parsed.interfaces.items) |interface| {
@@ -32,6 +33,14 @@ pub fn lower(allocator: std.mem.Allocator, module: *module_ir.Module, parsed: *P
         try interfaces.put(allocator, interface.name, id);
     }
 
+    for (parsed.resources.items) |resource| {
+        if (resources.contains(resource.name))
+            return error.DuplicateName;
+
+        const id = try builder.addResource(resource.ty, resource.kind, resource.set, resource.binding, resource.name);
+        try resources.put(allocator, resource.name, id);
+    }
+
     for (parsed.constants.items, 0..) |constant, constant_index| {
         const value: constant_ir.ConstantValue = switch (constant.value) {
             .boolean => |item| .{ .boolean = item },
@@ -94,7 +103,7 @@ pub fn lower(allocator: std.mem.Allocator, module: *module_ir.Module, parsed: *P
 
         for (function.blocks.items) |block| {
             for (block.instructions.items) |instruction| {
-                const lowered = try lowerOperation(allocator, module, &values, &interfaces, &functions, instruction.operation);
+                const lowered = try lowerOperation(allocator, module, &values, &interfaces, &resources, &functions, instruction.operation);
                 const result_type = instruction.result_type orelse lowered.inferred_type;
 
                 if (instruction.printed_result != null and result_type == null)
@@ -125,6 +134,7 @@ fn lowerOperation(
     module: *module_ir.Module,
     values: *const std.StringHashMapUnmanaged(ids.ValueId),
     interfaces: *const std.StringHashMapUnmanaged(ids.InterfaceVariableId),
+    resources: *const std.StringHashMapUnmanaged(ids.ResourceId),
     functions: *const std.StringHashMapUnmanaged(ids.FunctionId),
     parsed: ParsedOperation,
 ) !LoweredOperation {
@@ -249,6 +259,32 @@ fn lowerOperation(
                 .inferred_type = null,
             };
         },
+        .load_buffer => |op| blk: {
+            const resource_id = resources.get(op.resource_name) orelse return error.UnknownResource;
+            const byte_offset = resolveValue(values, op.byte_offset) orelse return error.UnknownValue;
+
+            break :blk .{
+                .operation = .{ .load_buffer = .{
+                    .resource = resource_id,
+                    .byte_offset = byte_offset,
+                } },
+                .inferred_type = null,
+            };
+        },
+        .store_buffer => |op| blk: {
+            const resource_id = resources.get(op.resource_name) orelse return error.UnknownResource;
+            const byte_offset = resolveValue(values, op.byte_offset) orelse return error.UnknownValue;
+            const value = resolveValue(values, op.value) orelse return error.UnknownValue;
+
+            break :blk .{
+                .operation = .{ .store_buffer = .{
+                    .resource = resource_id,
+                    .byte_offset = byte_offset,
+                    .value = value,
+                } },
+                .inferred_type = null,
+            };
+        },
         .call => |op| blk: {
             const function_id = functions.get(op.function_name) orelse return error.UnknownFunction;
             var arguments: std.ArrayList(ids.ValueId) = .empty;

src/compiler/ir/parser/parser.zig

@@ -26,6 +26,7 @@ pub const Error = error{
     UnknownConstant,
     UnknownFunction,
     UnknownInterface,
+    UnknownResource,
     UnknownValue,
 };
 
@@ -34,6 +35,7 @@ pub const max_file_size = 64 * 1024 * 1024;
 const ValueRef = ast.ValueRef;
 const ParsedModule = ast.ParsedModule;
 const ParsedInterface = ast.ParsedInterface;
+const ParsedResource = ast.ParsedResource;
 const ParsedConstantValue = ast.ParsedConstantValue;
 const ParsedConstant = ast.ParsedConstant;
 const ParsedParameter = ast.ParsedParameter;
@@ -46,22 +48,27 @@ const ParsedOperation = ast.ParsedOperation;
 const Token = Lexer.Token;
 const TokenTag = Lexer.TokenTag;
 
+const ParsedDeclaration = union(enum) {
+    interface: ParsedInterface,
+    resource: ParsedResource,
+};
+
 const Parser = struct {
     lexer: Lexer,
     allocator: std.mem.Allocator,
     module: ?*module_ir.Module = null,
 
-    fn parseInterface(self: *Parser) !ParsedInterface {
+    fn parseDeclaration(self: *Parser) !ParsedDeclaration {
         const name = (try self.expect(.at_name)).text;
 
         try self.expectDiscard(.colon);
         const ty = try self.parseType();
         try self.expectDiscard(.equal);
 
-        const direction_token = try self.expect(.identifier);
+        const kind_token = try self.expect(.identifier);
-        const direction = std.meta.stringToEnum(module_ir.InterfaceDirection, direction_token.text) orelse return Error.InvalidSemantic;
         try self.expectDiscard(.left_square);
 
+        if (std.meta.stringToEnum(module_ir.InterfaceDirection, kind_token.text)) |direction| {
             const semantic_name = (try self.expect(.identifier)).text;
             const semantic: module_ir.InterfaceSemantic = if (std.mem.eql(u8, semantic_name, "location")) blk: {
                 try self.expectDiscard(.left_paren);
@@ -98,12 +105,34 @@ const Parser = struct {
 
             try self.expectDiscard(.right_square);
 
-        return .{
+            return .{ .interface = .{
                 .direction = direction,
                 .name = name,
                 .ty = ty,
                 .semantic = semantic,
-        };
+            } };
+        }
+
+        const kind = std.meta.stringToEnum(type_ir.ResourceKind, kind_token.text) orelse return Error.InvalidSemantic;
+        try self.expectIdentifier("set");
+        try self.expectDiscard(.left_paren);
+        const set = try self.parseUnsigned(u32, .number);
+        try self.expectDiscard(.right_paren);
+
+        try self.expectDiscard(.comma);
+        try self.expectIdentifier("binding");
+        try self.expectDiscard(.left_paren);
+        const binding = try self.parseUnsigned(u32, .number);
+        try self.expectDiscard(.right_paren);
+        try self.expectDiscard(.right_square);
+
+        return .{ .resource = .{
+            .kind = kind,
+            .name = name,
+            .ty = ty,
+            .set = set,
+            .binding = binding,
+        } };
     }
 
     fn parseConstant(self: *Parser) !ParsedConstant {
@@ -362,6 +391,27 @@ const Parser = struct {
             };
         }
 
+        if (std.mem.eql(u8, name, "load_buffer")) {
+            const resource_name = (try self.expect(.at_name)).text;
+            try self.expectDiscard(.comma);
+            return .{ .load_buffer = .{
+                .resource_name = resource_name,
+                .byte_offset = try self.parseValueRef(),
+            } };
+        }
+
+        if (std.mem.eql(u8, name, "store_buffer")) {
+            const resource_name = (try self.expect(.at_name)).text;
+            try self.expectDiscard(.comma);
+            const byte_offset = try self.parseValueRef();
+            try self.expectDiscard(.comma);
+            return .{ .store_buffer = .{
+                .resource_name = resource_name,
+                .byte_offset = byte_offset,
+                .value = try self.parseValueRef(),
+            } };
+        }
+
         if (std.mem.eql(u8, name, "call")) {
             const function_name = (try self.expect(.at_name)).text;
             try self.expectDiscard(.left_paren);
@@ -730,7 +780,10 @@ pub fn parseString(backing_allocator: std.mem.Allocator, source: []const u8) !mo
         const token = try parser.peek();
         switch (token.tag) {
             .value_ref => try parsed.constants.append(temporary_allocator, try parser.parseConstant()),
-            .at_name => try parsed.interfaces.append(temporary_allocator, try parser.parseInterface()),
+            .at_name => switch (try parser.parseDeclaration()) {
+                .interface => |interface| try parsed.interfaces.append(temporary_allocator, interface),
+                .resource => |resource| try parsed.resources.append(temporary_allocator, resource),
+            },
             .identifier => {
                 if (std.mem.eql(u8, token.text, "fn")) {
                     try parsed.functions.append(temporary_allocator, try parser.parseFunction());
@@ -794,6 +847,50 @@ test "Parser: interface" {
     try std.testing.expect(std.mem.indexOf(u8, printed, "@position: vec4[f32] = output[builtin(position)]") != null);
 }
 
+test "Parser: resources and buffer operations" {
+    const printer = @import("../printer.zig");
+
+    const source =
+        \\ shader compute @main
+        \\ {
+        \\     @uniforms: u32 = uniform_buffer[set(0), binding(1)]
+        \\     @storage: struct[u32, f32] = storage_buffer[set(2), binding(3)]
+        \\     @texture: vec4[f32] = sampled_image[set(4), binding(5)]
+        \\     @image: vec4[f32] = storage_image[set(6), binding(7)]
+        \\     @linear_sampler: resourceHandle[sampler] = sampler[set(8), binding(9)]
+        \\     %offset: constant u32 = 4
+        \\     %value: constant u32 = 7
+        \\
+        \\     fn @main() -> void
+        \\     {
+        \\         .entry():
+        \\             %storage_value: vec2[f32] = load_buffer @storage, %offset
+        \\             store_buffer @storage, %offset, %value
+        \\             return
+        \\     }
+        \\ }
+    ;
+
+    var module = try parseString(std.testing.allocator, source);
+    defer module.deinit();
+    const printed = try printer.allocPrint(std.testing.allocator, &module);
+    defer std.testing.allocator.free(printed);
+
+    try std.testing.expect(std.mem.indexOf(u8, printed, "@uniforms: u32 = uniform_buffer[set(0), binding(1)]") != null);
+    try std.testing.expect(std.mem.indexOf(u8, printed, "@storage: struct[u32, f32] = storage_buffer[set(2), binding(3)]") != null);
+    try std.testing.expect(std.mem.indexOf(u8, printed, "@texture: vec4[f32] = sampled_image[set(4), binding(5)]") != null);
+    try std.testing.expect(std.mem.indexOf(u8, printed, "@image: vec4[f32] = storage_image[set(6), binding(7)]") != null);
+    try std.testing.expect(std.mem.indexOf(u8, printed, "@linear_sampler: resourceHandle[sampler] = sampler[set(8), binding(9)]") != null);
+    try std.testing.expect(std.mem.indexOf(u8, printed, "%storage_value: vec2[f32] = load_buffer @storage, %offset") != null);
+    try std.testing.expect(std.mem.indexOf(u8, printed, "store_buffer @storage, %offset, %value") != null);
+
+    var reparsed = try parseString(std.testing.allocator, printed);
+    defer reparsed.deinit();
+    const printed_again = try printer.allocPrint(std.testing.allocator, &reparsed);
+    defer std.testing.allocator.free(printed_again);
+    try std.testing.expectEqualStrings(printed, printed_again);
+}
+
 test "Parser: types, operations, calls, terminators" {
     const printer = @import("../printer.zig");
 

src/compiler/ir/printer.zig

@@ -32,6 +32,16 @@ pub fn write(module: *const module_ir.Module, writer: *std.Io.Writer) std.Io.Wri
         try writer.writeAll("]\n");
     }
 
+    for (module.resources.entries.items, 0..) |entry, index| {
+        const resource = entry orelse continue;
+
+        try writer.writeAll(indent);
+        try writeNamedRef(writer, resource.name, "resource", index);
+        try writer.writeAll(": ");
+        try writeType(module, writer, resource.type);
+        try writer.print(" = {t}[set({d}), binding({d})]\n", .{ resource.kind, resource.set, resource.binding });
+    }
+
     for (module.constants.entries.items, 0..) |entry, constant_index| {
         const constant = entry orelse continue;
         const value_id = constantValueId(module, ids.ConstantId.fromIndex(constant_index)) orelse continue;
@@ -214,6 +224,22 @@ fn writeOperation(module: *const module_ir.Module, writer: *std.Io.Writer, opera
             try writer.writeAll(", ");
             try writeValueRef(module, writer, op.value);
         },
+        .load_buffer => |op| {
+            try writer.writeAll("load_buffer ");
+            const resource = module.resources.get(op.resource);
+            try writeNamedRef(writer, if (resource) |r| r.name else null, "resource", op.resource.index());
+            try writer.writeAll(", ");
+            try writeValueRef(module, writer, op.byte_offset);
+        },
+        .store_buffer => |op| {
+            try writer.writeAll("store_buffer ");
+            const resource = module.resources.get(op.resource);
+            try writeNamedRef(writer, if (resource) |r| r.name else null, "resource", op.resource.index());
+            try writer.writeAll(", ");
+            try writeValueRef(module, writer, op.byte_offset);
+            try writer.writeAll(", ");
+            try writeValueRef(module, writer, op.value);
+        },
         .call => |op| {
             try writer.writeAll("call ");
             try writeFunctionRef(module, writer, op.function);

src/compiler/ir/transformers/inline_all_functions.zig

@@ -381,6 +381,15 @@ fn remapOperation(
             .value = try mappedValue(module, value_map, op.value),
             .element_index = if (op.element_index) |index| try mappedValue(module, value_map, index) else null,
         } },
+        .load_buffer => |op| .{ .load_buffer = .{
+            .resource = op.resource,
+            .byte_offset = try mappedValue(module, value_map, op.byte_offset),
+        } },
+        .store_buffer => |op| .{ .store_buffer = .{
+            .resource = op.resource,
+            .byte_offset = try mappedValue(module, value_map, op.byte_offset),
+            .value = try mappedValue(module, value_map, op.value),
+        } },
         .call => Error.InvalidModule,
     };
 }

src/compiler/ir/validator/validator.zig

@@ -28,6 +28,7 @@ pub const ValidationError = error{
     WrongOperandType,
     WrongParameterIndex,
     WrongParent,
+    WrongResourceKind,
     WrongResultPresence,
     WrongResultType,
     WrongReturnType,
@@ -329,6 +330,32 @@ fn validateOperation(module: *const module_ir.Module, function_id: ids.FunctionI
             if (op.element_index) |index|
                 _ = try operandType(module, function_id, index);
         },
+        .load_buffer => |op| {
+            const resource = module.resources.get(op.resource) orelse return ValidationError.InvalidValue;
+            if (resource.kind != .storage_buffer)
+                return ValidationError.WrongResourceKind;
+
+            if (!isUnsignedInteger(module, try operandType(module, function_id, op.byte_offset)))
+                return ValidationError.WrongOperandType;
+
+            const result = result_type orelse return ValidationError.WrongResultPresence;
+            if (!isBufferAccessibleType(module, result))
+                return ValidationError.WrongResultType;
+        },
+        .store_buffer => |op| {
+            if (result_type != null)
+                return ValidationError.WrongResultPresence;
+
+            const resource = module.resources.get(op.resource) orelse return ValidationError.InvalidValue;
+            if (resource.kind != .storage_buffer)
+                return ValidationError.WrongResourceKind;
+
+            if (!isUnsignedInteger(module, try operandType(module, function_id, op.byte_offset)))
+                return ValidationError.WrongOperandType;
+
+            if (!isBufferAccessibleType(module, try operandType(module, function_id, op.value)))
+                return ValidationError.WrongOperandType;
+        },
         .call => |op| {
             const callee = module.functions.get(op.function) orelse return ValidationError.InvalidFunction;
 
@@ -459,6 +486,26 @@ fn isBoolean(module: *const module_ir.Module, type_id: ids.TypeId) bool {
     return ty.* == .boolean;
 }
 
+fn isUnsignedInteger(module: *const module_ir.Module, type_id: ids.TypeId) bool {
+    const ty = module.types.get(type_id) orelse return false;
+    return switch (ty.*) {
+        .integer => |integer| integer.signedness == .unsigned,
+        else => false,
+    };
+}
+
+fn isBufferAccessibleType(module: *const module_ir.Module, type_id: ids.TypeId) bool {
+    const ty = module.types.get(type_id) orelse return false;
+    return switch (ty.*) {
+        .integer, .floating => true,
+        .vector => |vector| {
+            const element_type = module.types.get(vector.element_type) orelse return false;
+            return element_type.* == .integer or element_type.* == .floating;
+        },
+        else => false,
+    };
+}
+
 fn targetsBlock(terminator: module_ir.Terminator, target: ids.BlockId) bool {
     return switch (terminator) {
         .branch => |edge| edge.target == target,
@@ -809,6 +856,109 @@ test "Validator: check interface direction and value types" {
     );
 }
 
+test "Validator: check buffer resources, offsets, and value types" {
+    try expectValidationError(Error.WrongResourceKind,
+        \\shader compute @main
+        \\{
+        \\    @uniforms: u32 = uniform_buffer[set(0), binding(0)]
+        \\    %offset: constant u32 = 0
+        \\    fn @main() -> void
+        \\    {
+        \\        .entry():
+        \\            %value: u32 = load_buffer @uniforms, %offset
+        \\            return
+        \\    }
+        \\}
+    );
+
+    try expectValidationError(Error.WrongResourceKind,
+        \\shader compute @main
+        \\{
+        \\    @uniforms: u32 = uniform_buffer[set(0), binding(0)]
+        \\    %offset: constant u32 = 0
+        \\    %value: constant u32 = 1
+        \\    fn @main() -> void
+        \\    {
+        \\        .entry():
+        \\            store_buffer @uniforms, %offset, %value
+        \\            return
+        \\    }
+        \\}
+    );
+
+    try expectValidationError(Error.WrongOperandType,
+        \\shader compute @main
+        \\{
+        \\    @storage: u32 = storage_buffer[set(0), binding(0)]
+        \\    %offset: constant i32 = 0
+        \\    fn @main() -> void
+        \\    {
+        \\        .entry():
+        \\            %value: u32 = load_buffer @storage, %offset
+        \\            return
+        \\    }
+        \\}
+    );
+
+    try expectValidationError(Error.WrongResultType,
+        \\shader compute @main
+        \\{
+        \\    @storage: struct[u32, f32] = storage_buffer[set(0), binding(0)]
+        \\    %offset: constant u32 = 0
+        \\    fn @main() -> void
+        \\    {
+        \\        .entry():
+        \\            %value: struct[u32, f32] = load_buffer @storage, %offset
+        \\            return
+        \\    }
+        \\}
+    );
+
+    try expectValidationError(Error.WrongOperandType,
+        \\shader compute @main
+        \\{
+        \\    @storage: struct[u32, f32] = storage_buffer[set(0), binding(0)]
+        \\    %offset: constant u32 = 0
+        \\    %value: constant ptr[private, u32] = null
+        \\    fn @main() -> void
+        \\    {
+        \\        .entry():
+        \\            store_buffer @storage, %offset, %value
+        \\            return
+        \\    }
+        \\}
+    );
+
+    try expectValidationError(Error.WrongResultPresence,
+        \\shader compute @main
+        \\{
+        \\    @storage: struct[u32, f32] = storage_buffer[set(0), binding(0)]
+        \\    %offset: constant u32 = 0
+        \\    fn @main() -> void
+        \\    {
+        \\        .entry():
+        \\            load_buffer @storage, %offset
+        \\            return
+        \\    }
+        \\}
+    );
+
+    try expectValidationError(Error.WrongResultPresence,
+        \\shader compute @main
+        \\{
+        \\    @storage: u32 = storage_buffer[set(0), binding(0)]
+        \\    %offset: constant u32 = 0
+        \\    %value: constant u32 = 1
+        \\    fn @main() -> void
+        \\    {
+        \\        .entry():
+        \\            %result: u32 = store_buffer @storage, %offset, %value
+        \\            return
+        \\    }
+        \\}
+    );
+}
+
 test "Validator: check function calls" {
     try expectValidationError(Error.WrongOperandType,
         \\shader compute @main

src/compiler/spirv/spirv.zig

@@ -5,13 +5,18 @@ pub const header_word_count: usize = 5;
 pub const Opcode = enum(u32) {
     nop = 0,
     undef = 1,
+    source_continued = 2,
+    source = 3,
+    source_extension = 4,
     name = 5,
     member_name = 6,
     string = 7,
     line = 8,
+
     extension = 10,
     ext_inst_import = 11,
     ext_inst = 12,
+
     memory_model = 14,
     entry_point = 15,
     execution_mode = 16,
@@ -32,12 +37,20 @@ pub const Opcode = enum(u32) {
     type_opaque = 31,
     type_pointer = 32,
     type_function = 33,
+    type_event = 34,
+    type_device_event = 35,
+    type_reserve_id = 36,
+    type_queue = 37,
+    type_pipe = 38,
+    type_forward_pointer = 39,
 
     constant_true = 41,
     constant_false = 42,
     constant = 43,
     constant_composite = 44,
+    constant_sampler = 45,
     constant_null = 46,
+
     spec_constant_true = 48,
     spec_constant_false = 49,
     spec_constant = 50,
@@ -48,18 +61,56 @@ pub const Opcode = enum(u32) {
     function_parameter = 55,
     function_end = 56,
     function_call = 57,
+
     variable = 59,
+    image_texel_pointer = 60,
     load = 61,
     store = 62,
+    copy_memory = 63,
+    copy_memory_sized = 64,
     access_chain = 65,
-
+    in_bounds_access_chain = 66,
+    ptr_access_chain = 67,
+    array_length = 68,
+    generic_ptr_mem_semantics = 69,
+    in_bounds_ptr_access_chain = 70,
     decorate = 71,
     member_decorate = 72,
+    decoration_group = 73,
+    group_decorate = 74,
+    group_member_decorate = 75,
+
+    vector_extract_dynamic = 77,
+    vector_insert_dynamic = 78,
     vector_shuffle = 79,
     composite_construct = 80,
     composite_extract = 81,
     composite_insert = 82,
     copy_object = 83,
+    transpose = 84,
+
+    sampled_image = 86,
+    image_sample_implicit_lod = 87,
+    image_sample_explicit_lod = 88,
+    image_sample_dref_implicit_lod = 89,
+    image_sample_dref_explicit_lod = 90,
+    image_sample_proj_implicit_lod = 91,
+    image_sample_proj_explicit_lod = 92,
+    image_sample_proj_dref_implicit_lod = 93,
+    image_sample_proj_dref_explicit_lod = 94,
+    image_fetch = 95,
+    image_gather = 96,
+    image_dref_gather = 97,
+    image_read = 98,
+    image_write = 99,
+    image = 100,
+    image_query_format = 101,
+    image_query_order = 102,
+    image_query_size_lod = 103,
+    image_query_size = 104,
+    image_query_lod = 105,
+    image_query_levels = 106,
+    image_query_samples = 107,
 
     convert_f_to_u = 109,
     convert_f_to_s = 110,
@@ -68,6 +119,14 @@ pub const Opcode = enum(u32) {
     u_convert = 113,
     s_convert = 114,
     f_convert = 115,
+    quantize_to_f16 = 116,
+    convert_ptr_to_u = 117,
+    sat_convert_s_to_u = 118,
+    sat_convert_u_to_s = 119,
+    convert_u_to_ptr = 120,
+    ptr_cast_to_generic = 121,
+    generic_cast_to_ptr = 122,
+    generic_cast_to_ptr_explicit = 123,
     bitcast = 124,
 
     s_negate = 126,
@@ -86,12 +145,28 @@ pub const Opcode = enum(u32) {
     s_mod = 139,
     f_rem = 140,
     f_mod = 141,
-    shift_right_logical = 194,
+    vector_times_scalar = 142,
-    shift_right_arithmetic = 195,
+    matrix_times_scalar = 143,
-    shift_left_logical = 196,
+    vector_times_matrix = 144,
-    bitwise_or = 197,
+    matrix_times_vector = 145,
-    bitwise_xor = 198,
+    matrix_times_matrix = 146,
-    bitwise_and = 199,
+    outer_product = 147,
+    dot = 148,
+    i_add_carry = 149,
+    i_sub_borrow = 150,
+    u_mul_extended = 151,
+    s_mul_extended = 152,
+
+    any = 154,
+    all = 155,
+    is_nan = 156,
+    is_inf = 157,
+    is_finite = 158,
+    is_normal = 159,
+    sign_bit_set = 160,
+    less_or_greater = 161,
+    ordered = 162,
+    unordered = 163,
     logical_equal = 164,
     logical_not_equal = 165,
     logical_or = 166,
@@ -100,14 +175,74 @@ pub const Opcode = enum(u32) {
     select = 169,
     i_equal = 170,
     i_not_equal = 171,
+    u_greater_than = 172,
+    s_greater_than = 173,
+    u_greater_than_equal = 174,
+    s_greater_than_equal = 175,
     u_less_than = 176,
     s_less_than = 177,
+    u_less_than_equal = 178,
+    s_less_than_equal = 179,
     f_ord_equal = 180,
     f_unord_equal = 181,
     f_ord_not_equal = 182,
     f_unord_not_equal = 183,
     f_ord_less_than = 184,
     f_unord_less_than = 185,
+    f_ord_greater_than = 186,
+    f_unord_greater_than = 187,
+    f_ord_less_than_equal = 188,
+    f_unord_less_than_equal = 189,
+    f_ord_greater_than_equal = 190,
+    f_unord_greater_than_equal = 191,
+
+    shift_right_logical = 194,
+    shift_right_arithmetic = 195,
+    shift_left_logical = 196,
+    bitwise_or = 197,
+    bitwise_xor = 198,
+    bitwise_and = 199,
+    not = 200,
+    bit_field_insert = 201,
+    bit_field_s_extract = 202,
+    bit_field_u_extract = 203,
+    bit_reverse = 204,
+    bit_count = 205,
+
+    d_pdx = 207,
+    d_pdy = 208,
+    fwidth = 209,
+    d_pdx_fine = 210,
+    d_pdy_fine = 211,
+    fwidth_fine = 212,
+    d_pdx_coarse = 213,
+    d_pdy_coarse = 214,
+    fwidth_coarse = 215,
+
+    emit_vertex = 218,
+    end_primitive = 219,
+    emit_stream_vertex = 220,
+    end_stream_primitive = 221,
+
+    control_barrier = 224,
+    memory_barrier = 225,
+
+    atomic_load = 227,
+    atomic_store = 228,
+    atomic_exchange = 229,
+    atomic_compare_exchange = 230,
+    atomic_compare_exchange_weak = 231,
+    atomic_i_increment = 232,
+    atomic_i_decrement = 233,
+    atomic_i_add = 234,
+    atomic_i_sub = 235,
+    atomic_s_min = 236,
+    atomic_u_min = 237,
+    atomic_s_max = 238,
+    atomic_u_max = 239,
+    atomic_and = 240,
+    atomic_or = 241,
+    atomic_xor = 242,
 
     phi = 245,
     loop_merge = 246,
@@ -120,7 +255,728 @@ pub const Opcode = enum(u32) {
     return_ = 253,
     return_value = 254,
     @"unreachable" = 255,
+    lifetime_start = 256,
+    lifetime_stop = 257,
+
+    group_async_copy = 259,
+    group_wait_events = 260,
+    group_all = 261,
+    group_any = 262,
+    group_broadcast = 263,
+    group_i_add = 264,
+    group_f_add = 265,
+    group_f_min = 266,
+    group_u_min = 267,
+    group_s_min = 268,
+    group_f_max = 269,
+    group_u_max = 270,
+    group_s_max = 271,
+
+    read_pipe = 274,
+    write_pipe = 275,
+    reserved_read_pipe = 276,
+    reserved_write_pipe = 277,
+    reserve_read_pipe_packets = 278,
+    reserve_write_pipe_packets = 279,
+    commit_read_pipe = 280,
+    commit_write_pipe = 281,
+    is_valid_reserve_id = 282,
+    get_num_pipe_packets = 283,
+    get_max_pipe_packets = 284,
+    group_reserve_read_pipe_packets = 285,
+    group_reserve_write_pipe_packets = 286,
+    group_commit_read_pipe = 287,
+    group_commit_write_pipe = 288,
+
+    enqueue_marker = 291,
+    enqueue_kernel = 292,
+    get_kernel_n_drange_sub_group_count = 293,
+    get_kernel_n_drange_max_sub_group_size = 294,
+    get_kernel_work_group_size = 295,
+    get_kernel_preferred_work_group_size_multiple = 296,
+    retain_event = 297,
+    release_event = 298,
+    create_user_event = 299,
+    is_valid_event = 300,
+    set_user_event_status = 301,
+    capture_event_profiling_info = 302,
+    get_default_queue = 303,
+    build_nd_range = 304,
+    image_sparse_sample_implicit_lod = 305,
+    image_sparse_sample_explicit_lod = 306,
+    image_sparse_sample_dref_implicit_lod = 307,
+    image_sparse_sample_dref_explicit_lod = 308,
+    image_sparse_sample_proj_implicit_lod = 309,
+    image_sparse_sample_proj_explicit_lod = 310,
+    image_sparse_sample_proj_dref_implicit_lod = 311,
+    image_sparse_sample_proj_dref_explicit_lod = 312,
+    image_sparse_fetch = 313,
+    image_sparse_gather = 314,
+    image_sparse_dref_gather = 315,
+    image_sparse_texels_resident = 316,
     no_line = 317,
+    atomic_flag_test_and_set = 318,
+    atomic_flag_clear = 319,
+    image_sparse_read = 320,
+    size_of = 321,
+    type_pipe_storage = 322,
+    constant_pipe_storage = 323,
+    create_pipe_from_pipe_storage = 324,
+    get_kernel_local_size_for_subgroup_count = 325,
+    get_kernel_max_num_subgroups = 326,
+    type_named_barrier = 327,
+    named_barrier_initialize = 328,
+    memory_named_barrier = 329,
+    module_processed = 330,
+    execution_mode_id = 331,
+    decorate_id = 332,
+    group_non_uniform_elect = 333,
+    group_non_uniform_all = 334,
+    group_non_uniform_any = 335,
+    group_non_uniform_all_equal = 336,
+    group_non_uniform_broadcast = 337,
+    group_non_uniform_broadcast_first = 338,
+    group_non_uniform_ballot = 339,
+    group_non_uniform_inverse_ballot = 340,
+    group_non_uniform_ballot_bit_extract = 341,
+    group_non_uniform_ballot_bit_count = 342,
+    group_non_uniform_ballot_find_lsb = 343,
+    group_non_uniform_ballot_find_msb = 344,
+    group_non_uniform_shuffle = 345,
+    group_non_uniform_shuffle_xor = 346,
+    group_non_uniform_shuffle_up = 347,
+    group_non_uniform_shuffle_down = 348,
+    group_non_uniform_i_add = 349,
+    group_non_uniform_f_add = 350,
+    group_non_uniform_i_mul = 351,
+    group_non_uniform_f_mul = 352,
+    group_non_uniform_s_min = 353,
+    group_non_uniform_u_min = 354,
+    group_non_uniform_f_min = 355,
+    group_non_uniform_s_max = 356,
+    group_non_uniform_u_max = 357,
+    group_non_uniform_f_max = 358,
+    group_non_uniform_bitwise_and = 359,
+    group_non_uniform_bitwise_or = 360,
+    group_non_uniform_bitwise_xor = 361,
+    group_non_uniform_logical_and = 362,
+    group_non_uniform_logical_or = 363,
+    group_non_uniform_logical_xor = 364,
+    group_non_uniform_quad_broadcast = 365,
+    group_non_uniform_quad_swap = 366,
+
+    copy_logical = 400,
+    ptr_equal = 401,
+    ptr_not_equal = 402,
+    ptr_diff = 403,
+
+    color_attachment_read_ext = 4160,
+    depth_attachment_read_ext = 4161,
+    stencil_attachment_read_ext = 4162,
+    type_tensor_arm = 4163,
+    tensor_read_arm = 4164,
+    tensor_write_arm = 4165,
+    tensor_query_size_arm = 4166,
+
+    graph_constant_arm = 4181,
+    graph_entry_point_arm = 4182,
+    graph_arm = 4183,
+    graph_input_arm = 4184,
+    graph_set_output_arm = 4185,
+    graph_end_arm = 4186,
+
+    type_graph_arm = 4190,
+
+    bitcast_extract_ext = 4195,
+
+    terminate_invocation = 4416,
+    type_untyped_pointer_khr = 4417,
+    untyped_variable_khr = 4418,
+    untyped_access_chain_khr = 4419,
+    untyped_in_bounds_access_chain_khr = 4420,
+    subgroup_ballot_khr = 4421,
+    subgroup_first_invocation_khr = 4422,
+    untyped_ptr_access_chain_khr = 4423,
+    untyped_in_bounds_ptr_access_chain_khr = 4424,
+    untyped_array_length_khr = 4425,
+    untyped_prefetch_khr = 4426,
+    fma_khr = 4427,
+    subgroup_all_khr = 4428,
+    subgroup_any_khr = 4429,
+    subgroup_all_equal_khr = 4430,
+    group_non_uniform_rotate_khr = 4431,
+    subgroup_read_invocation_khr = 4432,
+    ext_inst_with_forward_refs_khr = 4433,
+    untyped_group_async_copy_khr = 4434,
+
+    trace_ray_khr = 4445,
+    execute_callable_khr = 4446,
+    convert_u_to_acceleration_structure_khr = 4447,
+    ignore_intersection_khr = 4448,
+    terminate_ray_khr = 4449,
+    s_dot = 4450,
+    u_dot = 4451,
+    su_dot = 4452,
+    s_dot_acc_sat = 4453,
+    u_dot_acc_sat = 4454,
+    su_dot_acc_sat = 4455,
+    type_cooperative_matrix_khr = 4456,
+    cooperative_matrix_load_khr = 4457,
+    cooperative_matrix_store_khr = 4458,
+    cooperative_matrix_mul_add_khr = 4459,
+    cooperative_matrix_length_khr = 4460,
+    constant_composite_replicate_ext = 4461,
+    spec_constant_composite_replicate_ext = 4462,
+    composite_construct_replicate_ext = 4463,
+
+    type_ray_query_khr = 4472,
+    ray_query_initialize_khr = 4473,
+    ray_query_terminate_khr = 4474,
+    ray_query_generate_intersection_khr = 4475,
+    ray_query_confirm_intersection_khr = 4476,
+    ray_query_proceed_khr = 4477,
+
+    ray_query_get_intersection_type_khr = 4479,
+    image_sample_weighted_qcom = 4480,
+    image_box_filter_qcom = 4481,
+    image_block_match_ssd_qcom = 4482,
+    image_block_match_sad_qcom = 4483,
+
+    bit_cast_array_qcom = 4497,
+
+    image_block_match_window_ssd_qcom = 4500,
+    image_block_match_window_sad_qcom = 4501,
+    image_block_match_gather_ssd_qcom = 4502,
+    image_block_match_gather_sad_qcom = 4503,
+
+    composite_construct_coop_mat_qcom = 4540,
+    composite_extract_coop_mat_qcom = 4541,
+    extract_sub_array_qcom = 4542,
+
+    image_gather_qcom = 4545,
+
+    group_i_add_non_uniform_amd = 5000,
+    group_f_add_non_uniform_amd = 5001,
+    group_f_min_non_uniform_amd = 5002,
+    group_u_min_non_uniform_amd = 5003,
+    group_s_min_non_uniform_amd = 5004,
+    group_f_max_non_uniform_amd = 5005,
+    group_u_max_non_uniform_amd = 5006,
+    group_s_max_non_uniform_amd = 5007,
+
+    fragment_mask_fetch_amd = 5011,
+    fragment_fetch_amd = 5012,
+
+    read_clock_khr = 5056,
+
+    allocate_node_payloads_amdx = 5074,
+    enqueue_node_payloads_amdx = 5075,
+    type_node_payload_array_amdx = 5076,
+
+    finish_writing_node_payload_amdx = 5078,
+
+    node_payload_array_length_amdx = 5090,
+
+    is_node_payload_valid_amdx = 5101,
+
+    constant_string_amdx = 5103,
+    spec_constant_string_amdx = 5104,
+
+    group_non_uniform_quad_all_khr = 5110,
+    group_non_uniform_quad_any_khr = 5111,
+
+    type_buffer_ext = 5115,
+
+    buffer_pointer_ext = 5119,
+
+    abort_khr = 5121,
+
+    untyped_image_texel_pointer_ext = 5126,
+    member_decorate_id_ext = 5127,
+
+    constant_size_of_ext = 5129,
+
+    constant_data_khr = 5147,
+    spec_constant_data_khr = 5148,
+
+    poison_khr = 5158,
+    freeze_khr = 5159,
+
+    hit_object_record_hit_motion_nv = 5249,
+    hit_object_record_hit_with_index_motion_nv = 5250,
+    hit_object_record_miss_motion_nv = 5251,
+    hit_object_get_world_to_object_nv = 5252,
+    hit_object_get_object_to_world_nv = 5253,
+    hit_object_get_object_ray_direction_nv = 5254,
+    hit_object_get_object_ray_origin_nv = 5255,
+    hit_object_trace_ray_motion_nv = 5256,
+    hit_object_get_shader_record_buffer_handle_nv = 5257,
+    hit_object_get_shader_binding_table_record_index_nv = 5258,
+    hit_object_record_empty_nv = 5259,
+    hit_object_trace_ray_nv = 5260,
+    hit_object_record_hit_nv = 5261,
+    hit_object_record_hit_with_index_nv = 5262,
+    hit_object_record_miss_nv = 5263,
+    hit_object_execute_shader_nv = 5264,
+    hit_object_get_current_time_nv = 5265,
+    hit_object_get_attributes_nv = 5266,
+    hit_object_get_hit_kind_nv = 5267,
+    hit_object_get_primitive_index_nv = 5268,
+    hit_object_get_geometry_index_nv = 5269,
+    hit_object_get_instance_id_nv = 5270,
+    hit_object_get_instance_custom_index_nv = 5271,
+    hit_object_get_world_ray_direction_nv = 5272,
+    hit_object_get_world_ray_origin_nv = 5273,
+    hit_object_get_ray_t_max_nv = 5274,
+    hit_object_get_ray_t_min_nv = 5275,
+    hit_object_is_empty_nv = 5276,
+    hit_object_is_hit_nv = 5277,
+    hit_object_is_miss_nv = 5278,
+    reorder_thread_with_hit_object_nv = 5279,
+    reorder_thread_with_hint_nv = 5280,
+    type_hit_object_nv = 5281,
+
+    image_sample_footprint_nv = 5283,
+
+    type_vector_id_ext = 5288,
+    cooperative_vector_matrix_mul_nv = 5289,
+    cooperative_vector_outer_product_accumulate_nv = 5290,
+    cooperative_vector_reduce_sum_accumulate_nv = 5291,
+    cooperative_vector_matrix_mul_add_nv = 5292,
+    cooperative_matrix_convert_use_ext = 5293,
+    emit_mesh_tasks_ext = 5294,
+    set_mesh_outputs_ext = 5295,
+    group_non_uniform_partition_ext = 5296,
+
+    write_packed_primitive_indices4x8_nv = 5299,
+    fetch_micro_triangle_vertex_position_nv = 5300,
+    fetch_micro_triangle_vertex_barycentric_nv = 5301,
+    cooperative_vector_load_nv = 5302,
+    cooperative_vector_store_nv = 5303,
+    hit_object_record_from_query_ext = 5304,
+    hit_object_record_miss_ext = 5305,
+    hit_object_record_miss_motion_ext = 5306,
+    hit_object_get_intersection_triangle_vertex_positions_ext = 5307,
+    hit_object_get_ray_flags_ext = 5308,
+    hit_object_set_shader_binding_table_record_index_ext = 5309,
+    hit_object_reorder_execute_shader_ext = 5310,
+    hit_object_trace_reorder_execute_ext = 5311,
+    hit_object_trace_motion_reorder_execute_ext = 5312,
+    type_hit_object_ext = 5313,
+    reorder_thread_with_hint_ext = 5314,
+    reorder_thread_with_hit_object_ext = 5315,
+    hit_object_trace_ray_ext = 5316,
+    hit_object_trace_ray_motion_ext = 5317,
+    hit_object_record_empty_ext = 5318,
+    hit_object_execute_shader_ext = 5319,
+    hit_object_get_current_time_ext = 5320,
+    hit_object_get_attributes_ext = 5321,
+    hit_object_get_hit_kind_ext = 5322,
+    hit_object_get_primitive_index_ext = 5323,
+    hit_object_get_geometry_index_ext = 5324,
+    hit_object_get_instance_id_ext = 5325,
+    hit_object_get_instance_custom_index_ext = 5326,
+    hit_object_get_object_ray_origin_ext = 5327,
+    hit_object_get_object_ray_direction_ext = 5328,
+    hit_object_get_world_ray_direction_ext = 5329,
+    hit_object_get_world_ray_origin_ext = 5330,
+    hit_object_get_object_to_world_ext = 5331,
+    hit_object_get_world_to_object_ext = 5332,
+    hit_object_get_ray_t_max_ext = 5333,
+    report_intersection_khr = 5334,
+    ignore_intersection_nv = 5335,
+    terminate_ray_nv = 5336,
+    trace_nv = 5337,
+    trace_motion_nv = 5338,
+    trace_ray_motion_nv = 5339,
+    ray_query_get_intersection_triangle_vertex_positions_khr = 5340,
+    type_acceleration_structure_khr = 5341,
+
+    execute_callable_nv = 5344,
+    ray_query_get_intersection_cluster_id_nv = 5345,
+    hit_object_get_cluster_id_nv = 5346,
+    hit_object_get_ray_t_min_ext = 5347,
+    hit_object_get_shader_binding_table_record_index_ext = 5348,
+    hit_object_get_shader_record_buffer_handle_ext = 5349,
+    hit_object_is_empty_ext = 5350,
+    hit_object_is_hit_ext = 5351,
+    hit_object_is_miss_ext = 5352,
+
+    type_cooperative_matrix_nv = 5358,
+    cooperative_matrix_load_nv = 5359,
+    cooperative_matrix_store_nv = 5360,
+    cooperative_matrix_mul_add_nv = 5361,
+    cooperative_matrix_length_nv = 5362,
+    cooperative_matrix_get_coordinate_ext = 5363,
+    begin_invocation_interlock_ext = 5364,
+    end_invocation_interlock_ext = 5365,
+    cooperative_matrix_reduce_ext = 5366,
+    cooperative_matrix_load_tensor_nv = 5367,
+    cooperative_matrix_store_tensor_nv = 5368,
+    cooperative_matrix_per_element_op_ext = 5369,
+    type_tensor_layout_nv = 5370,
+    type_tensor_view_nv = 5371,
+    create_tensor_layout_nv = 5372,
+    tensor_layout_set_dimension_nv = 5373,
+    tensor_layout_set_stride_nv = 5374,
+    tensor_layout_slice_nv = 5375,
+    tensor_layout_set_clamp_value_nv = 5376,
+    create_tensor_view_nv = 5377,
+    tensor_view_set_dimension_nv = 5378,
+    tensor_view_set_stride_nv = 5379,
+    demote_to_helper_invocation = 5380,
+    is_helper_invocation_ext = 5381,
+    tensor_view_set_clip_nv = 5382,
+
+    tensor_layout_set_block_size_nv = 5384,
+
+    cooperative_matrix_transpose_nv = 5390,
+    convert_u_to_image_nv = 5391,
+    convert_u_to_sampler_nv = 5392,
+    convert_image_to_u_nv = 5393,
+    convert_sampler_to_u_nv = 5394,
+    convert_u_to_sampled_image_nv = 5395,
+    convert_sampled_image_to_u_nv = 5396,
+    sampler_image_addressing_mode_nv = 5397,
+    raw_access_chain_nv = 5398,
+
+    ray_query_get_intersection_sphere_position_nv = 5427,
+    ray_query_get_intersection_sphere_radius_nv = 5428,
+    ray_query_get_intersection_lss_positions_nv = 5429,
+    ray_query_get_intersection_lss_radii_nv = 5430,
+    ray_query_get_intersection_lss_hit_value_nv = 5431,
+    hit_object_get_sphere_position_nv = 5432,
+    hit_object_get_sphere_radius_nv = 5433,
+    hit_object_get_lss_positions_nv = 5434,
+    hit_object_get_lss_radii_nv = 5435,
+    hit_object_is_sphere_hit_nv = 5436,
+    hit_object_is_lss_hit_nv = 5437,
+    ray_query_is_sphere_hit_nv = 5438,
+    ray_query_is_lss_hit_nv = 5439,
+
+    subgroup_shuffle_intel = 5571,
+    subgroup_shuffle_down_intel = 5572,
+    subgroup_shuffle_up_intel = 5573,
+    subgroup_shuffle_xor_intel = 5574,
+    subgroup_block_read_intel = 5575,
+    subgroup_block_write_intel = 5576,
+    subgroup_image_block_read_intel = 5577,
+    subgroup_image_block_write_intel = 5578,
+
+    subgroup_image_media_block_read_intel = 5580,
+    subgroup_image_media_block_write_intel = 5581,
+
+    u_count_leading_zeros_intel = 5585,
+    u_count_trailing_zeros_intel = 5586,
+    abs_i_sub_intel = 5587,
+    abs_u_sub_intel = 5588,
+    i_add_sat_intel = 5589,
+    u_add_sat_intel = 5590,
+    i_average_intel = 5591,
+    u_average_intel = 5592,
+    i_average_rounded_intel = 5593,
+    u_average_rounded_intel = 5594,
+    i_sub_sat_intel = 5595,
+    u_sub_sat_intel = 5596,
+    i_mul32x16_intel = 5597,
+    u_mul32x16_intel = 5598,
+
+    constant_function_pointer_intel = 5600,
+    function_pointer_call_intel = 5601,
+
+    asm_target_intel = 5609,
+    asm_intel = 5610,
+    asm_call_intel = 5611,
+
+    atomic_f_min_ext = 5614,
+    atomic_f_max_ext = 5615,
+
+    assume_true_khr = 5630,
+    expect_khr = 5631,
+    decorate_string = 5632,
+    member_decorate_string = 5633,
+
+    vme_image_intel = 5699,
+    type_vme_image_intel = 5700,
+    type_avc_ime_payload_intel = 5701,
+    type_avc_ref_payload_intel = 5702,
+    type_avc_sic_payload_intel = 5703,
+    type_avc_mce_payload_intel = 5704,
+    type_avc_mce_result_intel = 5705,
+    type_avc_ime_result_intel = 5706,
+    type_avc_ime_result_single_reference_streamout_intel = 5707,
+    type_avc_ime_result_dual_reference_streamout_intel = 5708,
+    type_avc_ime_single_reference_streamin_intel = 5709,
+    type_avc_ime_dual_reference_streamin_intel = 5710,
+    type_avc_ref_result_intel = 5711,
+    type_avc_sic_result_intel = 5712,
+    subgroup_avc_mce_get_default_inter_base_multi_reference_penalty_intel = 5713,
+    subgroup_avc_mce_set_inter_base_multi_reference_penalty_intel = 5714,
+    subgroup_avc_mce_get_default_inter_shape_penalty_intel = 5715,
+    subgroup_avc_mce_set_inter_shape_penalty_intel = 5716,
+    subgroup_avc_mce_get_default_inter_direction_penalty_intel = 5717,
+    subgroup_avc_mce_set_inter_direction_penalty_intel = 5718,
+    subgroup_avc_mce_get_default_intra_luma_shape_penalty_intel = 5719,
+    subgroup_avc_mce_get_default_inter_motion_vector_cost_table_intel = 5720,
+    subgroup_avc_mce_get_default_high_penalty_cost_table_intel = 5721,
+    subgroup_avc_mce_get_default_medium_penalty_cost_table_intel = 5722,
+    subgroup_avc_mce_get_default_low_penalty_cost_table_intel = 5723,
+    subgroup_avc_mce_set_motion_vector_cost_function_intel = 5724,
+    subgroup_avc_mce_get_default_intra_luma_mode_penalty_intel = 5725,
+    subgroup_avc_mce_get_default_non_dc_luma_intra_penalty_intel = 5726,
+    subgroup_avc_mce_get_default_intra_chroma_mode_base_penalty_intel = 5727,
+    subgroup_avc_mce_set_ac_only_haar_intel = 5728,
+    subgroup_avc_mce_set_source_interlaced_field_polarity_intel = 5729,
+    subgroup_avc_mce_set_single_reference_interlaced_field_polarity_intel = 5730,
+    subgroup_avc_mce_set_dual_reference_interlaced_field_polarities_intel = 5731,
+    subgroup_avc_mce_convert_to_ime_payload_intel = 5732,
+    subgroup_avc_mce_convert_to_ime_result_intel = 5733,
+    subgroup_avc_mce_convert_to_ref_payload_intel = 5734,
+    subgroup_avc_mce_convert_to_ref_result_intel = 5735,
+    subgroup_avc_mce_convert_to_sic_payload_intel = 5736,
+    subgroup_avc_mce_convert_to_sic_result_intel = 5737,
+    subgroup_avc_mce_get_motion_vectors_intel = 5738,
+    subgroup_avc_mce_get_inter_distortions_intel = 5739,
+    subgroup_avc_mce_get_best_inter_distortions_intel = 5740,
+    subgroup_avc_mce_get_inter_major_shape_intel = 5741,
+    subgroup_avc_mce_get_inter_minor_shape_intel = 5742,
+    subgroup_avc_mce_get_inter_directions_intel = 5743,
+    subgroup_avc_mce_get_inter_motion_vector_count_intel = 5744,
+    subgroup_avc_mce_get_inter_reference_ids_intel = 5745,
+    subgroup_avc_mce_get_inter_reference_interlaced_field_polarities_intel = 5746,
+    subgroup_avc_ime_initialize_intel = 5747,
+    subgroup_avc_ime_set_single_reference_intel = 5748,
+    subgroup_avc_ime_set_dual_reference_intel = 5749,
+    subgroup_avc_ime_ref_window_size_intel = 5750,
+    subgroup_avc_ime_adjust_ref_offset_intel = 5751,
+    subgroup_avc_ime_convert_to_mce_payload_intel = 5752,
+    subgroup_avc_ime_set_max_motion_vector_count_intel = 5753,
+    subgroup_avc_ime_set_unidirectional_mix_disable_intel = 5754,
+    subgroup_avc_ime_set_early_search_termination_threshold_intel = 5755,
+    subgroup_avc_ime_set_weighted_sad_intel = 5756,
+    subgroup_avc_ime_evaluate_with_single_reference_intel = 5757,
+    subgroup_avc_ime_evaluate_with_dual_reference_intel = 5758,
+    subgroup_avc_ime_evaluate_with_single_reference_streamin_intel = 5759,
+    subgroup_avc_ime_evaluate_with_dual_reference_streamin_intel = 5760,
+    subgroup_avc_ime_evaluate_with_single_reference_streamout_intel = 5761,
+    subgroup_avc_ime_evaluate_with_dual_reference_streamout_intel = 5762,
+    subgroup_avc_ime_evaluate_with_single_reference_streaminout_intel = 5763,
+    subgroup_avc_ime_evaluate_with_dual_reference_streaminout_intel = 5764,
+    subgroup_avc_ime_convert_to_mce_result_intel = 5765,
+    subgroup_avc_ime_get_single_reference_streamin_intel = 5766,
+    subgroup_avc_ime_get_dual_reference_streamin_intel = 5767,
+    subgroup_avc_ime_strip_single_reference_streamout_intel = 5768,
+    subgroup_avc_ime_strip_dual_reference_streamout_intel = 5769,
+    subgroup_avc_ime_get_streamout_single_reference_major_shape_motion_vectors_intel = 5770,
+    subgroup_avc_ime_get_streamout_single_reference_major_shape_distortions_intel = 5771,
+    subgroup_avc_ime_get_streamout_single_reference_major_shape_reference_ids_intel = 5772,
+    subgroup_avc_ime_get_streamout_dual_reference_major_shape_motion_vectors_intel = 5773,
+    subgroup_avc_ime_get_streamout_dual_reference_major_shape_distortions_intel = 5774,
+    subgroup_avc_ime_get_streamout_dual_reference_major_shape_reference_ids_intel = 5775,
+    subgroup_avc_ime_get_border_reached_intel = 5776,
+    subgroup_avc_ime_get_truncated_search_indication_intel = 5777,
+    subgroup_avc_ime_get_unidirectional_early_search_termination_intel = 5778,
+    subgroup_avc_ime_get_weighting_pattern_minimum_motion_vector_intel = 5779,
+    subgroup_avc_ime_get_weighting_pattern_minimum_distortion_intel = 5780,
+    subgroup_avc_fme_initialize_intel = 5781,
+    subgroup_avc_bme_initialize_intel = 5782,
+    subgroup_avc_ref_convert_to_mce_payload_intel = 5783,
+    subgroup_avc_ref_set_bidirectional_mix_disable_intel = 5784,
+    subgroup_avc_ref_set_bilinear_filter_enable_intel = 5785,
+    subgroup_avc_ref_evaluate_with_single_reference_intel = 5786,
+    subgroup_avc_ref_evaluate_with_dual_reference_intel = 5787,
+    subgroup_avc_ref_evaluate_with_multi_reference_intel = 5788,
+    subgroup_avc_ref_evaluate_with_multi_reference_interlaced_intel = 5789,
+    subgroup_avc_ref_convert_to_mce_result_intel = 5790,
+    subgroup_avc_sic_initialize_intel = 5791,
+    subgroup_avc_sic_configure_skc_intel = 5792,
+    subgroup_avc_sic_configure_ipe_luma_intel = 5793,
+    subgroup_avc_sic_configure_ipe_luma_chroma_intel = 5794,
+    subgroup_avc_sic_get_motion_vector_mask_intel = 5795,
+    subgroup_avc_sic_convert_to_mce_payload_intel = 5796,
+    subgroup_avc_sic_set_intra_luma_shape_penalty_intel = 5797,
+    subgroup_avc_sic_set_intra_luma_mode_cost_function_intel = 5798,
+    subgroup_avc_sic_set_intra_chroma_mode_cost_function_intel = 5799,
+    subgroup_avc_sic_set_bilinear_filter_enable_intel = 5800,
+    subgroup_avc_sic_set_skc_forward_transform_enable_intel = 5801,
+    subgroup_avc_sic_set_block_based_raw_skip_sad_intel = 5802,
+    subgroup_avc_sic_evaluate_ipe_intel = 5803,
+    subgroup_avc_sic_evaluate_with_single_reference_intel = 5804,
+    subgroup_avc_sic_evaluate_with_dual_reference_intel = 5805,
+    subgroup_avc_sic_evaluate_with_multi_reference_intel = 5806,
+    subgroup_avc_sic_evaluate_with_multi_reference_interlaced_intel = 5807,
+    subgroup_avc_sic_convert_to_mce_result_intel = 5808,
+    subgroup_avc_sic_get_ipe_luma_shape_intel = 5809,
+    subgroup_avc_sic_get_best_ipe_luma_distortion_intel = 5810,
+    subgroup_avc_sic_get_best_ipe_chroma_distortion_intel = 5811,
+    subgroup_avc_sic_get_packed_ipe_luma_modes_intel = 5812,
+    subgroup_avc_sic_get_ipe_chroma_mode_intel = 5813,
+    subgroup_avc_sic_get_packed_skc_luma_count_threshold_intel = 5814,
+    subgroup_avc_sic_get_packed_skc_luma_sum_threshold_intel = 5815,
+    subgroup_avc_sic_get_inter_raw_sads_intel = 5816,
+
+    variable_length_array_intel = 5818,
+    save_memory_intel = 5819,
+    restore_memory_intel = 5820,
+
+    arbitrary_float_sin_cos_pi_altera = 5840,
+    arbitrary_float_cast_altera = 5841,
+    arbitrary_float_cast_from_int_altera = 5842,
+    arbitrary_float_cast_to_int_altera = 5843,
+
+    arbitrary_float_add_altera = 5846,
+    arbitrary_float_sub_altera = 5847,
+    arbitrary_float_mul_altera = 5848,
+    arbitrary_float_div_altera = 5849,
+    arbitrary_float_gt_altera = 5850,
+    arbitrary_float_ge_altera = 5851,
+    arbitrary_float_lt_altera = 5852,
+    arbitrary_float_le_altera = 5853,
+    arbitrary_float_eq_altera = 5854,
+    arbitrary_float_recip_altera = 5855,
+    arbitrary_float_r_sqrt_altera = 5856,
+    arbitrary_float_cbrt_altera = 5857,
+    arbitrary_float_hypot_altera = 5858,
+    arbitrary_float_sqrt_altera = 5859,
+    arbitrary_float_log_intel = 5860,
+    arbitrary_float_log2_intel = 5861,
+    arbitrary_float_log10_intel = 5862,
+    arbitrary_float_log1p_intel = 5863,
+    arbitrary_float_exp_intel = 5864,
+    arbitrary_float_exp2_intel = 5865,
+    arbitrary_float_exp10_intel = 5866,
+    arbitrary_float_expm1_intel = 5867,
+    arbitrary_float_sin_intel = 5868,
+    arbitrary_float_cos_intel = 5869,
+    arbitrary_float_sin_cos_intel = 5870,
+    arbitrary_float_sin_pi_intel = 5871,
+    arbitrary_float_cos_pi_intel = 5872,
+    arbitrary_float_a_sin_intel = 5873,
+    arbitrary_float_a_sin_pi_intel = 5874,
+    arbitrary_float_a_cos_intel = 5875,
+    arbitrary_float_a_cos_pi_intel = 5876,
+    arbitrary_float_a_tan_intel = 5877,
+    arbitrary_float_a_tan_pi_intel = 5878,
+    arbitrary_float_a_tan2_intel = 5879,
+    arbitrary_float_pow_intel = 5880,
+    arbitrary_float_pow_r_intel = 5881,
+    arbitrary_float_pow_n_intel = 5882,
+
+    loop_control_intel = 5887,
+
+    alias_domain_decl_intel = 5911,
+    alias_scope_decl_intel = 5912,
+    alias_scope_list_decl_intel = 5913,
+
+    fixed_sqrt_altera = 5923,
+    fixed_recip_altera = 5924,
+    fixed_rsqrt_altera = 5925,
+    fixed_sin_altera = 5926,
+    fixed_cos_altera = 5927,
+    fixed_sin_cos_altera = 5928,
+    fixed_sin_pi_altera = 5929,
+    fixed_cos_pi_altera = 5930,
+    fixed_sin_cos_pi_altera = 5931,
+    fixed_log_altera = 5932,
+    fixed_exp_altera = 5933,
+    ptr_cast_to_cross_workgroup_altera = 5934,
+
+    cross_workgroup_cast_to_ptr_altera = 5938,
+
+    read_pipe_blocking_altera = 5946,
+    write_pipe_blocking_altera = 5947,
+
+    fpga_reg_altera = 5949,
+
+    ray_query_get_ray_t_min_khr = 6016,
+    ray_query_get_ray_flags_khr = 6017,
+    ray_query_get_intersection_t_khr = 6018,
+    ray_query_get_intersection_instance_custom_index_khr = 6019,
+    ray_query_get_intersection_instance_id_khr = 6020,
+    ray_query_get_intersection_instance_shader_binding_table_record_offset_khr = 6021,
+    ray_query_get_intersection_geometry_index_khr = 6022,
+    ray_query_get_intersection_primitive_index_khr = 6023,
+    ray_query_get_intersection_barycentrics_khr = 6024,
+    ray_query_get_intersection_front_face_khr = 6025,
+    ray_query_get_intersection_candidate_aabb_opaque_khr = 6026,
+    ray_query_get_intersection_object_ray_direction_khr = 6027,
+    ray_query_get_intersection_object_ray_origin_khr = 6028,
+    ray_query_get_world_ray_direction_khr = 6029,
+    ray_query_get_world_ray_origin_khr = 6030,
+    ray_query_get_intersection_object_to_world_khr = 6031,
+    ray_query_get_intersection_world_to_object_khr = 6032,
+
+    atomic_f_add_ext = 6035,
+
+    type_buffer_surface_intel = 6086,
+
+    type_struct_continued_intel = 6090,
+    constant_composite_continued_intel = 6091,
+    spec_constant_composite_continued_intel = 6092,
+
+    composite_construct_continued_intel = 6096,
+
+    convert_f_to_bf16_intel = 6116,
+    convert_bf16_to_f_intel = 6117,
+
+    control_barrier_arrive_ext = 6142,
+    control_barrier_wait_ext = 6143,
+
+    arithmetic_fence_ext = 6145,
+
+    task_sequence_create_altera = 6163,
+    task_sequence_async_altera = 6164,
+    task_sequence_get_altera = 6165,
+    task_sequence_release_altera = 6166,
+
+    type_task_sequence_altera = 6199,
+
+    subgroup_block_prefetch_intel = 6221,
+
+    subgroup2_d_block_load_intel = 6231,
+    subgroup2_d_block_load_transform_intel = 6232,
+    subgroup2_d_block_load_transpose_intel = 6233,
+    subgroup2_d_block_prefetch_intel = 6234,
+    subgroup2_d_block_store_intel = 6235,
+
+    subgroup_matrix_multiply_accumulate_intel = 6237,
+
+    bitwise_function_intel = 6242,
+
+    untyped_variable_length_array_intel = 6244,
+
+    conditional_extension_intel = 6248,
+    conditional_entry_point_intel = 6249,
+    conditional_capability_intel = 6250,
+    spec_constant_target_intel = 6251,
+    spec_constant_architecture_intel = 6252,
+    spec_constant_capabilities_intel = 6253,
+    conditional_copy_object_intel = 6254,
+
+    predicated_load_intel = 6258,
+    predicated_store_intel = 6259,
+
+    group_i_mul_khr = 6401,
+    group_f_mul_khr = 6402,
+    group_bitwise_and_khr = 6403,
+    group_bitwise_or_khr = 6404,
+    group_bitwise_xor_khr = 6405,
+    group_logical_and_khr = 6406,
+    group_logical_or_khr = 6407,
+    group_logical_xor_khr = 6408,
+
+    round_f_to_tf32_intel = 6426,
+
+    masked_gather_intel = 6428,
+    masked_scatter_intel = 6429,
+
+    convert_handle_to_image_intel = 6529,
+    convert_handle_to_sampler_intel = 6530,
+    convert_handle_to_sampled_image_intel = 6531,
+
+    f_dot2_mix_acc32_valve = 6916,
+    f_dot2_mix_acc16_valve = 6917,
+    f_dot4_mix_acc32_valve = 6918,
 
     _,
 };
@@ -162,9 +1018,15 @@ pub const ExecutionMode = enum(u32) {
 
 pub const Decoration = enum(u32) {
     spec_id = 1,
+    block = 2,
+    buffer_block = 3,
+    array_stride = 6,
     built_in = 11,
     location = 30,
     component = 31,
     index = 32,
+    binding = 33,
+    descriptor_set = 34,
+    offset = 35,
     _,
 };

src/compiler/spirv/translator.zig

@@ -50,6 +50,28 @@ const Decorations = struct {
     component: u8 = 0,
     index: u8 = 0,
     builtin: ?u32 = null,
+    binding: ?u32 = null,
+    descriptor_set: ?u32 = null,
+    array_stride: ?u32 = null,
+    block: bool = false,
+    buffer_block: bool = false,
+};
+
+const MemberOffset = struct {
+    structure_id: u32,
+    member: u32,
+    offset: u32,
+};
+
+const BufferAddress = struct {
+    resource: ir.id.ResourceId,
+    byte_offset: ?ir.id.ValueId,
+    pointee_type: u32,
+};
+
+const LocalVariable = struct {
+    spv_id: u32,
+    type: ir.id.TypeId,
 };
 
 const PhiInfo = struct {
@@ -75,8 +97,18 @@ const Context = struct {
     values: []?ir.id.ValueId,
     blocks: []?ir.id.BlockId,
     interfaces: []?ir.id.InterfaceVariableId,
+    resources: []?ir.id.ResourceId,
+    buffer_addresses: []?BufferAddress,
+    member_offsets: std.ArrayList(MemberOffset) = .empty,
     phi_infos: std.ArrayList(PhiInfo) = .empty,
 
+    local_indices: []?usize,
+    locals: std.ArrayList(LocalVariable) = .empty,
+    block_local_inputs: []?ir.id.ValueId,
+    block_local_outputs: []?ir.id.ValueId,
+    current_locals: []?ir.id.ValueId,
+    entry_label: ?u32 = null,
+
     fn idIndex(self: *const Context, id: u32) TranslationError!usize {
         if (id == 0 or id >= self.bound)
             return error.InvalidId;
@@ -345,6 +377,22 @@ const Context = struct {
         const index = try self.idIndex(spv_id);
         return self.interfaces[index] orelse error.UnsupportedOpcode;
     }
+
+    fn bufferAddress(self: *const Context, spv_id: u32) TranslationError!?BufferAddress {
+        const index = try self.idIndex(spv_id);
+        return self.buffer_addresses[index];
+    }
+
+    fn localIndex(self: *const Context, spv_id: u32) TranslationError!?usize {
+        const index = try self.idIndex(spv_id);
+        return self.local_indices[index];
+    }
+
+    fn blockLocalIndex(self: *const Context, label: u32, local_index: usize) TranslationError!usize {
+        const label_index = try self.idIndex(label);
+        const base = std.math.mul(usize, label_index, self.locals.items.len) catch return error.InvalidInstruction;
+        return std.math.add(usize, base, local_index) catch return error.InvalidInstruction;
+    }
 };
 
 /// Translates one entry point from a retained SPIR-V source into an independent
@@ -379,12 +427,21 @@ pub fn instantiate(allocator: std.mem.Allocator, source: *const SourceModule, op
         .values = try allocOptional(ir.id.ValueId, scratch, bound),
         .blocks = try allocOptional(ir.id.BlockId, scratch, bound),
         .interfaces = try allocOptional(ir.id.InterfaceVariableId, scratch, bound),
+        .resources = try allocOptional(ir.id.ResourceId, scratch, bound),
+        .buffer_addresses = try allocOptional(BufferAddress, scratch, bound),
+        .local_indices = try allocOptional(usize, scratch, bound),
+        .block_local_inputs = try allocOptional(ir.id.ValueId, scratch, 0),
+        .block_local_outputs = try allocOptional(ir.id.ValueId, scratch, 0),
+        .current_locals = try allocOptional(ir.id.ValueId, scratch, 0),
     };
     @memset(context.decorations, .{});
+    defer context.member_offsets.deinit(scratch);
     defer context.phi_infos.deinit(scratch);
+    defer context.locals.deinit(scratch);
 
     try collectDeclarations(&context);
     try translateInterfaces(&context, entry_point.interface_ids);
+    try translateResources(&context);
     try applyExecutionModes(&context, entry_point.function_id);
     try translateFunction(&context, entry_point.function_id, options.entry_point);
     try ir.validator.validate(&module);
@@ -437,6 +494,10 @@ fn collectDeclarations(context: *Context) !void {
         }
         if (instruction.opcode == .decorate) {
             try collectDecoration(context, operands);
+            continue;
+        }
+        if (instruction.opcode == .member_decorate) {
+            try collectMemberDecoration(context, operands);
         }
     }
 }
@@ -470,10 +531,47 @@ fn collectDecoration(context: *Context, operands: []const u32) !void {
             if (operands[2] > std.math.maxInt(u8)) return error.InvalidInstruction;
             context.decorations[index].index = @intCast(operands[2]);
         },
+        .binding => {
+            try expectOperandCount(operands, 3);
+            context.decorations[index].binding = operands[2];
+        },
+        .descriptor_set => {
+            try expectOperandCount(operands, 3);
+            context.decorations[index].descriptor_set = operands[2];
+        },
+        .array_stride => {
+            try expectOperandCount(operands, 3);
+            context.decorations[index].array_stride = operands[2];
+        },
+        .block => {
+            try expectOperandCount(operands, 2);
+            context.decorations[index].block = true;
+        },
+        .buffer_block => {
+            try expectOperandCount(operands, 2);
+            context.decorations[index].buffer_block = true;
+        },
         else => {},
     }
 }
 
+fn collectMemberDecoration(context: *Context, operands: []const u32) !void {
+    if (operands.len < 3)
+        return error.InvalidInstruction;
+
+    const decoration: spirv.Decoration = @enumFromInt(operands[2]);
+    if (decoration != .offset)
+        return;
+
+    try expectOperandCount(operands, 4);
+    _ = try context.idIndex(operands[0]);
+    try context.member_offsets.append(context.scratch, .{
+        .structure_id = operands[0],
+        .member = operands[1],
+        .offset = operands[3],
+    });
+}
+
 fn translateInterfaces(context: *Context, interface_ids: []const u32) !void {
     for (interface_ids) |spv_id| {
         const index = try context.idIndex(spv_id);
@@ -529,6 +627,52 @@ fn translateInterfaces(context: *Context, interface_ids: []const u32) !void {
     }
 }
 
+fn translateResources(context: *Context) !void {
+    for (context.variable_defs, 0..) |optional_variable, spv_index| {
+        const variable = optional_variable orelse continue;
+        if (variable.operands.len < 3 or variable.operands.len > 4)
+            return error.InvalidInstruction;
+
+        const storage_class: spirv.StorageClass = @enumFromInt(variable.operands[2]);
+        if (storage_class != .uniform and storage_class != .storage_buffer)
+            continue;
+
+        const pointer = context.type_defs[try context.idIndex(variable.operands[0])] orelse return error.MissingDefinition;
+        if (pointer.opcode != .type_pointer)
+            return error.InvalidInstruction;
+        try expectOperandCount(pointer.operands, 3);
+        if (pointer.operands[1] != variable.operands[2])
+            return error.InvalidInstruction;
+
+        const pointee_id = pointer.operands[2];
+        const pointee_decoration = context.decorations[try context.idIndex(pointee_id)];
+        const kind: ir.types.ResourceKind = if (storage_class == .storage_buffer or pointee_decoration.buffer_block)
+            .storage_buffer
+        else if (pointee_decoration.block)
+            .uniform_buffer
+        else
+            continue;
+
+        const variable_decoration = context.decorations[spv_index];
+        const resource = try context.builder.addResource(
+            try context.translateType(pointee_id),
+            kind,
+            variable_decoration.descriptor_set orelse return error.InvalidInstruction,
+            variable_decoration.binding orelse return error.InvalidInstruction,
+            context.nameOf(@intCast(spv_index)),
+        );
+        context.resources[spv_index] = resource;
+        context.buffer_addresses[spv_index] = .{
+            .resource = resource,
+            .byte_offset = null,
+            .pointee_type = pointee_id,
+        };
+    }
+
+    if (context.module.resources.entries.items.len != 0)
+        context.module.properties.explicit_resource_offsets = true;
+}
+
 fn findEntryPoint(parser: Parser, requested_name: []const u8, requested_stage: ?ir.module.Stage) !EntryPoint {
     var found: ?EntryPoint = null;
     var iterator = parser.iterator();
@@ -608,11 +752,58 @@ fn translateFunction(context: *Context, spv_function: u32, entry_name: []const u
     );
     context.builder.setEntryPoint(function);
 
+    try collectFunctionLocals(context, spv_function);
     try predeclareFunction(context, spv_function, function, function_type.operands[2..]);
     try translateFunctionInstructions(context, spv_function);
     try translateFunctionControlFlow(context, spv_function);
 }
 
+fn collectFunctionLocals(context: *Context, spv_function: u32) !void {
+    var active = false;
+    var iterator = context.parser.iterator();
+    while (try iterator.next()) |instruction| {
+        if (instruction.opcode == .function) {
+            active = instruction.operands.len >= 2 and instruction.operands[1] == spv_function;
+            continue;
+        }
+        if (!active)
+            continue;
+        if (instruction.opcode == .function_end)
+            break;
+        if (instruction.opcode != .variable)
+            continue;
+
+        try expectOperandCount(instruction.operands, 3);
+        const storage_class: spirv.StorageClass = @enumFromInt(instruction.operands[2]);
+        if (storage_class != .function)
+            return error.UnsupportedOpcode;
+
+        const pointer = context.type_defs[try context.idIndex(instruction.operands[0])] orelse return error.MissingDefinition;
+        if (pointer.opcode != .type_pointer)
+            return error.InvalidInstruction;
+        try expectOperandCount(pointer.operands, 3);
+        if (pointer.operands[1] != instruction.operands[2])
+            return error.InvalidInstruction;
+
+        const result_id = instruction.operands[1];
+        const result_index = try context.idIndex(result_id);
+        if (context.local_indices[result_index] != null)
+            return error.DuplicateId;
+
+        context.local_indices[result_index] = context.locals.items.len;
+        try context.locals.append(context.scratch, .{
+            .spv_id = result_id,
+            .type = try context.translateType(pointer.operands[2]),
+        });
+    }
+
+    const matrix_len = std.math.mul(usize, context.bound, context.locals.items.len) catch return error.InvalidInstruction;
+    context.block_local_inputs = try allocOptional(ir.id.ValueId, context.scratch, matrix_len);
+    context.block_local_outputs = try allocOptional(ir.id.ValueId, context.scratch, matrix_len);
+    context.current_locals = try allocOptional(ir.id.ValueId, context.scratch, context.locals.items.len);
+    context.module.properties.no_local_memory = true;
+}
+
 fn predeclareFunction(context: *Context, spv_function: u32, function: ir.id.FunctionId, parameter_types: []const u32) !void {
     var active = false;
     var parameter_index: usize = 0;
@@ -652,6 +843,18 @@ fn predeclareFunction(context: *Context, spv_function: u32, function: ir.id.Func
                     return error.DuplicateId;
 
                 context.blocks[index] = try context.builder.addBlock(function, context.nameOf(label_id));
+                if (context.entry_label == null) {
+                    context.entry_label = label_id;
+                } else {
+                    for (context.locals.items, 0..) |local, local_index| {
+                        const value = try context.builder.addBlockParameter(
+                            context.blocks[index].?,
+                            local.type,
+                            context.nameOf(local.spv_id),
+                        );
+                        context.block_local_inputs[try context.blockLocalIndex(label_id, local_index)] = value;
+                    }
+                }
                 current_label = label_id;
             },
             .phi => {
@@ -682,6 +885,7 @@ fn predeclareFunction(context: *Context, spv_function: u32, function: ir.id.Func
 
 fn translateFunctionInstructions(context: *Context, spv_function: u32) !void {
     var active = false;
+    var current_label: ?u32 = null;
     var current_block: ?ir.id.BlockId = null;
     var iterator = context.parser.iterator();
 
@@ -695,7 +899,21 @@ fn translateFunctionInstructions(context: *Context, spv_function: u32) !void {
             continue;
 
         switch (instruction.opcode) {
-            .label => current_block = try context.block(instruction.operands[0]),
+            .label => {
+                try expectOperandCount(instruction.operands, 1);
+                if (current_label) |label|
+                    try saveBlockLocals(context, label);
+
+                const label = instruction.operands[0];
+                current_label = label;
+                current_block = try context.block(label);
+                for (context.current_locals, 0..) |*current, local_index| {
+                    current.* = if (label == context.entry_label.?)
+                        null
+                    else
+                        context.block_local_inputs[try context.blockLocalIndex(label, local_index)];
+                }
+            },
 
             .function_parameter,
             .phi,
@@ -709,7 +927,13 @@ fn translateFunctionInstructions(context: *Context, spv_function: u32) !void {
             .@"unreachable",
             => {},
 
-            .function_end => break,
+            .function_end => {
+                if (current_label) |label|
+                    try saveBlockLocals(context, label);
+                break;
+            },
+
+            .variable => {},
 
             .nop,
             .line,
@@ -721,6 +945,11 @@ fn translateFunctionInstructions(context: *Context, spv_function: u32) !void {
     }
 }
 
+fn saveBlockLocals(context: *Context, label: u32) !void {
+    for (context.current_locals, 0..) |value, local_index|
+        context.block_local_outputs[try context.blockLocalIndex(label, local_index)] = value;
+}
+
 fn translateInstruction(context: *Context, block: ir.id.BlockId, instruction: Parser.Instruction) !void {
     const operands = instruction.operands;
     switch (instruction.opcode) {
@@ -741,24 +970,58 @@ fn translateInstruction(context: *Context, block: ir.id.BlockId, instruction: Pa
             if (operands.len < 3)
                 return error.InvalidInstruction;
 
-            const result = (try context.builder.appendInstruction(block, try context.translateType(operands[0]), .{
+            const result_type = try context.translateType(operands[0]);
+            if (try context.localIndex(operands[2])) |local_index| {
+                const value = context.current_locals[local_index] orelse return error.InvalidInstruction;
+                if (context.module.typeOf(value) != result_type)
+                    return error.InvalidInstruction;
+                try context.setValue(operands[1], value);
+            } else if (try context.bufferAddress(operands[2])) |address| {
+                const result = (try context.builder.appendInstruction(block, result_type, .{
+                    .load_buffer = .{
+                        .resource = address.resource,
+                        .byte_offset = try bufferByteOffset(context, address),
+                    },
+                }, context.nameOf(operands[1]))).?;
+                try context.setValue(operands[1], result);
+            } else {
+                const result = (try context.builder.appendInstruction(block, result_type, .{
                     .load_interface = .{ .variable = try context.interfaceVariable(operands[2]) },
                 }, context.nameOf(operands[1]))).?;
-
                 try context.setValue(operands[1], result);
+            }
         },
         .store => {
             if (operands.len < 2)
                 return error.InvalidInstruction;
 
+            const value = try context.resolveValue(operands[1]);
+            if (try context.localIndex(operands[0])) |local_index| {
+                if (context.module.typeOf(value) != context.locals.items[local_index].type)
+                    return error.InvalidInstruction;
+                context.current_locals[local_index] = value;
+            } else if (try context.bufferAddress(operands[0])) |address| {
+                _ = try context.builder.appendInstruction(block, null, .{
+                    .store_buffer = .{
+                        .resource = address.resource,
+                        .byte_offset = try bufferByteOffset(context, address),
+                        .value = value,
+                    },
+                }, null);
+            } else {
                 _ = try context.builder.appendInstruction(block, null, .{
                     .store_interface = .{
                         .variable = try context.interfaceVariable(operands[0]),
-                    .value = try context.resolveValue(operands[1]),
+                        .value = value,
                     },
                 }, null);
+            }
         },
-        .s_negate, .f_negate, .logical_not => {
+        .access_chain => try translateAccessChain(context, block, operands),
+        .s_negate,
+        .f_negate,
+        .logical_not,
+        => {
             try expectOperandCount(operands, 3);
 
             const result = (try context.builder.appendInstruction(block, try context.translateType(operands[0]), .{
@@ -882,8 +1145,133 @@ fn translateInstruction(context: *Context, block: ir.id.BlockId, instruction: Pa
         },
         .function_call => return error.UnsupportedOpcode,
 
-        else => return error.UnsupportedOpcode,
+        else => {
+            if (std.enums.tagName(spirv.Opcode, instruction.opcode)) |opcode| {
+                std.log.scoped(.spirv_translator).err("unsupported opcode {s}", .{opcode});
+            } else {
+                std.log.scoped(.spirv_translator).err("unsupported opcode {d}", .{instruction.opcode});
             }
+            return error.UnsupportedOpcode;
+        },
+    }
+}
+
+fn translateAccessChain(context: *Context, block: ir.id.BlockId, operands: []const u32) !void {
+    if (operands.len < 4)
+        return error.InvalidInstruction;
+
+    const base = (try context.bufferAddress(operands[2])) orelse return error.UnsupportedOpcode;
+    var current_type = base.pointee_type;
+    var byte_offset = base.byte_offset;
+
+    for (operands[3..]) |index_id| {
+        const type_definition = context.type_defs[try context.idIndex(current_type)] orelse return error.MissingDefinition;
+        switch (type_definition.opcode) {
+            .type_struct => {
+                const member = try constantIndex(context, index_id);
+                if (member + 1 >= type_definition.operands.len)
+                    return error.InvalidInstruction;
+
+                const member_offset = try findMemberOffset(context, current_type, member);
+                if (member_offset != 0) {
+                    const offset_value = try context.builder.internConstant(try unsigned32Type(context), .{ .integer_bits = member_offset });
+                    byte_offset = try addByteOffset(context, block, byte_offset, offset_value);
+                }
+                current_type = type_definition.operands[member + 1];
+            },
+            .type_array => {
+                try expectOperandCount(type_definition.operands, 3);
+                const stride = context.decorations[try context.idIndex(current_type)].array_stride orelse return error.InvalidInstruction;
+                const index = try unsignedOffsetValue(context, block, index_id);
+                const stride_value = try context.builder.internConstant(try unsigned32Type(context), .{ .integer_bits = stride });
+                const term = (try context.builder.appendInstruction(block, try unsigned32Type(context), .{
+                    .binary = .{
+                        .opcode = .integer_multiply,
+                        .lhs = index,
+                        .rhs = stride_value,
+                    },
+                }, null)).?;
+                byte_offset = try addByteOffset(context, block, byte_offset, term);
+                current_type = type_definition.operands[1];
+            },
+            else => return error.UnsupportedType,
+        }
+    }
+
+    const result_pointer = context.type_defs[try context.idIndex(operands[0])] orelse return error.MissingDefinition;
+    if (result_pointer.opcode != .type_pointer)
+        return error.InvalidInstruction;
+    try expectOperandCount(result_pointer.operands, 3);
+    if (result_pointer.operands[2] != current_type)
+        return error.InvalidInstruction;
+
+    const result_index = try context.idIndex(operands[1]);
+    if (context.buffer_addresses[result_index] != null)
+        return error.DuplicateId;
+    context.buffer_addresses[result_index] = .{
+        .resource = base.resource,
+        .byte_offset = byte_offset,
+        .pointee_type = current_type,
+    };
+}
+
+fn unsigned32Type(context: *Context) !ir.id.TypeId {
+    return context.builder.internType(.{ .integer = .{ .bits = 32, .signedness = .unsigned } });
+}
+
+fn unsignedOffsetValue(context: *Context, block: ir.id.BlockId, spv_id: u32) !ir.id.ValueId {
+    const value = try context.resolveValue(spv_id);
+    const type_id = context.module.typeOf(value) orelse return error.InvalidId;
+    const ty = context.module.types.get(type_id) orelse return error.InvalidId;
+    const integer = switch (ty.*) {
+        .integer => |integer| integer,
+        else => return error.UnsupportedType,
+    };
+    if (integer.bits != 32)
+        return error.UnsupportedType;
+    if (integer.signedness == .unsigned)
+        return value;
+
+    return (try context.builder.appendInstruction(block, try unsigned32Type(context), .{
+        .bitcast = value,
+    }, null)).?;
+}
+
+fn addByteOffset(context: *Context, block: ir.id.BlockId, current: ?ir.id.ValueId, term: ir.id.ValueId) !ir.id.ValueId {
+    const lhs = current orelse return term;
+    return (try context.builder.appendInstruction(block, try unsigned32Type(context), .{
+        .binary = .{
+            .opcode = .integer_add,
+            .lhs = lhs,
+            .rhs = term,
+        },
+    }, null)).?;
+}
+
+fn bufferByteOffset(context: *Context, address: BufferAddress) !ir.id.ValueId {
+    return address.byte_offset orelse context.builder.internConstant(try unsigned32Type(context), .{ .integer_bits = 0 });
+}
+
+fn constantIndex(context: *Context, spv_id: u32) !u32 {
+    const value = context.module.values.get(try context.resolveValue(spv_id)) orelse return error.InvalidId;
+    if (value.definition != .constant)
+        return error.InvalidInstruction;
+    const constant = context.module.constants.get(value.definition.constant) orelse return error.InvalidId;
+    if (constant.value != .integer_bits or constant.value.integer_bits > std.math.maxInt(u32))
+        return error.InvalidInstruction;
+    return @intCast(constant.value.integer_bits);
+}
+
+fn findMemberOffset(context: *const Context, structure_id: u32, member: u32) !u32 {
+    var found: ?u32 = null;
+    for (context.member_offsets.items) |entry| {
+        if (entry.structure_id != structure_id or entry.member != member)
+            continue;
+        if (found != null)
+            return error.InvalidInstruction;
+        found = entry.offset;
+    }
+    return found orelse error.InvalidInstruction;
 }
 
 fn translateFunctionControlFlow(context: *Context, spv_function: u32) !void {
@@ -971,6 +1359,13 @@ fn makeEdge(context: *Context, predecessor_label: u32, target_label: u32) !ir.mo
     var arguments: std.ArrayList(ir.id.ValueId) = .empty;
     defer arguments.deinit(context.scratch);
 
+    if (target_label != context.entry_label.?) {
+        for (context.locals.items, 0..) |_, local_index| {
+            const value = context.block_local_outputs[try context.blockLocalIndex(predecessor_label, local_index)] orelse return error.InvalidInstruction;
+            try arguments.append(context.scratch, value);
+        }
+    }
+
     for (context.phi_infos.items) |phi| {
         if (phi.target_label != target_label)
             continue;
@@ -1308,6 +1703,106 @@ test "SPIR-V: decorated vertex interfaces and load-store operations" {
     try std.testing.expect(std.mem.indexOf(u8, text, "store_interface @out_color") != null);
 }
 
+test "SPIR-V: storage buffers and promoted function locals" {
+    const assembly =
+        \\OpCapability Shader
+        \\OpMemoryModel Logical GLSL450
+        \\OpEntryPoint GLCompute %main "main"
+        \\OpExecutionMode %main LocalSize 1 1 1
+        \\OpName %index "index"
+        \\OpName %source "source"
+        \\OpName %destination "destination"
+        \\OpDecorate %source_array ArrayStride 16
+        \\OpDecorate %Source BufferBlock
+        \\OpMemberDecorate %Source 0 Offset 0
+        \\OpDecorate %source Binding 0
+        \\OpDecorate %source DescriptorSet 0
+        \\OpDecorate %destination_array ArrayStride 16
+        \\OpDecorate %Destination BufferBlock
+        \\OpMemberDecorate %Destination 0 Offset 0
+        \\OpDecorate %destination Binding 1
+        \\OpDecorate %destination DescriptorSet 0
+        \\%void = OpTypeVoid
+        \\%fn_void = OpTypeFunction %void
+        \\%int = OpTypeInt 32 1
+        \\%uint = OpTypeInt 32 0
+        \\%bool = OpTypeBool
+        \\%vec4 = OpTypeVector %uint 4
+        \\%uint_4 = OpConstant %uint 4
+        \\%source_array = OpTypeArray %vec4 %uint_4
+        \\%destination_array = OpTypeArray %vec4 %uint_4
+        \\%Source = OpTypeStruct %source_array
+        \\%Destination = OpTypeStruct %destination_array
+        \\%ptr_uniform_source = OpTypePointer Uniform %Source
+        \\%ptr_uniform_destination = OpTypePointer Uniform %Destination
+        \\%ptr_uniform_vec4 = OpTypePointer Uniform %vec4
+        \\%ptr_function_int = OpTypePointer Function %int
+        \\%int_0 = OpConstant %int 0
+        \\%int_1 = OpConstant %int 1
+        \\%int_4 = OpConstant %int 4
+        \\%source = OpVariable %ptr_uniform_source Uniform
+        \\%destination = OpVariable %ptr_uniform_destination Uniform
+        \\%main = OpFunction %void None %fn_void
+        \\    %entry = OpLabel
+        \\    %index = OpVariable %ptr_function_int Function
+        \\    OpStore %index %int_0
+        \\    OpBranch %header
+        \\    %header = OpLabel
+        \\    OpLoopMerge %exit %continue None
+        \\    OpBranch %condition
+        \\    %condition = OpLabel
+        \\    %current = OpLoad %int %index
+        \\    %less = OpSLessThan %bool %current %int_4
+        \\    OpBranchConditional %less %body %exit
+        \\    %body = OpLabel
+        \\    %source_index = OpLoad %int %index
+        \\    %source_ptr = OpAccessChain %ptr_uniform_vec4 %source %int_0 %source_index
+        \\    %value = OpLoad %vec4 %source_ptr
+        \\    %destination_index = OpLoad %int %index
+        \\    %destination_ptr = OpAccessChain %ptr_uniform_vec4 %destination %int_0 %destination_index
+        \\    OpStore %destination_ptr %value
+        \\    OpBranch %continue
+        \\    %continue = OpLabel
+        \\    %old_index = OpLoad %int %index
+        \\    %next_index = OpIAdd %int %old_index %int_1
+        \\    OpStore %index %next_index
+        \\    OpBranch %header
+        \\    %exit = OpLabel
+        \\    OpReturn
+        \\OpFunctionEnd
+    ;
+    const words = try assembleSpirv(std.testing.allocator, assembly);
+    defer std.testing.allocator.free(words);
+
+    var module = try translate(std.testing.allocator, words, .{ .entry_point = "main" });
+    defer module.deinit();
+
+    try std.testing.expectEqual(@as(usize, 2), module.resources.entries.items.len);
+    try std.testing.expect(module.properties.explicit_resource_offsets);
+    try std.testing.expect(module.properties.no_local_memory);
+
+    const source = module.resources.get(ir.id.ResourceId.fromIndex(0)).?;
+    const destination = module.resources.get(ir.id.ResourceId.fromIndex(1)).?;
+    try std.testing.expectEqual(ir.types.ResourceKind.storage_buffer, source.kind);
+    try std.testing.expectEqual(@as(u32, 0), source.binding);
+    try std.testing.expectEqual(@as(u32, 1), destination.binding);
+
+    const function = module.functions.get(module.entry_point.?).?;
+    try std.testing.expectEqual(@as(usize, 6), function.blocks.items.len);
+    try std.testing.expectEqual(@as(usize, 0), module.blocks.get(function.blocks.items[0]).?.parameters.items.len);
+    for (function.blocks.items[1..]) |block_id|
+        try std.testing.expectEqual(@as(usize, 1), module.blocks.get(block_id).?.parameters.items.len);
+
+    const text = try ir.printer.allocPrint(std.testing.allocator, &module);
+    defer std.testing.allocator.free(text);
+    try std.testing.expect(std.mem.indexOf(u8, text, "@source: struct[array[vec4[u32], 4]] = storage_buffer[set(0), binding(0)]") != null);
+    try std.testing.expect(std.mem.indexOf(u8, text, "load_buffer @source") != null);
+    try std.testing.expect(std.mem.indexOf(u8, text, "store_buffer @destination") != null);
+
+    var parsed = try ir.parser.parseString(std.testing.allocator, text);
+    defer parsed.deinit();
+}
+
 test "SPIR-V: fragment execution modes and translated properties" {
     const assembly =
         \\OpCapability Shader
@@ -1570,6 +2065,34 @@ test "SPIR-V: operation mappings to backend-agnostic IR" {
     try std.testing.expectEqualSlices(u32, &.{1}, extract.operation.composite_extract.indices);
 }
 
+test "SPIR-V: unknown opcode reports an error without formatting the enum" {
+    const assembly =
+        \\OpCapability Shader
+        \\OpMemoryModel Logical GLSL450
+        \\OpEntryPoint GLCompute %main "main"
+        \\OpExecutionMode %main LocalSize 1 1 1
+        \\%void = OpTypeVoid
+        \\%fn_void = OpTypeFunction %void
+        \\%main = OpFunction %void None %fn_void
+        \\    %entry = OpLabel
+        \\    OpNop
+        \\    OpReturn
+        \\OpFunctionEnd
+    ;
+    const words = try assembleSpirv(std.testing.allocator, assembly);
+    defer std.testing.allocator.free(words);
+
+    const nop_word: u32 = (@as(u32, 1) << 16) | @intFromEnum(spirv.Opcode.nop);
+    for (words[spirv.header_word_count..]) |*word| {
+        if (word.* != nop_word)
+            continue;
+        word.* = (@as(u32, 1) << 16) | 999;
+        break;
+    } else return error.MissingNop;
+
+    try std.testing.expectError(error.UnsupportedOpcode, translate(std.testing.allocator, words, .{ .entry_point = "main" }));
+}
+
 test "SPIR-V: structured loop and OpPhi back edge" {
     const assembly =
         \\OpCapability Shader

src/intel/compiler/lower/lower.zig

@@ -408,6 +408,7 @@ const LoweringState = struct {
             .store_interface => |operation| try self.lowerStoreInterface(block_id, source_instruction.result, operation),
             .composite_construct => |operation| try self.lowerCompositeConstruct(source_instruction.result, operation),
             .composite_extract => |operation| try self.lowerCompositeExtract(source_instruction.result, operation),
+            .load_buffer, .store_buffer => return Error.UnsupportedOperation,
             .call => return Error.UnsanitizedModule,
         }
     }
@@ -1494,6 +1495,20 @@ test "[ir] Lower: unsupported operations" {
         \\    }
         \\}
     , Error.UnsupportedType);
+
+    try expectLoweringError(
+        \\shader vertex @main
+        \\{
+        \\    @storage: u32 = storage_buffer[set(0), binding(0)]
+        \\    %offset: constant u32 = 0
+        \\    fn @main() -> void
+        \\    {
+        \\        .entry():
+        \\            %value: u32 = load_buffer @storage, %offset
+        \\            return
+        \\    }
+        \\}
+    , Error.UnsupportedOperation);
 }
 
 test "[ir] Lower: unreachable terminator" {

src/software/SoftDescriptorSet.zig

@@ -393,7 +393,9 @@ pub fn write(interface: *Interface, write_data: vk.WriteDescriptorSet) VkError!v
                     const buffer = try NonDispatchable(Buffer).fromHandleObject(buffer_info.buffer);
                     desc.object = @as(*SoftBuffer, @alignCast(@fieldParentPtr("interface", buffer)));
                     if (desc.size == vk.WHOLE_SIZE) {
-                        desc.size = if (buffer.memory) |memory| memory.size - desc.offset else return VkError.InvalidDeviceMemoryDrv;
+                        if (desc.offset > buffer.size)
+                            return VkError.ValidationFailed;
+                        desc.size = buffer.size - desc.offset;
                     }
                 }
             }

src/software/device/ComputeDispatcher.zig

@@ -1,351 +0,0 @@
-const std = @import("std");
-const base = @import("base");
-const spv = @import("spv");
-
-const ExecutionDevice = @import("Device.zig");
-const PipelineState = ExecutionDevice.PipelineState;
-
-const SoftDevice = @import("../SoftDevice.zig");
-const SoftPipeline = @import("../SoftPipeline.zig");
-const ir_compute = @import("../interpreter/compute.zig");
-
-const VkError = base.VkError;
-const SpvRuntimeError = spv.Runtime.RuntimeError;
-
-const Self = @This();
-
-const RunData = struct {
-    self: *Self,
-    batch_id: usize,
-    group_count: usize,
-    base_group_x: usize,
-    base_group_y: usize,
-    base_group_z: usize,
-    group_count_x: usize,
-    group_count_y: usize,
-    group_count_z: usize,
-    invocations_per_workgroup: usize,
-    local_size: @Vector(3, u32),
-    pipeline: *SoftPipeline,
-};
-
-device: *SoftDevice,
-state: *PipelineState,
-batch_size: usize,
-
-invocation_index: std.atomic.Value(usize),
-
-early_dump: ?u32,
-final_dump: ?u32,
-
-pub fn init(device: *SoftDevice, state: *PipelineState) Self {
-    return .{
-        .device = device,
-        .state = state,
-        .batch_size = 0,
-        .invocation_index = .init(0),
-        .early_dump = base.config.soft_compute_dump_early_results_table,
-        .final_dump = base.config.soft_compute_dump_final_results_table,
-    };
-}
-
-pub fn dispatch(self: *Self, group_count_x: u32, group_count_y: u32, group_count_z: u32) VkError!void {
-    try self.dispatchBase(0, 0, 0, group_count_x, group_count_y, group_count_z);
-}
-
-fn getLocalSize(rt: *spv.Runtime, allocator: std.mem.Allocator, spv_module: *const spv.Module) VkError!@Vector(3, u32) {
-    if (rt.getWorkgroupSize(allocator) catch return VkError.ValidationFailed) |workgroup_size| {
-        return workgroup_size;
-    }
-
-    return .{
-        spv_module.reflection_infos.local_size_x,
-        spv_module.reflection_infos.local_size_y,
-        spv_module.reflection_infos.local_size_z,
-    };
-}
-
-pub fn dispatchBase(self: *Self, base_group_x: u32, base_group_y: u32, base_group_z: u32, group_count_x: u32, group_count_y: u32, group_count_z: u32) VkError!void {
-    const group_count_xy = std.math.mul(usize, group_count_x, group_count_y) catch return VkError.ValidationFailed;
-    const group_count = std.math.mul(usize, group_count_xy, group_count_z) catch return VkError.ValidationFailed;
-
-    const pipeline = self.state.pipeline orelse return VkError.InvalidPipelineDrv;
-    const shader = pipeline.stages.getPtr(.compute) orelse return VkError.InvalidPipelineDrv;
-
-    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 using {s} interpreter", .{ ms / 1000, if (comptime base.config.soft_ir_interpreter) "IR" else "SPIR-V" });
-    };
-
-    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;
-
-        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;
-    }
-}
-
-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 (comptime base.config.logs == .verbose) {
-            if (@errorReturnTrace()) |trace| {
-                std.debug.dumpErrorReturnTrace(trace);
-            }
-        }
-    };
-}
-
-inline fn run(data: RunData) !void {
-    const allocator = data.self.device.interface.device_allocator.allocator();
-    const io = data.self.device.interface.io();
-
-    const shader = data.pipeline.stages.getPtrAssertContains(.compute);
-    const rt = &shader.runtimes[data.batch_id].rt;
-
-    const entry = try rt.getEntryPointByName(shader.entry);
-    const uses_control_barrier = rt.mod.reflection_infos.has_control_barriers or rt.mod.reflection_infos.has_atomics;
-
-    var barrier_runtimes: []spv.Runtime = &.{};
-    var barrier_statuses: []spv.Runtime.EntryPointStatus = &.{};
-    var initialized_barrier_runtimes: usize = 0;
-    defer {
-        for (barrier_runtimes[0..initialized_barrier_runtimes]) |*barrier_rt| {
-            barrier_rt.resetInvocation(allocator);
-            barrier_rt.deinit(allocator);
-        }
-        allocator.free(barrier_runtimes);
-        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;
-
-        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;
-
-        const group_count_vec = @Vector(3, u32){
-            @as(u32, @intCast(data.group_count_x)),
-            @as(u32, @intCast(data.group_count_y)),
-            @as(u32, @intCast(data.group_count_z)),
-        };
-        const group_id_vec = @Vector(3, u32){
-            @as(u32, @intCast(data.base_group_x + group_x)),
-            @as(u32, @intCast(data.base_group_y + group_y)),
-            @as(u32, @intCast(data.base_group_z + group_z)),
-        };
-
-        if (uses_control_barrier) {
-            try runBarrierWorkgroup(data, barrier_runtimes, barrier_statuses, entry, group_count_vec, group_id_vec);
-            continue;
-        }
-
-        const workgroup_memory = try rt.createWorkgroupMemory(allocator);
-        defer rt.destroyWorkgroupMemory(allocator, workgroup_memory);
-
-        rt.resetInvocation(allocator);
-        try rt.bindWorkgroupMemory(workgroup_memory);
-        try setupWorkgroupBuiltins(data.self, rt, data.local_size, group_count_vec, group_id_vec);
-
-        for (0..data.invocations_per_workgroup) |i| {
-            rt.resetInvocation(allocator);
-
-            const invocation_index = data.self.invocation_index.fetchAdd(1, .monotonic);
-
-            try setupSubgroupBuiltins(data.self, rt, data.local_size, .{
-                @as(u32, @intCast(data.base_group_x + group_x)),
-                @as(u32, @intCast(data.base_group_y + group_y)),
-                @as(u32, @intCast(data.base_group_z + group_z)),
-            }, i);
-
-            if (data.self.early_dump != null and data.self.early_dump.? == invocation_index) {
-                @branchHint(.cold);
-                try dumpResultsTable(allocator, io, rt, true);
-            }
-
-            rt.callEntryPoint(allocator, entry) catch |err| switch (err) {
-                // Some errors can be ignored
-                SpvRuntimeError.OutOfBounds => {},
-                SpvRuntimeError.Killed => continue,
-                else => return err,
-            };
-            try flushWorkgroupMemory(rt, workgroup_memory);
-            try rt.flushDescriptorSets(allocator);
-
-            if (data.self.final_dump != null and data.self.final_dump.? == invocation_index) {
-                @branchHint(.cold);
-                try dumpResultsTable(allocator, io, rt, false);
-            }
-        }
-    }
-}
-
-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,
-    statuses: []spv.Runtime.EntryPointStatus,
-    entry: spv.SpvWord,
-    group_count: @Vector(3, u32),
-    group_id: @Vector(3, u32),
-) !void {
-    const allocator = data.self.device.interface.device_allocator.allocator();
-
-    const workgroup_memory = try runtimes[0].createWorkgroupMemory(allocator);
-    defer runtimes[0].destroyWorkgroupMemory(allocator, workgroup_memory);
-    for (runtimes, 0..) |*rt, i| {
-        rt.resetInvocation(allocator);
-        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);
-        statuses[i] = try rt.beginEntryPoint(allocator, entry);
-        try flushWorkgroupMemory(rt, workgroup_memory);
-        try rt.flushDescriptorSets(allocator);
-    }
-
-    while (true) {
-        var pending = false;
-        for (statuses) |status| {
-            if (status == .barrier) {
-                pending = true;
-                break;
-            }
-        }
-        if (!pending)
-            break;
-
-        for (runtimes, 0..) |*rt, i| {
-            if (statuses[i] == .completed)
-                continue;
-            try rt.bindWorkgroupMemory(workgroup_memory);
-            statuses[i] = try rt.continueEntryPoint(allocator);
-            try flushWorkgroupMemory(rt, workgroup_memory);
-            try rt.flushDescriptorSets(allocator);
-        }
-    }
-}
-
-fn flushWorkgroupMemory(rt: *spv.Runtime, workgroup_memory: []const spv.Runtime.WorkgroupMemory) spv.Runtime.RuntimeError!void {
-    for (workgroup_memory) |memory| {
-        _ = try (try rt.results[memory.result].getValue()).read(memory.bytes);
-    }
-}
-
-fn dumpResultsTable(allocator: std.mem.Allocator, io: std.Io, rt: *spv.Runtime, comptime is_early: bool) !void {
-    @branchHint(.cold);
-    const file = try std.Io.Dir.cwd().createFile(
-        io,
-        std.fmt.comptimePrint("{s}_compute_result_table_dump.txt", .{if (is_early) "early" else "final"}),
-        .{ .truncate = true },
-    );
-    defer file.close(io);
-    var buffer = [_]u8{0} ** 1024;
-    var writer = file.writer(io, buffer[0..]);
-    try rt.dumpResultsTable(allocator, &writer.interface);
-}
-
-fn setupWorkgroupBuiltins(self: *Self, rt: *spv.Runtime, local_size: @Vector(3, u32), group_count: @Vector(3, u32), group_id: @Vector(3, u32)) spv.Runtime.RuntimeError!void {
-    const allocator = self.device.interface.device_allocator.allocator();
-
-    rt.writeBuiltIn(allocator, std.mem.asBytes(&local_size), .WorkgroupSize) catch |err| switch (err) {
-        SpvRuntimeError.NotFound => {},
-        else => return err,
-    };
-    rt.writeBuiltIn(allocator, std.mem.asBytes(&group_count), .NumWorkgroups) catch |err| switch (err) {
-        SpvRuntimeError.NotFound => {},
-        else => return err,
-    };
-    rt.writeBuiltIn(allocator, std.mem.asBytes(&group_id), .WorkgroupId) catch |err| switch (err) {
-        SpvRuntimeError.NotFound => {},
-        else => return err,
-    };
-}
-
-fn setupSubgroupBuiltins(self: *Self, rt: *spv.Runtime, local_size: @Vector(3, u32), group_id: @Vector(3, u32), local_invocation_index: usize) spv.Runtime.RuntimeError!void {
-    const allocator = self.device.interface.device_allocator.allocator();
-
-    const local_base = local_size * group_id;
-    var local_invocation = @Vector(3, u32){ 0, 0, 0 };
-
-    var idx: u32 = @intCast(local_invocation_index);
-    local_invocation[2] = @divTrunc(idx, local_size[0] * local_size[1]);
-    idx -= local_invocation[2] * local_size[0] * local_size[1];
-    local_invocation[1] = @divTrunc(idx, local_size[0]);
-    idx -= local_invocation[1] * local_size[0];
-    local_invocation[0] = idx;
-
-    const global_invocation_index = local_base + local_invocation;
-    const local_invocation_index_u32: u32 = @intCast(local_invocation_index);
-
-    rt.writeBuiltIn(allocator, std.mem.asBytes(&local_invocation), .LocalInvocationId) catch |err| switch (err) {
-        SpvRuntimeError.NotFound => {},
-        else => return err,
-    };
-    rt.writeBuiltIn(allocator, std.mem.asBytes(&local_invocation_index_u32), .LocalInvocationIndex) catch |err| switch (err) {
-        SpvRuntimeError.NotFound => {},
-        else => return err,
-    };
-    rt.writeBuiltIn(allocator, std.mem.asBytes(&global_invocation_index), .GlobalInvocationId) catch |err| switch (err) {
-        SpvRuntimeError.NotFound => {},
-        else => return err,
-    };
-}

src/software/device/Device.zig

@@ -4,11 +4,13 @@ const base = @import("base");
 const lib = @import("../lib.zig");
 const spv = @import("spv");
 
+const VkError = base.VkError;
+
 const SoftDescriptorSet = @import("../SoftDescriptorSet.zig");
 const SoftDevice = @import("../SoftDevice.zig");
 const SoftPipeline = @import("../SoftPipeline.zig");
 
-const ComputeDispatcher = @import("ComputeDispatcher.zig");
+const ComputeDispatcher = @import("compute/ComputeDispatcher.zig");
 const Renderer = @import("Renderer.zig");
 
 const Self = @This();
@@ -36,6 +38,46 @@ pub const PipelineState = struct {
     },
 };
 
+pub fn mapStorageBuffer(state: *const PipelineState, set: u32, binding: u32) VkError!?[]u8 {
+    const set_index: usize = set;
+    if (set_index >= state.sets.len)
+        return null;
+    const descriptor_set = state.sets[set_index] orelse return null;
+
+    const binding_index: usize = binding;
+    if (binding_index >= descriptor_set.descriptors.len or binding_index >= descriptor_set.interface.layout.bindings.len)
+        return null;
+
+    const binding_layout = descriptor_set.interface.layout.bindings[binding_index];
+    const dynamic_offset: vk.DeviceSize = switch (binding_layout.descriptor_type) {
+        .storage_buffer_dynamic => blk: {
+            if (binding_layout.dynamic_index >= state.dynamic_offsets[set_index].len)
+                return VkError.ValidationFailed;
+            break :blk state.dynamic_offsets[set_index][binding_layout.dynamic_index];
+        },
+        .storage_buffer => 0,
+        else => return null,
+    };
+
+    const descriptors = switch (descriptor_set.descriptors[binding_index]) {
+        .buffer => |descriptors| descriptors,
+        else => return null,
+    };
+    if (descriptors.len == 0)
+        return null;
+
+    const descriptor = descriptors[0];
+    const buffer = descriptor.object orelse return null;
+    const effective_offset = std.math.add(vk.DeviceSize, descriptor.offset, dynamic_offset) catch return VkError.ValidationFailed;
+    if (effective_offset > buffer.interface.size)
+        return VkError.ValidationFailed;
+
+    const logical_remaining = buffer.interface.size - effective_offset;
+    if (descriptor.size > logical_remaining)
+        return VkError.ValidationFailed;
+    return try buffer.mapAsSliceWithAddedOffset(u8, effective_offset, descriptor.size);
+}
+
 compute: ComputeDispatcher,
 renderer: Renderer,
 

src/software/device/compute/ComputeDispatcher.zig

@@ -0,0 +1,177 @@
+const std = @import("std");
+const base = @import("base");
+const spv = @import("spv");
+
+const ExecutionDevice = @import("../Device.zig");
+const PipelineState = ExecutionDevice.PipelineState;
+
+const SoftDevice = @import("../../SoftDevice.zig");
+const ir_interpreter = @import("ir_interpreter.zig");
+const spirv_interpreter = @import("spirv_interpreter.zig");
+
+const VkError = base.VkError;
+
+const Self = @This();
+
+pub const Batch = struct {
+    worker_index: usize,
+    worker_count: usize,
+    total_groups: usize,
+    base_group: [3]usize,
+    group_count: [3]usize,
+
+    pub fn groupId(self: Batch, linear_index: usize) [3]usize {
+        const groups_xy = self.group_count[0] * self.group_count[1];
+        const group_z = linear_index / groups_xy;
+        const remainder = linear_index - group_z * groups_xy;
+        const group_y = remainder / self.group_count[0];
+        const group_x = remainder - group_y * self.group_count[0];
+        return .{
+            self.base_group[0] + group_x,
+            self.base_group[1] + group_y,
+            self.base_group[2] + group_z,
+        };
+    }
+};
+
+const BackendContext = if (base.config.soft_ir_interpreter) ir_interpreter.Context else spirv_interpreter.SpvContext;
+
+device: *SoftDevice,
+state: *PipelineState,
+
+invocation_index: std.atomic.Value(usize),
+
+early_dump: ?u32,
+final_dump: ?u32,
+
+pub fn init(device: *SoftDevice, state: *PipelineState) Self {
+    return .{
+        .device = device,
+        .state = state,
+        .invocation_index = .init(0),
+        .early_dump = base.config.soft_compute_dump_early_results_table,
+        .final_dump = base.config.soft_compute_dump_final_results_table,
+    };
+}
+
+pub fn dispatch(self: *Self, group_count_x: u32, group_count_y: u32, group_count_z: u32) VkError!void {
+    try self.dispatchBase(0, 0, 0, group_count_x, group_count_y, group_count_z);
+}
+
+fn dispatchBatches(
+    io: std.Io,
+    context: anytype,
+    worker_count: usize,
+    total_groups: usize,
+    base_group: [3]usize,
+    group_count: [3]usize,
+    comptime worker: anytype,
+) !void {
+    if (total_groups == 0)
+        return;
+    if (worker_count == 0)
+        return error.NoWorkers;
+
+    const active_workers = @min(worker_count, total_groups);
+    var group: std.Io.Group = .init;
+    for (0..active_workers) |worker_index| {
+        group.async(io, worker, .{ context, Batch{
+            .worker_index = worker_index,
+            .worker_count = active_workers,
+            .total_groups = total_groups,
+            .base_group = base_group,
+            .group_count = group_count,
+        } });
+    }
+    try group.await(io);
+}
+
+fn getLocalSize(rt: *spv.Runtime, allocator: std.mem.Allocator, spv_module: *const spv.Module) VkError!@Vector(3, u32) {
+    if (rt.getWorkgroupSize(allocator) catch return VkError.ValidationFailed) |workgroup_size| {
+        return workgroup_size;
+    }
+
+    return .{
+        spv_module.reflection_infos.local_size_x,
+        spv_module.reflection_infos.local_size_y,
+        spv_module.reflection_infos.local_size_z,
+    };
+}
+
+pub fn dispatchBase(self: *Self, base_group_x: u32, base_group_y: u32, base_group_z: u32, group_count_x: u32, group_count_y: u32, group_count_z: u32) VkError!void {
+    const group_count_xy = std.math.mul(usize, group_count_x, group_count_y) catch return VkError.ValidationFailed;
+    const group_count = std.math.mul(usize, group_count_xy, group_count_z) catch return VkError.ValidationFailed;
+
+    const pipeline = self.state.pipeline orelse return VkError.InvalidPipelineDrv;
+    const shader = pipeline.stages.getPtr(.compute) orelse return VkError.InvalidPipelineDrv;
+
+    const io = self.device.interface.io();
+    const allocator = self.device.interface.device_allocator.allocator();
+    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 using {s} interpreter", .{ ms / 1000, if (comptime base.config.soft_ir_interpreter) "IR" else "SPIR-V" });
+    };
+
+    var context: BackendContext = if (comptime base.config.soft_ir_interpreter)
+        try ir_interpreter.prepare(allocator, shader, self.state, io)
+    else blk: {
+        if (shader.runtimes.len == 0)
+            return VkError.InvalidPipelineDrv;
+        const spv_module = &shader.module.module;
+        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);
+        break :blk .{
+            .dispatcher = self,
+            .pipeline = pipeline,
+            .invocations_per_workgroup = invocations_per_workgroup,
+            .local_size = local_size,
+        };
+    };
+    defer if (comptime base.config.soft_ir_interpreter)
+        context.deinit(allocator);
+
+    const worker_count = if (comptime base.config.soft_ir_interpreter)
+        shader.runtimes.len
+    else if (shader.module.module.reflection_infos.has_atomics)
+        1
+    else
+        shader.runtimes.len;
+
+    dispatchBatches(
+        io,
+        context,
+        worker_count,
+        group_count,
+        .{ base_group_x, base_group_y, base_group_z },
+        .{ group_count_x, group_count_y, group_count_z },
+        runWrapper,
+    ) catch |err| switch (err) {
+        error.NoWorkers => return VkError.InvalidPipelineDrv,
+        else => return VkError.DeviceLost,
+    };
+}
+
+fn runWrapper(context: BackendContext, batch: Batch) void {
+    @call(.always_inline, run, .{ context, batch }) catch |err| {
+        std.log.scoped(.ComputeDispatcher).err("{s} interpreter runtime caught a '{s}'", .{
+            if (comptime base.config.soft_ir_interpreter) "IR" else "SPIR-V",
+            @errorName(err),
+        });
+        if (comptime base.config.logs == .verbose) {
+            if (@errorReturnTrace()) |trace|
+                std.debug.dumpErrorReturnTrace(trace);
+        }
+    };
+}
+
+inline fn run(context: BackendContext, batch: Batch) !void {
+    if (comptime base.config.soft_ir_interpreter) {
+        return ir_interpreter.runBatch(context, batch);
+    } else {
+        return spirv_interpreter.runBatch(context, batch);
+    }
+}

src/software/device/compute/ir_interpreter.zig

@@ -0,0 +1,96 @@
+const std = @import("std");
+const base = @import("base");
+const shader_ir = @import("shader_ir");
+
+const ExecutionDevice = @import("../Device.zig");
+const PipelineState = ExecutionDevice.PipelineState;
+const Batch = @import("ComputeDispatcher.zig").Batch;
+const Shader = @import("../../interpreter/Shader.zig");
+
+const VkError = base.VkError;
+const ir = shader_ir.ir;
+
+pub const Context = struct {
+    shader: *Shader,
+    io: std.Io,
+    local_size: [3]u32,
+    local_xy: usize,
+    local_count: usize,
+    global_id: ?ir.id.InterfaceVariableId,
+    resource_buffers: []?[]u8,
+
+    pub fn deinit(self: *Context, allocator: std.mem.Allocator) void {
+        allocator.free(self.resource_buffers);
+        self.* = undefined;
+    }
+};
+
+pub fn prepare(allocator: std.mem.Allocator, shader: *Shader, state: *const PipelineState, io: std.Io) VkError!Context {
+    const local_size = shader.workgroup_size orelse return VkError.ValidationFailed;
+    const local_xy = std.math.mul(usize, local_size[0], local_size[1]) catch return VkError.ValidationFailed;
+    const local_count = std.math.mul(usize, local_xy, local_size[2]) catch return VkError.ValidationFailed;
+    if (shader.runtimes.len == 0)
+        return VkError.InvalidPipelineDrv;
+
+    const resource_buffers = allocator.alloc(?[]u8, shader.program.resources.len) catch return VkError.OutOfDeviceMemory;
+    errdefer allocator.free(resource_buffers);
+    @memset(resource_buffers, null);
+    for (shader.program.resources, resource_buffers) |optional_resource, *buffer| {
+        const resource = optional_resource orelse continue;
+        if (resource.kind == .storage_buffer)
+            buffer.* = try ExecutionDevice.mapStorageBuffer(state, resource.set, resource.binding);
+    }
+
+    return .{
+        .shader = shader,
+        .io = io,
+        .local_size = local_size,
+        .local_xy = local_xy,
+        .local_count = local_count,
+        .global_id = findGlobalInvocationId(&shader.program),
+        .resource_buffers = resource_buffers,
+    };
+}
+
+pub fn runBatch(context: Context, batch: Batch) !void {
+    const shader = context.shader;
+    if (batch.worker_index >= shader.runtimes.len)
+        return VkError.InvalidPipelineDrv;
+
+    const slot = &shader.runtimes[batch.worker_index];
+    slot.mutex.lock(context.io) catch return VkError.DeviceLost;
+    defer slot.mutex.unlock(context.io);
+
+    const runtime = &slot.runtime;
+    var group_index = batch.worker_index;
+    while (group_index < batch.total_groups) : (group_index += batch.worker_count) {
+        const group_id = batch.groupId(group_index);
+        const group_x = std.math.cast(u32, group_id[0]) orelse return VkError.ValidationFailed;
+        const group_y = std.math.cast(u32, group_id[1]) orelse return VkError.ValidationFailed;
+        const group_z = std.math.cast(u32, group_id[2]) orelse return VkError.ValidationFailed;
+
+        for (0..context.local_count) |local_index| {
+            if (context.global_id) |variable| {
+                const local_z = local_index / context.local_xy;
+                const local_remainder = local_index - local_z * context.local_xy;
+                const local_y = local_remainder / context.local_size[0];
+                const local_x = local_remainder - local_y * context.local_size[0];
+                try runtime.writeInput(&shader.program, variable, &.{
+                    group_x * context.local_size[0] + @as(u32, @intCast(local_x)),
+                    group_y * context.local_size[1] + @as(u32, @intCast(local_y)),
+                    group_z * context.local_size[2] + @as(u32, @intCast(local_z)),
+                });
+            }
+            _ = try runtime.run(&shader.program, .{ .resource_buffers = context.resource_buffers });
+        }
+    }
+}
+
+fn findGlobalInvocationId(program: *const @import("../../interpreter/Program.zig")) ?ir.id.InterfaceVariableId {
+    for (program.interfaces, 0..) |optional_binding, index| {
+        const binding = optional_binding orelse continue;
+        if (binding.direction == .input and binding.semantic == .builtin and binding.semantic.builtin == .global_invocation_id)
+            return ir.id.InterfaceVariableId.fromIndex(index);
+    }
+    return null;
+}

src/software/device/compute/spirv_interpreter.zig

@@ -0,0 +1,227 @@
+const std = @import("std");
+const spv = @import("spv");
+
+const SpvRuntimeError = spv.Runtime.RuntimeError;
+
+const ExecutionDevice = @import("../Device.zig");
+const SoftPipeline = @import("../../SoftPipeline.zig");
+const Dispatcher = @import("ComputeDispatcher.zig");
+const Batch = Dispatcher.Batch;
+
+pub const SpvContext = struct {
+    dispatcher: *Dispatcher,
+    pipeline: *SoftPipeline,
+    invocations_per_workgroup: usize,
+    local_size: @Vector(3, u32),
+};
+
+pub fn runBatch(context: SpvContext, batch: Batch) !void {
+    const dispatcher = context.dispatcher;
+    const allocator = dispatcher.device.interface.device_allocator.allocator();
+    const io = dispatcher.device.interface.io();
+
+    const shader = context.pipeline.stages.getPtrAssertContains(.compute);
+    const rt = &shader.runtimes[batch.worker_index].rt;
+
+    const entry = try rt.getEntryPointByName(shader.entry);
+    const uses_control_barrier = rt.mod.reflection_infos.has_control_barriers or rt.mod.reflection_infos.has_atomics;
+
+    var barrier_runtimes: []spv.Runtime = &.{};
+    var barrier_statuses: []spv.Runtime.EntryPointStatus = &.{};
+    var initialized_barrier_runtimes: usize = 0;
+    defer {
+        for (barrier_runtimes[0..initialized_barrier_runtimes]) |*barrier_rt| {
+            barrier_rt.resetInvocation(allocator);
+            barrier_rt.deinit(allocator);
+        }
+        allocator.free(barrier_runtimes);
+        allocator.free(barrier_statuses);
+    }
+
+    if (uses_control_barrier) {
+        barrier_runtimes = try allocator.alloc(spv.Runtime, context.invocations_per_workgroup);
+        barrier_statuses = try allocator.alloc(spv.Runtime.EntryPointStatus, context.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(dispatcher, barrier_rt);
+        }
+    } else {
+        try prepareRuntime(dispatcher, rt);
+    }
+
+    const group_count_vec = @Vector(3, u32){
+        @intCast(batch.group_count[0]),
+        @intCast(batch.group_count[1]),
+        @intCast(batch.group_count[2]),
+    };
+    var group_index = batch.worker_index;
+    while (group_index < batch.total_groups) : (group_index += batch.worker_count) {
+        const group_id = batch.groupId(group_index);
+        const group_id_vec = @Vector(3, u32){
+            @intCast(group_id[0]),
+            @intCast(group_id[1]),
+            @intCast(group_id[2]),
+        };
+
+        if (uses_control_barrier) {
+            try runBarrierWorkgroup(context, barrier_runtimes, barrier_statuses, entry, group_count_vec, group_id_vec);
+            continue;
+        }
+
+        const workgroup_memory = try rt.createWorkgroupMemory(allocator);
+        defer rt.destroyWorkgroupMemory(allocator, workgroup_memory);
+
+        rt.resetInvocation(allocator);
+        try rt.bindWorkgroupMemory(workgroup_memory);
+        try setupWorkgroupBuiltins(dispatcher, rt, context.local_size, group_count_vec, group_id_vec);
+
+        for (0..context.invocations_per_workgroup) |i| {
+            rt.resetInvocation(allocator);
+
+            const invocation_index = dispatcher.invocation_index.fetchAdd(1, .monotonic);
+
+            try setupSubgroupBuiltins(dispatcher, rt, context.local_size, group_id_vec, i);
+
+            if (dispatcher.early_dump != null and dispatcher.early_dump.? == invocation_index) {
+                @branchHint(.cold);
+                try dumpResultsTable(allocator, io, rt, true);
+            }
+
+            rt.callEntryPoint(allocator, entry) catch |err| switch (err) {
+                SpvRuntimeError.OutOfBounds => {},
+                SpvRuntimeError.Killed => continue,
+                else => return err,
+            };
+            try flushWorkgroupMemory(rt, workgroup_memory);
+            try rt.flushDescriptorSets(allocator);
+
+            if (dispatcher.final_dump != null and dispatcher.final_dump.? == invocation_index) {
+                @branchHint(.cold);
+                try dumpResultsTable(allocator, io, rt, false);
+            }
+        }
+    }
+}
+
+fn prepareRuntime(dispatcher: *Dispatcher, rt: *spv.Runtime) !void {
+    const allocator = dispatcher.device.interface.device_allocator.allocator();
+
+    rt.resetInvocation(allocator);
+    if (rt.specialization_constants.count() != 0)
+        try rt.applySpecializationInvocationLayout(allocator);
+    try ExecutionDevice.writeDescriptorSets(dispatcher.state, rt);
+    try rt.populatePushConstants(dispatcher.state.push_constant_blob[0..]);
+}
+
+fn runBarrierWorkgroup(
+    context: SpvContext,
+    runtimes: []spv.Runtime,
+    statuses: []spv.Runtime.EntryPointStatus,
+    entry: spv.SpvWord,
+    group_count: @Vector(3, u32),
+    group_id: @Vector(3, u32),
+) !void {
+    const dispatcher = context.dispatcher;
+    const allocator = dispatcher.device.interface.device_allocator.allocator();
+
+    const workgroup_memory = try runtimes[0].createWorkgroupMemory(allocator);
+    defer runtimes[0].destroyWorkgroupMemory(allocator, workgroup_memory);
+    for (runtimes, 0..) |*rt, i| {
+        rt.resetInvocation(allocator);
+        try rt.bindWorkgroupMemory(workgroup_memory);
+        try setupWorkgroupBuiltins(dispatcher, rt, context.local_size, group_count, group_id);
+        try setupSubgroupBuiltins(dispatcher, rt, context.local_size, group_id, i);
+        statuses[i] = try rt.beginEntryPoint(allocator, entry);
+        try flushWorkgroupMemory(rt, workgroup_memory);
+        try rt.flushDescriptorSets(allocator);
+    }
+
+    while (true) {
+        var pending = false;
+        for (statuses) |status| {
+            if (status == .barrier) {
+                pending = true;
+                break;
+            }
+        }
+        if (!pending)
+            break;
+
+        for (runtimes, 0..) |*rt, i| {
+            if (statuses[i] == .completed)
+                continue;
+            try rt.bindWorkgroupMemory(workgroup_memory);
+            statuses[i] = try rt.continueEntryPoint(allocator);
+            try flushWorkgroupMemory(rt, workgroup_memory);
+            try rt.flushDescriptorSets(allocator);
+        }
+    }
+}
+
+fn flushWorkgroupMemory(rt: *spv.Runtime, workgroup_memory: []const spv.Runtime.WorkgroupMemory) spv.Runtime.RuntimeError!void {
+    for (workgroup_memory) |memory| {
+        _ = try (try rt.results[memory.result].getValue()).read(memory.bytes);
+    }
+}
+
+fn dumpResultsTable(allocator: std.mem.Allocator, io: std.Io, rt: *spv.Runtime, comptime is_early: bool) !void {
+    @branchHint(.cold);
+    const file = try std.Io.Dir.cwd().createFile(
+        io,
+        std.fmt.comptimePrint("{s}_compute_result_table_dump.txt", .{if (is_early) "early" else "final"}),
+        .{ .truncate = true },
+    );
+    defer file.close(io);
+    var buffer = [_]u8{0} ** 1024;
+    var writer = file.writer(io, buffer[0..]);
+    try rt.dumpResultsTable(allocator, &writer.interface);
+}
+
+fn setupWorkgroupBuiltins(dispatcher: *Dispatcher, rt: *spv.Runtime, local_size: @Vector(3, u32), group_count: @Vector(3, u32), group_id: @Vector(3, u32)) spv.Runtime.RuntimeError!void {
+    const allocator = dispatcher.device.interface.device_allocator.allocator();
+
+    rt.writeBuiltIn(allocator, std.mem.asBytes(&local_size), .WorkgroupSize) catch |err| switch (err) {
+        SpvRuntimeError.NotFound => {},
+        else => return err,
+    };
+    rt.writeBuiltIn(allocator, std.mem.asBytes(&group_count), .NumWorkgroups) catch |err| switch (err) {
+        SpvRuntimeError.NotFound => {},
+        else => return err,
+    };
+    rt.writeBuiltIn(allocator, std.mem.asBytes(&group_id), .WorkgroupId) catch |err| switch (err) {
+        SpvRuntimeError.NotFound => {},
+        else => return err,
+    };
+}
+
+fn setupSubgroupBuiltins(dispatcher: *Dispatcher, rt: *spv.Runtime, local_size: @Vector(3, u32), group_id: @Vector(3, u32), local_invocation_index: usize) spv.Runtime.RuntimeError!void {
+    const allocator = dispatcher.device.interface.device_allocator.allocator();
+
+    const local_base = local_size * group_id;
+    var local_invocation = @Vector(3, u32){ 0, 0, 0 };
+
+    var idx: u32 = @intCast(local_invocation_index);
+    local_invocation[2] = @divTrunc(idx, local_size[0] * local_size[1]);
+    idx -= local_invocation[2] * local_size[0] * local_size[1];
+    local_invocation[1] = @divTrunc(idx, local_size[0]);
+    idx -= local_invocation[1] * local_size[0];
+    local_invocation[0] = idx;
+
+    const global_invocation_index = local_base + local_invocation;
+    const local_invocation_index_u32: u32 = @intCast(local_invocation_index);
+
+    rt.writeBuiltIn(allocator, std.mem.asBytes(&local_invocation), .LocalInvocationId) catch |err| switch (err) {
+        SpvRuntimeError.NotFound => {},
+        else => return err,
+    };
+    rt.writeBuiltIn(allocator, std.mem.asBytes(&local_invocation_index_u32), .LocalInvocationIndex) catch |err| switch (err) {
+        SpvRuntimeError.NotFound => {},
+        else => return err,
+    };
+    rt.writeBuiltIn(allocator, std.mem.asBytes(&global_invocation_index), .GlobalInvocationId) catch |err| switch (err) {
+        SpvRuntimeError.NotFound => {},
+        else => return err,
+    };
+}

src/software/interpreter/Program.zig

@@ -25,6 +25,12 @@ pub const InterfaceBinding = struct {
     span: bc.Span,
 };
 
+pub const ResourceBinding = struct {
+    kind: ir.types.ResourceKind,
+    set: u32,
+    binding: u32,
+};
+
 pub const RegisterInit = struct {
     register: bc.Register,
     value: u32,
@@ -43,6 +49,7 @@ copies: []const bc.Copy,
 branches: []const bc.Branch,
 initializers: []const RegisterInit,
 interfaces: []const ?InterfaceBinding,
+resources: []const ?ResourceBinding,
 
 pub fn compile(backing_allocator: std.mem.Allocator, module: *const module_ir.Module) !Self {
     try ir.validator.validate(module);
@@ -65,6 +72,7 @@ pub fn compile(backing_allocator: std.mem.Allocator, module: *const module_ir.Mo
         .branches = lowerer.branches.items,
         .initializers = lowerer.initializers.items,
         .interfaces = lowerer.interfaces,
+        .resources = lowerer.resources,
     };
 }
 
@@ -80,6 +88,13 @@ pub fn interfaceBinding(self: *const Self, variable: ids.InterfaceVariableId) ?I
     return self.interfaces[variable.index()];
 }
 
+pub fn resourceBinding(self: *const Self, resource: ids.ResourceId) ?ResourceBinding {
+    if (resource.index() >= self.resources.len)
+        return null;
+
+    return self.resources[resource.index()];
+}
+
 const Lowerer = struct {
     allocator: std.mem.Allocator,
     module: *const module_ir.Module,
@@ -87,6 +102,7 @@ const Lowerer = struct {
     entry_block: ids.BlockId,
     values: []?bc.Span,
     interfaces: []?InterfaceBinding,
+    resources: []?ResourceBinding,
     block_pcs: []?u32,
     register_count: usize = 0,
     scratch_count: usize = 0,
@@ -109,6 +125,14 @@ const Lowerer = struct {
         @memset(values, null);
         const interfaces = try allocator.alloc(?InterfaceBinding, module.interface_variables.entries.items.len);
         @memset(interfaces, null);
+        const resources = try allocator.alloc(?ResourceBinding, module.resources.entries.items.len);
+        for (module.resources.entries.items, resources) |entry, *binding| {
+            binding.* = if (entry) |resource| .{
+                .kind = resource.kind,
+                .set = resource.set,
+                .binding = resource.binding,
+            } else null;
+        }
         const block_pcs = try allocator.alloc(?u32, module.blocks.entries.items.len);
         @memset(block_pcs, null);
 
@@ -119,6 +143,7 @@ const Lowerer = struct {
             .entry_block = entry_block,
             .values = values,
             .interfaces = interfaces,
+            .resources = resources,
             .block_pcs = block_pcs,
         };
     }
@@ -347,10 +372,40 @@ const Lowerer = struct {
 
                 try self.emitCopy(binding.span, src);
             },
+            .load_buffer => |op| {
+                const dst = result orelse return CompileError.InvalidOperation;
+                const byte_offset = try self.bufferOffset(op.byte_offset);
+                _ = try self.storageBuffer(op.resource);
+                try self.emit(.load_buffer, dst.components, dst.base, byte_offset, bc.invalid_register, bc.invalid_register, @intFromEnum(op.resource));
+            },
+            .store_buffer => |op| {
+                if (result != null)
+                    return CompileError.InvalidOperation;
+                const src = try self.span(op.value);
+                const byte_offset = try self.bufferOffset(op.byte_offset);
+                _ = try self.storageBuffer(op.resource);
+                try self.emit(.store_buffer, src.components, src.base, byte_offset, bc.invalid_register, bc.invalid_register, @intFromEnum(op.resource));
+            },
             .call => return CompileError.UnsupportedOperation,
         }
     }
 
+    fn bufferOffset(self: *const Lowerer, id: ids.ValueId) !bc.Register {
+        const byte_offset = try self.span(id);
+        if (byte_offset.components != 1 or byte_offset.kind != .unsigned_integer)
+            return CompileError.InvalidOperation;
+        return byte_offset.base;
+    }
+
+    fn storageBuffer(self: *const Lowerer, id: ids.ResourceId) !ResourceBinding {
+        if (id.index() >= self.resources.len)
+            return CompileError.InvalidOperation;
+        const resource = self.resources[id.index()] orelse return CompileError.InvalidOperation;
+        if (resource.kind != .storage_buffer)
+            return CompileError.InvalidOperation;
+        return resource;
+    }
+
     fn lowerTerminator(self: *Lowerer, terminator: module_ir.Terminator) !void {
         switch (terminator) {
             .branch => |edge| try self.emit(.jump_edge, 1, bc.invalid_register, bc.invalid_register, bc.invalid_register, bc.invalid_register, try self.addEdge(edge)),

src/software/interpreter/Runtime.zig

@@ -6,10 +6,13 @@ const Program = @import("Program.zig");
 const ids = shader_ir.ir.id;
 
 pub const RuntimeError = error{
+    BufferOutOfBounds,
     DivisionByZero,
     IntegerOverflow,
     InvalidBytecode,
     InvalidInterface,
+    InvalidResource,
+    ResourceNotBound,
     ShiftOutOfRange,
     StepLimitExceeded,
     UnreachableExecuted,
@@ -24,6 +27,7 @@ pub const Outcome = enum {
 
 pub const RunOptions = struct {
     max_steps: usize = 1_000_000,
+    resource_buffers: []const ?[]u8 = &.{},
 };
 
 const Self = @This();
@@ -126,6 +130,8 @@ pub fn run(self: *Self, program: *const Program, options: RunOptions) RuntimeErr
             .compare_ordered_float_less => self.compareFloat(instruction, .ordered_less),
             .compare_unordered_float_less => self.compareFloat(instruction, .unordered_less),
             .select => self.select(instruction),
+            .load_buffer => try self.loadBuffer(program, options.resource_buffers, instruction),
+            .store_buffer => try self.storeBuffer(program, options.resource_buffers, instruction),
             .jump_edge => pc = try self.applyEdge(program, instruction.immediate),
             .branch => {
                 if (instruction.immediate >= program.branches.len)
@@ -278,6 +284,52 @@ fn select(self: *Self, instruction: bc.Instruction) void {
         self.registers[@as(usize, instruction.a) + component] = self.registers[@as(usize, selected) + component];
 }
 
+fn loadBuffer(self: *Self, program: *const Program, resource_buffers: []const ?[]u8, instruction: bc.Instruction) RuntimeError!void {
+    try self.validateRegisterSpan(instruction);
+    const buffer = try resourceBuffer(program, resource_buffers, instruction.immediate);
+    const bytes = try self.bufferRange(buffer, instruction);
+    for (0..instruction.components) |component| {
+        const offset = component * @sizeOf(u32);
+        self.registers[@as(usize, instruction.a) + component] = std.mem.readInt(u32, bytes[offset..][0..@sizeOf(u32)], .little);
+    }
+}
+
+fn storeBuffer(self: *const Self, program: *const Program, resource_buffers: []const ?[]u8, instruction: bc.Instruction) RuntimeError!void {
+    try self.validateRegisterSpan(instruction);
+    const buffer = try resourceBuffer(program, resource_buffers, instruction.immediate);
+    const bytes = try self.bufferRange(buffer, instruction);
+    for (0..instruction.components) |component| {
+        const offset = component * @sizeOf(u32);
+        std.mem.writeInt(u32, bytes[offset..][0..@sizeOf(u32)], self.registers[@as(usize, instruction.a) + component], .little);
+    }
+}
+
+fn validateRegisterSpan(self: *const Self, instruction: bc.Instruction) RuntimeError!void {
+    const register_end = std.math.add(usize, instruction.a, instruction.components) catch return RuntimeError.InvalidBytecode;
+    if (register_end > self.registers.len)
+        return RuntimeError.InvalidBytecode;
+}
+
+fn bufferRange(self: *const Self, buffer: []u8, instruction: bc.Instruction) RuntimeError![]u8 {
+    if (instruction.b >= self.registers.len)
+        return RuntimeError.InvalidBytecode;
+
+    const byte_offset: usize = self.registers[instruction.b];
+    const byte_count = std.math.mul(usize, instruction.components, @sizeOf(u32)) catch return RuntimeError.BufferOutOfBounds;
+    const end = std.math.add(usize, byte_offset, byte_count) catch return RuntimeError.BufferOutOfBounds;
+    if (end > buffer.len)
+        return RuntimeError.BufferOutOfBounds;
+    return buffer[byte_offset..end];
+}
+
+fn resourceBuffer(program: *const Program, resource_buffers: []const ?[]u8, resource_index: u32) RuntimeError![]u8 {
+    const resource = ids.ResourceId.fromIndex(resource_index);
+    _ = program.resourceBinding(resource) orelse return RuntimeError.InvalidResource;
+    if (resource.index() >= resource_buffers.len)
+        return RuntimeError.ResourceNotBound;
+    return resource_buffers[resource.index()] orelse RuntimeError.ResourceNotBound;
+}
+
 fn applyEdge(self: *Self, program: *const Program, edge_index: u32) RuntimeError!u32 {
     if (edge_index >= program.edges.len)
         return RuntimeError.InvalidBytecode;

src/software/interpreter/bytecode.zig

@@ -72,6 +72,8 @@ pub const Opcode = enum(u16) {
     compare_ordered_float_less,
     compare_unordered_float_less,
     select,
+    load_buffer,
+    store_buffer,
     jump_edge,
     branch,
     return_void,

src/software/interpreter/compute.zig

@@ -1,48 +0,0 @@
-const std = @import("std");
-const base = @import("base");
-const shader_ir = @import("shader_ir");
-
-const Shader = @import("Shader.zig");
-
-const VkError = base.VkError;
-const ir = shader_ir.ir;
-
-pub fn dispatch(shader: *Shader, base_group_x: u32, base_group_y: u32, base_group_z: u32, group_count_x: u32, group_count_y: u32, group_count_z: u32) VkError!void {
-    const local_size = shader.workgroup_size orelse return VkError.ValidationFailed;
-    const local_xy = std.math.mul(usize, local_size[0], local_size[1]) catch return VkError.ValidationFailed;
-    const local_count = std.math.mul(usize, local_xy, local_size[2]) catch return VkError.ValidationFailed;
-    if (shader.runtimes.len == 0)
-        return VkError.InvalidPipelineDrv;
-
-    const global_id = findGlobalInvocationId(&shader.program);
-    var runtime = &shader.runtimes[0].runtime;
-    for (0..group_count_z) |group_z| {
-        for (0..group_count_y) |group_y| {
-            for (0..group_count_x) |group_x| {
-                for (0..local_count) |local_index| {
-                    if (global_id) |variable| {
-                        const local_z = local_index / local_xy;
-                        const local_remainder = local_index - local_z * local_xy;
-                        const local_y = local_remainder / local_size[0];
-                        const local_x = local_remainder - local_y * local_size[0];
-                        runtime.writeInput(&shader.program, variable, &.{
-                            (base_group_x + @as(u32, @intCast(group_x))) * local_size[0] + @as(u32, @intCast(local_x)),
-                            (base_group_y + @as(u32, @intCast(group_y))) * local_size[1] + @as(u32, @intCast(local_y)),
-                            (base_group_z + @as(u32, @intCast(group_z))) * local_size[2] + @as(u32, @intCast(local_z)),
-                        }) catch return VkError.Unknown;
-                    }
-                    _ = runtime.run(&shader.program, .{}) catch return VkError.Unknown;
-                }
-            }
-        }
-    }
-}
-
-fn findGlobalInvocationId(program: *const @import("Program.zig")) ?ir.id.InterfaceVariableId {
-    for (program.interfaces, 0..) |optional_binding, index| {
-        const binding = optional_binding orelse continue;
-        if (binding.direction == .input and binding.semantic == .builtin and binding.semantic.builtin == .global_invocation_id)
-            return ir.id.InterfaceVariableId.fromIndex(index);
-    }
-    return null;
-}

src/software/interpreter/test/storage_buffers.zig

@@ -0,0 +1,127 @@
+const std = @import("std");
+const shader_ir = @import("shader_ir");
+
+const Program = @import("../Program.zig");
+const Runtime = @import("../Runtime.zig");
+
+const ir = shader_ir.ir;
+
+const copy_shader =
+    \\ shader compute @main
+    \\ {
+    \\     @source: vec4[u32] = storage_buffer[set(2), binding(3)]
+    \\     @destination: vec4[u32] = storage_buffer[set(4), binding(5)]
+    \\
+    \\     %source_offset: constant u32 = 1
+    \\     %destination_offset: constant u32 = 2
+    \\
+    \\     fn @main() -> void
+    \\     {
+    \\         .entry():
+    \\             %value: vec4[u32] = load_buffer @source, %source_offset
+    \\             store_buffer @destination, %destination_offset, %value
+    \\             return
+    \\     }
+    \\ }
+;
+
+const scalar_shader =
+    \\ shader compute @main
+    \\ {
+    \\     @source: u32 = storage_buffer[set(0), binding(0)]
+    \\     @destination: u32 = storage_buffer[set(0), binding(1)]
+    \\
+    \\     %zero: constant u32 = 0
+    \\     %one: constant u32 = 1
+    \\
+    \\     fn @main() -> void
+    \\     {
+    \\         .entry():
+    \\             %loaded: u32 = load_buffer @source, %zero
+    \\             %value: u32 = integer_add %loaded, %one
+    \\             store_buffer @destination, %zero, %value
+    \\             return
+    \\     }
+    \\ }
+;
+
+const bounds_shader =
+    \\ shader compute @main
+    \\ {
+    \\     @buffer: vec2[u32] = storage_buffer[set(0), binding(0)]
+    \\
+    \\     %offset: constant u32 = 1
+    \\     %first: constant u32 = bits(0x11223344)
+    \\     %second: constant u32 = bits(0x55667788)
+    \\
+    \\     fn @main() -> void
+    \\     {
+    \\         .entry():
+    \\             %value: vec2[u32] = composite_construct %first, %second
+    \\             store_buffer @buffer, %offset, %value
+    \\             return
+    \\     }
+    \\ }
+;
+
+test "[interpreter] storage-buffer vector load and store use portable little-endian words" {
+    var module = try ir.parser.parseString(std.testing.allocator, copy_shader);
+    defer module.deinit();
+
+    var program = try Program.compile(std.testing.allocator, &module);
+    defer program.deinit();
+    var runtime = try Runtime.init(std.testing.allocator, &program);
+    defer runtime.deinit();
+
+    const source_id = ir.id.ResourceId.fromIndex(0);
+    const destination_id = ir.id.ResourceId.fromIndex(1);
+    try std.testing.expectEqual(@as(u32, 2), program.resourceBinding(source_id).?.set);
+    try std.testing.expectEqual(@as(u32, 3), program.resourceBinding(source_id).?.binding);
+    try std.testing.expectEqual(@as(u32, 4), program.resourceBinding(destination_id).?.set);
+    try std.testing.expectEqual(@as(u32, 5), program.resourceBinding(destination_id).?.binding);
+
+    var source = [_]u8{ 0xff, 0x78, 0x56, 0x34, 0x12, 0xef, 0xcd, 0xab, 0x90, 0x04, 0x03, 0x02, 0x01, 0xdd, 0xcc, 0xbb, 0xaa };
+    var destination = [_]u8{0xcc} ** 20;
+    const resources = [_]?[]u8{ source[0..], destination[0..] };
+
+    try std.testing.expectEqual(Runtime.Outcome.returned, try runtime.run(&program, .{ .resource_buffers = &resources }));
+    try std.testing.expectEqualSlices(u8, source[1..17], destination[2..18]);
+    try std.testing.expectEqual(@as(u8, 0xcc), destination[0]);
+    try std.testing.expectEqual(@as(u8, 0xcc), destination[1]);
+    try std.testing.expectEqual(@as(u8, 0xcc), destination[18]);
+    try std.testing.expectEqual(@as(u8, 0xcc), destination[19]);
+}
+
+test "[interpreter] storage-buffer scalar load and store interpret little-endian words" {
+    var module = try ir.parser.parseString(std.testing.allocator, scalar_shader);
+    defer module.deinit();
+
+    var program = try Program.compile(std.testing.allocator, &module);
+    defer program.deinit();
+    var runtime = try Runtime.init(std.testing.allocator, &program);
+    defer runtime.deinit();
+
+    var source = [_]u8{ 0x78, 0x56, 0x34, 0x12 };
+    var destination = [_]u8{0} ** 4;
+    const resources = [_]?[]u8{ source[0..], destination[0..] };
+    try std.testing.expectEqual(Runtime.Outcome.returned, try runtime.run(&program, .{ .resource_buffers = &resources }));
+    try std.testing.expectEqualSlices(u8, &[_]u8{ 0x79, 0x56, 0x34, 0x12 }, &destination);
+}
+
+test "[interpreter] storage-buffer accesses report unbound and out-of-bounds resources" {
+    var module = try ir.parser.parseString(std.testing.allocator, bounds_shader);
+    defer module.deinit();
+
+    var program = try Program.compile(std.testing.allocator, &module);
+    defer program.deinit();
+    var runtime = try Runtime.init(std.testing.allocator, &program);
+    defer runtime.deinit();
+
+    try std.testing.expectError(Runtime.RuntimeError.ResourceNotBound, runtime.run(&program, .{}));
+
+    var buffer = [_]u8{0xa5} ** 8;
+    const resources = [_]?[]u8{buffer[0..]};
+    try std.testing.expectError(Runtime.RuntimeError.BufferOutOfBounds, runtime.run(&program, .{ .resource_buffers = &resources }));
+    const unchanged = [_]u8{0xa5} ** 8;
+    try std.testing.expectEqualSlices(u8, &unchanged, &buffer);
+}

src/software/interpreter/test/test.zig

@@ -10,5 +10,6 @@ comptime {
     _ = @import("arithmetic.zig");
     _ = @import("branching.zig");
     _ = @import("loops.zig");
+    _ = @import("storage_buffers.zig");
     _ = @import("termination.zig");
 }

test/test_runner.zig

@@ -189,7 +189,7 @@ const SlowTracker = struct {
         var slowest = self.slowest;
         const count = slowest.count();
         Printer.fmt("Slowest {d} test{s}: \n", .{ count, if (count != 1) "s" else "" });
-        while (slowest.popMin()) |info| {
+        while (slowest.popMax()) |info| {
             const ms = @as(f64, @floatFromInt(info.ns)) / 1_000_000.0;
             Printer.fmt("  {d:.2}ms\t{s}\n", .{ ms, info.name });
         }