From 5afb3d313f28bdd55b7284ce9044387d86495857 Mon Sep 17 00:00:00 2001
From: Hans-Kristian Arntzen <post@arntzen-software.no>
Date: Thu, 21 Oct 2021 21:53:15 +0200
Subject: [PATCH] MSL: Fix some trivial bugs not caught by CI when adding ray
 query.

---
 .../comp/ray-query.spv14.vk.ios.msl24.comp    | 92 +++++++++++++++++++
 .../comp/ray-query.spv14.vk.ios.msl24.comp    |  2 +-
 spirv_msl.cpp                                 |  2 +-
 3 files changed, 94 insertions(+), 2 deletions(-)
 create mode 100644 reference/opt/shaders-msl/comp/ray-query.spv14.vk.ios.msl24.comp

diff --git a/reference/opt/shaders-msl/comp/ray-query.spv14.vk.ios.msl24.comp b/reference/opt/shaders-msl/comp/ray-query.spv14.vk.ios.msl24.comp
new file mode 100644
index 00000000..84159295
--- /dev/null
+++ b/reference/opt/shaders-msl/comp/ray-query.spv14.vk.ios.msl24.comp
@@ -0,0 +1,92 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+#pragma clang diagnostic ignored "-Wmissing-braces"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+#if __METAL_VERSION__ >= 230
+#include <metal_raytracing>
+using namespace metal::raytracing;
+#endif
+
+using namespace metal;
+
+template<typename T, size_t Num>
+struct spvUnsafeArray
+{
+    T elements[Num ? Num : 1];
+    
+    thread T& operator [] (size_t pos) thread
+    {
+        return elements[pos];
+    }
+    constexpr const thread T& operator [] (size_t pos) const thread
+    {
+        return elements[pos];
+    }
+    
+    device T& operator [] (size_t pos) device
+    {
+        return elements[pos];
+    }
+    constexpr const device T& operator [] (size_t pos) const device
+    {
+        return elements[pos];
+    }
+    
+    constexpr const constant T& operator [] (size_t pos) const constant
+    {
+        return elements[pos];
+    }
+    
+    threadgroup T& operator [] (size_t pos) threadgroup
+    {
+        return elements[pos];
+    }
+    constexpr const threadgroup T& operator [] (size_t pos) const threadgroup
+    {
+        return elements[pos];
+    }
+};
+
+struct Params
+{
+    uint ray_flags;
+    uint cull_mask;
+    char _m2_pad[8];
+    packed_float3 origin;
+    float tmin;
+    packed_float3 dir;
+    float tmax;
+    float thit;
+};
+
+kernel void main0(constant Params& _18 [[buffer(1)]], acceleration_structure<instancing> AS0 [[buffer(0)]], acceleration_structure<instancing> AS1 [[buffer(2)]])
+{
+    intersection_query<instancing, triangle_data> q;
+    q.reset(ray(_18.origin, _18.dir, _18.tmin, _18.tmax), AS0, intersection_params());
+    spvUnsafeArray<intersection_query<instancing, triangle_data>, 2> q2;
+    q2[1].reset(ray(_18.origin, _18.dir, _18.tmin, _18.tmax), AS1, intersection_params());
+    bool _63 = q.next();
+    q2[0].abort();
+    q.commit_bounding_box_intersection(_18.thit);
+    q2[1].commit_triangle_intersection();
+    float _71 = q.get_ray_min_distance();
+    float3 _74 = q.get_world_space_ray_origin();
+    float3 _75 = q.get_world_space_ray_direction();
+    uint _80 = uint(q2[1].get_committed_intersection_type());
+    uint _83 = uint(q2[0].get_candidate_intersection_type()) - 1;
+    bool _85 = q2[1].is_candidate_non_opaque_bounding_box();
+    float _87 = q2[1].get_committed_distance();
+    float _89 = q2[1].get_candidate_triangle_distance();
+    int _92 = q.get_committed_user_instance_id();
+    int _94 = q2[0].get_candidate_instance_id();
+    int _96 = q2[1].get_candidate_geometry_id();
+    int _97 = q.get_committed_primitive_id();
+    float2 _100 = q2[0].get_candidate_triangle_barycentric_coord();
+    bool _103 = q.is_committed_triangle_front_facing();
+    float3 _104 = q.get_candidate_ray_direction();
+    float3 _106 = q2[0].get_committed_ray_origin();
+    float4x3 _110 = q.get_candidate_object_to_world_transform();
+    float4x3 _112 = q2[1].get_committed_world_to_object_transform();
+}
+
diff --git a/reference/shaders-msl/comp/ray-query.spv14.vk.ios.msl24.comp b/reference/shaders-msl/comp/ray-query.spv14.vk.ios.msl24.comp
index f73d4911..019393ef 100644
--- a/reference/shaders-msl/comp/ray-query.spv14.vk.ios.msl24.comp
+++ b/reference/shaders-msl/comp/ray-query.spv14.vk.ios.msl24.comp
@@ -70,7 +70,7 @@ kernel void main0(constant Params& _18 [[buffer(1)]], acceleration_structure<ins
     bool res = _63;
     q2[0].abort();
     q.commit_bounding_box_intersection(_18.thit);
-    _14.commit_triangle_intersection();
+    q2[1].commit_triangle_intersection();
     float _71 = q.get_ray_min_distance();
     float fval = _71;
     float3 _74 = q.get_world_space_ray_origin();
diff --git a/spirv_msl.cpp b/spirv_msl.cpp
index fed295a2..2e72deef 100644
--- a/spirv_msl.cpp
+++ b/spirv_msl.cpp
@@ -8460,7 +8460,7 @@ void CompilerMSL::emit_instruction(const Instruction &instruction)
 	}
 	case OpRayQueryConfirmIntersectionKHR:
 		flush_variable_declaration(ops[0]);
-		statement(to_expression(ops[2]), ".commit_triangle_intersection();");
+		statement(to_expression(ops[0]), ".commit_triangle_intersection();");
 		break;
 	case OpRayQueryGenerateIntersectionKHR:
 		flush_variable_declaration(ops[0]);