Skip to content

Commit f6d9a1d

Browse files
VecvecVecvec
authored
[msl-out] Switch to intersection queries (#9304)
1 parent cc44c1f commit f6d9a1d

8 files changed

Lines changed: 538 additions & 308 deletions

File tree

‎CHANGELOG.md‎

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -117,6 +117,10 @@ By @teoxoy in [#9351](https://github.com/gfx-rs/wgpu/pull/9351).
117117

118118
- Prefix `FeatureLevel` and `ShaderModel` enum variants with `V` instead of `_`. By @teoxoy in [#9337](https://github.com/gfx-rs/wgpu/pull/9337).
119119

120+
#### naga
121+
122+
- Switched from using an `intersector` to using an `intersection_query` on metal so AABBs and non-opaque triangles can be handled. By @Vecvec in [#9304](https://github.com/gfx-rs/wgpu/pull/9304).
123+
120124
### Bug Fixes
121125

122126
#### General

‎naga/src/back/msl/mod.rs‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -79,6 +79,7 @@ use crate::{arena::Handle, back::TaskDispatchLimits, ir, proc::index, valid::Mod
7979

8080
mod keywords;
8181
mod mesh_shader;
82+
mod ray;
8283
pub mod sampler;
8384
mod writer;
8485

‎naga/src/back/msl/ray.rs‎

Lines changed: 269 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,269 @@
1+
use alloc::{
2+
format,
3+
string::{String, ToString},
4+
};
5+
use core::fmt::Write;
6+
7+
use crate::{
8+
back::{
9+
self,
10+
msl::{
11+
writer::{StatementContext, TypeContext, WrappedFunction},
12+
BackendResult, Error, Writer,
13+
},
14+
Baked,
15+
},
16+
Handle,
17+
};
18+
19+
pub(super) const RT_NAMESPACE: &str = "metal::raytracing";
20+
21+
/// The ray query type, needs to be a function so it can format the constants.
22+
pub(super) fn metal_intersector_ty() -> String {
23+
format!("{RT_NAMESPACE}::intersection_query<{RT_NAMESPACE}::instancing, {RT_NAMESPACE}::triangle_data>")
24+
}
25+
26+
pub(super) const INTERSECTION_FUNCTION_NAME: &str = "ray_query_get_intersection";
27+
28+
impl<W: Write> Writer<W> {
29+
/// Writes a function to get the current intersection from the ray query
30+
///
31+
/// Like other backends, this is needed to have a single branch for constructing
32+
/// the parts of the intersection that need to be checked whether they do or don't
33+
/// hit.
34+
pub(super) fn write_rq_get_intersection_function(
35+
&mut self,
36+
module: &crate::Module,
37+
committed: bool,
38+
) -> BackendResult {
39+
let wrapped = WrappedFunction::RayQueryGetIntersection { committed };
40+
if !self.wrapped_functions.insert(wrapped) {
41+
return Ok(());
42+
}
43+
44+
let ty = if committed { "committed" } else { "candidate" };
45+
let intersection = TypeContext {
46+
handle: module
47+
.special_types
48+
.ray_intersection
49+
.expect("intersection ty should be there for intersection function"),
50+
gctx: module.to_ctx(),
51+
names: &self.names,
52+
access: crate::StorageAccess::empty(),
53+
first_time: false,
54+
};
55+
let level = back::Level(1);
56+
writeln!(
57+
self.out,
58+
"{intersection} {INTERSECTION_FUNCTION_NAME}_{committed}({} intersector) {{",
59+
metal_intersector_ty()
60+
)?;
61+
// Initialize the intersection to its default values (which should be zero).
62+
writeln!(
63+
self.out,
64+
"{level}{intersection} intersection = {intersection} {{}};"
65+
)?;
66+
writeln!(self.out, "{level}{RT_NAMESPACE}::intersection_type ty = intersector.get_{ty}_intersection_type();")?;
67+
// If the ray hit a triangle, call all methods that require that and set the intersection type.
68+
writeln!(
69+
self.out,
70+
"{level}if (ty == {RT_NAMESPACE}::intersection_type::triangle) {{"
71+
)?;
72+
writeln!(
73+
self.out,
74+
"{level}{level}intersection.kind = {};",
75+
crate::RayQueryIntersection::Triangle as u32
76+
)?;
77+
if !committed {
78+
writeln!(
79+
self.out,
80+
"{level}{level}intersection.t = intersector.get_candidate_triangle_distance();"
81+
)?;
82+
}
83+
writeln!(self.out, "{level}{level}intersection.barycentrics = intersector.get_{ty}_triangle_barycentric_coord();")?;
84+
writeln!(
85+
self.out,
86+
"{level}{level}intersection.front_face = intersector.is_{ty}_triangle_front_facing();"
87+
)?;
88+
// Otherwise, if the ray hit an AABB (called a bounding box in metal) set the intersection type
89+
// (which depends on whether this is a committed or candidate intersection).
90+
writeln!(
91+
self.out,
92+
"{level}}} else if (ty == {RT_NAMESPACE}::intersection_type::bounding_box) {{"
93+
)?;
94+
if committed {
95+
writeln!(
96+
self.out,
97+
"{level}{level}intersection.kind = {};",
98+
crate::RayQueryIntersection::Generated as u32
99+
)?;
100+
} else {
101+
writeln!(
102+
self.out,
103+
"{level}{level}intersection.kind = {};",
104+
crate::RayQueryIntersection::Aabb as u32
105+
)?;
106+
}
107+
writeln!(self.out, "{level}}}")?;
108+
109+
// If the ray hit anything at all, call all methods that require that.
110+
writeln!(
111+
self.out,
112+
"{level}if (ty != {RT_NAMESPACE}::intersection_type::none) {{"
113+
)?;
114+
if committed {
115+
writeln!(
116+
self.out,
117+
"{level}{level}intersection.t = intersector.get_committed_distance();"
118+
)?;
119+
}
120+
writeln!(self.out, "{level}{level}intersection.instance_custom_data = intersector.get_{ty}_user_instance_id();")?;
121+
writeln!(
122+
self.out,
123+
"{level}{level}intersection.instance_index = intersector.get_{ty}_instance_id();"
124+
)?;
125+
// Metal does not appear to support obtaining the intersection offset from a ray query.
126+
//writeln!(self.out, "{level}{level}intersection.sbt_record_offset = intersector.get_{ty}_user_instance_id();")?;
127+
writeln!(
128+
self.out,
129+
"{level}{level}intersection.geometry_index = intersector.get_{ty}_geometry_id();"
130+
)?;
131+
writeln!(
132+
self.out,
133+
"{level}{level}intersection.primitive_index = intersector.get_{ty}_primitive_id();"
134+
)?;
135+
writeln!(self.out, "{level}{level}intersection.object_to_world = intersector.get_{ty}_object_to_world_transform();")?;
136+
writeln!(self.out, "{level}{level}intersection.world_to_object = intersector.get_{ty}_world_to_object_transform();")?;
137+
writeln!(self.out, "{level}}}")?;
138+
writeln!(self.out, "{level}return intersection;")?;
139+
writeln!(self.out, "}}")?;
140+
141+
Ok(())
142+
}
143+
144+
pub(super) fn write_ray_query_stmt(
145+
&mut self,
146+
level: back::Level,
147+
context: &StatementContext,
148+
query: Handle<crate::Expression>,
149+
fun: &crate::RayQueryFunction,
150+
) -> BackendResult {
151+
if context.expression.lang_version < (2, 4) {
152+
return Err(Error::UnsupportedRayTracing);
153+
}
154+
155+
// TODO: check for misuse.
156+
match *fun {
157+
crate::RayQueryFunction::Initialize {
158+
acceleration_structure,
159+
descriptor,
160+
} => {
161+
//TODO: how to deal with winding? Is it by default the same as the other APIs?
162+
163+
// Put everything in a block so that the variable names
164+
// do not conflict with user variable names
165+
writeln!(self.out, "{level}{{")?;
166+
167+
let inner_level = level.next();
168+
169+
let naga_ray_desc_ty = TypeContext {
170+
handle: context
171+
.expression
172+
.module
173+
.special_types
174+
.ray_desc
175+
.expect("ray desc is required as an argument so should be there"),
176+
gctx: context.expression.module.to_ctx(),
177+
names: &self.names,
178+
access: crate::StorageAccess::empty(),
179+
first_time: false,
180+
};
181+
182+
write!(self.out, "{inner_level}{naga_ray_desc_ty} desc = ")?;
183+
self.put_expression(descriptor, &context.expression, false)?;
184+
writeln!(self.out, ";")?;
185+
186+
// Set up intersection parameters
187+
writeln!(
188+
self.out,
189+
"{inner_level}{RT_NAMESPACE}::intersection_params params;"
190+
)?;
191+
192+
{
193+
// Determine whether or not to cull opaque/non-opaques
194+
let f_opaque = back::RayFlag::CULL_OPAQUE.bits();
195+
let f_no_opaque = back::RayFlag::CULL_NO_OPAQUE.bits();
196+
writeln!(
197+
self.out,
198+
"{inner_level}params.set_opacity_cull_mode(
199+
{inner_level} (desc.flags & {f_opaque}) != 0 ? {RT_NAMESPACE}::opacity_cull_mode::opaque : (
200+
{inner_level} (desc.flags & {f_no_opaque}) != 0 ? {RT_NAMESPACE}::opacity_cull_mode::non_opaque : {RT_NAMESPACE}::opacity_cull_mode::none
201+
{inner_level} )
202+
{inner_level});"
203+
)?;
204+
}
205+
{
206+
// Determine whether to force a particular opacity
207+
let f_opaque = back::RayFlag::OPAQUE.bits();
208+
let f_no_opaque = back::RayFlag::NO_OPAQUE.bits();
209+
writeln!(self.out, "{inner_level}params.force_opacity(
210+
{inner_level} (desc.flags & {f_opaque}) != 0 ? {RT_NAMESPACE}::forced_opacity::opaque : (
211+
{inner_level} (desc.flags & {f_no_opaque}) != 0 ? {RT_NAMESPACE}::forced_opacity::non_opaque : {RT_NAMESPACE}::forced_opacity::none
212+
{inner_level} )
213+
{inner_level});")?;
214+
}
215+
{
216+
let flag = back::RayFlag::TERMINATE_ON_FIRST_HIT.bits();
217+
writeln!(
218+
self.out,
219+
"{inner_level}params.accept_any_intersection((desc.flags & {flag}) != 0);"
220+
)?;
221+
}
222+
223+
writeln!(
224+
self.out,
225+
"{inner_level}{RT_NAMESPACE}::ray ray = {RT_NAMESPACE}::ray(desc.origin, desc.dir, desc.tmin, desc.tmax);"
226+
)?;
227+
228+
write!(self.out, "{inner_level}")?;
229+
// A ray query can by initialized in metal by either using a "non-default constructor"
230+
// or by calling reset. Ray queries cannot be assigned to in metal, so reset needs to
231+
// be called.
232+
self.put_expression(query, &context.expression, true)?;
233+
write!(self.out, ".reset(ray,")?;
234+
self.put_expression(acceleration_structure, &context.expression, true)?;
235+
writeln!(self.out, ", desc.cull_mask, params);")?;
236+
writeln!(self.out, "{level}}}")?;
237+
}
238+
crate::RayQueryFunction::Proceed { result } => {
239+
write!(self.out, "{level}")?;
240+
let name = Baked(result).to_string();
241+
self.start_baking_expression(result, &context.expression, &name)?;
242+
self.named_expressions.insert(result, name);
243+
self.put_expression(query, &context.expression, true)?;
244+
writeln!(self.out, ".next();")?;
245+
}
246+
crate::RayQueryFunction::GenerateIntersection { hit_t } => {
247+
write!(self.out, "{level}")?;
248+
self.put_expression(query, &context.expression, true)?;
249+
write!(self.out, ".commit_bounding_box_intersection(")?;
250+
self.put_expression(hit_t, &context.expression, true)?;
251+
writeln!(self.out, ");")?;
252+
}
253+
crate::RayQueryFunction::ConfirmIntersection => {
254+
write!(self.out, "{level}")?;
255+
self.put_expression(query, &context.expression, true)?;
256+
writeln!(self.out, ".commit_triangle_intersection();")?;
257+
}
258+
crate::RayQueryFunction::Terminate => {
259+
write!(self.out, "{level}")?;
260+
self.put_expression(query, &context.expression, true)?;
261+
// Terminate appears to map to abort in spirv-cross, but metal only documents
262+
// the existence of this method, not what it does.
263+
writeln!(self.out, ".abort();")?;
264+
}
265+
}
266+
267+
Ok(())
268+
}
269+
}

0 commit comments

Comments
 (0)