AuroraMiddleware/SPIRV-Cross
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

Merge pull request #1773 from congyue1977/master

Browse Source 
Support Metal 2.4 Intersection Query, Implement GL_EXT_ray_query.
This commit is contained in:
Hans-Kristian Arntzen 2021-10-21 15:56:15 +02:00 committed by GitHub
commit 345a7d171c
No known key found for this signature in database
GPG Key ID: 4AEE18F83AFDEB23
6 changed files with 394 additions and 1 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

91
reference/opt/shaders-msl/comp/ray-query.nocompat.spv14.vk.comp Normal file
View File

@ -0,0 +1,91 @@
#pragma clang diagnostic ignored "-Wmissing-prototypes"
#pragma clang diagnostic ignored "-Wmissing-braces"
#include <metal_stdlib>
#include <simd/simd.h>
using namespace metal::raytracing;
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;
intersection_params _intersection_params_;
q.reset(ray(_18.origin, _18.dir, _18.tmin, _18.tmax), AS0, _intersection_params_);
spvUnsafeArray<intersection_query<instancing, triangle_data>, 2> q2;
intersection_params _intersection_params_;
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);
_14.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_committed_intersection_type();
bool _85 = q2[1].is_candidate_non_opaque_bounding_box();
float _87 = q2[1].get_committed_distance();
float _89 = q2[1].get_committed_distance();
int _92 = q.get_committed_user_instance_id();
int _94 = q2[0].get_committed_instance_id();
int _96 = q2[1].get_committed_geometry_id();
int _97 = q.get_committed_primitive_id();
float2 _100 = q2[0].get_committed_triangle_barycentric_coord();
bool _103 = q.is_committed_triangle_front_facing();
float3 _104 = q.get_committed_ray_direction();
float3 _106 = q2[0].get_committed_ray_origin();
float4x3 _110 = q.get_committed_object_to_world_transform();
float4x3 _112 = q2[1].get_committed_world_to_object_transform();
}

111
reference/shaders-msl/comp/ray-query.spv14.vk.ios.msl24.comp Normal file
View File

@ -0,0 +1,111 @@
#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();
bool res = _63;
q2[0].abort();
q.commit_bounding_box_intersection(_18.thit);
_14.commit_triangle_intersection();
float _71 = q.get_ray_min_distance();
float fval = _71;
float3 _74 = q.get_world_space_ray_origin();
float3 fvals = _74;
float3 _75 = q.get_world_space_ray_direction();
fvals = _75;
uint _80 = uint(q2[1].get_committed_intersection_type());
uint type = _80;
uint _83 = uint(q2[0].get_candidate_intersection_type()) - 1;
type = _83;
bool _85 = q2[1].is_candidate_non_opaque_bounding_box();
res = _85;
float _87 = q2[1].get_committed_distance();
fval = _87;
float _89 = q2[1].get_candidate_triangle_distance();
fval = _89;
int _92 = q.get_committed_user_instance_id();
int ival = _92;
int _94 = q2[0].get_candidate_instance_id();
ival = _94;
int _96 = q2[1].get_candidate_geometry_id();
ival = _96;
int _97 = q.get_committed_primitive_id();
ival = _97;
float2 _100 = q2[0].get_candidate_triangle_barycentric_coord();
fvals = float3(_100.x, _100.y, fvals.z);
bool _103 = q.is_committed_triangle_front_facing();
res = _103;
float3 _104 = q.get_candidate_ray_direction();
fvals = _104;
float3 _106 = q2[0].get_committed_ray_origin();
fvals = _106;
float4x3 _110 = q.get_candidate_object_to_world_transform();
float4x3 matrices = _110;
float4x3 _112 = q2[1].get_committed_world_to_object_transform();
matrices = _112;
}

58
shaders-msl/comp/ray-query.spv14.vk.ios.msl24.comp Normal file
View File

