blender
/
blender
136
404
Fork 586
Code Issues 6.2k Pull Requests 763 Projects 19 Wiki Activity

Cycles: simplify handling of ray distance in GPU rendering 

All our intersections functions now work with unnormalized ray direction,
which means we no longer need to transform ray distance between world and
object space, they can all remain in world space.

There doesn't seem to be any real performance difference one way or the
other, but it does simplify the code.
This commit is contained in:
Brecht Van Lommel 2022-07-21 16:37:38 +02:00
parent 023eb2ea7c
commit 484ad31653
11 changed files with 153 additions and 393 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

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

@ -475,12 +475,7 @@ ccl_device_intersect bool scene_intersect_local(KernelGlobals kg,
float3 P = ray->P;
float3 dir = ray->D;
float3 idir = ray->D;
Transform ob_itfm;
rtc_ray.tfar = ray->tmax *
bvh_instance_motion_push(kg, local_object, ray, &P, &dir, &idir, &ob_itfm);
/* bvh_instance_motion_push() returns the inverse transform but
* it's not needed here. */
(void)ob_itfm;
bvh_instance_motion_push(kg, local_object, ray, &P, &dir, &idir);
rtc_ray.org_x = P.x;
rtc_ray.org_y = P.y;
@ -488,6 +483,8 @@ ccl_device_intersect bool scene_intersect_local(KernelGlobals kg,
rtc_ray.dir_x = dir.x;
rtc_ray.dir_y = dir.y;
rtc_ray.dir_z = dir.z;
rtc_ray.tnear = ray->tmin;
rtc_ray.tfar = ray->tmax;
RTCScene scene = (RTCScene)rtcGetGeometryUserData(geom);
kernel_assert(scene);
if (scene) {

8
intern/cycles/kernel/bvh/local.h
View File

@ -59,14 +59,10 @@ ccl_device_inline
const int object_flag = kernel_data_fetch(object_flag, local_object);
if (!(object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
#if BVH_FEATURE(BVH_MOTION)
Transform ob_itfm;
const float t_world_to_instance = bvh_instance_motion_push(
kg, local_object, ray, &P, &dir, &idir, &ob_itfm);
bvh_instance_motion_push(kg, local_object, ray, &P, &dir, &idir);
#else
const float t_world_to_instance = bvh_instance_push(kg, local_object, ray, &P, &dir, &idir);
bvh_instance_push(kg, local_object, ray, &P, &dir, &idir);
#endif
isect_t *= t_world_to_instance;
tmin *= t_world_to_instance;
object = local_object;
}

81
intern/cycles/kernel/bvh/shadow_all.h
View File

@ -53,23 +53,11 @@ ccl_device_inline
int object = OBJECT_NONE;
uint num_hits = 0;
#if BVH_FEATURE(BVH_MOTION)
Transform ob_itfm;
#endif
/* Max distance in world space. May be dynamically reduced when max number of
* recorded hits is exceeded and we no longer need to find hits beyond the max
* distance found. */
float t_max_world = ray->tmax;
/* Current maximum distance to the intersection.
* Is calculated as a ray length, transformed to an object space when entering
* instance node. */
float t_max_current = ray->tmax;
/* Conversion from world to local space for the current instance if any, 1.0
* otherwise. */
float t_world_to_instance = 1.0f;
const float tmax = ray->tmax;
float tmax_hits = tmax;
*r_num_recorded_hits = 0;
*r_throughput = 1.0f;
@ -90,7 +78,7 @@ ccl_device_inline
#endif
idir,
tmin,
t_max_current,
tmax,
node_addr,
visibility,
dist);
@ -158,16 +146,8 @@ ccl_device_inline
switch (type & PRIMITIVE_ALL) {
case PRIMITIVE_TRIANGLE: {
hit = triangle_intersect(kg,
&isect,
P,
dir,
tmin,
t_max_current,
visibility,
prim_object,
prim,
prim_addr);
hit = triangle_intersect(
kg, &isect, P, dir, tmin, tmax, visibility, prim_object, prim, prim_addr);
break;
}
#if BVH_FEATURE(BVH_MOTION)
@ -177,7 +157,7 @@ ccl_device_inline
P,
dir,
tmin,
t_max_current,
tmax,
ray->time,
visibility,
prim_object,
@ -200,16 +180,8 @@ ccl_device_inline
}
const int curve_type = kernel_data_fetch(prim_type, prim_addr);
hit = curve_intersect(kg,
&isect,
P,
dir,
tmin,
t_max_current,
prim_object,
prim,
ray->time,
curve_type);
hit = curve_intersect(
kg, &isect, P, dir, tmin, tmax, prim_object, prim, ray->time, curve_type);
break;
}
@ -226,16 +198,8 @@ ccl_device_inline
}
const int point_type = kernel_data_fetch(prim_type, prim_addr);
hit = point_intersect(kg,
&isect,
P,
dir,
tmin,
t_max_current,
prim_object,
prim,
ray->time,
point_type);
hit = point_intersect(
kg, &isect, P, dir, tmin, tmax, prim_object, prim, ray->time, point_type);
break;
}
#endif /* BVH_FEATURE(BVH_POINTCLOUD) */
@ -247,9 +211,6 @@ ccl_device_inline
/* shadow ray early termination */
if (hit) {
/* Convert intersection distance to world space. */
isect.t /= t_world_to_instance;
/* detect if this surface has a shader with transparent shadows */
/* todo: optimize so primitive visibility flag indicates if
* the primitive has a transparent shadow shader? */
@ -281,7 +242,7 @@ ccl_device_inline
if (record_intersection) {
/* Test if we need to record this transparent intersection. */
const uint max_record_hits = min(max_hits, INTEGRATOR_SHADOW_ISECT_SIZE);
if (*r_num_recorded_hits < max_record_hits || isect.t < t_max_world) {
if (*r_num_recorded_hits < max_record_hits || isect.t < tmax_hits) {
/* If maximum number of hits was reached, replace the intersection with the
* highest distance. We want to find the N closest intersections. */
const uint num_recorded_hits = min(*r_num_recorded_hits, max_record_hits);
@ -303,7 +264,7 @@ ccl_device_inline
}
/* Limit the ray distance and stop counting hits beyond this. */
t_max_world = max(isect.t, max_t);
tmax_hits = max(isect.t, max_t);
}
integrator_state_write_shadow_isect(state, &isect, isect_index);
@ -321,16 +282,11 @@ ccl_device_inline
object = kernel_data_fetch(prim_object, -prim_addr - 1);
#if BVH_FEATURE(BVH_MOTION)
t_world_to_instance = bvh_instance_motion_push(
kg, object, ray, &P, &dir, &idir, &ob_itfm);
bvh_instance_motion_push(kg, object, ray, &P, &dir, &idir);
#else
t_world_to_instance = bvh_instance_push(kg, object, ray, &P, &dir, &idir);
bvh_instance_push(kg, object, ray, &P, &dir, &idir);
#endif
/* Convert intersection to object space. */
t_max_current *= t_world_to_instance;
tmin *= t_world_to_instance;
++stack_ptr;
kernel_assert(stack_ptr < BVH_STACK_SIZE);
traversal_stack[stack_ptr] = ENTRYPOINT_SENTINEL;
@ -345,17 +301,12 @@ ccl_device_inline
/* Instance pop. */
#if BVH_FEATURE(BVH_MOTION)
bvh_instance_motion_pop(kg, object, ray, &P, &dir, &idir, FLT_MAX, &ob_itfm);
bvh_instance_motion_pop(kg, object, ray, &P, &dir, &idir);
#else
bvh_instance_pop(kg, object, ray, &P, &dir, &idir, FLT_MAX);
bvh_instance_pop(kg, object, ray, &P, &dir, &idir);
#endif
/* Restore world space ray length. */
tmin = ray->tmin;
t_max_current = ray->tmax;
object = OBJECT_NONE;
t_world_to_instance = 1.0f;
node_addr = traversal_stack[stack_ptr];
--stack_ptr;
}

19
intern/cycles/kernel/bvh/traversal.h
View File

@ -43,13 +43,9 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg,
float3 P = ray->P;
float3 dir = bvh_clamp_direction(ray->D);
float3 idir = bvh_inverse_direction(dir);
float tmin = ray->tmin;
const float tmin = ray->tmin;
int object = OBJECT_NONE;
#if BVH_FEATURE(BVH_MOTION)
Transform ob_itfm;
#endif
isect->t = ray->tmax;
isect->u = 0.0f;
isect->v = 0.0f;
@ -223,15 +219,11 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg,
object = kernel_data_fetch(prim_object, -prim_addr - 1);
#if BVH_FEATURE(BVH_MOTION)
const float t_world_to_instance = bvh_instance_motion_push(
kg, object, ray, &P, &dir, &idir, &ob_itfm);
bvh_instance_motion_push(kg, object, ray, &P, &dir, &idir);
#else
const float t_world_to_instance = bvh_instance_push(kg, object, ray, &P, &dir, &idir);
bvh_instance_push(kg, object, ray, &P, &dir, &idir);
#endif
isect->t *= t_world_to_instance;
tmin *= t_world_to_instance;
++stack_ptr;
kernel_assert(stack_ptr < BVH_STACK_SIZE);
traversal_stack[stack_ptr] = ENTRYPOINT_SENTINEL;
@ -246,11 +238,10 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg,
/* instance pop */
#if BVH_FEATURE(BVH_MOTION)
isect->t = bvh_instance_motion_pop(kg, object, ray, &P, &dir, &idir, isect->t, &ob_itfm);
bvh_instance_motion_pop(kg, object, ray, &P, &dir, &idir);
#else
isect->t = bvh_instance_pop(kg, object, ray, &P, &dir, &idir, isect->t);
bvh_instance_pop(kg, object, ray, &P, &dir, &idir);
#endif
tmin = ray->tmin;
object = OBJECT_NONE;
node_addr = traversal_stack[stack_ptr];

20
intern/cycles/kernel/bvh/volume.h
View File

@ -46,13 +46,9 @@ ccl_device_inline
float3 P = ray->P;
float3 dir = bvh_clamp_direction(ray->D);
float3 idir = bvh_inverse_direction(dir);
float tmin = ray->tmin;
const float tmin = ray->tmin;
int object = OBJECT_NONE;
#if BVH_FEATURE(BVH_MOTION)
Transform ob_itfm;
#endif
isect->t = ray->tmax;
isect->u = 0.0f;
isect->v = 0.0f;
@ -189,15 +185,11 @@ ccl_device_inline
int object_flag = kernel_data_fetch(object_flag, object);
if (object_flag & SD_OBJECT_HAS_VOLUME) {
#if BVH_FEATURE(BVH_MOTION)
const float t_world_to_instance = bvh_instance_motion_push(
kg, object, ray, &P, &dir, &idir, &ob_itfm);
bvh_instance_motion_push(kg, object, ray, &P, &dir, &idir);
#else
const float t_world_to_instance = bvh_instance_push(kg, object, ray, &P, &dir, &idir);
bvh_instance_push(kg, object, ray, &P, &dir, &idir);
#endif
isect->t *= t_world_to_instance;
tmin *= t_world_to_instance;
++stack_ptr;
kernel_assert(stack_ptr < BVH_STACK_SIZE);
traversal_stack[stack_ptr] = ENTRYPOINT_SENTINEL;
@ -219,13 +211,11 @@ ccl_device_inline
/* instance pop */
#if BVH_FEATURE(BVH_MOTION)
isect->t = bvh_instance_motion_pop(kg, object, ray, &P, &dir, &idir, isect->t, &ob_itfm);
bvh_instance_motion_pop(kg, object, ray, &P, &dir, &idir);
#else
isect->t = bvh_instance_pop(kg, object, ray, &P, &dir, &idir, isect->t);
bvh_instance_pop(kg, object, ray, &P, &dir, &idir);
#endif
tmin = ray->tmin;
object = OBJECT_NONE;
node_addr = traversal_stack[stack_ptr];
--stack_ptr;

60
intern/cycles/kernel/bvh/volume_all.h
View File

@ -47,14 +47,10 @@ ccl_device_inline
float3 P = ray->P;
float3 dir = bvh_clamp_direction(ray->D);
float3 idir = bvh_inverse_direction(dir);
float tmin = ray->tmin;
const float tmin = ray->tmin;
int object = OBJECT_NONE;
float isect_t = ray->tmax;
#if BVH_FEATURE(BVH_MOTION)
Transform ob_itfm;
#endif
int num_hits_in_instance = 0;
uint num_hits = 0;
@ -159,18 +155,6 @@ ccl_device_inline
num_hits_in_instance++;
isect_array->t = isect_t;
if (num_hits == max_hits) {
if (object != OBJECT_NONE) {
#if BVH_FEATURE(BVH_MOTION)
float t_fac = 1.0f / len(transform_direction(&ob_itfm, dir));
#else
Transform itfm = object_fetch_transform(
kg, object, OBJECT_INVERSE_TRANSFORM);
float t_fac = 1.0f / len(transform_direction(&itfm, dir));
#endif
for (int i = 0; i < num_hits_in_instance; i++) {
(isect_array - i - 1)->t *= t_fac;
}
}
return num_hits;
}
}
@ -212,18 +196,6 @@ ccl_device_inline
num_hits_in_instance++;
isect_array->t = isect_t;
if (num_hits == max_hits) {
if (object != OBJECT_NONE) {
# if BVH_FEATURE(BVH_MOTION)
float t_fac = 1.0f / len(transform_direction(&ob_itfm, dir));
# else
Transform itfm = object_fetch_transform(
kg, object, OBJECT_INVERSE_TRANSFORM);
float t_fac = 1.0f / len(transform_direction(&itfm, dir));
# endif
for (int i = 0; i < num_hits_in_instance; i++) {
(isect_array - i - 1)->t *= t_fac;
}
}
return num_hits;
}
}
@ -242,15 +214,11 @@ ccl_device_inline
int object_flag = kernel_data_fetch(object_flag, object);
if (object_flag & SD_OBJECT_HAS_VOLUME) {
#if BVH_FEATURE(BVH_MOTION)
const float t_world_to_instance = bvh_instance_motion_push(
kg, object, ray, &P, &dir, &idir, &ob_itfm);
bvh_instance_motion_push(kg, object, ray, &P, &dir, &idir);
#else
const float t_world_to_instance = bvh_instance_push(kg, object, ray, &P, &dir, &idir);
bvh_instance_push(kg, object, ray, &P, &dir, &idir);
#endif
isect_t *= t_world_to_instance;
tmin *= t_world_to_instance;
num_hits_in_instance = 0;
isect_array->t = isect_t;
@ -274,29 +242,11 @@ ccl_device_inline
kernel_assert(object != OBJECT_NONE);
/* Instance pop. */
if (num_hits_in_instance) {
float t_fac;
#if BVH_FEATURE(BVH_MOTION)
bvh_instance_motion_pop_factor(kg, object, ray, &P, &dir, &idir, &t_fac, &ob_itfm);
bvh_instance_motion_pop(kg, object, ray, &P, &dir, &idir);
#else
bvh_instance_pop_factor(kg, object, ray, &P, &dir, &idir, &t_fac);
bvh_instance_pop(kg, object, ray, &P, &dir, &idir);
#endif
/* Scale isect->t to adjust for instancing. */
for (int i = 0; i < num_hits_in_instance; i++) {
(isect_array - i - 1)->t *= t_fac;
}
}
else {
#if BVH_FEATURE(BVH_MOTION)
bvh_instance_motion_pop(kg, object, ray, &P, &dir, &idir, FLT_MAX, &ob_itfm);
#else
bvh_instance_pop(kg, object, ray, &P, &dir, &idir, FLT_MAX);
#endif
}
tmin = ray->tmin;
isect_t = ray->tmax;
isect_array->t = isect_t;
object = OBJECT_NONE;
node_addr = traversal_stack[stack_ptr];

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

@ -407,8 +407,8 @@ void metalrt_intersection_curve(constant KernelParamsMetal &launch_params_metal,
const uint object,
const uint prim,
const uint type,
const float3 ray_origin,
const float3 ray_direction,
const float3 ray_P,
const float3 ray_D,
float time,
const float ray_tmin,
const float ray_tmax,
@ -421,25 +421,15 @@ void metalrt_intersection_curve(constant KernelParamsMetal &launch_params_metal,
}
# endif
float3 P = ray_origin;
float3 dir = ray_direction;
/* The direction is not normalized by default, but the curve intersection routine expects that */
float len;
dir = normalize_len(dir, &len);
Intersection isect;
isect.t = ray_tmax;
/* Transform maximum distance into object space. */
if (isect.t != FLT_MAX)
isect.t *= len;
MetalKernelContext context(launch_params_metal);
if (context.curve_intersect(NULL, &isect, P, dir, ray_tmin, isect.t, object, prim, time, type)) {
if (context.curve_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 / len;
result.distance = isect.t;
payload.u = isect.u;
payload.v = isect.v;
payload.prim = prim;
@ -454,8 +444,6 @@ void metalrt_intersection_curve_shadow(constant KernelParamsMetal &launch_params
const uint object,
const uint prim,
const uint type,
const float3 ray_origin,
const float3 ray_direction,
float time,
const float ray_tmin,
const float ray_tmax,
@ -463,28 +451,14 @@ void metalrt_intersection_curve_shadow(constant KernelParamsMetal &launch_params
{
const uint visibility = payload.visibility;
float3 P = ray_origin;
float3 dir = ray_direction;
/* The direction is not normalized by default, but the curve intersection routine expects that */
float len;
dir = normalize_len(dir, &len);
Intersection isect;
isect.t = ray_tmax;
/* Transform maximum distance into object space */
if (isect.t != FLT_MAX)
isect.t *= len;
MetalKernelContext context(launch_params_metal);
if (context.curve_intersect(NULL, &isect, P, dir, ray_tmin, isect.t, object, prim, time, type)) {
if (context.curve_intersect(NULL, &isect, ray_P, ray_D, ray_tmin, isect.t, object, prim, time, type)) {
result.continue_search = metalrt_shadow_all_hit<METALRT_HIT_BOUNDING_BOX>(
launch_params_metal, payload, object, prim, float2(isect.u, isect.v), ray_tmax);
result.accept = !result.continue_search;
if (result.accept) {
result.distance = isect.t / len;
}
}
}
@ -494,8 +468,8 @@ __intersection__curve_ribbon(constant KernelParamsMetal &launch_params_metal [[b
ray_data MetalKernelContext::MetalRTIntersectionPayload &payload [[payload]],
const uint object [[user_instance_id]],
const uint primitive_id [[primitive_id]],
const float3 ray_origin [[origin]],
const float3 ray_direction [[direction]],
const float3 ray_P [[origin]],
const float3 ray_D [[direction]],
const float ray_tmin [[min_distance]],
const float ray_tmax [[max_distance]])
{
@ -508,7 +482,7 @@ __intersection__curve_ribbon(constant KernelParamsMetal &launch_params_metal [[b
result.distance = ray_tmax;
if (segment.type & PRIMITIVE_CURVE_RIBBON) {
metalrt_intersection_curve(launch_params_metal, payload, object, segment.prim, segment.type, ray_origin, ray_direction,
metalrt_intersection_curve(launch_params_metal, payload, object, segment.prim, segment.type, ray_P, ray_D,
# if defined(__METALRT_MOTION__)
payload.time,
# else
@ -526,8 +500,8 @@ __intersection__curve_ribbon_shadow(constant KernelParamsMetal &launch_params_me
ray_data MetalKernelContext::MetalRTIntersectionShadowPayload &payload [[payload]],
const uint object [[user_instance_id]],
const uint primitive_id [[primitive_id]],
const float3 ray_origin [[origin]],
const float3 ray_direction [[direction]],
const float3 ray_P [[origin]],
const float3 ray_D [[direction]],
const float ray_tmin [[min_distance]],
const float ray_tmax [[max_distance]])
{
@ -540,7 +514,7 @@ __intersection__curve_ribbon_shadow(constant KernelParamsMetal &launch_params_me
result.distance = ray_tmax;
if (segment.type & PRIMITIVE_CURVE_RIBBON) {
metalrt_intersection_curve_shadow(launch_params_metal, payload, object, segment.prim, segment.type, ray_origin, ray_direction,
metalrt_intersection_curve_shadow(launch_params_metal, payload, object, segment.prim, segment.type, ray_P, ray_D,
# if defined(__METALRT_MOTION__)
payload.time,
# else
@ -558,8 +532,8 @@ __intersection__curve_all(constant KernelParamsMetal &launch_params_metal [[buff
ray_data MetalKernelContext::MetalRTIntersectionPayload &payload [[payload]],
const uint object [[user_instance_id]],
const uint primitive_id [[primitive_id]],
const float3 ray_origin [[origin]],
const float3 ray_direction [[direction]],
const float3 ray_P [[origin]],
const float3 ray_D [[direction]],
const float ray_tmin [[min_distance]],
const float ray_tmax [[max_distance]])
{
@ -570,7 +544,7 @@ __intersection__curve_all(constant KernelParamsMetal &launch_params_metal [[buff
result.accept = false;
result.continue_search = true;
result.distance = ray_tmax;
metalrt_intersection_curve(launch_params_metal, payload, object, segment.prim, segment.type, ray_origin, ray_direction,
metalrt_intersection_curve(launch_params_metal, payload, object, segment.prim, segment.type, ray_P, ray_D,
# if defined(__METALRT_MOTION__)
payload.time,
# else
@ -587,8 +561,8 @@ __intersection__curve_all_shadow(constant KernelParamsMetal &launch_params_metal
ray_data MetalKernelContext::MetalRTIntersectionShadowPayload &payload [[payload]],
const uint object [[user_instance_id]],
const uint primitive_id [[primitive_id]],
const float3 ray_origin [[origin]],
const float3 ray_direction [[direction]],
const float3 ray_P [[origin]],
const float3 ray_D [[direction]],
const float ray_tmin [[min_distance]],
const float ray_tmax [[max_distance]])
{
@ -600,7 +574,7 @@ __intersection__curve_all_shadow(constant KernelParamsMetal &launch_params_metal
result.continue_search = true;
result.distance = ray_tmax;
metalrt_intersection_curve_shadow(launch_params_metal, payload, object, segment.prim, segment.type, ray_origin, ray_direction,
metalrt_intersection_curve_shadow(launch_params_metal, payload, object, segment.prim, segment.type, ray_P, ray_D,
# if defined(__METALRT_MOTION__)
payload.time,
# else
@ -619,8 +593,8 @@ void metalrt_intersection_point(constant KernelParamsMetal &launch_params_metal,
const uint object,
const uint prim,
const uint type,
const float3 ray_origin,
const float3 ray_direction,
const float3 ray_P,
const float3 ray_D,
float time,
const float ray_tmin,
const float ray_tmax,
@ -633,25 +607,15 @@ void metalrt_intersection_point(constant KernelParamsMetal &launch_params_metal,
}
# endif
float3 P = ray_origin;
float3 dir = ray_direction;
/* The direction is not normalized by default, but the point intersection routine expects that */
float len;
dir = normalize_len(dir, &len);
Intersection isect;
isect.t = ray_tmax;
/* Transform maximum distance into object space. */
if (isect.t != FLT_MAX)
isect.t *= len;
MetalKernelContext context(launch_params_metal);
if (context.point_intersect(NULL, &isect, P, dir, ray_tmin, isect.t, object, prim, time, type)) {
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 / len;
result.distance = isect.t;
payload.u = isect.u;
payload.v = isect.v;
payload.prim = prim;
@ -666,8 +630,8 @@ void metalrt_intersection_point_shadow(constant KernelParamsMetal &launch_params
const uint object,
const uint prim,
const uint type,
const float3 ray_origin,
const float3 ray_direction,
const float3 ray_P,
const float3 ray_D,
float time,
const float ray_tmin,
const float ray_tmax,
@ -675,27 +639,17 @@ void metalrt_intersection_point_shadow(constant KernelParamsMetal &launch_params
{
const uint visibility = payload.visibility;
float3 P = ray_origin;
float3 dir = ray_direction;
/* The direction is not normalized by default, but the point intersection routine expects that */
float len;
dir = normalize_len(dir, &len);
Intersection isect;
isect.t = ray_tmax;
/* Transform maximum distance into object space */
if (isect.t != FLT_MAX)
isect.t *= len;
MetalKernelContext context(launch_params_metal);
if (context.point_intersect(NULL, &isect, P, dir, ray_tmin, isect.t, object, prim, time, type)) {
if (context.point_intersect(NULL, &isect, ray_P, ray_D, ray_tmin, isect.t, object, prim, time, type)) {
result.continue_search = metalrt_shadow_all_hit<METALRT_HIT_BOUNDING_BOX>(
launch_params_metal, payload, object, prim, float2(isect.u, isect.v), ray_tmax);
result.accept = !result.continue_search;
if (result.accept) {
result.distance = isect.t / len;
result.distance = isect.t;
}
}
}

35
intern/cycles/kernel/device/optix/kernel.cu
View File

@ -410,13 +410,9 @@ ccl_device_inline void optix_intersection_curve(const int prim, const int type)
}
# endif
float3 P = optixGetObjectRayOrigin();
float3 dir = optixGetObjectRayDirection();
float tmin = optixGetRayTmin();
/* The direction is not normalized by default, but the curve intersection routine expects that */
float len;
dir = normalize_len(dir, &len);
const float3 ray_P = optixGetObjectRayOrigin();
const float3 ray_D = optixGetObjectRayDirection();
const float ray_tmin = optixGetRayTmin();
# ifdef __OBJECT_MOTION__
const float time = optixGetRayTime();
@ -426,13 +422,10 @@ ccl_device_inline void optix_intersection_curve(const int prim, const int type)
Intersection isect;
isect.t = optixGetRayTmax();
/* Transform maximum distance into object space. */
if (isect.t != FLT_MAX)
isect.t *= len;
if (curve_intersect(NULL, &isect, P, dir, tmin, isect.t, object, prim, time, type)) {
if (curve_intersect(NULL, &isect, ray_P, ray_D, ray_tmin, isect.t, object, prim, time, type)) {
static_assert(PRIMITIVE_ALL < 128, "Values >= 128 are reserved for OptiX internal use");
optixReportIntersection(isect.t / len,
optixReportIntersection(isect.t,
type & PRIMITIVE_ALL,
__float_as_int(isect.u), /* Attribute_0 */
__float_as_int(isect.v)); /* Attribute_1 */
@ -465,13 +458,9 @@ extern "C" __global__ void __intersection__point()
}
# endif
float3 P = optixGetObjectRayOrigin();
float3 dir = optixGetObjectRayDirection();
float tmin = optixGetRayTmin();
/* The direction is not normalized by default, the point intersection routine expects that. */
float len;
dir = normalize_len(dir, &len);
const float3 ray_P = optixGetObjectRayOrigin();
const float3 ray_D = optixGetObjectRayDirection();
const float ray_tmin = optixGetRayTmin();
# ifdef __OBJECT_MOTION__
const float time = optixGetRayTime();
@ -481,14 +470,10 @@ extern "C" __global__ void __intersection__point()
Intersection isect;
isect.t = optixGetRayTmax();
/* Transform maximum distance into object space. */
if (isect.t != FLT_MAX) {
isect.t *= len;
}
if (point_intersect(NULL, &isect, P, dir, tmin, isect.t, object, prim, time, type)) {
if (point_intersect(NULL, &isect, ray_P, ray_D, ray_tmin, isect.t, object, prim, time, type)) {
static_assert(PRIMITIVE_ALL < 128, "Values >= 128 are reserved for OptiX internal use");
optixReportIntersection(isect.t / len, type & PRIMITIVE_ALL);
optixReportIntersection(isect.t, type & PRIMITIVE_ALL);
}
}
#endif

77
intern/cycles/kernel/geom/curve_intersect.h
View File

@ -72,7 +72,7 @@ ccl_device_inline float sqr_point_to_line_distance(const float3 PmQ0, const floa
ccl_device_inline bool cylinder_intersect(const float3 cylinder_start,
const float3 cylinder_end,
const float cylinder_radius,
const float3 ray_dir,
const float3 ray_D,
ccl_private float2 *t_o,
ccl_private float *u0_o,
ccl_private float3 *Ng0_o,
@ -82,7 +82,7 @@ ccl_device_inline bool cylinder_intersect(const float3 cylinder_start,
/* Calculate quadratic equation to solve. */
const float rl = 1.0f / len(cylinder_end - cylinder_start);
const float3 P0 = cylinder_start, dP = (cylinder_end - cylinder_start) * rl;
const float3 O = -P0, dO = ray_dir;
const float3 O = -P0, dO = ray_D;
const float dOdO = dot(dO, dO);
const float OdO = dot(dO, O);
@ -123,7 +123,7 @@ ccl_device_inline bool cylinder_intersect(const float3 cylinder_start,
/* Calculates u and Ng for near hit. */
{
*u0_o = (t0 * dOz + Oz) * rl;
const float3 Pr = t0 * ray_dir;
const float3 Pr = t0 * ray_D;
const float3 Pl = (*u0_o) * (cylinder_end - cylinder_start) + cylinder_start;
*Ng0_o = Pr - Pl;
}
@ -131,7 +131,7 @@ ccl_device_inline bool cylinder_intersect(const float3 cylinder_start,
/* Calculates u and Ng for far hit. */
{
*u1_o = (t1 * dOz + Oz) * rl;
const float3 Pr = t1 * ray_dir;
const float3 Pr = t1 * ray_D;
const float3 Pl = (*u1_o) * (cylinder_end - cylinder_start) + cylinder_start;
*Ng1_o = Pr - Pl;
}
@ -141,10 +141,10 @@ ccl_device_inline bool cylinder_intersect(const float3 cylinder_start,
return true;
}
ccl_device_inline float2 half_plane_intersect(const float3 P, const float3 N, const float3 ray_dir)
ccl_device_inline float2 half_plane_intersect(const float3 P, const float3 N, const float3 ray_D)
{
const float3 O = -P;
const float3 D = ray_dir;
const float3 D = ray_D;
const float ON = dot(O, N);
const float DN = dot(D, N);
const float min_rcp_input = 1e-18f;
@ -155,7 +155,7 @@ ccl_device_inline float2 half_plane_intersect(const float3 P, const float3 N, co
return make_float2(lower, upper);
}
ccl_device bool curve_intersect_iterative(const float3 ray_dir,
ccl_device bool curve_intersect_iterative(const float3 ray_D,
const float ray_tmin,
ccl_private float *ray_tmax,
const float dt,
@ -165,7 +165,7 @@ ccl_device bool curve_intersect_iterative(const float3 ray_dir,
const bool use_backfacing,
ccl_private Intersection *isect)
{
const float length_ray_dir = len(ray_dir);
const float length_ray_D = len(ray_D);
/* Error of curve evaluations is proportional to largest coordinate. */
const float4 box_min = min(min(curve[0], curve[1]), min(curve[2], curve[3]));
@ -176,9 +176,9 @@ ccl_device bool curve_intersect_iterative(const float3 ray_dir,
const float radius_max = box_max.w;
for (int i = 0; i < CURVE_NUM_JACOBIAN_ITERATIONS; i++) {
const float3 Q = ray_dir * t;
const float3 dQdt = ray_dir;
const float Q_err = 16.0f * FLT_EPSILON * length_ray_dir * t;
const float3 Q = ray_D * t;
const float3 dQdt = ray_D;
const float Q_err = 16.0f * FLT_EPSILON * length_ray_D * t;
const float4 P4 = catmull_rom_basis_eval(curve, u);
const float4 dPdu4 = catmull_rom_basis_derivative(curve, u);
@ -233,7 +233,7 @@ ccl_device bool curve_intersect_iterative(const float3 ray_dir,
const float3 U = dradiusdu * R + dPdu;
const float3 V = cross(dPdu, R);
const float3 Ng = cross(V, U);
if (!use_backfacing && dot(ray_dir, Ng) > 0.0f) {
if (!use_backfacing && dot(ray_D, Ng) > 0.0f) {
return false;
}
@ -249,8 +249,8 @@ ccl_device bool curve_intersect_iterative(const float3 ray_dir,
return false;
}
ccl_device bool curve_intersect_recursive(const float3 ray_orig,
const float3 ray_dir,
ccl_device bool curve_intersect_recursive(const float3 ray_P,
const float3 ray_D,
const float ray_tmin,
float ray_tmax,
float4 curve[4],
@ -258,8 +258,8 @@ ccl_device bool curve_intersect_recursive(const float3 ray_orig,
{
/* Move ray closer to make intersection stable. */
const float3 center = float4_to_float3(0.25f * (curve[0] + curve[1] + curve[2] + curve[3]));
const float dt = dot(center - ray_orig, ray_dir) / dot(ray_dir, ray_dir);
const float3 ref = ray_orig + ray_dir * dt;
const float dt = dot(center - ray_P, ray_D) / dot(ray_D, ray_D);
const float3 ref = ray_P + ray_D * dt;
const float4 ref4 = make_float4(ref.x, ref.y, ref.z, 0.0f);
curve[0] -= ref4;
curve[1] -= ref4;
@ -322,7 +322,7 @@ ccl_device bool curve_intersect_recursive(const float3 ray_orig,
valid = cylinder_intersect(float4_to_float3(P0),
float4_to_float3(P3),
r_outer,
ray_dir,
ray_D,
&tc_outer,
&u_outer0,
&Ng_outer0,
@ -335,11 +335,10 @@ ccl_device bool curve_intersect_recursive(const float3 ray_orig,
/* Intersect with cap-planes. */
float2 tp = make_float2(ray_tmin - dt, ray_tmax - dt);
tp = make_float2(max(tp.x, tc_outer.x), min(tp.y, tc_outer.y));
const float2 h0 = half_plane_intersect(
float4_to_float3(P0), float4_to_float3(dP0du), ray_dir);
const float2 h0 = half_plane_intersect(float4_to_float3(P0), float4_to_float3(dP0du), ray_D);
tp = make_float2(max(tp.x, h0.x), min(tp.y, h0.y));
const float2 h1 = half_plane_intersect(
float4_to_float3(P3), -float4_to_float3(dP3du), ray_dir);
float4_to_float3(P3), -float4_to_float3(dP3du), ray_D);
tp = make_float2(max(tp.x, h1.x), min(tp.y, h1.y));
valid = tp.x <= tp.y;
if (!valid) {
@ -359,7 +358,7 @@ ccl_device bool curve_intersect_recursive(const float3 ray_orig,
const bool valid_inner = cylinder_intersect(float4_to_float3(P0),
float4_to_float3(P3),
r_inner,
ray_dir,
ray_D,
&tc_inner,
&u_inner0,
&Ng_inner0,
@ -369,9 +368,9 @@ ccl_device bool curve_intersect_recursive(const float3 ray_orig,
/* At the unstable area we subdivide deeper. */
# if 0
const bool unstable0 = (!valid_inner) |
(fabsf(dot(normalize(ray_dir), normalize(Ng_inner0))) < 0.3f);
(fabsf(dot(normalize(ray_D), normalize(Ng_inner0))) < 0.3f);
const bool unstable1 = (!valid_inner) |
(fabsf(dot(normalize(ray_dir), normalize(Ng_inner1))) < 0.3f);
(fabsf(dot(normalize(ray_D), normalize(Ng_inner1))) < 0.3f);
# else
/* On the GPU appears to be a little faster if always enabled. */
(void)valid_inner;
@ -396,7 +395,7 @@ ccl_device bool curve_intersect_recursive(const float3 ray_orig,
CURVE_NUM_BEZIER_SUBDIVISIONS;
if (depth >= termDepth) {
found |= curve_intersect_iterative(
ray_dir, ray_tmin, &ray_tmax, dt, curve, u_outer0, tp0.x, use_backfacing, isect);
ray_D, ray_tmin, &ray_tmax, dt, curve, u_outer0, tp0.x, use_backfacing, isect);
}
else {
recurse = true;
@ -409,7 +408,7 @@ ccl_device bool curve_intersect_recursive(const float3 ray_orig,
CURVE_NUM_BEZIER_SUBDIVISIONS;
if (depth >= termDepth) {
found |= curve_intersect_iterative(
ray_dir, ray_tmin, &ray_tmax, dt, curve, u_outer1, tp1.y, use_backfacing, isect);
ray_D, ray_tmin, &ray_tmax, dt, curve, u_outer1, tp1.y, use_backfacing, isect);
}
else {
recurse = true;
@ -519,13 +518,16 @@ ccl_device_inline bool ribbon_intersect_quad(const float ray_tmin,
return true;
}
ccl_device_inline void ribbon_ray_space(const float3 ray_dir, float3 ray_space[3])
ccl_device_inline void ribbon_ray_space(const float3 ray_D,
const float ray_D_invlen,
float3 ray_space[3])
{
const float3 dx0 = make_float3(0, ray_dir.z, -ray_dir.y);
const float3 dx1 = make_float3(-ray_dir.z, 0, ray_dir.x);
const float3 D = ray_D * ray_D_invlen;
const float3 dx0 = make_float3(0, D.z, -D.y);
const float3 dx1 = make_float3(-D.z, 0, D.x);
ray_space[0] = normalize(dot(dx0, dx0) > dot(dx1, dx1) ? dx0 : dx1);
ray_space[1] = normalize(cross(ray_dir, ray_space[0]));
ray_space[2] = ray_dir;
ray_space[1] = normalize(cross(D, ray_space[0]));
ray_space[2] = D * ray_D_invlen;
}
ccl_device_inline float4 ribbon_to_ray_space(const float3 ray_space[3],
@ -537,7 +539,7 @@ ccl_device_inline float4 ribbon_to_ray_space(const float3 ray_space[3],
}
ccl_device_inline bool ribbon_intersect(const float3 ray_org,
const float3 ray_dir,
const float3 ray_D,
const float ray_tmin,
float ray_tmax,
const int N,
@ -545,8 +547,9 @@ ccl_device_inline bool ribbon_intersect(const float3 ray_org,
ccl_private Intersection *isect)
{
/* Transform control points into ray space. */
const float ray_D_invlen = 1.0f / len(ray_D);
float3 ray_space[3];
ribbon_ray_space(ray_dir, ray_space);
ribbon_ray_space(ray_D, ray_D_invlen, ray_space);
curve[0] = ribbon_to_ray_space(ray_space, ray_org, curve[0]);
curve[1] = ribbon_to_ray_space(ray_space, ray_org, curve[1]);
@ -594,7 +597,7 @@ ccl_device_inline bool ribbon_intersect(const float3 ray_org,
const float avoidance_factor = 2.0f;
if (avoidance_factor != 0.0f) {
float r = mix(p0.w, p1.w, vu);
valid0 = vt > avoidance_factor * r;
valid0 = vt > avoidance_factor * r * ray_D_invlen;
}
if (valid0) {
@ -619,8 +622,8 @@ ccl_device_inline bool ribbon_intersect(const float3 ray_org,
ccl_device_forceinline bool curve_intersect(KernelGlobals kg,
ccl_private Intersection *isect,
const float3 P,
const float3 dir,
const float3 ray_P,
const float3 ray_D,
const float tmin,
const float tmax,
int object,
@ -651,7 +654,7 @@ ccl_device_forceinline bool curve_intersect(KernelGlobals kg,
if (type & PRIMITIVE_CURVE_RIBBON) {
/* todo: adaptive number of subdivisions could help performance here. */
const int subdivisions = kernel_data.bvh.curve_subdivisions;
if (ribbon_intersect(P, dir, tmin, tmax, subdivisions, curve, isect)) {
if (ribbon_intersect(ray_P, ray_D, tmin, tmax, subdivisions, curve, isect)) {
isect->prim = prim;
isect->object = object;
isect->type = type;
@ -661,7 +664,7 @@ ccl_device_forceinline bool curve_intersect(KernelGlobals kg,
return false;
}
else {
if (curve_intersect_recursive(P, dir, tmin, tmax, curve, isect)) {
if (curve_intersect_recursive(ray_P, ray_D, tmin, tmax, curve, isect)) {
isect->prim = prim;
isect->object = object;
isect->type = type;

111
intern/cycles/kernel/geom/object.h
View File

@ -488,59 +488,30 @@ ccl_device_inline float3 bvh_inverse_direction(float3 dir)
/* Transform ray into object space to enter static object in BVH */
ccl_device_inline float bvh_instance_push(KernelGlobals kg,
int object,
ccl_private const Ray *ray,
ccl_private float3 *P,
ccl_private float3 *dir,
ccl_private float3 *idir)
ccl_device_inline void bvh_instance_push(KernelGlobals kg,
int object,
ccl_private const Ray *ray,
ccl_private float3 *P,
ccl_private float3 *dir,
ccl_private float3 *idir)
{
Transform tfm = object_fetch_transform(kg, object, OBJECT_INVERSE_TRANSFORM);
*P = transform_point(&tfm, ray->P);
float len;
*dir = bvh_clamp_direction(normalize_len(transform_direction(&tfm, ray->D), &len));
*dir = bvh_clamp_direction(transform_direction(&tfm, ray->D));
*idir = bvh_inverse_direction(*dir);
return len;
}
/* Transform ray to exit static object in BVH. */
ccl_device_inline float bvh_instance_pop(KernelGlobals kg,
int object,
ccl_private const Ray *ray,
ccl_private float3 *P,
ccl_private float3 *dir,
ccl_private float3 *idir,
float t)
ccl_device_inline void bvh_instance_pop(KernelGlobals kg,
int object,
ccl_private const Ray *ray,
ccl_private float3 *P,
ccl_private float3 *dir,
ccl_private float3 *idir)
{
if (t != FLT_MAX) {
Transform tfm = object_fetch_transform(kg, object, OBJECT_INVERSE_TRANSFORM);
t /= len(transform_direction(&tfm, ray->D));
}
*P = ray->P;
*dir = bvh_clamp_direction(ray->D);
*idir = bvh_inverse_direction(*dir);
return t;
}
/* Same as above, but returns scale factor to apply to multiple intersection distances */
ccl_device_inline void bvh_instance_pop_factor(KernelGlobals kg,
int object,
ccl_private const Ray *ray,
ccl_private float3 *P,
ccl_private float3 *dir,
ccl_private float3 *idir,
ccl_private float *t_fac)
{
Transform tfm = object_fetch_transform(kg, object, OBJECT_INVERSE_TRANSFORM);
*t_fac = 1.0f / len(transform_direction(&tfm, ray->D));
*P = ray->P;
*dir = bvh_clamp_direction(ray->D);
*idir = bvh_inverse_direction(*dir);
@ -549,59 +520,31 @@ ccl_device_inline void bvh_instance_pop_factor(KernelGlobals kg,
#ifdef __OBJECT_MOTION__
/* Transform ray into object space to enter motion blurred object in BVH */
ccl_device_inline float bvh_instance_motion_push(KernelGlobals kg,
int object,
ccl_private const Ray *ray,
ccl_private float3 *P,
ccl_private float3 *dir,
ccl_private float3 *idir,
ccl_private Transform *itfm)
{
object_fetch_transform_motion_test(kg, object, ray->time, itfm);
*P = transform_point(itfm, ray->P);
float len;
*dir = bvh_clamp_direction(normalize_len(transform_direction(itfm, ray->D), &len));
*idir = bvh_inverse_direction(*dir);
return len;
}
/* Transform ray to exit motion blurred object in BVH. */
ccl_device_inline float bvh_instance_motion_pop(KernelGlobals kg,
ccl_device_inline void bvh_instance_motion_push(KernelGlobals kg,
int object,
ccl_private const Ray *ray,
ccl_private float3 *P,
ccl_private float3 *dir,
ccl_private float3 *idir,
float t,
ccl_private Transform *itfm)
ccl_private float3 *idir)
{
if (t != FLT_MAX) {
t /= len(transform_direction(itfm, ray->D));
}
Transform tfm;
object_fetch_transform_motion_test(kg, object, ray->time, &tfm);
*P = ray->P;
*dir = bvh_clamp_direction(ray->D);
*P = transform_point(&tfm, ray->P);
*dir = bvh_clamp_direction(transform_direction(&tfm, ray->D));
*idir = bvh_inverse_direction(*dir);
return t;
}
/* Same as above, but returns scale factor to apply to multiple intersection distances */
/* Transform ray to exit motion blurred object in BVH. */
ccl_device_inline void bvh_instance_motion_pop_factor(KernelGlobals kg,
int object,
ccl_private const Ray *ray,
ccl_private float3 *P,
ccl_private float3 *dir,
ccl_private float3 *idir,
ccl_private float *t_fac,
ccl_private Transform *itfm)
ccl_device_inline void bvh_instance_motion_pop(KernelGlobals kg,
int object,
ccl_private const Ray *ray,
ccl_private float3 *P,
ccl_private float3 *dir,
ccl_private float3 *idir)
{
*t_fac = 1.0f / len(transform_direction(itfm, ray->D));
*P = ray->P;
*dir = bvh_clamp_direction(ray->D);
*idir = bvh_inverse_direction(*dir);

30
intern/cycles/kernel/geom/point_intersect.h
View File

@ -10,20 +10,20 @@ CCL_NAMESPACE_BEGIN
#ifdef __POINTCLOUD__
ccl_device_forceinline bool point_intersect_test(const float4 point,
const float3 P,
const float3 dir,
const float tmin,
const float tmax,
const float3 ray_P,
const float3 ray_D,
const float ray_tmin,
const float ray_tmax,
ccl_private float *t)
{
const float3 center = float4_to_float3(point);
const float radius = point.w;
const float rd2 = 1.0f / dot(dir, dir);
const float rd2 = 1.0f / dot(ray_D, ray_D);
const float3 c0 = center - P;
const float projC0 = dot(c0, dir) * rd2;
const float3 perp = c0 - projC0 * dir;
const float3 c0 = center - ray_P;
const float projC0 = dot(c0, ray_D) * rd2;
const float3 perp = c0 - projC0 * ray_D;
const float l2 = dot(perp, perp);
const float r2 = radius * radius;
if (!(l2 <= r2)) {
@ -32,12 +32,12 @@ ccl_device_forceinline bool point_intersect_test(const float4 point,
const float td = sqrt((r2 - l2) * rd2);
const float t_front = projC0 - td;
const bool valid_front = (tmin <= t_front) & (t_front <= tmax);
const bool valid_front = (ray_tmin <= t_front) & (t_front <= ray_tmax);
/* Always back-face culling for now. */
# if 0
const float t_back = projC0 + td;
const bool valid_back = (tmin <= t_back) & (t_back <= tmax);
const bool valid_back = (ray_tmin <= t_back) & (t_back <= ray_tmax);
/* check if there is a first hit */
const bool valid_first = valid_front | valid_back;
@ -58,10 +58,10 @@ ccl_device_forceinline bool point_intersect_test(const float4 point,
ccl_device_forceinline bool point_intersect(KernelGlobals kg,
ccl_private Intersection *isect,
const float3 P,
const float3 dir,
const float tmin,
const float tmax,
const float3 ray_P,
const float3 ray_D,
const float ray_tmin,
const float ray_tmax,
const int object,
const int prim,
const float time,
@ -70,7 +70,7 @@ ccl_device_forceinline bool point_intersect(KernelGlobals kg,
const float4 point = (type & PRIMITIVE_MOTION) ? motion_point(kg, object, prim, time) :
kernel_data_fetch(points, prim);
if (!point_intersect_test(point, P, dir, tmin, tmax, &isect->t)) {
if (!point_intersect_test(point, ray_P, ray_D, ray_tmin, ray_tmax, &isect->t)) {
return false;
}