diff options
Diffstat (limited to 'intern/cycles/kernel/bvh')
| -rw-r--r-- | intern/cycles/kernel/bvh/bvh.h | 312 | ||||
| -rw-r--r-- | intern/cycles/kernel/bvh/metal.h | 47 | ||||
| -rw-r--r-- | intern/cycles/kernel/bvh/shadow_all.h | 29 | ||||
| -rw-r--r-- | intern/cycles/kernel/bvh/traversal.h | 28 | ||||
| -rw-r--r-- | intern/cycles/kernel/bvh/types.h | 1 | ||||
| -rw-r--r-- | intern/cycles/kernel/bvh/util.h | 24 |
6 files changed, 417 insertions, 24 deletions
diff --git a/intern/cycles/kernel/bvh/bvh.h b/intern/cycles/kernel/bvh/bvh.h index 0e083812355..67804fb1d0d 100644 --- a/intern/cycles/kernel/bvh/bvh.h +++ b/intern/cycles/kernel/bvh/bvh.h @@ -31,6 +31,10 @@ # include "kernel/bvh/embree.h" #endif +#ifdef __METALRT__ +# include "kernel/bvh/metal.h" +#endif + #include "kernel/bvh/types.h" #include "kernel/bvh/util.h" @@ -38,31 +42,31 @@ CCL_NAMESPACE_BEGIN -#ifndef __KERNEL_OPTIX__ +#if !defined(__KERNEL_GPU_RAYTRACING__) /* Regular BVH traversal */ # include "kernel/bvh/nodes.h" # define BVH_FUNCTION_NAME bvh_intersect -# define BVH_FUNCTION_FEATURES 0 +# define BVH_FUNCTION_FEATURES BVH_POINTCLOUD # include "kernel/bvh/traversal.h" # if defined(__HAIR__) # define BVH_FUNCTION_NAME bvh_intersect_hair -# define BVH_FUNCTION_FEATURES BVH_HAIR +# define BVH_FUNCTION_FEATURES BVH_HAIR | BVH_POINTCLOUD # include "kernel/bvh/traversal.h" # endif # if defined(__OBJECT_MOTION__) # define BVH_FUNCTION_NAME bvh_intersect_motion -# define BVH_FUNCTION_FEATURES BVH_MOTION +# define BVH_FUNCTION_FEATURES BVH_MOTION | BVH_POINTCLOUD # include "kernel/bvh/traversal.h" # endif # if defined(__HAIR__) && defined(__OBJECT_MOTION__) # define BVH_FUNCTION_NAME bvh_intersect_hair_motion -# define BVH_FUNCTION_FEATURES BVH_HAIR | BVH_MOTION +# define BVH_FUNCTION_FEATURES BVH_HAIR | BVH_MOTION | BVH_POINTCLOUD # include "kernel/bvh/traversal.h" # endif @@ -98,26 +102,27 @@ CCL_NAMESPACE_BEGIN # if defined(__SHADOW_RECORD_ALL__) # define BVH_FUNCTION_NAME bvh_intersect_shadow_all -# define BVH_FUNCTION_FEATURES 0 +# define BVH_FUNCTION_FEATURES BVH_POINTCLOUD # include "kernel/bvh/shadow_all.h" # if defined(__HAIR__) # define BVH_FUNCTION_NAME bvh_intersect_shadow_all_hair -# define BVH_FUNCTION_FEATURES BVH_HAIR +# define BVH_FUNCTION_FEATURES BVH_HAIR | BVH_POINTCLOUD # include "kernel/bvh/shadow_all.h" # endif # if defined(__OBJECT_MOTION__) # define BVH_FUNCTION_NAME bvh_intersect_shadow_all_motion -# define BVH_FUNCTION_FEATURES BVH_MOTION +# define BVH_FUNCTION_FEATURES BVH_MOTION | BVH_POINTCLOUD # include "kernel/bvh/shadow_all.h" # endif # if defined(__HAIR__) && defined(__OBJECT_MOTION__) # define BVH_FUNCTION_NAME bvh_intersect_shadow_all_hair_motion -# define BVH_FUNCTION_FEATURES BVH_HAIR | BVH_MOTION +# define BVH_FUNCTION_FEATURES BVH_HAIR | BVH_MOTION | BVH_POINTCLOUD # include "kernel/bvh/shadow_all.h" # endif + # endif /* __SHADOW_RECORD_ALL__ */ /* Record all intersections - Volume BVH traversal. */ @@ -139,7 +144,7 @@ CCL_NAMESPACE_BEGIN # undef BVH_NAME_EVAL # undef BVH_FUNCTION_FULL_NAME -#endif /* __KERNEL_OPTIX__ */ +#endif /* !defined(__KERNEL_GPU_RAYTRACING__) */ ccl_device_inline bool scene_intersect_valid(ccl_private const Ray *ray) { @@ -205,7 +210,95 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg, isect->type = p5; return p5 != PRIMITIVE_NONE; -#else /* __KERNEL_OPTIX__ */ +#elif defined(__METALRT__) + + if (!scene_intersect_valid(ray)) { + isect->t = ray->t; + isect->type = PRIMITIVE_NONE; + return false; + } + +# if defined(__KERNEL_DEBUG__) + if (is_null_instance_acceleration_structure(metal_ancillaries->accel_struct)) { + isect->t = ray->t; + isect->type = PRIMITIVE_NONE; + kernel_assert(!"Invalid metal_ancillaries->accel_struct pointer"); + return false; + } + + if (is_null_intersection_function_table(metal_ancillaries->ift_default)) { + isect->t = ray->t; + isect->type = PRIMITIVE_NONE; + kernel_assert(!"Invalid ift_default"); + return false; + } +# endif + + metal::raytracing::ray r(ray->P, ray->D, 0.0f, ray->t); + metalrt_intersector_type metalrt_intersect; + + if (!kernel_data.bvh.have_curves) { + metalrt_intersect.assume_geometry_type(metal::raytracing::geometry_type::triangle); + } + + MetalRTIntersectionPayload payload; + payload.u = 0.0f; + payload.v = 0.0f; + payload.visibility = visibility; + + typename metalrt_intersector_type::result_type intersection; + + uint ray_mask = visibility & 0xFF; + if (0 == ray_mask && (visibility & ~0xFF) != 0) { + ray_mask = 0xFF; + /* No further intersector setup required: Default MetalRT behavior is any-hit. */ + } + else if (visibility & PATH_RAY_SHADOW_OPAQUE) { + /* No further intersector setup required: Shadow ray early termination is controlled by the + * intersection handler */ + } + +# if defined(__METALRT_MOTION__) + payload.time = ray->time; + intersection = metalrt_intersect.intersect(r, + metal_ancillaries->accel_struct, + ray_mask, + ray->time, + metal_ancillaries->ift_default, + payload); +# else + intersection = metalrt_intersect.intersect( + r, metal_ancillaries->accel_struct, ray_mask, metal_ancillaries->ift_default, payload); +# endif + + if (intersection.type == intersection_type::none) { + isect->t = ray->t; + isect->type = PRIMITIVE_NONE; + + return false; + } + + isect->t = intersection.distance; + + isect->prim = payload.prim; + isect->type = payload.type; + isect->object = intersection.user_instance_id; + + isect->t = intersection.distance; + if (intersection.type == intersection_type::triangle) { + isect->u = 1.0f - intersection.triangle_barycentric_coord.y - + intersection.triangle_barycentric_coord.x; + isect->v = intersection.triangle_barycentric_coord.x; + } + else { + isect->u = payload.u; + isect->v = payload.v; + } + + return isect->type != PRIMITIVE_NONE; + +#else + if (!scene_intersect_valid(ray)) { return false; } @@ -289,7 +382,69 @@ ccl_device_intersect bool scene_intersect_local(KernelGlobals kg, p5); return p5; -# else /* __KERNEL_OPTIX__ */ +# elif defined(__METALRT__) + if (!scene_intersect_valid(ray)) { + if (local_isect) { + local_isect->num_hits = 0; + } + return false; + } + +# if defined(__KERNEL_DEBUG__) + if (is_null_instance_acceleration_structure(metal_ancillaries->accel_struct)) { + if (local_isect) { + local_isect->num_hits = 0; + } + kernel_assert(!"Invalid metal_ancillaries->accel_struct pointer"); + return false; + } + + if (is_null_intersection_function_table(metal_ancillaries->ift_local)) { + if (local_isect) { + local_isect->num_hits = 0; + } + kernel_assert(!"Invalid ift_local"); + return false; + } +# endif + + metal::raytracing::ray r(ray->P, ray->D, 0.0f, ray->t); + metalrt_intersector_type metalrt_intersect; + + metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque); + if (!kernel_data.bvh.have_curves) { + metalrt_intersect.assume_geometry_type(metal::raytracing::geometry_type::triangle); + } + + MetalRTIntersectionLocalPayload payload; + payload.local_object = local_object; + payload.max_hits = max_hits; + payload.local_isect.num_hits = 0; + if (lcg_state) { + payload.has_lcg_state = true; + payload.lcg_state = *lcg_state; + } + payload.result = false; + + typename metalrt_intersector_type::result_type intersection; + +# if defined(__METALRT_MOTION__) + intersection = metalrt_intersect.intersect( + r, metal_ancillaries->accel_struct, 0xFF, ray->time, metal_ancillaries->ift_local, payload); +# else + intersection = metalrt_intersect.intersect( + r, metal_ancillaries->accel_struct, 0xFF, metal_ancillaries->ift_local, payload); +# endif + + if (lcg_state) { + *lcg_state = payload.lcg_state; + } + *local_isect = payload.local_isect; + + return payload.result; + +# else + if (!scene_intersect_valid(ray)) { if (local_isect) { local_isect->num_hits = 0; @@ -406,7 +561,67 @@ ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg, *throughput = __uint_as_float(p1); return p5; -# else /* __KERNEL_OPTIX__ */ +# elif defined(__METALRT__) + + if (!scene_intersect_valid(ray)) { + return false; + } + +# if defined(__KERNEL_DEBUG__) + if (is_null_instance_acceleration_structure(metal_ancillaries->accel_struct)) { + kernel_assert(!"Invalid metal_ancillaries->accel_struct pointer"); + return false; + } + + if (is_null_intersection_function_table(metal_ancillaries->ift_shadow)) { + kernel_assert(!"Invalid ift_shadow"); + return false; + } +# endif + + metal::raytracing::ray r(ray->P, ray->D, 0.0f, ray->t); + metalrt_intersector_type metalrt_intersect; + + metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque); + if (!kernel_data.bvh.have_curves) { + metalrt_intersect.assume_geometry_type(metal::raytracing::geometry_type::triangle); + } + + MetalRTIntersectionShadowPayload payload; + payload.visibility = visibility; + payload.max_hits = max_hits; + payload.num_hits = 0; + payload.num_recorded_hits = 0; + payload.throughput = 1.0f; + payload.result = false; + payload.state = state; + + uint ray_mask = visibility & 0xFF; + if (0 == ray_mask && (visibility & ~0xFF) != 0) { + ray_mask = 0xFF; + } + + typename metalrt_intersector_type::result_type intersection; + +# if defined(__METALRT_MOTION__) + payload.time = ray->time; + intersection = metalrt_intersect.intersect(r, + metal_ancillaries->accel_struct, + ray_mask, + ray->time, + metal_ancillaries->ift_shadow, + payload); +# else + intersection = metalrt_intersect.intersect( + r, metal_ancillaries->accel_struct, ray_mask, metal_ancillaries->ift_shadow, payload); +# endif + + *num_recorded_hits = payload.num_recorded_hits; + *throughput = payload.throughput; + + return payload.result; + +# else if (!scene_intersect_valid(ray)) { *num_recorded_hits = 0; *throughput = 1.0f; @@ -503,7 +718,76 @@ ccl_device_intersect bool scene_intersect_volume(KernelGlobals kg, isect->type = p5; return p5 != PRIMITIVE_NONE; -# else /* __KERNEL_OPTIX__ */ +# elif defined(__METALRT__) + + if (!scene_intersect_valid(ray)) { + return false; + } +# if defined(__KERNEL_DEBUG__) + if (is_null_instance_acceleration_structure(metal_ancillaries->accel_struct)) { + kernel_assert(!"Invalid metal_ancillaries->accel_struct pointer"); + return false; + } + + if (is_null_intersection_function_table(metal_ancillaries->ift_default)) { + kernel_assert(!"Invalid ift_default"); + return false; + } +# endif + + metal::raytracing::ray r(ray->P, ray->D, 0.0f, ray->t); + metalrt_intersector_type metalrt_intersect; + + metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque); + if (!kernel_data.bvh.have_curves) { + metalrt_intersect.assume_geometry_type(metal::raytracing::geometry_type::triangle); + } + + MetalRTIntersectionPayload payload; + payload.visibility = visibility; + + typename metalrt_intersector_type::result_type intersection; + + uint ray_mask = visibility & 0xFF; + if (0 == ray_mask && (visibility & ~0xFF) != 0) { + ray_mask = 0xFF; + } + +# if defined(__METALRT_MOTION__) + payload.time = ray->time; + intersection = metalrt_intersect.intersect(r, + metal_ancillaries->accel_struct, + ray_mask, + ray->time, + metal_ancillaries->ift_default, + payload); +# else + intersection = metalrt_intersect.intersect( + r, metal_ancillaries->accel_struct, ray_mask, metal_ancillaries->ift_default, payload); +# endif + + if (intersection.type == intersection_type::none) { + return false; + } + + isect->prim = payload.prim; + isect->type = payload.type; + isect->object = intersection.user_instance_id; + + isect->t = intersection.distance; + if (intersection.type == intersection_type::triangle) { + isect->u = 1.0f - intersection.triangle_barycentric_coord.y - + intersection.triangle_barycentric_coord.x; + isect->v = intersection.triangle_barycentric_coord.x; + } + else { + isect->u = payload.u; + isect->v = payload.v; + } + + return isect->type != PRIMITIVE_NONE; + +# else if (!scene_intersect_valid(ray)) { return false; } diff --git a/intern/cycles/kernel/bvh/metal.h b/intern/cycles/kernel/bvh/metal.h new file mode 100644 index 00000000000..55456d15f50 --- /dev/null +++ b/intern/cycles/kernel/bvh/metal.h @@ -0,0 +1,47 @@ +/* + * Copyright 2021 Blender Foundation + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +struct MetalRTIntersectionPayload { + uint visibility; + float u, v; + int prim; + int type; +#if defined(__METALRT_MOTION__) + float time; +#endif +}; + +struct MetalRTIntersectionLocalPayload { + uint local_object; + uint lcg_state; + short max_hits; + bool has_lcg_state; + bool result; + LocalIntersection local_isect; +}; + +struct MetalRTIntersectionShadowPayload { + uint visibility; +#if defined(__METALRT_MOTION__) + float time; +#endif + int state; + float throughput; + short max_hits; + short num_hits; + short num_recorded_hits; + bool result; +}; diff --git a/intern/cycles/kernel/bvh/shadow_all.h b/intern/cycles/kernel/bvh/shadow_all.h index 049c6a03fe0..caca85aac1a 100644 --- a/intern/cycles/kernel/bvh/shadow_all.h +++ b/intern/cycles/kernel/bvh/shadow_all.h @@ -28,6 +28,7 @@ * without new features slowing things down. * * BVH_HAIR: hair curve rendering + * BVH_POINTCLOUD: point cloud rendering * BVH_MOTION: motion blur rendering */ @@ -199,6 +200,34 @@ ccl_device_inline break; } #endif +#if BVH_FEATURE(BVH_POINTCLOUD) + case PRIMITIVE_POINT: + case PRIMITIVE_MOTION_POINT: { + if ((type & PRIMITIVE_ALL_MOTION) && kernel_data.bvh.use_bvh_steps) { + const float2 prim_time = kernel_tex_fetch(__prim_time, prim_addr); + if (ray->time < prim_time.x || ray->time > prim_time.y) { + hit = false; + break; + } + } + + const int point_object = (object == OBJECT_NONE) ? + kernel_tex_fetch(__prim_object, prim_addr) : + object; + const int point_prim = kernel_tex_fetch(__prim_index, prim_addr); + const int point_type = kernel_tex_fetch(__prim_type, prim_addr); + hit = point_intersect(kg, + &isect, + P, + dir, + t_max_current, + point_object, + point_prim, + ray->time, + point_type); + break; + } +#endif /* BVH_FEATURE(BVH_POINTCLOUD) */ default: { hit = false; break; diff --git a/intern/cycles/kernel/bvh/traversal.h b/intern/cycles/kernel/bvh/traversal.h index 1c17ebf767f..180f19d11c5 100644 --- a/intern/cycles/kernel/bvh/traversal.h +++ b/intern/cycles/kernel/bvh/traversal.h @@ -28,6 +28,7 @@ * without new features slowing things down. * * BVH_HAIR: hair curve rendering + * BVH_POINTCLOUD: point cloud rendering * BVH_MOTION: motion blur rendering */ @@ -188,6 +189,33 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg, break; } #endif /* BVH_FEATURE(BVH_HAIR) */ +#if BVH_FEATURE(BVH_POINTCLOUD) + case PRIMITIVE_POINT: + case PRIMITIVE_MOTION_POINT: { + for (; prim_addr < prim_addr2; prim_addr++) { + if ((type & PRIMITIVE_ALL_MOTION) && kernel_data.bvh.use_bvh_steps) { + const float2 prim_time = kernel_tex_fetch(__prim_time, prim_addr); + if (ray->time < prim_time.x || ray->time > prim_time.y) { + continue; + } + } + + const int point_object = (object == OBJECT_NONE) ? + kernel_tex_fetch(__prim_object, prim_addr) : + object; + const int point_prim = kernel_tex_fetch(__prim_index, prim_addr); + const int point_type = kernel_tex_fetch(__prim_type, prim_addr); + const bool hit = point_intersect( + kg, isect, P, dir, isect->t, point_object, point_prim, ray->time, point_type); + if (hit) { + /* shadow ray early termination */ + if (visibility & PATH_RAY_SHADOW_OPAQUE) + return true; + } + } + break; + } +#endif /* BVH_FEATURE(BVH_POINTCLOUD) */ } } else { diff --git a/intern/cycles/kernel/bvh/types.h b/intern/cycles/kernel/bvh/types.h index 6039e707fc3..f16f43333f8 100644 --- a/intern/cycles/kernel/bvh/types.h +++ b/intern/cycles/kernel/bvh/types.h @@ -34,6 +34,7 @@ CCL_NAMESPACE_BEGIN #define BVH_MOTION 1 #define BVH_HAIR 2 +#define BVH_POINTCLOUD 4 #define BVH_NAME_JOIN(x, y) x##_##y #define BVH_NAME_EVAL(x, y) BVH_NAME_JOIN(x, y) diff --git a/intern/cycles/kernel/bvh/util.h b/intern/cycles/kernel/bvh/util.h index 26ba136dd79..57593e42a88 100644 --- a/intern/cycles/kernel/bvh/util.h +++ b/intern/cycles/kernel/bvh/util.h @@ -118,14 +118,16 @@ ccl_device_forceinline int intersection_get_shader_flags(KernelGlobals kg, { int shader = 0; -#ifdef __HAIR__ - if (type & PRIMITIVE_ALL_TRIANGLE) -#endif - { + if (type & PRIMITIVE_ALL_TRIANGLE) { shader = kernel_tex_fetch(__tri_shader, prim); } +#ifdef __POINTCLOUD__ + else if (type & PRIMITIVE_ALL_POINT) { + shader = kernel_tex_fetch(__points_shader, prim); + } +#endif #ifdef __HAIR__ - else { + else if (type & PRIMITIVE_ALL_CURVE) { shader = kernel_tex_fetch(__curves, prim).shader_id; } #endif @@ -139,14 +141,16 @@ ccl_device_forceinline int intersection_get_shader_from_isect_prim(KernelGlobals { int shader = 0; -#ifdef __HAIR__ - if (isect_type & PRIMITIVE_ALL_TRIANGLE) -#endif - { + if (isect_type & PRIMITIVE_ALL_TRIANGLE) { shader = kernel_tex_fetch(__tri_shader, prim); } +#ifdef __POINTCLOUD__ + else if (isect_type & PRIMITIVE_ALL_POINT) { + shader = kernel_tex_fetch(__points_shader, prim); + } +#endif #ifdef __HAIR__ - else { + else if (isect_type & PRIMITIVE_ALL_CURVE) { shader = kernel_tex_fetch(__curves, prim).shader_id; } #endif |