diff options
author | Hans-Kristian Arntzen <post@arntzen-software.no> | 2021-10-21 22:53:15 +0300 |
---|---|---|
committer | Hans-Kristian Arntzen <post@arntzen-software.no> | 2021-10-21 22:53:41 +0300 |
commit | 5afb3d313f28bdd55b7284ce9044387d86495857 (patch) | |
tree | ff0ef9b391b2e8874c988e662b0acba479794d7d | |
parent | 345a7d171c43b9f79f27652d48b386578a70980a (diff) |
MSL: Fix some trivial bugs not caught by CI when adding ray query.
-rw-r--r-- | reference/opt/shaders-msl/comp/ray-query.spv14.vk.ios.msl24.comp | 92 | ||||
-rw-r--r-- | reference/shaders-msl/comp/ray-query.spv14.vk.ios.msl24.comp | 2 | ||||
-rw-r--r-- | spirv_msl.cpp | 2 |
3 files changed, 94 insertions, 2 deletions
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]); |