@ -0,0 +1,58 @@
#version 460
#extension GL_EXT_ray_query : require
#extension GL_EXT_ray_tracing : require
#extension GL_EXT_ray_flags_primitive_culling : require
layout(primitive_culling);
layout(set = 0, binding = 0) uniform accelerationStructureEXT AS0;
layout(set = 0, binding = 1) uniform accelerationStructureEXT AS1;
layout(set = 0, binding = 2) uniform Params
{
uint ray_flags;
uint cull_mask;
vec3 origin;
float tmin;
vec3 dir;
float tmax;
float thit;
};
rayQueryEXT q2[2];
void main()
{
rayQueryEXT q;
bool res;
uint type;
float fval;
vec3 fvals;
int ival;
mat4x3 matrices;
rayQueryInitializeEXT(q, AS0, ray_flags, cull_mask, origin, tmin, dir, tmax);
rayQueryInitializeEXT(q2[1], AS1, ray_flags, cull_mask, origin, tmin, dir, tmax);
res = rayQueryProceedEXT(q);
rayQueryTerminateEXT(q2[0]);
rayQueryGenerateIntersectionEXT(q, thit);
rayQueryConfirmIntersectionEXT(q2[1]);
fval = rayQueryGetRayTMinEXT(q);
fvals = rayQueryGetWorldRayDirectionEXT(q);
fvals = rayQueryGetWorldRayOriginEXT(q);
type = rayQueryGetIntersectionTypeEXT(q2[1], true);
type = rayQueryGetIntersectionTypeEXT(q2[0], false);
res = rayQueryGetIntersectionCandidateAABBOpaqueEXT(q2[1]);
fval = rayQueryGetIntersectionTEXT(q2[1], true);
fval = rayQueryGetIntersectionTEXT(q2[1], false);
ival = rayQueryGetIntersectionInstanceCustomIndexEXT(q, true);
ival = rayQueryGetIntersectionInstanceIdEXT(q2[0], false);
ival = rayQueryGetIntersectionGeometryIndexEXT(q2[1], false);
ival = rayQueryGetIntersectionPrimitiveIndexEXT(q, true);
fvals.xy = rayQueryGetIntersectionBarycentricsEXT(q2[0], false);
res = rayQueryGetIntersectionFrontFaceEXT(q, true);
fvals = rayQueryGetIntersectionObjectRayDirectionEXT(q, false);
fvals = rayQueryGetIntersectionObjectRayOriginEXT(q2[0], true);
matrices = rayQueryGetIntersectionObjectToWorldEXT(q, false);
matrices = rayQueryGetIntersectionWorldToObjectEXT(q2[1], true);
}

123
spirv_msl.cpp
View File

