whoops

.gitea/workflows/build.yml

@@ -19,13 +19,13 @@ jobs:
     - uses: https://codeberg.org/mlugg/setup-zig@v2
 
     - name: Building
-      run: zig build -Dno-example=true
+      run: zig build
 
     - name: Building FFI
-      run: zig build ffi-c -Dno-example=true
+      run: zig build ffi-c
 
     - name: Generating docs
-      run: zig build docs -Dno-example=true
+      run: zig build docs
 
     - name: Deploying docs
       if: ${{ github.event_name == 'push' && github.ref == 'refs/heads/master' }}

.gitea/workflows/test.yml

@@ -20,4 +20,4 @@ jobs:
     - uses: https://codeberg.org/mlugg/setup-zig@v2
 
     - name: Test
-      run: zig build test -Dno-example=true --release=fast
+      run: zig build test --release=fast

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,
+    SpvReadImageInfo info,
-    SpvDim dim,
-    int x,
-    int y,
-    int z,
     SpvVec4f* dst)
 {
-    (void)driver_image;
+    (void)info.driver_image;
-    (void)dim;
+    (void)info.dim;
-    (void)x;
+    (void)info.x;
-    (void)y;
+    (void)info.y;
-    (void)z;
+    (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,
+    SpvSampleImageInfo info,
-    void* driver_sampler,
-    SpvDim dim,
-    float x,
-    float y,
-    float z,
-    SpvBool has_lod,
-    float lod,
-    SpvImageOffset offset,
     SpvVec4f* dst)
 {
-    (void)driver_image;
+    (void)info.driver_image;
-    (void)driver_sampler;
+    (void)info.driver_sampler;
-    (void)dim;
+    (void)info.dim;
-    (void)has_lod;
+    (void)info.has_lod;
-    (void)lod;
+    (void)info.lod;
-    (void)offset;
+    (void)info.offset;
 
-    dst->x = x;
+    dst->x = (float)info.x;
-    dst->y = y;
+    dst->y = (float)info.y;
-    dst->z = z;
+    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,
+    SpvSampleImageInfo info,
-    void* driver_sampler,
-    SpvDim dim,
-    float x,
-    float y,
-    float z,
     float dref,
-    SpvBool has_lod,
-    float lod,
-    SpvImageOffset offset,
     float* dst)
 {
-    (void)driver_image;
+    (void)info.driver_image;
-    (void)driver_sampler;
+    (void)info.driver_sampler;
-    (void)dim;
+    (void)info.dim;
-    (void)x;
+    (void)info.x;
-    (void)y;
+    (void)info.y;
-    (void)z;
+    (void)info.z;
-    (void)has_lod;
+    (void)info.has_lod;
-    (void)lod;
+    (void)info.lod;
-    (void)offset;
+    (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

@@ -27,121 +27,13 @@ pub fn build(b: *std.Build) void {
         .use_llvm = use_llvm,
     });
 
-    const install_spv_lib = b.addInstallArtifact(spv_lib, .{});
+    b.installArtifact(spv_lib);
 
-    addSandbox(b, target, optimize, use_llvm, spv_mod, &install_spv_lib.step);
-    addExample(b, target, optimize, use_llvm, spv_mod, &install_spv_lib.step);
     addZigTests(b, target, optimize, use_llvm, spv_mod, zmath);
     addCffi(b, target, optimize, use_llvm, spv_mod);
     addDocs(b, spv_mod);
 }
 
-fn addExample(
-    b: *std.Build,
-    target: std.Build.ResolvedTarget,
-    optimize: std.builtin.OptimizeMode,
-    use_llvm: bool,
-    spv_mod: *std.Build.Module,
-    install_spv_lib_step: *std.Build.Step,
-) void {
-    const no_example = b.option(bool, "no-example", "Skip example build") orelse false;
-    if (!no_example) {
-        const sdl3 = b.lazyDependency("sdl3", .{
-            .target = target,
-            .optimize = optimize,
-        }) orelse return;
-
-        const exe = b.addExecutable(.{
-            .name = "spirv_interpreter_example",
-            .root_module = b.createModule(.{
-                .root_source_file = b.path("example/main.zig"),
-                .target = target,
-                .optimize = optimize,
-                .imports = &.{
-                    .{ .name = "spv", .module = spv_mod },
-                    .{ .name = "sdl3", .module = sdl3.module("sdl3") },
-                },
-            }),
-            .use_llvm = use_llvm,
-        });
-
-        const install_exe = b.addInstallArtifact(exe, .{});
-        install_exe.step.dependOn(install_spv_lib_step);
-
-        const run_exe = b.addRunArtifact(exe);
-        run_exe.step.dependOn(&install_exe.step);
-
-        const run_step = b.step("example", "Run the example");
-        run_step.dependOn(&run_exe.step);
-
-        addShaderCompileStep(
-            b,
-            "example-shader",
-            "Compile example shader using nzslc",
-            "example/shader.nzsl",
-            "example",
-        );
-    }
-}
-
-fn addSandbox(
-    b: *std.Build,
-    target: std.Build.ResolvedTarget,
-    optimize: std.builtin.OptimizeMode,
-    use_llvm: bool,
-    spv_mod: *std.Build.Module,
-    install_spv_lib_step: *std.Build.Step,
-) void {
-    const exe = b.addExecutable(.{
-        .name = "spirv_interpreter_sandbox",
-        .root_module = b.createModule(.{
-            .root_source_file = b.path("sandbox/main.zig"),
-            .target = target,
-            .optimize = optimize,
-            .imports = &.{
-                .{ .name = "spv", .module = spv_mod },
-            },
-        }),
-        .use_llvm = use_llvm,
-    });
-
-    const install_exe = b.addInstallArtifact(exe, .{});
-    install_exe.step.dependOn(install_spv_lib_step);
-
-    const run_exe = b.addRunArtifact(exe);
-    run_exe.step.dependOn(&install_exe.step);
-
-    const run_step = b.step("sandbox", "Run the sandbox");
-    run_step.dependOn(&run_exe.step);
-
-    addShaderCompileStep(
-        b,
-        "sandbox-shader",
-        "Compile sandbox shader using nzslc",
-        "sandbox/shader.nzsl",
-        "sandbox",
-    );
-}
-
-fn addShaderCompileStep(
-    b: *std.Build,
-    step_name: []const u8,
-    description: []const u8,
-    shader_path: []const u8,
-    output_dir: []const u8,
-) void {
-    const cmd = b.addSystemCommand(&.{
-        "nzslc",
-        shader_path,
-        "--compile=spv,spv-dis",
-        "-o",
-        output_dir,
-    });
-
-    const step = b.step(step_name, description);
-    step.dependOn(&cmd.step);
-}
-
 fn addZigTests(
     b: *std.Build,
     target: std.Build.ResolvedTarget,

build.zig.zon

@@ -1,6 +1,10 @@
 .{
     .name = .SPIRV_Interpreter,
-    .version = "0.0.1",
+    .version = "1.0.0",
+
+    .fingerprint = 0xd87ec8ab9fa9396a,
+    .minimum_zig_version = "0.16.0",
+
     .dependencies = .{
         .zmath = .{
             .url = "git+https://github.com/zig-gamedev/zmath.git#3a5955b2b72cd081563fbb084eff05bffd1e3fbb",
@@ -11,22 +15,15 @@
             .hash = "NZSL-1.1.5-N0xSVHV7AAB168H2nqMDYY0hjUXtXJ9_eQGNHasSr9r8",
             .lazy = true,
         },
-        .sdl3 = .{
-            .url = "git+https://codeberg.org/7Games/zig-sdl3?ref=v0.2.0#40c2e4b579aa556db37a502c936426aa1c8b5c95",
-            .hash = "sdl3-0.2.0-NmT1Q0mFJwBi9kZmArzh2rfJ_mFshydV0zPGULVlpACc",
-            .lazy = true,
-        },
         .pretty = .{
             .url = "git+https://github.com/Kbz-8/pretty.git#f91d534d033277ca1ae7fcd598a070e8b3ddc532",
             .hash = "pretty-0.10.6-Tm65r7FTAQA5BEL8tcIcF-Wp4XRC7J7BhuRI0KUnFj2X",
         },
     },
-    .minimum_zig_version = "0.16.0",
     .paths = .{
         "build.zig",
         "build.zig.zon",
         "src",
         "LICENSE",
     },
-    .fingerprint = 0xd87ec8ab9fa9396a,
 }

example/main.zig

@@ -1,152 +0,0 @@
-const std = @import("std");
-const sdl3 = @import("sdl3");
-const spv = @import("spv");
-
-const shader_source = @embedFile("shader.spv");
-
-const screen_width = 400;
-const screen_height = 240;
-
-pub fn main() !void {
-    {
-        //var gpa: std.heap.DebugAllocator(.{}) = .init;
-        //defer _ = gpa.deinit();
-
-        const allocator = std.heap.smp_allocator;
-
-        var threaded: std.Io.Threaded = .init(allocator, .{
-            .async_limit = @enumFromInt(screen_height),
-        });
-        defer threaded.deinit();
-        const io = threaded.io();
-
-        defer sdl3.shutdown();
-        const init_flags = sdl3.InitFlags{ .video = true, .events = true };
-        try sdl3.init(init_flags);
-        defer sdl3.quit(init_flags);
-
-        const window = try sdl3.video.Window.init("Hello triangle", screen_width, screen_height, .{});
-        defer window.deinit();
-
-        const surface = try window.getSurface();
-
-        var module = try spv.Module.init(allocator, @ptrCast(@alignCast(shader_source)), .{});
-        defer module.deinit(allocator);
-
-        const fragment_count = screen_height * screen_width;
-
-        const batch_size = switch (threaded.async_limit) {
-            .nothing => 1,
-            .unlimited => std.Thread.getCpuCount() catch 1, // If we cannot get the CPU count, fallback on single runtime
-            else => |count| @intFromEnum(count),
-        };
-
-        const runners: []Runner = try allocator.alloc(Runner, batch_size);
-        defer {
-            for (runners) |*runner| {
-                runner.rt.deinit(allocator);
-            }
-            allocator.free(runners);
-        }
-
-        for (runners) |*runner| {
-            var rt = try spv.Runtime.init(allocator, &module, undefined);
-            runner.* = .{
-                .allocator = allocator,
-                .surface = surface,
-                .rt = rt,
-                .entry = try rt.getEntryPointByName("main"),
-                .color = try rt.getResultByName("color"),
-                .time = try rt.getResultByName("time"),
-                .pos = try rt.getResultByName("pos"),
-                .res = try rt.getResultByName("res"),
-                .invocation_count = fragment_count,
-                .batch_size = batch_size,
-            };
-        }
-        const timer = std.Io.Timestamp.now(io, .real);
-
-        var wg: std.Io.Group = .init;
-        var quit = false;
-        while (!quit) {
-            try surface.clear(.{ .r = 0.0, .g = 0.0, .b = 0.0, .a = 1.0 });
-
-            while (sdl3.events.poll()) |event|
-                switch (event) {
-                    .quit => quit = true,
-                    .terminating => quit = true,
-                    else => {},
-                };
-
-            {
-                try surface.lock();
-                defer surface.unlock();
-
-                const pixel_map: [*]u32 = @as([*]u32, @ptrCast(@alignCast((surface.getPixels() orelse return).ptr)));
-
-                const duration = timer.untilNow(io, .real);
-                const delta: f32 = @as(f32, @floatFromInt(duration.toNanoseconds())) / std.time.ns_per_s;
-
-                for (0..@min(batch_size, fragment_count)) |batch_id| {
-                    wg.async(io, Runner.runWrapper, .{ &runners[batch_id], batch_id, pixel_map, delta });
-                }
-                try wg.await(io);
-            }
-
-            try window.updateSurface();
-        }
-    }
-    std.log.info("Successfully executed", .{});
-}
-
-const Runner = struct {
-    const Self = @This();
-
-    allocator: std.mem.Allocator,
-    surface: sdl3.surface.Surface,
-    rt: spv.Runtime,
-    entry: spv.SpvWord,
-    color: spv.SpvWord,
-    time: spv.SpvWord,
-    pos: spv.SpvWord,
-    res: spv.SpvWord,
-    invocation_count: usize,
-    batch_size: usize,
-
-    fn runWrapper(self: *Self, batch_id: usize, pixel_map: [*]u32, timer: f32) void {
-        @call(.always_inline, Self.run, .{ self, batch_id, pixel_map, timer }) catch |err| {
-            std.log.err("{s}", .{@errorName(err)});
-            if (@errorReturnTrace()) |trace| {
-                std.debug.dumpErrorReturnTrace(trace);
-            }
-            std.process.abort();
-        };
-    }
-
-    fn run(self: *Self, batch_id: usize, pixel_map: [*]u32, timer: f32) !void {
-        var rt = self.rt; // Copy to avoid pointer access of `self` at runtime. Okay as Runtime contains only pointers and trivially copyable fields
-
-        var output: [4]f32 = undefined;
-
-        var invocation_index: usize = batch_id;
-        while (invocation_index < self.invocation_count) : (invocation_index += self.batch_size) {
-            const y = @divTrunc(invocation_index, screen_width);
-            const x = @mod(invocation_index, screen_width);
-
-            try rt.writeInput(std.mem.asBytes(&timer), self.time);
-            try rt.writeInput(std.mem.asBytes(&[_]f32{ @floatFromInt(screen_width), @floatFromInt(screen_height) }), self.res);
-            try rt.writeInput(std.mem.asBytes(&[_]f32{ @floatFromInt(x), @floatFromInt(y) }), self.pos);
-            try rt.callEntryPoint(self.allocator, self.entry);
-            try rt.readOutput(std.mem.asBytes(output[0..]), self.color);
-
-            const rgba = self.surface.mapRgba(
-                @intCast(std.math.clamp(@as(i32, @intFromFloat(output[0] * 255.0)), 0, 255)),
-                @intCast(std.math.clamp(@as(i32, @intFromFloat(output[1] * 255.0)), 0, 255)),
-                @intCast(std.math.clamp(@as(i32, @intFromFloat(output[2] * 255.0)), 0, 255)),
-                @intCast(std.math.clamp(@as(i32, @intFromFloat(output[3] * 255.0)), 0, 255)),
-            );
-
-            pixel_map[invocation_index] = rgba.value;
-        }
-    }
-};

example/mangohud.conf

@@ -1,6 +0,0 @@
-gpu_stats=0
-font_size=16
-resolution
-hud_compact
-background_alpha=0
-width=140

example/shader.nzsl

@@ -1,68 +0,0 @@
-[nzsl_version("1.1")]
-module;
-
-struct FragIn
-{
-    [location(0)] time: f32,
-    [location(1)] res: vec2[f32],
-    [location(2)] pos: vec2[f32],
-}
-
-struct FragOut
-{
-    [location(0)] color: vec4[f32]
-}
-
-[entry(frag)]
-fn main(input: FragIn) -> FragOut
-{
-    const I: i32 = 128;
-    const A: f32 = 7.5;
-    const MA: f32 = 100.0;
-    const MI: f32 = 0.001;
-
-    let uv0 = input.pos / input.res * 2.0 - vec2[f32](1.0, 1.0);
-    let uv  = vec2[f32](uv0.x * (input.res.x / input.res.y), uv0.y);
-
-    let col = vec3[f32](0.0, 0.0, 0.0);  
-    let ro  = vec3[f32](0.0, 0.0, -2.0); 
-    let rd  = vec3[f32](uv.x, uv.y, 1.0);
-    let dt  = 0.0;
-    let ds  = 0.0;
-    let dm  = -1.0;
-    let p   = ro;
-    let c   = vec3[f32](0.0, 0.0, 0.0);
-
-    let l = vec3[f32](0.0, sin(input.time * 0.2) * 4.0, cos(input.time * 0.2) * 4.0);
-
-    for i in 0 -> I
-    {
-        p = ro + rd * dt;
-        ds = length(c - p) - 1.0;
-        dt += ds;
-
-        if (dm == -1.0 || ds < dm)
-            dm = ds;
-
-        if (ds <= MI)
-        {
-            let value = max(dot(normalize(c - p), normalize(p - l)) - 0.35, 0.0);
-            col = vec3[f32](value, value, value);
-            break;
-        }
-
-        if (ds >= MA)
-        {
-            if (dot(normalize(rd), normalize(l - ro)) <= 1.0)
-            {
-                let value = max(dot(normalize(rd), normalize(l - ro)) + 0.15, 0.05)/ 1.15 * (1.0 - dm * A);
-                col = vec3[f32](value, value, value);
-            }
-            break;
-        }
-    }
-
-    let output: FragOut;
-    output.color = vec4[f32](col.x, col.y, col.z, 1.0);
-    return output;
-}

example/shader.spv.txt

@@ -1,281 +0,0 @@
-Version 1.0
-Generator: 2560130
-Bound: 203
-Schema: 0
-       OpCapability Capability(Shader)
- %42 = OpExtInstImport "GLSL.std.450"
-       OpMemoryModel AddressingModel(Logical) MemoryModel(GLSL450)
-       OpEntryPoint ExecutionModel(Fragment) %43 "main" %5 %11 %14 %20
-       OpExecutionMode %43 ExecutionMode(OriginUpperLeft)
-       OpSource SourceLanguage(NZSL) 4198400
-       OpSourceExtension "Version: 1.1"
-       OpName %16 "FragIn"
-       OpMemberName %16 0 "time"
-       OpMemberName %16 1 "res"
-       OpMemberName %16 2 "pos"
-       OpName %21 "FragOut"
-       OpMemberName %21 0 "color"
-       OpName %5 "time"
-       OpName %11 "res"
-       OpName %14 "pos"
-       OpName %20 "color"
-       OpName %43 "main"
-       OpDecorate %5 Decoration(Location) 0
-       OpDecorate %11 Decoration(Location) 1
-       OpDecorate %14 Decoration(Location) 2
-       OpDecorate %20 Decoration(Location) 0
-       OpMemberDecorate %16 0 Decoration(Offset) 0
-       OpMemberDecorate %16 1 Decoration(Offset) 8
-       OpMemberDecorate %16 2 Decoration(Offset) 16
-       OpMemberDecorate %21 0 Decoration(Offset) 0
-  %1 = OpTypeVoid
-  %2 = OpTypeFunction %1
-  %3 = OpTypeFloat 32
-  %4 = OpTypePointer StorageClass(Input) %3
-  %6 = OpTypeInt 32 1
-  %7 = OpConstant %6 i32(0)
-  %8 = OpTypePointer StorageClass(Function) %3
-  %9 = OpTypeVector %3 2
- %10 = OpTypePointer StorageClass(Input) %9
- %12 = OpConstant %6 i32(1)
- %13 = OpTypePointer StorageClass(Function) %9
- %15 = OpConstant %6 i32(2)
- %16 = OpTypeStruct %3 %9 %9
- %17 = OpTypePointer StorageClass(Function) %16
- %18 = OpTypeVector %3 4
- %19 = OpTypePointer StorageClass(Output) %18
- %21 = OpTypeStruct %18
- %22 = OpConstant %3 f32(2)
- %23 = OpConstant %3 f32(1)
- %24 = OpConstant %3 f32(0)
- %25 = OpTypeVector %3 3
- %26 = OpTypePointer StorageClass(Function) %25
- %27 = OpConstant %3 f32(-2)
- %28 = OpConstant %3 f32(-1)
- %29 = OpConstant %3 f32(0.2)
- %30 = OpConstant %3 f32(4)
- %31 = OpTypePointer StorageClass(Function) %6
- %32 = OpConstant %6 i32(128)
- %33 = OpTypeBool
- %34 = OpConstant %3 f32(0.001)
- %35 = OpConstant %3 f32(0.35)
- %36 = OpConstant %3 f32(100)
- %37 = OpConstant %3 f32(0.15)
- %38 = OpConstant %3 f32(0.05)
- %39 = OpConstant %3 f32(1.15)
- %40 = OpConstant %3 f32(7.5)
- %41 = OpTypePointer StorageClass(Function) %21
-%200 = OpTypePointer StorageClass(Function) %18
-  %5 = OpVariable %4 StorageClass(Input)
- %11 = OpVariable %10 StorageClass(Input)
- %14 = OpVariable %10 StorageClass(Input)
- %20 = OpVariable %19 StorageClass(Output)
- %43 = OpFunction %1 FunctionControl(0) %2
- %44 = OpLabel
- %45 = OpVariable %13 StorageClass(Function)
- %46 = OpVariable %13 StorageClass(Function)
- %47 = OpVariable %26 StorageClass(Function)
- %48 = OpVariable %26 StorageClass(Function)
- %49 = OpVariable %26 StorageClass(Function)
- %50 = OpVariable %8 StorageClass(Function)
- %51 = OpVariable %8 StorageClass(Function)
- %52 = OpVariable %8 StorageClass(Function)
- %53 = OpVariable %26 StorageClass(Function)
- %54 = OpVariable %26 StorageClass(Function)
- %55 = OpVariable %26 StorageClass(Function)
- %56 = OpVariable %31 StorageClass(Function)
- %57 = OpVariable %31 StorageClass(Function)
- %58 = OpVariable %8 StorageClass(Function)
- %59 = OpVariable %8 StorageClass(Function)
- %60 = OpVariable %41 StorageClass(Function)
- %61 = OpVariable %17 StorageClass(Function)
- %62 = OpAccessChain %8 %61 %7
-       OpCopyMemory %62 %5
- %63 = OpAccessChain %13 %61 %12
-       OpCopyMemory %63 %11
- %64 = OpAccessChain %13 %61 %15
-       OpCopyMemory %64 %14
- %65 = OpAccessChain %13 %61 %15
- %66 = OpLoad %9 %65
- %67 = OpAccessChain %13 %61 %12
- %68 = OpLoad %9 %67
- %69 = OpFDiv %9 %66 %68
- %70 = OpVectorTimesScalar %9 %69 %22
- %71 = OpCompositeConstruct %9 %23 %23
- %72 = OpFSub %9 %70 %71
-       OpStore %45 %72
- %73 = OpLoad %9 %45
- %74 = OpCompositeExtract %3 %73 0
- %75 = OpAccessChain %13 %61 %12
- %76 = OpLoad %9 %75
- %77 = OpCompositeExtract %3 %76 0
- %78 = OpAccessChain %13 %61 %12
- %79 = OpLoad %9 %78
- %80 = OpCompositeExtract %3 %79 1
- %81 = OpFDiv %3 %77 %80
- %82 = OpFMul %3 %74 %81
- %83 = OpLoad %9 %45
- %84 = OpCompositeExtract %3 %83 1
- %85 = OpCompositeConstruct %9 %82 %84
-       OpStore %46 %85
- %86 = OpCompositeConstruct %25 %24 %24 %24
-       OpStore %47 %86
- %87 = OpCompositeConstruct %25 %24 %24 %27
-       OpStore %48 %87
- %88 = OpLoad %9 %46
- %89 = OpCompositeExtract %3 %88 0
- %90 = OpLoad %9 %46
- %91 = OpCompositeExtract %3 %90 1
- %92 = OpCompositeConstruct %25 %89 %91 %23
-       OpStore %49 %92
-       OpStore %50 %24
-       OpStore %51 %24
-       OpStore %52 %28
- %93 = OpLoad %25 %48
-       OpStore %53 %93
- %94 = OpCompositeConstruct %25 %24 %24 %24
-       OpStore %54 %94
- %95 = OpAccessChain %8 %61 %7
- %96 = OpLoad %3 %95
- %97 = OpFMul %3 %96 %29
- %98 = OpExtInst %3 GLSLstd450 Sin %97
- %99 = OpFMul %3 %98 %30
-%100 = OpAccessChain %8 %61 %7
-%101 = OpLoad %3 %100
-%102 = OpFMul %3 %101 %29
-%103 = OpExtInst %3 GLSLstd450 Cos %102
-%104 = OpFMul %3 %103 %30
-%105 = OpCompositeConstruct %25 %24 %99 %104
-       OpStore %55 %105
-       OpStore %56 %7
-       OpStore %57 %32
-       OpBranch %106
-%106 = OpLabel
-%110 = OpLoad %6 %56
-%111 = OpLoad %6 %57
-%112 = OpSLessThan %33 %110 %111
-       OpLoopMerge %108 %109 LoopControl(0)
-       OpBranchConditional %112 %107 %108
-%107 = OpLabel
-%113 = OpLoad %25 %48
-%114 = OpLoad %25 %49
-%115 = OpLoad %3 %50
-%116 = OpVectorTimesScalar %25 %114 %115
-%117 = OpFAdd %25 %113 %116
-       OpStore %53 %117
-%118 = OpLoad %25 %54
-%119 = OpLoad %25 %53
-%120 = OpFSub %25 %118 %119
-%121 = OpExtInst %3 GLSLstd450 Length %120
-%122 = OpFSub %3 %121 %23
-       OpStore %51 %122
-%123 = OpLoad %3 %50
-%124 = OpLoad %3 %51
-%125 = OpFAdd %3 %123 %124
-       OpStore %50 %125
-%129 = OpLoad %3 %52
-%130 = OpFOrdEqual %33 %129 %28
-%131 = OpLoad %3 %51
-%132 = OpLoad %3 %52
-%133 = OpFOrdLessThan %33 %131 %132
-%134 = OpLogicalOr %33 %130 %133
-       OpSelectionMerge %126 SelectionControl(0)
-       OpBranchConditional %134 %127 %128
-%127 = OpLabel
-%135 = OpLoad %3 %51
-       OpStore %52 %135
-       OpBranch %126
-%128 = OpLabel
-       OpBranch %126
-%126 = OpLabel
-%139 = OpLoad %3 %51
-%140 = OpFOrdLessThanEqual %33 %139 %34
-       OpSelectionMerge %136 SelectionControl(0)
-       OpBranchConditional %140 %137 %138
-%137 = OpLabel
-%141 = OpLoad %25 %54
-%142 = OpLoad %25 %53
-%143 = OpFSub %25 %141 %142
-%144 = OpExtInst %25 GLSLstd450 Normalize %143
-%145 = OpLoad %25 %53
-%146 = OpLoad %25 %55
-%147 = OpFSub %25 %145 %146
-%148 = OpExtInst %25 GLSLstd450 Normalize %147
-%149 = OpDot %3 %144 %148
-%150 = OpFSub %3 %149 %35
-%151 = OpExtInst %3 GLSLstd450 FMax %150 %24
-       OpStore %58 %151
-%152 = OpLoad %3 %58
-%153 = OpLoad %3 %58
-%154 = OpLoad %3 %58
-%155 = OpCompositeConstruct %25 %152 %153 %154
-       OpStore %47 %155
-       OpBranch %108
-%138 = OpLabel
-       OpBranch %136
-%136 = OpLabel
-%159 = OpLoad %3 %51
-%160 = OpFOrdGreaterThanEqual %33 %159 %36
-       OpSelectionMerge %156 SelectionControl(0)
-       OpBranchConditional %160 %157 %158
-%157 = OpLabel
-%164 = OpLoad %25 %49
-%165 = OpExtInst %25 GLSLstd450 Normalize %164
-%166 = OpLoad %25 %55
-%167 = OpLoad %25 %48
-%168 = OpFSub %25 %166 %167
-%169 = OpExtInst %25 GLSLstd450 Normalize %168
-%170 = OpDot %3 %165 %169
-%171 = OpFOrdLessThanEqual %33 %170 %23
-       OpSelectionMerge %161 SelectionControl(0)
-       OpBranchConditional %171 %162 %163
-%162 = OpLabel
-%172 = OpLoad %25 %49
-%173 = OpExtInst %25 GLSLstd450 Normalize %172
-%174 = OpLoad %25 %55
-%175 = OpLoad %25 %48
-%176 = OpFSub %25 %174 %175
-%177 = OpExtInst %25 GLSLstd450 Normalize %176
-%178 = OpDot %3 %173 %177
-%179 = OpFAdd %3 %178 %37
-%180 = OpExtInst %3 GLSLstd450 FMax %179 %38
-%181 = OpFDiv %3 %180 %39
-%182 = OpLoad %3 %52
-%183 = OpFMul %3 %182 %40
-%184 = OpFSub %3 %23 %183
-%185 = OpFMul %3 %181 %184
-       OpStore %59 %185
-%186 = OpLoad %3 %59
-%187 = OpLoad %3 %59
-%188 = OpLoad %3 %59
-%189 = OpCompositeConstruct %25 %186 %187 %188
-       OpStore %47 %189
-       OpBranch %161
-%163 = OpLabel
-       OpBranch %161
-%161 = OpLabel
-       OpBranch %108
-%158 = OpLabel
-       OpBranch %156
-%156 = OpLabel
-%190 = OpLoad %6 %56
-%191 = OpIAdd %6 %190 %12
-       OpStore %56 %191
-       OpBranch %109
-%109 = OpLabel
-       OpBranch %106
-%108 = OpLabel
-%192 = OpLoad %25 %47
-%193 = OpCompositeExtract %3 %192 0
-%194 = OpLoad %25 %47
-%195 = OpCompositeExtract %3 %194 1
-%196 = OpLoad %25 %47
-%197 = OpCompositeExtract %3 %196 2
-%198 = OpCompositeConstruct %18 %193 %195 %197 %23
-%199 = OpAccessChain %200 %60 %7
-       OpStore %199 %198
-%201 = OpLoad %21 %60
-%202 = OpCompositeExtract %18 %201 0
-       OpStore %20 %202
-       OpReturn
-       OpFunctionEnd

ffi/SpirvInterpreter.h

@@ -413,25 +413,46 @@ 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,
 	SPV_RESULT_BARRIER = 1,
 	SPV_RESULT_KILLED = 2,
-	SPV_RESULT_DIVISION_BY_ZERO = -1,
+	SPV_RESULT_INVALID_ENTRY_POINT = -1,
-	SPV_RESULT_INVALID_ENTRY_POINT = -2,
+	SPV_RESULT_INVALID_SPIRV = -2,
-	SPV_RESULT_INVALID_SPIRV = -3,
+	SPV_RESULT_INVALID_VALUE_TYPE = -3,
-	SPV_RESULT_INVALID_VALUE_TYPE = -4,
+	SPV_RESULT_NOT_FOUND = -4,
-	SPV_RESULT_NOT_FOUND = -5,
+	SPV_RESULT_OUT_OF_MEMORY = -5,
-	SPV_RESULT_OUT_OF_MEMORY = -6,
+	SPV_RESULT_OUT_OF_BOUNDS = -6,
-	SPV_RESULT_OUT_OF_BOUNDS = -7,
+	SPV_RESULT_TODO = -7,
-	SPV_RESULT_TODO = -8,
+	SPV_RESULT_UNREACHABLE = -8,
-	SPV_RESULT_UNREACHABLE = -9,
+	SPV_RESULT_UNSUPPORTED_SPIRV = -9,
-	SPV_RESULT_UNSUPPORTED_SPIRV = -10,
+	SPV_RESULT_UNSUPPORTED_EXTENSION = -10,
-	SPV_RESULT_UNSUPPORTED_EXTENSION = -11,
+	SPV_RESULT_UNSUPPORTED_ENDIANNESS = -11,
-	SPV_RESULT_UNSUPPORTED_ENDIANNESS = -12,
+	SPV_RESULT_INVALID_MAGIC = -12,
-	SPV_RESULT_INVALID_MAGIC = -13,
+	SPV_RESULT_UNKNOWN = -13
-	SPV_RESULT_UNKNOWN = -14
 } SpvResult;
 
 typedef struct
@@ -452,6 +473,8 @@ typedef struct
 
     SpvBool needs_derivatives;
     SpvBool has_control_barriers;
+    SpvBool has_atomics;
+    SpvBool early_fragment_tests;
 } SpvModuleReflectionInfos;
 
 typedef struct
