kannonier8/blender
1
0
Fork 0
forked from bartvdbraak/blender
Code Pull Requests Activity

Cycles: MetalRT optimisations (scene_intersect_shadow + random_walk)

Browse Source 
This PR contains optimisations and a general tidy-up of the MetalRT backend.

- Currently `scene_intersect` is used for both normal and (opaque) shadow rays, however the usage patterns are different enough to warrant specialisation. Shadow intersection tests (flagged with `PATH_RAY_SHADOW_OPAQUE`) only need a bool result, but need a larger "self" payload in order to exclude hits against target lights. By specialising we can minimise the payload size in each case (which is helps performance) and avoid some dynamic branching. This PR introduces a new `scene_intersect_shadow` function which is specialised in Metal, and currently redirects to `scene_intersect` in the other backends.

- Currently `scene_intersect_local` is implemented for worst-case payload requirements as demanded by `subsurface_disk` (where `max_hits` is 4). The random_walk case only demands 1 hit result which we can retrieve directly from the intersector object (rather than stashing it in the payload). By specialising, we significantly reduce the payload size for random_walk queries, which has a big impact on performance. Additionally, we only need to use a custom intersection function for the first ray test in a random walk (for self-primitive filtering), so this PR forces faster `opaque` intersection testing for all but the first random walk test.

- Currently `scene_intersect_volume` has a lot of redundant code to handle non-triangle primitives despite volumes only being enclosed by trimeshes. This PR removes this code.

Additionally, this PR tidies up the convoluted intersection function linking code, removes some redundant intersection handlers, and uses more consistent naming of intersection functions.

On a M3 MacBook Pro, these changes give 2-3% performance increase on typical scenes with opaque trimesh materials (e.g. barbershop, classroom junkshop), but can give over 15% performance increase for certain scenes using random walk SSS (e.g. monster).

Pull Request: https://projects.blender.org/blender/blender/pulls/121397
This commit is contained in:
Michael Jones 2024-05-10 16:38:02 +02:00 committed by Michael Jones (Apple)
parent 71d92735bf
commit 5508b41a40
14 changed files with 598 additions and 614 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

15
intern/cycles/device/metal/device_impl.mm
View File

@ -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];

25
intern/cycles/device/metal/kernel.h
View File

@ -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<MTLFunction> make_intersection_function(const char *function_name);
string error_str;
API_AVAILABLE(macos(11.0))
id<MTLIntersectionFunctionTable> intersection_func_table[METALRT_TABLE_NUM] = {nil};
id<MTLFunction> rt_intersection_function[METALRT_FUNC_NUM] = {nil};
};
/* Cache of Metal kernels for each DeviceKernel. */

176
intern/cycles/device/metal/kernel.mm
View File

@ -459,6 +459,35 @@ static MTLFunctionConstantValues *GetConstantValues(KernelData const *data = nul
return constant_values;
}
id<MTLFunction> 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<MTLFunction> 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<MTLFunction> 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<MTLFunction> curve_intersect_default = nil;
id<MTLFunction> curve_intersect_shadow = nil;
id<MTLFunction> point_intersect_default = nil;
id<MTLFunction> 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<MTLFunction> f1, id<MTLFunction> 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<MTLFunction> f1, id<MTLFunction> 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;

22
intern/cycles/device/metal/queue.mm
View File

@ -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<MTLAccelerationStructure> 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

9
intern/cycles/kernel/bvh/bvh.h
View File

@ -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<bool single_hit = false>
ccl_device_intersect bool scene_intersect_local(KernelGlobals kg,
ccl_private const Ray *ray,
ccl_private LocalIntersection *local_isect,

9
intern/cycles/kernel/device/hiprt/bvh.h
View File

@ -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<bool single_hit = false>
ccl_device_intersect bool scene_intersect_local(KernelGlobals kg,
ccl_private const Ray *ray,
ccl_private LocalIntersection *local_isect,

406
intern/cycles/kernel/device/metal/bvh.h
View File

@ -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<bool single_hit = false>
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

7
intern/cycles/kernel/device/metal/compat.h
View File

@ -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
};

515
intern/cycles/kernel/device/metal/kernel.metal
View File

@ -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<typename TReturn, uint intersection_type>
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<PrimitiveIntersectionResult, METALRT_HIT_TRIANGLE>(
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<PrimitiveIntersectionResult, METALRT_HIT_TRIANGLE>(
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<uint intersection_type>
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<typename TReturnType, uint intersection_type>
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<typename TReturnType, uint intersection_type>
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<PrimitiveIntersectionResult, METALRT_HIT_TRIANGLE>(
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<PrimitiveIntersectionResult, METALRT_HIT_TRIANGLE>(
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<PrimitiveIntersectionResult, METALRT_HIT_CURVE>(
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<BoundingBoxIntersectionResult, METALRT_HIT_BOUNDING_BOX>(
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<BoundingBoxIntersectionResult, METALRT_HIT_BOUNDING_BOX>(
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<BoundingBoxIntersectionResult, METALRT_HIT_BOUNDING_BOX>(
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__ */

9
intern/cycles/kernel/device/optix/bvh.h
View File

@ -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<bool single_hit = false>
ccl_device_intersect bool scene_intersect_local(KernelGlobals kg,
ccl_private const Ray *ray,
ccl_private LocalIntersection *local_isect,

3
intern/cycles/kernel/integrator/intersect_shadow.h
View File

@ -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

11
intern/cycles/kernel/integrator/shade_surface.h
View File

@ -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);
}
}

2
intern/cycles/kernel/integrator/subsurface_random_walk.h
View File

@ -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<true>(kg, &ray, &ss_isect, object, NULL, 1);
hit = (ss_isect.num_hits > 0);
if (hit) {

3
intern/cycles/kernel/svm/ao.h
View File

@ -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++;
}
}