diff options
Diffstat (limited to 'intern')
| -rw-r--r-- | intern/cycles/kernel/bvh/bvh.h | 9 | ||||
| -rw-r--r-- | intern/cycles/kernel/bvh/local.h | 8 | ||||
| -rw-r--r-- | intern/cycles/kernel/bvh/shadow_all.h | 81 | ||||
| -rw-r--r-- | intern/cycles/kernel/bvh/traversal.h | 19 | ||||
| -rw-r--r-- | intern/cycles/kernel/bvh/volume.h | 20 | ||||
| -rw-r--r-- | intern/cycles/kernel/bvh/volume_all.h | 60 | ||||
| -rw-r--r-- | intern/cycles/kernel/device/metal/kernel.metal | 96 | ||||
| -rw-r--r-- | intern/cycles/kernel/device/optix/kernel.cu | 35 | ||||
| -rw-r--r-- | intern/cycles/kernel/geom/curve_intersect.h | 77 | ||||
| -rw-r--r-- | intern/cycles/kernel/geom/object.h | 111 | ||||
| -rw-r--r-- | intern/cycles/kernel/geom/point_intersect.h | 30 |
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; } |