@ -1512,6 +1512,14 @@ void CompilerMSL::preprocess_op_codes()
(is_sample_rate() && (active_input_builtins.get(BuiltInFragCoord) ||
(need_subpass_input && !msl_options.use_framebuffer_fetch_subpasses))))
needs_sample_id = true;
if (is_intersection_query())
{
add_header_line("#if __METAL_VERSION__ >= 230");
add_header_line("#include <metal_raytracing>");
add_header_line("using namespace metal::raytracing;");
add_header_line("#endif");
}
}
// Move the Private and Workgroup global variables to the entry function.
@ -8373,6 +8381,100 @@ void CompilerMSL::emit_instruction(const Instruction &instruction)
SPIRV_CROSS_THROW("Raster order groups require MSL 2.0.");
break; // Nothing to do in the body
case OpConvertUToAccelerationStructureKHR:
SPIRV_CROSS_THROW("ConvertUToAccelerationStructure is not supported in MSL.");
break; // Nothing to do in the body
case OpRayQueryGetIntersectionInstanceShaderBindingTableRecordOffsetKHR:
SPIRV_CROSS_THROW("BindingTableRecordOffset is not supported in MSL.");
break; // Nothing to do in the body
case OpRayQueryInitializeKHR:
{
flush_variable_declaration(ops[0]);
statement(to_expression(ops[0]), ".reset(", "ray(", to_expression(ops[4]), ", ", to_expression(ops[6]), ", ",
to_expression(ops[5]), ", ", to_expression(ops[7]), "), ", to_expression(ops[1]),
", intersection_params());");
break;
}
case OpRayQueryProceedKHR:
{
flush_variable_declaration(ops[0]);
emit_op(ops[0], ops[1], join(to_expression(ops[2]), ".next()"), false);
break;
}
#define MSL_RAY_QUERY_IS_CANDIDATE get<SPIRConstant>(ops[3]).scalar_i32() == 0
#define MSL_RAY_QUERY_GET_OP(op, msl_op) \
case OpRayQueryGet##op##KHR: \
flush_variable_declaration(ops[2]); \
emit_op(ops[0], ops[1], join(to_expression(ops[2]), ".get_" #msl_op "()"), false); \
break
#define MSL_RAY_QUERY_OP_INNER2(op, msl_prefix, msl_op) \
case OpRayQueryGet##op##KHR: \
flush_variable_declaration(ops[2]); \
if (MSL_RAY_QUERY_IS_CANDIDATE) \
emit_op(ops[0], ops[1], join(to_expression(ops[2]), #msl_prefix "_candidate_" #msl_op "()"), false); \
else \
emit_op(ops[0], ops[1], join(to_expression(ops[2]), #msl_prefix "_committed_" #msl_op "()"), false); \
break
#define MSL_RAY_QUERY_GET_OP2(op, msl_op) MSL_RAY_QUERY_OP_INNER2(op, .get, msl_op)
#define MSL_RAY_QUERY_IS_OP2(op, msl_op) MSL_RAY_QUERY_OP_INNER2(op, .is, msl_op)
MSL_RAY_QUERY_GET_OP(RayTMin, ray_min_distance);
MSL_RAY_QUERY_GET_OP(WorldRayOrigin, world_space_ray_direction);
MSL_RAY_QUERY_GET_OP(WorldRayDirection, world_space_ray_origin);
MSL_RAY_QUERY_GET_OP2(IntersectionInstanceId, instance_id);
MSL_RAY_QUERY_GET_OP2(IntersectionInstanceCustomIndex, user_instance_id);
MSL_RAY_QUERY_GET_OP2(IntersectionBarycentrics, triangle_barycentric_coord);
MSL_RAY_QUERY_GET_OP2(IntersectionPrimitiveIndex, primitive_id);
MSL_RAY_QUERY_GET_OP2(IntersectionGeometryIndex, geometry_id);
MSL_RAY_QUERY_GET_OP2(IntersectionObjectRayOrigin, ray_origin);
MSL_RAY_QUERY_GET_OP2(IntersectionObjectRayDirection, ray_direction);
MSL_RAY_QUERY_GET_OP2(IntersectionObjectToWorld, object_to_world_transform);
MSL_RAY_QUERY_GET_OP2(IntersectionWorldToObject, world_to_object_transform);
MSL_RAY_QUERY_IS_OP2(IntersectionFrontFace, triangle_front_facing);
case OpRayQueryGetIntersectionTypeKHR:
flush_variable_declaration(ops[2]);
if (MSL_RAY_QUERY_IS_CANDIDATE)
emit_op(ops[0], ops[1], join("uint(", to_expression(ops[2]), ".get_candidate_intersection_type()) - 1"),
false);
else
emit_op(ops[0], ops[1], join("uint(", to_expression(ops[2]), ".get_committed_intersection_type())"), false);
break;
case OpRayQueryGetIntersectionTKHR:
flush_variable_declaration(ops[2]);
if (MSL_RAY_QUERY_IS_CANDIDATE)
emit_op(ops[0], ops[1], join(to_expression(ops[2]), ".get_candidate_triangle_distance()"), false);
else
emit_op(ops[0], ops[1], join(to_expression(ops[2]), ".get_committed_distance()"), false);
break;
case OpRayQueryGetIntersectionCandidateAABBOpaqueKHR:
{
flush_variable_declaration(ops[0]);
emit_op(ops[0], ops[1], join(to_expression(ops[2]), ".is_candidate_non_opaque_bounding_box()"), false);
break;
}
case OpRayQueryConfirmIntersectionKHR:
flush_variable_declaration(ops[0]);
statement(to_expression(ops[2]), ".commit_triangle_intersection();");
break;
case OpRayQueryGenerateIntersectionKHR:
flush_variable_declaration(ops[0]);
statement(to_expression(ops[0]), ".commit_bounding_box_intersection(", to_expression(ops[1]), ");");
break;
case OpRayQueryTerminateKHR:
flush_variable_declaration(ops[0]);
statement(to_expression(ops[0]), ".abort();");
break;
#undef MSL_RAY_QUERY_GET_OP
#undef MSL_RAY_QUERY_IS_CANDIDATE
#undef MSL_RAY_QUERY_IS_OP2
#undef MSL_RAY_QUERY_GET_OP2
#undef MSL_RAY_QUERY_OP_INNER2
default:
CompilerGLSL::emit_instruction(instruction);
break;
@ -11295,6 +11397,12 @@ bool CompilerMSL::is_sample_rate() const
(msl_options.use_framebuffer_fetch_subpasses && need_subpass_input));
}
bool CompilerMSL::is_intersection_query() const
{
auto &caps = get_declared_capabilities();
return std::find(caps.begin(), caps.end(), CapabilityRayQueryKHR) != caps.end();
}
void CompilerMSL::entry_point_args_builtin(string &ep_args)
{
// Builtin variables
@ -11773,6 +11881,10 @@ void CompilerMSL::entry_point_args_discrete_descriptors(string &ep_args)
}
break;
}
case SPIRType::AccelerationStructure:
ep_args += ", " + type_to_glsl(type, var_id) + " " + r.name;
ep_args += " [[buffer(" + convert_to_string(r.index) + ")]]";
break;
default:
if (!ep_args.empty())
ep_args += ", ";
@ -13283,6 +13395,17 @@ string CompilerMSL::type_to_glsl(const SPIRType &type, uint32_t id)
case SPIRType::Double:
type_name = "double"; // Currently unsupported
break;
case SPIRType::AccelerationStructure:
if (msl_options.supports_msl_version(2, 4))
type_name = "acceleration_structure<instancing>";
else if (msl_options.supports_msl_version(2, 3))
type_name = "instance_acceleration_structure";
else
SPIRV_CROSS_THROW("Acceleration Structure Type is supported in MSL 2.3 and above.");
break;
case SPIRType::RayQuery:
type_name = "intersection_query<instancing, triangle_data>";
break;
default:
return "unknown_type";

1
spirv_msl.hpp
View File

@ -864,6 +864,7 @@ protected:
std::string to_swizzle_expression(uint32_t id);
std::string to_buffer_size_expression(uint32_t id);
bool is_sample_rate() const;
bool is_intersection_query() const;
bool is_direct_input_builtin(spv::BuiltIn builtin);
std::string builtin_qualifier(spv::BuiltIn builtin);
std::string builtin_type_decl(spv::BuiltIn builtin, uint32_t id = 0);

11
test_shaders.py
View File

@ -132,6 +132,8 @@ def path_to_msl_standard(shader):
return '-std=ios-metal2.2'
elif '.msl23.' in shader:
return '-std=ios-metal2.3'
elif '.msl24.' in shader:
return '-std=ios-metal2.4'
elif '.msl11.' in shader:
return '-std=ios-metal1.1'
elif '.msl10.' in shader:
@ -147,6 +149,8 @@ def path_to_msl_standard(shader):
return '-std=macos-metal2.2'
elif '.msl23.' in shader:
return '-std=macos-metal2.3'
elif '.msl24.' in shader:
return '-std=macos-metal2.4'
elif '.msl11.' in shader:
return '-std=macos-metal1.1'
else:
@ -161,6 +165,8 @@ def path_to_msl_standard_cli(shader):
return '20200'
elif '.msl23.' in shader:
return '20300'
elif '.msl24.' in shader:
return '20400'
elif '.msl11.' in shader:
return '10100'
else:
@ -768,7 +774,8 @@ def test_shader_msl(stats, shader, args, paths):
shader_is_msl22 = 'msl22' in joined_path
shader_is_msl23 = 'msl23' in joined_path
skip_validation = (shader_is_msl22 and (not args.msl22)) or (shader_is_msl23 and (not args.msl23))
shader_is_msl24 = 'msl24' in joined_path
skip_validation = (shader_is_msl22 and (not args.msl22)) or (shader_is_msl23 and (not args.msl23)) or (shader_is_msl24 and (not args.msl24))
if '.invalid.' in joined_path:
skip_validation = True
@ -917,10 +924,12 @@ def main():
args.msl22 = False
args.msl23 = False
args.msl24 = False
if args.msl:
print_msl_compiler_version()
args.msl22 = msl_compiler_supports_version('2.2')
args.msl23 = msl_compiler_supports_version('2.3')
args.msl24 = msl_compiler_supports_version('2.4')
backend = 'glsl'
if (args.msl or args.metal):