From 589d30168f38e85e993763a59971d64288c70e35 Mon Sep 17 00:00:00 2001 From: Kbz-8 Date: Sat, 4 Jul 2026 20:59:03 +0200 Subject: [PATCH] implementing missing functions in FFI --- Usage.md | 208 ++++++++++++++++++++++++++++++++--------- build.zig.zon | 2 +- ffi/SpirvInterpreter.h | 58 ++++++++++-- ffi/module.zig | 4 + ffi/runtime.zig | 111 ++++++++++++++++++++-- src/Runtime.zig | 45 --------- src/opcodes.zig | 2 +- test/api.zig | 2 - test_c/main.c | 49 +++++++++- 9 files changed, 368 insertions(+), 113 deletions(-) diff --git a/Usage.md b/Usage.md index d0dfc01..7414f55 100644 --- a/Usage.md +++ b/Usage.md @@ -373,6 +373,41 @@ int main(void) } ``` +## Module reflection from C + +`SpvModuleGetReflectionInfos` returns shader-wide metadata collected while parsing the module: + +```c +SpvModuleReflectionInfos infos = SpvModuleGetReflectionInfos(module); + +if (infos.needs_derivatives) +{ + /* Provide derivative data before executing derivative operations. */ +} + +if (infos.has_control_barriers) +{ + /* Use SpvBeginEntryPoint/SpvContinueEntryPoint and synchronize invocations. */ +} + +if (infos.has_atomics) +{ + /* Descriptor or workgroup-backed storage may be modified atomically. */ +} + +if (infos.early_fragment_tests) +{ + /* Fragment shader requested early fragment tests. */ +} +``` + +Compute and geometry execution metadata is also available: + +```c +SpvWord local_x = infos.local_size_x; +SpvWord geometry_output_count = infos.geometry_output_count; +``` + ## Writing inputs from C Write by result id: @@ -426,6 +461,24 @@ SpvWord result = 0; SpvGetResultByLocationComponent(runtime, 0, 1, SPV_LOCATION_OUTPUT, &result); ``` +For built-ins: + +```c +SpvWord workgroup_id_result = 0; + +if (SpvGetBuiltinResult(runtime, SpvBuiltInWorkgroupId, &workgroup_id_result) == SPV_RESULT_SUCCESS) +{ + unsigned int workgroup_id[3] = {0}; + SpvReadBuiltIn(runtime, (SpvByte*)workgroup_id, sizeof(workgroup_id), SpvBuiltInWorkgroupId); +} +``` + +If a result's layout depends on specialization constants or runtime-array sizes, refresh it before querying or reading by result id: + +```c +SpvRefreshResultValueLayout(runtime, result); +``` + ## Push constants from C ```c @@ -485,12 +538,49 @@ SpvAddSpecializationInfo(runtime, entry, (const SpvByte*)&value, sizeof(value)); Add specialization constants before calling the entry point. +If specialization constants affect array sizes, workgroup size, or other layout-dependent metadata, apply the specialization layout after adding the values: + +```c +SpvApplySpecializationLayout(runtime); +``` + +For a fresh invocation layout, use: + +```c +SpvApplySpecializationInvocationLayout(runtime); +``` + To duplicate specialization constants from another runtime: ```c SpvCopySpecializationConstantsFrom(runtime, source_runtime); ``` +## Entry points from C + +Entry points can be looked up by name: + +```c +SpvWord entry = 0; +SpvGetEntryPointByName(runtime, "main", &entry); +``` + +When a module has multiple entry points with the same name, include the execution model: + +```c +SpvGetEntryPointByNameAndExecutionModel( + runtime, + "main", + SpvExecutionModelGLCompute, + &entry); +``` + +Selecting an entry point filters interface lookups such as input and output locations to that entry point: + +```c +SpvSelectEntryPoint(runtime, entry); +``` + ## Derivatives from C ```c @@ -531,24 +621,63 @@ while (status == SPV_ENTRY_POINT_BARRIER) } ``` +## Workgroup data from C + +Compute shaders may declare a static workgroup size. Query it with: + +```c +SpvWorkgroupSize size; + +if (SpvGetWorkgroupSize(runtime, &size) == SPV_RESULT_SUCCESS && size.has_size) +{ + printf("workgroup size = %lu, %lu, %lu\n", size.x, size.y, size.z); +} +``` + +Workgroup storage is allocated separately so multiple runtime invocations can bind the same shared memory during barrier-based execution: + +```c +SpvWorkgroupMemory memory; + +if (SpvCreateWorkgroupMemory(runtime, &memory) != SPV_RESULT_SUCCESS) + return 1; + +SpvSize count = SpvGetWorkgroupMemoryCount(memory); + +for (SpvSize i = 0; i < count; ++i) +{ + SpvWorkgroupMemoryItem item; + + if (SpvGetWorkgroupMemoryItem(memory, i, &item) == SPV_RESULT_SUCCESS) + { + /* item.result is the Workgroup variable result id. + item.data points to item.size bytes owned by memory. */ + } +} + +SpvBindWorkgroupMemory(runtime, memory); + +/* Run one or more invocations that share this workgroup memory. */ + +SpvDestroyWorkgroupMemory(runtime, memory); +``` + ## Image API from C Shaders that execute image operations must provide callbacks in `SpvImageAPI`. ```c static SpvResult ReadImageFloat4( - void* driver_image, - SpvDim dim, - int x, - int y, - int z, + SpvReadImageInfo info, SpvVec4f* dst) { - (void)driver_image; - (void)dim; - (void)x; - (void)y; - (void)z; + (void)info.driver_image; + (void)info.dim; + (void)info.x; + (void)info.y; + (void)info.z; + (void)info.has_lod; + (void)info.lod; dst->x = 0.0f; dst->y = 0.0f; @@ -563,27 +692,19 @@ Sampling callbacks include an explicit-LOD flag/value and an offset: ```c static SpvResult SampleImageFloat4( - void* driver_image, - void* driver_sampler, - SpvDim dim, - float x, - float y, - float z, - SpvBool has_lod, - float lod, - SpvImageOffset offset, + SpvSampleImageInfo info, SpvVec4f* dst) { - (void)driver_image; - (void)driver_sampler; - (void)dim; - (void)has_lod; - (void)lod; - (void)offset; + (void)info.driver_image; + (void)info.driver_sampler; + (void)info.dim; + (void)info.has_lod; + (void)info.lod; + (void)info.offset; - dst->x = x; - dst->y = y; - dst->z = z; + dst->x = (float)info.x; + dst->y = (float)info.y; + dst->z = (float)info.z; dst->w = 1.0f; return SPV_RESULT_SUCCESS; } @@ -593,27 +714,19 @@ Depth-comparison samplers add `dref` and write one `float`: ```c static SpvResult SampleImageDref( - void* driver_image, - void* driver_sampler, - SpvDim dim, - float x, - float y, - float z, + SpvSampleImageInfo info, float dref, - SpvBool has_lod, - float lod, - SpvImageOffset offset, float* dst) { - (void)driver_image; - (void)driver_sampler; - (void)dim; - (void)x; - (void)y; - (void)z; - (void)has_lod; - (void)lod; - (void)offset; + (void)info.driver_image; + (void)info.driver_sampler; + (void)info.dim; + (void)info.x; + (void)info.y; + (void)info.z; + (void)info.has_lod; + (void)info.lod; + (void)info.offset; *dst = dref; return SPV_RESULT_SUCCESS; @@ -632,6 +745,9 @@ SpvImageAPI image_api = { .SpvSampleImageInt4 = SampleImageInt4, .SpvSampleImageDref = SampleImageDref, .SpvQueryImageSize = QueryImageSize, + .SpvQueryImageLevels = QueryImageLevels, + .SpvQueryImageSamples = QueryImageSamples, + .SpvQueryImageLod = QueryImageLod, }; ``` diff --git a/build.zig.zon b/build.zig.zon index a6d2f9d..e1ca61d 100644 --- a/build.zig.zon +++ b/build.zig.zon @@ -1,6 +1,6 @@ .{ .name = .SPIRV_Interpreter, - .version = "0.0.1", + .version = "1.0.0", .dependencies = .{ .zmath = .{ .url = "git+https://github.com/zig-gamedev/zmath.git#3a5955b2b72cd081563fbb084eff05bffd1e3fbb", diff --git a/ffi/SpirvInterpreter.h b/ffi/SpirvInterpreter.h index d8c17b9..01e236c 100644 --- a/ffi/SpirvInterpreter.h +++ b/ffi/SpirvInterpreter.h @@ -413,6 +413,28 @@ typedef unsigned char SpvByte; typedef unsigned long SpvWord; typedef unsigned long SpvSize; +typedef enum +{ + SpvExecutionModelVertex = 0, + SpvExecutionModelTessellationControl = 1, + SpvExecutionModelTessellationEvaluation = 2, + SpvExecutionModelGeometry = 3, + SpvExecutionModelFragment = 4, + SpvExecutionModelGLCompute = 5, + SpvExecutionModelKernel = 6, + SpvExecutionModelTaskNV = 5267, + SpvExecutionModelMeshNV = 5268, + SpvExecutionModelRayGeneration = 5313, + SpvExecutionModelIntersection = 5314, + SpvExecutionModelAnyHit = 5315, + SpvExecutionModelClosestHit = 5316, + SpvExecutionModelMiss = 5317, + SpvExecutionModelCallable = 5318, + SpvExecutionModelTaskEXT = 5364, + SpvExecutionModelMeshEXT = 5365, + SpvExecutionModelMax = 0x7fffffff +} SpvExecutionModel; + typedef enum { SPV_RESULT_SUCCESS = 0, @@ -451,6 +473,8 @@ typedef struct SpvBool needs_derivatives; SpvBool has_control_barriers; + SpvBool has_atomics; + SpvBool early_fragment_tests; } SpvModuleReflectionInfos; typedef struct @@ -472,13 +496,20 @@ typedef enum SPV_ENTRY_POINT_BARRIER = 1 } SpvEntryPointStatus; -typedef enum +typedef struct { - SPV_PRIMITIVE_BOOL = 0, - SPV_PRIMITIVE_FLOAT = 1, - SPV_PRIMITIVE_SINT = 2, - SPV_PRIMITIVE_UINT = 3 -} SpvPrimitiveType; + SpvBool has_size; + SpvWord x; + SpvWord y; + SpvWord z; +} SpvWorkgroupSize; + +typedef struct +{ + SpvWord result; + SpvByte* data; + SpvSize size; +} SpvWorkgroupMemoryItem; typedef struct { @@ -589,6 +620,7 @@ typedef struct typedef void* SpvModule; typedef void* SpvRuntime; +typedef void* SpvWorkgroupMemory; SPV_API SpvResult SpvInitModule(SpvModule* module, const SpvWord* source, SpvSize source_len, SpvModuleOptions options); SPV_API void SpvDeinitModule(SpvModule module); @@ -604,20 +636,31 @@ SPV_API SpvResult SpvFlushDescriptorSets(SpvRuntime runtime); SPV_API SpvResult SpvAddSpecializationInfo(SpvRuntime runtime, SpvRuntimeSpecializationEntry entry, const SpvByte* data, SpvSize data_size); SPV_API SpvResult SpvCopySpecializationConstantsFrom(SpvRuntime runtime, SpvRuntime other); +SPV_API SpvResult SpvApplySpecializationLayout(SpvRuntime runtime); +SPV_API SpvResult SpvApplySpecializationInvocationLayout(SpvRuntime runtime); SPV_API SpvResult SpvSetDerivativeFromMemory(SpvRuntime runtime, SpvWord result, const SpvByte* dx, SpvSize dx_size, const SpvByte* dy, SpvSize dy_size); SPV_API void SpvClearDerivative(SpvRuntime runtime, SpvWord result); SPV_API SpvResult SpvCopyDerivative(SpvRuntime runtime, SpvWord dst, SpvWord src); SPV_API SpvResult SpvPopulatePushConstants(SpvRuntime runtime, const SpvByte* data, SpvSize data_size); +SPV_API SpvResult SpvRefreshResultValueLayout(SpvRuntime runtime, SpvWord result); SPV_API SpvResult SpvGetResultByName(SpvRuntime runtime, const char* name, SpvWord* result); SPV_API SpvResult SpvGetResultLocation(SpvRuntime runtime, SpvWord location, SpvLocationType type, SpvWord* result); SPV_API SpvResult SpvGetResultByLocation(SpvRuntime runtime, SpvWord location, SpvLocationType type, SpvWord* result); SPV_API SpvResult SpvGetResultByLocationComponent(SpvRuntime runtime, SpvWord location, SpvWord component, SpvLocationType type, SpvWord* result); SPV_API SpvResult SpvGetEntryPointByName(SpvRuntime runtime, const char* name, SpvWord* result); +SPV_API SpvResult SpvGetEntryPointByNameAndExecutionModel(SpvRuntime runtime, const char* name, SpvExecutionModel execution_model, SpvWord* result); +SPV_API SpvResult SpvSelectEntryPoint(SpvRuntime runtime, SpvWord entry_point_index); SPV_API SpvResult SpvGetResultMemorySize(SpvRuntime runtime, SpvWord result, SpvSize* size); SPV_API SpvResult SpvGetInputLocationMemorySize(SpvRuntime runtime, SpvWord location, SpvSize* size); -SPV_API SpvResult SpvGetResultPrimitiveType(SpvRuntime runtime, SpvWord result, SpvPrimitiveType* primitive_type); SPV_API SpvBool SpvHasResultDecoration(SpvRuntime runtime, SpvWord result, SpvDecoration decoration); +SPV_API SpvResult SpvGetWorkgroupSize(SpvRuntime runtime, SpvWorkgroupSize* size); + +SPV_API SpvResult SpvCreateWorkgroupMemory(SpvRuntime runtime, SpvWorkgroupMemory* workgroup_memory); +SPV_API void SpvDestroyWorkgroupMemory(SpvRuntime runtime, SpvWorkgroupMemory workgroup_memory); +SPV_API SpvResult SpvBindWorkgroupMemory(SpvRuntime runtime, SpvWorkgroupMemory workgroup_memory); +SPV_API SpvSize SpvGetWorkgroupMemoryCount(SpvWorkgroupMemory workgroup_memory); +SPV_API SpvResult SpvGetWorkgroupMemoryItem(SpvWorkgroupMemory workgroup_memory, SpvSize index, SpvWorkgroupMemoryItem* item); SPV_API SpvResult SpvCallEntryPoint(SpvRuntime runtime, SpvWord entry_point_index); SPV_API SpvResult SpvBeginEntryPoint(SpvRuntime runtime, SpvWord entry_point_index, SpvEntryPointStatus* status); @@ -626,6 +669,7 @@ SPV_API void SpvResetInvocation(SpvRuntime runtime); SPV_API SpvResult SpvReadOutput(SpvRuntime runtime, SpvByte* output, SpvSize output_size, SpvWord result); SPV_API SpvResult SpvReadBuiltIn(SpvRuntime runtime, SpvByte* output, SpvSize output_size, SpvBuiltIn builtin); +SPV_API SpvResult SpvGetBuiltinResult(SpvRuntime runtime, SpvBuiltIn builtin, SpvWord* result); SPV_API SpvResult SpvWriteInput(SpvRuntime runtime, const SpvByte* input, SpvSize input_size, SpvWord result); SPV_API SpvResult SpvWriteInputLocation(SpvRuntime runtime, const SpvByte* input, SpvSize input_size, SpvWord location); diff --git a/ffi/module.zig b/ffi/module.zig index c9510e1..76f3316 100644 --- a/ffi/module.zig +++ b/ffi/module.zig @@ -18,6 +18,8 @@ const ReflectionInfos = extern struct { needs_derivatives: ffi.SpvCBool, has_control_barriers: ffi.SpvCBool, + has_atomics: ffi.SpvCBool, + early_fragment_tests: ffi.SpvCBool, }; fn toCResult(err: spv.Module.ModuleError) ffi.Result { @@ -66,6 +68,8 @@ export fn SpvModuleGetReflectionInfos(module: *spv.Module) callconv(.c) Reflecti .geometry_output = module.reflection_infos.geometry_output, .needs_derivatives = if (module.reflection_infos.needs_derivatives) 1 else 0, .has_control_barriers = if (module.reflection_infos.has_control_barriers) 1 else 0, + .has_atomics = if (module.reflection_infos.has_atomics) 1 else 0, + .early_fragment_tests = if (module.reflection_infos.early_fragment_tests) 1 else 0, }; } diff --git a/ffi/runtime.zig b/ffi/runtime.zig index 9b46e0d..90b23c9 100644 --- a/ffi/runtime.zig +++ b/ffi/runtime.zig @@ -18,6 +18,19 @@ const EntryPointStatus = enum(c_int) { barrier = 1, }; +const WorkgroupSize = extern struct { + has_size: ffi.SpvCBool, + x: ffi.SpvCWord, + y: ffi.SpvCWord, + z: ffi.SpvCWord, +}; + +const WorkgroupMemoryItem = extern struct { + result: ffi.SpvCWord, + data: ?[*]u8, + size: ffi.SpvCSize, +}; + const PrimitiveType = enum(c_int) { bool = 0, float = 1, @@ -392,6 +405,10 @@ const RuntimeWrapper = struct { image_api: ImageAPI, }; +const WorkgroupMemoryWrapper = struct { + memories: []spv.Runtime.WorkgroupMemory, +}; + export fn SpvInitRuntime(rt: **RuntimeWrapper, module: *spv.Module, image_api: ImageAPI) callconv(.c) ffi.Result { const allocator = std.heap.c_allocator; @@ -482,6 +499,18 @@ export fn SpvCopySpecializationConstantsFrom(rt: *RuntimeWrapper, other: *const return .Success; } +export fn SpvApplySpecializationLayout(rt: *RuntimeWrapper) callconv(.c) ffi.Result { + const allocator = std.heap.c_allocator; + rt.rt.applySpecializationLayout(allocator) catch |err| return toCResult(err); + return .Success; +} + +export fn SpvApplySpecializationInvocationLayout(rt: *RuntimeWrapper) callconv(.c) ffi.Result { + const allocator = std.heap.c_allocator; + rt.rt.applySpecializationInvocationLayout(allocator) catch |err| return toCResult(err); + return .Success; +} + export fn SpvSetDerivativeFromMemory(rt: *RuntimeWrapper, result: spv.SpvWord, dx: [*]const u8, dx_size: c_ulong, dy: [*]const u8, dy_size: c_ulong) callconv(.c) ffi.Result { const allocator = std.heap.c_allocator; rt.rt.setDerivativeFromMemory(allocator, result, dx[0..dx_size], dy[0..dy_size]) catch |err| return toCResult(err); @@ -504,11 +533,26 @@ export fn SpvPopulatePushConstants(rt: *RuntimeWrapper, data: [*]const u8, data_ return .Success; } +export fn SpvRefreshResultValueLayout(rt: *RuntimeWrapper, result: spv.SpvWord) callconv(.c) ffi.Result { + rt.rt.refreshResultValueLayout(result) catch |err| return toCResult(err); + return .Success; +} + export fn SpvGetEntryPointByName(rt: *RuntimeWrapper, name: [*:0]const u8, result: *spv.SpvWord) callconv(.c) ffi.Result { result.* = rt.rt.getEntryPointByName(std.mem.span(name)) catch |err| return toCResult(err); return .Success; } +export fn SpvGetEntryPointByNameAndExecutionModel(rt: *RuntimeWrapper, name: [*:0]const u8, execution_model: spv.spv.SpvExecutionModel, result: *spv.SpvWord) callconv(.c) ffi.Result { + result.* = rt.rt.getEntryPointByNameAndExecutionModel(std.mem.span(name), execution_model) catch |err| return toCResult(err); + return .Success; +} + +export fn SpvSelectEntryPoint(rt: *RuntimeWrapper, entry_point: spv.SpvWord) callconv(.c) ffi.Result { + rt.rt.selectEntryPoint(entry_point) catch |err| return toCResult(err); + return .Success; +} + export fn SpvGetResultByLocation(rt: *RuntimeWrapper, location: spv.SpvWord, kind: LocationType, result: *spv.SpvWord) callconv(.c) ffi.Result { result.* = rt.rt.getResultByLocation(location, switch (kind) { .input => .input, @@ -544,18 +588,64 @@ export fn SpvGetInputLocationMemorySize(rt: *RuntimeWrapper, location: spv.SpvWo return .Success; } -export fn SpvGetResultPrimitiveType(rt: *RuntimeWrapper, result: spv.SpvWord, primitive_type: *PrimitiveType) callconv(.c) ffi.Result { - primitive_type.* = switch (rt.rt.getResultPrimitiveType(result) catch |err| return toCResult(err)) { - .Bool => .bool, - .Float => .float, - .SInt => .sint, - .UInt => .uint, +export fn SpvHasResultDecoration(rt: *RuntimeWrapper, result: spv.SpvWord, decoration: spv.spv.SpvDecoration) callconv(.c) c_int { + return if (rt.rt.hasResultDecoration(result, decoration)) 1 else 0; +} + +export fn SpvGetWorkgroupSize(rt: *RuntimeWrapper, size: *WorkgroupSize) callconv(.c) ffi.Result { + const allocator = std.heap.c_allocator; + if (rt.rt.getWorkgroupSize(allocator) catch |err| return toCResult(err)) |workgroup_size| { + size.* = .{ + .has_size = 1, + .x = workgroup_size[0], + .y = workgroup_size[1], + .z = workgroup_size[2], + }; + } else { + size.* = .{ + .has_size = 0, + .x = 0, + .y = 0, + .z = 0, + }; + } + return .Success; +} + +export fn SpvCreateWorkgroupMemory(rt: *RuntimeWrapper, workgroup_memory: **WorkgroupMemoryWrapper) callconv(.c) ffi.Result { + const allocator = std.heap.c_allocator; + workgroup_memory.* = allocator.create(WorkgroupMemoryWrapper) catch return .OutOfMemory; + workgroup_memory.*.memories = rt.rt.createWorkgroupMemory(allocator) catch |err| { + allocator.destroy(workgroup_memory.*); + return toCResult(err); }; return .Success; } -export fn SpvHasResultDecoration(rt: *RuntimeWrapper, result: spv.SpvWord, decoration: spv.spv.SpvDecoration) callconv(.c) c_int { - return if (rt.rt.hasResultDecoration(result, decoration)) 1 else 0; +export fn SpvDestroyWorkgroupMemory(rt: *RuntimeWrapper, workgroup_memory: *WorkgroupMemoryWrapper) callconv(.c) void { + const allocator = std.heap.c_allocator; + rt.rt.destroyWorkgroupMemory(allocator, workgroup_memory.memories); + allocator.destroy(workgroup_memory); +} + +export fn SpvBindWorkgroupMemory(rt: *RuntimeWrapper, workgroup_memory: *const WorkgroupMemoryWrapper) callconv(.c) ffi.Result { + rt.rt.bindWorkgroupMemory(workgroup_memory.memories) catch |err| return toCResult(err); + return .Success; +} + +export fn SpvGetWorkgroupMemoryCount(workgroup_memory: *const WorkgroupMemoryWrapper) callconv(.c) ffi.SpvCSize { + return workgroup_memory.memories.len; +} + +export fn SpvGetWorkgroupMemoryItem(workgroup_memory: *const WorkgroupMemoryWrapper, index: ffi.SpvCSize, item: *WorkgroupMemoryItem) callconv(.c) ffi.Result { + if (index >= workgroup_memory.memories.len) return .OutOfBounds; + const memory = workgroup_memory.memories[index]; + item.* = .{ + .result = memory.result, + .data = memory.bytes.ptr, + .size = memory.bytes.len, + }; + return .Success; } export fn SpvCallEntryPoint(rt: *RuntimeWrapper, entry_point: spv.SpvWord) callconv(.c) ffi.Result { @@ -607,6 +697,11 @@ export fn SpvReadBuiltIn(rt: *RuntimeWrapper, output: [*]u8, output_size: c_ulon return .Success; } +export fn SpvGetBuiltinResult(rt: *RuntimeWrapper, builtin: spv.spv.SpvBuiltIn, result: *spv.SpvWord) callconv(.c) ffi.Result { + result.* = rt.rt.getBuiltinResult(builtin) orelse return .NotFound; + return .Success; +} + export fn SpvWriteInput(rt: *RuntimeWrapper, input: [*]const u8, input_size: c_ulong, result: spv.SpvWord) callconv(.c) ffi.Result { const allocator = std.heap.c_allocator; rt.rt.writeInput(allocator, input[0..input_size], result) catch |err| return toCResult(err); diff --git a/src/Runtime.zig b/src/Runtime.zig index 5b0d25b..d655075 100644 --- a/src/Runtime.zig +++ b/src/Runtime.zig @@ -465,10 +465,6 @@ pub fn snapshotPhiValuesForBranch(self: *Self, allocator: std.mem.Allocator, tar } } -pub fn getPhiValueSnapshot(self: *Self, id: SpvWord) ?*const Value { - return self.phi_values.getPtr(id); -} - pub fn setDerivative(self: *Self, allocator: std.mem.Allocator, result: SpvWord, dx: *const Value, dy: *const Value) RuntimeError!void { const derivative: Derivative = .{ .dx = try dx.dupe(allocator), @@ -694,12 +690,6 @@ fn resultComponent(self: *const Self, result: SpvWord) SpvWord { return 0; } -pub fn getResultPrimitiveType(self: *const Self, result: SpvWord) RuntimeError!PrimitiveType { - if (result >= self.results.len) - return RuntimeError.OutOfBounds; - return (try self.results[result].getConstValue()).resolvePrimitiveType(); -} - pub fn getWorkgroupSize(self: *Self, allocator: std.mem.Allocator) RuntimeError!?@Vector(3, u32) { try self.pass(allocator, .initMany(&.{ .SpecConstantTrue, @@ -1275,41 +1265,6 @@ pub fn hasResultDecoration(self: *const Self, result: SpvWord, decoration: spv.S return false; } -pub fn hasResultOrMemberDecoration(self: *const Self, result: SpvWord, decoration: spv.SpvDecoration) bool { - if (self.hasResultDecoration(result, decoration)) - return true; - - if (result >= self.results.len) - return false; - - const type_word = switch ((self.results[result].variant orelse return false)) { - .Variable => |variable| variable.type_word, - else => return false, - }; - const target_type_word = switch ((self.results[type_word].variant orelse return false)) { - .Type => |t| switch (t) { - .Pointer => |ptr| ptr.target, - else => type_word, - }, - else => return false, - }; - const target_type = self.results[target_type_word].variant orelse return false; - switch (target_type) { - .Type => |t| switch (t) { - .Structure => { - for (self.results[target_type_word].decorations.items) |member_decoration| { - if (member_decoration.rtype == decoration) - return true; - } - }, - else => {}, - }, - else => {}, - } - - return false; -} - pub fn resetInvocation(self: *Self, allocator: std.mem.Allocator) void { self.derivative_sign_x = 1.0; self.derivative_sign_y = 1.0; diff --git a/src/opcodes.zig b/src/opcodes.zig index 27fc5ee..2dff77a 100644 --- a/src/opcodes.zig +++ b/src/opcodes.zig @@ -6201,7 +6201,7 @@ fn opPhi(allocator: std.mem.Allocator, word_count: SpvWord, rt: *Runtime) Runtim const parent_label_id = try rt.it.next(); if (parent_label_id == predecessor) { - const value = rt.getPhiValueSnapshot(value_id) orelse try rt.results[value_id].getValue(); + const value = rt.phi_values.getPtr(value_id) orelse try rt.results[value_id].getValue(); try copyValue(try rt.results[id].getValue(), value); try rt.copyDerivative(allocator, id, value_id); return; diff --git a/test/api.zig b/test/api.zig index de29d8c..70bdaa8 100644 --- a/test/api.zig +++ b/test/api.zig @@ -141,7 +141,6 @@ test "Runtime API lifecycle" { try std.testing.expectEqual(@as(usize, 16), try rt.getInputLocationMemorySize(0)); try std.testing.expectEqual(@as(usize, 16), try rt.getResultMemorySize(output_result)); try std.testing.expectEqual(@as(usize, 16), try rt.getResultMemorySize(output_location_result)); - try std.testing.expectEqual(.Float, try rt.getResultPrimitiveType(output_result)); try std.testing.expect(rt.hasResultDecoration(output_result, .Location)); const input = [_]f32{ 10.0, 20.0, 30.0, 40.0 }; @@ -319,7 +318,6 @@ test "Integer output metadata" { defer rt.deinit(allocator); const output_result = try rt.getResultByLocation(0, .output); - try std.testing.expectEqual(.UInt, try rt.getResultPrimitiveType(output_result)); try std.testing.expectEqual(@as(usize, 4), try rt.getResultMemorySize(output_result)); try rt.callEntryPoint(allocator, try rt.getEntryPointByName("main")); diff --git a/test_c/main.c b/test_c/main.c index 53c8594..0506f24 100644 --- a/test_c/main.c +++ b/test_c/main.c @@ -146,6 +146,12 @@ int main(void) fprintf(stderr, "Unexpected binding lookup result\n"); return -1; } + SpvModuleReflectionInfos reflection_infos = SpvModuleGetReflectionInfos(module); + if (reflection_infos.has_atomics || reflection_infos.early_fragment_tests) + { + fprintf(stderr, "Unexpected reflection flags\n"); + return -1; + } SpvImageAPI image_api = { .SpvReadImageFloat4 = ReadImageFloat4, @@ -187,18 +193,55 @@ int main(void) SpvWord main_entry_index; CHECK_RESULT(SpvGetEntryPointByName(runtime, "main", &main_entry_index)); + SpvWord fragment_entry_index; + CHECK_RESULT(SpvGetEntryPointByNameAndExecutionModel(runtime, "main", SpvExecutionModelFragment, &fragment_entry_index)); + if (fragment_entry_index != main_entry_index) + { + fprintf(stderr, "Unexpected execution-model entry point lookup\n"); + return -1; + } + CHECK_RESULT(SpvSelectEntryPoint(runtime, main_entry_index)); + CHECK_RESULT(SpvApplySpecializationLayout(runtime)); + CHECK_RESULT(SpvApplySpecializationInvocationLayout(second_runtime)); + SpvWorkgroupSize workgroup_size; + CHECK_RESULT(SpvGetWorkgroupSize(runtime, &workgroup_size)); + if (workgroup_size.has_size) + { + fprintf(stderr, "Unexpected workgroup size\n"); + return -1; + } + SpvWorkgroupMemory workgroup_memory; + CHECK_RESULT(SpvCreateWorkgroupMemory(runtime, &workgroup_memory)); + if (SpvGetWorkgroupMemoryCount(workgroup_memory) != 0) + { + fprintf(stderr, "Unexpected workgroup memory count\n"); + return -1; + } + SpvWorkgroupMemoryItem workgroup_memory_item; + if (SpvGetWorkgroupMemoryItem(workgroup_memory, 0, &workgroup_memory_item) != SPV_RESULT_OUT_OF_BOUNDS) + { + fprintf(stderr, "Unexpected workgroup memory lookup result\n"); + return -1; + } + CHECK_RESULT(SpvBindWorkgroupMemory(runtime, workgroup_memory)); + SpvDestroyWorkgroupMemory(runtime, workgroup_memory); CHECK_RESULT(SpvCallEntryPoint(runtime, main_entry_index)); float output[4]; SpvWord output_result; CHECK_RESULT(SpvGetResultByName(runtime, "color", &output_result)); + CHECK_RESULT(SpvRefreshResultValueLayout(runtime, output_result)); CHECK_RESULT(SpvReadOutput(runtime, (SpvByte*)output, sizeof(output), output_result)); + SpvWord builtin_result; + if (SpvGetBuiltinResult(runtime, SpvBuiltInWorkgroupId, &builtin_result) != SPV_RESULT_NOT_FOUND) + { + fprintf(stderr, "Unexpected builtin lookup result\n"); + return -1; + } SpvSize output_size = 0; - SpvPrimitiveType primitive_type; CHECK_RESULT(SpvGetResultMemorySize(runtime, output_result, &output_size)); - CHECK_RESULT(SpvGetResultPrimitiveType(runtime, output_result, &primitive_type)); - if (output_size != sizeof(output) || primitive_type != SPV_PRIMITIVE_FLOAT || primitive_type == SPV_PRIMITIVE_SINT || primitive_type == SPV_PRIMITIVE_UINT) + if (output_size != sizeof(output)) { fprintf(stderr, "Unexpected output metadata\n"); SpvDeinitRuntime(second_runtime);