implementing missing functions in FFI

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,
 };
 ```
 

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",

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);

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,
     };
 }
 

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);

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;

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;

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"));

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);