Welcome to mirror list, hosted at ThFree Co, Russian Federation.

git.blender.org/blender.git - Unnamed repository; edit this file 'description' to name the repository.
summaryrefslogtreecommitdiff
path: root/intern
diff options
context:
space:
mode:
authorBrecht Van Lommel <brecht@blender.org>2022-07-21 17:37:38 +0300
committerBrecht Van Lommel <brecht@blender.org>2022-07-25 14:27:40 +0300
commit484ad3165307391aa5c55656b876b3ff7d615e80 (patch)
treebec9ffec4502fb93981cfae7d8b7b82852c0233d /intern
parent023eb2ea7c16a00272f83d564145e28aeb9ed2b7 (diff)
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.
Diffstat (limited to 'intern')
-rw-r--r--intern/cycles/kernel/bvh/bvh.h9
-rw-r--r--intern/cycles/kernel/bvh/local.h8
-rw-r--r--intern/cycles/kernel/bvh/shadow_all.h81
-rw-r--r--intern/cycles/kernel/bvh/traversal.h19
-rw-r--r--intern/cycles/kernel/bvh/volume.h20
-rw-r--r--intern/cycles/kernel/bvh/volume_all.h60
-rw-r--r--intern/cycles/kernel/device/metal/kernel.metal96
-rw-r--r--intern/cycles/kernel/device/optix/kernel.cu35
-rw-r--r--intern/cycles/kernel/geom/curve_intersect.h77
-rw-r--r--intern/cycles/kernel/geom/object.h111
-rw-r--r--intern/cycles/kernel/geom/point_intersect.h30
11 files changed, 153 insertions, 393 deletions
diff --git a/intern/cycles/kernel/bvh/bvh.h b/intern/cycles/kernel/bvh/bvh.h
index 9972de86c47..387e74b9885 100644
--- a/intern/cycles/kernel/bvh/bvh.h
+++ b/intern/cycles/kernel/bvh/bvh.h
@@ -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) {
diff --git a/intern/cycles/kernel/bvh/local.h b/intern/cycles/kernel/bvh/local.h
index 017a241ef4a..add61adc126 100644
--- a/intern/cycles/kernel/bvh/local.h
+++ b/intern/cycles/kernel/bvh/local.h
@@ -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;
}
diff --git a/intern/cycles/kernel/bvh/shadow_all.h b/intern/cycles/kernel/bvh/shadow_all.h
index db3c91569aa..f37af2a1e65 100644
--- a/intern/cycles/kernel/bvh/shadow_all.h
+++ b/intern/cycles/kernel/bvh/shadow_all.h
@@ -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;
}
diff --git a/intern/cycles/kernel/bvh/traversal.h b/intern/cycles/kernel/bvh/traversal.h
index 0ff38bf02de..9069d16912b 100644
--- a/intern/cycles/kernel/bvh/traversal.h
+++ b/intern/cycles/kernel/bvh/traversal.h
@@ -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];
diff --git a/intern/cycles/kernel/bvh/volume.h b/intern/cycles/kernel/bvh/volume.h
index bd4e508ecac..cc3915b4bf7 100644
--- a/intern/cycles/kernel/bvh/volume.h
+++ b/intern/cycles/kernel/bvh/volume.h
@@ -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;
diff --git a/intern/cycles/kernel/bvh/volume_all.h b/intern/cycles/kernel/bvh/volume_all.h
index c6eeb07a14d..5cdea3e354c 100644
--- a/intern/cycles/kernel/bvh/volume_all.h
+++ b/intern/cycles/kernel/bvh/volume_all.h
@@ -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];
diff --git a/intern/cycles/kernel/device/metal/kernel.metal b/intern/cycles/kernel/device/metal/kernel.metal
index 764c26dbe8f..8c6f2e1df5e 100644
--- a/intern/cycles/kernel/device/metal/kernel.metal
+++ b/intern/cycles/kernel/device/metal/kernel.metal
@@ -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;
}
}
}
diff --git a/intern/cycles/kernel/device/optix/kernel.cu b/intern/cycles/kernel/device/optix/kernel.cu
index 510f7cca5d6..204aa8182a1 100644
--- a/intern/cycles/kernel/device/optix/kernel.cu
+++ b/intern/cycles/kernel/device/optix/kernel.cu
@@ -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
diff --git a/intern/cycles/kernel/geom/curve_intersect.h b/intern/cycles/kernel/geom/curve_intersect.h
index 9770105dd81..97644aacaa8 100644
--- a/intern/cycles/kernel/geom/curve_intersect.h
+++ b/intern/cycles/kernel/geom/curve_intersect.h
@@ -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;
diff --git a/intern/cycles/kernel/geom/object.h b/intern/cycles/kernel/geom/object.h
index bef7d710159..badfd311985 100644
--- a/intern/cycles/kernel/geom/object.h
+++ b/intern/cycles/kernel/geom/object.h
@@ -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);
- *idir = bvh_inverse_direction(*dir);
+ *P = transform_point(&tfm, ray->P);
- return t;
+ *dir = bvh_clamp_direction(transform_direction(&tfm, ray->D));
+ *idir = bvh_inverse_direction(*dir);
}
-/* 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);
diff --git a/intern/cycles/kernel/geom/point_intersect.h b/intern/cycles/kernel/geom/point_intersect.h
index ee5a564947b..15fb814c58d 100644
--- a/intern/cycles/kernel/geom/point_intersect.h
+++ b/intern/cycles/kernel/geom/point_intersect.h
@@ -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;
}