diff --git a/CHANGELOG.md b/CHANGELOG.md index 0f81b26bf..2e7eccd1f 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -40,6 +40,14 @@ Bottom level categories: ## Unreleased +## 23.1.0 (2024-12-11) + +### New Features + +#### Naga + +- Expose Ray Query flags as constants in WGSL. Implement candidate intersections. By @kvark in [#5429](https://github.com/gfx-rs/wgpu/pull/5429) + ## 23.0.1 (2024-11-25) This release includes patches for `wgpu`, `wgpu-core` and `wgpu-hal`. All other crates remain at [23.0.0](https://github.com/gfx-rs/wgpu/releases/tag/v23.0.0). diff --git a/Cargo.lock b/Cargo.lock index 97dcaed7f..7a6d7f640 100644 --- a/Cargo.lock +++ b/Cargo.lock @@ -1875,7 +1875,7 @@ dependencies = [ [[package]] name = "naga" -version = "23.0.0" +version = "23.1.0" dependencies = [ "arbitrary", "arrayvec", diff --git a/examples/src/ray_scene/shader.wgsl b/examples/src/ray_scene/shader.wgsl new file mode 100644 index 000000000..4e16bd945 --- /dev/null +++ b/examples/src/ray_scene/shader.wgsl @@ -0,0 +1,164 @@ +struct VertexOutput { + @builtin(position) position: vec4, + @location(0) tex_coords: vec2, +}; + +@vertex +fn vs_main(@builtin(vertex_index) vertex_index: u32) -> VertexOutput { + var result: VertexOutput; + let x = i32(vertex_index) / 2; + let y = i32(vertex_index) & 1; + let tc = vec2( + f32(x) * 2.0, + f32(y) * 2.0 + ); + result.position = vec4( + tc.x * 2.0 - 1.0, + 1.0 - tc.y * 2.0, + 0.0, 1.0 + ); + result.tex_coords = tc; + return result; +} + +/* +The contents of the RayQuery struct are roughly as follows +let RAY_FLAG_NONE = 0x00u; +let RAY_FLAG_OPAQUE = 0x01u; +let RAY_FLAG_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_index: u32, + instance_id: u32, + sbt_record_offset: u32, + geometry_index: u32, + primitive_index: u32, + barycentrics: vec2, + front_face: bool, + object_to_world: mat4x3, + world_to_object: mat4x3, +} +*/ + +struct Uniforms { + view_inv: mat4x4, + proj_inv: mat4x4, +}; + +struct Vertex { + pos: vec3, + normal: vec3, + uv: vec2, +}; + + +struct Instance { + first_vertex: u32, + first_geometry: u32, + last_geometry: u32, + _pad: u32 +}; + +struct Material{ + roughness_exponent: f32, + metalness: f32, + specularity: f32, + albedo: vec3 +} + +struct Geometry { + first_index: u32, + material: Material, +}; + + +@group(0) @binding(0) +var uniforms: Uniforms; + +@group(0) @binding(1) +var vertices: array; + +@group(0) @binding(2) +var indices: array; + +@group(0) @binding(3) +var geometries: array; + +@group(0) @binding(4) +var instances: array; + +@group(0) @binding(5) +var acc_struct: acceleration_structure; + +@fragment +fn fs_main(vertex: VertexOutput) -> @location(0) vec4 { + + var color = vec4(vertex.tex_coords, 0.0, 1.0); + + let d = vertex.tex_coords * 2.0 - 1.0; + + let origin = (uniforms.view_inv * vec4(0.0,0.0,0.0,1.0)).xyz; + let temp = uniforms.proj_inv * vec4(d.x, d.y, 1.0, 1.0); + let direction = (uniforms.view_inv * vec4(normalize(temp.xyz), 0.0)).xyz; + + var rq: ray_query; + rayQueryInitialize(&rq, acc_struct, RayDesc(0u, 0xFFu, 0.1, 200.0, origin, direction)); + rayQueryProceed(&rq); + + let intersection = rayQueryGetCommittedIntersection(&rq); + if (intersection.kind != RAY_QUERY_INTERSECTION_NONE) { + let instance = instances[intersection.instance_custom_index]; + let geometry = geometries[intersection.geometry_index + instance.first_geometry]; + + let index_offset = geometry.first_index; + let vertex_offset = instance.first_vertex; + + let first_index_index = intersection.primitive_index * 3u + index_offset; + + let v_0 = vertices[vertex_offset+indices[first_index_index+0u]]; + let v_1 = vertices[vertex_offset+indices[first_index_index+1u]]; + let v_2 = vertices[vertex_offset+indices[first_index_index+2u]]; + + let bary = vec3(1.0 - intersection.barycentrics.x - intersection.barycentrics.y, intersection.barycentrics); + + let pos = v_0.pos * bary.x + v_1.pos * bary.y + v_2.pos * bary.z; + let normal_raw = v_0.normal * bary.x + v_1.normal * bary.y + v_2.normal * bary.z; + let uv = v_0.uv * bary.x + v_1.uv * bary.y + v_2.uv * bary.z; + + let normal = normalize(normal_raw); + + let material = geometry.material; + + color = vec4(material.albedo, 1.0); + + if(intersection.instance_custom_index == 1u){ + color = vec4(normal, 1.0); + } + } + + return color; +} diff --git a/naga/Cargo.toml b/naga/Cargo.toml index 16682e526..17be61b68 100644 --- a/naga/Cargo.toml +++ b/naga/Cargo.toml @@ -1,6 +1,6 @@ [package] name = "naga" -version = "23.0.0" +version = "23.1.0" authors = ["gfx-rs developers"] edition = "2021" description = "Shader translation infrastructure" diff --git a/naga/src/back/msl/writer.rs b/naga/src/back/msl/writer.rs index 83d937d48..4cc363b05 100644 --- a/naga/src/back/msl/writer.rs +++ b/naga/src/back/msl/writer.rs @@ -2237,14 +2237,14 @@ impl Writer { write!(self.out, ")")?; } } - crate::Expression::RayQueryGetIntersection { query, committed } => { + crate::Expression::RayQueryGetIntersection { + query, + committed: _, + } => { if context.lang_version < (2, 4) { return Err(Error::UnsupportedRayTracing); } - if !committed { - unimplemented!() - } let ty = context.module.special_types.ray_intersection.unwrap(); let type_name = &self.names[&NameKey::Type(ty)]; write!(self.out, "{type_name} {{{RAY_QUERY_FUN_MAP_INTERSECTION}(")?; diff --git a/naga/src/back/spv/block.rs b/naga/src/back/spv/block.rs index 65c48d9e7..8330537cf 100644 --- a/naga/src/back/spv/block.rs +++ b/naga/src/back/spv/block.rs @@ -1730,10 +1730,7 @@ impl<'w> BlockContext<'w> { } crate::Expression::ArrayLength(expr) => self.write_runtime_array_length(expr, block)?, crate::Expression::RayQueryGetIntersection { query, committed } => { - if !committed { - return Err(Error::FeatureNotImplemented("candidate intersection")); - } - self.write_ray_query_get_intersection(query, block) + self.write_ray_query_get_intersection(query, block, committed) } }; diff --git a/naga/src/back/spv/ray.rs b/naga/src/back/spv/ray.rs index c2daf4b3f..7f16f803e 100644 --- a/naga/src/back/spv/ray.rs +++ b/naga/src/back/spv/ray.rs @@ -106,23 +106,60 @@ impl<'w> BlockContext<'w> { &mut self, query: Handle, block: &mut Block, + is_committed: bool, ) -> spirv::Word { let query_id = self.cached[query]; - let intersection_id = self.writer.get_constant_scalar(crate::Literal::U32( - spirv::RayQueryIntersection::RayQueryCommittedIntersectionKHR as _, - )); + let intersection_id = + self.writer + .get_constant_scalar(crate::Literal::U32(if is_committed { + spirv::RayQueryIntersection::RayQueryCommittedIntersectionKHR + } else { + spirv::RayQueryIntersection::RayQueryCandidateIntersectionKHR + } as _)); let flag_type_id = self.get_type_id(LookupType::Local(LocalType::Numeric( NumericType::Scalar(crate::Scalar::U32), ))); - let kind_id = self.gen_id(); + let raw_kind_id = self.gen_id(); block.body.push(Instruction::ray_query_get_intersection( spirv::Op::RayQueryGetIntersectionTypeKHR, flag_type_id, - kind_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 + } else { + // Remap from the candidate kind to IR + let condition_id = self.gen_id(); + let committed_triangle_kind_id = self.writer.get_constant_scalar(crate::Literal::U32( + spirv::RayQueryCandidateIntersectionType::RayQueryCandidateIntersectionTriangleKHR + as _, + )); + block.body.push(Instruction::binary( + spirv::Op::IEqual, + self.writer.get_bool_type_id(), + condition_id, + raw_kind_id, + committed_triangle_kind_id, + )); + let kind_id = self.gen_id(); + block.body.push(Instruction::select( + flag_type_id, + kind_id, + condition_id, + self.writer.get_constant_scalar(crate::Literal::U32( + crate::RayQueryIntersection::Triangle as _, + )), + self.writer.get_constant_scalar(crate::Literal::U32( + crate::RayQueryIntersection::Aabb as _, + )), + )); + kind_id + }; + let instance_custom_index_id = self.gen_id(); block.body.push(Instruction::ray_query_get_intersection( spirv::Op::RayQueryGetIntersectionInstanceCustomIndexKHR, @@ -201,6 +238,8 @@ impl<'w> BlockContext<'w> { query_id, intersection_id, )); + //Note: there is also `OpRayQueryGetIntersectionCandidateAABBOpaqueKHR`, + // but it's not a property of an intersection. let transform_type_id = self.get_type_id(LookupType::Local(LocalType::Numeric(NumericType::Matrix { diff --git a/naga/src/front/wgsl/lower/mod.rs b/naga/src/front/wgsl/lower/mod.rs index 78e81350b..86bdc170e 100644 --- a/naga/src/front/wgsl/lower/mod.rs +++ b/naga/src/front/wgsl/lower/mod.rs @@ -2563,12 +2563,22 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { args.finish()?; let _ = ctx.module.generate_ray_intersection_type(); - crate::Expression::RayQueryGetIntersection { query, committed: true, } } + "rayQueryGetCandidateIntersection" => { + let mut args = ctx.prepare_args(arguments, 1, span); + let query = self.ray_query_pointer(args.next()?, ctx)?; + args.finish()?; + + let _ = ctx.module.generate_ray_intersection_type(); + crate::Expression::RayQueryGetIntersection { + query, + committed: false, + } + } "RayDesc" => { let ty = ctx.module.generate_ray_desc_type(); let handle = self.construct( diff --git a/naga/src/front/wgsl/parse/mod.rs b/naga/src/front/wgsl/parse/mod.rs index fcfcc3775..f380f406a 100644 --- a/naga/src/front/wgsl/parse/mod.rs +++ b/naga/src/front/wgsl/parse/mod.rs @@ -651,6 +651,14 @@ impl Parser { ctx: &mut ExpressionContext<'a, '_, '_>, ) -> Result>, Error<'a>> { self.push_rule_span(Rule::PrimaryExpr, lexer); + const fn literal_ray_flag<'b>(flag: crate::RayFlag) -> ast::Expression<'b> { + ast::Expression::Literal(ast::Literal::Number(Number::U32(flag.bits()))) + } + const fn literal_ray_intersection<'b>( + intersection: crate::RayQueryIntersection, + ) -> ast::Expression<'b> { + ast::Expression::Literal(ast::Literal::Number(Number::U32(intersection as u32))) + } let expr = match lexer.peek() { (Token::Paren('('), _) => { @@ -683,15 +691,63 @@ impl Parser { } (Token::Word("RAY_FLAG_NONE"), _) => { let _ = lexer.next(); - ast::Expression::Literal(ast::Literal::Number(Number::U32(0))) + literal_ray_flag(crate::RayFlag::empty()) + } + (Token::Word("RAY_FLAG_FORCE_OPAQUE"), _) => { + let _ = lexer.next(); + literal_ray_flag(crate::RayFlag::FORCE_OPAQUE) + } + (Token::Word("RAY_FLAG_FORCE_NO_OPAQUE"), _) => { + let _ = lexer.next(); + literal_ray_flag(crate::RayFlag::FORCE_NO_OPAQUE) } (Token::Word("RAY_FLAG_TERMINATE_ON_FIRST_HIT"), _) => { let _ = lexer.next(); - ast::Expression::Literal(ast::Literal::Number(Number::U32(4))) + literal_ray_flag(crate::RayFlag::TERMINATE_ON_FIRST_HIT) + } + (Token::Word("RAY_FLAG_SKIP_CLOSEST_HIT_SHADER"), _) => { + let _ = lexer.next(); + literal_ray_flag(crate::RayFlag::SKIP_CLOSEST_HIT_SHADER) + } + (Token::Word("RAY_FLAG_CULL_BACK_FACING"), _) => { + let _ = lexer.next(); + literal_ray_flag(crate::RayFlag::CULL_BACK_FACING) + } + (Token::Word("RAY_FLAG_CULL_FRONT_FACING"), _) => { + let _ = lexer.next(); + literal_ray_flag(crate::RayFlag::CULL_FRONT_FACING) + } + (Token::Word("RAY_FLAG_CULL_OPAQUE"), _) => { + let _ = lexer.next(); + literal_ray_flag(crate::RayFlag::CULL_OPAQUE) + } + (Token::Word("RAY_FLAG_CULL_NO_OPAQUE"), _) => { + let _ = lexer.next(); + literal_ray_flag(crate::RayFlag::CULL_NO_OPAQUE) + } + (Token::Word("RAY_FLAG_SKIP_TRIANGLES"), _) => { + let _ = lexer.next(); + literal_ray_flag(crate::RayFlag::SKIP_TRIANGLES) + } + (Token::Word("RAY_FLAG_SKIP_AABBS"), _) => { + let _ = lexer.next(); + literal_ray_flag(crate::RayFlag::SKIP_AABBS) } (Token::Word("RAY_QUERY_INTERSECTION_NONE"), _) => { let _ = lexer.next(); - ast::Expression::Literal(ast::Literal::Number(Number::U32(0))) + literal_ray_intersection(crate::RayQueryIntersection::None) + } + (Token::Word("RAY_QUERY_INTERSECTION_TRIANGLE"), _) => { + let _ = lexer.next(); + literal_ray_intersection(crate::RayQueryIntersection::Triangle) + } + (Token::Word("RAY_QUERY_INTERSECTION_GENERATED"), _) => { + let _ = lexer.next(); + literal_ray_intersection(crate::RayQueryIntersection::Generated) + } + (Token::Word("RAY_QUERY_INTERSECTION_AABB"), _) => { + let _ = lexer.next(); + literal_ray_intersection(crate::RayQueryIntersection::Aabb) } (Token::Word(word), span) => { let start = lexer.start_byte_offset(); diff --git a/naga/src/lib.rs b/naga/src/lib.rs index 145b95f66..3450c0d20 100644 --- a/naga/src/lib.rs +++ b/naga/src/lib.rs @@ -2271,3 +2271,59 @@ pub struct Module { /// Entry points. pub entry_points: Vec, } + +bitflags::bitflags! { + /// Ray flags used when casting rays. + /// Matching vulkan constants can be found in + /// https://github.com/KhronosGroup/SPIRV-Registry/blob/main/extensions/KHR/ray_common/ray_flags_section.txt + #[cfg_attr(feature = "serialize", derive(Serialize))] + #[cfg_attr(feature = "deserialize", derive(Deserialize))] + #[cfg_attr(feature = "arbitrary", derive(Arbitrary))] + #[derive(Clone, Copy, Debug, Default, Eq, Hash, Ord, PartialEq, PartialOrd)] + pub struct RayFlag: u32 { + /// Force all intersections to be treated as opaque. + const FORCE_OPAQUE = 0x1; + /// Force all intersections to be treated as non-opaque. + const FORCE_NO_OPAQUE = 0x2; + /// Stop traversal after the first hit. + const TERMINATE_ON_FIRST_HIT = 0x4; + /// Don't execute the closest hit shader. + const SKIP_CLOSEST_HIT_SHADER = 0x8; + /// Cull back facing geometry. + const CULL_BACK_FACING = 0x10; + /// Cull front facing geometry. + const CULL_FRONT_FACING = 0x20; + /// Cull opaque geometry. + const CULL_OPAQUE = 0x40; + /// Cull non-opaque geometry. + const CULL_NO_OPAQUE = 0x80; + /// Skip triangular geometry. + const SKIP_TRIANGLES = 0x100; + /// Skip axis-aligned bounding boxes. + const SKIP_AABBS = 0x200; + } +} + +/// Type of a ray query intersection. +/// Matching vulkan constants can be found in +/// +/// but the actual values are different for candidate intersections. +#[cfg_attr(feature = "serialize", derive(Serialize))] +#[cfg_attr(feature = "deserialize", derive(Deserialize))] +#[cfg_attr(feature = "arbitrary", derive(Arbitrary))] +#[derive(Clone, Copy, Debug, Default, Eq, Hash, Ord, PartialEq, PartialOrd)] +pub enum RayQueryIntersection { + /// No intersection found. + /// Matches `RayQueryCommittedIntersectionNoneKHR`. + #[default] + None = 0, + /// Intersecting with triangles. + /// Matches `RayQueryCommittedIntersectionTriangleKHR` and `RayQueryCandidateIntersectionTriangleKHR`. + Triangle = 1, + /// Intersecting with generated primitives. + /// Matches `RayQueryCommittedIntersectionGeneratedKHR`. + Generated = 2, + /// Intersecting with Axis Aligned Bounding Boxes. + /// Matches `RayQueryCandidateIntersectionAABBKHR`. + Aabb = 3, +} diff --git a/naga/tests/in/ray-query.wgsl b/naga/tests/in/ray-query.wgsl index 0af8c7c95..9f94356b8 100644 --- a/naga/tests/in/ray-query.wgsl +++ b/naga/tests/in/ray-query.wgsl @@ -1,7 +1,7 @@ /* let RAY_FLAG_NONE = 0x00u; -let RAY_FLAG_OPAQUE = 0x01u; -let RAY_FLAG_NO_OPAQUE = 0x02u; +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; @@ -14,7 +14,7 @@ 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 = 4u; +let RAY_QUERY_INTERSECTION_AABB = 3u; struct RayDesc { flags: u32, @@ -78,3 +78,14 @@ fn main() { 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); + output.visible = u32(intersection.kind == RAY_QUERY_INTERSECTION_AABB); +} diff --git a/naga/tests/out/msl/ray-query.msl b/naga/tests/out/msl/ray-query.msl index 129ad108a..b8d5bb340 100644 --- a/naga/tests/out/msl/ray-query.msl +++ b/naga/tests/out/msl/ray-query.msl @@ -46,23 +46,23 @@ RayIntersection query_loop( metal::float3 dir, metal::raytracing::instance_acceleration_structure acs ) { - _RayQuery rq = {}; + _RayQuery rq_1 = {}; RayDesc _e8 = RayDesc {4u, 255u, 0.1, 100.0, pos, dir}; - rq.intersector.assume_geometry_type(metal::raytracing::geometry_type::triangle); - rq.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.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.intersector.accept_any_intersection((_e8.flags & 4) != 0); - rq.intersection = rq.intersector.intersect(metal::raytracing::ray(_e8.origin, _e8.dir, _e8.tmin, _e8.tmax), acs, _e8.cull_mask); rq.ready = true; + 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; #define LOOP_IS_REACHABLE if (volatile bool unpredictable_jump_over_loop = true; unpredictable_jump_over_loop) LOOP_IS_REACHABLE while(true) { - bool _e9 = rq.ready; - rq.ready = false; + bool _e9 = rq_1.ready; + rq_1.ready = false; if (_e9) { } else { break; } } - return 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}; + 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( @@ -87,3 +87,22 @@ kernel void main_( output.normal = _e18; return; } + + +kernel void main_candidate( + metal::raytracing::instance_acceleration_structure acc_struct [[user(fake0)]] +, device Output& output [[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}; + output.visible = static_cast(intersection_1.kind == 3u); + return; +} diff --git a/naga/tests/out/spv/ray-query.spvasm b/naga/tests/out/spv/ray-query.spvasm index 8b784f2fa..5279bfc2e 100644 --- a/naga/tests/out/spv/ray-query.spvasm +++ b/naga/tests/out/spv/ray-query.spvasm @@ -1,14 +1,16 @@ ; SPIR-V ; Version: 1.4 ; Generator: rspirv -; Bound: 104 +; Bound: 136 OpCapability Shader OpCapability RayQueryKHR OpExtension "SPV_KHR_ray_query" %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 OpEntryPoint GLCompute %84 "main" %15 %17 +OpEntryPoint GLCompute %105 "main_candidate" %15 %17 OpExecutionMode %84 LocalSize 1 1 1 +OpExecutionMode %105 LocalSize 1 1 1 OpMemberDecorate %10 0 Offset 0 OpMemberDecorate %10 1 Offset 4 OpMemberDecorate %10 2 Offset 8 @@ -74,6 +76,8 @@ OpMemberDecorate %18 0 Offset 0 %91 = OpConstantComposite %4 %70 %68 %70 %94 = OpTypePointer StorageBuffer %6 %99 = OpTypePointer StorageBuffer %4 +%108 = OpConstantComposite %12 %27 %28 %29 %30 %90 %91 +%109 = OpConstant %6 3 %25 = OpFunction %10 None %26 %21 = OpFunctionParameter %4 %22 = OpFunctionParameter %4 @@ -161,4 +165,39 @@ OpStore %98 %97 %103 = OpAccessChain %99 %89 %50 OpStore %103 %102 OpReturn +OpFunctionEnd +%105 = OpFunction %2 None %85 +%104 = OpLabel +%110 = OpVariable %32 Function +%106 = OpLoad %5 %15 +%107 = OpAccessChain %87 %17 %88 +OpBranch %111 +%111 = OpLabel +%112 = OpCompositeExtract %6 %108 0 +%113 = OpCompositeExtract %6 %108 1 +%114 = OpCompositeExtract %3 %108 2 +%115 = OpCompositeExtract %3 %108 3 +%116 = OpCompositeExtract %4 %108 4 +%117 = OpCompositeExtract %4 %108 5 +OpRayQueryInitializeKHR %110 %106 %112 %113 %116 %114 %117 %115 +%118 = OpRayQueryGetIntersectionTypeKHR %6 %110 %88 +%119 = OpIEqual %8 %118 %88 +%120 = OpSelect %6 %119 %50 %109 +%121 = OpRayQueryGetIntersectionInstanceCustomIndexKHR %6 %110 %88 +%122 = OpRayQueryGetIntersectionInstanceIdKHR %6 %110 %88 +%123 = OpRayQueryGetIntersectionInstanceShaderBindingTableRecordOffsetKHR %6 %110 %88 +%124 = OpRayQueryGetIntersectionGeometryIndexKHR %6 %110 %88 +%125 = OpRayQueryGetIntersectionPrimitiveIndexKHR %6 %110 %88 +%126 = OpRayQueryGetIntersectionTKHR %3 %110 %88 +%127 = OpRayQueryGetIntersectionBarycentricsKHR %7 %110 %88 +%128 = OpRayQueryGetIntersectionFrontFaceKHR %8 %110 %88 +%129 = OpRayQueryGetIntersectionObjectToWorldKHR %9 %110 %88 +%130 = OpRayQueryGetIntersectionWorldToObjectKHR %9 %110 %88 +%131 = OpCompositeConstruct %10 %120 %126 %121 %122 %123 %124 %125 %127 %128 %129 %130 +%132 = OpCompositeExtract %6 %131 0 +%133 = OpIEqual %8 %132 %109 +%134 = OpSelect %6 %133 %50 %88 +%135 = OpAccessChain %94 %107 %88 +OpStore %135 %134 +OpReturn OpFunctionEnd \ No newline at end of file