diff --git a/intern/cycles/device/metal/device_impl.mm b/intern/cycles/device/metal/device_impl.mm index 4a7fdd9202c..588fb6ede2d 100644 --- a/intern/cycles/device/metal/device_impl.mm +++ b/intern/cycles/device/metal/device_impl.mm @@ -217,20 +217,27 @@ MetalDevice::MetalDevice(const DeviceInfo &info, Stats &stats, Profiler &profile arg_desc_as.index = index++; [ancillary_desc addObject:[arg_desc_as copy]]; /* accel_struct */ + + /* Intersection function tables */ arg_desc_ift.index = index++; [ancillary_desc addObject:[arg_desc_ift copy]]; /* ift_default */ arg_desc_ift.index = index++; [ancillary_desc addObject:[arg_desc_ift copy]]; /* ift_shadow */ arg_desc_ift.index = index++; + [ancillary_desc addObject:[arg_desc_ift copy]]; /* ift_shadow_all */ + arg_desc_ift.index = index++; [ancillary_desc addObject:[arg_desc_ift copy]]; /* ift_volume */ arg_desc_ift.index = index++; [ancillary_desc addObject:[arg_desc_ift copy]]; /* ift_local */ arg_desc_ift.index = index++; - [ancillary_desc addObject:[arg_desc_ift copy]]; /* ift_local_prim */ + [ancillary_desc addObject:[arg_desc_ift copy]]; /* ift_local_mblur */ + arg_desc_ift.index = index++; + [ancillary_desc addObject:[arg_desc_ift copy]]; /* ift_local_single_hit */ + arg_desc_ift.index = index++; + [ancillary_desc addObject:[arg_desc_ift copy]]; /* ift_local_single_hit_mblur */ + arg_desc_ptrs.index = index++; - [ancillary_desc addObject:[arg_desc_ptrs copy]]; /* blas array */ - arg_desc_ptrs.index = index++; - [ancillary_desc addObject:[arg_desc_ptrs copy]]; /* look up table for blas */ + [ancillary_desc addObject:[arg_desc_ptrs copy]]; /* blas_accel_structs */ [arg_desc_ift release]; [arg_desc_as release]; diff --git a/intern/cycles/device/metal/kernel.h b/intern/cycles/device/metal/kernel.h index c5e3b354fc1..cc343c1b4e4 100644 --- a/intern/cycles/device/metal/kernel.h +++ b/intern/cycles/device/metal/kernel.h @@ -13,30 +13,15 @@ CCL_NAMESPACE_BEGIN class MetalDevice; -enum { - METALRT_FUNC_DEFAULT_TRI, - METALRT_FUNC_DEFAULT_BOX, - METALRT_FUNC_SHADOW_TRI, - METALRT_FUNC_SHADOW_BOX, - METALRT_FUNC_VOLUME_TRI, - METALRT_FUNC_VOLUME_BOX, - METALRT_FUNC_LOCAL_TRI, - METALRT_FUNC_LOCAL_BOX, - METALRT_FUNC_LOCAL_TRI_PRIM, - METALRT_FUNC_LOCAL_BOX_PRIM, - METALRT_FUNC_CURVE, - METALRT_FUNC_CURVE_SHADOW, - METALRT_FUNC_POINT, - METALRT_FUNC_POINT_SHADOW, - METALRT_FUNC_NUM -}; - enum { METALRT_TABLE_DEFAULT, METALRT_TABLE_SHADOW, + METALRT_TABLE_SHADOW_ALL, METALRT_TABLE_VOLUME, METALRT_TABLE_LOCAL, - METALRT_TABLE_LOCAL_PRIM, + METALRT_TABLE_LOCAL_MBLUR, + METALRT_TABLE_LOCAL_SINGLE_HIT, + METALRT_TABLE_LOCAL_SINGLE_HIT_MBLUR, METALRT_TABLE_NUM }; @@ -94,12 +79,12 @@ struct MetalKernelPipeline { int num_threads_per_block = 0; bool should_use_binary_archive() const; + id make_intersection_function(const char *function_name); string error_str; API_AVAILABLE(macos(11.0)) id intersection_func_table[METALRT_TABLE_NUM] = {nil}; - id rt_intersection_function[METALRT_FUNC_NUM] = {nil}; }; /* Cache of Metal kernels for each DeviceKernel. */ diff --git a/intern/cycles/device/metal/kernel.mm b/intern/cycles/device/metal/kernel.mm index 125c7129de0..38430b8032a 100644 --- a/intern/cycles/device/metal/kernel.mm +++ b/intern/cycles/device/metal/kernel.mm @@ -459,6 +459,35 @@ static MTLFunctionConstantValues *GetConstantValues(KernelData const *data = nul return constant_values; } +id MetalKernelPipeline::make_intersection_function(const char *function_name) +{ + MTLFunctionDescriptor *desc = [MTLIntersectionFunctionDescriptor functionDescriptor]; + desc.name = [@(function_name) copy]; + + if (pso_type != PSO_GENERIC) { + desc.constantValues = GetConstantValues(&kernel_data_); + } + else { + desc.constantValues = GetConstantValues(); + } + + NSError *error = NULL; + id rt_intersection_function = [mtlLibrary newFunctionWithDescriptor:desc + error:&error]; + + if (rt_intersection_function == nil) { + NSString *err = [error localizedDescription]; + string errors = [err UTF8String]; + + error_str = string_printf( + "Error getting intersection function \"%s\": %s", function_name, errors.c_str()); + } + else { + rt_intersection_function.label = [@(function_name) copy]; + } + return rt_intersection_function; +} + void MetalKernelPipeline::compile() { const std::string function_name = std::string("cycles_metal_") + @@ -487,117 +516,50 @@ void MetalKernelPipeline::compile() function.label = [@(function_name.c_str()) copy]; - if (use_metalrt) { - /* create the id for each intersection function */ - const char *function_names[] = { - "__anyhit__cycles_metalrt_visibility_test_tri", - "__anyhit__cycles_metalrt_visibility_test_box", - "__anyhit__cycles_metalrt_shadow_all_hit_tri", - "__anyhit__cycles_metalrt_shadow_all_hit_box", - "__anyhit__cycles_metalrt_volume_test_tri", - "__anyhit__cycles_metalrt_volume_test_box", - "__anyhit__cycles_metalrt_local_hit_tri", - "__anyhit__cycles_metalrt_local_hit_box", - "__anyhit__cycles_metalrt_local_hit_tri_prim", - "__anyhit__cycles_metalrt_local_hit_box_prim", - "__intersection__curve", - "__intersection__curve_shadow", - "__intersection__point", - "__intersection__point_shadow", - }; - assert(sizeof(function_names) / sizeof(function_names[0]) == METALRT_FUNC_NUM); - - MTLFunctionDescriptor *desc = [MTLIntersectionFunctionDescriptor functionDescriptor]; - for (int i = 0; i < METALRT_FUNC_NUM; i++) { - const char *function_name = function_names[i]; - desc.name = [@(function_name) copy]; - - if (pso_type != PSO_GENERIC) { - desc.constantValues = GetConstantValues(&kernel_data_); - } - else { - desc.constantValues = GetConstantValues(); - } - - NSError *error = NULL; - rt_intersection_function[i] = [mtlLibrary newFunctionWithDescriptor:desc error:&error]; - - if (rt_intersection_function[i] == nil) { - NSString *err = [error localizedDescription]; - string errors = [err UTF8String]; - - error_str = string_printf( - "Error getting intersection function \"%s\": %s", function_name, errors.c_str()); - break; - } - - rt_intersection_function[i].label = [@(function_name) copy]; - } - } - NSArray *table_functions[METALRT_TABLE_NUM] = {nil}; NSArray *linked_functions = nil; - if (use_metalrt) { - id curve_intersect_default = nil; - id curve_intersect_shadow = nil; - id point_intersect_default = nil; - id point_intersect_shadow = nil; - if (kernel_features & KERNEL_FEATURE_HAIR) { - curve_intersect_default = rt_intersection_function[METALRT_FUNC_CURVE]; - curve_intersect_shadow = rt_intersection_function[METALRT_FUNC_CURVE_SHADOW]; - } - if (kernel_features & KERNEL_FEATURE_POINTCLOUD) { - point_intersect_default = rt_intersection_function[METALRT_FUNC_POINT]; - point_intersect_shadow = rt_intersection_function[METALRT_FUNC_POINT_SHADOW]; - } - table_functions[METALRT_TABLE_DEFAULT] = [NSArray - arrayWithObjects:rt_intersection_function[METALRT_FUNC_DEFAULT_TRI], - curve_intersect_default ? - curve_intersect_default : - rt_intersection_function[METALRT_FUNC_DEFAULT_BOX], - point_intersect_default ? - point_intersect_default : - rt_intersection_function[METALRT_FUNC_DEFAULT_BOX], - nil]; - table_functions[METALRT_TABLE_SHADOW] = [NSArray - arrayWithObjects:rt_intersection_function[METALRT_FUNC_SHADOW_TRI], - curve_intersect_shadow ? - curve_intersect_shadow : - rt_intersection_function[METALRT_FUNC_SHADOW_BOX], - point_intersect_shadow ? - point_intersect_shadow : - rt_intersection_function[METALRT_FUNC_SHADOW_BOX], - nil]; - table_functions[METALRT_TABLE_VOLUME] = [NSArray - arrayWithObjects:rt_intersection_function[METALRT_FUNC_VOLUME_TRI], - rt_intersection_function[METALRT_FUNC_VOLUME_BOX], - rt_intersection_function[METALRT_FUNC_VOLUME_BOX], - nil]; - table_functions[METALRT_TABLE_LOCAL] = [NSArray - arrayWithObjects:rt_intersection_function[METALRT_FUNC_LOCAL_TRI], - rt_intersection_function[METALRT_FUNC_LOCAL_BOX], - rt_intersection_function[METALRT_FUNC_LOCAL_BOX], - nil]; - table_functions[METALRT_TABLE_LOCAL_PRIM] = [NSArray - arrayWithObjects:rt_intersection_function[METALRT_FUNC_LOCAL_TRI_PRIM], - rt_intersection_function[METALRT_FUNC_LOCAL_BOX_PRIM], - rt_intersection_function[METALRT_FUNC_LOCAL_BOX_PRIM], - nil]; + if (use_metalrt && device_kernel_has_intersection(device_kernel)) { NSMutableSet *unique_functions = [[NSMutableSet alloc] init]; - [unique_functions addObjectsFromArray:table_functions[METALRT_TABLE_DEFAULT]]; - [unique_functions addObjectsFromArray:table_functions[METALRT_TABLE_SHADOW]]; - [unique_functions addObjectsFromArray:table_functions[METALRT_TABLE_VOLUME]]; - [unique_functions addObjectsFromArray:table_functions[METALRT_TABLE_LOCAL]]; - [unique_functions addObjectsFromArray:table_functions[METALRT_TABLE_LOCAL_PRIM]]; - if (device_kernel_has_intersection(device_kernel)) { - linked_functions = [[NSArray arrayWithArray:[unique_functions allObjects]] - sortedArrayUsingComparator:^NSComparisonResult(id f1, id f2) { - return [f1.label compare:f2.label]; - }]; - } + auto add_intersection_functions = [&](int table_index, + const char *tri_fn, + const char *curve_fn = nullptr, + const char *point_fn = nullptr) { + table_functions[table_index] = [NSArray + arrayWithObjects:make_intersection_function(tri_fn), + curve_fn ? make_intersection_function(curve_fn) : nil, + point_fn ? make_intersection_function(point_fn) : nil, + nil]; + + [unique_functions addObjectsFromArray:table_functions[table_index]]; + }; + + add_intersection_functions(METALRT_TABLE_DEFAULT, + "__intersection__tri", + "__intersection__curve", + "__intersection__point"); + add_intersection_functions(METALRT_TABLE_SHADOW, + "__intersection__tri_shadow", + "__intersection__curve_shadow", + "__intersection__point_shadow"); + add_intersection_functions(METALRT_TABLE_SHADOW_ALL, + "__intersection__tri_shadow_all", + "__intersection__curve_shadow_all", + "__intersection__point_shadow_all"); + add_intersection_functions(METALRT_TABLE_VOLUME, "__intersection__volume_tri"); + add_intersection_functions(METALRT_TABLE_LOCAL, "__intersection__local_tri"); + add_intersection_functions(METALRT_TABLE_LOCAL_MBLUR, "__intersection__local_tri_mblur"); + add_intersection_functions(METALRT_TABLE_LOCAL_SINGLE_HIT, + "__intersection__local_tri_single_hit"); + add_intersection_functions(METALRT_TABLE_LOCAL_SINGLE_HIT_MBLUR, + "__intersection__local_tri_single_hit_mblur"); + + linked_functions = [[NSArray arrayWithArray:[unique_functions allObjects]] + sortedArrayUsingComparator:^NSComparisonResult(id f1, id f2) { + return [f1.label compare:f2.label]; + }]; unique_functions = nil; } @@ -619,8 +581,8 @@ void MetalKernelPipeline::compile() computePipelineStateDescriptor.linkedFunctions.functions = linked_functions; } computePipelineStateDescriptor.maxCallStackDepth = 1; - if (use_metalrt) { - computePipelineStateDescriptor.maxCallStackDepth = 8; + if (use_metalrt && device_kernel_has_intersection(device_kernel)) { + computePipelineStateDescriptor.maxCallStackDepth = 2; } MTLPipelineOption pipelineOptions = MTLPipelineOptionNone; diff --git a/intern/cycles/device/metal/queue.mm b/intern/cycles/device/metal/queue.mm index f1cec1b0263..7faa2c2f9dd 100644 --- a/intern/cycles/device/metal/queue.mm +++ b/intern/cycles/device/metal/queue.mm @@ -495,13 +495,13 @@ bool MetalDeviceQueue::enqueue(DeviceKernel kernel, atIndex:2]; if (@available(macos 12.0, *)) { - if (metal_device_->use_metalrt) { + if (metal_device_->use_metalrt && device_kernel_has_intersection(kernel)) { if (metal_device_->bvhMetalRT) { id accel_struct = metal_device_->bvhMetalRT->accel_struct; [metal_device_->mtlAncillaryArgEncoder setAccelerationStructure:accel_struct atIndex:3]; [metal_device_->mtlAncillaryArgEncoder setBuffer:metal_device_->blas_buffer offset:0 - atIndex:9]; + atIndex:(METALRT_TABLE_NUM + 4)]; } for (int table = 0; table < METALRT_TABLE_NUM; table++) { @@ -532,24 +532,10 @@ bool MetalDeviceQueue::enqueue(DeviceKernel kernel, [mtlComputeCommandEncoder setBuffer:arg_buffer offset:globals_offsets atIndex:1]; [mtlComputeCommandEncoder setBuffer:arg_buffer offset:metal_offsets atIndex:2]; - if (metal_device_->use_metalrt) { + if (metal_device_->use_metalrt && device_kernel_has_intersection(kernel)) { if (@available(macos 12.0, *)) { - auto bvhMetalRT = metal_device_->bvhMetalRT; - switch (kernel) { - case DEVICE_KERNEL_INTEGRATOR_INTERSECT_CLOSEST: - case DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW: - case DEVICE_KERNEL_INTEGRATOR_INTERSECT_SUBSURFACE: - case DEVICE_KERNEL_INTEGRATOR_INTERSECT_VOLUME_STACK: - case DEVICE_KERNEL_INTEGRATOR_INTERSECT_DEDICATED_LIGHT: - case DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE: - case DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_MNEE: - break; - default: - bvhMetalRT = nil; - break; - } - + BVHMetal *bvhMetalRT = metal_device_->bvhMetalRT; if (bvhMetalRT && bvhMetalRT->accel_struct) { /* Mark all Accelerations resources as used */ [mtlComputeCommandEncoder useResource:bvhMetalRT->accel_struct diff --git a/intern/cycles/kernel/bvh/bvh.h b/intern/cycles/kernel/bvh/bvh.h index 8f4dea3dc80..18da0b75f5c 100644 --- a/intern/cycles/kernel/bvh/bvh.h +++ b/intern/cycles/kernel/bvh/bvh.h @@ -134,6 +134,14 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg, return false; } +ccl_device_intersect bool scene_intersect_shadow(KernelGlobals kg, + ccl_private const Ray *ray, + const uint visibility) +{ + Intersection isect; + return scene_intersect(kg, ray, visibility, &isect); +} + /* Single object BVH traversal, for SSS/AO/bevel. */ # ifdef __BVH_LOCAL__ @@ -148,6 +156,7 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg, # include "kernel/bvh/local.h" # endif +template ccl_device_intersect bool scene_intersect_local(KernelGlobals kg, ccl_private const Ray *ray, ccl_private LocalIntersection *local_isect, diff --git a/intern/cycles/kernel/device/hiprt/bvh.h b/intern/cycles/kernel/device/hiprt/bvh.h index a7242d99c7b..b7c818d44a7 100644 --- a/intern/cycles/kernel/device/hiprt/bvh.h +++ b/intern/cycles/kernel/device/hiprt/bvh.h @@ -64,7 +64,16 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg, return false; } +ccl_device_intersect bool scene_intersect_shadow(KernelGlobals kg, + ccl_private const Ray *ray, + const uint visibility) +{ + Intersection isect; + return scene_intersect(kg, ray, visibility, &isect); +} + #ifdef __BVH_LOCAL__ +template ccl_device_intersect bool scene_intersect_local(KernelGlobals kg, ccl_private const Ray *ray, ccl_private LocalIntersection *local_isect, diff --git a/intern/cycles/kernel/device/metal/bvh.h b/intern/cycles/kernel/device/metal/bvh.h index d783e13dae1..69bc2ba6a2d 100644 --- a/intern/cycles/kernel/device/metal/bvh.h +++ b/intern/cycles/kernel/device/metal/bvh.h @@ -11,26 +11,41 @@ CCL_NAMESPACE_BEGIN -/* Payload types. */ +/* Payload types. + * + * Best practice is to minimize the size of MetalRT payloads to avoid heavy spilling during + * intersection tests. + */ struct MetalRTIntersectionPayload { - RaySelfPrimitives self; - uint visibility; + int self_prim; + int self_object; +}; + +struct MetalRTIntersectionLocalPayload_single_hit { + int self_prim; }; struct MetalRTIntersectionLocalPayload { - RaySelfPrimitives self; - uint local_object; + int self_prim; uint lcg_state; - short max_hits; - bool has_lcg_state; - bool result; - LocalIntersection local_isect; + uint hit_prim[LOCAL_MAX_HITS]; + float hit_t[LOCAL_MAX_HITS]; + float hit_u[LOCAL_MAX_HITS]; + float hit_v[LOCAL_MAX_HITS]; + uint max_hits : 3; + uint num_hits : 3; + uint has_lcg_state : 1; }; +static_assert(LOCAL_MAX_HITS < 8, + "MetalRTIntersectionLocalPayload max_hits & num_hits bitfields are too small"); struct MetalRTIntersectionShadowPayload { RaySelfPrimitives self; - uint visibility; +}; + +struct MetalRTIntersectionShadowAllPayload { + RaySelfPrimitives self; int state; float throughput; short max_hits; @@ -140,28 +155,6 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg, const uint visibility, ccl_private Intersection *isect) { - if (!intersection_ray_valid(ray)) { - isect->t = ray->tmax; - isect->type = PRIMITIVE_NONE; - return false; - } - -#if defined(WITH_CYCLES_DEBUG) - if (is_null_instance_acceleration_structure(metal_ancillaries->accel_struct)) { - isect->t = ray->tmax; - 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->tmax; - isect->type = PRIMITIVE_NONE; - kernel_assert(!"Invalid ift_default"); - return false; - } -#endif - metal::raytracing::ray r(ray->P, ray->D, ray->tmin, ray->tmax); metalrt_intersector_type metalrt_intersect; metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque); @@ -172,15 +165,11 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg, (kernel_data.bvh.have_points ? metal::raytracing::geometry_type::bounding_box : metal::raytracing::geometry_type::none)); - if (visibility & PATH_RAY_SHADOW_OPAQUE) { - metalrt_intersect.accept_any_intersection(true); - } + typename metalrt_intersector_type::result_type intersection; MetalRTIntersectionPayload payload; - payload.self = ray->self; - payload.visibility = visibility; - - typename metalrt_intersector_type::result_type intersection; + payload.self_prim = ray->self.prim; + payload.self_object = ray->self.object; #if defined(__METALRT_MOTION__) intersection = metalrt_intersect.intersect(r, @@ -272,7 +261,43 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg, return true; } +ccl_device_intersect bool scene_intersect_shadow(KernelGlobals kg, + ccl_private const Ray *ray, + const uint visibility) +{ + metal::raytracing::ray r(ray->P, ray->D, ray->tmin, ray->tmax); + metalrt_intersector_type metalrt_intersect; + metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque); + metalrt_intersect.assume_geometry_type( + metal::raytracing::geometry_type::triangle | + (kernel_data.bvh.have_curves ? metal::raytracing::geometry_type::curve : + metal::raytracing::geometry_type::none) | + (kernel_data.bvh.have_points ? metal::raytracing::geometry_type::bounding_box : + metal::raytracing::geometry_type::none)); + + typename metalrt_intersector_type::result_type intersection; + + metalrt_intersect.accept_any_intersection(true); + + MetalRTIntersectionShadowPayload payload; + payload.self = ray->self; + +#if defined(__METALRT_MOTION__) + intersection = metalrt_intersect.intersect(r, + metal_ancillaries->accel_struct, + visibility, + ray->time, + metal_ancillaries->ift_shadow, + payload); +#else + intersection = metalrt_intersect.intersect( + r, metal_ancillaries->accel_struct, visibility, metal_ancillaries->ift_shadow, payload); +#endif + return (intersection.type != intersection_type::none); +} + #ifdef __BVH_LOCAL__ +template ccl_device_intersect bool scene_intersect_local(KernelGlobals kg, ccl_private const Ray *ray, ccl_private LocalIntersection *local_isect, @@ -280,48 +305,7 @@ ccl_device_intersect bool scene_intersect_local(KernelGlobals kg, ccl_private uint *lcg_state, int max_hits) { - if (!intersection_ray_valid(ray)) { - if (local_isect) { - local_isect->num_hits = 0; - } - return false; - } - -# if defined(WITH_CYCLES_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; - } - if (is_null_intersection_function_table(metal_ancillaries->ift_local_prim)) { - if (local_isect) { - local_isect->num_hits = 0; - } - kernel_assert(!"Invalid ift_local_prim"); - return false; - } -# endif - - MetalRTIntersectionLocalPayload payload; - payload.self = ray->self; - 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; + uint primitive_id_offset = kernel_data_fetch(object_prim_offset, local_object); metal::raytracing::ray r(ray->P, ray->D, ray->tmin, ray->tmax); @@ -331,48 +315,130 @@ ccl_device_intersect bool scene_intersect_local(KernelGlobals kg, # else metalrt_blas_intersector_type metalrt_intersect; typename metalrt_blas_intersector_type::result_type intersection; -# endif - metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque); - metalrt_intersect.assume_geometry_type( - metal::raytracing::geometry_type::triangle | - (kernel_data.bvh.have_curves ? metal::raytracing::geometry_type::curve : - metal::raytracing::geometry_type::none) | - (kernel_data.bvh.have_points ? metal::raytracing::geometry_type::bounding_box : - metal::raytracing::geometry_type::none)); - - // if we know we are going to get max one hit, like for random-sss-walk we can - // optimize and accept the first hit - if (max_hits == 1) { - metalrt_intersect.accept_any_intersection(true); - } - -# if defined(__METALRT_MOTION__) - intersection = metalrt_intersect.intersect( - r, metal_ancillaries->accel_struct, ~0, ray->time, metal_ancillaries->ift_local, payload); -# else if (!(kernel_data_fetch(object_flag, local_object) & SD_OBJECT_TRANSFORM_APPLIED)) { - // transform the ray into object's local space + /* Transform the ray into object's local space. */ Transform itfm = kernel_data_fetch(objects, local_object).itfm; r.origin = transform_point(&itfm, r.origin); r.direction = transform_direction(&itfm, r.direction); } - - intersection = metalrt_intersect.intersect( - r, - metal_ancillaries->blas_accel_structs[local_object].blas, - metal_ancillaries->ift_local_prim, - payload); # endif - if (lcg_state) { - *lcg_state = payload.lcg_state; - } - if (local_isect) { - *local_isect = payload.local_isect; - } + metalrt_intersect.assume_geometry_type(metal::raytracing::geometry_type::triangle); - return payload.result; + if (single_hit) { + MetalRTIntersectionLocalPayload_single_hit payload; + payload.self_prim = ray->self.prim - primitive_id_offset; + + /* We know we are going to get max one hit, so we can optimize and accept the first hit. */ + metalrt_intersect.accept_any_intersection(true); + + /* We only need custom intersection filtering (i.e. non_opaque) if we are performing a + * self-primitive intersection check. */ + metalrt_intersect.force_opacity((ray->self.prim == PRIM_NONE) ? + metal::raytracing::forced_opacity::opaque : + metal::raytracing::forced_opacity::non_opaque); + +# if defined(__METALRT_MOTION__) + intersection = metalrt_intersect.intersect(r, + metal_ancillaries->accel_struct, + ~0, + ray->time, + metal_ancillaries->ift_local_single_hit_mblur, + payload); +# else + intersection = metalrt_intersect.intersect( + r, + metal_ancillaries->blas_accel_structs[local_object].blas, + metal_ancillaries->ift_local_single_hit, + payload); +# endif + + if (intersection.type == intersection_type::none) { + local_isect->num_hits = 0; + return false; + } + + uint prim = intersection.primitive_id + primitive_id_offset; + int prim_type = kernel_data_fetch(objects, local_object).primitive_type; + + local_isect->num_hits = 1; + local_isect->hits[0].prim = prim; + local_isect->hits[0].type = prim_type; + local_isect->hits[0].object = local_object; + local_isect->hits[0].u = intersection.triangle_barycentric_coord.x; + local_isect->hits[0].v = intersection.triangle_barycentric_coord.y; + local_isect->hits[0].t = intersection.distance; + + const packed_uint3 tri_vindex = kernel_data_fetch(tri_vindex, prim); + const float3 tri_a = float3(kernel_data_fetch(tri_verts, tri_vindex.x)); + const float3 tri_b = float3(kernel_data_fetch(tri_verts, tri_vindex.y)); + const float3 tri_c = float3(kernel_data_fetch(tri_verts, tri_vindex.z)); + local_isect->Ng[0] = normalize(cross(tri_b - tri_a, tri_c - tri_a)); + return true; + } + else { + MetalRTIntersectionLocalPayload payload; + payload.self_prim = ray->self.prim - primitive_id_offset; + payload.max_hits = max_hits; + payload.num_hits = 0; + if (lcg_state) { + payload.has_lcg_state = 1; + payload.lcg_state = *lcg_state; + } + + metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque); + +# if defined(__METALRT_MOTION__) + intersection = metalrt_intersect.intersect(r, + metal_ancillaries->accel_struct, + ~0, + ray->time, + metal_ancillaries->ift_local_mblur, + payload); +# else + intersection = metalrt_intersect.intersect( + r, + metal_ancillaries->blas_accel_structs[local_object].blas, + metal_ancillaries->ift_local, + payload); +# endif + + if (max_hits == 0) { + /* Special case for when no hit information is requested, just report that something was hit + */ + return (intersection.type != intersection_type::none); + } + + if (lcg_state) { + *lcg_state = payload.lcg_state; + } + + const int num_hits = payload.num_hits; + if (local_isect) { + + /* Record geometric normal */ + int prim_type = kernel_data_fetch(objects, local_object).primitive_type; + + local_isect->num_hits = num_hits; + for (int hit = 0; hit < num_hits; hit++) { + uint prim = payload.hit_prim[hit] + primitive_id_offset; + local_isect->hits[hit].prim = prim; + local_isect->hits[hit].t = payload.hit_t[hit]; + local_isect->hits[hit].u = payload.hit_u[hit]; + local_isect->hits[hit].v = payload.hit_v[hit]; + local_isect->hits[hit].object = local_object; + local_isect->hits[hit].type = prim_type; + + const packed_uint3 tri_vindex = kernel_data_fetch(tri_vindex, prim); + const float3 tri_a = float3(kernel_data_fetch(tri_verts, tri_vindex.x)); + const float3 tri_b = float3(kernel_data_fetch(tri_verts, tri_vindex.y)); + const float3 tri_c = float3(kernel_data_fetch(tri_verts, tri_vindex.z)); + local_isect->Ng[hit] = normalize(cross(tri_b - tri_a, tri_c - tri_a)); + } + } + return num_hits > 0; + } } #endif @@ -385,22 +451,6 @@ ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg, ccl_private uint *num_recorded_hits, ccl_private float *throughput) { - if (!intersection_ray_valid(ray)) { - return false; - } - -# if defined(WITH_CYCLES_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, ray->tmin, ray->tmax); metalrt_intersector_type metalrt_intersect; metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque); @@ -411,9 +461,8 @@ ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg, (kernel_data.bvh.have_points ? metal::raytracing::geometry_type::bounding_box : metal::raytracing::geometry_type::none)); - MetalRTIntersectionShadowPayload payload; + MetalRTIntersectionShadowAllPayload payload; payload.self = ray->self; - payload.visibility = visibility; payload.max_hits = max_hits; payload.num_hits = 0; payload.num_recorded_hits = 0; @@ -428,11 +477,11 @@ ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg, metal_ancillaries->accel_struct, visibility, ray->time, - metal_ancillaries->ift_shadow, + metal_ancillaries->ift_shadow_all, payload); # else intersection = metalrt_intersect.intersect( - r, metal_ancillaries->accel_struct, visibility, metal_ancillaries->ift_shadow, payload); + r, metal_ancillaries->accel_struct, visibility, metal_ancillaries->ift_shadow_all, payload); # endif *num_recorded_hits = payload.num_recorded_hits; @@ -448,25 +497,11 @@ ccl_device_intersect bool scene_intersect_volume(KernelGlobals kg, ccl_private Intersection *isect, const uint visibility) { - if (!intersection_ray_valid(ray)) { - return false; - } - -# if defined(WITH_CYCLES_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_volume)) { - kernel_assert(!"Invalid ift_volume"); - return false; - } -# endif - metal::raytracing::ray r(ray->P, ray->D, ray->tmin, ray->tmax); metalrt_intersector_type metalrt_intersect; metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque); + metalrt_intersect.set_geometry_cull_mode(metal::raytracing::geometry_cull_mode::bounding_box | + metal::raytracing::geometry_cull_mode::curve); metalrt_intersect.assume_geometry_type( metal::raytracing::geometry_type::triangle | (kernel_data.bvh.have_curves ? metal::raytracing::geometry_type::curve : @@ -474,9 +509,8 @@ ccl_device_intersect bool scene_intersect_volume(KernelGlobals kg, (kernel_data.bvh.have_points ? metal::raytracing::geometry_type::bounding_box : metal::raytracing::geometry_type::none)); - MetalRTIntersectionPayload payload; + MetalRTIntersectionShadowPayload payload; payload.self = ray->self; - payload.visibility = visibility; typename metalrt_intersector_type::result_type intersection; @@ -492,78 +526,16 @@ ccl_device_intersect bool scene_intersect_volume(KernelGlobals kg, r, metal_ancillaries->accel_struct, visibility, metal_ancillaries->ift_volume, payload); # endif - if (intersection.type == intersection_type::none) { - return false; - } - else if (intersection.type == intersection_type::triangle) { + if (intersection.type == intersection_type::triangle) { isect->prim = intersection.primitive_id + intersection.user_instance_id; isect->type = kernel_data_fetch(objects, intersection.instance_id).primitive_type; isect->u = intersection.triangle_barycentric_coord.x; isect->v = intersection.triangle_barycentric_coord.y; isect->object = intersection.instance_id; isect->t = intersection.distance; + return true; } -# ifdef __HAIR__ - else if (kernel_data.bvh.have_curves && intersection.type == intersection_type::curve) { - int prim = intersection.primitive_id + intersection.user_instance_id; - const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim); - isect->prim = segment.prim; - isect->type = segment.type; - isect->u = intersection.curve_parameter; - - if (segment.type & PRIMITIVE_CURVE_RIBBON) { - isect->v = curve_ribbon_v(kg, - intersection.curve_parameter, - intersection.distance, - ray, - intersection.instance_id, - segment.prim, - segment.type); - } - else { - isect->v = 0.0f; - } - } -# endif -# ifdef __POINTCLOUD__ - else if (kernel_data.bvh.have_points && intersection.type == intersection_type::bounding_box) { - const int object = intersection.instance_id; - const uint prim = intersection.primitive_id + intersection.user_instance_id; - const int prim_type = kernel_data_fetch(objects, intersection.instance_id).primitive_type; - - isect->object = object; - - if (!(kernel_data_fetch(object_flag, object) & SD_OBJECT_TRANSFORM_APPLIED)) { - float3 idir; -# if defined(__METALRT_MOTION__) - bvh_instance_motion_push(NULL, object, ray, &r.origin, &r.direction, &idir); -# else - bvh_instance_push(NULL, object, ray, &r.origin, &r.direction, &idir); -# endif - } - - if (prim_type & PRIMITIVE_POINT) { - if (!point_intersect(NULL, - isect, - r.origin, - r.direction, - ray->tmin, - ray->tmax, - intersection.instance_id, - prim, - ray->time, - prim_type)) - { - /* Shouldn't get here */ - kernel_assert(!"Intersection mismatch"); - return false; - } - return true; - } - } -# endif - - return true; + return false; } #endif diff --git a/intern/cycles/kernel/device/metal/compat.h b/intern/cycles/kernel/device/metal/compat.h index 35a90642acf..6370dfe6aa2 100644 --- a/intern/cycles/kernel/device/metal/compat.h +++ b/intern/cycles/kernel/device/metal/compat.h @@ -359,9 +359,12 @@ struct MetalAncillaries { metalrt_as_type accel_struct; metalrt_ift_type ift_default; metalrt_ift_type ift_shadow; + metalrt_ift_type ift_shadow_all; metalrt_ift_type ift_volume; - metalrt_ift_type ift_local; - metalrt_blas_ift_type ift_local_prim; + metalrt_blas_ift_type ift_local; + metalrt_ift_type ift_local_mblur; + metalrt_blas_ift_type ift_local_single_hit; + metalrt_ift_type ift_local_single_hit_mblur; constant MetalRTBlasWrapper *blas_accel_structs; #endif }; diff --git a/intern/cycles/kernel/device/metal/kernel.metal b/intern/cycles/kernel/device/metal/kernel.metal index 906f9839f98..073573d07b0 100644 --- a/intern/cycles/kernel/device/metal/kernel.metal +++ b/intern/cycles/kernel/device/metal/kernel.metal @@ -36,10 +36,34 @@ enum { METALRT_HIT_TRIANGLE, METALRT_HIT_CURVE, METALRT_HIT_BOUNDING_BOX }; /* Hit functions. */ +[[intersection(triangle, triangle_data, curve_data)]] PrimitiveIntersectionResult +__intersection__local_tri_single_hit( + ray_data MetalKernelContext::MetalRTIntersectionLocalPayload_single_hit &payload [[payload]], + uint primitive_id [[primitive_id]]) +{ + PrimitiveIntersectionResult result; + result.continue_search = true; + result.accept = (payload.self_prim != primitive_id); + return result; +} + +[[intersection(triangle, + triangle_data, + curve_data, + METALRT_TAGS, + extended_limits)]] PrimitiveIntersectionResult +__intersection__local_tri_single_hit_mblur( + ray_data MetalKernelContext::MetalRTIntersectionLocalPayload_single_hit &payload [[payload]], + uint primitive_id [[primitive_id]]) +{ + PrimitiveIntersectionResult result; + result.continue_search = true; + result.accept = (payload.self_prim != primitive_id); + return result; +} + template -TReturn metalrt_local_hit(constant KernelParamsMetal &launch_params_metal, - ray_data MetalKernelContext::MetalRTIntersectionLocalPayload &payload, - const uint object, +TReturn metalrt_local_hit(ray_data MetalKernelContext::MetalRTIntersectionLocalPayload &payload, const uint prim, const float2 barycentrics, const float ray_tmax) @@ -47,10 +71,7 @@ TReturn metalrt_local_hit(constant KernelParamsMetal &launch_params_metal, TReturn result; # ifdef __BVH_LOCAL__ - MetalKernelContext context(launch_params_metal); - - if ((object != payload.local_object) || context.intersection_skip_self_local(payload.self, prim)) - { + if (payload.self_prim == prim) { /* Only intersect with matching object and skip self-intersection. */ result.accept = false; result.continue_search = true; @@ -60,7 +81,6 @@ TReturn metalrt_local_hit(constant KernelParamsMetal &launch_params_metal, const short max_hits = payload.max_hits; if (max_hits == 0) { /* Special case for when no hit information is requested, just report that something was hit */ - payload.result = true; result.accept = true; result.continue_search = false; return result; @@ -68,18 +88,20 @@ TReturn metalrt_local_hit(constant KernelParamsMetal &launch_params_metal, int hit = 0; if (payload.has_lcg_state) { - for (short i = min(max_hits, short(payload.local_isect.num_hits)) - 1; i >= 0; --i) { - if (ray_tmax == payload.local_isect.hits[i].t) { + for (short i = min(max_hits, short(payload.num_hits)) - 1; i >= 0; --i) { + if (ray_tmax == payload.hit_t[i]) { result.accept = false; result.continue_search = true; return result; } } - hit = payload.local_isect.num_hits++; - - if (payload.local_isect.num_hits > max_hits) { - hit = lcg_step_uint(&payload.lcg_state) % payload.local_isect.num_hits; + hit = payload.num_hits; + if (hit < max_hits) { + payload.num_hits++; + } + else { + hit = lcg_step_uint(&payload.lcg_state) % payload.num_hits; if (hit >= max_hits) { result.accept = false; result.continue_search = true; @@ -88,7 +110,7 @@ TReturn metalrt_local_hit(constant KernelParamsMetal &launch_params_metal, } } else { - if (payload.local_isect.num_hits && ray_tmax > payload.local_isect.hits[0].t) { + if (payload.num_hits && ray_tmax > payload.hit_t[0]) { /* Record closest intersection only. Do not terminate ray here, since there is no guarantee * about distance ordering in any-hit */ result.accept = false; @@ -96,107 +118,61 @@ TReturn metalrt_local_hit(constant KernelParamsMetal &launch_params_metal, return result; } - payload.local_isect.num_hits = 1; + payload.num_hits = 1; } - ray_data Intersection *isect = &payload.local_isect.hits[hit]; - isect->t = ray_tmax; - isect->prim = prim; - isect->object = object; - isect->type = kernel_data_fetch(objects, object).primitive_type; - - isect->u = barycentrics.x; - isect->v = barycentrics.y; - - /* Record geometric normal */ - const packed_uint3 tri_vindex = kernel_data_fetch(tri_vindex, isect->prim); - const float3 tri_a = float3(kernel_data_fetch(tri_verts, tri_vindex.x)); - const float3 tri_b = float3(kernel_data_fetch(tri_verts, tri_vindex.y)); - const float3 tri_c = float3(kernel_data_fetch(tri_verts, tri_vindex.z)); - payload.local_isect.Ng[hit] = normalize(cross(tri_b - tri_a, tri_c - tri_a)); + payload.hit_prim[hit] = prim; + payload.hit_t[hit] = ray_tmax; + payload.hit_u[hit] = barycentrics.x; + payload.hit_v[hit] = barycentrics.y; /* Continue tracing (without this the trace call would return after the first hit) */ result.accept = false; result.continue_search = true; - return result; # endif + return result; } [[intersection(triangle, triangle_data, curve_data)]] PrimitiveIntersectionResult -__anyhit__cycles_metalrt_local_hit_tri_prim( - constant KernelParamsMetal &launch_params_metal [[buffer(1)]], - ray_data MetalKernelContext::MetalRTIntersectionLocalPayload &payload [[payload]], - uint primitive_id [[primitive_id]], - float2 barycentrics [[barycentric_coord]], - float ray_tmax [[distance]]) +__intersection__local_tri(ray_data MetalKernelContext::MetalRTIntersectionLocalPayload &payload + [[payload]], + uint primitive_id [[primitive_id]], + float2 barycentrics [[barycentric_coord]], + float ray_tmax [[distance]]) { - uint prim = primitive_id + kernel_data_fetch(object_prim_offset, payload.local_object); - /* instance_id, aka the user_id has been removed. If we take this function we optimized the * SSS for starting traversal from a primitive acceleration structure instead of the root of the * global AS. this means we will always be intersecting the correct object no need for the * user-id to check */ return metalrt_local_hit( - launch_params_metal, payload, payload.local_object, prim, barycentrics, ray_tmax); + payload, primitive_id, barycentrics, ray_tmax); } + [[intersection(triangle, triangle_data, curve_data, METALRT_TAGS, extended_limits)]] PrimitiveIntersectionResult -__anyhit__cycles_metalrt_local_hit_tri( - constant KernelParamsMetal &launch_params_metal [[buffer(1)]], +__intersection__local_tri_mblur( ray_data MetalKernelContext::MetalRTIntersectionLocalPayload &payload [[payload]], - uint instance_id [[instance_id]], uint primitive_id [[primitive_id]], - uint primitive_id_offset [[user_instance_id]], float2 barycentrics [[barycentric_coord]], float ray_tmax [[distance]]) { return metalrt_local_hit( - launch_params_metal, - payload, - instance_id, - primitive_id + primitive_id_offset, - barycentrics, - ray_tmax); -} - -[[intersection(bounding_box, - triangle_data, - curve_data, - METALRT_TAGS, - extended_limits)]] BoundingBoxIntersectionResult -__anyhit__cycles_metalrt_local_hit_box(const float ray_tmax [[max_distance]]) -{ - /* unused function */ - BoundingBoxIntersectionResult result; - result.distance = ray_tmax; - result.accept = false; - result.continue_search = false; - return result; -} - -[[intersection(bounding_box, triangle_data, curve_data)]] BoundingBoxIntersectionResult -__anyhit__cycles_metalrt_local_hit_box_prim(const float ray_tmax [[max_distance]]) -{ - /* unused function */ - BoundingBoxIntersectionResult result; - result.distance = ray_tmax; - result.accept = false; - result.continue_search = false; - return result; + payload, primitive_id, barycentrics, ray_tmax); } template -bool metalrt_shadow_all_hit(constant KernelParamsMetal &launch_params_metal, - ray_data MetalKernelContext::MetalRTIntersectionShadowPayload &payload, - uint object, - uint prim, - const float2 barycentrics, - const float ray_tmax, - const float t = 0.0f, - ccl_private const Ray *ray = NULL) +bool metalrt_shadow_all_hit( + constant KernelParamsMetal &launch_params_metal, + ray_data MetalKernelContext::MetalRTIntersectionShadowAllPayload &payload, + uint object, + uint prim, + const float2 barycentrics, + const float ray_tmax, + const float t = 0.0f, + ccl_private const Ray *ray = NULL) { # ifdef __SHADOW_RECORD_ALL__ float u = barycentrics.x; @@ -338,9 +314,9 @@ bool metalrt_shadow_all_hit(constant KernelParamsMetal &launch_params_metal, curve_data, METALRT_TAGS, extended_limits)]] PrimitiveIntersectionResult -__anyhit__cycles_metalrt_shadow_all_hit_tri( +__intersection__tri_shadow_all( constant KernelParamsMetal &launch_params_metal [[buffer(1)]], - ray_data MetalKernelContext::MetalRTIntersectionShadowPayload &payload [[payload]], + ray_data MetalKernelContext::MetalRTIntersectionShadowAllPayload &payload [[payload]], const unsigned int object [[instance_id]], const unsigned int primitive_id [[primitive_id]], const uint primitive_id_offset [[user_instance_id]], @@ -356,43 +332,21 @@ __anyhit__cycles_metalrt_shadow_all_hit_tri( return result; } -[[intersection(bounding_box, - triangle_data, - curve_data, - METALRT_TAGS, - extended_limits)]] BoundingBoxIntersectionResult -__anyhit__cycles_metalrt_shadow_all_hit_box(const float ray_tmax [[max_distance]]) -{ - /* unused function */ - BoundingBoxIntersectionResult result; - result.distance = ray_tmax; - result.accept = false; - result.continue_search = false; - return result; -} - [[intersection(triangle, triangle_data, curve_data, METALRT_TAGS, extended_limits)]] PrimitiveIntersectionResult -__anyhit__cycles_metalrt_volume_test_tri( - constant KernelParamsMetal &launch_params_metal [[buffer(1)]], - ray_data MetalKernelContext::MetalRTIntersectionPayload &payload [[payload]], - const unsigned int object [[instance_id]], - const unsigned int primitive_id [[primitive_id]], - const uint primitive_id_offset [[user_instance_id]]) +__intersection__volume_tri(constant KernelParamsMetal &launch_params_metal [[buffer(1)]], + ray_data MetalKernelContext::MetalRTIntersectionShadowPayload &payload + [[payload]], + const unsigned int object [[instance_id]], + const unsigned int primitive_id [[primitive_id]], + const uint primitive_id_offset [[user_instance_id]]) { PrimitiveIntersectionResult result; result.continue_search = true; -# ifdef __VISIBILITY_FLAG__ - if ((kernel_data_fetch(objects, object).visibility & payload.visibility) == 0) { - result.accept = false; - return result; - } -# endif - if ((kernel_data_fetch(object_flag, object) & SD_OBJECT_HAS_VOLUME) == 0) { result.accept = false; return result; @@ -409,21 +363,6 @@ __anyhit__cycles_metalrt_volume_test_tri( return result; } -[[intersection(bounding_box, - triangle_data, - curve_data, - METALRT_TAGS, - extended_limits)]] BoundingBoxIntersectionResult -__anyhit__cycles_metalrt_volume_test_box(const float ray_tmax [[max_distance]]) -{ - /* unused function */ - BoundingBoxIntersectionResult result; - result.distance = ray_tmax; - result.accept = false; - result.continue_search = false; - return result; -} - template inline TReturnType metalrt_visibility_test( constant KernelParamsMetal &launch_params_metal, @@ -460,37 +399,72 @@ inline TReturnType metalrt_visibility_test( } # endif - uint visibility = payload.visibility; + if (payload.self_object == object && payload.self_prim == prim) { + result.accept = false; + result.continue_search = true; + return result; + } + result.accept = true; + result.continue_search = true; + return result; +} + +template +inline TReturnType metalrt_visibility_test_shadow( + constant KernelParamsMetal &launch_params_metal, + ray_data MetalKernelContext::MetalRTIntersectionShadowPayload &payload, + const uint object, + uint prim, + const float u, + const float t = 0.0f, + ccl_private const Ray *ray = NULL) +{ + TReturnType result; + +# ifdef __HAIR__ + if constexpr (intersection_type == METALRT_HIT_CURVE) { + /* Filter out curve end-caps. */ + if (u == 0.0f || u == 1.0f) { + result.accept = false; + result.continue_search = true; + return result; + } + + const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim); + int type = segment.type; + prim = segment.prim; + + if (type & PRIMITIVE_CURVE_RIBBON) { + MetalKernelContext context(launch_params_metal); + if (!context.curve_ribbon_accept(NULL, u, t, ray, object, prim, type)) { + result.accept = false; + result.continue_search = true; + return result; + } + } + } +# endif MetalKernelContext context(launch_params_metal); /* Shadow ray early termination. */ - if (visibility & PATH_RAY_SHADOW_OPAQUE) { # ifdef __SHADOW_LINKING__ - if (context.intersection_skip_shadow_link(nullptr, payload.self, object)) { - result.accept = false; - result.continue_search = true; - return result; - } + if (context.intersection_skip_shadow_link(nullptr, payload.self, object)) { + result.accept = false; + result.continue_search = true; + return result; + } # endif - if (context.intersection_skip_self_shadow(payload.self, object, prim)) { - result.accept = false; - result.continue_search = true; - return result; - } - else { - result.accept = true; - result.continue_search = false; - return result; - } + if (context.intersection_skip_self_shadow(payload.self, object, prim)) { + result.accept = false; + result.continue_search = true; + return result; } else { - if (context.intersection_skip_self(payload.self, object, prim)) { - result.accept = false; - result.continue_search = true; - return result; - } + result.accept = true; + result.continue_search = false; + return result; } result.accept = true; @@ -503,32 +477,35 @@ inline TReturnType metalrt_visibility_test( curve_data, METALRT_TAGS, extended_limits)]] PrimitiveIntersectionResult -__anyhit__cycles_metalrt_visibility_test_tri( - constant KernelParamsMetal &launch_params_metal [[buffer(1)]], - ray_data MetalKernelContext::MetalRTIntersectionPayload &payload [[payload]], - const unsigned int object [[instance_id]], - const uint primitive_id_offset [[user_instance_id]], - const unsigned int primitive_id [[primitive_id]]) +__intersection__tri(constant KernelParamsMetal &launch_params_metal [[buffer(1)]], + ray_data MetalKernelContext::MetalRTIntersectionPayload &payload [[payload]], + const unsigned int object [[instance_id]], + const uint primitive_id_offset [[user_instance_id]], + const unsigned int primitive_id [[primitive_id]]) { - uint prim = primitive_id + primitive_id_offset; - PrimitiveIntersectionResult result = - metalrt_visibility_test( - launch_params_metal, payload, object, prim, 0.0f); + PrimitiveIntersectionResult result; + result.continue_search = true; + result.accept = (payload.self_object != object || + payload.self_prim != (primitive_id + primitive_id_offset)); return result; } -[[intersection(bounding_box, +[[intersection(triangle, triangle_data, curve_data, METALRT_TAGS, - extended_limits)]] BoundingBoxIntersectionResult -__anyhit__cycles_metalrt_visibility_test_box(const float ray_tmax [[max_distance]]) + extended_limits)]] PrimitiveIntersectionResult +__intersection__tri_shadow(constant KernelParamsMetal &launch_params_metal [[buffer(1)]], + ray_data MetalKernelContext::MetalRTIntersectionShadowPayload &payload + [[payload]], + const unsigned int object [[instance_id]], + const uint primitive_id_offset [[user_instance_id]], + const unsigned int primitive_id [[primitive_id]]) { - /* Unused function */ - BoundingBoxIntersectionResult result; - result.accept = false; - result.continue_search = true; - result.distance = ray_tmax; + uint prim = primitive_id + primitive_id_offset; + PrimitiveIntersectionResult result = + metalrt_visibility_test_shadow( + launch_params_metal, payload, object, prim, 0.0f); return result; } @@ -577,15 +554,51 @@ __intersection__curve_shadow(constant KernelParamsMetal &launch_params_metal [[b const uint object [[instance_id]], const uint primitive_id [[primitive_id]], const uint primitive_id_offset [[user_instance_id]], + float distance [[distance]], const float3 ray_P [[origin]], const float3 ray_D [[direction]], float u [[curve_parameter]], - float t [[distance]], -# if defined(__METALRT_MOTION__) - const float time [[time]], -# endif const float ray_tmin [[min_distance]], - const float ray_tmax [[max_distance]]) + const float ray_tmax [[max_distance]] +# if defined(__METALRT_MOTION__) + , + const float time [[time]] +# endif +) +{ + uint prim = primitive_id + primitive_id_offset; + + Ray ray; + ray.P = ray_P; + ray.D = ray_D; +# if defined(__METALRT_MOTION__) + ray.time = time; +# endif + + PrimitiveIntersectionResult result = + metalrt_visibility_test_shadow( + launch_params_metal, payload, object, prim, u, distance, &ray); + + return result; +} + +[[intersection( + curve, triangle_data, curve_data, METALRT_TAGS, extended_limits)]] PrimitiveIntersectionResult +__intersection__curve_shadow_all( + constant KernelParamsMetal &launch_params_metal [[buffer(1)]], + ray_data MetalKernelContext::MetalRTIntersectionShadowAllPayload &payload [[payload]], + const uint object [[instance_id]], + const uint primitive_id [[primitive_id]], + const uint primitive_id_offset [[user_instance_id]], + const float3 ray_P [[origin]], + const float3 ray_D [[direction]], + float u [[curve_parameter]], + float t [[distance]], +# if defined(__METALRT_MOTION__) + const float time [[time]], +# endif + const float ray_tmin [[min_distance]], + const float ray_tmax [[max_distance]]) { uint prim = primitive_id + primitive_id_offset; @@ -606,37 +619,9 @@ __intersection__curve_shadow(constant KernelParamsMetal &launch_params_metal [[b } # ifdef __POINTCLOUD__ -ccl_device_inline void metalrt_intersection_point( +ccl_device_inline void metalrt_intersection_point_shadow_all( constant KernelParamsMetal &launch_params_metal, - ray_data MetalKernelContext::MetalRTIntersectionPayload &payload, - const uint object, - const uint prim, - const uint type, - const float3 ray_P, - const float3 ray_D, - float time, - const float ray_tmin, - const float ray_tmax, - thread BoundingBoxIntersectionResult &result) -{ - Intersection isect; - isect.t = ray_tmax; - - MetalKernelContext context(launch_params_metal); - if (context.point_intersect( - NULL, &isect, ray_P, ray_D, ray_tmin, isect.t, object, prim, time, type)) - { - result = metalrt_visibility_test( - launch_params_metal, payload, object, prim, isect.u); - if (result.accept) { - result.distance = isect.t; - } - } -} - -ccl_device_inline void metalrt_intersection_point_shadow( - constant KernelParamsMetal &launch_params_metal, - ray_data MetalKernelContext::MetalRTIntersectionShadowPayload &payload, + ray_data MetalKernelContext::MetalRTIntersectionShadowAllPayload &payload, const uint object, const uint prim, const uint type, @@ -664,8 +649,6 @@ ccl_device_inline void metalrt_intersection_point_shadow( } } -# endif /* __POINTCLOUD__ */ - [[intersection(bounding_box, triangle_data, curve_data, @@ -678,9 +661,9 @@ __intersection__point(constant KernelParamsMetal &launch_params_metal [[buffer(1 const uint primitive_id_offset [[user_instance_id]], const float3 ray_origin [[origin]], const float3 ray_direction [[direction]], -# if defined(__METALRT_MOTION__) +# if defined(__METALRT_MOTION__) const float time [[time]], -# endif +# endif const float ray_tmin [[min_distance]], const float ray_tmax [[max_distance]]) { @@ -692,29 +675,28 @@ __intersection__point(constant KernelParamsMetal &launch_params_metal [[buffer(1 result.continue_search = true; result.distance = ray_tmax; -# ifdef __POINTCLOUD__ + Intersection isect; + isect.t = ray_tmax; - metalrt_intersection_point(launch_params_metal, - payload, - object, - prim, - type, - ray_origin, - ray_direction, -# if defined(__METALRT_MOTION__) - time, -# else - 0.0f, +# ifndef __METALRT_MOTION__ + const float time = 0.0f; # endif - ray_tmin, - ray_tmax, - result); - -# endif /* __POINTCLOUD__ */ + MetalKernelContext context(launch_params_metal); + if (context.point_intersect( + NULL, &isect, ray_origin, ray_direction, ray_tmin, isect.t, object, prim, time, type)) + { + result = metalrt_visibility_test( + launch_params_metal, payload, object, prim, isect.u); + if (result.accept) { + result.distance = isect.t; + } + } return result; } +# endif /* __POINTCLOUD__ */ + [[intersection(bounding_box, triangle_data, curve_data, @@ -744,21 +726,74 @@ __intersection__point_shadow(constant KernelParamsMetal &launch_params_metal [[b # ifdef __POINTCLOUD__ - metalrt_intersection_point_shadow(launch_params_metal, - payload, - object, - prim, - type, - ray_origin, - ray_direction, -# if defined(__METALRT_MOTION__) - time, -# else - 0.0f, + Intersection isect; + isect.t = ray_tmax; + +# ifndef __METALRT_MOTION__ + const float time = 0.0f; # endif - ray_tmin, - ray_tmax, - result); + + MetalKernelContext context(launch_params_metal); + if (context.point_intersect( + NULL, &isect, ray_origin, ray_direction, ray_tmin, isect.t, object, prim, time, type)) + { + result = + metalrt_visibility_test_shadow( + launch_params_metal, payload, object, prim, isect.u); + if (result.accept) { + result.distance = isect.t; + } + } + +# endif /* __POINTCLOUD__ */ + + return result; +} + +[[intersection(bounding_box, + triangle_data, + curve_data, + METALRT_TAGS, + extended_limits)]] BoundingBoxIntersectionResult +__intersection__point_shadow_all( + constant KernelParamsMetal &launch_params_metal [[buffer(1)]], + ray_data MetalKernelContext::MetalRTIntersectionShadowAllPayload &payload [[payload]], + const uint object [[instance_id]], + const uint primitive_id [[primitive_id]], + const uint primitive_id_offset [[user_instance_id]], + const float3 ray_origin [[origin]], + const float3 ray_direction [[direction]], +# if defined(__METALRT_MOTION__) + const float time [[time]], +# endif + const float ray_tmin [[min_distance]], + const float ray_tmax [[max_distance]]) +{ + const uint prim = primitive_id + primitive_id_offset; + const int type = kernel_data_fetch(objects, object).primitive_type; + + BoundingBoxIntersectionResult result; + result.accept = false; + result.continue_search = true; + result.distance = ray_tmax; + +# ifdef __POINTCLOUD__ + + metalrt_intersection_point_shadow_all(launch_params_metal, + payload, + object, + prim, + type, + ray_origin, + ray_direction, +# if defined(__METALRT_MOTION__) + time, +# else + 0.0f, +# endif + ray_tmin, + ray_tmax, + result); # endif /* __POINTCLOUD__ */ diff --git a/intern/cycles/kernel/device/optix/bvh.h b/intern/cycles/kernel/device/optix/bvh.h index 7e0cca78789..ce0ec9520e9 100644 --- a/intern/cycles/kernel/device/optix/bvh.h +++ b/intern/cycles/kernel/device/optix/bvh.h @@ -521,7 +521,16 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg, return p5 != PRIMITIVE_NONE; } +ccl_device_intersect bool scene_intersect_shadow(KernelGlobals kg, + ccl_private const Ray *ray, + const uint visibility) +{ + Intersection isect; + return scene_intersect(kg, ray, visibility, &isect); +} + #ifdef __BVH_LOCAL__ +template ccl_device_intersect bool scene_intersect_local(KernelGlobals kg, ccl_private const Ray *ray, ccl_private LocalIntersection *local_isect, diff --git a/intern/cycles/kernel/integrator/intersect_shadow.h b/intern/cycles/kernel/integrator/intersect_shadow.h index 546f7818ead..82f16cf8852 100644 --- a/intern/cycles/kernel/integrator/intersect_shadow.h +++ b/intern/cycles/kernel/integrator/intersect_shadow.h @@ -32,8 +32,7 @@ ccl_device bool integrate_intersect_shadow_opaque(KernelGlobals kg, constexpr const uint opaque_mask = SHADOW_CATCHER_VISIBILITY_SHIFT(PATH_RAY_SHADOW_OPAQUE) | PATH_RAY_SHADOW_OPAQUE; - Intersection isect; - const bool opaque_hit = scene_intersect(kg, ray, visibility & opaque_mask, &isect); + const bool opaque_hit = scene_intersect_shadow(kg, ray, visibility & opaque_mask); /* Only record the number of hits if nothing was hit, so that the shadow shading kernel does not * consider any intersections. There is no need to write anything to the state if the hit is diff --git a/intern/cycles/kernel/integrator/shade_surface.h b/intern/cycles/kernel/integrator/shade_surface.h index 9a094b8612d..185ec4a427d 100644 --- a/intern/cycles/kernel/integrator/shade_surface.h +++ b/intern/cycles/kernel/integrator/shade_surface.h @@ -59,6 +59,13 @@ ccl_device_forceinline float3 integrate_surface_ray_offset(KernelGlobals kg, * or dot(sd->Ng, ray_D) is small. Detect such cases and skip test? * - Instead of ray offset, can we tweak P to lie within the triangle? */ + +#ifndef __METALRT__ + /* MetalRT and Cycles triangle tests aren't numerically identical, meaning this method + * isn't robust for MetalRT. In this case, just applying the ray offset uniformly gives + * identical looking results. + */ + float3 verts[3]; if (sd->type == PRIMITIVE_TRIANGLE) { triangle_vertices(kg, sd->prim, verts); @@ -80,7 +87,9 @@ ccl_device_forceinline float3 integrate_surface_ray_offset(KernelGlobals kg, if (ray_triangle_intersect_self(local_ray_P, local_ray_D, verts)) { return ray_P; } - else { + else +#endif + { return ray_offset(ray_P, sd->Ng); } } diff --git a/intern/cycles/kernel/integrator/subsurface_random_walk.h b/intern/cycles/kernel/integrator/subsurface_random_walk.h index 4da558be4b0..d0dc7082aa9 100644 --- a/intern/cycles/kernel/integrator/subsurface_random_walk.h +++ b/intern/cycles/kernel/integrator/subsurface_random_walk.h @@ -359,7 +359,7 @@ ccl_device_inline bool subsurface_random_walk(KernelGlobals kg, ray.self.object = OBJECT_NONE; ray.self.prim = PRIM_NONE; } - scene_intersect_local(kg, &ray, &ss_isect, object, NULL, 1); + scene_intersect_local(kg, &ray, &ss_isect, object, NULL, 1); hit = (ss_isect.num_hits > 0); if (hit) { diff --git a/intern/cycles/kernel/svm/ao.h b/intern/cycles/kernel/svm/ao.h index cfa674a7a92..263aa081c01 100644 --- a/intern/cycles/kernel/svm/ao.h +++ b/intern/cycles/kernel/svm/ao.h @@ -77,8 +77,7 @@ ccl_device float svm_ao( } } else { - Intersection isect; - if (!scene_intersect(kg, &ray, PATH_RAY_SHADOW_OPAQUE, &isect)) { + if (!scene_intersect_shadow(kg, &ray, PATH_RAY_SHADOW_OPAQUE)) { unoccluded++; } }