diff --git a/CHANGELOG.md b/CHANGELOG.md index 35b9d3c7128..610e82e8ac1 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -90,6 +90,10 @@ SamplerDescriptor { - Validation errors from `CommandEncoder::finish()` will report the label of the invalid encoder. By @kpreid in [#8449](https://github.com/gfx-rs/wgpu/pull/8449). - Corrected documentation of the minimum alignment of the *end* of a mapped range of a buffer (it is 4, not 8). By @kpreid in [#8450](https://github.com/gfx-rs/wgpu/pull/8450). +#### Naga + +- Prevent UB with invalid ray query calls on spirv. By @Vecvec in [#8390](https://github.com/gfx-rs/wgpu/pull/8390). + ### Bug Fixes #### naga diff --git a/naga-test/src/lib.rs b/naga-test/src/lib.rs index 078aa27c405..769026edf5e 100644 --- a/naga-test/src/lib.rs +++ b/naga-test/src/lib.rs @@ -114,6 +114,7 @@ pub struct SpirvOutParameters { pub separate_entry_points: bool, #[serde(deserialize_with = "deserialize_binding_map")] pub binding_map: naga::back::spv::BindingMap, + pub ray_query_initialization_tracking: bool, pub use_storage_input_output_16: bool, } impl Default for SpirvOutParameters { @@ -126,6 +127,7 @@ impl Default for SpirvOutParameters { force_point_size: false, clamp_frag_depth: false, separate_entry_points: false, + ray_query_initialization_tracking: true, use_storage_input_output_16: true, binding_map: naga::back::spv::BindingMap::default(), } @@ -159,6 +161,7 @@ impl SpirvOutParameters { binding_map: self.binding_map.clone(), zero_initialize_workgroup_memory: spv::ZeroInitializeWorkgroupMemoryMode::Polyfill, force_loop_bounding: true, + ray_query_initialization_tracking: true, debug_info, use_storage_input_output_16: self.use_storage_input_output_16, } diff --git a/naga/src/back/spv/block.rs b/naga/src/back/spv/block.rs index d0556acdc53..cc9280d5778 100644 --- a/naga/src/back/spv/block.rs +++ b/naga/src/back/spv/block.rs @@ -203,7 +203,7 @@ impl Writer { )); let clamp_id = self.id_gen.next(); - body.push(Instruction::ext_inst( + body.push(Instruction::ext_inst_gl_op( self.gl450_ext_inst_id, spirv::GLOp::FClamp, float_type_id, @@ -1026,7 +1026,7 @@ impl BlockContext<'_> { }; let max_id = self.gen_id(); - block.body.push(Instruction::ext_inst( + block.body.push(Instruction::ext_inst_gl_op( self.writer.gl450_ext_inst_id, max_op, result_type_id, @@ -1034,7 +1034,7 @@ impl BlockContext<'_> { &[arg0_id, arg1_id], )); - MathOp::Custom(Instruction::ext_inst( + MathOp::Custom(Instruction::ext_inst_gl_op( self.writer.gl450_ext_inst_id, min_op, result_type_id, @@ -1068,7 +1068,7 @@ impl BlockContext<'_> { arg2_id = self.writer.get_constant_composite(ty, &self.temp_list); } - MathOp::Custom(Instruction::ext_inst( + MathOp::Custom(Instruction::ext_inst_gl_op( self.writer.gl450_ext_inst_id, spirv::GLOp::FClamp, result_type_id, @@ -1282,7 +1282,7 @@ impl BlockContext<'_> { &self.temp_list, )); - MathOp::Custom(Instruction::ext_inst( + MathOp::Custom(Instruction::ext_inst_gl_op( self.writer.gl450_ext_inst_id, spirv::GLOp::FMix, result_type_id, @@ -1339,7 +1339,7 @@ impl BlockContext<'_> { }; let lsb_id = self.gen_id(); - block.body.push(Instruction::ext_inst( + block.body.push(Instruction::ext_inst_gl_op( self.writer.gl450_ext_inst_id, spirv::GLOp::FindILsb, result_type_id, @@ -1347,7 +1347,7 @@ impl BlockContext<'_> { &[arg0_id], )); - MathOp::Custom(Instruction::ext_inst( + MathOp::Custom(Instruction::ext_inst_gl_op( self.writer.gl450_ext_inst_id, spirv::GLOp::UMin, result_type_id, @@ -1388,7 +1388,7 @@ impl BlockContext<'_> { }; let msb_id = self.gen_id(); - block.body.push(Instruction::ext_inst( + block.body.push(Instruction::ext_inst_gl_op( self.writer.gl450_ext_inst_id, if width != 4 { spirv::GLOp::FindILsb @@ -1445,7 +1445,7 @@ impl BlockContext<'_> { // o = min(offset, w) let offset_id = self.gen_id(); - block.body.push(Instruction::ext_inst( + block.body.push(Instruction::ext_inst_gl_op( self.writer.gl450_ext_inst_id, spirv::GLOp::UMin, u32_type, @@ -1465,7 +1465,7 @@ impl BlockContext<'_> { // c = min(count, tmp) let count_id = self.gen_id(); - block.body.push(Instruction::ext_inst( + block.body.push(Instruction::ext_inst_gl_op( self.writer.gl450_ext_inst_id, spirv::GLOp::UMin, u32_type, @@ -1495,7 +1495,7 @@ impl BlockContext<'_> { // o = min(offset, w) let offset_id = self.gen_id(); - block.body.push(Instruction::ext_inst( + block.body.push(Instruction::ext_inst_gl_op( self.writer.gl450_ext_inst_id, spirv::GLOp::UMin, u32_type, @@ -1515,7 +1515,7 @@ impl BlockContext<'_> { // c = min(count, tmp) let count_id = self.gen_id(); - block.body.push(Instruction::ext_inst( + block.body.push(Instruction::ext_inst_gl_op( self.writer.gl450_ext_inst_id, spirv::GLOp::UMin, u32_type, @@ -1610,7 +1610,7 @@ impl BlockContext<'_> { }; block.body.push(match math_op { - MathOp::Ext(op) => Instruction::ext_inst( + MathOp::Ext(op) => Instruction::ext_inst_gl_op( self.writer.gl450_ext_inst_id, op, result_type_id, @@ -1621,7 +1621,27 @@ impl BlockContext<'_> { }); id } - crate::Expression::LocalVariable(variable) => self.function.variables[&variable].id, + crate::Expression::LocalVariable(variable) => { + if let Some(rq_tracker) = self + .function + .ray_query_initialization_tracker_variables + .get(&variable) + { + self.ray_query_tracker_expr.insert( + expr_handle, + super::RayQueryTrackers { + initialized_tracker: rq_tracker.id, + t_max_tracker: self + .function + .ray_query_t_max_tracker_variables + .get(&variable) + .expect("Both trackers are set at the same time.") + .id, + }, + ); + } + self.function.variables[&variable].id + } crate::Expression::Load { pointer } => { self.write_checked_load(pointer, block, AccessTypeAdjustment::None, result_type_id)? } @@ -1772,6 +1792,10 @@ impl BlockContext<'_> { crate::Expression::ArrayLength(expr) => self.write_runtime_array_length(expr, block)?, crate::Expression::RayQueryGetIntersection { query, committed } => { let query_id = self.cached[query]; + let init_tracker_id = *self + .ray_query_tracker_expr + .get(&query) + .expect("not a cached ray query"); let func_id = self .writer .write_ray_query_get_intersection_function(committed, self.ir_module); @@ -1782,7 +1806,7 @@ impl BlockContext<'_> { intersection_type_id, id, func_id, - &[query_id], + &[query_id, init_tracker_id.initialized_tracker], )); id } @@ -2008,7 +2032,7 @@ impl BlockContext<'_> { let max_const_id = maybe_splat_const(self.writer, max_const_id); let clamp_id = self.gen_id(); - block.body.push(Instruction::ext_inst( + block.body.push(Instruction::ext_inst_gl_op( self.writer.gl450_ext_inst_id, spirv::GLOp::FClamp, expr_type_id, @@ -2671,7 +2695,7 @@ impl BlockContext<'_> { }); let clamp_id = self.gen_id(); - block.body.push(Instruction::ext_inst( + block.body.push(Instruction::ext_inst_gl_op( self.writer.gl450_ext_inst_id, clamp_op, wide_vector_type_id, @@ -2765,7 +2789,7 @@ impl BlockContext<'_> { let [min, max] = [min, max].map(|lit| self.writer.get_constant_scalar(lit)); let clamp_id = self.gen_id(); - block.body.push(Instruction::ext_inst( + block.body.push(Instruction::ext_inst_gl_op( self.writer.gl450_ext_inst_id, clamp_op, result_type_id, diff --git a/naga/src/back/spv/image.rs b/naga/src/back/spv/image.rs index 3aec1333f0c..78d7c79edfb 100644 --- a/naga/src/back/spv/image.rs +++ b/naga/src/back/spv/image.rs @@ -446,7 +446,7 @@ impl BlockContext<'_> { // and negative values in a single instruction: negative values of // `input_id` get treated as very large positive values. let restricted_id = self.gen_id(); - block.body.push(Instruction::ext_inst( + block.body.push(Instruction::ext_inst_gl_op( self.writer.gl450_ext_inst_id, spirv::GLOp::UMin, type_id, @@ -580,7 +580,7 @@ impl BlockContext<'_> { // and negative values in a single instruction: negative values of // `coordinates` get treated as very large positive values. let restricted_coordinates_id = self.gen_id(); - block.body.push(Instruction::ext_inst( + block.body.push(Instruction::ext_inst_gl_op( self.writer.gl450_ext_inst_id, spirv::GLOp::UMin, coordinates.type_id, @@ -923,7 +923,7 @@ impl BlockContext<'_> { // Clamp the coords to the calculated margins let clamped_coords_id = self.gen_id(); - block.body.push(Instruction::ext_inst( + block.body.push(Instruction::ext_inst_gl_op( self.writer.gl450_ext_inst_id, spirv::GLOp::NClamp, vec2f_type_id, diff --git a/naga/src/back/spv/index.rs b/naga/src/back/spv/index.rs index 3a15ee88060..85caa9457f0 100644 --- a/naga/src/back/spv/index.rs +++ b/naga/src/back/spv/index.rs @@ -366,7 +366,7 @@ impl BlockContext<'_> { // One or the other of the index or length is dynamic, so emit code for // BoundsCheckPolicy::Restrict. let restricted_index_id = self.gen_id(); - block.body.push(Instruction::ext_inst( + block.body.push(Instruction::ext_inst_gl_op( self.writer.gl450_ext_inst_id, spirv::GLOp::UMin, self.writer.get_u32_type_id(), diff --git a/naga/src/back/spv/instructions.rs b/naga/src/back/spv/instructions.rs index 788c3bc119a..a76ba72ed88 100644 --- a/naga/src/back/spv/instructions.rs +++ b/naga/src/back/spv/instructions.rs @@ -156,18 +156,28 @@ impl super::Instruction { instruction } - pub(super) fn ext_inst( + pub(super) fn ext_inst_gl_op( set_id: Word, op: spirv::GLOp, result_type_id: Word, id: Word, operands: &[Word], + ) -> Self { + Self::ext_inst(set_id, op as u32, result_type_id, id, operands) + } + + pub(super) fn ext_inst( + set_id: Word, + op: u32, + result_type_id: Word, + id: Word, + operands: &[Word], ) -> Self { let mut instruction = Self::new(Op::ExtInst); instruction.set_type(result_type_id); instruction.set_result(id); instruction.add_operand(set_id); - instruction.add_operand(op as u32); + instruction.add_operand(op); for operand in operands { instruction.add_operand(*operand) } @@ -824,6 +834,14 @@ impl super::Instruction { instruction } + pub(super) fn ray_query_get_t_min(result_type_id: Word, id: Word, query: Word) -> Self { + let mut instruction = Self::new(Op::RayQueryGetRayTMinKHR); + instruction.set_type(result_type_id); + instruction.set_result(id); + instruction.add_operand(query); + instruction + } + // // Conversion Instructions // diff --git a/naga/src/back/spv/mod.rs b/naga/src/back/spv/mod.rs index 4690dc71951..3d356112b2d 100644 --- a/naga/src/back/spv/mod.rs +++ b/naga/src/back/spv/mod.rs @@ -151,6 +151,12 @@ struct Function { signature: Option, parameters: Vec, variables: crate::FastHashMap, LocalVariable>, + /// Map from a local variable that is a ray query to its u32 tracker. + ray_query_initialization_tracker_variables: + crate::FastHashMap, LocalVariable>, + /// Map from a local variable that is a ray query to its tracker for the t max. + ray_query_t_max_tracker_variables: + crate::FastHashMap, LocalVariable>, /// List of local variables used as a counters to ensure that all loops are bounded. force_loop_bounding_vars: Vec, @@ -445,6 +451,16 @@ struct LookupFunctionType { return_type_id: Word, } +#[derive(Debug, PartialEq, Clone, Hash, Eq)] +enum LookupRayQueryFunction { + Initialize, + Proceed, + GenerateIntersection, + ConfirmIntersection, + GetVertexPositions { committed: bool }, + GetIntersection { committed: bool }, +} + #[derive(Debug)] enum Dimension { Scalar, @@ -685,6 +701,21 @@ struct BlockContext<'w> { expression_constness: ExpressionConstnessTracker, force_loop_bounding: bool, + + /// Hash from an expression whose type is a ray query / pointer to a ray query to its tracker. + /// Note: this is sparse, so can't be a handle vec + ray_query_tracker_expr: crate::FastHashMap, RayQueryTrackers>, +} + +#[derive(Clone, Copy)] +struct RayQueryTrackers { + // Initialization tracker + initialized_tracker: Word, + // Tracks the t max from ray query initialize. + // Unlike HLSL, spir-v's equivalent getter for the current committed t has UB (instead of just + // returning t_max) if there was no previous hit (though in some places it treats the behaviour as + // defined), therefore we must track the tmax inputted into ray query initialize. + t_max_tracker: Word, } impl BlockContext<'_> { @@ -741,6 +772,7 @@ pub struct Writer { /// The set of spirv extensions used. extensions_used: crate::FastIndexSet<&'static str>, + debug_strings: Vec, debugs: Vec, annotations: Vec, flags: WriterFlags, @@ -773,12 +805,15 @@ pub struct Writer { // Just a temporary list of SPIR-V ids temp_list: Vec, - ray_get_committed_intersection_function: Option, - ray_get_candidate_intersection_function: Option, + ray_query_functions: crate::FastHashMap, /// F16 I/O polyfill manager for handling `f16` input/output variables /// when `StorageInputOutput16` capability is not available. io_f16_polyfills: f16_polyfill::F16IoPolyfill, + + /// Non semantic debug printf extension `OpExtInstImport` + debug_printf: Option, + pub(crate) ray_query_initialization_tracking: bool, } bitflags::bitflags! { @@ -810,6 +845,26 @@ bitflags::bitflags! { /// /// [`BuiltIn::FragDepth`]: crate::BuiltIn::FragDepth const CLAMP_FRAG_DEPTH = 0x10; + + /// Instead of silently failing if the arguments to generate a ray query are + /// invalid, uses debug printf extension to print to the command line + /// + /// Note: VK_KHR_shader_non_semantic_info must be enabled. This will have no + /// effect if `options.ray_query_initialization_tracking` is set to false. + const PRINT_ON_RAY_QUERY_INITIALIZATION_FAIL = 0x20; + } +} + +bitflags::bitflags! { + /// How far through a ray query are we + #[derive(Clone, Copy, Debug, Eq, PartialEq)] + pub(super) struct RayQueryPoint: u32 { + /// Ray query has been successfully initialized. + const INITIALIZED = 1 << 0; + /// Proceed has been called on ray query. + const PROCEED = 1 << 1; + /// Proceed has returned false (have finished traversal). + const FINISHED_TRAVERSAL = 1 << 2; } } @@ -867,6 +922,10 @@ pub struct Options<'a> { /// to think the number of iterations is bounded. pub force_loop_bounding: bool, + /// if set, ray queries will get a variable to track their state to prevent + /// misuse. + pub ray_query_initialization_tracking: bool, + /// Whether to use the `StorageInputOutput16` capability for `f16` shader I/O. /// When false, `f16` I/O is polyfilled using `f32` types with conversions. pub use_storage_input_output_16: bool, @@ -891,6 +950,7 @@ impl Default for Options<'_> { bounds_check_policies: BoundsCheckPolicies::default(), zero_initialize_workgroup_memory: ZeroInitializeWorkgroupMemoryMode::Polyfill, force_loop_bounding: true, + ray_query_initialization_tracking: true, use_storage_input_output_16: true, debug_info: None, } diff --git a/naga/src/back/spv/ray.rs b/naga/src/back/spv/ray.rs index 05a55c78d83..405b29b8347 100644 --- a/naga/src/back/spv/ray.rs +++ b/naga/src/back/spv/ray.rs @@ -2,27 +2,123 @@ Generating SPIR-V for ray query operations. */ -use alloc::vec; +use alloc::{vec, vec::Vec}; use super::{ Block, BlockContext, Function, FunctionArgument, Instruction, LookupFunctionType, NumericType, Writer, }; -use crate::arena::Handle; +use crate::{arena::Handle, back::spv::LookupRayQueryFunction}; + +/// helper function to check if a particular flag is set in a u32. +fn write_ray_flags_contains_flags( + writer: &mut Writer, + block: &mut Block, + id: spirv::Word, + flag: u32, +) -> spirv::Word { + let bit_id = writer.get_constant_scalar(crate::Literal::U32(flag)); + let zero_id = writer.get_constant_scalar(crate::Literal::U32(0)); + let u32_type_id = writer.get_u32_type_id(); + let bool_ty = writer.get_bool_type_id(); + + let and_id = writer.id_gen.next(); + block.body.push(Instruction::binary( + spirv::Op::BitwiseAnd, + u32_type_id, + and_id, + id, + bit_id, + )); + + let eq_id = writer.id_gen.next(); + block.body.push(Instruction::binary( + spirv::Op::INotEqual, + bool_ty, + eq_id, + and_id, + zero_id, + )); + + eq_id +} impl Writer { + /// writes a logical and of two scalar booleans + fn write_logical_and( + &mut self, + block: &mut Block, + one: spirv::Word, + two: spirv::Word, + ) -> spirv::Word { + let id = self.id_gen.next(); + let bool_id = self.get_bool_type_id(); + block.body.push(Instruction::binary( + spirv::Op::LogicalAnd, + bool_id, + id, + one, + two, + )); + id + } + + fn write_reduce_and(&mut self, block: &mut Block, mut bools: Vec) -> spirv::Word { + // The combined `and`ed together of all of the bools up to this point. + let mut current_combined = bools.pop().unwrap(); + for boolean in bools { + current_combined = self.write_logical_and(block, current_combined, boolean) + } + current_combined + } + + // returns the id of the function, the function, and ids for its arguments. + fn write_function_signature( + &mut self, + arg_types: &[spirv::Word], + return_ty: spirv::Word, + ) -> (spirv::Word, Function, Vec) { + let func_ty = self.get_function_type(LookupFunctionType { + parameter_type_ids: Vec::from(arg_types), + return_type_id: return_ty, + }); + + let mut function = Function::default(); + let func_id = self.id_gen.next(); + function.signature = Some(Instruction::function( + return_ty, + func_id, + spirv::FunctionControl::empty(), + func_ty, + )); + + let mut arg_ids = Vec::with_capacity(arg_types.len()); + + for (idx, &arg_ty) in arg_types.iter().enumerate() { + let id = self.id_gen.next(); + let instruction = Instruction::function_parameter(arg_ty, id); + function.parameters.push(FunctionArgument { + instruction, + handle_id: idx as u32, + }); + arg_ids.push(id); + } + (func_id, function, arg_ids) + } + pub(super) fn write_ray_query_get_intersection_function( &mut self, is_committed: bool, ir_module: &crate::Module, ) -> spirv::Word { - if is_committed { - if let Some(func_id) = self.ray_get_committed_intersection_function { - return func_id; - } - } else if let Some(func_id) = self.ray_get_candidate_intersection_function { - return func_id; - }; + if let Some(&word) = + self.ray_query_functions + .get(&LookupRayQueryFunction::GetIntersection { + committed: is_committed, + }) + { + return word; + } let ray_intersection = ir_module.special_types.ray_intersection.unwrap(); let intersection_type_id = self.get_handle_type_id(ray_intersection); let intersection_pointer_type_id = @@ -56,31 +152,20 @@ impl Writer { let argument_type_id = self.get_ray_query_pointer_id(); - let func_ty = self.get_function_type(LookupFunctionType { - parameter_type_ids: vec![argument_type_id], - return_type_id: intersection_type_id, - }); - - let mut function = Function::default(); - let func_id = self.id_gen.next(); - function.signature = Some(Instruction::function( + let (func_id, mut function, arg_ids) = self.write_function_signature( + &[argument_type_id, flag_pointer_type_id], intersection_type_id, - func_id, - spirv::FunctionControl::empty(), - func_ty, - )); - let blank_intersection = self.get_constant_null(intersection_type_id); - let query_id = self.id_gen.next(); - let instruction = Instruction::function_parameter(argument_type_id, query_id); - function.parameters.push(FunctionArgument { - instruction, - handle_id: 0, - }); + ); + + let query_id = arg_ids[0]; + let intersection_tracker_id = arg_ids[1]; let label_id = self.id_gen.next(); let mut block = Block::new(label_id); + let blank_intersection = self.get_constant_null(intersection_type_id); let blank_intersection_id = self.id_gen.next(); + // This must be before everything else in the function. block.body.push(Instruction::variable( intersection_pointer_type_id, blank_intersection_id, @@ -93,14 +178,67 @@ impl Writer { } else { spirv::RayQueryIntersection::RayQueryCandidateIntersectionKHR } as _)); - let raw_kind_id = self.id_gen.next(); - block.body.push(Instruction::ray_query_get_intersection( - spirv::Op::RayQueryGetIntersectionTypeKHR, + + let loaded_ray_query_tracker_id = self.id_gen.next(); + block.body.push(Instruction::load( flag_type_id, - raw_kind_id, - query_id, - intersection_id, + loaded_ray_query_tracker_id, + intersection_tracker_id, + None, + )); + let proceeded_id = write_ray_flags_contains_flags( + self, + &mut block, + loaded_ray_query_tracker_id, + super::RayQueryPoint::PROCEED.bits(), + ); + let finished_proceed_id = write_ray_flags_contains_flags( + self, + &mut block, + loaded_ray_query_tracker_id, + super::RayQueryPoint::FINISHED_TRAVERSAL.bits(), + ); + let proceed_finished_correct_id = if is_committed { + finished_proceed_id + } else { + let not_finished_id = self.id_gen.next(); + block.body.push(Instruction::unary( + spirv::Op::LogicalNot, + bool_type_id, + not_finished_id, + finished_proceed_id, + )); + not_finished_id + }; + + let is_valid_id = + self.write_logical_and(&mut block, proceed_finished_correct_id, proceeded_id); + + let valid_id = self.id_gen.next(); + let mut valid_block = Block::new(valid_id); + + let final_label_id = self.id_gen.next(); + let mut final_block = Block::new(final_label_id); + + block.body.push(Instruction::selection_merge( + final_label_id, + spirv::SelectionControl::NONE, )); + function.consume( + block, + Instruction::branch_conditional(is_valid_id, valid_id, final_label_id), + ); + + let raw_kind_id = self.id_gen.next(); + valid_block + .body + .push(Instruction::ray_query_get_intersection( + spirv::Op::RayQueryGetIntersectionTypeKHR, + flag_type_id, + raw_kind_id, + query_id, + intersection_id, + )); let kind_id = if is_committed { // Nothing to do: the IR value matches `spirv::RayQueryCommittedIntersectionType` raw_kind_id @@ -111,7 +249,7 @@ impl Writer { spirv::RayQueryCandidateIntersectionType::RayQueryCandidateIntersectionTriangleKHR as _, )); - block.body.push(Instruction::binary( + valid_block.body.push(Instruction::binary( spirv::Op::IEqual, self.get_bool_type_id(), condition_id, @@ -119,7 +257,7 @@ impl Writer { committed_triangle_kind_id, )); let kind_id = self.id_gen.next(); - block.body.push(Instruction::select( + valid_block.body.push(Instruction::select( flag_type_id, kind_id, condition_id, @@ -134,20 +272,20 @@ impl Writer { }; let idx_id = self.get_index_constant(0); let access_idx = self.id_gen.next(); - block.body.push(Instruction::access_chain( + valid_block.body.push(Instruction::access_chain( flag_pointer_type_id, access_idx, blank_intersection_id, &[idx_id], )); - block + valid_block .body .push(Instruction::store(access_idx, kind_id, None)); let not_none_comp_id = self.id_gen.next(); let none_id = self.get_constant_scalar(crate::Literal::U32(crate::RayQueryIntersection::None as _)); - block.body.push(Instruction::binary( + valid_block.body.push(Instruction::binary( spirv::Op::INotEqual, self.get_bool_type_id(), not_none_comp_id, @@ -158,16 +296,20 @@ impl Writer { let not_none_label_id = self.id_gen.next(); let mut not_none_block = Block::new(not_none_label_id); - let final_label_id = self.id_gen.next(); - let mut final_block = Block::new(final_label_id); + let outer_merge_label_id = self.id_gen.next(); + let outer_merge_block = Block::new(outer_merge_label_id); - block.body.push(Instruction::selection_merge( - final_label_id, + valid_block.body.push(Instruction::selection_merge( + outer_merge_label_id, spirv::SelectionControl::NONE, )); function.consume( - block, - Instruction::branch_conditional(not_none_comp_id, not_none_label_id, final_label_id), + valid_block, + Instruction::branch_conditional( + not_none_comp_id, + not_none_label_id, + outer_merge_label_id, + ), ); let instance_custom_index_id = self.id_gen.next(); @@ -426,7 +568,8 @@ impl Writer { .body .push(Instruction::store(access_idx, front_face_id, None)); function.consume(tri_block, Instruction::branch(merge_label_id)); - function.consume(merge_block, Instruction::branch(final_label_id)); + function.consume(merge_block, Instruction::branch(outer_merge_label_id)); + function.consume(outer_merge_block, Instruction::branch(final_label_id)); let loaded_blank_intersection_id = self.id_gen.next(); final_block.body.push(Instruction::load( @@ -441,151 +584,1321 @@ impl Writer { ); function.to_words(&mut self.logical_layout.function_definitions); - if is_committed { - self.ray_get_committed_intersection_function = Some(func_id); - } else { - self.ray_get_candidate_intersection_function = Some(func_id); - } + self.ray_query_functions.insert( + LookupRayQueryFunction::GetIntersection { + committed: is_committed, + }, + func_id, + ); func_id } -} -impl BlockContext<'_> { - pub(super) fn write_ray_query_function( - &mut self, - query: Handle, - function: &crate::RayQueryFunction, - block: &mut Block, - ) { - let query_id = self.cached[query]; - match *function { - crate::RayQueryFunction::Initialize { - acceleration_structure, - descriptor, - } => { - //Note: composite extract indices and types must match `generate_ray_desc_type` - let desc_id = self.cached[descriptor]; - let acc_struct_id = self.get_handle_id(acceleration_structure); + fn write_ray_query_initialize(&mut self, ir_module: &crate::Module) -> spirv::Word { + if let Some(&word) = self + .ray_query_functions + .get(&LookupRayQueryFunction::Initialize) + { + return word; + } - let flag_type_id = - self.get_numeric_type_id(NumericType::Scalar(crate::Scalar::U32)); - let ray_flags_id = self.gen_id(); - block.body.push(Instruction::composite_extract( - flag_type_id, - ray_flags_id, - desc_id, - &[0], - )); - let cull_mask_id = self.gen_id(); - block.body.push(Instruction::composite_extract( - flag_type_id, - cull_mask_id, - desc_id, - &[1], - )); + let ray_query_type_id = self.get_ray_query_pointer_id(); + let acceleration_structure_type_id = + self.get_localtype_id(super::LocalType::AccelerationStructure); + let ray_desc_type_id = self.get_handle_type_id( + ir_module + .special_types + .ray_desc + .expect("ray desc should be set if ray queries are being initialized"), + ); - let scalar_type_id = - self.get_numeric_type_id(NumericType::Scalar(crate::Scalar::F32)); - let tmin_id = self.gen_id(); - block.body.push(Instruction::composite_extract( - scalar_type_id, - tmin_id, - desc_id, - &[2], - )); - let tmax_id = self.gen_id(); - block.body.push(Instruction::composite_extract( - scalar_type_id, - tmax_id, - desc_id, - &[3], - )); + let u32_ty = self.get_u32_type_id(); + let u32_ptr_ty = self.get_pointer_type_id(u32_ty, spirv::StorageClass::Function); - let vector_type_id = self.get_numeric_type_id(NumericType::Vector { - size: crate::VectorSize::Tri, - scalar: crate::Scalar::F32, - }); - let ray_origin_id = self.gen_id(); - block.body.push(Instruction::composite_extract( - vector_type_id, - ray_origin_id, - desc_id, - &[4], - )); - let ray_dir_id = self.gen_id(); - block.body.push(Instruction::composite_extract( - vector_type_id, - ray_dir_id, - desc_id, - &[5], - )); + let f32_type_id = self.get_f32_type_id(); + let f32_ptr_ty = self.get_pointer_type_id(f32_type_id, spirv::StorageClass::Function); - block.body.push(Instruction::ray_query_initialize( - query_id, - acc_struct_id, - ray_flags_id, - cull_mask_id, - ray_origin_id, - tmin_id, - ray_dir_id, - tmax_id, + let bool_type_id = self.get_bool_type_id(); + let bool_vec3_type_id = self.get_vec3_bool_type_id(); + + let (func_id, mut function, arg_ids) = self.write_function_signature( + &[ + ray_query_type_id, + acceleration_structure_type_id, + ray_desc_type_id, + u32_ptr_ty, + f32_ptr_ty, + ], + self.void_type, + ); + + let query_id = arg_ids[0]; + let acceleration_structure_id = arg_ids[1]; + let desc_id = arg_ids[2]; + let init_tracker_id = arg_ids[3]; + let t_max_tracker_id = arg_ids[4]; + + let label_id = self.id_gen.next(); + let mut block = Block::new(label_id); + + let flag_type_id = self.get_numeric_type_id(NumericType::Scalar(crate::Scalar::U32)); + + //Note: composite extract indices and types must match `generate_ray_desc_type` + let ray_flags_id = self.id_gen.next(); + block.body.push(Instruction::composite_extract( + flag_type_id, + ray_flags_id, + desc_id, + &[0], + )); + let cull_mask_id = self.id_gen.next(); + block.body.push(Instruction::composite_extract( + flag_type_id, + cull_mask_id, + desc_id, + &[1], + )); + + let tmin_id = self.id_gen.next(); + block.body.push(Instruction::composite_extract( + f32_type_id, + tmin_id, + desc_id, + &[2], + )); + let tmax_id = self.id_gen.next(); + block.body.push(Instruction::composite_extract( + f32_type_id, + tmax_id, + desc_id, + &[3], + )); + block + .body + .push(Instruction::store(t_max_tracker_id, tmax_id, None)); + + let vector_type_id = self.get_numeric_type_id(NumericType::Vector { + size: crate::VectorSize::Tri, + scalar: crate::Scalar::F32, + }); + let ray_origin_id = self.id_gen.next(); + block.body.push(Instruction::composite_extract( + vector_type_id, + ray_origin_id, + desc_id, + &[4], + )); + let ray_dir_id = self.id_gen.next(); + block.body.push(Instruction::composite_extract( + vector_type_id, + ray_dir_id, + desc_id, + &[5], + )); + + let valid_id = self.ray_query_initialization_tracking.then(||{ + let tmin_le_tmax_id = self.id_gen.next(); + // Check both that tmin is less than or equal to tmax (https://docs.vulkan.org/spec/latest/appendices/spirvenv.html#VUID-RuntimeSpirv-OpRayQueryInitializeKHR-06350) + // and implicitly that neither tmin or tmax are NaN (https://docs.vulkan.org/spec/latest/appendices/spirvenv.html#VUID-RuntimeSpirv-OpRayQueryInitializeKHR-06351) + // because this checks if tmin and tmax are ordered too (i.e: not NaN). + block.body.push(Instruction::binary( + spirv::Op::FOrdLessThanEqual, + bool_type_id, + tmin_le_tmax_id, + tmin_id, + tmax_id, + )); + + // Check that tmin is greater than or equal to 0 (and + // therefore also tmax is too because it is greater than + // or equal to tmin) (https://docs.vulkan.org/spec/latest/appendices/spirvenv.html#VUID-RuntimeSpirv-OpRayQueryInitializeKHR-06349). + let tmin_ge_zero_id = self.id_gen.next(); + let zero_id = self.get_constant_scalar(crate::Literal::F32(0.0)); + block.body.push(Instruction::binary( + spirv::Op::FOrdGreaterThanEqual, + bool_type_id, + tmin_ge_zero_id, + tmin_id, + zero_id, + )); + + // Check that ray origin is finite (https://docs.vulkan.org/spec/latest/appendices/spirvenv.html#VUID-RuntimeSpirv-OpRayQueryInitializeKHR-06348) + let ray_origin_infinite_id = self.id_gen.next(); + block.body.push(Instruction::unary( + spirv::Op::IsInf, + bool_vec3_type_id, + ray_origin_infinite_id, + ray_origin_id, + )); + let any_ray_origin_infinite_id = self.id_gen.next(); + block.body.push(Instruction::unary( + spirv::Op::Any, + bool_type_id, + any_ray_origin_infinite_id, + ray_origin_infinite_id, + )); + + let ray_origin_nan_id = self.id_gen.next(); + block.body.push(Instruction::unary( + spirv::Op::IsNan, + bool_vec3_type_id, + ray_origin_nan_id, + ray_origin_id, + )); + let any_ray_origin_nan_id = self.id_gen.next(); + block.body.push(Instruction::unary( + spirv::Op::Any, + bool_type_id, + any_ray_origin_nan_id, + ray_origin_nan_id, + )); + + let ray_origin_not_finite_id = self.id_gen.next(); + block.body.push(Instruction::binary( + spirv::Op::LogicalOr, + bool_type_id, + ray_origin_not_finite_id, + any_ray_origin_nan_id, + any_ray_origin_infinite_id, + )); + + let all_ray_origin_finite_id = self.id_gen.next(); + block.body.push(Instruction::unary( + spirv::Op::LogicalNot, + bool_type_id, + all_ray_origin_finite_id, + ray_origin_not_finite_id, + )); + + // Check that ray direction is finite (https://docs.vulkan.org/spec/latest/appendices/spirvenv.html#VUID-RuntimeSpirv-OpRayQueryInitializeKHR-06348) + let ray_dir_infinite_id = self.id_gen.next(); + block.body.push(Instruction::unary( + spirv::Op::IsInf, + bool_vec3_type_id, + ray_dir_infinite_id, + ray_dir_id, + )); + let any_ray_dir_infinite_id = self.id_gen.next(); + block.body.push(Instruction::unary( + spirv::Op::Any, + bool_type_id, + any_ray_dir_infinite_id, + ray_dir_infinite_id, + )); + + let ray_dir_nan_id = self.id_gen.next(); + block.body.push(Instruction::unary( + spirv::Op::IsNan, + bool_vec3_type_id, + ray_dir_nan_id, + ray_dir_id, + )); + let any_ray_dir_nan_id = self.id_gen.next(); + block.body.push(Instruction::unary( + spirv::Op::Any, + bool_type_id, + any_ray_dir_nan_id, + ray_dir_nan_id, + )); + + let ray_dir_not_finite_id = self.id_gen.next(); + block.body.push(Instruction::binary( + spirv::Op::LogicalOr, + bool_type_id, + ray_dir_not_finite_id, + any_ray_dir_nan_id, + any_ray_dir_infinite_id, + )); + + let all_ray_dir_finite_id = self.id_gen.next(); + block.body.push(Instruction::unary( + spirv::Op::LogicalNot, + bool_type_id, + all_ray_dir_finite_id, + ray_dir_not_finite_id, + )); + + /// Writes spirv to check that less than two booleans are true + /// + /// For each boolean: removes it, `and`s it with all others (i.e for all possible combinations of two booleans in the list checks to see if both are true). + /// Then `or`s all of these checks together. This produces whether two or more booleans are true. + fn write_less_than_2_true( + writer: &mut Writer, + block: &mut Block, + mut bools: Vec, + ) -> spirv::Word { + assert!(bools.len() > 1, "Must have multiple booleans!"); + let bool_ty = writer.get_bool_type_id(); + let mut each_two_true = Vec::new(); + while let Some(last_bool) = bools.pop() { + for &bool in &bools { + let both_true_id = writer.write_logical_and( + block, + last_bool, + bool, + ); + each_two_true.push(both_true_id); + } + } + let mut all_or_id = each_two_true.pop().expect("since this must have multiple booleans, there must be at least one thing in `each_two_true`"); + for two_true in each_two_true { + let new_all_or_id = writer.id_gen.next(); + block.body.push(Instruction::binary( + spirv::Op::LogicalOr, + bool_ty, + new_all_or_id, + all_or_id, + two_true, + )); + all_or_id = new_all_or_id; + } + + let less_than_two_id = writer.id_gen.next(); + block.body.push(Instruction::unary( + spirv::Op::LogicalNot, + bool_ty, + less_than_two_id, + all_or_id, )); + less_than_two_id } - crate::RayQueryFunction::Proceed { result } => { - let id = self.gen_id(); - self.cached[result] = id; - let result_type_id = self.get_expression_type_id(&self.fun_info[result].ty); - block - .body - .push(Instruction::ray_query_proceed(result_type_id, id, query_id)); - } - crate::RayQueryFunction::GenerateIntersection { hit_t } => { - let hit_id = self.cached[hit_t]; - block - .body - .push(Instruction::ray_query_generate_intersection( - query_id, hit_id, - )); + // Check that at most one of skip triangles and skip AABBs is + // present (https://docs.vulkan.org/spec/latest/appendices/spirvenv.html#VUID-RuntimeSpirv-OpRayQueryInitializeKHR-06889) + let contains_skip_triangles = write_ray_flags_contains_flags( + self, + &mut block, + ray_flags_id, + crate::RayFlag::SKIP_TRIANGLES.bits(), + ); + let contains_skip_aabbs = write_ray_flags_contains_flags( + self, + &mut block, + ray_flags_id, + crate::RayFlag::SKIP_AABBS.bits(), + ); + + let not_contain_skip_triangles_aabbs = write_less_than_2_true( + self, + &mut block, + vec![contains_skip_triangles, contains_skip_aabbs], + ); + + // Check that at most one of skip triangles (taken from above check), + // cull back facing, and cull front face is present (https://docs.vulkan.org/spec/latest/appendices/spirvenv.html#VUID-RuntimeSpirv-OpRayQueryInitializeKHR-06890) + let contains_cull_back = write_ray_flags_contains_flags( + self, + &mut block, + ray_flags_id, + crate::RayFlag::CULL_BACK_FACING.bits(), + ); + let contains_cull_front = write_ray_flags_contains_flags( + self, + &mut block, + ray_flags_id, + crate::RayFlag::CULL_FRONT_FACING.bits(), + ); + + let not_contain_skip_triangles_cull = write_less_than_2_true( + self, + &mut block, + vec![ + contains_skip_triangles, + contains_cull_back, + contains_cull_front, + ], + ); + + // Check that at most one of force opaque, force not opaque, cull opaque, + // and cull not opaque are present (https://docs.vulkan.org/spec/latest/appendices/spirvenv.html#VUID-RuntimeSpirv-OpRayQueryInitializeKHR-06891) + let contains_opaque = write_ray_flags_contains_flags( + self, + &mut block, + ray_flags_id, + crate::RayFlag::FORCE_OPAQUE.bits(), + ); + let contains_no_opaque = write_ray_flags_contains_flags( + self, + &mut block, + ray_flags_id, + crate::RayFlag::FORCE_NO_OPAQUE.bits(), + ); + let contains_cull_opaque = write_ray_flags_contains_flags( + self, + &mut block, + ray_flags_id, + crate::RayFlag::CULL_OPAQUE.bits(), + ); + let contains_cull_no_opaque = write_ray_flags_contains_flags( + self, + &mut block, + ray_flags_id, + crate::RayFlag::CULL_NO_OPAQUE.bits(), + ); + + let not_contain_multiple_opaque = write_less_than_2_true( + self, + &mut block, + vec![ + contains_opaque, + contains_no_opaque, + contains_cull_opaque, + contains_cull_no_opaque, + ], + ); + + // Combine all checks into a single flag saying whether the call is valid or not. + self.write_reduce_and( + &mut block, + vec![ + tmin_le_tmax_id, + tmin_ge_zero_id, + all_ray_origin_finite_id, + all_ray_dir_finite_id, + not_contain_skip_triangles_aabbs, + not_contain_skip_triangles_cull, + not_contain_multiple_opaque, + ], + ) + }); + + let merge_label_id = self.id_gen.next(); + let merge_block = Block::new(merge_label_id); + + // NOTE: this block will be unreachable if initialization tracking is disabled. + let invalid_label_id = self.id_gen.next(); + let mut invalid_block = Block::new(invalid_label_id); + + let valid_label_id = self.id_gen.next(); + let mut valid_block = Block::new(valid_label_id); + + match valid_id { + Some(all_valid_id) => { + block.body.push(Instruction::selection_merge( + merge_label_id, + spirv::SelectionControl::NONE, + )); + function.consume( + block, + Instruction::branch_conditional(all_valid_id, valid_label_id, invalid_label_id), + ); } - crate::RayQueryFunction::ConfirmIntersection => { - block - .body - .push(Instruction::ray_query_confirm_intersection(query_id)); + None => { + function.consume(block, Instruction::branch(valid_label_id)); } - crate::RayQueryFunction::Terminate => {} } - } - pub(super) fn write_ray_query_return_vertex_position( - &mut self, - query: Handle, - block: &mut Block, - is_committed: bool, - ) -> spirv::Word { - let query_id = self.cached[query]; - let id = self.gen_id(); - let ray_vertex_return_ty = self - .ir_module - .special_types - .ray_vertex_return - .expect("type should have been populated"); - let ray_vertex_return_ty_id = self.writer.get_handle_type_id(ray_vertex_return_ty); - let intersection_id = - self.writer - .get_constant_scalar(crate::Literal::U32(if is_committed { - spirv::RayQueryIntersection::RayQueryCommittedIntersectionKHR - } else { - spirv::RayQueryIntersection::RayQueryCandidateIntersectionKHR - } as _)); - block + valid_block.body.push(Instruction::ray_query_initialize( + query_id, + acceleration_structure_id, + ray_flags_id, + cull_mask_id, + ray_origin_id, + tmin_id, + ray_dir_id, + tmax_id, + )); + + let const_initialized = self.get_constant_scalar(crate::Literal::U32( + super::RayQueryPoint::INITIALIZED.bits(), + )); + valid_block .body - .push(Instruction::ray_query_return_vertex_position( - ray_vertex_return_ty_id, - id, - query_id, - intersection_id, - )); - id + .push(Instruction::store(init_tracker_id, const_initialized, None)); + + function.consume(valid_block, Instruction::branch(merge_label_id)); + + if self + .flags + .contains(super::WriterFlags::PRINT_ON_RAY_QUERY_INITIALIZATION_FAIL) + { + self.write_debug_printf( + &mut invalid_block, + "Naga ignored invalid arguments to rayQueryInitialize with flags: %u t_min: %f t_max: %f origin: %v4f dir: %v4f", + &[ + ray_flags_id, + tmin_id, + tmax_id, + ray_origin_id, + ray_dir_id, + ], + ); + } + + function.consume(invalid_block, Instruction::branch(merge_label_id)); + + function.consume(merge_block, Instruction::return_void()); + + function.to_words(&mut self.logical_layout.function_definitions); + + self.ray_query_functions + .insert(LookupRayQueryFunction::Initialize, func_id); + func_id + } + + fn write_ray_query_proceed(&mut self) -> spirv::Word { + if let Some(&word) = self + .ray_query_functions + .get(&LookupRayQueryFunction::Proceed) + { + return word; + } + + let ray_query_type_id = self.get_ray_query_pointer_id(); + + let u32_ty = self.get_u32_type_id(); + let u32_ptr_ty = self.get_pointer_type_id(u32_ty, spirv::StorageClass::Function); + + let bool_type_id = self.get_bool_type_id(); + let bool_ptr_ty = self.get_pointer_type_id(bool_type_id, spirv::StorageClass::Function); + + let (func_id, mut function, arg_ids) = + self.write_function_signature(&[ray_query_type_id, u32_ptr_ty], bool_type_id); + + let query_id = arg_ids[0]; + let init_tracker_id = arg_ids[1]; + + let block_id = self.id_gen.next(); + let mut block = Block::new(block_id); + + // TODO: perhaps this could be replaced with an OpPhi? + let proceeded_id = self.id_gen.next(); + let const_false = self.get_constant_scalar(crate::Literal::Bool(false)); + block.body.push(Instruction::variable( + bool_ptr_ty, + proceeded_id, + spirv::StorageClass::Function, + Some(const_false), + )); + + let initialized_tracker_id = self.id_gen.next(); + block.body.push(Instruction::load( + u32_ty, + initialized_tracker_id, + init_tracker_id, + None, + )); + + let merge_id = self.id_gen.next(); + let mut merge_block = Block::new(merge_id); + + let valid_block_id = self.id_gen.next(); + let mut valid_block = Block::new(valid_block_id); + + let instruction = if self.ray_query_initialization_tracking { + let is_initialized = write_ray_flags_contains_flags( + self, + &mut block, + initialized_tracker_id, + super::RayQueryPoint::INITIALIZED.bits(), + ); + + block.body.push(Instruction::selection_merge( + merge_id, + spirv::SelectionControl::NONE, + )); + + Instruction::branch_conditional(is_initialized, valid_block_id, merge_id) + } else { + Instruction::branch(valid_block_id) + }; + + function.consume(block, instruction); + + let has_proceeded = self.id_gen.next(); + valid_block.body.push(Instruction::ray_query_proceed( + bool_type_id, + has_proceeded, + query_id, + )); + + valid_block + .body + .push(Instruction::store(proceeded_id, has_proceeded, None)); + + let add_flag_finished = self.get_constant_scalar(crate::Literal::U32( + (super::RayQueryPoint::PROCEED | super::RayQueryPoint::FINISHED_TRAVERSAL).bits(), + )); + let add_flag_continuing = + self.get_constant_scalar(crate::Literal::U32(super::RayQueryPoint::PROCEED.bits())); + + let add_flags_id = self.id_gen.next(); + valid_block.body.push(Instruction::select( + u32_ty, + add_flags_id, + has_proceeded, + add_flag_continuing, + add_flag_finished, + )); + let final_flags = self.id_gen.next(); + valid_block.body.push(Instruction::binary( + spirv::Op::BitwiseOr, + u32_ty, + final_flags, + initialized_tracker_id, + add_flags_id, + )); + valid_block + .body + .push(Instruction::store(init_tracker_id, final_flags, None)); + + function.consume(valid_block, Instruction::branch(merge_id)); + + let loaded_proceeded_id = self.id_gen.next(); + merge_block.body.push(Instruction::load( + bool_type_id, + loaded_proceeded_id, + proceeded_id, + None, + )); + + function.consume(merge_block, Instruction::return_value(loaded_proceeded_id)); + + function.to_words(&mut self.logical_layout.function_definitions); + + self.ray_query_functions + .insert(LookupRayQueryFunction::Proceed, func_id); + func_id + } + + fn write_ray_query_generate_intersection(&mut self) -> spirv::Word { + if let Some(&word) = self + .ray_query_functions + .get(&LookupRayQueryFunction::GenerateIntersection) + { + return word; + } + + let ray_query_type_id = self.get_ray_query_pointer_id(); + + let u32_ty = self.get_u32_type_id(); + let u32_ptr_ty = self.get_pointer_type_id(u32_ty, spirv::StorageClass::Function); + + let f32_type_id = self.get_f32_type_id(); + let f32_ptr_type_id = self.get_pointer_type_id(f32_type_id, spirv::StorageClass::Function); + + let bool_type_id = self.get_bool_type_id(); + + let (func_id, mut function, arg_ids) = self.write_function_signature( + &[ray_query_type_id, u32_ptr_ty, f32_type_id, f32_ptr_type_id], + self.void_type, + ); + + let query_id = arg_ids[0]; + let init_tracker_id = arg_ids[1]; + let depth_id = arg_ids[2]; + let t_max_tracker_id = arg_ids[3]; + + let block_id = self.id_gen.next(); + let mut block = Block::new(block_id); + + let current_t = self.id_gen.next(); + block.body.push(Instruction::variable( + f32_ptr_type_id, + current_t, + spirv::StorageClass::Function, + None, + )); + + let current_t = self.id_gen.next(); + block.body.push(Instruction::variable( + f32_ptr_type_id, + current_t, + spirv::StorageClass::Function, + None, + )); + + let valid_id = self.id_gen.next(); + let mut valid_block = Block::new(valid_id); + + let final_label_id = self.id_gen.next(); + let final_block = Block::new(final_label_id); + + let instruction = if self.ray_query_initialization_tracking { + let initialized_tracker_id = self.id_gen.next(); + block.body.push(Instruction::load( + u32_ty, + initialized_tracker_id, + init_tracker_id, + None, + )); + + let proceeded_id = write_ray_flags_contains_flags( + self, + &mut block, + initialized_tracker_id, + super::RayQueryPoint::PROCEED.bits(), + ); + let finished_proceed_id = write_ray_flags_contains_flags( + self, + &mut block, + initialized_tracker_id, + super::RayQueryPoint::FINISHED_TRAVERSAL.bits(), + ); + + // Can't find anything to suggest double calling this function is invalid. + + let not_finished_id = self.id_gen.next(); + block.body.push(Instruction::unary( + spirv::Op::LogicalNot, + bool_type_id, + not_finished_id, + finished_proceed_id, + )); + + let is_valid_id = self.write_logical_and(&mut block, not_finished_id, proceeded_id); + + block.body.push(Instruction::selection_merge( + final_label_id, + spirv::SelectionControl::NONE, + )); + + Instruction::branch_conditional(is_valid_id, valid_id, final_label_id) + } else { + Instruction::branch(valid_id) + }; + + function.consume(block, instruction); + + let intersection_id = self.get_constant_scalar(crate::Literal::U32( + spirv::RayQueryIntersection::RayQueryCandidateIntersectionKHR as _, + )); + let committed_intersection_id = self.get_constant_scalar(crate::Literal::U32( + spirv::RayQueryIntersection::RayQueryCommittedIntersectionKHR as _, + )); + let raw_kind_id = self.id_gen.next(); + valid_block + .body + .push(Instruction::ray_query_get_intersection( + spirv::Op::RayQueryGetIntersectionTypeKHR, + u32_ty, + raw_kind_id, + query_id, + intersection_id, + )); + + let candidate_aabb_id = self.get_constant_scalar(crate::Literal::U32( + spirv::RayQueryCandidateIntersectionType::RayQueryCandidateIntersectionAABBKHR as _, + )); + let intersection_aabb_id = self.id_gen.next(); + valid_block.body.push(Instruction::binary( + spirv::Op::IEqual, + bool_type_id, + intersection_aabb_id, + raw_kind_id, + candidate_aabb_id, + )); + + // Check that the provided t value is between t min and the current committed + // t value, (https://docs.vulkan.org/spec/latest/appendices/spirvenv.html#VUID-RuntimeSpirv-OpRayQueryGenerateIntersectionKHR-06353) + + // Get the tmin + let t_min_id = self.id_gen.next(); + valid_block.body.push(Instruction::ray_query_get_t_min( + f32_type_id, + t_min_id, + query_id, + )); + + // Get the current committed t, or tmax if no hit. + // Basically emulate HLSL's (easier) version + // Pseudo-code: + // ````wgsl + // // start of function + // var current_t:f32; + // ... + // let committed_type_id = RayQueryGetIntersectionTypeKHR(query_id); + // if committed_type_id == Committed_None { + // current_t = load(t_max_tracker); + // } else { + // current_t = RayQueryGetIntersectionTKHR(query_id); + // } + // ... + // ```` + + let committed_type_id = self.id_gen.next(); + valid_block + .body + .push(Instruction::ray_query_get_intersection( + spirv::Op::RayQueryGetIntersectionTypeKHR, + u32_ty, + committed_type_id, + query_id, + committed_intersection_id, + )); + + let no_committed = self.id_gen.next(); + valid_block.body.push(Instruction::binary( + spirv::Op::IEqual, + bool_type_id, + no_committed, + committed_type_id, + self.get_constant_scalar(crate::Literal::U32( + spirv::RayQueryCommittedIntersectionType::RayQueryCommittedIntersectionNoneKHR as _, + )), + )); + + let next_valid_block_id = self.id_gen.next(); + let no_committed_block_id = self.id_gen.next(); + let mut no_committed_block = Block::new(no_committed_block_id); + let committed_block_id = self.id_gen.next(); + let mut committed_block = Block::new(committed_block_id); + valid_block.body.push(Instruction::selection_merge( + next_valid_block_id, + spirv::SelectionControl::NONE, + )); + function.consume( + valid_block, + Instruction::branch_conditional( + no_committed, + no_committed_block_id, + committed_block_id, + ), + ); + + // Assign t_max to current_t + let t_max_id = self.id_gen.next(); + no_committed_block.body.push(Instruction::load( + f32_type_id, + t_max_id, + t_max_tracker_id, + None, + )); + no_committed_block + .body + .push(Instruction::store(current_t, t_max_id, None)); + function.consume(no_committed_block, Instruction::branch(next_valid_block_id)); + + // Assign t_current to current_t + let latest_t_id = self.id_gen.next(); + committed_block + .body + .push(Instruction::ray_query_get_intersection( + spirv::Op::RayQueryGetIntersectionTKHR, + f32_type_id, + latest_t_id, + query_id, + intersection_id, + )); + committed_block + .body + .push(Instruction::store(current_t, latest_t_id, None)); + function.consume(committed_block, Instruction::branch(next_valid_block_id)); + + let mut valid_block = Block::new(next_valid_block_id); + + let t_ge_t_min = self.id_gen.next(); + valid_block.body.push(Instruction::binary( + spirv::Op::FOrdGreaterThanEqual, + bool_type_id, + t_ge_t_min, + depth_id, + t_min_id, + )); + let t_current = self.id_gen.next(); + valid_block + .body + .push(Instruction::load(f32_type_id, t_current, current_t, None)); + let t_le_t_current = self.id_gen.next(); + valid_block.body.push(Instruction::binary( + spirv::Op::FOrdLessThanEqual, + bool_type_id, + t_le_t_current, + depth_id, + t_current, + )); + + let t_in_range = self.id_gen.next(); + valid_block.body.push(Instruction::binary( + spirv::Op::LogicalAnd, + bool_type_id, + t_in_range, + t_ge_t_min, + t_le_t_current, + )); + + let call_valid_id = self.id_gen.next(); + valid_block.body.push(Instruction::binary( + spirv::Op::LogicalAnd, + bool_type_id, + call_valid_id, + t_in_range, + intersection_aabb_id, + )); + + let generate_label_id = self.id_gen.next(); + let mut generate_block = Block::new(generate_label_id); + + let merge_label_id = self.id_gen.next(); + let merge_block = Block::new(merge_label_id); + + valid_block.body.push(Instruction::selection_merge( + merge_label_id, + spirv::SelectionControl::NONE, + )); + function.consume( + valid_block, + Instruction::branch_conditional(call_valid_id, generate_label_id, merge_label_id), + ); + + generate_block + .body + .push(Instruction::ray_query_generate_intersection( + query_id, depth_id, + )); + + function.consume(generate_block, Instruction::branch(merge_label_id)); + function.consume(merge_block, Instruction::branch(final_label_id)); + + function.consume(final_block, Instruction::return_void()); + + function.to_words(&mut self.logical_layout.function_definitions); + + self.ray_query_functions + .insert(LookupRayQueryFunction::GenerateIntersection, func_id); + func_id + } + + fn write_ray_query_confirm_intersection(&mut self) -> spirv::Word { + if let Some(&word) = self + .ray_query_functions + .get(&LookupRayQueryFunction::ConfirmIntersection) + { + return word; + } + + let ray_query_type_id = self.get_ray_query_pointer_id(); + + let u32_ty = self.get_u32_type_id(); + let u32_ptr_ty = self.get_pointer_type_id(u32_ty, spirv::StorageClass::Function); + + let bool_type_id = self.get_bool_type_id(); + + let (func_id, mut function, arg_ids) = + self.write_function_signature(&[ray_query_type_id, u32_ptr_ty], self.void_type); + + let query_id = arg_ids[0]; + let init_tracker_id = arg_ids[1]; + + let block_id = self.id_gen.next(); + let mut block = Block::new(block_id); + + let valid_id = self.id_gen.next(); + let mut valid_block = Block::new(valid_id); + + let final_label_id = self.id_gen.next(); + let final_block = Block::new(final_label_id); + + let instruction = if self.ray_query_initialization_tracking { + let initialized_tracker_id = self.id_gen.next(); + block.body.push(Instruction::load( + u32_ty, + initialized_tracker_id, + init_tracker_id, + None, + )); + + let proceeded_id = write_ray_flags_contains_flags( + self, + &mut block, + initialized_tracker_id, + super::RayQueryPoint::PROCEED.bits(), + ); + let finished_proceed_id = write_ray_flags_contains_flags( + self, + &mut block, + initialized_tracker_id, + super::RayQueryPoint::FINISHED_TRAVERSAL.bits(), + ); + // Although it seems strange to call this twice, I (Vecvec) can't find anything to suggest double calling this function is invalid. + let not_finished_id = self.id_gen.next(); + block.body.push(Instruction::unary( + spirv::Op::LogicalNot, + bool_type_id, + not_finished_id, + finished_proceed_id, + )); + + let is_valid_id = self.write_logical_and(&mut block, not_finished_id, proceeded_id); + + block.body.push(Instruction::selection_merge( + final_label_id, + spirv::SelectionControl::NONE, + )); + + Instruction::branch_conditional(is_valid_id, valid_id, final_label_id) + } else { + Instruction::branch(valid_id) + }; + + function.consume(block, instruction); + + let intersection_id = self.get_constant_scalar(crate::Literal::U32( + spirv::RayQueryIntersection::RayQueryCandidateIntersectionKHR as _, + )); + let raw_kind_id = self.id_gen.next(); + valid_block + .body + .push(Instruction::ray_query_get_intersection( + spirv::Op::RayQueryGetIntersectionTypeKHR, + u32_ty, + raw_kind_id, + query_id, + intersection_id, + )); + + let candidate_tri_id = self.get_constant_scalar(crate::Literal::U32( + spirv::RayQueryCandidateIntersectionType::RayQueryCandidateIntersectionTriangleKHR as _, + )); + let intersection_tri_id = self.id_gen.next(); + valid_block.body.push(Instruction::binary( + spirv::Op::IEqual, + bool_type_id, + intersection_tri_id, + raw_kind_id, + candidate_tri_id, + )); + + let generate_label_id = self.id_gen.next(); + let mut generate_block = Block::new(generate_label_id); + + let merge_label_id = self.id_gen.next(); + let merge_block = Block::new(merge_label_id); + + valid_block.body.push(Instruction::selection_merge( + merge_label_id, + spirv::SelectionControl::NONE, + )); + function.consume( + valid_block, + Instruction::branch_conditional(intersection_tri_id, generate_label_id, merge_label_id), + ); + + generate_block + .body + .push(Instruction::ray_query_confirm_intersection(query_id)); + + function.consume(generate_block, Instruction::branch(merge_label_id)); + function.consume(merge_block, Instruction::branch(final_label_id)); + + function.consume(final_block, Instruction::return_void()); + + self.ray_query_functions + .insert(LookupRayQueryFunction::ConfirmIntersection, func_id); + + function.to_words(&mut self.logical_layout.function_definitions); + + func_id + } + + fn write_ray_query_get_vertex_positions( + &mut self, + is_committed: bool, + ir_module: &crate::Module, + ) -> spirv::Word { + if let Some(&word) = + self.ray_query_functions + .get(&LookupRayQueryFunction::GetVertexPositions { + committed: is_committed, + }) + { + return word; + } + + let (committed_ty, committed_tri_ty) = if is_committed { + ( + spirv::RayQueryIntersection::RayQueryCommittedIntersectionKHR as u32, + spirv::RayQueryCommittedIntersectionType::RayQueryCommittedIntersectionTriangleKHR + as u32, + ) + } else { + ( + spirv::RayQueryIntersection::RayQueryCandidateIntersectionKHR as u32, + spirv::RayQueryCandidateIntersectionType::RayQueryCandidateIntersectionTriangleKHR + as u32, + ) + }; + + let ray_query_type_id = self.get_ray_query_pointer_id(); + + let u32_ty = self.get_u32_type_id(); + let u32_ptr_ty = self.get_pointer_type_id(u32_ty, spirv::StorageClass::Function); + + let rq_get_vertex_positions_ty_id = self.get_handle_type_id( + *ir_module + .special_types + .ray_vertex_return + .as_ref() + .expect("must be generated when reading in get vertex position"), + ); + let ptr_return_ty = + self.get_pointer_type_id(rq_get_vertex_positions_ty_id, spirv::StorageClass::Function); + + let bool_type_id = self.get_bool_type_id(); + + let (func_id, mut function, arg_ids) = self.write_function_signature( + &[ray_query_type_id, u32_ptr_ty], + rq_get_vertex_positions_ty_id, + ); + + let query_id = arg_ids[0]; + let init_tracker_id = arg_ids[1]; + + let block_id = self.id_gen.next(); + let mut block = Block::new(block_id); + + let return_id = self.id_gen.next(); + block.body.push(Instruction::variable( + ptr_return_ty, + return_id, + spirv::StorageClass::Function, + Some(self.get_constant_null(rq_get_vertex_positions_ty_id)), + )); + + let valid_id = self.id_gen.next(); + let mut valid_block = Block::new(valid_id); + + let final_label_id = self.id_gen.next(); + let mut final_block = Block::new(final_label_id); + + let instruction = if self.ray_query_initialization_tracking { + let initialized_tracker_id = self.id_gen.next(); + block.body.push(Instruction::load( + u32_ty, + initialized_tracker_id, + init_tracker_id, + None, + )); + + let proceeded_id = write_ray_flags_contains_flags( + self, + &mut block, + initialized_tracker_id, + super::RayQueryPoint::PROCEED.bits(), + ); + let finished_proceed_id = write_ray_flags_contains_flags( + self, + &mut block, + initialized_tracker_id, + super::RayQueryPoint::FINISHED_TRAVERSAL.bits(), + ); + + let correct_finish_id = if is_committed { + finished_proceed_id + } else { + let not_finished_id = self.id_gen.next(); + block.body.push(Instruction::unary( + spirv::Op::LogicalNot, + bool_type_id, + not_finished_id, + finished_proceed_id, + )); + not_finished_id + }; + + let is_valid_id = self.write_logical_and(&mut block, correct_finish_id, proceeded_id); + block.body.push(Instruction::selection_merge( + final_label_id, + spirv::SelectionControl::NONE, + )); + Instruction::branch_conditional(is_valid_id, valid_id, final_label_id) + } else { + Instruction::branch(valid_id) + }; + + function.consume(block, instruction); + + let intersection_id = self.get_constant_scalar(crate::Literal::U32(committed_ty)); + let raw_kind_id = self.id_gen.next(); + valid_block + .body + .push(Instruction::ray_query_get_intersection( + spirv::Op::RayQueryGetIntersectionTypeKHR, + u32_ty, + raw_kind_id, + query_id, + intersection_id, + )); + + let candidate_tri_id = self.get_constant_scalar(crate::Literal::U32(committed_tri_ty)); + let intersection_tri_id = self.id_gen.next(); + valid_block.body.push(Instruction::binary( + spirv::Op::IEqual, + bool_type_id, + intersection_tri_id, + raw_kind_id, + candidate_tri_id, + )); + + let generate_label_id = self.id_gen.next(); + let mut vertex_return_block = Block::new(generate_label_id); + + let merge_label_id = self.id_gen.next(); + let merge_block = Block::new(merge_label_id); + + valid_block.body.push(Instruction::selection_merge( + merge_label_id, + spirv::SelectionControl::NONE, + )); + function.consume( + valid_block, + Instruction::branch_conditional(intersection_tri_id, generate_label_id, merge_label_id), + ); + + let vertices_id = self.id_gen.next(); + vertex_return_block + .body + .push(Instruction::ray_query_return_vertex_position( + rq_get_vertex_positions_ty_id, + vertices_id, + query_id, + intersection_id, + )); + vertex_return_block + .body + .push(Instruction::store(return_id, vertices_id, None)); + + function.consume(vertex_return_block, Instruction::branch(merge_label_id)); + function.consume(merge_block, Instruction::branch(final_label_id)); + + let loaded_pos_id = self.id_gen.next(); + final_block.body.push(Instruction::load( + rq_get_vertex_positions_ty_id, + loaded_pos_id, + return_id, + None, + )); + + function.consume(final_block, Instruction::return_value(loaded_pos_id)); + + self.ray_query_functions.insert( + LookupRayQueryFunction::GetVertexPositions { + committed: is_committed, + }, + func_id, + ); + + function.to_words(&mut self.logical_layout.function_definitions); + + func_id + } +} + +impl BlockContext<'_> { + pub(super) fn write_ray_query_function( + &mut self, + query: Handle, + function: &crate::RayQueryFunction, + block: &mut Block, + ) { + let query_id = self.cached[query]; + let tracker_ids = *self + .ray_query_tracker_expr + .get(&query) + .expect("not a cached ray query"); + + match *function { + crate::RayQueryFunction::Initialize { + acceleration_structure, + descriptor, + } => { + let desc_id = self.cached[descriptor]; + let acc_struct_id = self.get_handle_id(acceleration_structure); + + let func = self.writer.write_ray_query_initialize(self.ir_module); + + let func_id = self.gen_id(); + block.body.push(Instruction::function_call( + self.writer.void_type, + func_id, + func, + &[ + query_id, + acc_struct_id, + desc_id, + tracker_ids.initialized_tracker, + tracker_ids.t_max_tracker, + ], + )); + } + crate::RayQueryFunction::Proceed { result } => { + let id = self.gen_id(); + self.cached[result] = id; + + let bool_ty = self.writer.get_bool_type_id(); + + let func_id = self.writer.write_ray_query_proceed(); + block.body.push(Instruction::function_call( + bool_ty, + id, + func_id, + &[query_id, tracker_ids.initialized_tracker], + )); + } + crate::RayQueryFunction::GenerateIntersection { hit_t } => { + let hit_id = self.cached[hit_t]; + + let func_id = self.writer.write_ray_query_generate_intersection(); + + let func_call_id = self.gen_id(); + block.body.push(Instruction::function_call( + self.writer.void_type, + func_call_id, + func_id, + &[ + query_id, + tracker_ids.initialized_tracker, + hit_id, + tracker_ids.t_max_tracker, + ], + )); + } + crate::RayQueryFunction::ConfirmIntersection => { + let func_id = self.writer.write_ray_query_confirm_intersection(); + + let func_call_id = self.gen_id(); + block.body.push(Instruction::function_call( + self.writer.void_type, + func_call_id, + func_id, + &[query_id, tracker_ids.initialized_tracker], + )); + } + crate::RayQueryFunction::Terminate => {} + } + } + + pub(super) fn write_ray_query_return_vertex_position( + &mut self, + query: Handle, + block: &mut Block, + is_committed: bool, + ) -> spirv::Word { + let fn_id = self + .writer + .write_ray_query_get_vertex_positions(is_committed, self.ir_module); + + let query_id = self.cached[query]; + let tracker_id = *self + .ray_query_tracker_expr + .get(&query) + .expect("not a cached ray query"); + + let rq_get_vertex_positions_ty_id = self.get_handle_type_id( + *self + .ir_module + .special_types + .ray_vertex_return + .as_ref() + .expect("must be generated when reading in get vertex position"), + ); + + let func_call_id = self.gen_id(); + block.body.push(Instruction::function_call( + rq_get_vertex_positions_ty_id, + func_call_id, + fn_id, + &[query_id, tracker_id.initialized_tracker], + )); + func_call_id } } diff --git a/naga/src/back/spv/writer.rs b/naga/src/back/spv/writer.rs index 1beb86577c8..f58f9dc04a4 100644 --- a/naga/src/back/spv/writer.rs +++ b/naga/src/back/spv/writer.rs @@ -35,6 +35,12 @@ impl Function { for local_var in self.variables.values() { local_var.instruction.to_words(sink); } + for local_var in self.ray_query_initialization_tracker_variables.values() { + local_var.instruction.to_words(sink); + } + for local_var in self.ray_query_t_max_tracker_variables.values() { + local_var.instruction.to_words(sink); + } for local_var in self.force_loop_bounding_vars.iter() { local_var.instruction.to_words(sink); } @@ -71,12 +77,14 @@ impl Writer { capabilities_available: options.capabilities.clone(), capabilities_used, extensions_used: crate::FastIndexSet::default(), + debug_strings: vec![], debugs: vec![], annotations: vec![], flags: options.flags, bounds_check_policies: options.bounds_check_policies, zero_initialize_workgroup_memory: options.zero_initialize_workgroup_memory, force_loop_bounding: options.force_loop_bounding, + ray_query_initialization_tracking: options.ray_query_initialization_tracking, use_storage_input_output_16: options.use_storage_input_output_16, void_type, lookup_type: crate::FastHashMap::default(), @@ -91,11 +99,11 @@ impl Writer { saved_cached: CachedExpressions::default(), gl450_ext_inst_id, temp_list: Vec::new(), - ray_get_committed_intersection_function: None, - ray_get_candidate_intersection_function: None, + ray_query_functions: crate::FastHashMap::default(), io_f16_polyfills: super::f16_polyfill::F16IoPolyfill::new( options.use_storage_input_output_16, ), + debug_printf: None, }) } @@ -147,6 +155,7 @@ impl Writer { bounds_check_policies: self.bounds_check_policies, zero_initialize_workgroup_memory: self.zero_initialize_workgroup_memory, force_loop_bounding: self.force_loop_bounding, + ray_query_initialization_tracking: self.ray_query_initialization_tracking, use_storage_input_output_16: self.use_storage_input_output_16, capabilities_available: take(&mut self.capabilities_available), fake_missing_bindings: self.fake_missing_bindings, @@ -162,6 +171,7 @@ impl Writer { extensions_used: take(&mut self.extensions_used).recycle(), physical_layout: self.physical_layout.clone().recycle(), logical_layout: take(&mut self.logical_layout).recycle(), + debug_strings: take(&mut self.debug_strings).recycle(), debugs: take(&mut self.debugs).recycle(), annotations: take(&mut self.annotations).recycle(), lookup_type: take(&mut self.lookup_type).recycle(), @@ -173,9 +183,9 @@ impl Writer { global_variables: take(&mut self.global_variables).recycle(), saved_cached: take(&mut self.saved_cached).recycle(), temp_list: take(&mut self.temp_list).recycle(), - ray_get_candidate_intersection_function: None, - ray_get_committed_intersection_function: None, + ray_query_functions: take(&mut self.ray_query_functions).recycle(), io_f16_polyfills: take(&mut self.io_f16_polyfills).recycle(), + debug_printf: None, }; *self = fresh; @@ -1022,6 +1032,7 @@ impl Writer { expression_constness: super::ExpressionConstnessTracker::from_arena( &ir_function.expressions, ), + ray_query_tracker_expr: crate::FastHashMap::default(), }; // fill up the pre-emitted and const expressions @@ -1063,6 +1074,58 @@ impl Writer { .function .variables .insert(handle, LocalVariable { id, instruction }); + + if let crate::TypeInner::RayQuery { .. } = ir_module.types[variable.ty].inner { + // Don't refactor this into a struct: Although spirv itself allows opaque types in structs, + // the vulkan environment for spirv does not. Putting ray queries into structs can cause + // confusing bugs. + let u32_type_id = context.writer.get_u32_type_id(); + let ptr_u32_type_id = context + .writer + .get_pointer_type_id(u32_type_id, spirv::StorageClass::Function); + let tracker_id = context.gen_id(); + let tracker_init_id = context + .writer + .get_constant_scalar(crate::Literal::U32(super::RayQueryPoint::empty().bits())); + let tracker_instruction = Instruction::variable( + ptr_u32_type_id, + tracker_id, + spirv::StorageClass::Function, + Some(tracker_init_id), + ); + + context + .function + .ray_query_initialization_tracker_variables + .insert( + handle, + LocalVariable { + id: tracker_id, + instruction: tracker_instruction, + }, + ); + let f32_type_id = context.writer.get_f32_type_id(); + let ptr_f32_type_id = context + .writer + .get_pointer_type_id(f32_type_id, spirv::StorageClass::Function); + let t_max_tracker_id = context.gen_id(); + let t_max_tracker_init_id = + context.writer.get_constant_scalar(crate::Literal::F32(0.0)); + let t_max_tracker_instruction = Instruction::variable( + ptr_f32_type_id, + t_max_tracker_id, + spirv::StorageClass::Function, + Some(t_max_tracker_init_id), + ); + + context.function.ray_query_t_max_tracker_variables.insert( + handle, + LocalVariable { + id: t_max_tracker_id, + instruction: t_max_tracker_instruction, + }, + ); + } } for (handle, expr) in ir_function.expressions.iter() { @@ -2651,6 +2714,10 @@ impl Writer { Instruction::memory_model(addressing_model, memory_model) .to_words(&mut self.logical_layout.memory_model); + for debug_string in self.debug_strings.iter() { + debug_string.to_words(&mut self.logical_layout.debugs); + } + if self.flags.contains(WriterFlags::DEBUG) { for debug in self.debugs.iter() { debug.to_words(&mut self.logical_layout.debugs); @@ -2710,6 +2777,41 @@ impl Writer { pub(super) fn needs_f16_polyfill(&self, ty_inner: &crate::TypeInner) -> bool { self.io_f16_polyfills.needs_polyfill(ty_inner) } + + #[allow(dead_code)] + pub(super) fn write_debug_printf( + &mut self, + block: &mut Block, + string: &str, + format_params: &[Word], + ) { + if self.debug_printf.is_none() { + self.use_extension("SPV_KHR_non_semantic_info"); + let import_id = self.id_gen.next(); + Instruction::ext_inst_import(import_id, "NonSemantic.DebugPrintf") + .to_words(&mut self.logical_layout.ext_inst_imports); + self.debug_printf = Some(import_id) + } + + let import_id = self.debug_printf.unwrap(); + + let string_id = self.id_gen.next(); + self.debug_strings + .push(Instruction::string(string, string_id)); + + let mut operands = Vec::with_capacity(1 + format_params.len()); + operands.push(string_id); + operands.extend(format_params.iter()); + + let print_id = self.id_gen.next(); + block.body.push(Instruction::ext_inst( + import_id, + 1, + self.void_type, + print_id, + &operands, + )); + } } #[test] diff --git a/naga/tests/in/wgsl/ray-query-no-init-tracking.toml b/naga/tests/in/wgsl/ray-query-no-init-tracking.toml new file mode 100644 index 00000000000..e2602b7b4d2 --- /dev/null +++ b/naga/tests/in/wgsl/ray-query-no-init-tracking.toml @@ -0,0 +1,19 @@ +god_mode = true +targets = "SPIRV | METAL | HLSL" + +[msl] +fake_missing_bindings = true +lang_version = [2, 4] +spirv_cross_compatibility = false +zero_initialize_workgroup_memory = false + +[hlsl] +shader_model = "V6_5" +fake_missing_bindings = true +zero_initialize_workgroup_memory = true +# Not yet implemented +# ray_query_initialization_tracking = false + +[spv] +version = [1, 4] +ray_query_initialization_tracking = false diff --git a/naga/tests/in/wgsl/ray-query-no-init-tracking.wgsl b/naga/tests/in/wgsl/ray-query-no-init-tracking.wgsl new file mode 100644 index 00000000000..e8fabb0208c --- /dev/null +++ b/naga/tests/in/wgsl/ray-query-no-init-tracking.wgsl @@ -0,0 +1,97 @@ +/* +let RAY_FLAG_NONE = 0x00u; +let RAY_FLAG_FORCE_OPAQUE = 0x01u; +let RAY_FLAG_FORCE_NO_OPAQUE = 0x02u; +let RAY_FLAG_TERMINATE_ON_FIRST_HIT = 0x04u; +let RAY_FLAG_SKIP_CLOSEST_HIT_SHADER = 0x08u; +let RAY_FLAG_CULL_BACK_FACING = 0x10u; +let RAY_FLAG_CULL_FRONT_FACING = 0x20u; +let RAY_FLAG_CULL_OPAQUE = 0x40u; +let RAY_FLAG_CULL_NO_OPAQUE = 0x80u; +let RAY_FLAG_SKIP_TRIANGLES = 0x100u; +let RAY_FLAG_SKIP_AABBS = 0x200u; + +let RAY_QUERY_INTERSECTION_NONE = 0u; +let RAY_QUERY_INTERSECTION_TRIANGLE = 1u; +let RAY_QUERY_INTERSECTION_GENERATED = 2u; +let RAY_QUERY_INTERSECTION_AABB = 3u; + +struct RayDesc { + flags: u32, + cull_mask: u32, + t_min: f32, + t_max: f32, + origin: vec3, + dir: vec3, +} + +struct RayIntersection { + kind: u32, + t: f32, + instance_custom_data: u32, + instance_index: u32, + sbt_record_offset: u32, + geometry_index: u32, + primitive_index: u32, + barycentrics: vec2, + front_face: bool, + object_to_world: mat4x3, + world_to_object: mat4x3, +} +*/ + +fn query_loop(pos: vec3, dir: vec3, acs: acceleration_structure) -> RayIntersection { + var rq: ray_query; + rayQueryInitialize(&rq, acs, RayDesc(RAY_FLAG_TERMINATE_ON_FIRST_HIT, 0xFFu, 0.1, 100.0, pos, dir)); + + while (rayQueryProceed(&rq)) {} + + return rayQueryGetCommittedIntersection(&rq); +} + +@group(0) @binding(0) +var acc_struct: acceleration_structure; + +struct Output { + visible: u32, + normal: vec3, +} + +@group(0) @binding(1) +var output: Output; + +fn get_torus_normal(world_point: vec3, intersection: RayIntersection) -> vec3 { + let local_point = intersection.world_to_object * vec4(world_point, 1.0); + let point_on_guiding_line = normalize(local_point.xy) * 2.4; + let world_point_on_guiding_line = intersection.object_to_world * vec4(point_on_guiding_line, 0.0, 1.0); + return normalize(world_point - world_point_on_guiding_line); +} + + + +@compute @workgroup_size(1) +fn main() { + let pos = vec3(0.0); + let dir = vec3(0.0, 1.0, 0.0); + let intersection = query_loop(pos, dir, acc_struct); + + output.visible = u32(intersection.kind == RAY_QUERY_INTERSECTION_NONE); + output.normal = get_torus_normal(dir * intersection.t, intersection); +} + +@compute @workgroup_size(1) +fn main_candidate() { + let pos = vec3(0.0); + let dir = vec3(0.0, 1.0, 0.0); + + var rq: ray_query; + rayQueryInitialize(&rq, acc_struct, RayDesc(RAY_FLAG_TERMINATE_ON_FIRST_HIT, 0xFFu, 0.1, 100.0, pos, dir)); + let intersection = rayQueryGetCandidateIntersection(&rq); + if (intersection.kind == RAY_QUERY_INTERSECTION_AABB) { + rayQueryGenerateIntersection(&rq, 10.0); + } else if (intersection.kind == RAY_QUERY_INTERSECTION_TRIANGLE) { + rayQueryConfirmIntersection(&rq); + } else { + rayQueryTerminate(&rq); + } +} diff --git a/naga/tests/out/hlsl/wgsl-ray-query-no-init-tracking.hlsl b/naga/tests/out/hlsl/wgsl-ray-query-no-init-tracking.hlsl new file mode 100644 index 00000000000..68be09b6b01 --- /dev/null +++ b/naga/tests/out/hlsl/wgsl-ray-query-no-init-tracking.hlsl @@ -0,0 +1,165 @@ +struct RayIntersection { + uint kind; + float t; + uint instance_custom_data; + uint instance_index; + uint sbt_record_offset; + uint geometry_index; + uint primitive_index; + float2 barycentrics; + bool front_face; + int _pad9_0; + int _pad9_1; + row_major float4x3 object_to_world; + int _pad10_0; + row_major float4x3 world_to_object; + int _end_pad_0; +}; + +struct RayDesc_ { + uint flags; + uint cull_mask; + float tmin; + float tmax; + float3 origin; + int _pad5_0; + float3 dir; + int _end_pad_0; +}; + +struct Output { + uint visible; + int _pad1_0; + int _pad1_1; + int _pad1_2; + float3 normal; + int _end_pad_0; +}; + +RayDesc RayDescFromRayDesc_(RayDesc_ arg0) { + RayDesc ret = (RayDesc)0; + ret.Origin = arg0.origin; + ret.TMin = arg0.tmin; + ret.Direction = arg0.dir; + ret.TMax = arg0.tmax; + return ret; +} + +RaytracingAccelerationStructure acc_struct : register(t0); +RWByteAddressBuffer output : register(u1); + +RayDesc_ ConstructRayDesc_(uint arg0, uint arg1, float arg2, float arg3, float3 arg4, float3 arg5) { + RayDesc_ ret = (RayDesc_)0; + ret.flags = arg0; + ret.cull_mask = arg1; + ret.tmin = arg2; + ret.tmax = arg3; + ret.origin = arg4; + ret.dir = arg5; + return ret; +} + +RayIntersection GetCommittedIntersection(RayQuery rq) { + RayIntersection ret = (RayIntersection)0; + ret.kind = rq.CommittedStatus(); + if( rq.CommittedStatus() == COMMITTED_NOTHING) {} else { + ret.t = rq.CommittedRayT(); + ret.instance_custom_data = rq.CommittedInstanceID(); + ret.instance_index = rq.CommittedInstanceIndex(); + ret.sbt_record_offset = rq.CommittedInstanceContributionToHitGroupIndex(); + ret.geometry_index = rq.CommittedGeometryIndex(); + ret.primitive_index = rq.CommittedPrimitiveIndex(); + if( rq.CommittedStatus() == COMMITTED_TRIANGLE_HIT ) { + ret.barycentrics = rq.CommittedTriangleBarycentrics(); + ret.front_face = rq.CommittedTriangleFrontFace(); + } + ret.object_to_world = rq.CommittedObjectToWorld4x3(); + ret.world_to_object = rq.CommittedWorldToObject4x3(); + } + return ret; +} + +RayIntersection query_loop(float3 pos, float3 dir, RaytracingAccelerationStructure acs) +{ + RayQuery rq_1; + + rq_1.TraceRayInline(acs, ConstructRayDesc_(4u, 255u, 0.1, 100.0, pos, dir).flags, ConstructRayDesc_(4u, 255u, 0.1, 100.0, pos, dir).cull_mask, RayDescFromRayDesc_(ConstructRayDesc_(4u, 255u, 0.1, 100.0, pos, dir))); + uint2 loop_bound = uint2(4294967295u, 4294967295u); + while(true) { + if (all(loop_bound == uint2(0u, 0u))) { break; } + loop_bound -= uint2(loop_bound.y == 0u, 1u); + const bool _e9 = rq_1.Proceed(); + if (_e9) { + } else { + break; + } + { + } + } + const RayIntersection rayintersection = GetCommittedIntersection(rq_1); + return rayintersection; +} + +float3 get_torus_normal(float3 world_point, RayIntersection intersection) +{ + float3 local_point = mul(float4(world_point, 1.0), intersection.world_to_object); + float2 point_on_guiding_line = (normalize(local_point.xy) * 2.4); + float3 world_point_on_guiding_line = mul(float4(point_on_guiding_line, 0.0, 1.0), intersection.object_to_world); + return normalize((world_point - world_point_on_guiding_line)); +} + +[numthreads(1, 1, 1)] +void main() +{ + float3 pos_1 = (0.0).xxx; + float3 dir_1 = float3(0.0, 1.0, 0.0); + const RayIntersection _e7 = query_loop(pos_1, dir_1, acc_struct); + output.Store(0, asuint(uint((_e7.kind == 0u)))); + const float3 _e18 = get_torus_normal((dir_1 * _e7.t), _e7); + output.Store3(16, asuint(_e18)); + return; +} + +RayIntersection GetCandidateIntersection(RayQuery rq) { + RayIntersection ret = (RayIntersection)0; + CANDIDATE_TYPE kind = rq.CandidateType(); + if (kind == CANDIDATE_NON_OPAQUE_TRIANGLE) { + ret.kind = 1; + ret.t = rq.CandidateTriangleRayT(); + ret.barycentrics = rq.CandidateTriangleBarycentrics(); + ret.front_face = rq.CandidateTriangleFrontFace(); + } else { + ret.kind = 3; + } + ret.instance_custom_data = rq.CandidateInstanceID(); + ret.instance_index = rq.CandidateInstanceIndex(); + ret.sbt_record_offset = rq.CandidateInstanceContributionToHitGroupIndex(); + ret.geometry_index = rq.CandidateGeometryIndex(); + ret.primitive_index = rq.CandidatePrimitiveIndex(); + ret.object_to_world = rq.CandidateObjectToWorld4x3(); + ret.world_to_object = rq.CandidateWorldToObject4x3(); + return ret; +} + +[numthreads(1, 1, 1)] +void main_candidate() +{ + RayQuery rq; + + float3 pos_2 = (0.0).xxx; + float3 dir_2 = float3(0.0, 1.0, 0.0); + rq.TraceRayInline(acc_struct, ConstructRayDesc_(4u, 255u, 0.1, 100.0, pos_2, dir_2).flags, ConstructRayDesc_(4u, 255u, 0.1, 100.0, pos_2, dir_2).cull_mask, RayDescFromRayDesc_(ConstructRayDesc_(4u, 255u, 0.1, 100.0, pos_2, dir_2))); + RayIntersection intersection_1 = GetCandidateIntersection(rq); + if ((intersection_1.kind == 3u)) { + rq.CommitProceduralPrimitiveHit(10.0); + return; + } else { + if ((intersection_1.kind == 1u)) { + rq.CommitNonOpaqueTriangleHit(); + return; + } else { + rq.Abort(); + return; + } + } +} diff --git a/naga/tests/out/hlsl/wgsl-ray-query-no-init-tracking.ron b/naga/tests/out/hlsl/wgsl-ray-query-no-init-tracking.ron new file mode 100644 index 00000000000..a31e1db125a --- /dev/null +++ b/naga/tests/out/hlsl/wgsl-ray-query-no-init-tracking.ron @@ -0,0 +1,16 @@ +( + vertex:[ + ], + fragment:[ + ], + compute:[ + ( + entry_point:"main", + target_profile:"cs_6_5", + ), + ( + entry_point:"main_candidate", + target_profile:"cs_6_5", + ), + ], +) diff --git a/naga/tests/out/msl/wgsl-ray-query-no-init-tracking.msl b/naga/tests/out/msl/wgsl-ray-query-no-init-tracking.msl new file mode 100644 index 00000000000..55840c10920 --- /dev/null +++ b/naga/tests/out/msl/wgsl-ray-query-no-init-tracking.msl @@ -0,0 +1,116 @@ +// language: metal2.4 +#include +#include + +using metal::uint; +struct _RayQuery { + metal::raytracing::intersector intersector; + metal::raytracing::intersector::result_type intersection; + bool ready = false; +}; +constexpr metal::uint _map_intersection_type(const metal::raytracing::intersection_type ty) { + return ty==metal::raytracing::intersection_type::triangle ? 1 : + ty==metal::raytracing::intersection_type::bounding_box ? 4 : 0; +} + +struct RayIntersection { + uint kind; + float t; + uint instance_custom_data; + uint instance_index; + uint sbt_record_offset; + uint geometry_index; + uint primitive_index; + metal::float2 barycentrics; + bool front_face; + char _pad9[11]; + metal::float4x3 object_to_world; + metal::float4x3 world_to_object; +}; +struct RayDesc { + uint flags; + uint cull_mask; + float tmin; + float tmax; + metal::float3 origin; + metal::float3 dir; +}; +struct Output { + uint visible; + char _pad1[12]; + metal::float3 normal; +}; + +RayIntersection query_loop( + metal::float3 pos, + metal::float3 dir, + metal::raytracing::instance_acceleration_structure acs +) { + _RayQuery rq_1 = {}; + RayDesc _e8 = RayDesc {4u, 255u, 0.1, 100.0, pos, dir}; + rq_1.intersector.assume_geometry_type(metal::raytracing::geometry_type::triangle); + rq_1.intersector.set_opacity_cull_mode((_e8.flags & 64) != 0 ? metal::raytracing::opacity_cull_mode::opaque : (_e8.flags & 128) != 0 ? metal::raytracing::opacity_cull_mode::non_opaque : metal::raytracing::opacity_cull_mode::none); + rq_1.intersector.force_opacity((_e8.flags & 1) != 0 ? metal::raytracing::forced_opacity::opaque : (_e8.flags & 2) != 0 ? metal::raytracing::forced_opacity::non_opaque : metal::raytracing::forced_opacity::none); + rq_1.intersector.accept_any_intersection((_e8.flags & 4) != 0); + rq_1.intersection = rq_1.intersector.intersect(metal::raytracing::ray(_e8.origin, _e8.dir, _e8.tmin, _e8.tmax), acs, _e8.cull_mask); rq_1.ready = true; + uint2 loop_bound = uint2(4294967295u); + while(true) { + if (metal::all(loop_bound == uint2(0u))) { break; } + loop_bound -= uint2(loop_bound.y == 0u, 1u); + bool _e9 = rq_1.ready; + if (_e9) { + } else { + break; + } + } + return RayIntersection {_map_intersection_type(rq_1.intersection.type), rq_1.intersection.distance, rq_1.intersection.user_instance_id, rq_1.intersection.instance_id, {}, rq_1.intersection.geometry_id, rq_1.intersection.primitive_id, rq_1.intersection.triangle_barycentric_coord, rq_1.intersection.triangle_front_facing, {}, rq_1.intersection.object_to_world_transform, rq_1.intersection.world_to_object_transform}; +} + +metal::float3 get_torus_normal( + metal::float3 world_point, + RayIntersection intersection +) { + metal::float3 local_point = intersection.world_to_object * metal::float4(world_point, 1.0); + metal::float2 point_on_guiding_line = metal::normalize(local_point.xy) * 2.4; + metal::float3 world_point_on_guiding_line = intersection.object_to_world * metal::float4(point_on_guiding_line, 0.0, 1.0); + return metal::normalize(world_point - world_point_on_guiding_line); +} + +kernel void main_( + metal::raytracing::instance_acceleration_structure acc_struct [[user(fake0)]] +, device Output& output [[user(fake0)]] +) { + metal::float3 pos_1 = metal::float3(0.0); + metal::float3 dir_1 = metal::float3(0.0, 1.0, 0.0); + RayIntersection _e7 = query_loop(pos_1, dir_1, acc_struct); + output.visible = static_cast(_e7.kind == 0u); + metal::float3 _e18 = get_torus_normal(dir_1 * _e7.t, _e7); + output.normal = _e18; + return; +} + + +kernel void main_candidate( + metal::raytracing::instance_acceleration_structure acc_struct [[user(fake0)]] +) { + _RayQuery rq = {}; + metal::float3 pos_2 = metal::float3(0.0); + metal::float3 dir_2 = metal::float3(0.0, 1.0, 0.0); + RayDesc _e12 = RayDesc {4u, 255u, 0.1, 100.0, pos_2, dir_2}; + rq.intersector.assume_geometry_type(metal::raytracing::geometry_type::triangle); + rq.intersector.set_opacity_cull_mode((_e12.flags & 64) != 0 ? metal::raytracing::opacity_cull_mode::opaque : (_e12.flags & 128) != 0 ? metal::raytracing::opacity_cull_mode::non_opaque : metal::raytracing::opacity_cull_mode::none); + rq.intersector.force_opacity((_e12.flags & 1) != 0 ? metal::raytracing::forced_opacity::opaque : (_e12.flags & 2) != 0 ? metal::raytracing::forced_opacity::non_opaque : metal::raytracing::forced_opacity::none); + rq.intersector.accept_any_intersection((_e12.flags & 4) != 0); + rq.intersection = rq.intersector.intersect(metal::raytracing::ray(_e12.origin, _e12.dir, _e12.tmin, _e12.tmax), acc_struct, _e12.cull_mask); rq.ready = true; + RayIntersection intersection_1 = RayIntersection {_map_intersection_type(rq.intersection.type), rq.intersection.distance, rq.intersection.user_instance_id, rq.intersection.instance_id, {}, rq.intersection.geometry_id, rq.intersection.primitive_id, rq.intersection.triangle_barycentric_coord, rq.intersection.triangle_front_facing, {}, rq.intersection.object_to_world_transform, rq.intersection.world_to_object_transform}; + if (intersection_1.kind == 3u) { + return; + } else { + if (intersection_1.kind == 1u) { + return; + } else { + rq.ready = false; + return; + } + } +} diff --git a/naga/tests/out/spv/wgsl-aliased-ray-query.spvasm b/naga/tests/out/spv/wgsl-aliased-ray-query.spvasm index b095e8b8e83..75b7f16bbd9 100644 --- a/naga/tests/out/spv/wgsl-aliased-ray-query.spvasm +++ b/naga/tests/out/spv/wgsl-aliased-ray-query.spvasm @@ -1,7 +1,7 @@ ; SPIR-V ; Version: 1.4 ; Generator: rspirv -; Bound: 102 +; Bound: 244 OpCapability Shader OpCapability RayQueryKHR OpExtension "SPV_KHR_ray_query" @@ -59,109 +59,284 @@ OpDecorate %13 Binding 0 %29 = OpConstant %5 10 %30 = OpConstant %7 1 %32 = OpTypePointer Function %3 -%40 = OpTypePointer Function %12 -%41 = OpTypePointer Function %7 -%42 = OpTypePointer Function %11 -%43 = OpTypePointer Function %9 -%44 = OpTypePointer Function %10 -%45 = OpTypePointer Function %5 -%46 = OpTypeFunction %12 %32 -%48 = OpConstantNull %12 -%52 = OpConstant %7 0 -%67 = OpConstant %7 2 -%71 = OpConstant %7 5 -%73 = OpConstant %7 6 -%75 = OpConstant %7 9 -%77 = OpConstant %7 10 -%86 = OpConstant %7 7 -%88 = OpConstant %7 8 -%47 = OpFunction %12 None %46 -%49 = OpFunctionParameter %32 -%50 = OpLabel -%51 = OpVariable %40 Function %48 -%53 = OpRayQueryGetIntersectionTypeKHR %7 %49 %52 -%54 = OpIEqual %10 %53 %52 -%55 = OpSelect %7 %54 %30 %28 -%56 = OpAccessChain %41 %51 %52 -OpStore %56 %55 -%57 = OpINotEqual %10 %55 %52 -OpSelectionMerge %59 None -OpBranchConditional %57 %58 %59 -%58 = OpLabel -%60 = OpRayQueryGetIntersectionInstanceCustomIndexKHR %7 %49 %52 -%61 = OpRayQueryGetIntersectionInstanceIdKHR %7 %49 %52 -%62 = OpRayQueryGetIntersectionInstanceShaderBindingTableRecordOffsetKHR %7 %49 %52 -%63 = OpRayQueryGetIntersectionGeometryIndexKHR %7 %49 %52 -%64 = OpRayQueryGetIntersectionPrimitiveIndexKHR %7 %49 %52 -%65 = OpRayQueryGetIntersectionObjectToWorldKHR %11 %49 %52 -%66 = OpRayQueryGetIntersectionWorldToObjectKHR %11 %49 %52 -%68 = OpAccessChain %41 %51 %67 -OpStore %68 %60 -%69 = OpAccessChain %41 %51 %28 -OpStore %69 %61 -%70 = OpAccessChain %41 %51 %23 -OpStore %70 %62 -%72 = OpAccessChain %41 %51 %71 -OpStore %72 %63 -%74 = OpAccessChain %41 %51 %73 -OpStore %74 %64 -%76 = OpAccessChain %42 %51 %75 -OpStore %76 %65 -%78 = OpAccessChain %42 %51 %77 -OpStore %78 %66 -%79 = OpIEqual %10 %55 %30 -OpSelectionMerge %81 None -OpBranchConditional %57 %80 %81 -%80 = OpLabel -%82 = OpRayQueryGetIntersectionTKHR %5 %49 %52 -%83 = OpAccessChain %45 %51 %30 -OpStore %83 %82 -%84 = OpRayQueryGetIntersectionBarycentricsKHR %9 %49 %52 -%85 = OpRayQueryGetIntersectionFrontFaceKHR %10 %49 %52 -%87 = OpAccessChain %43 %51 %86 -OpStore %87 %84 -%89 = OpAccessChain %44 %51 %88 -OpStore %89 %85 -OpBranch %81 -%81 = OpLabel -OpBranch %59 -%59 = OpLabel -%90 = OpLoad %12 %51 -OpReturnValue %90 +%33 = OpTypePointer Function %7 +%35 = OpConstant %7 0 +%36 = OpTypePointer Function %5 +%39 = OpTypeVector %10 3 +%40 = OpTypeFunction %2 %32 %4 %8 %33 %36 +%68 = OpConstant %7 256 +%71 = OpConstant %7 512 +%76 = OpConstant %7 16 +%79 = OpConstant %7 32 +%90 = OpConstant %7 2 +%93 = OpConstant %7 64 +%96 = OpConstant %7 128 +%121 = OpTypePointer Function %12 +%122 = OpTypePointer Function %11 +%123 = OpTypePointer Function %9 +%124 = OpTypePointer Function %10 +%125 = OpTypeFunction %12 %32 %33 +%130 = OpConstantNull %12 +%158 = OpConstant %7 5 +%160 = OpConstant %7 6 +%162 = OpConstant %7 9 +%164 = OpConstant %7 10 +%173 = OpConstant %7 7 +%175 = OpConstant %7 8 +%184 = OpTypeFunction %2 %32 %33 %5 %36 +%225 = OpTypeFunction %2 %32 %33 +%41 = OpFunction %2 None %40 +%42 = OpFunctionParameter %32 +%43 = OpFunctionParameter %4 +%44 = OpFunctionParameter %8 +%45 = OpFunctionParameter %33 +%46 = OpFunctionParameter %36 +%47 = OpLabel +%48 = OpCompositeExtract %7 %44 0 +%49 = OpCompositeExtract %7 %44 1 +%50 = OpCompositeExtract %5 %44 2 +%51 = OpCompositeExtract %5 %44 3 +OpStore %46 %51 +%52 = OpCompositeExtract %6 %44 4 +%53 = OpCompositeExtract %6 %44 5 +%54 = OpFOrdLessThanEqual %10 %50 %51 +%55 = OpFOrdGreaterThanEqual %10 %50 %19 +%56 = OpIsInf %39 %52 +%57 = OpAny %10 %56 +%58 = OpIsNan %39 %52 +%59 = OpAny %10 %58 +%60 = OpLogicalOr %10 %59 %57 +%61 = OpLogicalNot %10 %60 +%62 = OpIsInf %39 %53 +%63 = OpAny %10 %62 +%64 = OpIsNan %39 %53 +%65 = OpAny %10 %64 +%66 = OpLogicalOr %10 %65 %63 +%67 = OpLogicalNot %10 %66 +%69 = OpBitwiseAnd %7 %48 %68 +%70 = OpINotEqual %10 %69 %35 +%72 = OpBitwiseAnd %7 %48 %71 +%73 = OpINotEqual %10 %72 %35 +%74 = OpLogicalAnd %10 %73 %70 +%75 = OpLogicalNot %10 %74 +%77 = OpBitwiseAnd %7 %48 %76 +%78 = OpINotEqual %10 %77 %35 +%80 = OpBitwiseAnd %7 %48 %79 +%81 = OpINotEqual %10 %80 %35 +%82 = OpLogicalAnd %10 %81 %70 +%83 = OpLogicalAnd %10 %81 %78 +%84 = OpLogicalAnd %10 %78 %70 +%85 = OpLogicalOr %10 %84 %82 +%86 = OpLogicalOr %10 %85 %83 +%87 = OpLogicalNot %10 %86 +%88 = OpBitwiseAnd %7 %48 %30 +%89 = OpINotEqual %10 %88 %35 +%91 = OpBitwiseAnd %7 %48 %90 +%92 = OpINotEqual %10 %91 %35 +%94 = OpBitwiseAnd %7 %48 %93 +%95 = OpINotEqual %10 %94 %35 +%97 = OpBitwiseAnd %7 %48 %96 +%98 = OpINotEqual %10 %97 %35 +%99 = OpLogicalAnd %10 %98 %89 +%100 = OpLogicalAnd %10 %98 %92 +%101 = OpLogicalAnd %10 %98 %95 +%102 = OpLogicalAnd %10 %95 %89 +%103 = OpLogicalAnd %10 %95 %92 +%104 = OpLogicalAnd %10 %92 %89 +%105 = OpLogicalOr %10 %104 %99 +%106 = OpLogicalOr %10 %105 %100 +%107 = OpLogicalOr %10 %106 %101 +%108 = OpLogicalOr %10 %107 %102 +%109 = OpLogicalOr %10 %108 %103 +%110 = OpLogicalNot %10 %109 +%111 = OpLogicalAnd %10 %110 %54 +%112 = OpLogicalAnd %10 %111 %55 +%113 = OpLogicalAnd %10 %112 %61 +%114 = OpLogicalAnd %10 %113 %67 +%115 = OpLogicalAnd %10 %114 %75 +%116 = OpLogicalAnd %10 %115 %87 +OpSelectionMerge %117 None +OpBranchConditional %116 %119 %118 +%119 = OpLabel +OpRayQueryInitializeKHR %42 %43 %48 %49 %52 %50 %53 %51 +OpStore %45 %30 +OpBranch %117 +%118 = OpLabel +OpBranch %117 +%117 = OpLabel +OpReturn +OpFunctionEnd +%126 = OpFunction %12 None %125 +%127 = OpFunctionParameter %32 +%128 = OpFunctionParameter %33 +%129 = OpLabel +%131 = OpVariable %121 Function %130 +%132 = OpLoad %7 %128 +%133 = OpBitwiseAnd %7 %132 %90 +%134 = OpINotEqual %10 %133 %35 +%135 = OpBitwiseAnd %7 %132 %23 +%136 = OpINotEqual %10 %135 %35 +%137 = OpLogicalNot %10 %136 +%138 = OpLogicalAnd %10 %137 %134 +OpSelectionMerge %140 None +OpBranchConditional %138 %139 %140 +%139 = OpLabel +%141 = OpRayQueryGetIntersectionTypeKHR %7 %127 %35 +%142 = OpIEqual %10 %141 %35 +%143 = OpSelect %7 %142 %30 %28 +%144 = OpAccessChain %33 %131 %35 +OpStore %144 %143 +%145 = OpINotEqual %10 %143 %35 +OpSelectionMerge %147 None +OpBranchConditional %145 %146 %147 +%146 = OpLabel +%148 = OpRayQueryGetIntersectionInstanceCustomIndexKHR %7 %127 %35 +%149 = OpRayQueryGetIntersectionInstanceIdKHR %7 %127 %35 +%150 = OpRayQueryGetIntersectionInstanceShaderBindingTableRecordOffsetKHR %7 %127 %35 +%151 = OpRayQueryGetIntersectionGeometryIndexKHR %7 %127 %35 +%152 = OpRayQueryGetIntersectionPrimitiveIndexKHR %7 %127 %35 +%153 = OpRayQueryGetIntersectionObjectToWorldKHR %11 %127 %35 +%154 = OpRayQueryGetIntersectionWorldToObjectKHR %11 %127 %35 +%155 = OpAccessChain %33 %131 %90 +OpStore %155 %148 +%156 = OpAccessChain %33 %131 %28 +OpStore %156 %149 +%157 = OpAccessChain %33 %131 %23 +OpStore %157 %150 +%159 = OpAccessChain %33 %131 %158 +OpStore %159 %151 +%161 = OpAccessChain %33 %131 %160 +OpStore %161 %152 +%163 = OpAccessChain %122 %131 %162 +OpStore %163 %153 +%165 = OpAccessChain %122 %131 %164 +OpStore %165 %154 +%166 = OpIEqual %10 %143 %30 +OpSelectionMerge %168 None +OpBranchConditional %145 %167 %168 +%167 = OpLabel +%169 = OpRayQueryGetIntersectionTKHR %5 %127 %35 +%170 = OpAccessChain %36 %131 %30 +OpStore %170 %169 +%171 = OpRayQueryGetIntersectionBarycentricsKHR %9 %127 %35 +%172 = OpRayQueryGetIntersectionFrontFaceKHR %10 %127 %35 +%174 = OpAccessChain %123 %131 %173 +OpStore %174 %171 +%176 = OpAccessChain %124 %131 %175 +OpStore %176 %172 +OpBranch %168 +%168 = OpLabel +OpBranch %147 +%147 = OpLabel +OpBranch %140 +%140 = OpLabel +%177 = OpLoad %12 %131 +OpReturnValue %177 +OpFunctionEnd +%185 = OpFunction %2 None %184 +%186 = OpFunctionParameter %32 +%187 = OpFunctionParameter %33 +%188 = OpFunctionParameter %5 +%189 = OpFunctionParameter %36 +%190 = OpLabel +%191 = OpVariable %36 Function +%192 = OpVariable %36 Function +%195 = OpLoad %7 %187 +%196 = OpBitwiseAnd %7 %195 %90 +%197 = OpINotEqual %10 %196 %35 +%198 = OpBitwiseAnd %7 %195 %23 +%199 = OpINotEqual %10 %198 %35 +%200 = OpLogicalNot %10 %199 +%201 = OpLogicalAnd %10 %200 %197 +OpSelectionMerge %194 None +OpBranchConditional %201 %193 %194 +%193 = OpLabel +%202 = OpRayQueryGetIntersectionTypeKHR %7 %186 %35 +%203 = OpIEqual %10 %202 %30 +%204 = OpRayQueryGetRayTMinKHR %5 %186 +%205 = OpRayQueryGetIntersectionTypeKHR %7 %186 %30 +%206 = OpIEqual %10 %205 %35 +OpSelectionMerge %207 None +OpBranchConditional %206 %208 %209 +%208 = OpLabel +%210 = OpLoad %5 %189 +OpStore %192 %210 +OpBranch %207 +%209 = OpLabel +%211 = OpRayQueryGetIntersectionTKHR %5 %186 %35 +OpStore %192 %211 +OpBranch %207 +%207 = OpLabel +%212 = OpFOrdGreaterThanEqual %10 %188 %204 +%213 = OpLoad %5 %192 +%214 = OpFOrdLessThanEqual %10 %188 %213 +%215 = OpLogicalAnd %10 %212 %214 +%216 = OpLogicalAnd %10 %215 %203 +OpSelectionMerge %218 None +OpBranchConditional %216 %217 %218 +%217 = OpLabel +OpRayQueryGenerateIntersectionKHR %186 %188 +OpBranch %218 +%218 = OpLabel +OpBranch %194 +%194 = OpLabel +OpReturn +OpFunctionEnd +%226 = OpFunction %2 None %225 +%227 = OpFunctionParameter %32 +%228 = OpFunctionParameter %33 +%229 = OpLabel +%232 = OpLoad %7 %228 +%233 = OpBitwiseAnd %7 %232 %90 +%234 = OpINotEqual %10 %233 %35 +%235 = OpBitwiseAnd %7 %232 %23 +%236 = OpINotEqual %10 %235 %35 +%237 = OpLogicalNot %10 %236 +%238 = OpLogicalAnd %10 %237 %234 +OpSelectionMerge %231 None +OpBranchConditional %238 %230 %231 +%230 = OpLabel +%239 = OpRayQueryGetIntersectionTypeKHR %7 %227 %35 +%240 = OpIEqual %10 %239 %35 +OpSelectionMerge %242 None +OpBranchConditional %240 %241 %242 +%241 = OpLabel +OpRayQueryConfirmIntersectionKHR %227 +OpBranch %242 +%242 = OpLabel +OpBranch %231 +%231 = OpLabel +OpReturn OpFunctionEnd %16 = OpFunction %2 None %17 %15 = OpLabel %31 = OpVariable %32 Function +%34 = OpVariable %33 Function %35 +%37 = OpVariable %36 Function %19 %18 = OpLoad %4 %13 -OpBranch %33 -%33 = OpLabel -%34 = OpCompositeExtract %7 %27 0 -%35 = OpCompositeExtract %7 %27 1 -%36 = OpCompositeExtract %5 %27 2 -%37 = OpCompositeExtract %5 %27 3 -%38 = OpCompositeExtract %6 %27 4 -%39 = OpCompositeExtract %6 %27 5 -OpRayQueryInitializeKHR %31 %18 %34 %35 %38 %36 %39 %37 -%91 = OpFunctionCall %12 %47 %31 -%92 = OpCompositeExtract %7 %91 0 -%93 = OpIEqual %10 %92 %28 -OpSelectionMerge %94 None -OpBranchConditional %93 %95 %96 -%95 = OpLabel -OpRayQueryGenerateIntersectionKHR %31 %29 +OpBranch %38 +%38 = OpLabel +%120 = OpFunctionCall %2 %41 %31 %18 %27 %34 %37 +%178 = OpFunctionCall %12 %126 %31 %34 +%179 = OpCompositeExtract %7 %178 0 +%180 = OpIEqual %10 %179 %28 +OpSelectionMerge %181 None +OpBranchConditional %180 %182 %183 +%182 = OpLabel +%219 = OpFunctionCall %2 %185 %31 %34 %29 %37 OpReturn -%96 = OpLabel -%97 = OpCompositeExtract %7 %91 0 -%98 = OpIEqual %10 %97 %30 -OpSelectionMerge %99 None -OpBranchConditional %98 %100 %101 -%100 = OpLabel -OpRayQueryConfirmIntersectionKHR %31 +%183 = OpLabel +%220 = OpCompositeExtract %7 %178 0 +%221 = OpIEqual %10 %220 %30 +OpSelectionMerge %222 None +OpBranchConditional %221 %223 %224 +%223 = OpLabel +%243 = OpFunctionCall %2 %226 %31 %34 OpReturn -%101 = OpLabel +%224 = OpLabel OpReturn -%99 = OpLabel -OpBranch %94 -%94 = OpLabel +%222 = OpLabel +OpBranch %181 +%181 = OpLabel OpReturn OpFunctionEnd \ No newline at end of file diff --git a/naga/tests/out/spv/wgsl-overrides-ray-query.main.spvasm b/naga/tests/out/spv/wgsl-overrides-ray-query.main.spvasm index 34a8df87711..1a48f7a9f9a 100644 --- a/naga/tests/out/spv/wgsl-overrides-ray-query.main.spvasm +++ b/naga/tests/out/spv/wgsl-overrides-ray-query.main.spvasm @@ -1,7 +1,7 @@ ; SPIR-V ; Version: 1.4 ; Generator: rspirv -; Bound: 65 +; Bound: 164 OpCapability Shader OpCapability RayQueryKHR OpExtension "SPV_KHR_ray_query" @@ -40,61 +40,175 @@ OpDecorate %10 Binding 0 %25 = OpConstantComposite %7 %22 %23 %24 %26 = OpConstantComposite %8 %16 %17 %18 %19 %21 %25 %28 = OpTypePointer Function %5 -%40 = OpTypeVector %6 2 -%41 = OpTypePointer Function %40 -%42 = OpTypeBool -%43 = OpTypeVector %42 2 -%44 = OpConstant %6 0 -%45 = OpConstantComposite %40 %44 %44 -%46 = OpConstant %6 1 -%47 = OpConstant %6 4294967295 -%48 = OpConstantComposite %40 %47 %47 +%29 = OpTypePointer Function %6 +%31 = OpConstant %6 0 +%32 = OpTypePointer Function %3 +%34 = OpConstant %3 0 +%36 = OpTypeBool +%37 = OpTypeVector %36 3 +%38 = OpTypeFunction %2 %28 %4 %8 %29 %32 +%66 = OpConstant %6 256 +%69 = OpConstant %6 512 +%74 = OpConstant %6 16 +%77 = OpConstant %6 32 +%86 = OpConstant %6 1 +%89 = OpConstant %6 2 +%92 = OpConstant %6 64 +%95 = OpConstant %6 128 +%124 = OpTypeVector %6 2 +%125 = OpTypePointer Function %124 +%126 = OpTypeVector %36 2 +%127 = OpConstantComposite %124 %31 %31 +%128 = OpConstant %6 4294967295 +%129 = OpConstantComposite %124 %128 %128 +%142 = OpTypePointer Function %36 +%143 = OpTypeFunction %36 %28 %29 +%149 = OpConstantFalse %36 +%156 = OpConstant %6 6 +%39 = OpFunction %2 None %38 +%40 = OpFunctionParameter %28 +%41 = OpFunctionParameter %4 +%42 = OpFunctionParameter %8 +%43 = OpFunctionParameter %29 +%44 = OpFunctionParameter %32 +%45 = OpLabel +%46 = OpCompositeExtract %6 %42 0 +%47 = OpCompositeExtract %6 %42 1 +%48 = OpCompositeExtract %3 %42 2 +%49 = OpCompositeExtract %3 %42 3 +OpStore %44 %49 +%50 = OpCompositeExtract %7 %42 4 +%51 = OpCompositeExtract %7 %42 5 +%52 = OpFOrdLessThanEqual %36 %48 %49 +%53 = OpFOrdGreaterThanEqual %36 %48 %34 +%54 = OpIsInf %37 %50 +%55 = OpAny %36 %54 +%56 = OpIsNan %37 %50 +%57 = OpAny %36 %56 +%58 = OpLogicalOr %36 %57 %55 +%59 = OpLogicalNot %36 %58 +%60 = OpIsInf %37 %51 +%61 = OpAny %36 %60 +%62 = OpIsNan %37 %51 +%63 = OpAny %36 %62 +%64 = OpLogicalOr %36 %63 %61 +%65 = OpLogicalNot %36 %64 +%67 = OpBitwiseAnd %6 %46 %66 +%68 = OpINotEqual %36 %67 %31 +%70 = OpBitwiseAnd %6 %46 %69 +%71 = OpINotEqual %36 %70 %31 +%72 = OpLogicalAnd %36 %71 %68 +%73 = OpLogicalNot %36 %72 +%75 = OpBitwiseAnd %6 %46 %74 +%76 = OpINotEqual %36 %75 %31 +%78 = OpBitwiseAnd %6 %46 %77 +%79 = OpINotEqual %36 %78 %31 +%80 = OpLogicalAnd %36 %79 %68 +%81 = OpLogicalAnd %36 %79 %76 +%82 = OpLogicalAnd %36 %76 %68 +%83 = OpLogicalOr %36 %82 %80 +%84 = OpLogicalOr %36 %83 %81 +%85 = OpLogicalNot %36 %84 +%87 = OpBitwiseAnd %6 %46 %86 +%88 = OpINotEqual %36 %87 %31 +%90 = OpBitwiseAnd %6 %46 %89 +%91 = OpINotEqual %36 %90 %31 +%93 = OpBitwiseAnd %6 %46 %92 +%94 = OpINotEqual %36 %93 %31 +%96 = OpBitwiseAnd %6 %46 %95 +%97 = OpINotEqual %36 %96 %31 +%98 = OpLogicalAnd %36 %97 %88 +%99 = OpLogicalAnd %36 %97 %91 +%100 = OpLogicalAnd %36 %97 %94 +%101 = OpLogicalAnd %36 %94 %88 +%102 = OpLogicalAnd %36 %94 %91 +%103 = OpLogicalAnd %36 %91 %88 +%104 = OpLogicalOr %36 %103 %98 +%105 = OpLogicalOr %36 %104 %99 +%106 = OpLogicalOr %36 %105 %100 +%107 = OpLogicalOr %36 %106 %101 +%108 = OpLogicalOr %36 %107 %102 +%109 = OpLogicalNot %36 %108 +%110 = OpLogicalAnd %36 %109 %52 +%111 = OpLogicalAnd %36 %110 %53 +%112 = OpLogicalAnd %36 %111 %59 +%113 = OpLogicalAnd %36 %112 %65 +%114 = OpLogicalAnd %36 %113 %73 +%115 = OpLogicalAnd %36 %114 %85 +OpSelectionMerge %116 None +OpBranchConditional %115 %118 %117 +%118 = OpLabel +OpRayQueryInitializeKHR %40 %41 %46 %47 %50 %48 %51 %49 +OpStore %43 %86 +OpBranch %116 +%117 = OpLabel +OpBranch %116 +%116 = OpLabel +OpReturn +OpFunctionEnd +%144 = OpFunction %36 None %143 +%145 = OpFunctionParameter %28 +%146 = OpFunctionParameter %29 +%147 = OpLabel +%148 = OpVariable %142 Function %149 +%150 = OpLoad %6 %146 +%153 = OpBitwiseAnd %6 %150 %86 +%154 = OpINotEqual %36 %153 %31 +OpSelectionMerge %151 None +OpBranchConditional %154 %152 %151 +%152 = OpLabel +%155 = OpRayQueryProceedKHR %36 %145 +OpStore %148 %155 +%157 = OpSelect %6 %155 %89 %156 +%158 = OpBitwiseOr %6 %150 %157 +OpStore %146 %158 +OpBranch %151 +%151 = OpLabel +%159 = OpLoad %36 %148 +OpReturnValue %159 +OpFunctionEnd %13 = OpFunction %2 None %14 %12 = OpLabel %27 = OpVariable %28 Function -%49 = OpVariable %41 Function %48 +%30 = OpVariable %29 Function %31 +%33 = OpVariable %32 Function %34 +%130 = OpVariable %125 Function %129 %15 = OpLoad %4 %10 -OpBranch %29 -%29 = OpLabel -%30 = OpCompositeExtract %6 %26 0 -%31 = OpCompositeExtract %6 %26 1 -%32 = OpCompositeExtract %3 %26 2 -%33 = OpCompositeExtract %3 %26 3 -%34 = OpCompositeExtract %7 %26 4 -%35 = OpCompositeExtract %7 %26 5 -OpRayQueryInitializeKHR %27 %15 %30 %31 %34 %32 %35 %33 -OpBranch %36 -%36 = OpLabel -OpLoopMerge %37 %39 None -OpBranch %50 -%50 = OpLabel -%51 = OpLoad %40 %49 -%52 = OpIEqual %43 %45 %51 -%53 = OpAll %42 %52 -OpSelectionMerge %54 None -OpBranchConditional %53 %37 %54 -%54 = OpLabel -%55 = OpCompositeExtract %6 %51 1 -%56 = OpIEqual %42 %55 %44 -%57 = OpSelect %6 %56 %46 %44 -%58 = OpCompositeConstruct %40 %57 %46 -%59 = OpISub %40 %51 %58 -OpStore %49 %59 -OpBranch %38 -%38 = OpLabel -%60 = OpRayQueryProceedKHR %42 %27 -OpSelectionMerge %61 None -OpBranchConditional %60 %61 %62 -%62 = OpLabel -OpBranch %37 -%61 = OpLabel -OpBranch %63 -%63 = OpLabel -OpBranch %64 -%64 = OpLabel -OpBranch %39 -%39 = OpLabel -OpBranch %36 -%37 = OpLabel +OpBranch %35 +%35 = OpLabel +%119 = OpFunctionCall %2 %39 %27 %15 %26 %30 %33 +OpBranch %120 +%120 = OpLabel +OpLoopMerge %121 %123 None +OpBranch %131 +%131 = OpLabel +%132 = OpLoad %124 %130 +%133 = OpIEqual %126 %127 %132 +%134 = OpAll %36 %133 +OpSelectionMerge %135 None +OpBranchConditional %134 %121 %135 +%135 = OpLabel +%136 = OpCompositeExtract %6 %132 1 +%137 = OpIEqual %36 %136 %31 +%138 = OpSelect %6 %137 %86 %31 +%139 = OpCompositeConstruct %124 %138 %86 +%140 = OpISub %124 %132 %139 +OpStore %130 %140 +OpBranch %122 +%122 = OpLabel +%141 = OpFunctionCall %36 %144 %27 %30 +OpSelectionMerge %160 None +OpBranchConditional %141 %160 %161 +%161 = OpLabel +OpBranch %121 +%160 = OpLabel +OpBranch %162 +%162 = OpLabel +OpBranch %163 +%163 = OpLabel +OpBranch %123 +%123 = OpLabel +OpBranch %120 +%121 = OpLabel OpReturn OpFunctionEnd \ No newline at end of file diff --git a/naga/tests/out/spv/wgsl-ray-query-no-init-tracking.spvasm b/naga/tests/out/spv/wgsl-ray-query-no-init-tracking.spvasm new file mode 100644 index 00000000000..c52bec02ad8 --- /dev/null +++ b/naga/tests/out/spv/wgsl-ray-query-no-init-tracking.spvasm @@ -0,0 +1,542 @@ +; SPIR-V +; Version: 1.4 +; Generator: rspirv +; Bound: 382 +OpCapability Shader +OpCapability RayQueryKHR +OpExtension "SPV_KHR_ray_query" +%1 = OpExtInstImport "GLSL.std.450" +OpMemoryModel Logical GLSL450 +OpEntryPoint GLCompute %242 "main" %15 %17 +OpEntryPoint GLCompute %262 "main_candidate" %15 +OpExecutionMode %242 LocalSize 1 1 1 +OpExecutionMode %262 LocalSize 1 1 1 +OpMemberDecorate %10 0 Offset 0 +OpMemberDecorate %10 1 Offset 4 +OpMemberDecorate %10 2 Offset 8 +OpMemberDecorate %10 3 Offset 12 +OpMemberDecorate %10 4 Offset 16 +OpMemberDecorate %10 5 Offset 20 +OpMemberDecorate %10 6 Offset 24 +OpMemberDecorate %10 7 Offset 28 +OpMemberDecorate %10 8 Offset 36 +OpMemberDecorate %10 9 Offset 48 +OpMemberDecorate %10 9 ColMajor +OpMemberDecorate %10 9 MatrixStride 16 +OpMemberDecorate %10 10 Offset 112 +OpMemberDecorate %10 10 ColMajor +OpMemberDecorate %10 10 MatrixStride 16 +OpMemberDecorate %12 0 Offset 0 +OpMemberDecorate %12 1 Offset 4 +OpMemberDecorate %12 2 Offset 8 +OpMemberDecorate %12 3 Offset 12 +OpMemberDecorate %12 4 Offset 16 +OpMemberDecorate %12 5 Offset 32 +OpMemberDecorate %13 0 Offset 0 +OpMemberDecorate %13 1 Offset 16 +OpDecorate %15 DescriptorSet 0 +OpDecorate %15 Binding 0 +OpDecorate %17 DescriptorSet 0 +OpDecorate %17 Binding 1 +OpDecorate %18 Block +OpMemberDecorate %18 0 Offset 0 +%2 = OpTypeVoid +%3 = OpTypeFloat 32 +%4 = OpTypeVector %3 3 +%5 = OpTypeAccelerationStructureKHR +%6 = OpTypeInt 32 0 +%7 = OpTypeVector %3 2 +%8 = OpTypeBool +%9 = OpTypeMatrix %4 4 +%10 = OpTypeStruct %6 %3 %6 %6 %6 %6 %6 %7 %8 %9 %9 +%11 = OpTypeRayQueryKHR +%12 = OpTypeStruct %6 %6 %3 %3 %4 %4 +%13 = OpTypeStruct %6 %4 +%14 = OpTypeVector %3 4 +%16 = OpTypePointer UniformConstant %5 +%15 = OpVariable %16 UniformConstant +%18 = OpTypeStruct %13 +%19 = OpTypePointer StorageBuffer %18 +%17 = OpVariable %19 StorageBuffer +%26 = OpTypeFunction %10 %4 %4 %16 +%27 = OpConstant %6 4 +%28 = OpConstant %6 255 +%29 = OpConstant %3 0.1 +%30 = OpConstant %3 100 +%32 = OpTypePointer Function %11 +%33 = OpTypePointer Function %6 +%35 = OpConstant %6 0 +%36 = OpTypePointer Function %3 +%38 = OpConstant %3 0 +%41 = OpTypeVector %8 3 +%42 = OpTypeFunction %2 %32 %5 %12 %33 %36 +%70 = OpConstant %6 256 +%73 = OpConstant %6 512 +%78 = OpConstant %6 16 +%81 = OpConstant %6 32 +%90 = OpConstant %6 1 +%93 = OpConstant %6 2 +%96 = OpConstant %6 64 +%99 = OpConstant %6 128 +%128 = OpTypeVector %6 2 +%129 = OpTypePointer Function %128 +%130 = OpTypeVector %8 2 +%131 = OpConstantComposite %128 %35 %35 +%132 = OpConstant %6 4294967295 +%133 = OpConstantComposite %128 %132 %132 +%146 = OpTypePointer Function %8 +%147 = OpTypeFunction %8 %32 %33 +%153 = OpConstantFalse %8 +%160 = OpConstant %6 6 +%168 = OpTypePointer Function %10 +%169 = OpTypePointer Function %9 +%170 = OpTypePointer Function %7 +%171 = OpTypeFunction %10 %32 %33 +%176 = OpConstantNull %10 +%199 = OpConstant %6 3 +%202 = OpConstant %6 5 +%205 = OpConstant %6 9 +%207 = OpConstant %6 10 +%216 = OpConstant %6 7 +%218 = OpConstant %6 8 +%226 = OpTypeFunction %4 %4 %10 +%227 = OpConstant %3 1 +%228 = OpConstant %3 2.4 +%243 = OpTypeFunction %2 +%245 = OpTypePointer StorageBuffer %13 +%247 = OpConstantComposite %4 %38 %38 %38 +%248 = OpConstantComposite %4 %38 %227 %38 +%251 = OpTypePointer StorageBuffer %6 +%256 = OpTypePointer StorageBuffer %4 +%264 = OpConstantComposite %12 %27 %28 %29 %30 %247 %248 +%265 = OpConstant %3 10 +%322 = OpTypeFunction %2 %32 %33 %3 %36 +%363 = OpTypeFunction %2 %32 %33 +%43 = OpFunction %2 None %42 +%44 = OpFunctionParameter %32 +%45 = OpFunctionParameter %5 +%46 = OpFunctionParameter %12 +%47 = OpFunctionParameter %33 +%48 = OpFunctionParameter %36 +%49 = OpLabel +%50 = OpCompositeExtract %6 %46 0 +%51 = OpCompositeExtract %6 %46 1 +%52 = OpCompositeExtract %3 %46 2 +%53 = OpCompositeExtract %3 %46 3 +OpStore %48 %53 +%54 = OpCompositeExtract %4 %46 4 +%55 = OpCompositeExtract %4 %46 5 +%56 = OpFOrdLessThanEqual %8 %52 %53 +%57 = OpFOrdGreaterThanEqual %8 %52 %38 +%58 = OpIsInf %41 %54 +%59 = OpAny %8 %58 +%60 = OpIsNan %41 %54 +%61 = OpAny %8 %60 +%62 = OpLogicalOr %8 %61 %59 +%63 = OpLogicalNot %8 %62 +%64 = OpIsInf %41 %55 +%65 = OpAny %8 %64 +%66 = OpIsNan %41 %55 +%67 = OpAny %8 %66 +%68 = OpLogicalOr %8 %67 %65 +%69 = OpLogicalNot %8 %68 +%71 = OpBitwiseAnd %6 %50 %70 +%72 = OpINotEqual %8 %71 %35 +%74 = OpBitwiseAnd %6 %50 %73 +%75 = OpINotEqual %8 %74 %35 +%76 = OpLogicalAnd %8 %75 %72 +%77 = OpLogicalNot %8 %76 +%79 = OpBitwiseAnd %6 %50 %78 +%80 = OpINotEqual %8 %79 %35 +%82 = OpBitwiseAnd %6 %50 %81 +%83 = OpINotEqual %8 %82 %35 +%84 = OpLogicalAnd %8 %83 %72 +%85 = OpLogicalAnd %8 %83 %80 +%86 = OpLogicalAnd %8 %80 %72 +%87 = OpLogicalOr %8 %86 %84 +%88 = OpLogicalOr %8 %87 %85 +%89 = OpLogicalNot %8 %88 +%91 = OpBitwiseAnd %6 %50 %90 +%92 = OpINotEqual %8 %91 %35 +%94 = OpBitwiseAnd %6 %50 %93 +%95 = OpINotEqual %8 %94 %35 +%97 = OpBitwiseAnd %6 %50 %96 +%98 = OpINotEqual %8 %97 %35 +%100 = OpBitwiseAnd %6 %50 %99 +%101 = OpINotEqual %8 %100 %35 +%102 = OpLogicalAnd %8 %101 %92 +%103 = OpLogicalAnd %8 %101 %95 +%104 = OpLogicalAnd %8 %101 %98 +%105 = OpLogicalAnd %8 %98 %92 +%106 = OpLogicalAnd %8 %98 %95 +%107 = OpLogicalAnd %8 %95 %92 +%108 = OpLogicalOr %8 %107 %102 +%109 = OpLogicalOr %8 %108 %103 +%110 = OpLogicalOr %8 %109 %104 +%111 = OpLogicalOr %8 %110 %105 +%112 = OpLogicalOr %8 %111 %106 +%113 = OpLogicalNot %8 %112 +%114 = OpLogicalAnd %8 %113 %56 +%115 = OpLogicalAnd %8 %114 %57 +%116 = OpLogicalAnd %8 %115 %63 +%117 = OpLogicalAnd %8 %116 %69 +%118 = OpLogicalAnd %8 %117 %77 +%119 = OpLogicalAnd %8 %118 %89 +OpSelectionMerge %120 None +OpBranchConditional %119 %122 %121 +%122 = OpLabel +OpRayQueryInitializeKHR %44 %45 %50 %51 %54 %52 %55 %53 +OpStore %47 %90 +OpBranch %120 +%121 = OpLabel +OpBranch %120 +%120 = OpLabel +OpReturn +OpFunctionEnd +%148 = OpFunction %8 None %147 +%149 = OpFunctionParameter %32 +%150 = OpFunctionParameter %33 +%151 = OpLabel +%152 = OpVariable %146 Function %153 +%154 = OpLoad %6 %150 +%157 = OpBitwiseAnd %6 %154 %90 +%158 = OpINotEqual %8 %157 %35 +OpSelectionMerge %155 None +OpBranchConditional %158 %156 %155 +%156 = OpLabel +%159 = OpRayQueryProceedKHR %8 %149 +OpStore %152 %159 +%161 = OpSelect %6 %159 %93 %160 +%162 = OpBitwiseOr %6 %154 %161 +OpStore %150 %162 +OpBranch %155 +%155 = OpLabel +%163 = OpLoad %8 %152 +OpReturnValue %163 +OpFunctionEnd +%172 = OpFunction %10 None %171 +%173 = OpFunctionParameter %32 +%174 = OpFunctionParameter %33 +%175 = OpLabel +%177 = OpVariable %168 Function %176 +%178 = OpLoad %6 %174 +%179 = OpBitwiseAnd %6 %178 %93 +%180 = OpINotEqual %8 %179 %35 +%181 = OpBitwiseAnd %6 %178 %27 +%182 = OpINotEqual %8 %181 %35 +%183 = OpLogicalAnd %8 %182 %180 +OpSelectionMerge %185 None +OpBranchConditional %183 %184 %185 +%184 = OpLabel +%186 = OpRayQueryGetIntersectionTypeKHR %6 %173 %90 +%187 = OpAccessChain %33 %177 %35 +OpStore %187 %186 +%188 = OpINotEqual %8 %186 %35 +OpSelectionMerge %190 None +OpBranchConditional %188 %189 %190 +%189 = OpLabel +%191 = OpRayQueryGetIntersectionInstanceCustomIndexKHR %6 %173 %90 +%192 = OpRayQueryGetIntersectionInstanceIdKHR %6 %173 %90 +%193 = OpRayQueryGetIntersectionInstanceShaderBindingTableRecordOffsetKHR %6 %173 %90 +%194 = OpRayQueryGetIntersectionGeometryIndexKHR %6 %173 %90 +%195 = OpRayQueryGetIntersectionPrimitiveIndexKHR %6 %173 %90 +%196 = OpRayQueryGetIntersectionObjectToWorldKHR %9 %173 %90 +%197 = OpRayQueryGetIntersectionWorldToObjectKHR %9 %173 %90 +%198 = OpAccessChain %33 %177 %93 +OpStore %198 %191 +%200 = OpAccessChain %33 %177 %199 +OpStore %200 %192 +%201 = OpAccessChain %33 %177 %27 +OpStore %201 %193 +%203 = OpAccessChain %33 %177 %202 +OpStore %203 %194 +%204 = OpAccessChain %33 %177 %160 +OpStore %204 %195 +%206 = OpAccessChain %169 %177 %205 +OpStore %206 %196 +%208 = OpAccessChain %169 %177 %207 +OpStore %208 %197 +%209 = OpIEqual %8 %186 %90 +%212 = OpRayQueryGetIntersectionTKHR %3 %173 %90 +%213 = OpAccessChain %36 %177 %90 +OpStore %213 %212 +OpSelectionMerge %211 None +OpBranchConditional %188 %210 %211 +%210 = OpLabel +%214 = OpRayQueryGetIntersectionBarycentricsKHR %7 %173 %90 +%215 = OpRayQueryGetIntersectionFrontFaceKHR %8 %173 %90 +%217 = OpAccessChain %170 %177 %216 +OpStore %217 %214 +%219 = OpAccessChain %146 %177 %218 +OpStore %219 %215 +OpBranch %211 +%211 = OpLabel +OpBranch %190 +%190 = OpLabel +OpBranch %185 +%185 = OpLabel +%220 = OpLoad %10 %177 +OpReturnValue %220 +OpFunctionEnd +%25 = OpFunction %10 None %26 +%21 = OpFunctionParameter %4 +%22 = OpFunctionParameter %4 +%23 = OpFunctionParameter %16 +%20 = OpLabel +%31 = OpVariable %32 Function +%34 = OpVariable %33 Function %35 +%37 = OpVariable %36 Function %38 +%134 = OpVariable %129 Function %133 +%24 = OpLoad %5 %23 +OpBranch %39 +%39 = OpLabel +%40 = OpCompositeConstruct %12 %27 %28 %29 %30 %21 %22 +%123 = OpFunctionCall %2 %43 %31 %24 %40 %34 %37 +OpBranch %124 +%124 = OpLabel +OpLoopMerge %125 %127 None +OpBranch %135 +%135 = OpLabel +%136 = OpLoad %128 %134 +%137 = OpIEqual %130 %131 %136 +%138 = OpAll %8 %137 +OpSelectionMerge %139 None +OpBranchConditional %138 %125 %139 +%139 = OpLabel +%140 = OpCompositeExtract %6 %136 1 +%141 = OpIEqual %8 %140 %35 +%142 = OpSelect %6 %141 %90 %35 +%143 = OpCompositeConstruct %128 %142 %90 +%144 = OpISub %128 %136 %143 +OpStore %134 %144 +OpBranch %126 +%126 = OpLabel +%145 = OpFunctionCall %8 %148 %31 %34 +OpSelectionMerge %164 None +OpBranchConditional %145 %164 %165 +%165 = OpLabel +OpBranch %125 +%164 = OpLabel +OpBranch %166 +%166 = OpLabel +OpBranch %167 +%167 = OpLabel +OpBranch %127 +%127 = OpLabel +OpBranch %124 +%125 = OpLabel +%221 = OpFunctionCall %10 %172 %31 %34 +OpReturnValue %221 +OpFunctionEnd +%225 = OpFunction %4 None %226 +%223 = OpFunctionParameter %4 +%224 = OpFunctionParameter %10 +%222 = OpLabel +OpBranch %229 +%229 = OpLabel +%230 = OpCompositeExtract %9 %224 10 +%231 = OpCompositeConstruct %14 %223 %227 +%232 = OpMatrixTimesVector %4 %230 %231 +%233 = OpVectorShuffle %7 %232 %232 0 1 +%234 = OpExtInst %7 %1 Normalize %233 +%235 = OpVectorTimesScalar %7 %234 %228 +%236 = OpCompositeExtract %9 %224 9 +%237 = OpCompositeConstruct %14 %235 %38 %227 +%238 = OpMatrixTimesVector %4 %236 %237 +%239 = OpFSub %4 %223 %238 +%240 = OpExtInst %4 %1 Normalize %239 +OpReturnValue %240 +OpFunctionEnd +%242 = OpFunction %2 None %243 +%241 = OpLabel +%244 = OpLoad %5 %15 +%246 = OpAccessChain %245 %17 %35 +OpBranch %249 +%249 = OpLabel +%250 = OpFunctionCall %10 %25 %247 %248 %15 +%252 = OpCompositeExtract %6 %250 0 +%253 = OpIEqual %8 %252 %35 +%254 = OpSelect %6 %253 %90 %35 +%255 = OpAccessChain %251 %246 %35 +OpStore %255 %254 +%257 = OpCompositeExtract %3 %250 1 +%258 = OpVectorTimesScalar %4 %248 %257 +%259 = OpFunctionCall %4 %225 %258 %250 +%260 = OpAccessChain %256 %246 %90 +OpStore %260 %259 +OpReturn +OpFunctionEnd +%271 = OpFunction %10 None %171 +%272 = OpFunctionParameter %32 +%273 = OpFunctionParameter %33 +%274 = OpLabel +%275 = OpVariable %168 Function %176 +%276 = OpLoad %6 %273 +%277 = OpBitwiseAnd %6 %276 %93 +%278 = OpINotEqual %8 %277 %35 +%279 = OpBitwiseAnd %6 %276 %27 +%280 = OpINotEqual %8 %279 %35 +%281 = OpLogicalNot %8 %280 +%282 = OpLogicalAnd %8 %281 %278 +OpSelectionMerge %284 None +OpBranchConditional %282 %283 %284 +%283 = OpLabel +%285 = OpRayQueryGetIntersectionTypeKHR %6 %272 %35 +%286 = OpIEqual %8 %285 %35 +%287 = OpSelect %6 %286 %90 %199 +%288 = OpAccessChain %33 %275 %35 +OpStore %288 %287 +%289 = OpINotEqual %8 %287 %35 +OpSelectionMerge %291 None +OpBranchConditional %289 %290 %291 +%290 = OpLabel +%292 = OpRayQueryGetIntersectionInstanceCustomIndexKHR %6 %272 %35 +%293 = OpRayQueryGetIntersectionInstanceIdKHR %6 %272 %35 +%294 = OpRayQueryGetIntersectionInstanceShaderBindingTableRecordOffsetKHR %6 %272 %35 +%295 = OpRayQueryGetIntersectionGeometryIndexKHR %6 %272 %35 +%296 = OpRayQueryGetIntersectionPrimitiveIndexKHR %6 %272 %35 +%297 = OpRayQueryGetIntersectionObjectToWorldKHR %9 %272 %35 +%298 = OpRayQueryGetIntersectionWorldToObjectKHR %9 %272 %35 +%299 = OpAccessChain %33 %275 %93 +OpStore %299 %292 +%300 = OpAccessChain %33 %275 %199 +OpStore %300 %293 +%301 = OpAccessChain %33 %275 %27 +OpStore %301 %294 +%302 = OpAccessChain %33 %275 %202 +OpStore %302 %295 +%303 = OpAccessChain %33 %275 %160 +OpStore %303 %296 +%304 = OpAccessChain %169 %275 %205 +OpStore %304 %297 +%305 = OpAccessChain %169 %275 %207 +OpStore %305 %298 +%306 = OpIEqual %8 %287 %90 +OpSelectionMerge %308 None +OpBranchConditional %289 %307 %308 +%307 = OpLabel +%309 = OpRayQueryGetIntersectionTKHR %3 %272 %35 +%310 = OpAccessChain %36 %275 %90 +OpStore %310 %309 +%311 = OpRayQueryGetIntersectionBarycentricsKHR %7 %272 %35 +%312 = OpRayQueryGetIntersectionFrontFaceKHR %8 %272 %35 +%313 = OpAccessChain %170 %275 %216 +OpStore %313 %311 +%314 = OpAccessChain %146 %275 %218 +OpStore %314 %312 +OpBranch %308 +%308 = OpLabel +OpBranch %291 +%291 = OpLabel +OpBranch %284 +%284 = OpLabel +%315 = OpLoad %10 %275 +OpReturnValue %315 +OpFunctionEnd +%323 = OpFunction %2 None %322 +%324 = OpFunctionParameter %32 +%325 = OpFunctionParameter %33 +%326 = OpFunctionParameter %3 +%327 = OpFunctionParameter %36 +%328 = OpLabel +%329 = OpVariable %36 Function +%330 = OpVariable %36 Function +%333 = OpLoad %6 %325 +%334 = OpBitwiseAnd %6 %333 %93 +%335 = OpINotEqual %8 %334 %35 +%336 = OpBitwiseAnd %6 %333 %27 +%337 = OpINotEqual %8 %336 %35 +%338 = OpLogicalNot %8 %337 +%339 = OpLogicalAnd %8 %338 %335 +OpSelectionMerge %332 None +OpBranchConditional %339 %331 %332 +%331 = OpLabel +%340 = OpRayQueryGetIntersectionTypeKHR %6 %324 %35 +%341 = OpIEqual %8 %340 %90 +%342 = OpRayQueryGetRayTMinKHR %3 %324 +%343 = OpRayQueryGetIntersectionTypeKHR %6 %324 %90 +%344 = OpIEqual %8 %343 %35 +OpSelectionMerge %345 None +OpBranchConditional %344 %346 %347 +%346 = OpLabel +%348 = OpLoad %3 %327 +OpStore %330 %348 +OpBranch %345 +%347 = OpLabel +%349 = OpRayQueryGetIntersectionTKHR %3 %324 %35 +OpStore %330 %349 +OpBranch %345 +%345 = OpLabel +%350 = OpFOrdGreaterThanEqual %8 %326 %342 +%351 = OpLoad %3 %330 +%352 = OpFOrdLessThanEqual %8 %326 %351 +%353 = OpLogicalAnd %8 %350 %352 +%354 = OpLogicalAnd %8 %353 %341 +OpSelectionMerge %356 None +OpBranchConditional %354 %355 %356 +%355 = OpLabel +OpRayQueryGenerateIntersectionKHR %324 %326 +OpBranch %356 +%356 = OpLabel +OpBranch %332 +%332 = OpLabel +OpReturn +OpFunctionEnd +%364 = OpFunction %2 None %363 +%365 = OpFunctionParameter %32 +%366 = OpFunctionParameter %33 +%367 = OpLabel +%370 = OpLoad %6 %366 +%371 = OpBitwiseAnd %6 %370 %93 +%372 = OpINotEqual %8 %371 %35 +%373 = OpBitwiseAnd %6 %370 %27 +%374 = OpINotEqual %8 %373 %35 +%375 = OpLogicalNot %8 %374 +%376 = OpLogicalAnd %8 %375 %372 +OpSelectionMerge %369 None +OpBranchConditional %376 %368 %369 +%368 = OpLabel +%377 = OpRayQueryGetIntersectionTypeKHR %6 %365 %35 +%378 = OpIEqual %8 %377 %35 +OpSelectionMerge %380 None +OpBranchConditional %378 %379 %380 +%379 = OpLabel +OpRayQueryConfirmIntersectionKHR %365 +OpBranch %380 +%380 = OpLabel +OpBranch %369 +%369 = OpLabel +OpReturn +OpFunctionEnd +%262 = OpFunction %2 None %243 +%261 = OpLabel +%266 = OpVariable %32 Function +%267 = OpVariable %33 Function %35 +%268 = OpVariable %36 Function %38 +%263 = OpLoad %5 %15 +OpBranch %269 +%269 = OpLabel +%270 = OpFunctionCall %2 %43 %266 %263 %264 %267 %268 +%316 = OpFunctionCall %10 %271 %266 %267 +%317 = OpCompositeExtract %6 %316 0 +%318 = OpIEqual %8 %317 %199 +OpSelectionMerge %319 None +OpBranchConditional %318 %320 %321 +%320 = OpLabel +%357 = OpFunctionCall %2 %323 %266 %267 %265 %268 +OpReturn +%321 = OpLabel +%358 = OpCompositeExtract %6 %316 0 +%359 = OpIEqual %8 %358 %90 +OpSelectionMerge %360 None +OpBranchConditional %359 %361 %362 +%361 = OpLabel +%381 = OpFunctionCall %2 %364 %266 %267 +OpReturn +%362 = OpLabel +OpReturn +%360 = OpLabel +OpBranch %319 +%319 = OpLabel +OpReturn +OpFunctionEnd \ No newline at end of file diff --git a/naga/tests/out/spv/wgsl-ray-query.spvasm b/naga/tests/out/spv/wgsl-ray-query.spvasm index d49ae40b2f8..c52bec02ad8 100644 --- a/naga/tests/out/spv/wgsl-ray-query.spvasm +++ b/naga/tests/out/spv/wgsl-ray-query.spvasm @@ -1,16 +1,16 @@ ; SPIR-V ; Version: 1.4 ; Generator: rspirv -; Bound: 218 +; Bound: 382 OpCapability Shader OpCapability RayQueryKHR OpExtension "SPV_KHR_ray_query" %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint GLCompute %140 "main" %15 %17 -OpEntryPoint GLCompute %160 "main_candidate" %15 -OpExecutionMode %140 LocalSize 1 1 1 -OpExecutionMode %160 LocalSize 1 1 1 +OpEntryPoint GLCompute %242 "main" %15 %17 +OpEntryPoint GLCompute %262 "main_candidate" %15 +OpExecutionMode %242 LocalSize 1 1 1 +OpExecutionMode %262 LocalSize 1 1 1 OpMemberDecorate %10 0 Offset 0 OpMemberDecorate %10 1 Offset 4 OpMemberDecorate %10 2 Offset 8 @@ -64,93 +64,219 @@ OpMemberDecorate %18 0 Offset 0 %29 = OpConstant %3 0.1 %30 = OpConstant %3 100 %32 = OpTypePointer Function %11 -%45 = OpTypeVector %6 2 -%46 = OpTypePointer Function %45 -%47 = OpTypeVector %8 2 -%48 = OpConstant %6 0 -%49 = OpConstantComposite %45 %48 %48 -%50 = OpConstant %6 1 -%51 = OpConstant %6 4294967295 -%52 = OpConstantComposite %45 %51 %51 -%69 = OpTypePointer Function %10 -%70 = OpTypePointer Function %6 -%71 = OpTypePointer Function %9 -%72 = OpTypePointer Function %7 -%73 = OpTypePointer Function %8 -%74 = OpTypePointer Function %3 -%75 = OpTypeFunction %10 %32 -%77 = OpConstantNull %10 +%33 = OpTypePointer Function %6 +%35 = OpConstant %6 0 +%36 = OpTypePointer Function %3 +%38 = OpConstant %3 0 +%41 = OpTypeVector %8 3 +%42 = OpTypeFunction %2 %32 %5 %12 %33 %36 +%70 = OpConstant %6 256 +%73 = OpConstant %6 512 +%78 = OpConstant %6 16 +%81 = OpConstant %6 32 +%90 = OpConstant %6 1 %93 = OpConstant %6 2 -%95 = OpConstant %6 3 -%98 = OpConstant %6 5 -%100 = OpConstant %6 6 -%102 = OpConstant %6 9 -%104 = OpConstant %6 10 -%113 = OpConstant %6 7 -%115 = OpConstant %6 8 -%123 = OpTypeFunction %4 %4 %10 -%124 = OpConstant %3 1 -%125 = OpConstant %3 2.4 -%126 = OpConstant %3 0 -%141 = OpTypeFunction %2 -%143 = OpTypePointer StorageBuffer %13 -%145 = OpConstantComposite %4 %126 %126 %126 -%146 = OpConstantComposite %4 %126 %124 %126 -%149 = OpTypePointer StorageBuffer %6 -%154 = OpTypePointer StorageBuffer %4 -%162 = OpConstantComposite %12 %27 %28 %29 %30 %145 %146 -%163 = OpConstant %3 10 -%76 = OpFunction %10 None %75 -%78 = OpFunctionParameter %32 -%79 = OpLabel -%80 = OpVariable %69 Function %77 -%81 = OpRayQueryGetIntersectionTypeKHR %6 %78 %50 -%82 = OpAccessChain %70 %80 %48 -OpStore %82 %81 -%83 = OpINotEqual %8 %81 %48 -OpSelectionMerge %85 None -OpBranchConditional %83 %84 %85 -%84 = OpLabel -%86 = OpRayQueryGetIntersectionInstanceCustomIndexKHR %6 %78 %50 -%87 = OpRayQueryGetIntersectionInstanceIdKHR %6 %78 %50 -%88 = OpRayQueryGetIntersectionInstanceShaderBindingTableRecordOffsetKHR %6 %78 %50 -%89 = OpRayQueryGetIntersectionGeometryIndexKHR %6 %78 %50 -%90 = OpRayQueryGetIntersectionPrimitiveIndexKHR %6 %78 %50 -%91 = OpRayQueryGetIntersectionObjectToWorldKHR %9 %78 %50 -%92 = OpRayQueryGetIntersectionWorldToObjectKHR %9 %78 %50 -%94 = OpAccessChain %70 %80 %93 -OpStore %94 %86 -%96 = OpAccessChain %70 %80 %95 -OpStore %96 %87 -%97 = OpAccessChain %70 %80 %27 -OpStore %97 %88 -%99 = OpAccessChain %70 %80 %98 -OpStore %99 %89 -%101 = OpAccessChain %70 %80 %100 -OpStore %101 %90 -%103 = OpAccessChain %71 %80 %102 -OpStore %103 %91 -%105 = OpAccessChain %71 %80 %104 -OpStore %105 %92 -%106 = OpIEqual %8 %81 %50 -%109 = OpRayQueryGetIntersectionTKHR %3 %78 %50 -%110 = OpAccessChain %74 %80 %50 -OpStore %110 %109 -OpSelectionMerge %108 None -OpBranchConditional %83 %107 %108 -%107 = OpLabel -%111 = OpRayQueryGetIntersectionBarycentricsKHR %7 %78 %50 -%112 = OpRayQueryGetIntersectionFrontFaceKHR %8 %78 %50 -%114 = OpAccessChain %72 %80 %113 -OpStore %114 %111 -%116 = OpAccessChain %73 %80 %115 -OpStore %116 %112 -OpBranch %108 -%108 = OpLabel -OpBranch %85 -%85 = OpLabel -%117 = OpLoad %10 %80 -OpReturnValue %117 +%96 = OpConstant %6 64 +%99 = OpConstant %6 128 +%128 = OpTypeVector %6 2 +%129 = OpTypePointer Function %128 +%130 = OpTypeVector %8 2 +%131 = OpConstantComposite %128 %35 %35 +%132 = OpConstant %6 4294967295 +%133 = OpConstantComposite %128 %132 %132 +%146 = OpTypePointer Function %8 +%147 = OpTypeFunction %8 %32 %33 +%153 = OpConstantFalse %8 +%160 = OpConstant %6 6 +%168 = OpTypePointer Function %10 +%169 = OpTypePointer Function %9 +%170 = OpTypePointer Function %7 +%171 = OpTypeFunction %10 %32 %33 +%176 = OpConstantNull %10 +%199 = OpConstant %6 3 +%202 = OpConstant %6 5 +%205 = OpConstant %6 9 +%207 = OpConstant %6 10 +%216 = OpConstant %6 7 +%218 = OpConstant %6 8 +%226 = OpTypeFunction %4 %4 %10 +%227 = OpConstant %3 1 +%228 = OpConstant %3 2.4 +%243 = OpTypeFunction %2 +%245 = OpTypePointer StorageBuffer %13 +%247 = OpConstantComposite %4 %38 %38 %38 +%248 = OpConstantComposite %4 %38 %227 %38 +%251 = OpTypePointer StorageBuffer %6 +%256 = OpTypePointer StorageBuffer %4 +%264 = OpConstantComposite %12 %27 %28 %29 %30 %247 %248 +%265 = OpConstant %3 10 +%322 = OpTypeFunction %2 %32 %33 %3 %36 +%363 = OpTypeFunction %2 %32 %33 +%43 = OpFunction %2 None %42 +%44 = OpFunctionParameter %32 +%45 = OpFunctionParameter %5 +%46 = OpFunctionParameter %12 +%47 = OpFunctionParameter %33 +%48 = OpFunctionParameter %36 +%49 = OpLabel +%50 = OpCompositeExtract %6 %46 0 +%51 = OpCompositeExtract %6 %46 1 +%52 = OpCompositeExtract %3 %46 2 +%53 = OpCompositeExtract %3 %46 3 +OpStore %48 %53 +%54 = OpCompositeExtract %4 %46 4 +%55 = OpCompositeExtract %4 %46 5 +%56 = OpFOrdLessThanEqual %8 %52 %53 +%57 = OpFOrdGreaterThanEqual %8 %52 %38 +%58 = OpIsInf %41 %54 +%59 = OpAny %8 %58 +%60 = OpIsNan %41 %54 +%61 = OpAny %8 %60 +%62 = OpLogicalOr %8 %61 %59 +%63 = OpLogicalNot %8 %62 +%64 = OpIsInf %41 %55 +%65 = OpAny %8 %64 +%66 = OpIsNan %41 %55 +%67 = OpAny %8 %66 +%68 = OpLogicalOr %8 %67 %65 +%69 = OpLogicalNot %8 %68 +%71 = OpBitwiseAnd %6 %50 %70 +%72 = OpINotEqual %8 %71 %35 +%74 = OpBitwiseAnd %6 %50 %73 +%75 = OpINotEqual %8 %74 %35 +%76 = OpLogicalAnd %8 %75 %72 +%77 = OpLogicalNot %8 %76 +%79 = OpBitwiseAnd %6 %50 %78 +%80 = OpINotEqual %8 %79 %35 +%82 = OpBitwiseAnd %6 %50 %81 +%83 = OpINotEqual %8 %82 %35 +%84 = OpLogicalAnd %8 %83 %72 +%85 = OpLogicalAnd %8 %83 %80 +%86 = OpLogicalAnd %8 %80 %72 +%87 = OpLogicalOr %8 %86 %84 +%88 = OpLogicalOr %8 %87 %85 +%89 = OpLogicalNot %8 %88 +%91 = OpBitwiseAnd %6 %50 %90 +%92 = OpINotEqual %8 %91 %35 +%94 = OpBitwiseAnd %6 %50 %93 +%95 = OpINotEqual %8 %94 %35 +%97 = OpBitwiseAnd %6 %50 %96 +%98 = OpINotEqual %8 %97 %35 +%100 = OpBitwiseAnd %6 %50 %99 +%101 = OpINotEqual %8 %100 %35 +%102 = OpLogicalAnd %8 %101 %92 +%103 = OpLogicalAnd %8 %101 %95 +%104 = OpLogicalAnd %8 %101 %98 +%105 = OpLogicalAnd %8 %98 %92 +%106 = OpLogicalAnd %8 %98 %95 +%107 = OpLogicalAnd %8 %95 %92 +%108 = OpLogicalOr %8 %107 %102 +%109 = OpLogicalOr %8 %108 %103 +%110 = OpLogicalOr %8 %109 %104 +%111 = OpLogicalOr %8 %110 %105 +%112 = OpLogicalOr %8 %111 %106 +%113 = OpLogicalNot %8 %112 +%114 = OpLogicalAnd %8 %113 %56 +%115 = OpLogicalAnd %8 %114 %57 +%116 = OpLogicalAnd %8 %115 %63 +%117 = OpLogicalAnd %8 %116 %69 +%118 = OpLogicalAnd %8 %117 %77 +%119 = OpLogicalAnd %8 %118 %89 +OpSelectionMerge %120 None +OpBranchConditional %119 %122 %121 +%122 = OpLabel +OpRayQueryInitializeKHR %44 %45 %50 %51 %54 %52 %55 %53 +OpStore %47 %90 +OpBranch %120 +%121 = OpLabel +OpBranch %120 +%120 = OpLabel +OpReturn +OpFunctionEnd +%148 = OpFunction %8 None %147 +%149 = OpFunctionParameter %32 +%150 = OpFunctionParameter %33 +%151 = OpLabel +%152 = OpVariable %146 Function %153 +%154 = OpLoad %6 %150 +%157 = OpBitwiseAnd %6 %154 %90 +%158 = OpINotEqual %8 %157 %35 +OpSelectionMerge %155 None +OpBranchConditional %158 %156 %155 +%156 = OpLabel +%159 = OpRayQueryProceedKHR %8 %149 +OpStore %152 %159 +%161 = OpSelect %6 %159 %93 %160 +%162 = OpBitwiseOr %6 %154 %161 +OpStore %150 %162 +OpBranch %155 +%155 = OpLabel +%163 = OpLoad %8 %152 +OpReturnValue %163 +OpFunctionEnd +%172 = OpFunction %10 None %171 +%173 = OpFunctionParameter %32 +%174 = OpFunctionParameter %33 +%175 = OpLabel +%177 = OpVariable %168 Function %176 +%178 = OpLoad %6 %174 +%179 = OpBitwiseAnd %6 %178 %93 +%180 = OpINotEqual %8 %179 %35 +%181 = OpBitwiseAnd %6 %178 %27 +%182 = OpINotEqual %8 %181 %35 +%183 = OpLogicalAnd %8 %182 %180 +OpSelectionMerge %185 None +OpBranchConditional %183 %184 %185 +%184 = OpLabel +%186 = OpRayQueryGetIntersectionTypeKHR %6 %173 %90 +%187 = OpAccessChain %33 %177 %35 +OpStore %187 %186 +%188 = OpINotEqual %8 %186 %35 +OpSelectionMerge %190 None +OpBranchConditional %188 %189 %190 +%189 = OpLabel +%191 = OpRayQueryGetIntersectionInstanceCustomIndexKHR %6 %173 %90 +%192 = OpRayQueryGetIntersectionInstanceIdKHR %6 %173 %90 +%193 = OpRayQueryGetIntersectionInstanceShaderBindingTableRecordOffsetKHR %6 %173 %90 +%194 = OpRayQueryGetIntersectionGeometryIndexKHR %6 %173 %90 +%195 = OpRayQueryGetIntersectionPrimitiveIndexKHR %6 %173 %90 +%196 = OpRayQueryGetIntersectionObjectToWorldKHR %9 %173 %90 +%197 = OpRayQueryGetIntersectionWorldToObjectKHR %9 %173 %90 +%198 = OpAccessChain %33 %177 %93 +OpStore %198 %191 +%200 = OpAccessChain %33 %177 %199 +OpStore %200 %192 +%201 = OpAccessChain %33 %177 %27 +OpStore %201 %193 +%203 = OpAccessChain %33 %177 %202 +OpStore %203 %194 +%204 = OpAccessChain %33 %177 %160 +OpStore %204 %195 +%206 = OpAccessChain %169 %177 %205 +OpStore %206 %196 +%208 = OpAccessChain %169 %177 %207 +OpStore %208 %197 +%209 = OpIEqual %8 %186 %90 +%212 = OpRayQueryGetIntersectionTKHR %3 %173 %90 +%213 = OpAccessChain %36 %177 %90 +OpStore %213 %212 +OpSelectionMerge %211 None +OpBranchConditional %188 %210 %211 +%210 = OpLabel +%214 = OpRayQueryGetIntersectionBarycentricsKHR %7 %173 %90 +%215 = OpRayQueryGetIntersectionFrontFaceKHR %8 %173 %90 +%217 = OpAccessChain %170 %177 %216 +OpStore %217 %214 +%219 = OpAccessChain %146 %177 %218 +OpStore %219 %215 +OpBranch %211 +%211 = OpLabel +OpBranch %190 +%190 = OpLabel +OpBranch %185 +%185 = OpLabel +%220 = OpLoad %10 %177 +OpReturnValue %220 OpFunctionEnd %25 = OpFunction %10 None %26 %21 = OpFunctionParameter %4 @@ -158,179 +284,259 @@ OpFunctionEnd %23 = OpFunctionParameter %16 %20 = OpLabel %31 = OpVariable %32 Function -%53 = OpVariable %46 Function %52 +%34 = OpVariable %33 Function %35 +%37 = OpVariable %36 Function %38 +%134 = OpVariable %129 Function %133 %24 = OpLoad %5 %23 -OpBranch %33 -%33 = OpLabel -%34 = OpCompositeConstruct %12 %27 %28 %29 %30 %21 %22 -%35 = OpCompositeExtract %6 %34 0 -%36 = OpCompositeExtract %6 %34 1 -%37 = OpCompositeExtract %3 %34 2 -%38 = OpCompositeExtract %3 %34 3 -%39 = OpCompositeExtract %4 %34 4 -%40 = OpCompositeExtract %4 %34 5 -OpRayQueryInitializeKHR %31 %24 %35 %36 %39 %37 %40 %38 -OpBranch %41 -%41 = OpLabel -OpLoopMerge %42 %44 None -OpBranch %54 -%54 = OpLabel -%55 = OpLoad %45 %53 -%56 = OpIEqual %47 %49 %55 -%57 = OpAll %8 %56 -OpSelectionMerge %58 None -OpBranchConditional %57 %42 %58 -%58 = OpLabel -%59 = OpCompositeExtract %6 %55 1 -%60 = OpIEqual %8 %59 %48 -%61 = OpSelect %6 %60 %50 %48 -%62 = OpCompositeConstruct %45 %61 %50 -%63 = OpISub %45 %55 %62 -OpStore %53 %63 -OpBranch %43 -%43 = OpLabel -%64 = OpRayQueryProceedKHR %8 %31 -OpSelectionMerge %65 None -OpBranchConditional %64 %65 %66 -%66 = OpLabel -OpBranch %42 -%65 = OpLabel -OpBranch %67 -%67 = OpLabel -OpBranch %68 -%68 = OpLabel -OpBranch %44 -%44 = OpLabel -OpBranch %41 -%42 = OpLabel -%118 = OpFunctionCall %10 %76 %31 -OpReturnValue %118 -OpFunctionEnd -%122 = OpFunction %4 None %123 -%120 = OpFunctionParameter %4 -%121 = OpFunctionParameter %10 -%119 = OpLabel +OpBranch %39 +%39 = OpLabel +%40 = OpCompositeConstruct %12 %27 %28 %29 %30 %21 %22 +%123 = OpFunctionCall %2 %43 %31 %24 %40 %34 %37 +OpBranch %124 +%124 = OpLabel +OpLoopMerge %125 %127 None +OpBranch %135 +%135 = OpLabel +%136 = OpLoad %128 %134 +%137 = OpIEqual %130 %131 %136 +%138 = OpAll %8 %137 +OpSelectionMerge %139 None +OpBranchConditional %138 %125 %139 +%139 = OpLabel +%140 = OpCompositeExtract %6 %136 1 +%141 = OpIEqual %8 %140 %35 +%142 = OpSelect %6 %141 %90 %35 +%143 = OpCompositeConstruct %128 %142 %90 +%144 = OpISub %128 %136 %143 +OpStore %134 %144 +OpBranch %126 +%126 = OpLabel +%145 = OpFunctionCall %8 %148 %31 %34 +OpSelectionMerge %164 None +OpBranchConditional %145 %164 %165 +%165 = OpLabel +OpBranch %125 +%164 = OpLabel +OpBranch %166 +%166 = OpLabel +OpBranch %167 +%167 = OpLabel OpBranch %127 %127 = OpLabel -%128 = OpCompositeExtract %9 %121 10 -%129 = OpCompositeConstruct %14 %120 %124 -%130 = OpMatrixTimesVector %4 %128 %129 -%131 = OpVectorShuffle %7 %130 %130 0 1 -%132 = OpExtInst %7 %1 Normalize %131 -%133 = OpVectorTimesScalar %7 %132 %125 -%134 = OpCompositeExtract %9 %121 9 -%135 = OpCompositeConstruct %14 %133 %126 %124 -%136 = OpMatrixTimesVector %4 %134 %135 -%137 = OpFSub %4 %120 %136 -%138 = OpExtInst %4 %1 Normalize %137 -OpReturnValue %138 +OpBranch %124 +%125 = OpLabel +%221 = OpFunctionCall %10 %172 %31 %34 +OpReturnValue %221 OpFunctionEnd -%140 = OpFunction %2 None %141 -%139 = OpLabel -%142 = OpLoad %5 %15 -%144 = OpAccessChain %143 %17 %48 -OpBranch %147 -%147 = OpLabel -%148 = OpFunctionCall %10 %25 %145 %146 %15 -%150 = OpCompositeExtract %6 %148 0 -%151 = OpIEqual %8 %150 %48 -%152 = OpSelect %6 %151 %50 %48 -%153 = OpAccessChain %149 %144 %48 -OpStore %153 %152 -%155 = OpCompositeExtract %3 %148 1 -%156 = OpVectorTimesScalar %4 %146 %155 -%157 = OpFunctionCall %4 %122 %156 %148 -%158 = OpAccessChain %154 %144 %50 -OpStore %158 %157 +%225 = OpFunction %4 None %226 +%223 = OpFunctionParameter %4 +%224 = OpFunctionParameter %10 +%222 = OpLabel +OpBranch %229 +%229 = OpLabel +%230 = OpCompositeExtract %9 %224 10 +%231 = OpCompositeConstruct %14 %223 %227 +%232 = OpMatrixTimesVector %4 %230 %231 +%233 = OpVectorShuffle %7 %232 %232 0 1 +%234 = OpExtInst %7 %1 Normalize %233 +%235 = OpVectorTimesScalar %7 %234 %228 +%236 = OpCompositeExtract %9 %224 9 +%237 = OpCompositeConstruct %14 %235 %38 %227 +%238 = OpMatrixTimesVector %4 %236 %237 +%239 = OpFSub %4 %223 %238 +%240 = OpExtInst %4 %1 Normalize %239 +OpReturnValue %240 +OpFunctionEnd +%242 = OpFunction %2 None %243 +%241 = OpLabel +%244 = OpLoad %5 %15 +%246 = OpAccessChain %245 %17 %35 +OpBranch %249 +%249 = OpLabel +%250 = OpFunctionCall %10 %25 %247 %248 %15 +%252 = OpCompositeExtract %6 %250 0 +%253 = OpIEqual %8 %252 %35 +%254 = OpSelect %6 %253 %90 %35 +%255 = OpAccessChain %251 %246 %35 +OpStore %255 %254 +%257 = OpCompositeExtract %3 %250 1 +%258 = OpVectorTimesScalar %4 %248 %257 +%259 = OpFunctionCall %4 %225 %258 %250 +%260 = OpAccessChain %256 %246 %90 +OpStore %260 %259 OpReturn OpFunctionEnd -%172 = OpFunction %10 None %75 -%173 = OpFunctionParameter %32 -%174 = OpLabel -%175 = OpVariable %69 Function %77 -%176 = OpRayQueryGetIntersectionTypeKHR %6 %173 %48 -%177 = OpIEqual %8 %176 %48 -%178 = OpSelect %6 %177 %50 %95 -%179 = OpAccessChain %70 %175 %48 -OpStore %179 %178 -%180 = OpINotEqual %8 %178 %48 -OpSelectionMerge %182 None -OpBranchConditional %180 %181 %182 -%181 = OpLabel -%183 = OpRayQueryGetIntersectionInstanceCustomIndexKHR %6 %173 %48 -%184 = OpRayQueryGetIntersectionInstanceIdKHR %6 %173 %48 -%185 = OpRayQueryGetIntersectionInstanceShaderBindingTableRecordOffsetKHR %6 %173 %48 -%186 = OpRayQueryGetIntersectionGeometryIndexKHR %6 %173 %48 -%187 = OpRayQueryGetIntersectionPrimitiveIndexKHR %6 %173 %48 -%188 = OpRayQueryGetIntersectionObjectToWorldKHR %9 %173 %48 -%189 = OpRayQueryGetIntersectionWorldToObjectKHR %9 %173 %48 -%190 = OpAccessChain %70 %175 %93 -OpStore %190 %183 -%191 = OpAccessChain %70 %175 %95 -OpStore %191 %184 -%192 = OpAccessChain %70 %175 %27 -OpStore %192 %185 -%193 = OpAccessChain %70 %175 %98 -OpStore %193 %186 -%194 = OpAccessChain %70 %175 %100 -OpStore %194 %187 -%195 = OpAccessChain %71 %175 %102 -OpStore %195 %188 -%196 = OpAccessChain %71 %175 %104 -OpStore %196 %189 -%197 = OpIEqual %8 %178 %50 -OpSelectionMerge %199 None -OpBranchConditional %180 %198 %199 -%198 = OpLabel -%200 = OpRayQueryGetIntersectionTKHR %3 %173 %48 -%201 = OpAccessChain %74 %175 %50 -OpStore %201 %200 -%202 = OpRayQueryGetIntersectionBarycentricsKHR %7 %173 %48 -%203 = OpRayQueryGetIntersectionFrontFaceKHR %8 %173 %48 -%204 = OpAccessChain %72 %175 %113 -OpStore %204 %202 -%205 = OpAccessChain %73 %175 %115 -OpStore %205 %203 -OpBranch %199 -%199 = OpLabel -OpBranch %182 -%182 = OpLabel -%206 = OpLoad %10 %175 -OpReturnValue %206 +%271 = OpFunction %10 None %171 +%272 = OpFunctionParameter %32 +%273 = OpFunctionParameter %33 +%274 = OpLabel +%275 = OpVariable %168 Function %176 +%276 = OpLoad %6 %273 +%277 = OpBitwiseAnd %6 %276 %93 +%278 = OpINotEqual %8 %277 %35 +%279 = OpBitwiseAnd %6 %276 %27 +%280 = OpINotEqual %8 %279 %35 +%281 = OpLogicalNot %8 %280 +%282 = OpLogicalAnd %8 %281 %278 +OpSelectionMerge %284 None +OpBranchConditional %282 %283 %284 +%283 = OpLabel +%285 = OpRayQueryGetIntersectionTypeKHR %6 %272 %35 +%286 = OpIEqual %8 %285 %35 +%287 = OpSelect %6 %286 %90 %199 +%288 = OpAccessChain %33 %275 %35 +OpStore %288 %287 +%289 = OpINotEqual %8 %287 %35 +OpSelectionMerge %291 None +OpBranchConditional %289 %290 %291 +%290 = OpLabel +%292 = OpRayQueryGetIntersectionInstanceCustomIndexKHR %6 %272 %35 +%293 = OpRayQueryGetIntersectionInstanceIdKHR %6 %272 %35 +%294 = OpRayQueryGetIntersectionInstanceShaderBindingTableRecordOffsetKHR %6 %272 %35 +%295 = OpRayQueryGetIntersectionGeometryIndexKHR %6 %272 %35 +%296 = OpRayQueryGetIntersectionPrimitiveIndexKHR %6 %272 %35 +%297 = OpRayQueryGetIntersectionObjectToWorldKHR %9 %272 %35 +%298 = OpRayQueryGetIntersectionWorldToObjectKHR %9 %272 %35 +%299 = OpAccessChain %33 %275 %93 +OpStore %299 %292 +%300 = OpAccessChain %33 %275 %199 +OpStore %300 %293 +%301 = OpAccessChain %33 %275 %27 +OpStore %301 %294 +%302 = OpAccessChain %33 %275 %202 +OpStore %302 %295 +%303 = OpAccessChain %33 %275 %160 +OpStore %303 %296 +%304 = OpAccessChain %169 %275 %205 +OpStore %304 %297 +%305 = OpAccessChain %169 %275 %207 +OpStore %305 %298 +%306 = OpIEqual %8 %287 %90 +OpSelectionMerge %308 None +OpBranchConditional %289 %307 %308 +%307 = OpLabel +%309 = OpRayQueryGetIntersectionTKHR %3 %272 %35 +%310 = OpAccessChain %36 %275 %90 +OpStore %310 %309 +%311 = OpRayQueryGetIntersectionBarycentricsKHR %7 %272 %35 +%312 = OpRayQueryGetIntersectionFrontFaceKHR %8 %272 %35 +%313 = OpAccessChain %170 %275 %216 +OpStore %313 %311 +%314 = OpAccessChain %146 %275 %218 +OpStore %314 %312 +OpBranch %308 +%308 = OpLabel +OpBranch %291 +%291 = OpLabel +OpBranch %284 +%284 = OpLabel +%315 = OpLoad %10 %275 +OpReturnValue %315 OpFunctionEnd -%160 = OpFunction %2 None %141 -%159 = OpLabel -%164 = OpVariable %32 Function -%161 = OpLoad %5 %15 -OpBranch %165 -%165 = OpLabel -%166 = OpCompositeExtract %6 %162 0 -%167 = OpCompositeExtract %6 %162 1 -%168 = OpCompositeExtract %3 %162 2 -%169 = OpCompositeExtract %3 %162 3 -%170 = OpCompositeExtract %4 %162 4 -%171 = OpCompositeExtract %4 %162 5 -OpRayQueryInitializeKHR %164 %161 %166 %167 %170 %168 %171 %169 -%207 = OpFunctionCall %10 %172 %164 -%208 = OpCompositeExtract %6 %207 0 -%209 = OpIEqual %8 %208 %95 -OpSelectionMerge %210 None -OpBranchConditional %209 %211 %212 -%211 = OpLabel -OpRayQueryGenerateIntersectionKHR %164 %163 +%323 = OpFunction %2 None %322 +%324 = OpFunctionParameter %32 +%325 = OpFunctionParameter %33 +%326 = OpFunctionParameter %3 +%327 = OpFunctionParameter %36 +%328 = OpLabel +%329 = OpVariable %36 Function +%330 = OpVariable %36 Function +%333 = OpLoad %6 %325 +%334 = OpBitwiseAnd %6 %333 %93 +%335 = OpINotEqual %8 %334 %35 +%336 = OpBitwiseAnd %6 %333 %27 +%337 = OpINotEqual %8 %336 %35 +%338 = OpLogicalNot %8 %337 +%339 = OpLogicalAnd %8 %338 %335 +OpSelectionMerge %332 None +OpBranchConditional %339 %331 %332 +%331 = OpLabel +%340 = OpRayQueryGetIntersectionTypeKHR %6 %324 %35 +%341 = OpIEqual %8 %340 %90 +%342 = OpRayQueryGetRayTMinKHR %3 %324 +%343 = OpRayQueryGetIntersectionTypeKHR %6 %324 %90 +%344 = OpIEqual %8 %343 %35 +OpSelectionMerge %345 None +OpBranchConditional %344 %346 %347 +%346 = OpLabel +%348 = OpLoad %3 %327 +OpStore %330 %348 +OpBranch %345 +%347 = OpLabel +%349 = OpRayQueryGetIntersectionTKHR %3 %324 %35 +OpStore %330 %349 +OpBranch %345 +%345 = OpLabel +%350 = OpFOrdGreaterThanEqual %8 %326 %342 +%351 = OpLoad %3 %330 +%352 = OpFOrdLessThanEqual %8 %326 %351 +%353 = OpLogicalAnd %8 %350 %352 +%354 = OpLogicalAnd %8 %353 %341 +OpSelectionMerge %356 None +OpBranchConditional %354 %355 %356 +%355 = OpLabel +OpRayQueryGenerateIntersectionKHR %324 %326 +OpBranch %356 +%356 = OpLabel +OpBranch %332 +%332 = OpLabel OpReturn -%212 = OpLabel -%213 = OpCompositeExtract %6 %207 0 -%214 = OpIEqual %8 %213 %50 -OpSelectionMerge %215 None -OpBranchConditional %214 %216 %217 -%216 = OpLabel -OpRayQueryConfirmIntersectionKHR %164 +OpFunctionEnd +%364 = OpFunction %2 None %363 +%365 = OpFunctionParameter %32 +%366 = OpFunctionParameter %33 +%367 = OpLabel +%370 = OpLoad %6 %366 +%371 = OpBitwiseAnd %6 %370 %93 +%372 = OpINotEqual %8 %371 %35 +%373 = OpBitwiseAnd %6 %370 %27 +%374 = OpINotEqual %8 %373 %35 +%375 = OpLogicalNot %8 %374 +%376 = OpLogicalAnd %8 %375 %372 +OpSelectionMerge %369 None +OpBranchConditional %376 %368 %369 +%368 = OpLabel +%377 = OpRayQueryGetIntersectionTypeKHR %6 %365 %35 +%378 = OpIEqual %8 %377 %35 +OpSelectionMerge %380 None +OpBranchConditional %378 %379 %380 +%379 = OpLabel +OpRayQueryConfirmIntersectionKHR %365 +OpBranch %380 +%380 = OpLabel +OpBranch %369 +%369 = OpLabel OpReturn -%217 = OpLabel +OpFunctionEnd +%262 = OpFunction %2 None %243 +%261 = OpLabel +%266 = OpVariable %32 Function +%267 = OpVariable %33 Function %35 +%268 = OpVariable %36 Function %38 +%263 = OpLoad %5 %15 +OpBranch %269 +%269 = OpLabel +%270 = OpFunctionCall %2 %43 %266 %263 %264 %267 %268 +%316 = OpFunctionCall %10 %271 %266 %267 +%317 = OpCompositeExtract %6 %316 0 +%318 = OpIEqual %8 %317 %199 +OpSelectionMerge %319 None +OpBranchConditional %318 %320 %321 +%320 = OpLabel +%357 = OpFunctionCall %2 %323 %266 %267 %265 %268 OpReturn -%215 = OpLabel -OpBranch %210 -%210 = OpLabel +%321 = OpLabel +%358 = OpCompositeExtract %6 %316 0 +%359 = OpIEqual %8 %358 %90 +OpSelectionMerge %360 None +OpBranchConditional %359 %361 %362 +%361 = OpLabel +%381 = OpFunctionCall %2 %364 %266 %267 +OpReturn +%362 = OpLabel +OpReturn +%360 = OpLabel +OpBranch %319 +%319 = OpLabel OpReturn OpFunctionEnd \ No newline at end of file diff --git a/tests/tests/wgpu-gpu/ray_tracing/shader.rs b/tests/tests/wgpu-gpu/ray_tracing/shader.rs index fcd29af52e6..472ca3891e0 100644 --- a/tests/tests/wgpu-gpu/ray_tracing/shader.rs +++ b/tests/tests/wgpu-gpu/ray_tracing/shader.rs @@ -1,17 +1,19 @@ use crate::ray_tracing::{acceleration_structure_limits, AsBuildContext}; +use wgpu::util::{BufferInitDescriptor, DeviceExt}; use wgpu::{ - include_wgsl, BindGroupDescriptor, BindGroupEntry, BindingResource, BufferDescriptor, - CommandEncoderDescriptor, ComputePassDescriptor, ComputePipelineDescriptor, + include_wgsl, Backends, BindGroupDescriptor, BindGroupEntry, BindingResource, BufferDescriptor, + CommandEncoderDescriptor, ComputePassDescriptor, ComputePipelineDescriptor, InstanceFlags, }; use wgpu::{AccelerationStructureFlags, BufferUsages}; use wgpu_macros::gpu_test; -use wgpu_test::GpuTestInitializer; +use wgpu_test::{FailureCase, GpuTestInitializer}; use wgpu_test::{GpuTestConfiguration, TestParameters, TestingContext}; const STRUCT_SIZE: wgpu::BufferAddress = 176; pub fn all_tests(tests: &mut Vec) { tests.push(ACCESS_ALL_STRUCT_MEMBERS); + tests.push(PREVENT_INVALID_RAY_QUERY_CALLS); } #[gpu_test] @@ -103,3 +105,97 @@ fn access_all_struct_members(ctx: TestingContext) { ctx.queue.submit([encoder_compute.finish()]); } + +#[gpu_test] +static PREVENT_INVALID_RAY_QUERY_CALLS: GpuTestConfiguration = GpuTestConfiguration::new() + .parameters( + TestParameters::default() + .test_features_limits() + .limits(acceleration_structure_limits()) + .features(wgpu::Features::EXPERIMENTAL_RAY_QUERY) + // Otherwise, mistakes in the generated code won't be caught. + .instance_flags(InstanceFlags::GPU_BASED_VALIDATION) + // not yet implemented in directx12 + .skip(FailureCase::backend(Backends::DX12)), + ) + .run_sync(prevent_invalid_ray_query_calls); + +fn prevent_invalid_ray_query_calls(ctx: TestingContext) { + let invalid_values_buffer = ctx.device.create_buffer_init(&BufferInitDescriptor { + label: Some("invalid values buffer"), + contents: bytemuck::cast_slice(&[f32::NAN, f32::INFINITY]), + usage: BufferUsages::STORAGE, + }); + + // + // Create a clean `AsBuildContext` + // + + let as_ctx = AsBuildContext::new( + &ctx, + AccelerationStructureFlags::empty(), + AccelerationStructureFlags::empty(), + ); + + let mut encoder_build = ctx + .device + .create_command_encoder(&CommandEncoderDescriptor { + label: Some("Build"), + }); + + encoder_build.build_acceleration_structures([&as_ctx.blas_build_entry()], [&as_ctx.tlas]); + + ctx.queue.submit([encoder_build.finish()]); + + // + // Create shader + // + + let shader = ctx + .device + .create_shader_module(include_wgsl!("shader.wgsl")); + let compute_pipeline = ctx + .device + .create_compute_pipeline(&ComputePipelineDescriptor { + label: None, + layout: None, + module: &shader, + entry_point: Some("invalid_usages"), + compilation_options: Default::default(), + cache: None, + }); + + let bind_group = ctx.device.create_bind_group(&BindGroupDescriptor { + label: None, + layout: &compute_pipeline.get_bind_group_layout(0), + entries: &[ + BindGroupEntry { + binding: 0, + resource: BindingResource::AccelerationStructure(&as_ctx.tlas), + }, + BindGroupEntry { + binding: 1, + resource: BindingResource::Buffer(invalid_values_buffer.as_entire_buffer_binding()), + }, + ], + }); + + // + // Submit once to check for no issues + // + + let mut encoder_compute = ctx + .device + .create_command_encoder(&CommandEncoderDescriptor::default()); + { + let mut pass = encoder_compute.begin_compute_pass(&ComputePassDescriptor { + label: None, + timestamp_writes: None, + }); + pass.set_pipeline(&compute_pipeline); + pass.set_bind_group(0, Some(&bind_group), &[]); + pass.dispatch_workgroups(1, 1, 1) + } + + ctx.queue.submit([encoder_compute.finish()]); +} diff --git a/tests/tests/wgpu-gpu/ray_tracing/shader.wgsl b/tests/tests/wgpu-gpu/ray_tracing/shader.wgsl index 2130b8d9ae6..55a8f4b85d6 100644 --- a/tests/tests/wgpu-gpu/ray_tracing/shader.wgsl +++ b/tests/tests/wgpu-gpu/ray_tracing/shader.wgsl @@ -48,4 +48,78 @@ fn all_of_struct() { intersection.world_to_object, intersection.object_to_world, ); +} + +struct MaybeInvalidValues { + nan: f32, + inf: f32, +} + +@group(0) @binding(1) +var invalid_values: MaybeInvalidValues; + +@workgroup_size(1) +@compute +fn invalid_usages() { + { + var rq: ray_query; + // no initialize + rayQueryProceed(&rq); + let intersection = rayQueryGetCommittedIntersection(&rq); + } + { + var rq: ray_query; + rayQueryInitialize(&rq, acc_struct, RayDesc(0u, 0xFFu, 0.001, 100000.0, vec3f(0.0, 0.0, 0.0), vec3f(0.0, 0.0, 1.0))); + // no proceed + let intersection = rayQueryGetCommittedIntersection(&rq); + } + { + var rq: ray_query; + rayQueryInitialize(&rq, acc_struct, RayDesc(0u, 0xFFu, 0.001, 100000.0, vec3f(0.0, 0.0, 0.0), vec3f(0.0, 0.0, 1.0))); + rayQueryProceed(&rq); + // The acceleration structure has been set up to not generate an intersections, meaning it will be a committed intersection, not candidate. + let intersection = rayQueryGetCandidateIntersection(&rq); + } + { + var rq: ray_query; + // NaN in origin + rayQueryInitialize(&rq, acc_struct, RayDesc(0u, 0xFFu, 0.001, 100000.0, vec3f(0.0, invalid_values.nan, 0.0), vec3f(0.0, 0.0, 1.0))); + rayQueryProceed(&rq); + let intersection = rayQueryGetCommittedIntersection(&rq); + } + { + var rq: ray_query; + // Inf in origin + rayQueryInitialize(&rq, acc_struct, RayDesc(0u, 0xFFu, 0.001, 100000.0, vec3f(0.0, invalid_values.inf, 0.0), vec3f(0.0, 0.0, 1.0))); + rayQueryProceed(&rq); + let intersection = rayQueryGetCommittedIntersection(&rq); + } + { + var rq: ray_query; + // NaN in direction + rayQueryInitialize(&rq, acc_struct, RayDesc(0u, 0xFFu, 0.001, 100000.0, vec3f(0.0, 0.0, 0.0), vec3f(0.0, invalid_values.nan, 1.0))); + rayQueryProceed(&rq); + let intersection = rayQueryGetCommittedIntersection(&rq); + } + { + var rq: ray_query; + // Inf in direction + rayQueryInitialize(&rq, acc_struct, RayDesc(0u, 0xFFu, 0.001, 100000.0, vec3f(0.0, 0.0, 0.0), vec3f(0.0, invalid_values.inf, 1.0))); + rayQueryProceed(&rq); + let intersection = rayQueryGetCommittedIntersection(&rq); + } + { + var rq: ray_query; + // t_min greater than t_max + rayQueryInitialize(&rq, acc_struct, RayDesc(0u, 0xFFu, 100000.0, 0.1, vec3f(0.0, 0.0, 0.0), vec3f(0.0, 0.0, 1.0))); + rayQueryProceed(&rq); + let intersection = rayQueryGetCommittedIntersection(&rq); + } + { + var rq: ray_query; + // t_min less than 0 + rayQueryInitialize(&rq, acc_struct, RayDesc(0u, 0xFFu, -0.001, 100000.0, vec3f(0.0, 0.0, 0.0), vec3f(0.0, 0.0, 1.0))); + rayQueryProceed(&rq); + let intersection = rayQueryGetCommittedIntersection(&rq); + } } \ No newline at end of file diff --git a/wgpu-hal/src/dx12/device.rs b/wgpu-hal/src/dx12/device.rs index c7e69b63c13..626b390ece1 100644 --- a/wgpu-hal/src/dx12/device.rs +++ b/wgpu-hal/src/dx12/device.rs @@ -290,6 +290,7 @@ impl super::Device { || stage.module.runtime_checks.bounds_checks != layout.naga_options.restrict_indexing || stage.module.runtime_checks.force_loop_bounding != layout.naga_options.force_loop_bounding; + // Note: ray query initialization tracking not yet implemented let mut temp_options; let naga_options = if needs_temp_options { temp_options = layout.naga_options.clone(); diff --git a/wgpu-hal/src/vulkan/adapter.rs b/wgpu-hal/src/vulkan/adapter.rs index 45b78f4dc0e..3b27accda6e 100644 --- a/wgpu-hal/src/vulkan/adapter.rs +++ b/wgpu-hal/src/vulkan/adapter.rs @@ -2191,6 +2191,12 @@ impl super::Adapter { // But this requires cloning the `spv::Options` struct, which has heap allocations. true, // could check `super::Workarounds::SEPARATE_ENTRY_POINTS` ); + flags.set( + spv::WriterFlags::PRINT_ON_RAY_QUERY_INITIALIZATION_FAIL, + self.instance.flags.contains(wgt::InstanceFlags::DEBUG) + && (self.instance.instance_api_version >= vk::API_VERSION_1_3 + || enabled_extensions.contains(&khr::shader_non_semantic_info::NAME)), + ); if features.contains(wgt::Features::EXPERIMENTAL_RAY_QUERY) { capabilities.push(spv::Capability::RayQueryKHR); } @@ -2249,6 +2255,7 @@ impl super::Adapter { spv::ZeroInitializeWorkgroupMemoryMode::Polyfill }, force_loop_bounding: true, + ray_query_initialization_tracking: true, use_storage_input_output_16: features.contains(wgt::Features::SHADER_F16) && self.phd_features.supports_storage_input_output_16(), fake_missing_bindings: false, diff --git a/wgpu-hal/src/vulkan/device.rs b/wgpu-hal/src/vulkan/device.rs index 29d0bbf95a3..bddea14032c 100644 --- a/wgpu-hal/src/vulkan/device.rs +++ b/wgpu-hal/src/vulkan/device.rs @@ -766,6 +766,7 @@ impl super::Device { }; let needs_temp_options = !runtime_checks.bounds_checks || !runtime_checks.force_loop_bounding + || !runtime_checks.ray_query_initialization_tracking || !binding_map.is_empty() || naga_shader.debug_source.is_some() || !stage.zero_initialize_workgroup_memory; @@ -783,6 +784,9 @@ impl super::Device { if !runtime_checks.force_loop_bounding { temp_options.force_loop_bounding = false; } + if !runtime_checks.ray_query_initialization_tracking { + temp_options.ray_query_initialization_tracking = false; + } if !binding_map.is_empty() { temp_options.binding_map = binding_map.clone(); } diff --git a/wgpu-types/src/lib.rs b/wgpu-types/src/lib.rs index ba284ff841d..806b141d698 100644 --- a/wgpu-types/src/lib.rs +++ b/wgpu-types/src/lib.rs @@ -7928,6 +7928,20 @@ pub struct ShaderRuntimeChecks { /// conclusions about other safety-critical code paths. This option SHOULD NOT be disabled /// when running untrusted code. pub force_loop_bounding: bool, + /// If false, the caller **MUST** ensure that in all passed shaders every function operating + /// on a ray query must obey these rules (functions using wgsl naming) + /// - `rayQueryInitialize` must have called before `rayQueryProceed` + /// - `rayQueryProceed` must have been called, returned true and have hit an AABB before + /// `rayQueryGenerateIntersection` is called + /// - `rayQueryProceed` must have been called, returned true and have hit a triangle before + /// `rayQueryConfirmIntersection` is called + /// - `rayQueryProceed` must have been called and have returned true before `rayQueryTerminate`, + /// `getCandidateHitVertexPositions` or `rayQueryGetCandidateIntersection` is called + /// - `rayQueryProceed` must have been called and have returned false before `rayQueryGetCommittedIntersection` + /// or `getCommittedHitVertexPositions` are called + /// + /// It is the aim that these cases will not cause UB if this is set to true, but currently this will still happen on DX12 and Metal. + pub ray_query_initialization_tracking: bool, } impl ShaderRuntimeChecks { @@ -7960,6 +7974,7 @@ impl ShaderRuntimeChecks { Self { bounds_checks: all_checks, force_loop_bounding: all_checks, + ray_query_initialization_tracking: all_checks, } } }