Blender  V3.3
kernel/device/metal/bvh.h
Go to the documentation of this file.
1 /* SPDX-License-Identifier: Apache-2.0
2  * Copyright 2021-2022 Blender Foundation */
3 
4 /* MetalRT implementation of ray-scene intersection. */
5 
6 #pragma once
7 
8 #include "kernel/bvh/types.h"
9 #include "kernel/bvh/util.h"
10 
12 
13 /* Payload types. */
14 
18  float u, v;
19  int prim;
20  int type;
21 #if defined(__METALRT_MOTION__)
22  float time;
23 #endif
24 };
25 
30  short max_hits;
32  bool result;
34 };
35 
39 #if defined(__METALRT_MOTION__)
40  float time;
41 #endif
42  int state;
43  float throughput;
44  short max_hits;
45  short num_hits;
47  bool result;
48 };
49 
50 /* Scene intersection. */
51 
53  ccl_private const Ray *ray,
54  const uint visibility,
56 {
57  if (!intersection_ray_valid(ray)) {
58  isect->t = ray->tmax;
59  isect->type = PRIMITIVE_NONE;
60  return false;
61  }
62 
63 #if defined(__KERNEL_DEBUG__)
64  if (is_null_instance_acceleration_structure(metal_ancillaries->accel_struct)) {
65  isect->t = ray->tmax;
66  isect->type = PRIMITIVE_NONE;
67  kernel_assert(!"Invalid metal_ancillaries->accel_struct pointer");
68  return false;
69  }
70 
71  if (is_null_intersection_function_table(metal_ancillaries->ift_default)) {
72  isect->t = ray->tmax;
73  isect->type = PRIMITIVE_NONE;
74  kernel_assert(!"Invalid ift_default");
75  return false;
76  }
77 #endif
78 
79  metal::raytracing::ray r(ray->P, ray->D, ray->tmin, ray->tmax);
80  metalrt_intersector_type metalrt_intersect;
81 
82  bool triangle_only = !kernel_data.bvh.have_curves && !kernel_data.bvh.have_points;
83  if (triangle_only) {
84  metalrt_intersect.assume_geometry_type(metal::raytracing::geometry_type::triangle);
85  }
86 
88  payload.self = ray->self;
89  payload.u = 0.0f;
90  payload.v = 0.0f;
91  payload.visibility = visibility;
92 
93  typename metalrt_intersector_type::result_type intersection;
94 
95  uint ray_mask = visibility & 0xFF;
96  if (0 == ray_mask && (visibility & ~0xFF) != 0) {
97  ray_mask = 0xFF;
98  /* No further intersector setup required: Default MetalRT behavior is any-hit. */
99  }
100  else if (visibility & PATH_RAY_SHADOW_OPAQUE) {
101  /* No further intersector setup required: Shadow ray early termination is controlled by the
102  * intersection handler */
103  }
104 
105 #if defined(__METALRT_MOTION__)
106  payload.time = ray->time;
107  intersection = metalrt_intersect.intersect(r,
108  metal_ancillaries->accel_struct,
109  ray_mask,
110  ray->time,
111  metal_ancillaries->ift_default,
112  payload);
113 #else
114  intersection = metalrt_intersect.intersect(
115  r, metal_ancillaries->accel_struct, ray_mask, metal_ancillaries->ift_default, payload);
116 #endif
117 
119  isect->t = ray->tmax;
120  isect->type = PRIMITIVE_NONE;
121 
122  return false;
123  }
124 
125  isect->t = intersection.distance;
126 
127  isect->prim = payload.prim;
128  isect->type = payload.type;
129  isect->object = intersection.user_instance_id;
130 
131  isect->t = intersection.distance;
132  if (intersection.type == intersection_type::triangle) {
133  isect->u = intersection.triangle_barycentric_coord.x;
134  isect->v = intersection.triangle_barycentric_coord.y;
135  }
136  else {
137  isect->u = payload.u;
138  isect->v = payload.v;
139  }
140 
141  return isect->type != PRIMITIVE_NONE;
142 }
143 
144 #ifdef __BVH_LOCAL__
145 ccl_device_intersect bool scene_intersect_local(KernelGlobals kg,
146  ccl_private const Ray *ray,
147  ccl_private LocalIntersection *local_isect,
148  int local_object,
149  ccl_private uint *lcg_state,
150  int max_hits)
151 {
152  if (!intersection_ray_valid(ray)) {
153  if (local_isect) {
154  local_isect->num_hits = 0;
155  }
156  return false;
157  }
158 
159 # if defined(__KERNEL_DEBUG__)
160  if (is_null_instance_acceleration_structure(metal_ancillaries->accel_struct)) {
161  if (local_isect) {
162  local_isect->num_hits = 0;
163  }
164  kernel_assert(!"Invalid metal_ancillaries->accel_struct pointer");
165  return false;
166  }
167 
168  if (is_null_intersection_function_table(metal_ancillaries->ift_local)) {
169  if (local_isect) {
170  local_isect->num_hits = 0;
171  }
172  kernel_assert(!"Invalid ift_local");
173  return false;
174  }
175 # endif
176 
177  metal::raytracing::ray r(ray->P, ray->D, ray->tmin, ray->tmax);
178  metalrt_intersector_type metalrt_intersect;
179 
180  metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque);
181 
182  bool triangle_only = !kernel_data.bvh.have_curves && !kernel_data.bvh.have_points;
183  if (triangle_only) {
184  metalrt_intersect.assume_geometry_type(metal::raytracing::geometry_type::triangle);
185  }
186 
188  payload.self = ray->self;
189  payload.local_object = local_object;
190  payload.max_hits = max_hits;
191  payload.local_isect.num_hits = 0;
192  if (lcg_state) {
193  payload.has_lcg_state = true;
194  payload.lcg_state = *lcg_state;
195  }
196  payload.result = false;
197 
198  typename metalrt_intersector_type::result_type intersection;
199 
200 # if defined(__METALRT_MOTION__)
201  intersection = metalrt_intersect.intersect(
202  r, metal_ancillaries->accel_struct, 0xFF, ray->time, metal_ancillaries->ift_local, payload);
203 # else
204  intersection = metalrt_intersect.intersect(
205  r, metal_ancillaries->accel_struct, 0xFF, metal_ancillaries->ift_local, payload);
206 # endif
207 
208  if (lcg_state) {
209  *lcg_state = payload.lcg_state;
210  }
211  if (local_isect) {
212  *local_isect = payload.local_isect;
213  }
214 
215  return payload.result;
216 }
217 #endif
218 
219 #ifdef __SHADOW_RECORD_ALL__
220 ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg,
222  ccl_private const Ray *ray,
223  uint visibility,
224  uint max_hits,
225  ccl_private uint *num_recorded_hits,
226  ccl_private float *throughput)
227 {
228  if (!intersection_ray_valid(ray)) {
229  return false;
230  }
231 
232 # if defined(__KERNEL_DEBUG__)
233  if (is_null_instance_acceleration_structure(metal_ancillaries->accel_struct)) {
234  kernel_assert(!"Invalid metal_ancillaries->accel_struct pointer");
235  return false;
236  }
237 
238  if (is_null_intersection_function_table(metal_ancillaries->ift_shadow)) {
239  kernel_assert(!"Invalid ift_shadow");
240  return false;
241  }
242 # endif
243 
244  metal::raytracing::ray r(ray->P, ray->D, ray->tmin, ray->tmax);
245  metalrt_intersector_type metalrt_intersect;
246 
247  metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque);
248 
249  bool triangle_only = !kernel_data.bvh.have_curves && !kernel_data.bvh.have_points;
250  if (triangle_only) {
251  metalrt_intersect.assume_geometry_type(metal::raytracing::geometry_type::triangle);
252  }
253 
255  payload.self = ray->self;
256  payload.visibility = visibility;
257  payload.max_hits = max_hits;
258  payload.num_hits = 0;
259  payload.num_recorded_hits = 0;
260  payload.throughput = 1.0f;
261  payload.result = false;
262  payload.state = state;
263 
264  uint ray_mask = visibility & 0xFF;
265  if (0 == ray_mask && (visibility & ~0xFF) != 0) {
266  ray_mask = 0xFF;
267  }
268 
269  typename metalrt_intersector_type::result_type intersection;
270 
271 # if defined(__METALRT_MOTION__)
272  payload.time = ray->time;
273  intersection = metalrt_intersect.intersect(r,
274  metal_ancillaries->accel_struct,
275  ray_mask,
276  ray->time,
277  metal_ancillaries->ift_shadow,
278  payload);
279 # else
280  intersection = metalrt_intersect.intersect(
281  r, metal_ancillaries->accel_struct, ray_mask, metal_ancillaries->ift_shadow, payload);
282 # endif
283 
284  *num_recorded_hits = payload.num_recorded_hits;
285  *throughput = payload.throughput;
286 
287  return payload.result;
288 }
289 #endif
290 
291 #ifdef __VOLUME__
292 ccl_device_intersect bool scene_intersect_volume(KernelGlobals kg,
293  ccl_private const Ray *ray,
294  ccl_private Intersection *isect,
295  const uint visibility)
296 {
297  if (!intersection_ray_valid(ray)) {
298  return false;
299  }
300 
301 # if defined(__KERNEL_DEBUG__)
302  if (is_null_instance_acceleration_structure(metal_ancillaries->accel_struct)) {
303  kernel_assert(!"Invalid metal_ancillaries->accel_struct pointer");
304  return false;
305  }
306 
307  if (is_null_intersection_function_table(metal_ancillaries->ift_default)) {
308  kernel_assert(!"Invalid ift_default");
309  return false;
310  }
311 # endif
312 
313  metal::raytracing::ray r(ray->P, ray->D, ray->tmin, ray->tmax);
314  metalrt_intersector_type metalrt_intersect;
315 
316  metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque);
317 
318  bool triangle_only = !kernel_data.bvh.have_curves && !kernel_data.bvh.have_points;
319  if (triangle_only) {
320  metalrt_intersect.assume_geometry_type(metal::raytracing::geometry_type::triangle);
321  }
322 
324  payload.self = ray->self;
325  payload.visibility = visibility;
326 
327  typename metalrt_intersector_type::result_type intersection;
328 
329  uint ray_mask = visibility & 0xFF;
330  if (0 == ray_mask && (visibility & ~0xFF) != 0) {
331  ray_mask = 0xFF;
332  }
333 
334 # if defined(__METALRT_MOTION__)
335  payload.time = ray->time;
336  intersection = metalrt_intersect.intersect(r,
337  metal_ancillaries->accel_struct,
338  ray_mask,
339  ray->time,
340  metal_ancillaries->ift_default,
341  payload);
342 # else
343  intersection = metalrt_intersect.intersect(
344  r, metal_ancillaries->accel_struct, ray_mask, metal_ancillaries->ift_default, payload);
345 # endif
346 
348  return false;
349  }
350 
351  isect->prim = payload.prim;
352  isect->type = payload.type;
353  isect->object = intersection.user_instance_id;
354 
355  isect->t = intersection.distance;
356  if (intersection.type == intersection_type::triangle) {
357  isect->u = intersection.triangle_barycentric_coord.x;
358  isect->v = intersection.triangle_barycentric_coord.y;
359  }
360  else {
361  isect->u = payload.u;
362  isect->v = payload.v;
363  }
364 
365  return isect->type != PRIMITIVE_NONE;
366 }
367 #endif
368 
unsigned int uint
Definition: BLI_sys_types.h:67
_GL_VOID GLfloat value _GL_VOID_RET _GL_VOID const GLuint GLboolean *residences _GL_BOOL_RET _GL_VOID GLsizei GLfloat GLfloat GLfloat GLfloat const GLubyte *bitmap _GL_VOID_RET _GL_VOID GLenum const void *lists _GL_VOID_RET _GL_VOID const GLdouble *equation _GL_VOID_RET _GL_VOID GLdouble GLdouble blue _GL_VOID_RET _GL_VOID GLfloat GLfloat blue _GL_VOID_RET _GL_VOID GLint GLint blue _GL_VOID_RET _GL_VOID GLshort GLshort blue _GL_VOID_RET _GL_VOID GLubyte GLubyte blue _GL_VOID_RET _GL_VOID GLuint GLuint blue _GL_VOID_RET _GL_VOID GLushort GLushort blue _GL_VOID_RET _GL_VOID GLbyte GLbyte GLbyte alpha _GL_VOID_RET _GL_VOID GLdouble GLdouble GLdouble alpha _GL_VOID_RET _GL_VOID GLfloat GLfloat GLfloat alpha _GL_VOID_RET _GL_VOID GLint GLint GLint alpha _GL_VOID_RET _GL_VOID GLshort GLshort GLshort alpha _GL_VOID_RET _GL_VOID GLubyte GLubyte GLubyte alpha _GL_VOID_RET _GL_VOID GLuint GLuint GLuint alpha _GL_VOID_RET _GL_VOID GLushort GLushort GLushort alpha _GL_VOID_RET _GL_VOID GLenum mode _GL_VOID_RET _GL_VOID GLint GLsizei GLsizei GLenum type _GL_VOID_RET _GL_VOID GLsizei GLenum GLenum const void *pixels _GL_VOID_RET _GL_VOID const void *pointer _GL_VOID_RET _GL_VOID GLdouble v _GL_VOID_RET _GL_VOID GLfloat v _GL_VOID_RET _GL_VOID GLint GLint i2 _GL_VOID_RET _GL_VOID GLint j _GL_VOID_RET _GL_VOID GLfloat param _GL_VOID_RET _GL_VOID GLint param _GL_VOID_RET _GL_VOID GLdouble GLdouble GLdouble GLdouble GLdouble zFar _GL_VOID_RET _GL_UINT GLdouble *equation _GL_VOID_RET _GL_VOID GLenum GLint *params _GL_VOID_RET _GL_VOID GLenum GLfloat *v _GL_VOID_RET _GL_VOID GLenum GLfloat *params _GL_VOID_RET _GL_VOID GLfloat *values _GL_VOID_RET _GL_VOID GLushort *values _GL_VOID_RET _GL_VOID GLenum GLfloat *params _GL_VOID_RET _GL_VOID GLenum GLdouble *params _GL_VOID_RET _GL_VOID GLenum GLint *params _GL_VOID_RET _GL_VOID GLsizei const void *pointer _GL_VOID_RET _GL_VOID GLsizei const void *pointer _GL_VOID_RET _GL_BOOL GLfloat param _GL_VOID_RET _GL_VOID GLint param _GL_VOID_RET _GL_VOID GLenum GLfloat param _GL_VOID_RET _GL_VOID GLenum GLint param _GL_VOID_RET _GL_VOID GLushort pattern _GL_VOID_RET _GL_VOID GLdouble GLdouble GLint GLint const GLdouble *points _GL_VOID_RET _GL_VOID GLdouble GLdouble GLint GLint GLdouble GLdouble GLint GLint const GLdouble *points _GL_VOID_RET _GL_VOID GLdouble GLdouble u2 _GL_VOID_RET _GL_VOID GLdouble GLdouble GLint GLdouble GLdouble v2 _GL_VOID_RET _GL_VOID GLenum GLfloat param _GL_VOID_RET _GL_VOID GLenum GLint param _GL_VOID_RET _GL_VOID GLenum mode _GL_VOID_RET _GL_VOID GLdouble GLdouble nz _GL_VOID_RET _GL_VOID GLfloat GLfloat nz _GL_VOID_RET _GL_VOID GLint GLint nz _GL_VOID_RET _GL_VOID GLshort GLshort nz _GL_VOID_RET _GL_VOID GLsizei const void *pointer _GL_VOID_RET _GL_VOID GLsizei const GLfloat *values _GL_VOID_RET _GL_VOID GLsizei const GLushort *values _GL_VOID_RET _GL_VOID GLint param _GL_VOID_RET _GL_VOID const GLuint const GLclampf *priorities _GL_VOID_RET _GL_VOID GLdouble y _GL_VOID_RET _GL_VOID GLfloat y _GL_VOID_RET _GL_VOID GLint y _GL_VOID_RET _GL_VOID GLshort y _GL_VOID_RET _GL_VOID GLdouble GLdouble z _GL_VOID_RET _GL_VOID GLfloat GLfloat z _GL_VOID_RET _GL_VOID GLint GLint z _GL_VOID_RET _GL_VOID GLshort GLshort z _GL_VOID_RET _GL_VOID GLdouble GLdouble GLdouble w _GL_VOID_RET _GL_VOID GLfloat GLfloat GLfloat w _GL_VOID_RET _GL_VOID GLint GLint GLint w _GL_VOID_RET _GL_VOID GLshort GLshort GLshort w _GL_VOID_RET _GL_VOID GLdouble GLdouble GLdouble y2 _GL_VOID_RET _GL_VOID GLfloat GLfloat GLfloat y2 _GL_VOID_RET _GL_VOID GLint GLint GLint y2 _GL_VOID_RET _GL_VOID GLshort GLshort GLshort y2 _GL_VOID_RET _GL_VOID GLdouble GLdouble GLdouble z _GL_VOID_RET _GL_VOID GLdouble GLdouble z _GL_VOID_RET _GL_VOID GLuint *buffer _GL_VOID_RET _GL_VOID GLdouble t _GL_VOID_RET _GL_VOID GLfloat t _GL_VOID_RET _GL_VOID GLint t _GL_VOID_RET _GL_VOID GLshort t _GL_VOID_RET _GL_VOID GLdouble GLdouble r _GL_VOID_RET _GL_VOID GLfloat GLfloat r _GL_VOID_RET _GL_VOID GLint GLint r _GL_VOID_RET _GL_VOID GLshort GLshort r _GL_VOID_RET _GL_VOID GLdouble GLdouble r
__forceinline bool none(const avxb &b)
Definition: avxb.h:209
#define kernel_assert(cond)
Definition: cpu/compat.h:34
#define ccl_private
Definition: cuda/compat.h:48
#define CCL_NAMESPACE_END
Definition: cuda/compat.h:9
CCL_NAMESPACE_BEGIN ccl_device_inline bool intersection_ray_valid(ccl_private const Ray *ray)
double time
#define kernel_data
const KernelGlobalsCPU *ccl_restrict KernelGlobals
#define ccl_device_intersect
const int state
ccl_device_intersect bool scene_intersect(KernelGlobals kg, ccl_private const Ray *ray, const uint visibility, ccl_private Intersection *isect)
@ PRIMITIVE_NONE
Definition: kernel/types.h:550
@ PATH_RAY_SHADOW_OPAQUE
Definition: kernel/types.h:204
Intersection< segment > intersection
IntegratorShadowStateCPU *ccl_restrict IntegratorShadowState
Definition: state.h:149