From a0d6fa487e0d3b07bb505a11993b720a6943493f Mon Sep 17 00:00:00 2001 From: Kbz-8 Date: Wed, 12 Aug 2026 19:26:02 +0200 Subject: [PATCH] [IR] adding external resources managements [Soft] adding descriptor sets management --- build.zig | 6 +- src/compiler/README.md | 52 +- src/compiler/ir/Builder.zig | 10 + src/compiler/ir/instruction.zig | 30 +- src/compiler/ir/parser/ast.zig | 12 + src/compiler/ir/parser/lower.zig | 38 +- src/compiler/ir/parser/parser.zig | 171 +++- src/compiler/ir/printer.zig | 26 + .../ir/transformers/inline_all_functions.zig | 9 + src/compiler/ir/validator/validator.zig | 150 +++ src/compiler/spirv/spirv.zig | 876 +++++++++++++++++- src/compiler/spirv/translator.zig | 553 ++++++++++- src/intel/compiler/lower/lower.zig | 15 + src/software/SoftDescriptorSet.zig | 4 +- src/software/device/ComputeDispatcher.zig | 351 ------- src/software/device/Device.zig | 44 +- .../device/compute/ComputeDispatcher.zig | 177 ++++ .../device/compute/ir_interpreter.zig | 96 ++ .../device/compute/spirv_interpreter.zig | 227 +++++ src/software/interpreter/Program.zig | 55 ++ src/software/interpreter/Runtime.zig | 52 ++ src/software/interpreter/bytecode.zig | 2 + src/software/interpreter/compute.zig | 48 - .../interpreter/test/storage_buffers.zig | 127 +++ src/software/interpreter/test/test.zig | 1 + test/test_runner.zig | 2 +- 26 files changed, 2664 insertions(+), 470 deletions(-) delete mode 100644 src/software/device/ComputeDispatcher.zig create mode 100644 src/software/device/compute/ComputeDispatcher.zig create mode 100644 src/software/device/compute/ir_interpreter.zig create mode 100644 src/software/device/compute/spirv_interpreter.zig delete mode 100644 src/software/interpreter/compute.zig create mode 100644 src/software/interpreter/test/storage_buffers.zig diff --git a/build.zig b/build.zig index 1ddfd86..4e849cf 100644 --- a/build.zig +++ b/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); - (try addMultithreadedCTS(b, target, &impl, lib, mode)).dependOn(&lib_install.step); + if (addCTS(b, target, &impl, lib, mode) catch null) |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(.{ diff --git a/src/compiler/README.md b/src/compiler/README.md index 459a425..6906138 100644 --- a/src/compiler/README.md +++ b/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 @ { + fn @() -> @@ -67,8 +68,8 @@ shader @ } ``` -Execution modes, resources, source locations, and structured-control metadata -exist in memory, but the printer does not display them yet. +Execution modes, source locations, and structured-control metadata exist in +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 -treated as side-effect free by the rewriter. A block's terminator is stored +source location. Except for `store_interface`, `store_buffer`, and `call`, current +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 diff --git a/src/compiler/ir/Builder.zig b/src/compiler/ir/Builder.zig index 748dde2..84781e1 100644 --- a/src/compiler/ir/Builder.zig +++ b/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, diff --git a/src/compiler/ir/instruction.zig b/src/compiler/ir/instruction.zig index 53cdb6e..c7091e9 100644 --- a/src/compiler/ir/instruction.zig +++ b/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, }; } diff --git a/src/compiler/ir/parser/ast.zig b/src/compiler/ir/parser/ast.zig index d896ebc..f18446f 100644 --- a/src/compiler/ir/parser/ast.zig +++ b/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 }, }; diff --git a/src/compiler/ir/parser/lower.zig b/src/compiler/ir/parser/lower.zig index 2ee4f34..537af13 100644 --- a/src/compiler/ir/parser/lower.zig +++ b/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; diff --git a/src/compiler/ir/parser/parser.zig b/src/compiler/ir/parser/parser.zig index 742fe07..8eff048 100644 --- a/src/compiler/ir/parser/parser.zig +++ b/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,64 +48,91 @@ 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 direction = std.meta.stringToEnum(module_ir.InterfaceDirection, direction_token.text) orelse return Error.InvalidSemantic; + const kind_token = try self.expect(.identifier); try self.expectDiscard(.left_square); - 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); - const location = try self.parseUnsigned(u32, .number); - try self.expectDiscard(.right_paren); + 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); + const location = try self.parseUnsigned(u32, .number); + try self.expectDiscard(.right_paren); - try self.expectDiscard(.comma); - try self.expectIdentifier("component"); - try self.expectDiscard(.left_paren); - const component = try self.parseUnsigned(u8, .number); - try self.expectDiscard(.right_paren); + try self.expectDiscard(.comma); + try self.expectIdentifier("component"); + try self.expectDiscard(.left_paren); + const component = try self.parseUnsigned(u8, .number); + try self.expectDiscard(.right_paren); - try self.expectDiscard(.comma); - try self.expectIdentifier("index"); - try self.expectDiscard(.left_paren); - const index = try self.parseUnsigned(u8, .number); - try self.expectDiscard(.right_paren); + try self.expectDiscard(.comma); + try self.expectIdentifier("index"); + try self.expectDiscard(.left_paren); + const index = try self.parseUnsigned(u8, .number); + try self.expectDiscard(.right_paren); - break :blk .{ - .location = .{ - .location = location, - .component = component, - .index = index, - }, - }; - } else if (std.mem.eql(u8, semantic_name, "builtin")) blk: { - try self.expectDiscard(.left_paren); - const builtin_name = (try self.expect(.identifier)).text; - try self.expectDiscard(.right_paren); + break :blk .{ + .location = .{ + .location = location, + .component = component, + .index = index, + }, + }; + } else if (std.mem.eql(u8, semantic_name, "builtin")) blk: { + try self.expectDiscard(.left_paren); + const builtin_name = (try self.expect(.identifier)).text; + try self.expectDiscard(.right_paren); - const builtin = std.meta.stringToEnum(module_ir.Builtin, builtin_name) orelse return Error.InvalidSemantic; - break :blk .{ .builtin = builtin }; - } else return Error.InvalidSemantic; + const builtin = std.meta.stringToEnum(module_ir.Builtin, builtin_name) orelse return Error.InvalidSemantic; + break :blk .{ .builtin = builtin }; + } else return Error.InvalidSemantic; + try self.expectDiscard(.right_square); + + 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 .{ - .direction = direction, + return .{ .resource = .{ + .kind = kind, .name = name, .ty = ty, - .semantic = semantic, - }; + .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"); diff --git a/src/compiler/ir/printer.zig b/src/compiler/ir/printer.zig index 2350ca7..feccf18 100644 --- a/src/compiler/ir/printer.zig +++ b/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); diff --git a/src/compiler/ir/transformers/inline_all_functions.zig b/src/compiler/ir/transformers/inline_all_functions.zig index c8b1086..122b9d7 100644 --- a/src/compiler/ir/transformers/inline_all_functions.zig +++ b/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, }; } diff --git a/src/compiler/ir/validator/validator.zig b/src/compiler/ir/validator/validator.zig index 91b5e89..1d7de4e 100644 --- a/src/compiler/ir/validator/validator.zig +++ b/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 diff --git a/src/compiler/spirv/spirv.zig b/src/compiler/spirv/spirv.zig index 606c836..645333b 100644 --- a/src/compiler/spirv/spirv.zig +++ b/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, - shift_right_arithmetic = 195, - shift_left_logical = 196, - bitwise_or = 197, - bitwise_xor = 198, - bitwise_and = 199, + vector_times_scalar = 142, + matrix_times_scalar = 143, + vector_times_matrix = 144, + matrix_times_vector = 145, + matrix_times_matrix = 146, + 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, _, }; diff --git a/src/compiler/spirv/translator.zig b/src/compiler/spirv/translator.zig index b7ffc1e..56d4c5a 100644 --- a/src/compiler/spirv/translator.zig +++ b/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]), .{ - .load_interface = .{ .variable = try context.interfaceVariable(operands[2]) }, - }, context.nameOf(operands[1]))).?; - - try context.setValue(operands[1], result); + 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; - _ = try context.builder.appendInstruction(block, null, .{ - .store_interface = .{ - .variable = try context.interfaceVariable(operands[0]), - .value = try context.resolveValue(operands[1]), - }, - }, null); + 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 = 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,10 +1145,135 @@ 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 { var active = false; var current_label: ?u32 = null; @@ -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 diff --git a/src/intel/compiler/lower/lower.zig b/src/intel/compiler/lower/lower.zig index 0bf5424..23d23d9 100644 --- a/src/intel/compiler/lower/lower.zig +++ b/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" { diff --git a/src/software/SoftDescriptorSet.zig b/src/software/SoftDescriptorSet.zig index 91e96f5..4329720 100644 --- a/src/software/SoftDescriptorSet.zig +++ b/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; } } } diff --git a/src/software/device/ComputeDispatcher.zig b/src/software/device/ComputeDispatcher.zig deleted file mode 100644 index 9b5f431..0000000 --- a/src/software/device/ComputeDispatcher.zig +++ /dev/null @@ -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, - }; -} diff --git a/src/software/device/Device.zig b/src/software/device/Device.zig index 759116f..34e5825 100644 --- a/src/software/device/Device.zig +++ b/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, diff --git a/src/software/device/compute/ComputeDispatcher.zig b/src/software/device/compute/ComputeDispatcher.zig new file mode 100644 index 0000000..7d4236d --- /dev/null +++ b/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); + } +} diff --git a/src/software/device/compute/ir_interpreter.zig b/src/software/device/compute/ir_interpreter.zig new file mode 100644 index 0000000..2bc41f2 --- /dev/null +++ b/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; +} diff --git a/src/software/device/compute/spirv_interpreter.zig b/src/software/device/compute/spirv_interpreter.zig new file mode 100644 index 0000000..fae6e29 --- /dev/null +++ b/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, + }; +} diff --git a/src/software/interpreter/Program.zig b/src/software/interpreter/Program.zig index 7924cf9..2443f21 100644 --- a/src/software/interpreter/Program.zig +++ b/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)), diff --git a/src/software/interpreter/Runtime.zig b/src/software/interpreter/Runtime.zig index 84701be..11284b9 100644 --- a/src/software/interpreter/Runtime.zig +++ b/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; diff --git a/src/software/interpreter/bytecode.zig b/src/software/interpreter/bytecode.zig index cb48e3f..8d59754 100644 --- a/src/software/interpreter/bytecode.zig +++ b/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, diff --git a/src/software/interpreter/compute.zig b/src/software/interpreter/compute.zig deleted file mode 100644 index 070b8cb..0000000 --- a/src/software/interpreter/compute.zig +++ /dev/null @@ -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; -} diff --git a/src/software/interpreter/test/storage_buffers.zig b/src/software/interpreter/test/storage_buffers.zig new file mode 100644 index 0000000..f10a5b0 --- /dev/null +++ b/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); +} diff --git a/src/software/interpreter/test/test.zig b/src/software/interpreter/test/test.zig index e8eaaa2..ac2ab18 100644 --- a/src/software/interpreter/test/test.zig +++ b/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"); } diff --git a/test/test_runner.zig b/test/test_runner.zig index d1be8c9..e34e80f 100644 --- a/test/test_runner.zig +++ b/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 }); }