1#include random.wgsl
2#include brdf.wgsl
3#include sky.wgsl
4#include any_hit.wgsl
5#include utils.wgsl
6
7@group(0) @binding(0) var output_buffer : texture_storage_2d<rgba32float, write>;
8@group(0) @binding(1) var<storage, read> objects: Objects;
9@group(0) @binding(2) var<uniform> uniforms: UniformLayout;
10@group(0) @binding(5) var<storage, read> node_tree: BVH;
11@group(0) @binding(6) var<storage, read> tri_lut: ObjectIndices;
12@group(0) @binding(7) var<storage, read_write> input_buffer:array<vec3f>;
13
14@group(1) @binding(0) var<storage, read> materials:array<Material>;
15@group(1) @binding(1) var textures: texture_2d_array<f32>;
16@group(1) @binding(2) var t_sampler: sampler;
17@group(1) @binding(4) var<uniform> textureSizes: array<vec4<f32>, 128>;
18@group(1) @binding(6) var blueNoiseTexture : texture_storage_2d<rgba8unorm, read>;
19@group(1) @binding(7) var<storage, read> emissiveTriangleIndices : array<f32>;
20
21// @group(0) @binding(3) var skybox: texture_2d<f32>;
22// @group(0) @binding(4) var skybox_sampler: sampler;
23// @group(1) @binding(5) var skyboxCDF: texture_storage_2d<rg32float, read>;
24// @group(1) @binding(3) var<storage, read> areaLights:array<AreaLight>;
25
26struct Triangle {
27 corner_a: vec3<f32>,
28 corner_b: vec3<f32>,
29 corner_c: vec3<f32>,
30 normal_a: vec3<f32>,
31 normal_b: vec3<f32>,
32 normal_c: vec3<f32>,
33 material_idx: f32,
34 uv_a: vec2<f32>,
35 uv_b: vec2<f32>,
36 uv_c: vec2<f32>,
37 tangent_a: vec4f,
38 tangent_b: vec4f,
39 tangent_c: vec4f,
40}
41
42// struct AreaLight {
43// center: vec3<f32>,
44// u: vec3<f32>,
45// v: vec3<f32>,
46// normal: vec3<f32>,
47// emission: vec3<f32>,
48// };
49
50struct Ray {
51 direction: vec3<f32>,
52 origin: vec3<f32>,
53}
54
55struct HitInfo {
56 dist: f32,
57 hit: bool,
58 position: vec3<f32>,
59 normal: vec3<f32>,
60 material_idx: i32,
61 geo_normal: vec3f,
62 tri: Triangle,
63 uv: vec2f,
64 tangent: vec3<f32>,
65 bitangent: vec3<f32>,
66}
67
68struct UniformLayout {
69 position: vec3<f32>,
70 frame_idx: f32,
71 view: mat4x4<f32>,
72 inverse_view: mat4x4<f32>,
73 projection: mat4x4<f32>,
74 sun_direction: vec3<f32>,
75 sun_angular_size: f32,
76 sun_radiance: vec3<f32>,
77 sample_count: f32,
78 max_depth: f32,
79 aperture: f32,
80 focus_distance: f32,
81 emissive_triangle_count: f32,
82 thin_lens: f32,
83}
84
85struct Node {
86 min_corner: vec3<f32>,
87 left_child: f32,
88 max_corner: vec3<f32>,
89 primitive_count: f32,
90}
91
92struct BVH {
93 nodes: array<Node>,
94}
95
96struct ObjectIndices {
97 primitive_indices: array<f32>,
98}
99
100struct Objects {
101 triangles: array<Triangle>,
102}
103
104struct Material {
105 albedo: vec4<f32>,
106 metallic: f32,
107 alpha_mode: f32,
108 alpha_cutoff: f32,
109 double_sided: f32,
110 emission: vec3<f32>,
111 roughness: f32,
112 base_color_texture: f32,
113 normal_texture: f32,
114 metallic_roughness_texture: f32,
115 emissive_texture: f32,
116}
117
118
119
120const EPSILON :f32 = 0.00001f;
121const PI :f32 = 3.1415927f;
122// ray tracing gems part 1 chapter 6
123const FLOAT_SCALE = 1.0 / 65536.0;
124const INT_SCALE = 256.0;
125const ORIGIN = 1.0 / 32.0;
126
127// Slightly offsets a ray to prevent self intersection artifacts
128// Ray tracing gems part 1 chapter 6
129fn offset_ray(p: vec3<f32>, n: vec3<f32>) -> vec3<f32> {
130 let of_i = vec3<i32>(
131 i32(INT_SCALE * n.x),
132 i32(INT_SCALE * n.y),
133 i32(INT_SCALE * n.z)
134 );
135
136 let p_i = vec3<f32>(
137 int_to_float(float_to_int(p.x) + select(of_i.x, -of_i.x, p.x < 0.0)),
138 int_to_float(float_to_int(p.y) + select(of_i.y, -of_i.y, p.y < 0.0)),
139 int_to_float(float_to_int(p.z) + select(of_i.z, -of_i.z, p.z < 0.0))
140 );
141
142 return vec3<f32>(
143 select(p.x + FLOAT_SCALE * n.x, p_i.x, abs(p.x) >= ORIGIN),
144 select(p.y + FLOAT_SCALE * n.y, p_i.y, abs(p.y) >= ORIGIN),
145 select(p.z + FLOAT_SCALE * n.z, p_i.z, abs(p.z) >= ORIGIN)
146 );
147}
148
149fn sample_material_texture(uv: vec2<f32>, texture_index: u32) -> vec4<f32> {
150 let tex_size = textureSizes[texture_index].xy;
151 let max_tex_size = vec2<f32>(textureDimensions(textures).xy);
152 // let scaled_uv = uv * tex_size / max_tex_size;
153 // let clamped_uv = clamp(scaled_uv, vec2<f32>(0.0), vec2<f32>(1.0));
154 // compute the valid uv bounds inside the texture array
155 let tex_uv_min = vec2<f32>(0.0); // always starts at (0,0)
156 let tex_uv_max = tex_size / max_tex_size; // upper-right boundary in the atlas
157 // remap u_vs to this valid range
158 let mapped_uv = mix(tex_uv_min, tex_uv_max, uv);
159 return textureSampleLevel(textures, t_sampler, mapped_uv, texture_index, 1.0).rgba;
160}
161
162
163fn parse_textures(curr_material: Material, result: HitInfo) -> Material {
164 var material = curr_material;
165 if material.base_color_texture > -1.0 {
166 material.albedo *= sample_material_texture(result.uv, u32(curr_material.base_color_texture)).rgba;
167 }
168 if material.metallic_roughness_texture > -1.0 {
169 let metallic_roughness_texture = sample_material_texture(result.uv, u32(curr_material.metallic_roughness_texture));
170 material.roughness *= metallic_roughness_texture.g;
171 material.metallic *= metallic_roughness_texture.b;
172 }
173 if material.emissive_texture > -1.0 {
174 material.emission = sample_material_texture(result.uv, u32(curr_material.emissive_texture)).rgb;
175 }
176 return material;
177}
178
179
180fn point_in_unit_disk(u: vec2f) -> vec2f {
181 let r = sqrt(u.x);
182 let theta = 2f * PI * u.y;
183 return vec2f(r * cos(theta), r * sin(theta));
184}
185
186fn generate_pinhole_camera_ray(ndc: vec2<f32>, noise: vec2f) -> Ray {
187 var ray : Ray;
188 let aspect = uniforms.projection[1][1] / uniforms.projection[0][0]; // same as 1/tan_half_fov_y divided by 1/tan_half_fov_x
189 let tan_half_fov_y = 1.0 / uniforms.projection[1][1];
190
191 let x = ndc.x * aspect * tan_half_fov_y;
192 let y = ndc.y * tan_half_fov_y;
193
194 // camera basis vectors from the view matrix
195 let right = uniforms.inverse_view[0].xyz;
196 let up = uniforms.inverse_view[1].xyz;
197 let forward = -uniforms.inverse_view[2].xyz;
198 let origin = uniforms.position;
199
200 let pinhole_dir = normalize(x * right + y * up + forward);
201
202 let focus_dist = uniforms.focus_distance;
203 let aperture = uniforms.aperture;
204 let focus_point = origin + pinhole_dir * focus_dist;
205
206 // sample lens (in local right-up plane)
207 let lens_sample = point_in_unit_disk(noise) * aperture;
208 let lens_offset = lens_sample.x * right + lens_sample.y * up;
209
210 if (uniforms.thin_lens == 0.0){
211 ray.origin = origin;
212 ray.direction = pinhole_dir;
213 } else {
214 ray.origin = origin + lens_offset;
215 ray.direction = normalize(focus_point - ray.origin);
216 }
217 return ray;
218}
219
220
221@compute @workgroup_size(16, 16)
222fn main(
223 @builtin(global_invocation_id) GlobalInvocationID: vec3<u32>,
224 @builtin(local_invocation_id) LocalInvocationID: vec3<u32>,
225 @builtin(workgroup_id) GroupIndex: vec3<u32>) {
226 // https://www.w3.org/TR/webgpu/#coordinate-systems
227 let output_dimension: vec2<i32> = vec2<i32>(textureDimensions(output_buffer));
228 let pixel_position: vec2<i32> = vec2<i32>(i32(GlobalInvocationID.x), i32(GlobalInvocationID.y));
229 let pixel_idx: i32 = pixel_position.y * output_dimension.x + pixel_position.x;
230
231 let pixel_center: vec2<f32> = vec2<f32>(pixel_position) + vec2f(0.5);
232 let uv: vec2<f32> = pixel_center / vec2f(output_dimension);
233 let ndc: vec2<f32> = uv * 2.0 - vec2f(1.0);
234
235 let noise = animated_blue_noise(pixel_position, u32(uniforms.frame_idx), u32(64));
236 var rnd_state = u32(0);
237 init_random(&rnd_state, u32(uniforms.frame_idx));
238 init_random(&rnd_state, u32(pixel_position.x));
239 init_random(&rnd_state, u32(pixel_position.y));
240
241 let jitter_scale: f32 = 1;
242 // Apply blue noise instead of uniformFloat
243 let jitter_x: f32 = (noise.x - 0.5) / f32(output_dimension.x) * jitter_scale;
244 let jitter_y: f32 = (noise.y - 0.5) / f32(output_dimension.y) * jitter_scale;
245
246 let n2 = (ndc.x + jitter_x);
247 let n3 = ndc.y + jitter_y;
248 let ray = generate_pinhole_camera_ray(vec2f(n2, n3), noise);
249
250 var accumulated_color: vec3<f32> = vec3<f32>(0.0);
251 let frame_weight: f32 = 1.0 / (uniforms.frame_idx + 1);
252 let samples_per_pixel: i32 = i32(uniforms.sample_count);
253 for (var i: i32 = 0; i < samples_per_pixel; i ++) {
254 var pixel_color: vec3<f32> = shade_hit(ray, rnd_state, noise);
255 var r = pixel_color.x;
256 var g = pixel_color.y;
257 var b = pixel_color.z;
258 // lazy NaN catching
259 if (r != r){ pixel_color.r = 0.0;};
260 if (g != g){ pixel_color.g = 0.0;};
261 if (b != b){ pixel_color.b = 0.0;};
262 accumulated_color += pixel_color;
263 }
264
265 accumulated_color = accumulated_color / f32(samples_per_pixel);
266 var prev_color: vec3<f32> = input_buffer[pixel_idx];
267 var final_output : vec3f = (prev_color * uniforms.frame_idx + accumulated_color) / (uniforms.frame_idx + 1.0);
268 input_buffer[pixel_idx] = final_output;
269 textureStore(output_buffer, pixel_position, vec4f(final_output, 1.0));
270}
271
272fn trace(ray: Ray) -> HitInfo {
273 var render_state: HitInfo;
274 render_state.hit = false;
275 var nearest_hit: f32 = 999.0;
276
277 // set up for bvh traversal
278 var node: Node = node_tree.nodes[0];
279 var stack: array<Node, 32>;
280 var stack_location: i32 = 0;
281
282 while true {
283 var primitive_count: u32 = u32(node.primitive_count);
284 var contents: u32 = u32(node.left_child);
285
286 if primitive_count == 0 {
287 var child1: Node = node_tree.nodes[contents];
288 var child2: Node = node_tree.nodes[contents + 1];
289
290 var distance1: f32 = hit_aabb(ray, child1);
291 var distance2: f32 = hit_aabb(ray, child2);
292
293 if distance1 > distance2 {
294 var temp_dist: f32 = distance1;
295 distance1 = distance2;
296 distance2 = temp_dist;
297
298 var temp_child: Node = child1;
299 child1 = child2;
300 child2 = temp_child;
301 }
302
303 if distance1 > nearest_hit {
304 if stack_location == 0 {
305 break;
306 } else {
307 stack_location -= 1;
308 node = stack[stack_location];
309 }
310 } else {
311 node = child1;
312 if distance1 < nearest_hit {
313 stack[stack_location] = child2;
314 stack_location += 1;
315 }
316 }
317 } else {
318 for (var i: u32 = 0; i < primitive_count; i++) {
319 var new_render_state: HitInfo = hit_triangle(
320 ray,
321 objects.triangles[u32(tri_lut.primitive_indices[i + contents])],
322 0.001,
323 nearest_hit,
324 render_state,
325 );
326 if new_render_state.hit {
327 nearest_hit = new_render_state.dist;
328 render_state = new_render_state;
329 }
330 }
331 if stack_location == 0 {
332 break;
333 } else {
334 stack_location -= 1;
335 node = stack[stack_location];
336 }
337 }
338 }
339 return render_state;
340}
341
342fn shade_hit(ray: Ray, seed: u32, noise: vec2f) -> vec3<f32> {
343 var current_seed = seed;
344 var radiance = vec3f(0.0);
345 var throughput = vec3f(1.0);
346 var result: HitInfo;
347
348 var temp_ray = ray;
349 let bounces: u32 = u32(uniforms.max_depth);
350
351 var pdf: f32;
352 var env_pdf: f32;
353 var mis_weight : f32 = 1.0;
354
355 var sun_solid_angle = 2.0 * PI * (1.0 - cos(uniforms.sun_angular_size));
356 let sun_pdf = 1.0 / sun_solid_angle;
357 let sky_pdf = 1.0 / PI;
358
359 for (var bounce: u32 = 0; bounce < bounces; bounce++) {
360 result = trace(temp_ray);
361 if (!result.hit) {
362 // We hit the environment; skip the sun for now. Atleast till this rudimentry temporal accmulation exists.
363 // let to_sun = dot(temp_ray.direction, uniforms.sun_direction) > cos(uniforms.sun_angular_size);
364 // let sun_radiance = sun_glow(temp_ray.direction, uniforms.sun_direction);
365 // if (to_sun) {
366 // radianceOut += sun_radiance;
367 // }
368 // if (to_sun) {
369 // env_pdf_eval = 0.5 * sun_pdf;
370 // }
371 let viewZenith = abs(temp_ray.direction.z);
372 let extinction = exp(-2.0 * pow(1.0 - viewZenith, 3.0));
373 let skyRadiance = sky_glow(temp_ray.direction, uniforms.sun_direction) * extinction;
374 let radianceOut = skyRadiance;
375 if (bounce == 0) {
376 radiance += throughput * radianceOut;
377 break;
378 }
379 // bsdf generated ray carries the PDF forward to this bounce
380 var env_pdf_eval = 0.5 * sky_pdf;
381 let env_mis_weight = pdf / (pdf + env_pdf_eval);
382 radiance += clamp_hdr(throughput * radianceOut * env_mis_weight, 10.0);
383 break;
384 }
385
386 let rand = vec2f(uniform_float(¤t_seed), uniform_float(¤t_seed));
387 var material: Material = parse_textures(materials[result.material_idx], result);
388 if (material.emission.x > 0.0 || material.emission.y > 0.0 || material.emission.z > 0.0) {
389 radiance += throughput * material.emission;
390 // break;
391 }
392
393 // sun nee, mis weight based on prior bounce brdf
394 let env_dir = sample_sun_cone_dir(rand);
395 let env_color = sun_glow(env_dir, uniforms.sun_direction);
396 let env_pdf = sun_pdf;
397 let n_dot_env = dot(result.normal, env_dir);
398 if (n_dot_env > 0.0 && !is_occluded(result.position, result.geo_normal, env_dir, 99999.9)) {
399 let env_brdf = eval_brdf(result.normal, -temp_ray.direction, env_dir, material);
400 let diffuse_density = cosine_pdf(result.normal, env_dir);
401 let specular_density = ggx_pdf(-temp_ray.direction, result.normal, normalize(-temp_ray.direction + env_dir), material.roughness);
402 let bsdf_pdf = 0.5 * specular_density + 0.5 * diffuse_density;
403 let weight = env_pdf / (env_pdf + bsdf_pdf);
404 radiance += clamp_hdr(throughput * env_brdf * env_color * n_dot_env * weight / env_pdf, 10.0);
405 }
406
407 // TODO: Better selection, and also move this out.
408 // emissive nee, uniformly sample emissives
409 let light_index = min(u32(floor(rand.x * f32(uniforms.emissive_triangle_count))), u32(uniforms.emissive_triangle_count - 1.0));
410 let tri_index = emissiveTriangleIndices[light_index];
411 let tri = objects.triangles[i32(tri_index)];
412 // uniformly sample point on triangle
413 let u = 1.0 - rand.x;
414 let v = rand.x * (1.0 - rand.y);
415 let w = rand.x * rand.y;
416 let light_pos = u * tri.corner_a + v * tri.corner_b + w * tri.corner_c;
417 let light_normal = normalize(cross(tri.corner_b - tri.corner_a, tri.corner_c - tri.corner_a));
418 let to_light = light_pos - result.position;
419 let dist2 = dot(to_light, to_light);
420 let dist = sqrt(dist2);
421 let light_dir = to_light / dist;
422 let cos_surf = dot(result.normal, light_dir);
423 let cos_light = dot(light_normal, -light_dir);
424
425 if (cos_surf > 0.0 && cos_light > 0.0 && !is_occluded(result.position, result.geo_normal, light_dir, dist)) {
426 var mat = materials[i32(tri.material_idx)];
427 let direct_light_emissive_brdf = eval_brdf(result.normal, -temp_ray.direction, light_dir, material);
428 // compute area of the triangle
429 let edge1 = tri.corner_b - tri.corner_a;
430 let edge2 = tri.corner_c - tri.corner_a;
431 let area = 0.5 * length(cross(edge1, edge2));
432 let light_power = area * mat.emission;
433 // area to solid angle PDF conversion
434 let pdf_solid_angle = dist2 / ( area);
435
436 let diffuse_pdf = cosine_pdf(result.normal, light_dir);
437 let specular_pdf = ggx_pdf(-temp_ray.direction, result.normal, normalize(-temp_ray.direction + light_dir), material.roughness);
438 let bsdf_pdf = 0.5 * diffuse_pdf + 0.5 * specular_pdf;
439 let mis_weight = pdf_solid_angle / (pdf_solid_angle + bsdf_pdf);
440 let contrib = (throughput * direct_light_emissive_brdf * light_power * mis_weight) / pdf_solid_angle;
441 radiance += clamp_hdr(contrib, 10.0);
442 }
443
444 // rr
445 if (bounce > u32(2)) {
446 let rrProbability = min(0.9, luminance(throughput));
447 if (rrProbability < rand.y) {
448 break;
449 } else {
450 throughput /= rrProbability;
451 }
452 }
453
454 var view_dir = -temp_ray.direction;
455 var new_dir: vec3<f32>;
456 var specular_density: f32;
457 var diffuse_density: f32;
458
459 if (uniform_float(¤t_seed) < 0.5) {
460 new_dir = ggx_specular_sample(view_dir, result.normal, rand, material.roughness);
461 } else {
462 new_dir = cosine_hemisphere_sample(result.normal, vec2f(rand.y, rand.x));
463 }
464 let n_dot_l = dot(result.normal, new_dir);
465 if (n_dot_l <= 0.0) { break; }
466 specular_density = ggx_pdf(view_dir, result.normal, normalize(view_dir + new_dir), material.roughness);
467 diffuse_density = cosine_pdf(result.normal, normalize(new_dir));
468 pdf = 0.5 * specular_density + 0.5 * diffuse_density;
469
470 let indirect_brdf = eval_brdf(result.normal, view_dir, new_dir, material);
471 throughput *= (indirect_brdf * n_dot_l) / pdf;
472
473 temp_ray.origin = offset_ray(result.position, result.geo_normal);
474 temp_ray.direction = new_dir;
475 }
476
477 return radiance;
478}
479
480fn hit_triangle(ray: Ray, tri: Triangle, dist_min: f32, dist_max: f32, prevRay: HitInfo) -> HitInfo {
481 var hit: HitInfo;
482 hit.hit = false;
483
484 let edge1 = tri.corner_b - tri.corner_a;
485 let edge2 = tri.corner_c - tri.corner_a;
486
487 let pvec = cross(ray.direction, edge2);
488 let determinant = dot(edge1, pvec);
489
490 // reject nearly parallel rays.
491 if abs(determinant) < EPSILON {
492 return hit;
493 }
494
495 let inv_det = 1.0 / determinant;
496 let tvec = ray.origin - tri.corner_a;
497
498 // compute barycentric coordinate u.
499 let u = dot(tvec, pvec) * inv_det;
500 if (u < 0.0 || u > 1.0) {
501 return hit;
502 }
503
504 // compute barycentric coordinate v.
505 let qvec = cross(tvec, edge1);
506 let v = dot(ray.direction, qvec) * inv_det;
507 if (v < 0.0 || (u + v) > 1.0) {
508 return hit;
509 }
510
511 // calculate ray parameter (distance).
512 let dist = dot(edge2, qvec) * inv_det;
513 if (dist < dist_min || dist > dist_max) {
514 return hit;
515 }
516
517 // no early outs; valid hit
518 hit.hit = true;
519 hit.dist = dist;
520 hit.position = ray.origin + ray.direction * dist;
521 hit.tri = tri;
522 hit.material_idx = i32(tri.material_idx);
523
524 var geo_normal = normalize(cross(edge1, edge2));
525 var shading_normal = normalize((1.0 - u - v) * tri.normal_a + u * tri.normal_b + v * tri.normal_c);
526 let tangent = normalize((1.0 - u - v) * tri.tangent_a + u * tri.tangent_b + v * tri.tangent_c);
527
528 // shadow terminator fix: warp the hit position based on vertex normals
529 // normal aware EPSILON on hit position basically
530 let w = 1.0 - u - v;
531 let tmpu = hit.position - tri.corner_a;
532 let tmpv = hit.position - tri.corner_b;
533 let tmpw = hit.position - tri.corner_c;
534
535 let dotu = min(0.0, dot(tmpu, tri.normal_a));
536 let dotv = min(0.0, dot(tmpv, tri.normal_b));
537 let dotw = min(0.0, dot(tmpw, tri.normal_c));
538
539 let pu = tmpu - dotu * tri.normal_a;
540 let pv = tmpv - dotv * tri.normal_b;
541 let pw = tmpw - dotw * tri.normal_c;
542
543 let warped_offset = w * pu + u * pv + v * pw;
544 // Move the hit point slightly along the warped vector field
545 hit.position = hit.position + warped_offset;
546
547 // TBN
548 let T = normalize(tangent.xyz);
549 let N = normalize(shading_normal);
550 let B = normalize(cross(N, T)) * tangent.w;
551
552 hit.tangent = cross(B, N);
553 hit.normal = shading_normal;
554 hit.uv = (1.0 - u - v) * tri.uv_a + u * tri.uv_b + v * tri.uv_c;
555
556 // If a normal map is present, perturb the shading normal.
557 let material = materials[i32(tri.material_idx)];
558 if (material.normal_texture > -1.0) {
559 var normal_map = sample_material_texture(hit.uv, u32(material.normal_texture));
560 var normalized_map = normalize(normal_map * 2.0 - 1.0);
561 normalized_map.y = -normalized_map.y;
562 let world_normal = normalize(
563 normalized_map.x * T +
564 normalized_map.y * B +
565 normalized_map.z * N
566 );
567 hit.normal = world_normal;
568 }
569 var ray_dot_tri: f32 = dot(ray.direction, geo_normal);
570 if (ray_dot_tri > 0.0) {
571 hit.geo_normal = -hit.geo_normal;
572 hit.normal = -hit.normal;
573 }
574 return hit;
575}
576
577fn hit_aabb(ray: Ray, node: Node) -> f32 {
578 var reciprocal : vec3<f32> = vec3f(1.0) / ray.direction;
579 var t_near: vec3<f32> = (node.min_corner - ray.origin) * reciprocal;
580 var t_far: vec3<f32> = (node.max_corner - ray.origin) * reciprocal;
581 var t_min: vec3<f32> = min(t_near, t_far);
582 var t_max: vec3<f32> = max(t_near, t_far);
583
584 var min_intersection: f32 = max(max(t_min.x, t_min.y), t_min.z); // t0
585 var max_intersection: f32 = min(min(t_max.x, t_max.y), t_max.z); // t1
586
587 var mask: f32 = step(max_intersection, min_intersection) + step(max_intersection, 0.0);
588 if min_intersection > max_intersection || max_intersection < 0 {
589 return 9999.0;
590 } else {
591 return min_intersection;
592 }
593}