@@ -473,13 +496,20 @@ typedef enum
 	SPV_ENTRY_POINT_BARRIER = 1
 } SpvEntryPointStatus;
 
-typedef enum
+typedef struct
 {
-	SPV_PRIMITIVE_BOOL = 0,
+	SpvBool has_size;
-	SPV_PRIMITIVE_FLOAT = 1,
+	SpvWord x;
-	SPV_PRIMITIVE_SINT = 2,
+	SpvWord y;
-	SPV_PRIMITIVE_UINT = 3
+	SpvWord z;
-} SpvPrimitiveType;
+} SpvWorkgroupSize;
+
+typedef struct
+{
+	SpvWord result;
+	SpvByte* data;
+	SpvSize size;
+} SpvWorkgroupMemoryItem;
 
 typedef struct
 {
@@ -538,6 +568,7 @@ typedef struct
 	int x;
 	int y;
 	int z;
+	float w;
 	float lod;
 	SpvBool has_lod;
 	SpvImageOffset offset;
@@ -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/ffi.zig

@@ -8,20 +8,19 @@ pub const Result = enum(c_int) {
     Success = 0,
     Barrier = 1,
     Killed = 2,
-    DivisionByZero = -1,
+    InvalidEntryPoint = -1,
-    InvalidEntryPoint = -2,
+    InvalidSpirV = -2,
-    InvalidSpirV = -3,
+    InvalidValueType = -3,
-    InvalidValueType = -4,
+    NotFound = -4,
-    NotFound = -5,
+    OutOfMemory = -5,
-    OutOfMemory = -6,
+    OutOfBounds = -6,
-    OutOfBounds = -7,
+    ToDo = -7,
-    ToDo = -8,
+    Unreachable = -8,
-    Unreachable = -9,
+    UnsupportedSpirV = -9,
-    UnsupportedSpirV = -10,
+    UnsupportedExtension = -10,
-    UnsupportedExtension = -11,
+    UnsupportedEndianness = -11,
-    UnsupportedEndianness = -12,
+    InvalidMagic = -12,
-    InvalidMagic = -13,
+    Unknown = -13,
-    Unknown = -14,
 };
 
 comptime {

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,
@@ -75,6 +88,7 @@ const SampleImageInfo = extern struct {
     x: f32,
     y: f32,
     z: f32,
+    w: f32,
     lod: f32,
     has_lod: ffi.SpvCBool,
     offset: ImageOffset,
@@ -124,7 +138,6 @@ const ImageAPI = extern struct {
 fn toCResult(err: spv.Runtime.RuntimeError) ffi.Result {
     return switch (err) {
         spv.Runtime.RuntimeError.Barrier => ffi.Result.Barrier,
-        spv.Runtime.RuntimeError.DivisionByZero => ffi.Result.DivisionByZero,
         spv.Runtime.RuntimeError.InvalidEntryPoint => ffi.Result.InvalidEntryPoint,
         spv.Runtime.RuntimeError.InvalidSpirV => ffi.Result.InvalidSpirV,
         spv.Runtime.RuntimeError.InvalidValueType => ffi.Result.InvalidValueType,
@@ -143,7 +156,6 @@ fn toCResult(err: spv.Runtime.RuntimeError) ffi.Result {
 fn fromCResult(res: ffi.Result) spv.Runtime.RuntimeError!void {
     return switch (res) {
         ffi.Result.Barrier => spv.Runtime.RuntimeError.Barrier,
-        ffi.Result.DivisionByZero => spv.Runtime.RuntimeError.DivisionByZero,
         ffi.Result.InvalidEntryPoint => spv.Runtime.RuntimeError.InvalidEntryPoint,
         ffi.Result.InvalidSpirV => spv.Runtime.RuntimeError.InvalidSpirV,
         ffi.Result.InvalidValueType => spv.Runtime.RuntimeError.InvalidValueType,
@@ -205,7 +217,7 @@ const ImageAPIBridge = struct {
         };
     }
 
-    fn sampleImageInfo(driver_image: *anyopaque, driver_sampler: *anyopaque, dim: spv.spv.SpvDim, x: f32, y: f32, z: f32, lod: ?f32, offset: spv.Runtime.ImageOffset) SampleImageInfo {
+    fn sampleImageInfo(driver_image: *anyopaque, driver_sampler: *anyopaque, dim: spv.spv.SpvDim, x: f32, y: f32, z: f32, w: f32, lod: ?f32, offset: spv.Runtime.ImageOffset) SampleImageInfo {
         return .{
             .driver_image = driver_image,
             .driver_sampler = driver_sampler,
@@ -213,6 +225,7 @@ const ImageAPIBridge = struct {
             .x = x,
             .y = y,
             .z = z,
+            .w = w,
             .lod = lod orelse 0.0,
             .has_lod = if (lod == null) 0 else 1,
             .offset = toCImageOffset(offset),
@@ -279,7 +292,7 @@ const ImageAPIBridge = struct {
         const image_api = try getImageAPI();
 
         var dst: Vec4f = undefined;
-        const result = image_api.sampleImageFloat4(sampleImageInfo(driver_image, driver_sampler, dim, x, y, z, lod, offset), &dst);
+        const result = image_api.sampleImageFloat4(sampleImageInfo(driver_image, driver_sampler, dim, x, y, z, 1.0, lod, offset), &dst);
 
         try fromCResult(result);
 
@@ -295,7 +308,7 @@ const ImageAPIBridge = struct {
         const image_api = try getImageAPI();
 
         var dst: Vec4u = undefined;
-        const result = image_api.sampleImageInt4(sampleImageInfo(driver_image, driver_sampler, dim, x, y, z, lod, offset), &dst);
+        const result = image_api.sampleImageInt4(sampleImageInfo(driver_image, driver_sampler, dim, x, y, z, 1.0, lod, offset), &dst);
 
         try fromCResult(result);
 
@@ -307,11 +320,11 @@ const ImageAPIBridge = struct {
         };
     }
 
-    fn sampleImageDref(driver_image: *anyopaque, driver_sampler: *anyopaque, dim: spv.spv.SpvDim, x: f32, y: f32, z: f32, dref: f32, lod: ?f32, offset: spv.Runtime.ImageOffset) spv.Runtime.RuntimeError!f32 {
+    fn sampleImageDref(driver_image: *anyopaque, driver_sampler: *anyopaque, dim: spv.spv.SpvDim, x: f32, y: f32, z: f32, w: f32, dref: f32, lod: ?f32, offset: spv.Runtime.ImageOffset) spv.Runtime.RuntimeError!f32 {
         const image_api = try getImageAPI();
 
         var dst: f32 = undefined;
-        const result = image_api.sampleImageDref(sampleImageInfo(driver_image, driver_sampler, dim, x, y, z, lod, offset), dref, &dst);
+        const result = image_api.sampleImageDref(sampleImageInfo(driver_image, driver_sampler, dim, x, y, z, w, lod, offset), dref, &dst);
 
         try fromCResult(result);
 
@@ -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 {
+export fn SpvHasResultDecoration(rt: *RuntimeWrapper, result: spv.SpvWord, decoration: spv.spv.SpvDecoration) callconv(.c) c_int {
-    primitive_type.* = switch (rt.rt.getResultPrimitiveType(result) catch |err| return toCResult(err)) {
+    return if (rt.rt.hasResultDecoration(result, decoration)) 1 else 0;
-        .Bool => .bool,
+}
-        .Float => .float,
+
-        .SInt => .sint,
+export fn SpvGetWorkgroupSize(rt: *RuntimeWrapper, size: *WorkgroupSize) callconv(.c) ffi.Result {
-        .UInt => .uint,
+    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 {
+export fn SpvDestroyWorkgroupMemory(rt: *RuntimeWrapper, workgroup_memory: *WorkgroupMemoryWrapper) callconv(.c) void {
-    return if (rt.rt.hasResultDecoration(result, decoration)) 1 else 0;
+    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,8 +697,14 @@ 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 {
-    rt.rt.writeInput(input[0..input_size], result) catch |err| return toCResult(err);
+    const allocator = std.heap.c_allocator;
+    rt.rt.writeInput(allocator, input[0..input_size], result) catch |err| return toCResult(err);
     return .Success;
 }
 
@@ -618,7 +714,8 @@ export fn SpvWriteInputLocation(rt: *RuntimeWrapper, input: [*]const u8, input_s
 }
 
 export fn SpvWriteBuiltIn(rt: *RuntimeWrapper, input: [*]const u8, input_size: c_ulong, builtin: spv.spv.SpvBuiltIn) callconv(.c) ffi.Result {
-    rt.rt.writeBuiltIn(input[0..input_size], builtin) catch |err| return toCResult(err);
+    const allocator = std.heap.c_allocator;
+    rt.rt.writeBuiltIn(allocator, input[0..input_size], builtin) catch |err| return toCResult(err);
     return .Success;
 }
 

sandbox/main.zig

@@ -1,53 +0,0 @@
-const std = @import("std");
-const spv = @import("spv");
-
-const shader_source = @embedFile("shader.spv");
-
-const SSBO = struct {
-    value: [256]i32 = [_]i32{0} ** 256,
-};
-
-pub fn main() !void {
-    {
-        var gpa: std.heap.DebugAllocator(.{
-            .enable_memory_limit = true,
-        }) = .init;
-        defer _ = gpa.deinit();
-
-        const allocator = gpa.allocator();
-
-        var module = try spv.Module.init(allocator, @ptrCast(@alignCast(shader_source)), .{});
-        defer module.deinit(allocator);
-
-        var rt = try spv.Runtime.init(allocator, &module);
-        defer rt.deinit(allocator);
-
-        const entry = try rt.getEntryPointByName("main");
-
-        var ssbo: SSBO = .{};
-
-        try rt.writeDescriptorSet(std.mem.asBytes(&ssbo), 0, 0, 0);
-
-        for (0..16) |i| {
-            for (0..16) |x| {
-                const global_invocation_indices = [3]i32{
-                    @as(i32, @intCast(i * 16 + x)),
-                    1,
-                    1,
-                };
-
-                try rt.writeBuiltIn(std.mem.asBytes(&global_invocation_indices), .GlobalInvocationId);
-                rt.callEntryPoint(allocator, entry) catch |err| switch (err) {
-                    spv.Runtime.RuntimeError.OutOfBounds => continue,
-                    else => return err,
-                };
-                try rt.flushDescriptorSets(allocator);
-            }
-        }
-
-        std.log.info("Output: {any}", .{ssbo});
-
-        std.log.info("Total memory used: {d:.3} KB\n", .{@as(f32, @floatFromInt(gpa.total_requested_bytes)) / 1000.0});
-    }
-    std.log.info("Successfully executed", .{});
-}

sandbox/shader.nzsl

@@ -1,25 +0,0 @@
-[nzsl_version("1.1")]
-module;
-
-struct Input
-{
-    [builtin(global_invocation_indices)] indices: vec3[u32]
-}
-
-[layout(std430)]
-struct SSBO
-{
-    data: dyn_array[i32]
-}
-
-external
-{
-    [set(0), binding(0)] ssbo: storage[SSBO],
-}
-
-[entry(compute)]
-[workgroup(16, 1, 1)]
-fn main(input: Input)
-{
-	ssbo.data[input.indices.x] = i32(input.indices.x);
-}

sandbox/shader.spv.txt

@@ -1,59 +0,0 @@
-Version 1.0
-Generator: 2560130
-Bound: 32
-Schema: 0
-      OpCapability Capability(Shader)
-      OpMemoryModel AddressingModel(Logical) MemoryModel(GLSL450)
-      OpEntryPoint ExecutionModel(GLCompute) %17 "main" %11
-      OpExecutionMode %17 ExecutionMode(LocalSize) 16 1 1
-      OpSource SourceLanguage(NZSL) 4198400
-      OpSourceExtension "Version: 1.1"
-      OpName %3 "SSBO"
-      OpMemberName %3 0 "data"
-      OpName %14 "Input"
-      OpMemberName %14 0 "indices"
-      OpName %5 "ssbo"
-      OpName %11 "global_invocation_indices"
-      OpName %17 "main"
-      OpDecorate %5 Decoration(Binding) 0
-      OpDecorate %5 Decoration(DescriptorSet) 0
-      OpDecorate %11 Decoration(BuiltIn) BuiltIn(GlobalInvocationId)
-      OpDecorate %2 Decoration(ArrayStride) 4
-      OpDecorate %3 Decoration(BufferBlock)
-      OpMemberDecorate %3 0 Decoration(Offset) 0
-      OpMemberDecorate %14 0 Decoration(Offset) 0
- %1 = OpTypeInt 32 1
- %2 = OpTypeRuntimeArray %1
- %3 = OpTypeStruct %2
- %4 = OpTypePointer StorageClass(Uniform) %3
- %6 = OpTypeVoid
- %7 = OpTypeFunction %6
- %8 = OpTypeInt 32 0
- %9 = OpTypeVector %8 3
-%10 = OpTypePointer StorageClass(Input) %9
-%12 = OpConstant %1 i32(0)
-%13 = OpTypePointer StorageClass(Function) %9
-%14 = OpTypeStruct %9
-%15 = OpTypePointer StorageClass(Function) %14
-%16 = OpTypeRuntimeArray %1
-%26 = OpTypePointer StorageClass(Uniform) %2
-%31 = OpTypePointer StorageClass(Uniform) %1
- %5 = OpVariable %4 StorageClass(Uniform)
-%11 = OpVariable %10 StorageClass(Input)
-%17 = OpFunction %6 FunctionControl(0) %7
-%18 = OpLabel
-%19 = OpVariable %15 StorageClass(Function)
-%20 = OpAccessChain %13 %19 %12
-      OpCopyMemory %20 %11
-%21 = OpAccessChain %13 %19 %12
-%22 = OpLoad %9 %21
-%23 = OpCompositeExtract %8 %22 0
-%24 = OpBitcast %1 %23
-%25 = OpAccessChain %26 %5 %12
-%27 = OpAccessChain %13 %19 %12
-%28 = OpLoad %9 %27
-%29 = OpCompositeExtract %8 %28 0
-%30 = OpAccessChain %31 %25 %29
-      OpStore %30 %24
-      OpReturn
-      OpFunctionEnd

src/GLSL_std_450/opcodes.zig

@@ -258,10 +258,25 @@ fn readIntegralLaneAsI32(src: *const Value, lane_index: usize) RuntimeError!i32
     const lane_bits = try src.resolveLaneBitWidth();
     const sign = try src.resolveSign();
     return switch (lane_bits) {
-        inline 8, 16, 32, 64 => |bits| if (sign == .signed)
+        inline 8, 16, 32, 64 => |bits| blk: {
-            @intCast(try Value.readLane(.SInt, bits, src, lane_index))
+            if (sign == .signed) {
+                const value = try Value.readLane(.SInt, bits, src, lane_index);
+                if (bits > 32) {
+                    break :blk std.math.cast(i32, value) orelse if (value < 0)
+                        std.math.minInt(i32)
                     else
-            @intCast(try Value.readLane(.UInt, bits, src, lane_index)),
+                        std.math.maxInt(i32);
+                }
+                break :blk @intCast(value);
+            } else {
+                const value = try Value.readLane(.UInt, bits, src, lane_index);
+                const bits32: u32 = if (bits > 32)
+                    std.math.cast(u32, value) orelse std.math.maxInt(u32)
+                else
+                    @intCast(value);
+                break :blk @bitCast(bits32);
+            }
+        },
         else => RuntimeError.InvalidSpirV,
     };
 }
@@ -874,11 +889,8 @@ fn opModfStruct(_: std.mem.Allocator, _: SpvWord, id: SpvWord, _: SpvWord, rt: *
 }
 
 fn frexpScalar(comptime T: type, x: T) struct { sig: T, exp: i32 } {
-    if (x == 0.0)
+    const result = std.math.frexp(x);
-        return .{ .sig = 0.0, .exp = 0 };
+    return .{ .sig = result.significand, .exp = result.exponent };
-
-    const exp = @as(i32, @intFromFloat(@floor(@log2(@abs(x))))) + 1;
-    return .{ .sig = x / @exp2(@as(T, @floatFromInt(exp))), .exp = exp };
 }
 
 fn opFrexp(_: std.mem.Allocator, target_type_id: SpvWord, id: SpvWord, _: SpvWord, rt: *Runtime) RuntimeError!void {
@@ -936,8 +948,9 @@ fn opLdexp(_: std.mem.Allocator, target_type_id: SpvWord, id: SpvWord, _: SpvWor
         inline 16, 32, 64 => |bits| for (0..lane_count) |lane_index| {
             const exp_lane = if (exp_lane_count == 1) 0 else lane_index;
             const exponent = try readIntegralLaneAsI32(exp_value, exp_lane);
-            const scale = @exp2(@as(std.meta.Float(bits), @floatFromInt(exponent)));
+            const value = try Value.readLane(.Float, bits, x, lane_index);
-            try writeFloatLane(bits, dst, lane_index, try Value.readLane(.Float, bits, x, lane_index) * scale);
+            const result = if (value == 0.0) value else std.math.ldexp(value, exponent);
+            try writeFloatLane(bits, dst, lane_index, result);
         },
         else => return RuntimeError.InvalidSpirV,
     }
@@ -1106,7 +1119,51 @@ fn matrixSize(matrix: *const Value) RuntimeError!usize {
     };
 }
 
+fn determinant3Values(comptime T: type, a00: T, a01: T, a02: T, a10: T, a11: T, a12: T, a20: T, a21: T, a22: T) T {
+    return a00 * ((a11 * a22) - (a12 * a21)) -
+        a01 * ((a10 * a22) - (a12 * a20)) +
+        a02 * ((a10 * a21) - (a11 * a20));
+}
+
 fn determinant(comptime bits: u32, src: *const Value, n: usize) RuntimeError!std.meta.Float(bits) {
+    const FloatT = std.meta.Float(bits);
+    if (n == 2) {
+        const a00 = try readMatrixElement(bits, src, 0, 0);
+        const a01 = try readMatrixElement(bits, src, 0, 1);
+        const a10 = try readMatrixElement(bits, src, 1, 0);
+        const a11 = try readMatrixElement(bits, src, 1, 1);
+        return (a00 * a11) - (a10 * a01);
+    }
+
+    if (n == 3) {
+        return determinant3Values(
+            FloatT,
+            try readMatrixElement(bits, src, 0, 0),
+            try readMatrixElement(bits, src, 0, 1),
+            try readMatrixElement(bits, src, 0, 2),
+            try readMatrixElement(bits, src, 1, 0),
+            try readMatrixElement(bits, src, 1, 1),
+            try readMatrixElement(bits, src, 1, 2),
+            try readMatrixElement(bits, src, 2, 0),
+            try readMatrixElement(bits, src, 2, 1),
+            try readMatrixElement(bits, src, 2, 2),
+        );
+    }
+
+    if (n == 4) {
+        var m: [4][4]FloatT = undefined;
+        for (0..4) |row| {
+            for (0..4) |col| {
+                m[row][col] = try readMatrixElement(bits, src, row, col);
+            }
+        }
+
+        return m[0][0] * determinant3Values(FloatT, m[1][1], m[1][2], m[1][3], m[2][1], m[2][2], m[2][3], m[3][1], m[3][2], m[3][3]) -
+            m[0][1] * determinant3Values(FloatT, m[1][0], m[1][2], m[1][3], m[2][0], m[2][2], m[2][3], m[3][0], m[3][2], m[3][3]) +
+            m[0][2] * determinant3Values(FloatT, m[1][0], m[1][1], m[1][3], m[2][0], m[2][1], m[2][3], m[3][0], m[3][1], m[3][3]) -
+            m[0][3] * determinant3Values(FloatT, m[1][0], m[1][1], m[1][2], m[2][0], m[2][1], m[2][2], m[3][0], m[3][1], m[3][2]);
+    }
+
     var m: [16]std.meta.Float(bits) = undefined;
     for (0..n) |row| {
         for (0..n) |col| {
@@ -1171,6 +1228,19 @@ fn opMatrixInverse(_: std.mem.Allocator, target_type_id: SpvWord, id: SpvWord, _
 
     switch (lane_bits) {
         inline 16, 32, 64 => |bits| {
+            if (n == 2) {
+                const a00 = try readMatrixElement(bits, src, 0, 0);
+                const a01 = try readMatrixElement(bits, src, 0, 1);
+                const a10 = try readMatrixElement(bits, src, 1, 0);
+                const a11 = try readMatrixElement(bits, src, 1, 1);
+                const inv_det = 1.0 / ((a00 * a11) - (a10 * a01));
+                try writeMatrixElement(bits, dst, 0, 0, a11 * inv_det);
+                try writeMatrixElement(bits, dst, 0, 1, -a01 * inv_det);
+                try writeMatrixElement(bits, dst, 1, 0, -a10 * inv_det);
+                try writeMatrixElement(bits, dst, 1, 1, a00 * inv_det);
+                return;
+            }
+
             var aug: [32]std.meta.Float(bits) = undefined;
             const width = n * 2;
             for (0..n) |row| {

src/Module.zig

@@ -54,6 +54,8 @@ pub const ReflectionInfos = struct {
 
     needs_derivatives: bool,
     has_control_barriers: bool,
+    has_atomics: bool,
+    early_fragment_tests: bool,
 };
 
 options: ModuleOptions,
@@ -274,7 +276,69 @@ fn applyStructMemberInterfaceDecorations(self: *Self, storage_class: spv.SpvStor
 }
 
 fn applyDecorations(self: *Self, allocator: std.mem.Allocator) ModuleError!void {
-    for (self.results, 0..) |result, id| {
+    const helpers = struct {
+        fn applyValueLayout(results: []Result, value: *Value, type_word: SpvWord) ModuleError!void {
+            const type_variant = results[type_word].variant orelse return;
+            const resolved_type_word = switch (type_variant) {
+                .Type => |type_data| switch (type_data) {
+                    .Pointer => |ptr| ptr.target,
+                    else => type_word,
+                },
+                else => return,
+            };
+            const resolved_type_variant = results[resolved_type_word].variant orelse return;
+            const type_data = switch (resolved_type_variant) {
+                .Type => |type_data| type_data,
+                else => return,
+            };
+
+            switch (value.*) {
+                .Structure => |*structure| switch (type_data) {
+                    .Structure => |type_structure| {
+                        @memcpy(@constCast(structure.offsets), type_structure.members_offsets);
+                        @memcpy(@constCast(structure.matrix_strides), type_structure.members_matrix_strides);
+                        @memcpy(@constCast(structure.row_major), type_structure.members_row_major);
+
+                        for (structure.values, type_structure.members_type_word, 0..) |*member_value, member_type_word, member_index| {
+                            if (member_value.* == .RuntimeArray) {
+                                member_value.RuntimeArray.matrix_stride = type_structure.members_matrix_strides[member_index];
+                                member_value.RuntimeArray.row_major = type_structure.members_row_major[member_index];
+                            }
+                            try applyValueLayout(results, member_value, member_type_word);
+                        }
+                    },
+                    else => {},
+                },
+                .Array => |array| switch (type_data) {
+                    .Array => |type_array| {
+                        for (array.values) |*element| {
+                            try applyValueLayout(results, element, type_array.components_type_word);
+                        }
+                    },
+                    else => {},
+                },
+                .Matrix => |columns| switch (type_data) {
+                    .Matrix => |type_matrix| {
+                        for (columns) |*column| {
+                            try applyValueLayout(results, column, type_matrix.column_type_word);
+                        }
+                    },
+                    else => {},
+                },
+                .Vector => |elements| switch (type_data) {
+                    .Vector => |type_vector| {
+                        for (elements) |*element| {
+                            try applyValueLayout(results, element, type_vector.components_type_word);
+                        }
+                    },
+                    else => {},
+                },
+                else => {},
+            }
+        }
+    };
+
+    for (self.results, 0..) |*result, id| {
         if (result.variant == null)
             continue;
 
@@ -306,6 +370,8 @@ fn applyDecorations(self: *Self, allocator: std.mem.Allocator) ModuleError!void
                             switch (decoration.rtype) {
                                 .Offset => s.members_offsets[decoration.index] = decoration.literal_1,
                                 .MatrixStride => s.members_matrix_strides[decoration.index] = decoration.literal_1,
+                                .RowMajor => s.members_row_major[decoration.index] = true,
+                                .ColMajor => s.members_row_major[decoration.index] = false,
                                 else => {},
                             }
                         },
@@ -323,18 +389,7 @@ fn applyDecorations(self: *Self, allocator: std.mem.Allocator) ModuleError!void
                     .StorageBuffer,
                     .Uniform,
                     .PushConstant,
-                    => if (v.value == .Structure) {
+                    => try helpers.applyValueLayout(self.results, &v.value, v.type_word),
-                        if (self.results[v.type_word].variant) |type_variant| switch (type_variant) {
-                            .Type => |type_data| switch (type_data) {
-                                .Structure => |s| {
-                                    @memcpy(@constCast(v.value.Structure.offsets), s.members_offsets);
-                                    @memcpy(@constCast(v.value.Structure.matrix_strides), s.members_matrix_strides);
-                                },
-                                else => {},
-                            },
-                            else => {},
-                        };
-                    },
                     else => {},
                 }
             },

src/Result.zig

@@ -123,6 +123,7 @@ pub const TypeData = union(Type) {
         members_type_word: []const SpvWord,
         members_offsets: []?SpvWord,
         members_matrix_strides: []?SpvWord,
+        members_row_major: []bool,
         member_names: std.ArrayList([]const u8),
     },
     Function: struct {
@@ -154,7 +155,7 @@ pub const TypeData = union(Type) {
             .Int => |i| @divExact(i.bit_length, 8),
             .Float => |f| @divExact(f.bit_length, 8),
             .Vector => |v| results[v.components_type_word].variant.?.Type.getSize(results) * v.member_count,
-            .Array => |a| a.stride,
+            .Array => |a| a.stride * a.member_count,
             .Matrix => |m| results[m.column_type_word].variant.?.Type.getSize(results) * m.member_count,
             .RuntimeArray => |a| a.stride,
             .Structure => |s| blk: {
@@ -247,6 +248,7 @@ pub fn deinit(self: *Self, allocator: std.mem.Allocator) void {
                     }
                     allocator.free(data.members_offsets);
                     allocator.free(data.members_matrix_strides);
+                    allocator.free(data.members_row_major);
                     data.member_names.deinit(allocator);
                 },
                 else => {},
@@ -331,6 +333,7 @@ pub fn dupe(self: *const Self, allocator: std.mem.Allocator) RuntimeError!Self {
                                     .members_type_word = allocator.dupe(SpvWord, s.members_type_word) catch return RuntimeError.OutOfMemory,
                                     .members_offsets = allocator.dupe(?SpvWord, s.members_offsets) catch return RuntimeError.OutOfMemory,
                                     .members_matrix_strides = allocator.dupe(?SpvWord, s.members_matrix_strides) catch return RuntimeError.OutOfMemory,
+                                    .members_row_major = allocator.dupe(bool, s.members_row_major) catch return RuntimeError.OutOfMemory,
                                     .member_names = blk2: {
                                         const member_names = s.member_names.clone(allocator) catch return RuntimeError.OutOfMemory;
                                         for (member_names.items, s.member_names.items) |*new_name, name| {
@@ -479,7 +482,6 @@ pub fn getMemberCounts(self: *const Self) usize {
 
 pub inline fn flushPtr(self: *Self, allocator: std.mem.Allocator) RuntimeError!void {
     if (self.variant) |*variant| switch (variant.*) {
-        .Variable => |*v| try v.value.flushPtr(allocator),
         .AccessChain => |*a| {
             if (!std.mem.allEqual(u8, std.mem.asBytes(&a.value), 0xaa))
                 try a.value.flushPtr(allocator);

src/Runtime.zig

@@ -21,7 +21,6 @@ const Self = @This();
 
 pub const RuntimeError = error{
     Barrier,
-    DivisionByZero,
     InvalidEntryPoint,
     InvalidSpirV,
     InvalidValueType,
@@ -47,6 +46,11 @@ pub const SpecializationEntry = struct {
     size: usize,
 };
 
+pub const WorkgroupMemory = struct {
+    result: SpvWord,
+    bytes: []u8,
+};
+
 pub const Derivative = struct {
     dx: Value,
     dy: Value,
@@ -99,7 +103,7 @@ pub const ImageAPI = struct {
     writeImageInt4: *const fn (driver_image: *anyopaque, dim: spv.SpvDim, x: i32, y: i32, z: i32, pixel: Vec4(u32)) RuntimeError!void,
     sampleImageFloat4: *const fn (driver_image: *anyopaque, driver_sampler: *anyopaque, dim: spv.SpvDim, x: f32, y: f32, z: f32, lod: ?f32, offset: ImageOffset) RuntimeError!Vec4(f32),
     sampleImageInt4: *const fn (driver_image: *anyopaque, driver_sampler: *anyopaque, dim: spv.SpvDim, x: f32, y: f32, z: f32, lod: ?f32, offset: ImageOffset) RuntimeError!Vec4(u32),
-    sampleImageDref: *const fn (driver_image: *anyopaque, driver_sampler: *anyopaque, dim: spv.SpvDim, x: f32, y: f32, z: f32, dref: f32, lod: ?f32, offset: ImageOffset) RuntimeError!f32,
+    sampleImageDref: *const fn (driver_image: *anyopaque, driver_sampler: *anyopaque, dim: spv.SpvDim, x: f32, y: f32, z: f32, w: f32, dref: f32, lod: ?f32, offset: ImageOffset) RuntimeError!f32,
     queryImageSize: *const fn (driver_image: *anyopaque, dim: spv.SpvDim, arrayed: bool, lod: ?i32) RuntimeError!Vec4(u32),
     queryImageLevels: *const fn (driver_image: *anyopaque) RuntimeError!u32,
     queryImageSamples: *const fn (driver_image: *anyopaque) RuntimeError!u32,
@@ -118,14 +122,19 @@ function_stack: std.ArrayList(Function),
 
 current_label: ?SpvWord,
 previous_label: ?SpvWord,
+active_entry_point: ?SpvWord,
 
 specialization_constants: std.AutoHashMapUnmanaged(u32, []const u8),
 derivatives: std.AutoHashMapUnmanaged(SpvWord, Derivative),
+phi_values: std.AutoHashMapUnmanaged(SpvWord, Value),
+helper_invocation: bool,
+derivative_sign_x: f32,
+derivative_sign_y: f32,
 
 image_api: ImageAPI,
 
 pub fn init(allocator: std.mem.Allocator, module: *Module, image_api: ImageAPI) RuntimeError!Self {
-    return .{
+    var self: Self = .{
         .mod = module,
         .it = module.it,
         .results = blk: {
@@ -150,10 +159,18 @@ pub fn init(allocator: std.mem.Allocator, module: *Module, image_api: ImageAPI)
         .function_stack = .empty,
         .current_label = null,
         .previous_label = null,
+        .active_entry_point = null,
         .specialization_constants = .empty,
         .derivatives = .empty,
+        .phi_values = .empty,
+        .helper_invocation = false,
+        .derivative_sign_x = 1.0,
+        .derivative_sign_y = 1.0,
         .image_api = image_api,
     };
+    errdefer self.deinit(allocator);
+    try self.refreshValueLayouts();
+    return self;
 }
 
 pub fn initFrom(allocator: std.mem.Allocator, other: *const Self, image_api: ImageAPI) RuntimeError!Self {
@@ -180,13 +197,19 @@ pub fn initFrom(allocator: std.mem.Allocator, other: *const Self, image_api: Ima
         .function_stack = .empty,
         .current_label = null,
         .previous_label = null,
+        .active_entry_point = other.active_entry_point,
         .specialization_constants = .empty,
         .derivatives = .empty,
+        .phi_values = .empty,
+        .helper_invocation = other.helper_invocation,
+        .derivative_sign_x = other.derivative_sign_x,
+        .derivative_sign_y = other.derivative_sign_y,
         .image_api = image_api,
     };
     errdefer self.deinit(allocator);
 
     try self.copySpecializationConstantsFrom(allocator, other);
+    try self.refreshValueLayouts();
     return self;
 }
 
@@ -207,6 +230,9 @@ pub fn deinit(self: *Self, allocator: std.mem.Allocator) void {
         entry.value_ptr.deinit(allocator);
     }
     self.derivatives.deinit(allocator);
+
+    self.clearPhiValues(allocator);
+    self.phi_values.deinit(allocator);
 }
 
 pub fn addSpecializationInfo(self: *Self, allocator: std.mem.Allocator, entry: SpecializationEntry, data: []const u8) RuntimeError!void {
@@ -225,6 +251,220 @@ pub fn copySpecializationConstantsFrom(self: *Self, allocator: std.mem.Allocator
     }
 }
 
+fn applyValueLayout(results: []Result, value: *Value, type_word: SpvWord) RuntimeError!void {
+    const resolved_type_word = results[type_word].resolveTypeWordOrNull() orelse type_word;
+    const type_data = (try results[resolved_type_word].getConstVariant()).Type;
+
+    switch (value.*) {
+        .Structure => |*structure| switch (type_data) {
+            .Structure => |type_structure| {
+                @memcpy(@constCast(structure.offsets), type_structure.members_offsets);
+                @memcpy(@constCast(structure.matrix_strides), type_structure.members_matrix_strides);
+                @memcpy(@constCast(structure.row_major), type_structure.members_row_major);
+                for (structure.values, type_structure.members_type_word, 0..) |*member_value, member_type_word, member_index| {
+                    if (member_value.* == .RuntimeArray) {
+                        member_value.RuntimeArray.matrix_stride = type_structure.members_matrix_strides[member_index];
+                        member_value.RuntimeArray.row_major = type_structure.members_row_major[member_index];
+                    }
+                    try applyValueLayout(results, member_value, member_type_word);
+                }
+            },
+            else => {},
+        },
+        .Array => |array| switch (type_data) {
+            .Array => |type_array| for (array.values) |*element| {
+                try applyValueLayout(results, element, type_array.components_type_word);
+            },
+            else => {},
+        },
+        .Matrix => |columns| switch (type_data) {
+            .Matrix => |type_matrix| for (columns) |*column| {
+                try applyValueLayout(results, column, type_matrix.column_type_word);
+            },
+            else => {},
+        },
+        .Vector => |elements| switch (type_data) {
+            .Vector => |type_vector| for (elements) |*element| {
+                try applyValueLayout(results, element, type_vector.components_type_word);
+            },
+            else => {},
+        },
+        else => {},
+    }
+}
+
+fn refreshValueLayouts(self: *Self) RuntimeError!void {
+    for (self.results) |*result| {
+        if (result.variant) |*variant| switch (variant.*) {
+            .Variable => |*v| switch (v.storage_class) {
+                .StorageBuffer,
+                .Uniform,
+                .PushConstant,
+                .Workgroup,
+                => try applyValueLayout(self.results, &v.value, v.type_word),
+                else => {},
+            },
+            else => {},
+        };
+    }
+}
+
+pub fn refreshResultValueLayout(self: *Self, result: SpvWord) RuntimeError!void {
+    const type_word = try self.results[result].getValueTypeWord();
+    try applyValueLayout(self.results, try self.results[result].getValue(), type_word);
+}
+
+fn typePlainMemorySize(self: *const Self, type_word: SpvWord) RuntimeError!usize {
+    const resolved_word = self.results[type_word].resolveTypeWordOrNull() orelse type_word;
+    const target_type = (try self.results[resolved_word].getConstVariant()).Type;
+
+    return switch (target_type) {
+        .Array => |a| blk: {
+            const stride: usize = if (a.stride != 0)
+                @intCast(a.stride)
+            else
+                try self.typePlainMemorySize(a.components_type_word);
+            break :blk stride * @as(usize, @intCast(a.member_count));
+        },
+        .RuntimeArray => return RuntimeError.InvalidValueType,
+        .Structure => |s| blk: {
+            var size: usize = 0;
+            for (s.members_type_word, 0..) |member_type_word, i| {
+                const member_offset: usize = @intCast(s.members_offsets[i] orelse size);
+                size = @max(size, member_offset + try self.typePlainMemorySize(member_type_word));
+            }
+            break :blk size;
+        },
+        else => target_type.getSize(self.results),
+    };
+}
+
+pub fn createWorkgroupMemory(self: *Self, allocator: std.mem.Allocator) RuntimeError![]WorkgroupMemory {
+    var count: usize = 0;
+    for (self.results) |result| {
+        if (result.variant) |variant| switch (variant) {
+            .Variable => |v| {
+                if (v.storage_class == .Workgroup) count += 1;
+            },
+            else => {},
+        };
+    }
+
+    const memories = allocator.alloc(WorkgroupMemory, count) catch return RuntimeError.OutOfMemory;
+    errdefer allocator.free(memories);
+
+    var index: usize = 0;
+    for (self.results, 0..) |result, result_id| {
+        if (result.variant) |variant| switch (variant) {
+            .Variable => |v| {
+                if (v.storage_class == .Workgroup) {
+                    const size = try self.typePlainMemorySize(v.type_word);
+                    const bytes = allocator.alloc(u8, size) catch return RuntimeError.OutOfMemory;
+                    @memset(bytes, 0);
+                    memories[index] = .{
+                        .result = @intCast(result_id),
+                        .bytes = bytes,
+                    };
+                    index += 1;
+                }
+            },
+            else => {},
+        };
+    }
+
+    return memories;
+}
+
+pub fn destroyWorkgroupMemory(_: *Self, allocator: std.mem.Allocator, memories: []WorkgroupMemory) void {
+    for (memories) |memory| allocator.free(memory.bytes);
+    allocator.free(memories);
+}
+
+pub fn bindWorkgroupMemory(self: *Self, memories: []const WorkgroupMemory) RuntimeError!void {
+    for (memories) |memory| {
+        _ = try (try self.results[memory.result].getValue()).write(memory.bytes);
+    }
+}
+
+fn clearPhiValues(self: *Self, allocator: std.mem.Allocator) void {
+    var it = self.phi_values.iterator();
+    while (it.next()) |entry| {
+        entry.value_ptr.deinit(allocator);
+    }
+    self.phi_values.clearRetainingCapacity();
+}
+
+fn snapshotPhiValue(self: *Self, allocator: std.mem.Allocator, result_id: SpvWord) RuntimeError!void {
+    if (result_id >= self.results.len) return RuntimeError.InvalidSpirV;
+
+    const value = switch (self.results[result_id].variant orelse return) {
+        .Constant => |*constant| &constant.value,
+        .FunctionParameter => |*parameter| parameter.value_ptr orelse return,
+        else => return,
+    };
+    if (std.meta.activeTag(value.*) == .Pointer) return;
+
+    const snapshot = try value.dupe(allocator);
+    const gop = self.phi_values.getOrPut(allocator, result_id) catch {
+        var tmp = snapshot;
+        tmp.deinit(allocator);
+        return RuntimeError.OutOfMemory;
+    };
+    if (gop.found_existing) gop.value_ptr.deinit(allocator);
+    gop.value_ptr.* = snapshot;
+}
+
+pub fn snapshotPhiValuesForBranch(self: *Self, allocator: std.mem.Allocator, target_source_location: usize) RuntimeError!void {
+    self.clearPhiValues(allocator);
+
+    const predecessor = self.current_label orelse return;
+
+    var index = target_source_location;
+    if (index >= self.it.buffer.len) return RuntimeError.InvalidSpirV;
+
+    const label_opcode_data = self.it.buffer[index];
+    const label_word_count_with_header = (label_opcode_data & (~spv.SpvOpCodeMask)) >> spv.SpvWordCountShift;
+    if (label_word_count_with_header == 0) return RuntimeError.InvalidSpirV;
+    if ((label_opcode_data & spv.SpvOpCodeMask) != @intFromEnum(spv.SpvOp.Label))
+        return RuntimeError.InvalidSpirV;
+    index += label_word_count_with_header;
+
+    const phi_start = index;
+    var phi_end = phi_start;
+    while (index < self.it.buffer.len) {
+        const opcode_data = self.it.buffer[index];
+        const word_count_with_header = (opcode_data & (~spv.SpvOpCodeMask)) >> spv.SpvWordCountShift;
+        if (word_count_with_header == 0) return RuntimeError.InvalidSpirV;
+        const opcode = opcode_data & spv.SpvOpCodeMask;
+        if (opcode != @intFromEnum(spv.SpvOp.Phi)) break;
+        if (index + word_count_with_header > self.it.buffer.len) return RuntimeError.InvalidSpirV;
+
+        const word_count = word_count_with_header - 1;
+        if (word_count < 2 or ((word_count - 2) % 2) != 0) return RuntimeError.InvalidSpirV;
+
+        index += word_count_with_header;
+        phi_end = index;
+    }
+
+    index = phi_start;
+    while (index < phi_end) {
+        const opcode_data = self.it.buffer[index];
+        const word_count_with_header = (opcode_data & (~spv.SpvOpCodeMask)) >> spv.SpvWordCountShift;
+        var operand_index = index + 3;
+        const operand_end = index + word_count_with_header;
+        while (operand_index < operand_end) : (operand_index += 2) {
+            const value_id = self.it.buffer[operand_index];
+            const parent_label_id = self.it.buffer[operand_index + 1];
+            if (parent_label_id == predecessor) {
+                try self.snapshotPhiValue(allocator, value_id);
+                break;
+            }
+        }
+
+        index += word_count_with_header;
+    }
+}
+
 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),
@@ -247,6 +487,8 @@ pub fn setDerivativeFromMemory(self: *Self, allocator: std.mem.Allocator, result
 
     var dx_value = try Value.init(allocator, self.results, target_type, false);
     defer dx_value.deinit(allocator);
+    const memory_size = try dx_value.getPlainMemorySize();
+    if (dx.len < memory_size or dy.len < memory_size) return RuntimeError.OutOfBounds;
     _ = try dx_value.write(dx);
 
     var dy_value = try Value.init(allocator, self.results, target_type, false);
@@ -258,7 +500,13 @@ pub fn setDerivativeFromMemory(self: *Self, allocator: std.mem.Allocator, result
 
 fn getResultTargetTypeWord(self: *const Self, result: SpvWord) RuntimeError!SpvWord {
     return switch ((try self.results[result].getConstVariant()).*) {
-        .Variable => |v| v.type_word,
+        .Variable => |v| switch ((try self.results[v.type_word].getConstVariant()).*) {
+            .Type => |t| switch (t) {
+                .Pointer => |p| p.target,
+                else => v.type_word,
+            },
+            else => return RuntimeError.InvalidSpirV,
+        },
         .Constant => |c| c.type_word,
         .FunctionParameter => |p| p.type_word,
         .AccessChain => |a| a.target,
@@ -274,6 +522,8 @@ pub fn clearDerivative(self: *Self, allocator: std.mem.Allocator, result: SpvWor
 }
 
 pub fn copyDerivative(self: *Self, allocator: std.mem.Allocator, dst: SpvWord, src: SpvWord) RuntimeError!void {
+    if (self.derivatives.count() == 0) return;
+
     if (self.derivatives.get(src)) |derivative| {
         try self.setDerivative(allocator, dst, &derivative.dx, &derivative.dy);
     } else {
@@ -281,22 +531,73 @@ pub fn copyDerivative(self: *Self, allocator: std.mem.Allocator, dst: SpvWord, s
     }
 }
 
+pub fn applySpecializationLayout(self: *Self, allocator: std.mem.Allocator) RuntimeError!void {
+    self.reset();
+    try self.applySpecializationConstants(allocator);
+    try self.applySpecializationDependentLayout(allocator);
+}
+
+pub fn applySpecializationInvocationLayout(self: *Self, allocator: std.mem.Allocator) RuntimeError!void {
+    self.reset();
+    if (self.specialization_constants.count() != 0)
+        try self.applySpecializationConstants(allocator);
+    try self.applySpecializationDependentLayout(allocator);
+}
+
+fn applySpecializationConstants(self: *Self, allocator: std.mem.Allocator) RuntimeError!void {
+    try self.pass(allocator, .initMany(&.{
+        .SpecConstantTrue,
+        .SpecConstantFalse,
+        .SpecConstantComposite,
+        .ConstantNull,
+        .SpecConstant,
+        .SpecConstantOp,
+    }));
+}
+
+fn applySpecializationDependentLayout(self: *Self, allocator: std.mem.Allocator) RuntimeError!void {
+    self.reset();
+    try self.pass(allocator, .initMany(&.{
+        .TypeArray,
+        .Variable,
+    }));
+}
+
 pub fn getEntryPointByName(self: *const Self, name: []const u8) RuntimeError!SpvWord {
-    for (self.mod.entry_points.items, 0..) |entry_point, i| {
+    if (self.active_entry_point) |entry_point| {
-        if (blk: {
+        if (entryPointNameMatches(self.mod.entry_points.items[entry_point].name, name))
-            // Not using std.mem.eql as entry point names may have longer size than their content
+            return entry_point;
-            for (0..@min(name.len, entry_point.name.len)) |j| {
-                if (name[j] != entry_point.name[j])
-                    break :blk false;
     }
-            if (entry_point.name.len != name.len and entry_point.name[name.len] != 0)
+
-                break :blk false;
+    for (self.mod.entry_points.items, 0..) |entry_point, i| {
-            break :blk true;
+        if (entryPointNameMatches(entry_point.name, name)) return @intCast(i);
-        }) return @intCast(i);
     }
     return RuntimeError.NotFound;
 }
 
+pub fn getEntryPointByNameAndExecutionModel(self: *const Self, name: []const u8, execution_model: spv.SpvExecutionModel) RuntimeError!SpvWord {
+    for (self.mod.entry_points.items, 0..) |entry_point, i| {
+        if (entry_point.exec_model == execution_model and entryPointNameMatches(entry_point.name, name))
+            return @intCast(i);
+    }
+    return RuntimeError.NotFound;
+}
+
+fn entryPointNameMatches(entry_point_name: []const u8, name: []const u8) bool {
+    // Not using std.mem.eql as entry point names may have longer size than their content.
+    for (0..@min(name.len, entry_point_name.len)) |j| {
+        if (name[j] != entry_point_name[j])
+            return false;
+    }
+    return entry_point_name.len == name.len or entry_point_name[name.len] == 0;
+}
+
+pub fn selectEntryPoint(self: *Self, entry_point_index: SpvWord) RuntimeError!void {
+    if (entry_point_index >= self.mod.entry_points.items.len)
+        return RuntimeError.InvalidEntryPoint;
+    self.active_entry_point = entry_point_index;
+}
+
 pub fn getResultByName(self: *const Self, name: []const u8) RuntimeError!SpvWord {
     for (self.results, 0..) |result, i| {
         if (result.name) |result_name| {
@@ -319,21 +620,113 @@ pub inline fn getResultByLocation(self: *const Self, location: SpvWord, kind: Lo
 }
 
 pub fn getResultByLocationComponent(self: *const Self, location: SpvWord, component: SpvWord, kind: LocationKind) RuntimeError!SpvWord {
-    switch (kind) {
+    const locations = switch (kind) {
-        .input => if (location < self.mod.input_locations.len and component < 4 and self.mod.input_locations[location][component] != 0) {
+        .input => &self.mod.input_locations,
-            return self.mod.input_locations[location][component];
+        .output => &self.mod.output_locations,
-        },
+    };
-        .output => if (location < self.mod.output_locations.len and component < 4 and self.mod.output_locations[location][component] != 0) {
+    if (location >= locations.len or component >= 4)
-            return self.mod.output_locations[location][component];
+        return RuntimeError.NotFound;
-        },
+
+    const result = locations[location][component];
+    if (result != 0 and self.resultIsInActiveInterface(result))
+        return result;
+
+    if (self.active_entry_point) |entry_point_index| {
+        const entry_point = self.mod.entry_points.items[entry_point_index];
+        for (entry_point.globals) |global| {
+            if (global >= self.results.len)
+                continue;
+
+            const variant = self.results[global].variant orelse continue;
+            const variable = switch (variant) {
+                .Variable => |v| v,
+                else => continue,
+            };
+            const storage_class_matches = switch (kind) {
+                .input => variable.storage_class == .Input,
+                .output => variable.storage_class == .Output,
+            };
+            if (!storage_class_matches)
+                continue;
+
+            for (self.results[global].decorations.items) |decoration| {
+                if (decoration.rtype == .Location and
+                    decoration.literal_1 == location and
+                    self.resultComponent(global) == component)
+                {
+                    return global;
+                }
+            }
+        }
     }
     return RuntimeError.NotFound;
 }
 
-pub fn getResultPrimitiveType(self: *const Self, result: SpvWord) RuntimeError!PrimitiveType {
+fn resultIsInActiveInterface(self: *const Self, result: SpvWord) bool {
+    const entry_point_index = self.active_entry_point orelse return true;
+    const global = self.resultGlobal(result) orelse return false;
+    return std.mem.indexOfScalar(SpvWord, self.mod.entry_points.items[entry_point_index].globals, global) != null;
+}
+
+fn resultGlobal(self: *const Self, result: SpvWord) ?SpvWord {
     if (result >= self.results.len)
-        return RuntimeError.OutOfBounds;
+        return null;
-    return (try self.results[result].getConstValue()).resolvePrimitiveType();
+
+    const variant = self.results[result].variant orelse return result;
+    return switch (variant) {
+        .AccessChain => |access_chain| access_chain.base,
+        else => result,
+    };
+}
+
+fn resultComponent(self: *const Self, result: SpvWord) SpvWord {
+    if (result >= self.results.len)
+        return 0;
+
+    for (self.results[result].decorations.items) |decoration| {
+        if (decoration.rtype == .Component)
+            return decoration.literal_1;
+    }
+    return 0;
+}
+
+pub fn getWorkgroupSize(self: *Self, allocator: std.mem.Allocator) RuntimeError!?@Vector(3, u32) {
+    try self.pass(allocator, .initMany(&.{
+        .SpecConstantTrue,
+        .SpecConstantFalse,
+        .SpecConstantComposite,
+        .ConstantNull,
+        .SpecConstant,
+        .SpecConstantOp,
+    }));
+
+    for (self.results) |*result| {
+        for (result.decorations.items) |decoration| {
+            if (decoration.rtype != .BuiltIn or decoration.literal_1 != @intFromEnum(spv.SpvBuiltIn.WorkgroupSize))
+                continue;
+
+            const value = try result.getValue();
+            return switch (value.*) {
+                .Vector3u32 => |v| v,
+                .Vector => |values| blk: {
+                    if (values.len != 3)
+                        return RuntimeError.InvalidValueType;
+
+                    var result_value = @Vector(3, u32){ 0, 0, 0 };
+                    inline for (0..3) |i| {
+                        result_value[i] = switch (values[i]) {
+                            .Int => |int| int.value.uint32,
+                            else => return RuntimeError.InvalidValueType,
+                        };
+                    }
+                    break :blk result_value;
+                },
+                else => return RuntimeError.InvalidValueType,
+            };
+        }
+    }
+
+    return null;
 }
 
 pub fn dumpResultsTable(self: *Self, allocator: std.mem.Allocator, writer: *std.Io.Writer) RuntimeError!void {
@@ -359,13 +752,8 @@ pub fn beginEntryPoint(self: *Self, allocator: std.mem.Allocator, entry_point_in
         return RuntimeError.InvalidEntryPoint;
 
     // Spec constants pass
-    try self.pass(allocator, .initMany(&.{
+    if (self.specialization_constants.count() != 0)
-        .SpecConstantTrue,
+        try self.applySpecializationConstants(allocator);
-        .SpecConstantFalse,
-        .SpecConstantComposite,
-        .SpecConstant,
-        .SpecConstantOp,
-    }));
 
     {
         const entry_point_desc = &self.mod.entry_points.items[entry_point_index];
@@ -439,12 +827,16 @@ pub fn populatePushConstants(self: *Self, blob: []const u8) RuntimeError!void {
         if (variable.storage_class != .PushConstant)
             continue;
         _ = try variable.value.write(blob);
+        variable.value.clearExternalData();
     }
 }
 
 pub fn writeDescriptorSet(self: *const Self, input: []const u8, set: SpvWord, binding: SpvWord, descriptor_index: SpvWord) RuntimeError!void {
     const result = self.mod.getBindingResult(set, binding) orelse return RuntimeError.NotFound;
-    const value = &(self.results[result].variant orelse return).Variable.value;
+    const value = if (self.results[result].variant) |*variant| switch (variant.*) {
+        .Variable => |*variable| &variable.value,
+        else => return RuntimeError.InvalidSpirV,
+    } else return;
     switch (value.*) {
         .Array => |arr| {
             if (descriptor_index >= arr.values.len)
@@ -485,24 +877,132 @@ fn readResultValue(self: *const Self, output: []u8, result: SpvWord) RuntimeErro
     }
 }
 
-fn writeResultValue(self: *const Self, input: []const u8, result: SpvWord) RuntimeError!void {
+fn readConstantIndex(self: *const Self, result: SpvWord) RuntimeError!usize {
-    if (self.results[result].variant) |*variant| {
+    const value = try self.results[result].getConstValue();
-        switch (variant.*) {
+    return switch (value.*) {
-            .Variable => |*v| _ = try v.value.write(input),
+        .Int => |i| switch (i.bit_count) {
+            8 => if (i.is_signed) std.math.cast(usize, i.value.sint8) orelse RuntimeError.OutOfBounds else @as(usize, i.value.uint8),
+            16 => if (i.is_signed) std.math.cast(usize, i.value.sint16) orelse RuntimeError.OutOfBounds else @as(usize, i.value.uint16),
+            32 => if (i.is_signed) std.math.cast(usize, i.value.sint32) orelse RuntimeError.OutOfBounds else @as(usize, i.value.uint32),
+            64 => if (i.is_signed) std.math.cast(usize, i.value.sint64) orelse RuntimeError.OutOfBounds else std.math.cast(usize, i.value.uint64) orelse RuntimeError.OutOfBounds,
+            else => RuntimeError.InvalidSpirV,
+        },
+        else => RuntimeError.InvalidSpirV,
+    };
+}
+
+fn accessChainPrefixValue(self: *const Self, result: SpvWord, prefix_len: usize) RuntimeError!*Value {
+    const access_chain = switch ((self.results[result].variant orelse return RuntimeError.InvalidSpirV)) {
+        .AccessChain => |*a| a,
+        else => return RuntimeError.InvalidSpirV,
+    };
+    if (prefix_len > access_chain.indexes.len) return RuntimeError.OutOfBounds;
+
+    var value = switch ((self.results[access_chain.base].variant orelse return RuntimeError.InvalidSpirV)) {
+        .Variable => |*v| &v.value,
         .AccessChain => |*a| switch (a.value) {
             .Pointer => |ptr| switch (ptr.ptr) {
-                    .common => |value_ptr| _ = try value_ptr.write(input),
+                .common => |value_ptr| value_ptr,
+                else => return RuntimeError.InvalidSpirV,
+            },
+            else => &a.value,
+        },
+        else => return RuntimeError.InvalidSpirV,
+    };
+
+    for (access_chain.indexes[0..prefix_len]) |index_id| {
+        const index = try self.readConstantIndex(index_id);
+        value = switch (value.*) {
+            .Vector, .Matrix => |values| blk: {
+                if (index >= values.len) return RuntimeError.OutOfBounds;
+                break :blk &values[index];
+            },
+            .Array => |arr| blk: {
+                if (index >= arr.values.len) return RuntimeError.OutOfBounds;
+                break :blk &arr.values[index];
+            },
+            .Structure => |structure| blk: {
+                if (index >= structure.values.len) return RuntimeError.OutOfBounds;
+                break :blk &structure.values[index];
+            },
+            else => return RuntimeError.InvalidValueType,
+        };
+    }
+
+    return value;
+}
+
+fn writeResultValue(self: *const Self, allocator: std.mem.Allocator, input: []const u8, result: SpvWord) RuntimeError!void {
+    if (self.results[result].variant) |*variant| {
+        switch (variant.*) {
+            .Variable => |*v| switch (v.value) {
+                .Pointer => |*ptr| {
+                    if (ptr.owns_uniform_backing_value) {
+                        if (ptr.uniform_backing_value) |value_ptr| {
+                            _ = try value_ptr.write(input);
+                            return;
+                        }
+                    }
+
+                    const target_type = switch ((try self.results[v.type_word].getConstVariant()).*) {
+                        .Type => |t| switch (t) {
+                            .Pointer => |p| p.target,
+                            else => return RuntimeError.InvalidSpirV,
+                        },
+                        else => return RuntimeError.InvalidSpirV,
+                    };
+                    const value_ptr = allocator.create(Value) catch return RuntimeError.OutOfMemory;
+                    errdefer allocator.destroy(value_ptr);
+
+                    value_ptr.* = try Value.init(allocator, self.results, target_type, false);
+                    errdefer value_ptr.deinit(allocator);
+
+                    _ = try value_ptr.write(input);
+                    ptr.* = .{
+                        .ptr = .{ .common = value_ptr },
+                        .uniform_backing_value = value_ptr,
+                        .owns_uniform_backing_value = true,
+                    };
+                },
+                else => _ = try v.value.write(input),
+            },
+            .AccessChain => |*a| switch (a.value) {
+                .Pointer => |ptr| switch (ptr.ptr) {
+                    .common => |value_ptr| {
+                        const value_size = try value_ptr.getPlainMemorySize();
+                        if (a.indexes.len > 1 and input.len > value_size) {
+                            const parent = try self.accessChainPrefixValue(result, a.indexes.len - 1);
+                            _ = try parent.write(input);
+                        } else {
+                            _ = try value_ptr.write(input);
+                        }
+                    },
                     .f32_ptr => |value_ptr| {
                         if (input.len < @sizeOf(f32)) return RuntimeError.OutOfBounds;
+                        if (a.indexes.len > 1 and input.len > @sizeOf(f32)) {
+                            const parent = try self.accessChainPrefixValue(result, a.indexes.len - 1);
+                            _ = try parent.write(input);
+                        } else {
                             std.mem.copyForwards(u8, std.mem.asBytes(value_ptr), input[0..@sizeOf(f32)]);
+                        }
                     },
                     .i32_ptr => |value_ptr| {
                         if (input.len < @sizeOf(i32)) return RuntimeError.OutOfBounds;
+                        if (a.indexes.len > 1 and input.len > @sizeOf(i32)) {
+                            const parent = try self.accessChainPrefixValue(result, a.indexes.len - 1);
+                            _ = try parent.write(input);
+                        } else {
                             std.mem.copyForwards(u8, std.mem.asBytes(value_ptr), input[0..@sizeOf(i32)]);
+                        }
                     },
                     .u32_ptr => |value_ptr| {
                         if (input.len < @sizeOf(u32)) return RuntimeError.OutOfBounds;
+                        if (a.indexes.len > 1 and input.len > @sizeOf(u32)) {
+                            const parent = try self.accessChainPrefixValue(result, a.indexes.len - 1);
+                            _ = try parent.write(input);
+                        } else {
                             std.mem.copyForwards(u8, std.mem.asBytes(value_ptr), input[0..@sizeOf(u32)]);
+                        }
                     },
                 },
                 else => _ = try a.value.write(input),
@@ -516,32 +1016,96 @@ fn writeResultValue(self: *const Self, input: []const u8, result: SpvWord) Runti
 
 const InputLocationTarget = struct {
     result: SpvWord,
+    struct_member: ?usize = null,
+    array_element: ?usize = null,
     matrix_column: ?usize = null,
 };
 
 fn resolveInputLocationTarget(self: *const Self, location: SpvWord) RuntimeError!InputLocationTarget {
-    if (location < self.mod.input_locations.len and self.mod.input_locations[location][0] != 0) {
+    if (self.getResultByLocationComponent(location, 0, .input)) |result| {
-        const result = self.mod.input_locations[location][0];
+        if (self.results[result].variant == null)
+            return RuntimeError.NotFound;
+
         const value = try self.results[result].getConstValue();
         switch (value.*) {
+            .Array => |arr| {
+                if (arr.values.len == 0) return RuntimeError.OutOfBounds;
+                return .{ .result = result };
+            },
             .Matrix => return .{ .result = result, .matrix_column = 0 },
             else => return .{ .result = result },
         }
+    } else |err| switch (err) {
+        RuntimeError.NotFound => {},
+        else => return err,
+    }
+
+    for (self.results, 0..) |*result, id| {
+        const variant = result.variant orelse continue;
+        const variable = switch (variant) {
+            .Variable => |v| v,
+            else => continue,
+        };
+        if (variable.storage_class != .Input) continue;
+
+        const type_word = switch ((self.results[variable.type_word].variant orelse continue)) {
+            .Type => |t| switch (t) {
+                .Pointer => |ptr| ptr.target,
+                else => variable.type_word,
+            },
+            else => continue,
+        };
+        const type_result = &self.results[type_word];
+        const type_variant = type_result.variant orelse continue;
+        switch (type_variant) {
+            .Type => |t| switch (t) {
+                .Structure => {
+                    for (type_result.decorations.items) |decoration| {
+                        if (decoration.rtype == .Location and decoration.literal_1 == location) {
+                            return .{
+                                .result = @intCast(id),
+                                .struct_member = @intCast(decoration.index),
+                            };
+                        }
+                    }
+                },
+                else => {},
+            },
+            else => {},
+        }
     }
 
     var base_location = location;
     while (base_location > 0) {
         base_location -= 1;
 
-        const result = if (base_location < self.mod.input_locations.len)
+        const result = self.getResultByLocationComponent(base_location, 0, .input) catch |err| switch (err) {
-            self.mod.input_locations[base_location][0]
+            RuntimeError.NotFound => continue,
-        else
+            else => return err,
-            0;
+        };
-        if (result == 0) continue;
 
         const location_offset: usize = @intCast(location - base_location);
         const value = try self.results[result].getConstValue();
         switch (value.*) {
+            .Array => |arr| {
+                if (arr.values.len == 0) continue;
+
+                const element_locations = switch (arr.values[0]) {
+                    .Matrix => |columns| columns.len,
+                    else => 1,
+                };
+                if (element_locations == 0) continue;
+
+                const element_index = location_offset / element_locations;
+                if (element_index >= arr.values.len) continue;
+
+                const element_location = location_offset % element_locations;
+                return .{
+                    .result = result,
+                    .array_element = element_index,
+                    .matrix_column = if (std.meta.activeTag(arr.values[element_index]) == .Matrix) element_location else null,
+                };
+            },
             .Matrix => |columns| {
                 if (location_offset < columns.len) {
                     return .{
@@ -564,8 +1128,28 @@ fn getInputLocationTargetValue(self: *const Self, target: InputLocationTarget) R
         else => return RuntimeError.InvalidSpirV,
     };
 
-    if (target.matrix_column) |column| {
+    const element_value = if (target.array_element) |element| blk: {
         switch (value.*) {
+            .Array => |arr| {
+                if (element >= arr.values.len) return RuntimeError.OutOfBounds;
+                break :blk &arr.values[element];
+            },
+            else => return RuntimeError.InvalidValueType,
+        }
+    } else value;
+
+    const member_value = if (target.struct_member) |member| blk: {
+        switch (element_value.*) {
+            .Structure => |structure| {
+                if (member >= structure.values.len) return RuntimeError.OutOfBounds;
+                break :blk &structure.values[member];
+            },
+            else => return RuntimeError.InvalidValueType,
+        }
+    } else element_value;
+
+    if (target.matrix_column) |column| {
+        switch (member_value.*) {
             .Matrix => |columns| {
                 if (column >= columns.len) return RuntimeError.OutOfBounds;
                 return &columns[column];
@@ -574,7 +1158,7 @@ fn getInputLocationTargetValue(self: *const Self, target: InputLocationTarget) R
         }
     }
 
-    return value;
+    return member_value;
 }
 
 pub fn readOutput(self: *const Self, output: []u8, result: SpvWord) RuntimeError!void {
@@ -588,17 +1172,17 @@ pub fn readOutput(self: *const Self, output: []u8, result: SpvWord) RuntimeError
 }
 
 pub fn readBuiltIn(self: *const Self, output: []u8, builtin: spv.SpvBuiltIn) RuntimeError!void {
-    if (self.mod.builtins.get(builtin)) |result| {
+    if (self.getBuiltinResult(builtin)) |result| {
         try self.readResultValue(output, result);
     } else {
         return RuntimeError.NotFound;
     }
 }
 
-pub fn writeInput(self: *const Self, input: []const u8, result: SpvWord) RuntimeError!void {
+pub fn writeInput(self: *const Self, allocator: std.mem.Allocator, input: []const u8, result: SpvWord) RuntimeError!void {
     for (&self.mod.input_locations) |*location| {
         if (std.mem.indexOfScalar(SpvWord, location, result)) |_| {
-            try self.writeResultValue(input, result);
+            try self.writeResultValue(allocator, input, result);
             if (self.results[result].variant) |*variant| switch (variant.*) {
                 .Variable => |*v| v.value.clearExternalData(),
                 .AccessChain => |*a| a.value.clearExternalData(),
@@ -622,14 +1206,35 @@ pub fn writeInputLocation(self: *const Self, input: []const u8, location: SpvWor
     value.clearExternalData();
 }
 
-pub fn writeBuiltIn(self: *const Self, input: []const u8, builtin: spv.SpvBuiltIn) RuntimeError!void {
+pub fn writeBuiltIn(self: *const Self, allocator: std.mem.Allocator, input: []const u8, builtin: spv.SpvBuiltIn) RuntimeError!void {
-    if (self.mod.builtins.get(builtin)) |result| {
+    if (self.getBuiltinResult(builtin)) |result| {
-        try self.writeResultValue(input, result);
+        try self.writeResultValue(allocator, input, result);
     } else {
         return RuntimeError.NotFound;
     }
 }
 
+pub fn getBuiltinResult(self: *const Self, builtin: spv.SpvBuiltIn) ?SpvWord {
+    if (self.mod.builtins.get(builtin)) |result| {
+        if (self.resultIsInActiveInterface(result))
+            return result;
+    }
+
+    const entry_point_index = self.active_entry_point orelse return null;
+    const entry_point = self.mod.entry_points.items[entry_point_index];
+    for (entry_point.globals) |global| {
+        if (global >= self.results.len)
+            continue;
+
+        for (self.results[global].decorations.items) |decoration| {
+            if (decoration.rtype == .BuiltIn and decoration.literal_1 == @intFromEnum(builtin))
+                return global;
+        }
+    }
+
+    return null;
+}
+
 pub fn flushDescriptorSets(self: *const Self, allocator: std.mem.Allocator) RuntimeError!void {
     for (self.results) |*result| {
         try result.flushPtr(allocator);
@@ -637,8 +1242,19 @@ pub fn flushDescriptorSets(self: *const Self, allocator: std.mem.Allocator) Runt
 }
 
 pub fn getResultMemorySize(self: *const Self, result: SpvWord) RuntimeError!usize {
-    const value = try self.results[result].getConstValue();
+    const variant = self.results[result].variant orelse return RuntimeError.InvalidSpirV;
-    return value.getPlainMemorySize();
+    return switch (variant) {
+        .AccessChain => |a| switch (a.value) {
+            .Pointer => |ptr| switch (ptr.ptr) {
+                .common => |value_ptr| value_ptr.getPlainMemorySize(),
+                .f32_ptr => @sizeOf(f32),
+                .i32_ptr => @sizeOf(i32),
+                .u32_ptr => @sizeOf(u32),
+            },
+            else => a.value.getPlainMemorySize(),
+        },
+        else => (try self.results[result].getConstValue()).getPlainMemorySize(),
+    };
 }
 
 pub fn hasResultDecoration(self: *const Self, result: SpvWord, decoration: spv.SpvDecoration) bool {
@@ -650,11 +1266,15 @@ pub fn hasResultDecoration(self: *const Self, result: SpvWord, decoration: spv.S
 }
 
 pub fn resetInvocation(self: *Self, allocator: std.mem.Allocator) void {
+    self.derivative_sign_x = 1.0;
+    self.derivative_sign_y = 1.0;
+
     var derivatives = self.derivatives.iterator();
     while (derivatives.next()) |entry| {
         entry.value_ptr.deinit(allocator);
     }
     self.derivatives.clearRetainingCapacity();
+    self.clearPhiValues(allocator);
 
     for (self.results) |*result| {
         if (result.variant) |*variant| {
@@ -680,10 +1300,12 @@ pub fn resetInvocation(self: *Self, allocator: std.mem.Allocator) void {
 }
 
 fn reset(self: *Self) void {
-    self.it = self.mod.it;
+    self.it = WordIterator.init(self.mod.code);
+    self.it.index = 5;
     self.function_stack.clearRetainingCapacity();
     self.current_parameter_index = 0;
     self.current_function = null;
     self.current_label = null;
     self.previous_label = null;
+    self.helper_invocation = false;
 }

src/Value.zig

@@ -78,6 +78,8 @@ pub const Value = union(Type) {
         type_word: SpvWord,
         stride: SpvWord,
         data: []u8,
+        matrix_stride: ?SpvWord = null,
+        row_major: bool = false,
 
         pub inline fn createValueFromIndex(self: *const @This(), allocator: std.mem.Allocator, results: []const Result, index: usize) RuntimeError!*Value {
             const offset = try self.getCheckedOffsetOfIndex(index);
@@ -86,7 +88,34 @@ pub const Value = union(Type) {
 
             value.* = try Value.init(allocator, results, self.type_word, false);
             errdefer value.deinit(allocator);
-            _ = try value.write(self.data[offset .. offset + self.stride]);
+            _ = if (self.matrix_stride) |matrix_stride|
+                try value.writeWithMatrixLayout(self.data[offset .. offset + self.stride], matrix_stride, self.row_major)
+            else
+                try value.write(self.data[offset .. offset + self.stride]);
+
+            return value;
+        }
+
+        pub inline fn createRobustValueFromIndex(self: *const @This(), allocator: std.mem.Allocator, results: []const Result, index: usize) RuntimeError!*Value {
+            const value = allocator.create(Value) catch return RuntimeError.OutOfMemory;
+            errdefer allocator.destroy(value);
+
+            value.* = try Value.init(allocator, results, self.type_word, false);
+            errdefer value.deinit(allocator);
+
+            const bytes = allocator.alloc(u8, self.stride) catch return RuntimeError.OutOfMemory;
+            defer allocator.free(bytes);
+            @memset(bytes, 0);
+
+            if (self.getRobustOffsetOfIndex(index)) |offset| {
+                const available = @min(self.stride, self.data.len - offset);
+                @memcpy(bytes[0..available], self.data[offset .. offset + available]);
+            }
+
+            _ = if (self.matrix_stride) |matrix_stride|
+                try value.writeWithMatrixLayout(bytes, matrix_stride, self.row_major)
+            else
+                try value.write(bytes);
 
             return value;
         }
@@ -95,7 +124,10 @@ pub const Value = union(Type) {
             const offset = try self.getCheckedOffsetOfIndex(index);
             var value = try Value.init(allocator, results, self.type_word, false);
             errdefer value.deinit(allocator);
-            _ = try value.write(self.data[offset .. offset + self.stride]);
+            _ = if (self.matrix_stride) |matrix_stride|
+                try value.writeWithMatrixLayout(self.data[offset .. offset + self.stride], matrix_stride, self.row_major)
+            else
+                try value.write(self.data[offset .. offset + self.stride]);
             return value;
         }
 
@@ -103,6 +135,12 @@ pub const Value = union(Type) {
             return self.stride * index;
         }
 
+        pub inline fn getRobustOffsetOfIndex(self: *const @This(), index: usize) ?usize {
+            if (self.stride == 0) return null;
+            const offset = std.math.mul(usize, self.stride, index) catch return null;
+            return if (offset < self.data.len) offset else null;
+        }
+
         pub inline fn getCheckedOffsetOfIndex(self: *const @This(), index: usize) RuntimeError!usize {
             const offset = self.getOffsetOfIndex(index);
             if (self.stride == 0 or offset > self.data.len or self.stride > self.data.len - offset)
@@ -119,6 +157,7 @@ pub const Value = union(Type) {
         external_data: ?[]u8,
         offsets: []const ?SpvWord,
         matrix_strides: []const ?SpvWord,
+        row_major: []const bool,
         values: []Self,
     },
     Function: noreturn,
@@ -153,6 +192,8 @@ pub const Value = union(Type) {
         /// corresponds to this pointer. For a pointer to struct member N this
         /// starts at the member offset, not at the containing struct.
         uniform_slice_window: ?[]u8 = null,
+        uniform_root_window: ?[]u8 = null,
+        uniform_window_offset: usize = 0,
 
         /// Heap-owned value that backs a pointer into a materialized runtime
         /// array element. This may differ from ptr.common when the pointer is
@@ -160,6 +201,7 @@ pub const Value = union(Type) {
         uniform_backing_value: ?*Self = null,
         owns_uniform_backing_value: bool = false,
         matrix_stride: ?SpvWord = null,
+        matrix_row_major: bool = false,
     },
 
     pub inline fn getCompositeDataOrNull(self: *const Self) ?[]Self {
@@ -254,17 +296,6 @@ pub const Value = union(Type) {
                     break :blk self;
                 },
                 .Array => |a| blk: {
-                    // If an array is in externally visible storage we treat it as a runtime array
-                    if (is_externally_visible) {
-                        break :blk .{
-                            .RuntimeArray = .{
-                                .type_word = a.components_type_word,
-                                .stride = a.stride,
-                                .data = &.{},
-                            },
-                        };
-                    }
-
                     const self: Self = .{
                         .Array = .{
                             .stride = a.stride,
@@ -284,6 +315,7 @@ pub const Value = union(Type) {
                             .external_data = null,
                             .offsets = allocator.dupe(?SpvWord, s.members_offsets) catch return RuntimeError.OutOfMemory,
                             .matrix_strides = allocator.dupe(?SpvWord, s.members_matrix_strides) catch return RuntimeError.OutOfMemory,
+                            .row_major = allocator.dupe(bool, s.members_row_major) catch return RuntimeError.OutOfMemory,
                             .values = allocator.alloc(Self, member_count) catch return RuntimeError.OutOfMemory,
                         },
                     };
@@ -299,6 +331,8 @@ pub const Value = union(Type) {
                         .type_word = a.components_type_word,
                         .stride = a.stride,
                         .data = &.{},
+                        .matrix_stride = null,
+                        .row_major = false,
                     },
                 },
                 .Image => .{
@@ -369,10 +403,24 @@ pub const Value = union(Type) {
                         .external_data = s.external_data,
                         .offsets = allocator.dupe(?SpvWord, s.offsets) catch return RuntimeError.OutOfMemory,
                         .matrix_strides = allocator.dupe(?SpvWord, s.matrix_strides) catch return RuntimeError.OutOfMemory,
+                        .row_major = allocator.dupe(bool, s.row_major) catch return RuntimeError.OutOfMemory,
                         .values = values,
                     };
                 },
             },
+            .Pointer => |p| .{
+                .Pointer = .{
+                    .ptr = p.ptr,
+                    .image_texel = p.image_texel,
+                    .uniform_slice_window = p.uniform_slice_window,
+                    .uniform_root_window = p.uniform_root_window,
+                    .uniform_window_offset = p.uniform_window_offset,
+                    .uniform_backing_value = p.uniform_backing_value,
+                    .owns_uniform_backing_value = false,
+                    .matrix_stride = p.matrix_stride,
+                    .matrix_row_major = p.matrix_row_major,
+                },
+            },
             else => self.*,
         };
     }
@@ -463,7 +511,7 @@ pub const Value = union(Type) {
                 var offset: usize = 0;
                 for (arr.values) |v| {
                     _ = try v.read(output[offset..]);
-                    offset += arr.stride;
+                    offset += if (arr.stride != 0) arr.stride else try v.getPlainMemorySize();
                 }
                 return offset;
             },
@@ -472,7 +520,15 @@ pub const Value = union(Type) {
                 for (s.values, 0..) |v, i| {
                     const member_offset: usize = @intCast(s.offsets[i] orelse end_offset);
                     if (member_offset > output.len) return RuntimeError.OutOfBounds;
-                    const read_size = try v.read(output[member_offset..]);
+                    const member_output = if (v == .RuntimeArray and i + 1 < s.values.len and s.offsets[i + 1] != null) blk: {
+                        const next_offset: usize = @intCast(s.offsets[i + 1].?);
+                        if (next_offset < member_offset or next_offset > output.len) return RuntimeError.OutOfBounds;
+                        break :blk output[member_offset..next_offset];
+                    } else output[member_offset..];
+                    const read_size = if (s.matrix_strides[i]) |matrix_stride|
+                        try v.readWithMatrixLayout(member_output, matrix_stride, s.row_major[i])
+                    else
+                        try v.read(member_output);
                     end_offset = @max(end_offset, member_offset + read_size);
                 }
                 return end_offset;
@@ -483,7 +539,90 @@ pub const Value = union(Type) {
         return 0;
     }
 
-    pub fn readMatrixWithStride(self: *const Self, output: []u8, matrix_stride: SpvWord) RuntimeError!usize {
+    fn readVectorLane(self: *const Self, lane_index: usize, output: []u8) RuntimeError!usize {
+        switch (self.*) {
+            .Vector => |values| {
+                if (lane_index >= values.len) return RuntimeError.OutOfBounds;
+                return values[lane_index].read(output);
+            },
+            .Vector2f32, .Vector3f32, .Vector4f32 => {
+                const value = try readLane(.Float, 32, self, lane_index);
+                if (output.len < @sizeOf(f32)) return RuntimeError.OutOfBounds;
+                @memcpy(output[0..@sizeOf(f32)], std.mem.asBytes(&value));
+                return @sizeOf(f32);
+            },
+            .Vector2i32, .Vector3i32, .Vector4i32 => {
+                const value = try readLane(.SInt, 32, self, lane_index);
+                if (output.len < @sizeOf(i32)) return RuntimeError.OutOfBounds;
+                @memcpy(output[0..@sizeOf(i32)], std.mem.asBytes(&value));
+                return @sizeOf(i32);
+            },
+            .Vector2u32, .Vector3u32, .Vector4u32 => {
+                const value = try readLane(.UInt, 32, self, lane_index);
+                if (output.len < @sizeOf(u32)) return RuntimeError.OutOfBounds;
+                @memcpy(output[0..@sizeOf(u32)], std.mem.asBytes(&value));
+                return @sizeOf(u32);
+            },
+            else => return RuntimeError.InvalidValueType,
+        }
+    }
+
+    fn writeVectorLane(self: *Self, lane_index: usize, input: []const u8) RuntimeError!usize {
+        switch (self.*) {
+            .Vector => |values| {
+                if (lane_index >= values.len) return RuntimeError.OutOfBounds;
+                return values[lane_index].write(input);
+            },
+            .Vector2f32, .Vector3f32, .Vector4f32 => {
+                const value: f32 = if (input.len >= @sizeOf(f32)) std.mem.bytesToValue(f32, input[0..@sizeOf(f32)]) else 0;
+                try writeLane(.Float, 32, self, lane_index, value);
+                return @sizeOf(f32);
+            },
+            .Vector2i32, .Vector3i32, .Vector4i32 => {
+                const value: i32 = if (input.len >= @sizeOf(i32)) std.mem.bytesToValue(i32, input[0..@sizeOf(i32)]) else 0;
+                try writeLane(.SInt, 32, self, lane_index, value);
+                return @sizeOf(i32);
+            },
+            .Vector2u32, .Vector3u32, .Vector4u32 => {
+                const value: u32 = if (input.len >= @sizeOf(u32)) std.mem.bytesToValue(u32, input[0..@sizeOf(u32)]) else 0;
+                try writeLane(.UInt, 32, self, lane_index, value);
+                return @sizeOf(u32);
+            },
+            else => return RuntimeError.InvalidValueType,
+        }
+    }
+
+    pub fn readVectorWithStride(self: *const Self, output: []u8, stride: SpvWord) RuntimeError!usize {
+        const lane_count: usize = @intCast(try self.resolveLaneCount());
+        if (lane_count == 0) return 0;
+
+        const stride_usize: usize = @intCast(stride);
+        const scalar_size = @divExact(try self.getPlainMemorySize(), lane_count);
+        const vector_size = (lane_count - 1) * stride_usize + scalar_size;
+        if (output.len < vector_size) return RuntimeError.OutOfBounds;
+
+        for (0..lane_count) |index| {
+            _ = try self.readVectorLane(index, output[index * stride_usize ..]);
+        }
+        return vector_size;
+    }
+
+    pub fn writeVectorWithStride(self: *Self, input: []const u8, stride: SpvWord) RuntimeError!usize {
+        const lane_count: usize = @intCast(try self.resolveLaneCount());
+        if (lane_count == 0) return 0;
+
+        const stride_usize: usize = @intCast(stride);
+        const scalar_size = @divExact(try self.getPlainMemorySize(), lane_count);
+        const vector_size = (lane_count - 1) * stride_usize + scalar_size;
+        if (input.len < vector_size) return RuntimeError.OutOfBounds;
+
+        for (0..lane_count) |index| {
+            _ = try self.writeVectorLane(index, input[index * stride_usize ..]);
+        }
+        return vector_size;
+    }
+
+    pub fn readMatrixWithStride(self: *const Self, output: []u8, matrix_stride: SpvWord, row_major: bool) RuntimeError!usize {
         const columns = switch (self.*) {
             .Matrix => |columns| columns,
             else => return RuntimeError.InvalidValueType,
@@ -491,26 +630,104 @@ pub const Value = union(Type) {
         if (columns.len == 0) return 0;
 
         const matrix_stride_usize: usize = @intCast(matrix_stride);
-        const column_size = try columns[0].getPlainMemorySize();
+        const row_count: usize = @intCast(try columns[0].resolveLaneCount());
-        const matrix_size = (columns.len - 1) * matrix_stride_usize + column_size;
+        if (row_count == 0) return 0;
+
+        const scalar_size = @divExact(try columns[0].getPlainMemorySize(), row_count);
+        const matrix_size = if (row_major)
+            (row_count - 1) * matrix_stride_usize + columns.len * scalar_size
+        else
+            (columns.len - 1) * matrix_stride_usize + try columns[0].getPlainMemorySize();
         if (output.len < matrix_size) return RuntimeError.OutOfBounds;
 
+        if (row_major) {
+            for (columns, 0..) |column, column_index| {
+                for (0..row_count) |row_index| {
+                    _ = try column.readVectorLane(row_index, output[row_index * matrix_stride_usize + column_index * scalar_size ..]);
+                }
+            }
+        } else {
             for (columns, 0..) |column, column_index| {
                 _ = try column.read(output[column_index * matrix_stride_usize ..]);
             }
+        }
         return matrix_size;
     }
 
+    pub fn readWithMatrixLayout(self: *const Self, output: []u8, matrix_stride: SpvWord, row_major: bool) RuntimeError!usize {
+        return switch (self.*) {
+            .Matrix => try self.readMatrixWithStride(output, matrix_stride, row_major),
+            .Array => |arr| blk: {
+                var offset: usize = 0;
+                for (arr.values) |*value| {
+                    _ = try value.readWithMatrixLayout(output[offset..], matrix_stride, row_major);
+                    offset += if (arr.stride != 0) arr.stride else try value.getPlainMemorySize();
+                }
+                break :blk offset;
+            },
+            else => try self.read(output),
+        };
+    }
+
+    pub fn writeMatrixWithStride(self: *Self, input: []const u8, matrix_stride: SpvWord, row_major: bool) RuntimeError!usize {
+        const columns = switch (self.*) {
+            .Matrix => |columns| columns,
+            else => return RuntimeError.InvalidValueType,
+        };
+        if (columns.len == 0) return 0;
+
+        const matrix_stride_usize: usize = @intCast(matrix_stride);
+        const row_count: usize = @intCast(try columns[0].resolveLaneCount());
+        if (row_count == 0) return 0;
+
+        const scalar_size = @divExact(try columns[0].getPlainMemorySize(), row_count);
+        const matrix_size = if (row_major)
+            (row_count - 1) * matrix_stride_usize + columns.len * scalar_size
+        else
+            (columns.len - 1) * matrix_stride_usize + try columns[0].getPlainMemorySize();
+        if (row_major) {
+            for (columns, 0..) |*column, column_index| {
+                for (0..row_count) |row_index| {
+                    const offset = row_index * matrix_stride_usize + column_index * scalar_size;
+                    _ = try column.writeVectorLane(row_index, input[@min(offset, input.len)..]);
+                }
+            }
+        } else {
+            for (columns, 0..) |*column, column_index| {
+                const offset = column_index * matrix_stride_usize;
+                _ = try column.write(input[@min(offset, input.len)..]);
+            }
+        }
+        return matrix_size;
+    }
+
+    pub fn writeWithMatrixLayout(self: *Self, input: []const u8, matrix_stride: SpvWord, row_major: bool) RuntimeError!usize {
+        return switch (self.*) {
+            .Matrix => try self.writeMatrixWithStride(input, matrix_stride, row_major),
+            .Array => |*arr| blk: {
+                var offset: usize = 0;
+                for (arr.values) |*value| {
+                    _ = try value.writeWithMatrixLayout(input[@min(offset, input.len)..], matrix_stride, row_major);
+                    offset += if (arr.stride != 0) arr.stride else try value.getPlainMemorySize();
+                }
+                break :blk offset;
+            },
+            else => try self.write(input),
+        };
+    }
+
     pub fn write(self: *Self, input: []const u8) RuntimeError!usize {
         const vecRoutine = struct {
             inline fn routine(comptime T: type, vec: *T, in: []const u8) RuntimeError!usize {
                 const info = @typeInfo(T).vector;
                 const Lane = info.child;
                 const size = info.len * @sizeOf(Lane);
-                if (in.len < size) return RuntimeError.OutOfBounds;
                 inline for (0..info.len) |i| {
                     const offset = i * @sizeOf(Lane);
-                    vec[i] = std.mem.bytesToValue(Lane, in[offset..][0..@sizeOf(Lane)]);
+                    vec[i] = if (offset + @sizeOf(Lane) <= in.len)
+                        std.mem.bytesToValue(Lane, in[offset..][0..@sizeOf(Lane)])
+                    else
+                        0;
                 }
                 return size;
             }
@@ -518,27 +735,25 @@ pub const Value = union(Type) {
 
         switch (self.*) {
             .Bool => |*b| {
-                if (input.len < 1) return RuntimeError.OutOfBounds;
+                b.* = if (input.len >= 1 and input[0] != 0) true else false;
-                b.* = if (input[0] != 0) true else false;
                 return 1;
             },
             .Int => |*i| {
                 switch (i.bit_count) {
                     8 => {
-                        if (input.len < 1) return RuntimeError.OutOfBounds;
+                        i.value.uint8 = if (input.len >= 1) @bitCast(input[0]) else 0;
-                        i.value.uint8 = @bitCast(input[0]);
                     },
                     16 => {
-                        if (input.len < 2) return RuntimeError.OutOfBounds;
+                        i.value.uint16 = 0;
-                        @memcpy(std.mem.asBytes(&i.value.uint16), input[0..2]);
+                        if (input.len >= 2) @memcpy(std.mem.asBytes(&i.value.uint16), input[0..2]);
                     },
                     32 => {
-                        if (input.len < 4) return RuntimeError.OutOfBounds;
+                        i.value.uint32 = 0;
-                        @memcpy(std.mem.asBytes(&i.value.uint32), input[0..4]);
+                        if (input.len >= 4) @memcpy(std.mem.asBytes(&i.value.uint32), input[0..4]);
                     },
                     64 => {
-                        if (input.len < 8) return RuntimeError.OutOfBounds;
+                        i.value.uint64 = 0;
-                        @memcpy(std.mem.asBytes(&i.value.uint64), input[0..8]);
+                        if (input.len >= 8) @memcpy(std.mem.asBytes(&i.value.uint64), input[0..8]);
                     },
                     else => return RuntimeError.InvalidValueType,
                 }
@@ -547,16 +762,16 @@ pub const Value = union(Type) {
             .Float => |*f| {
                 switch (f.bit_count) {
                     16 => {
-                        if (input.len < 2) return RuntimeError.OutOfBounds;
+                        f.value.float16 = 0;
-                        @memcpy(std.mem.asBytes(&f.value.float16), input[0..2]);
+                        if (input.len >= 2) @memcpy(std.mem.asBytes(&f.value.float16), input[0..2]);
                     },
                     32 => {
-                        if (input.len < 4) return RuntimeError.OutOfBounds;
+                        f.value.float32 = 0;
-                        @memcpy(std.mem.asBytes(&f.value.float32), input[0..4]);
+                        if (input.len >= 4) @memcpy(std.mem.asBytes(&f.value.float32), input[0..4]);
                     },
                     64 => {
-                        if (input.len < 8) return RuntimeError.OutOfBounds;
+                        f.value.float64 = 0;
-                        @memcpy(std.mem.asBytes(&f.value.float64), input[0..8]);
+                        if (input.len >= 8) @memcpy(std.mem.asBytes(&f.value.float64), input[0..8]);
                     },
                     else => return RuntimeError.InvalidValueType,
                 }
@@ -578,22 +793,27 @@ pub const Value = union(Type) {
             .Vector => |*values| {
                 var offset: usize = 0;
                 for (values.*) |*v| {
-                    offset += try v.write(input[offset..]);
+                    offset += try v.write(input[@min(offset, input.len)..]);
                 }
                 return offset;
             },
             .Matrix => |*values| {
                 var offset: usize = 0;
                 for (values.*) |*v| {
-                    offset += try v.write(input[offset..]);
+                    offset += try v.write(input[@min(offset, input.len)..]);
                 }
                 return offset;
             },
             .Array => |*arr| {
                 var offset: usize = 0;
                 for (arr.values) |*v| {
-                    _ = try v.write(input[offset..]);
+                    if (v.* == .RuntimeArray) {
-                    offset += arr.stride;
+                        v.RuntimeArray.data = @constCast(input[@min(offset, input.len)..]);
+                        offset += input.len - @min(offset, input.len);
+                        continue;
+                    }
+                    _ = try v.write(input[@min(offset, input.len)..]);
+                    offset += if (arr.stride != 0) arr.stride else try v.getPlainMemorySize();
                 }
                 return offset;
             },
@@ -601,32 +821,35 @@ pub const Value = union(Type) {
                 var end_offset: usize = 0;
                 for (s.values, 0..) |*v, i| {
                     const member_offset: usize = @intCast(s.offsets[i] orelse end_offset);
-                    if (member_offset > input.len) return RuntimeError.OutOfBounds;
+                    if (member_offset > input.len) {
-                    const write_size = switch (v.*) {
+                        if (v.* == .RuntimeArray) {
-                        .Matrix => |*columns| blk: {
+                            v.RuntimeArray.data = &.{};
-                            const matrix_stride = s.matrix_strides[i] orelse break :blk try v.write(input[member_offset..]);
+                            continue;
-                            if (columns.len == 0) break :blk @as(usize, 0);
-                            const matrix_stride_usize: usize = @intCast(matrix_stride);
-
-                            const column_size = try columns.*[0].getPlainMemorySize();
-                            const matrix_size = (columns.len - 1) * matrix_stride_usize + column_size;
-                            if (input.len - member_offset < matrix_size)
-                                return RuntimeError.OutOfBounds;
-
-                            for (columns.*, 0..) |*column, column_index| {
-                                _ = try column.write(input[member_offset + (column_index * matrix_stride_usize) ..]);
                         }
-                            break :blk matrix_size;
+                    }
-                        },
+                    const member_input = if (v.* == .RuntimeArray and i + 1 < s.values.len and s.offsets[i + 1] != null) blk: {
-                        else => try v.write(input[member_offset..]),
+                        const next_offset: usize = @intCast(s.offsets[i + 1].?);
-                    };
+                        if (next_offset < member_offset) return RuntimeError.OutOfBounds;
+                        break :blk input[member_offset..@min(next_offset, input.len)];
+                    } else input[@min(member_offset, input.len)..];
+                    if (v.* == .RuntimeArray) {
+                        v.RuntimeArray.data = @constCast(member_input);
+                        end_offset = @max(end_offset, input.len);
+                        continue;
+                    }
+                    const write_size = if (s.matrix_strides[i]) |matrix_stride|
+                        try v.writeWithMatrixLayout(member_input, matrix_stride, s.row_major[i])
+                    else
+                        try v.write(member_input);
                     end_offset = @max(end_offset, member_offset + write_size);
                 }
-                if (end_offset > input.len) return RuntimeError.OutOfBounds;
+                s.external_data = @constCast(input[0..@min(end_offset, input.len)]);
-                s.external_data = @constCast(input[0..end_offset]);
                 return end_offset;
             },
-            .RuntimeArray => |*arr| arr.data = @constCast(input[0..]),
+            .RuntimeArray => |*arr| {
+                arr.data = @constCast(input[0..]);
+                return input.len;
+            },
             .Image => |*img| {
                 if (input.len < @sizeOf(usize)) return RuntimeError.OutOfBounds;
                 img.driver_image = @ptrFromInt(std.mem.bytesToValue(usize, input[0..@sizeOf(usize)]));
@@ -676,7 +899,13 @@ pub const Value = union(Type) {
                 }
                 break :blk size;
             },
-            .Array => |arr| arr.stride * arr.values.len,
+            .Array => |arr| blk: {
+                var size: usize = 0;
+                for (arr.values) |v| {
+                    size += if (arr.stride != 0) arr.stride else try v.getPlainMemorySize();
+                }
+                break :blk size;
+            },
             .Structure => |s| blk: {
                 var size: usize = 0;
                 for (s.values, 0..) |v, i| {
@@ -728,14 +957,20 @@ pub const Value = union(Type) {
     pub fn flushPtr(self: *Self, allocator: std.mem.Allocator) RuntimeError!void {
         _ = allocator;
         switch (self.*) {
+            .Structure => |*s| {
+                if (s.external_data) |data| {
+                    _ = try self.read(data);
+                }
+            },
             .Pointer => |*p| {
                 if (p.uniform_slice_window) |window| {
                     switch (p.ptr) {
                         .common => |ptr| {
-                            _ = if (p.matrix_stride) |matrix_stride|
+                            _ = if (p.matrix_stride) |matrix_stride| blk: {
-                                try ptr.readMatrixWithStride(window, matrix_stride)
+                                if (p.matrix_row_major and ptr.isVector())
-                            else
+                                    break :blk try ptr.readVectorWithStride(window, matrix_stride);
-                                try ptr.read(window);
+                                break :blk try ptr.readWithMatrixLayout(window, matrix_stride, p.matrix_row_major);
+                            } else try ptr.read(window);
                         },
                         .f32_ptr => |ptr| {
                             if (window.len < @sizeOf(f32)) return RuntimeError.OutOfBounds;
@@ -773,6 +1008,7 @@ pub const Value = union(Type) {
                 allocator.free(s.values);
                 allocator.free(s.offsets);
                 allocator.free(s.matrix_strides);
+                allocator.free(s.row_major);
             },
             .Pointer => |*p| {
                 if (p.owns_uniform_backing_value) {
@@ -788,14 +1024,54 @@ pub const Value = union(Type) {
         }
     }
 
+    fn readScalarLane(comptime T: PrimitiveType, comptime bits: u32, v: *const Value) RuntimeError!getPrimitiveFieldType(T, bits) {
+        const TT = getPrimitiveFieldType(T, bits);
+
+        return switch (v.*) {
+            .Bool => |b| blk: {
+                if (T != .Bool or bits != 8) return RuntimeError.InvalidSpirV;
+                break :blk @as(TT, b);
+            },
+            .Int => |i| blk: {
+                if (i.bit_count != bits) return RuntimeError.InvalidSpirV;
+                break :blk switch (T) {
+                    .SInt => switch (bits) {
+                        8 => @as(TT, if (i.is_signed) i.value.sint8 else @bitCast(i.value.uint8)),
+                        16 => @as(TT, if (i.is_signed) i.value.sint16 else @bitCast(i.value.uint16)),
+                        32 => @as(TT, if (i.is_signed) i.value.sint32 else @bitCast(i.value.uint32)),
+                        64 => @as(TT, if (i.is_signed) i.value.sint64 else @bitCast(i.value.uint64)),
+                        else => return RuntimeError.InvalidSpirV,
+                    },
+                    .UInt => switch (bits) {
+                        8 => @as(TT, if (i.is_signed) @bitCast(i.value.sint8) else i.value.uint8),
+                        16 => @as(TT, if (i.is_signed) @bitCast(i.value.sint16) else i.value.uint16),
+                        32 => @as(TT, if (i.is_signed) @bitCast(i.value.sint32) else i.value.uint32),
+                        64 => @as(TT, if (i.is_signed) @bitCast(i.value.sint64) else i.value.uint64),
+                        else => return RuntimeError.InvalidSpirV,
+                    },
+                    else => return RuntimeError.InvalidSpirV,
+                };
+            },
+            .Float => |f| blk: {
+                if (T != .Float or f.bit_count != bits) return RuntimeError.InvalidSpirV;
+                break :blk switch (bits) {
+                    16 => @as(TT, f.value.float16),
+                    32 => @as(TT, f.value.float32),
+                    64 => @as(TT, f.value.float64),
+                    else => return RuntimeError.InvalidSpirV,
+                };
+            },
+            else => return RuntimeError.InvalidSpirV,
+        };
+    }
+
     pub inline fn readLane(comptime T: PrimitiveType, comptime bits: u32, v: *const Value, lane_index: usize) RuntimeError!getPrimitiveFieldType(T, bits) {
         const TT = getPrimitiveFieldType(T, bits);
 
         return switch (v.*) {
-            .Int => (try getPrimitiveField(T, bits, @constCast(v))).*,
+            .Bool, .Int, .Float => try readScalarLane(T, bits, v),
-            .Float => (try getPrimitiveField(T, bits, @constCast(v))).*,
 
-            .Vector => |lanes| (try getPrimitiveField(T, bits, &lanes[lane_index])).*,
+            .Vector => |lanes| try readScalarLane(T, bits, &lanes[lane_index]),
 
             .Vector2f32 => |*vec| switch (lane_index) {
                 inline 0...1 => |i| blk: {
@@ -830,7 +1106,7 @@ pub const Value = union(Type) {
 
             .Vector2i32 => |*vec| switch (lane_index) {
                 inline 0...1 => |i| blk: {
-                    if (bits == 32) {
+                    if (comptime (T == .SInt or T == .UInt) and bits == 32) {
                         break :blk @as(TT, @bitCast(vec[i]));
                     } else {
                         return RuntimeError.InvalidSpirV;
@@ -840,7 +1116,7 @@ pub const Value = union(Type) {
             },
             .Vector3i32 => |*vec| switch (lane_index) {
                 inline 0...2 => |i| blk: {
-                    if (bits == 32) {
+                    if (comptime (T == .SInt or T == .UInt) and bits == 32) {
                         break :blk @as(TT, @bitCast(vec[i]));
                     } else {
                         return RuntimeError.InvalidSpirV;
@@ -850,7 +1126,7 @@ pub const Value = union(Type) {
             },
             .Vector4i32 => |*vec| switch (lane_index) {
                 inline 0...3 => |i| blk: {
-                    if (bits == 32) {
+                    if (comptime (T == .SInt or T == .UInt) and bits == 32) {
                         break :blk @as(TT, @bitCast(vec[i]));
                     } else {
                         return RuntimeError.InvalidSpirV;
@@ -861,7 +1137,7 @@ pub const Value = union(Type) {
 
             .Vector2u32 => |*vec| switch (lane_index) {
                 inline 0...1 => |i| blk: {
-                    if (bits == 32) {
+                    if (comptime (T == .SInt or T == .UInt) and bits == 32) {
                         break :blk @as(TT, @bitCast(vec[i]));
                     } else {
                         return RuntimeError.InvalidSpirV;
@@ -871,7 +1147,7 @@ pub const Value = union(Type) {
             },
             .Vector3u32 => |*vec| switch (lane_index) {
                 inline 0...2 => |i| blk: {
-                    if (bits == 32) {
+                    if (comptime (T == .SInt or T == .UInt) and bits == 32) {
                         break :blk @as(TT, @bitCast(vec[i]));
                     } else {
                         return RuntimeError.InvalidSpirV;
@@ -881,7 +1157,7 @@ pub const Value = union(Type) {
             },
             .Vector4u32 => |*vec| switch (lane_index) {
                 inline 0...3 => |i| blk: {
-                    if (bits == 32) {
+                    if (comptime (T == .SInt or T == .UInt) and bits == 32) {
                         break :blk @as(TT, @bitCast(vec[i]));
                     } else {
                         return RuntimeError.InvalidSpirV;
@@ -1014,7 +1290,7 @@ pub const Value = union(Type) {
                             else => unreachable,
                         },
                     } },
-                    else => return RuntimeError.InvalidSpirV,
+                    .Bool => .{ .Bool = v },
                 };
             },
             else => return RuntimeError.InvalidSpirV,

test/api.zig

@@ -48,7 +48,7 @@ fn sampleImageInt4(_: *anyopaque, _: *anyopaque, _: spv.spv.SpvDim, _: f32, _: f
     return spv.Runtime.RuntimeError.UnsupportedSpirV;
 }
 
-fn sampleImageDref(driver_image: *anyopaque, driver_sampler: *anyopaque, _: spv.spv.SpvDim, x: f32, y: f32, z: f32, dref: f32, lod: ?f32, offset: spv.Runtime.ImageOffset) spv.Runtime.RuntimeError!f32 {
+fn sampleImageDref(driver_image: *anyopaque, driver_sampler: *anyopaque, _: spv.spv.SpvDim, x: f32, y: f32, z: f32, _: f32, dref: f32, lod: ?f32, offset: spv.Runtime.ImageOffset) spv.Runtime.RuntimeError!f32 {
     const state: *ImageState = @ptrCast(@alignCast(driver_image));
     if (state.expected_sampler != driver_sampler) return spv.Runtime.RuntimeError.InvalidSpirV;
     state.dref_calls += 1;
@@ -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 };
@@ -283,7 +282,7 @@ test "Built-in inputs and outputs" {
     defer rt.deinit(allocator);
 
     const vertex_index: i32 = 7;
-    try rt.writeBuiltIn(std.mem.asBytes(&vertex_index), .VertexIndex);
+    try rt.writeBuiltIn(allocator, std.mem.asBytes(&vertex_index), .VertexIndex);
     try rt.callEntryPoint(allocator, try rt.getEntryPointByName("main"));
 
     var position: [4]f32 = undefined;
@@ -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"));
@@ -412,7 +410,7 @@ test "Derivative memory buffers" {
     rt.clearDerivative(allocator, output_result);
 
     const input = [_]f32{ 1.0, 2.0, 3.0 };
-    try rt.writeInput(std.mem.asBytes(&input), input_result);
+    try rt.writeInput(allocator, std.mem.asBytes(&input), input_result);
     try rt.callEntryPoint(allocator, try rt.getEntryPointByName("main"));
 
     var output: [4]f32 = undefined;

test/casts.zig

@@ -64,6 +64,107 @@ test "Primitives casts" {
     }
 }
 
+test "Vector float casts" {
+    const allocator = std.testing.allocator;
+    const cases = [_]struct {
+        len: usize,
+        source_type: []const u8,
+        target_type: []const u8,
+        values: []const f64,
+    }{
+        .{ .len = 2, .source_type = "f32", .target_type = "f64", .values = &.{ 1.25, -2.5 } },
+        .{ .len = 3, .source_type = "f32", .target_type = "f64", .values = &.{ 1.25, -2.5, 3.75 } },
+        .{ .len = 4, .source_type = "f32", .target_type = "f64", .values = &.{ 1.25, -2.5, 3.75, -4.5 } },
+        .{ .len = 2, .source_type = "f64", .target_type = "f32", .values = &.{ 1.25, -2.5 } },
+        .{ .len = 3, .source_type = "f64", .target_type = "f32", .values = &.{ 1.25, -2.5, 3.75 } },
+        .{ .len = 4, .source_type = "f64", .target_type = "f32", .values = &.{ 1.25, -2.5, 3.75, -4.5 } },
+    };
+
+    inline for (cases) |c| {
+        const constructor_values = try std.fmt.allocPrint(
+            allocator,
+            switch (c.len) {
+                2 => "{d}, {d}",
+                3 => "{d}, {d}, {d}",
+                4 => "{d}, {d}, {d}, {d}",
+                else => unreachable,
+            },
+            switch (c.len) {
+                2 => .{ c.values[0], c.values[1] },
+                3 => .{ c.values[0], c.values[1], c.values[2] },
+                4 => .{ c.values[0], c.values[1], c.values[2], c.values[3] },
+                else => unreachable,
+            },
+        );
+        defer allocator.free(constructor_values);
+
+        const output_values = switch (c.len) {
+            2 => "casted.x, casted.y, 0.0, 1.0",
+            3 => "casted.x, casted.y, casted.z, 1.0",
+            4 => "casted.x, casted.y, casted.z, casted.w",
+            else => unreachable,
+        };
+
+        const shader = try std.fmt.allocPrint(
+            allocator,
+            \\ [nzsl_version("1.1")]
+            \\ [feature(float64)]
+            \\ module;
+            \\ 
+            \\ struct FragOut
+            \\ {{
+            \\     [location(0)] color: vec4[{s}]
+            \\ }}
+            \\
+            \\ [entry(frag)]
+            \\ fn main() -> FragOut
+            \\ {{
+            \\     let base = vec{d}[{s}]({s});
+            \\     let casted = vec{d}[{s}](base);
+            \\
+            \\     let output: FragOut;
+            \\     output.color = vec4[{s}]({s});
+            \\     return output;
+            \\ }}
+        ,
+            .{
+                c.target_type,
+                c.len,
+                c.source_type,
+                constructor_values,
+                c.len,
+                c.target_type,
+                c.target_type,
+                output_values,
+            },
+        );
+        defer allocator.free(shader);
+
+        const code = try compileNzsl(allocator, shader);
+        defer allocator.free(code);
+
+        if (std.mem.eql(u8, c.target_type, "f32")) {
+            var expected = [_]f32{ 0.0, 0.0, 0.0, 1.0 };
+            for (c.values, 0..) |value, i| expected[i] = @floatCast(value);
+            try case.expect(.{
+                .source = code,
+                .expected_outputs = &.{
+                    std.mem.asBytes(&expected),
+                },
+            });
+        } else {
+            var expected = [_]f64{ 0.0, 0.0, 0.0, 1.0 };
+            for (c.values, 0..) |value, i| expected[i] = @floatCast(value);
+            try case.expect(.{
+                .source = code,
+                .expected_outputs = &.{
+                    std.mem.asBytes(&expected),
+                },
+            });
+        }
+    }
+}
+
 test "Primitives bitcasts" {
     const allocator = std.testing.allocator;
     const types = [_][2]type{

test/maths.zig

@@ -78,6 +78,9 @@ test "Maths primitives" {
                 .expected_outputs = &.{
                     std.mem.asBytes(&[_]T{ expected, expected, expected, expected }),
                 },
+                .expected_output_tolerances = if (@typeInfo(T) == .float) &.{
+                    case.FloatTolerance.default(T),
+                } else &.{},
             });
         }
     }
@@ -150,12 +153,47 @@ test "Maths vectors" {
                     .expected_outputs = &.{
                         std.mem.asBytes(&@as([L]T, expected)),
                     },
+                    .expected_output_tolerances = if (@typeInfo(T) == .float) &.{
+                        case.FloatTolerance.default(T),
+                    } else &.{},
                 });
             }
         }
     }
 }
 
+test "Unsigned vector negate wraps" {
+    const allocator = std.testing.allocator;
+    const shader =
+        \\ [nzsl_version("1.1")]
+        \\ module;
+        \\
+        \\ struct FragOut
+        \\ {
+        \\     [location(0)] color: vec4[u32]
+        \\ }
+        \\
+        \\ [entry(frag)]
+        \\ fn main() -> FragOut
+        \\ {
+        \\     let v = vec4[u32](1, 2, 2147483648, 4294967295);
+        \\     let negated = -v;
+        \\     let output: FragOut;
+        \\     output.color = negated;
+        \\     return output;
+        \\ }
+    ;
+    const code = try compileNzsl(allocator, shader);
+    defer allocator.free(code);
+
+    try case.expect(.{
+        .source = code,
+        .expected_outputs = &.{
+            std.mem.asBytes(&[_]u32{ 4294967295, 4294967294, 2147483648, 1 }),
+        },
+    });
+}
+
 // Tests all mathematical operation on vec2/3/4 with scalars with all NZSL supported primitive types
 test "Maths vectors with scalars" {
     const allocator = std.testing.allocator;
@@ -217,6 +255,9 @@ test "Maths vectors with scalars" {
                     .expected_outputs = &.{
                         std.mem.asBytes(&@as([L]T, expected)),
                     },
+                    .expected_output_tolerances = if (@typeInfo(T) == .float) &.{
+                        case.FloatTolerance.default(T),
+                    } else &.{},
                 });
             }
         }

test/root.zig

@@ -24,10 +24,37 @@ pub const case = struct {
         source: []const u32,
         inputs: []const []const u8 = &.{},
         expected_outputs: []const []const u8 = &.{},
+        expected_output_tolerances: []const ?FloatTolerance = &.{},
         descriptor_sets: []const []const []u8 = &.{},
         expected_descriptor_sets: []const []const []const u8 = &.{},
     };
 
+    const FloatKind = enum { f32, f64 };
+
+    pub const FloatTolerance = struct {
+        kind: FloatKind,
+        absolute: f64,
+
+        pub fn of(comptime T: type, absolute: f64) FloatTolerance {
+            return .{
+                .kind = switch (T) {
+                    f32 => .f32,
+                    f64 => .f64,
+                    else => @compileError("FloatTolerance only supports f32 and f64"),
+                },
+                .absolute = absolute,
+            };
+        }
+
+        pub fn default(comptime T: type) FloatTolerance {
+            return of(T, switch (T) {
+                f32 => 1e-6,
+                f64 => 1e-12,
+                else => @compileError("FloatTolerance only supports f32 and f64"),
+            });
+        }
+    };
+
     pub fn expect(config: Config) !void {
         const allocator = std.testing.allocator;
 
@@ -49,7 +76,7 @@ pub const case = struct {
             defer rt.deinit(allocator);
 
             for (config.inputs, 0..) |input, n| {
-                try rt.writeInput(input[0..], module.input_locations[n][0]);
+                try rt.writeInput(allocator, input[0..], module.input_locations[n][0]);
             }
 
             for (config.descriptor_sets, 0..) |descriptor_set, set_index| {
@@ -66,6 +93,12 @@ pub const case = struct {
                 defer allocator.free(output);
 
                 try rt.readOutput(output[0..], module.output_locations[n][0]);
+                if (n < config.expected_output_tolerances.len) {
+                    if (config.expected_output_tolerances[n]) |tolerance| {
+                        try expectApproxBytes(expected, output, tolerance);
+                        continue;
+                    }
+                }
                 try std.testing.expectEqualSlices(u8, expected, output);
             }
 
@@ -77,6 +110,30 @@ pub const case = struct {
         }
     }
 
+    fn expectApproxBytes(expected: []const u8, actual: []const u8, tolerance: FloatTolerance) !void {
+        try std.testing.expectEqual(expected.len, actual.len);
+
+        const stride: usize = switch (tolerance.kind) {
+            .f32 => @sizeOf(f32),
+            .f64 => @sizeOf(f64),
+        };
+        try std.testing.expectEqual(@as(usize, 0), expected.len % stride);
+
+        var i: usize = 0;
+        while (i < expected.len) : (i += stride) {
+            const expected_value = readFloat(expected[i .. i + stride], tolerance.kind);
+            const actual_value = readFloat(actual[i .. i + stride], tolerance.kind);
+            try std.testing.expectApproxEqAbs(expected_value, actual_value, tolerance.absolute);
+        }
+    }
+
+    fn readFloat(bytes: []const u8, kind: FloatKind) f64 {
+        return switch (kind) {
+            .f32 => @floatCast(@as(f32, @bitCast(std.mem.readInt(u32, bytes[0..@sizeOf(u32)], .little)))),
+            .f64 => @bitCast(std.mem.readInt(u64, bytes[0..@sizeOf(u64)], .little)),
+        };
+    }
+
     pub fn random(comptime T: type) T {
         var prng: std.Random.DefaultPrng = .init(@intCast(std.Io.Timestamp.now(std.testing.io, .real).toNanoseconds()));
         const rand = prng.random();

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))
-	if (output_size != sizeof(output) || primitive_type != SPV_PRIMITIVE_FLOAT || primitive_type == SPV_PRIMITIVE_SINT || primitive_type == SPV_PRIMITIVE_UINT)
 	{
 		fprintf(stderr, "Unexpected output metadata\n");
 		SpvDeinitRuntime(second_runtime);