diff options
| author | Brecht Van Lommel <brecht@blender.org> | 2022-07-13 17:54:53 +0300 |
|---|---|---|
| committer | Brecht Van Lommel <brecht@blender.org> | 2022-07-15 19:46:24 +0300 |
| commit | 5152c7c152e52d563cbd3ba3c792de3af0c2c14f (patch) | |
| tree | 45cad004066413ce309279ad52f2a363e5477150 /intern/cycles/kernel | |
| parent | bb376da6dfdd2476fc3738ce1fc89dac27825cef (diff) | |
Cycles: refactor rays to have start and end distance, fix precision issues
For transparency, volume and light intersection rays, adjust these distances
rather than the ray start position. This way we increment the start distance
by the smallest possible float increment to avoid self intersections, and be
sure it works as the distance compared to be will be exactly the same as
before, due to the ray start position and direction remaining the same.
Fix T98764, T96537, hair ray tracing precision issues.
Differential Revision: https://developer.blender.org/D15455
Diffstat (limited to 'intern/cycles/kernel')
41 files changed, 407 insertions, 277 deletions
diff --git a/intern/cycles/kernel/bvh/bvh.h b/intern/cycles/kernel/bvh/bvh.h index f375529a6f6..9972de86c47 100644 --- a/intern/cycles/kernel/bvh/bvh.h +++ b/intern/cycles/kernel/bvh/bvh.h @@ -175,8 +175,8 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg, optixTrace(scene_intersect_valid(ray) ? kernel_data.device_bvh : 0, ray->P, ray->D, - 0.0f, - ray->t, + ray->tmin, + ray->tmax, ray->time, ray_mask, ray_flags, @@ -203,28 +203,28 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg, #elif defined(__METALRT__) if (!scene_intersect_valid(ray)) { - isect->t = ray->t; + isect->t = ray->tmax; isect->type = PRIMITIVE_NONE; return false; } # if defined(__KERNEL_DEBUG__) if (is_null_instance_acceleration_structure(metal_ancillaries->accel_struct)) { - isect->t = ray->t; + isect->t = ray->tmax; isect->type = PRIMITIVE_NONE; kernel_assert(!"Invalid metal_ancillaries->accel_struct pointer"); return false; } if (is_null_intersection_function_table(metal_ancillaries->ift_default)) { - isect->t = ray->t; + isect->t = ray->tmax; isect->type = PRIMITIVE_NONE; kernel_assert(!"Invalid ift_default"); return false; } # endif - metal::raytracing::ray r(ray->P, ray->D, 0.0f, ray->t); + metal::raytracing::ray r(ray->P, ray->D, ray->tmin, ray->tmax); metalrt_intersector_type metalrt_intersect; if (!kernel_data.bvh.have_curves) { @@ -263,7 +263,7 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg, # endif if (intersection.type == intersection_type::none) { - isect->t = ray->t; + isect->t = ray->tmax; isect->type = PRIMITIVE_NONE; return false; @@ -296,7 +296,7 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg, # ifdef __EMBREE__ if (kernel_data.device_bvh) { - isect->t = ray->t; + isect->t = ray->tmax; CCLIntersectContext ctx(kg, CCLIntersectContext::RAY_REGULAR); IntersectContext rtc_ctx(&ctx); RTCRayHit ray_hit; @@ -360,8 +360,8 @@ ccl_device_intersect bool scene_intersect_local(KernelGlobals kg, optixTrace(scene_intersect_valid(ray) ? kernel_data.device_bvh : 0, ray->P, ray->D, - 0.0f, - ray->t, + ray->tmin, + ray->tmax, ray->time, 0xFF, /* Need to always call into __anyhit__kernel_optix_local_hit. */ @@ -405,7 +405,7 @@ ccl_device_intersect bool scene_intersect_local(KernelGlobals kg, } # endif - metal::raytracing::ray r(ray->P, ray->D, 0.0f, ray->t); + metal::raytracing::ray r(ray->P, ray->D, ray->tmin, ray->tmax); metalrt_intersector_type metalrt_intersect; metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque); @@ -476,7 +476,7 @@ ccl_device_intersect bool scene_intersect_local(KernelGlobals kg, float3 dir = ray->D; float3 idir = ray->D; Transform ob_itfm; - rtc_ray.tfar = ray->t * + 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. */ @@ -542,8 +542,8 @@ ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg, optixTrace(scene_intersect_valid(ray) ? kernel_data.device_bvh : 0, ray->P, ray->D, - 0.0f, - ray->t, + ray->tmin, + ray->tmax, ray->time, ray_mask, /* Need to always call into __anyhit__kernel_optix_shadow_all_hit. */ @@ -582,7 +582,7 @@ ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg, } # endif - metal::raytracing::ray r(ray->P, ray->D, 0.0f, ray->t); + metal::raytracing::ray r(ray->P, ray->D, ray->tmin, ray->tmax); metalrt_intersector_type metalrt_intersect; metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque); @@ -701,8 +701,8 @@ ccl_device_intersect bool scene_intersect_volume(KernelGlobals kg, optixTrace(scene_intersect_valid(ray) ? kernel_data.device_bvh : 0, ray->P, ray->D, - 0.0f, - ray->t, + ray->tmin, + ray->tmax, ray->time, ray_mask, /* Need to always call into __anyhit__kernel_optix_volume_test. */ @@ -744,7 +744,7 @@ ccl_device_intersect bool scene_intersect_volume(KernelGlobals kg, } # endif - metal::raytracing::ray r(ray->P, ray->D, 0.0f, ray->t); + metal::raytracing::ray r(ray->P, ray->D, ray->tmin, ray->tmax); metalrt_intersector_type metalrt_intersect; metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque); diff --git a/intern/cycles/kernel/bvh/embree.h b/intern/cycles/kernel/bvh/embree.h index 77eec2468f4..fecbccac2f8 100644 --- a/intern/cycles/kernel/bvh/embree.h +++ b/intern/cycles/kernel/bvh/embree.h @@ -83,8 +83,8 @@ ccl_device_inline void kernel_embree_setup_ray(const Ray &ray, rtc_ray.dir_x = ray.D.x; rtc_ray.dir_y = ray.D.y; rtc_ray.dir_z = ray.D.z; - rtc_ray.tnear = 0.0f; - rtc_ray.tfar = ray.t; + rtc_ray.tnear = ray.tmin; + rtc_ray.tfar = ray.tmax; rtc_ray.time = ray.time; rtc_ray.mask = visibility; } diff --git a/intern/cycles/kernel/bvh/local.h b/intern/cycles/kernel/bvh/local.h index 3b6b30ea93d..017a241ef4a 100644 --- a/intern/cycles/kernel/bvh/local.h +++ b/intern/cycles/kernel/bvh/local.h @@ -47,8 +47,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; int object = OBJECT_NONE; - float isect_t = ray->t; + float isect_t = ray->tmax; if (local_isect != NULL) { local_isect->num_hits = 0; @@ -59,10 +60,13 @@ ccl_device_inline if (!(object_flag & SD_OBJECT_TRANSFORM_APPLIED)) { #if BVH_FEATURE(BVH_MOTION) Transform ob_itfm; - isect_t *= bvh_instance_motion_push(kg, local_object, ray, &P, &dir, &idir, &ob_itfm); + const float t_world_to_instance = bvh_instance_motion_push( + kg, local_object, ray, &P, &dir, &idir, &ob_itfm); #else - isect_t *= bvh_instance_push(kg, local_object, ray, &P, &dir, &idir); + const float t_world_to_instance = 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,6 +85,7 @@ ccl_device_inline dir, #endif idir, + tmin, isect_t, node_addr, PATH_RAY_ALL_VISIBILITY, @@ -155,6 +160,7 @@ ccl_device_inline local_object, prim, prim_addr, + tmin, isect_t, lcg_state, max_hits)) { @@ -191,6 +197,7 @@ ccl_device_inline local_object, prim, prim_addr, + tmin, isect_t, lcg_state, max_hits)) { diff --git a/intern/cycles/kernel/bvh/nodes.h b/intern/cycles/kernel/bvh/nodes.h index c19dea9223b..e02841fad16 100644 --- a/intern/cycles/kernel/bvh/nodes.h +++ b/intern/cycles/kernel/bvh/nodes.h @@ -18,7 +18,8 @@ ccl_device_forceinline Transform bvh_unaligned_node_fetch_space(KernelGlobals kg ccl_device_forceinline int bvh_aligned_node_intersect(KernelGlobals kg, const float3 P, const float3 idir, - const float t, + const float tmin, + const float tmax, const int node_addr, const uint visibility, float dist[2]) @@ -39,8 +40,8 @@ ccl_device_forceinline int bvh_aligned_node_intersect(KernelGlobals kg, float c0hiy = (node1.z - P.y) * idir.y; float c0loz = (node2.x - P.z) * idir.z; float c0hiz = (node2.z - P.z) * idir.z; - float c0min = max4(0.0f, min(c0lox, c0hix), min(c0loy, c0hiy), min(c0loz, c0hiz)); - float c0max = min4(t, max(c0lox, c0hix), max(c0loy, c0hiy), max(c0loz, c0hiz)); + float c0min = max4(tmin, min(c0lox, c0hix), min(c0loy, c0hiy), min(c0loz, c0hiz)); + float c0max = min4(tmax, max(c0lox, c0hix), max(c0loy, c0hiy), max(c0loz, c0hiz)); float c1lox = (node0.y - P.x) * idir.x; float c1hix = (node0.w - P.x) * idir.x; @@ -48,8 +49,8 @@ ccl_device_forceinline int bvh_aligned_node_intersect(KernelGlobals kg, float c1hiy = (node1.w - P.y) * idir.y; float c1loz = (node2.y - P.z) * idir.z; float c1hiz = (node2.w - P.z) * idir.z; - float c1min = max4(0.0f, min(c1lox, c1hix), min(c1loy, c1hiy), min(c1loz, c1hiz)); - float c1max = min4(t, max(c1lox, c1hix), max(c1loy, c1hiy), max(c1loz, c1hiz)); + float c1min = max4(tmin, min(c1lox, c1hix), min(c1loy, c1hiy), min(c1loz, c1hiz)); + float c1max = min4(tmax, max(c1lox, c1hix), max(c1loy, c1hiy), max(c1loz, c1hiz)); dist[0] = c0min; dist[1] = c1min; @@ -66,7 +67,8 @@ ccl_device_forceinline int bvh_aligned_node_intersect(KernelGlobals kg, ccl_device_forceinline bool bvh_unaligned_node_intersect_child(KernelGlobals kg, const float3 P, const float3 dir, - const float t, + const float tmin, + const float tmax, int node_addr, int child, float dist[2]) @@ -83,8 +85,8 @@ ccl_device_forceinline bool bvh_unaligned_node_intersect_child(KernelGlobals kg, const float far_x = max(lower_xyz.x, upper_xyz.x); const float far_y = max(lower_xyz.y, upper_xyz.y); const float far_z = max(lower_xyz.z, upper_xyz.z); - const float tnear = max4(0.0f, near_x, near_y, near_z); - const float tfar = min4(t, far_x, far_y, far_z); + const float tnear = max4(tmin, near_x, near_y, near_z); + const float tfar = min4(tmax, far_x, far_y, far_z); *dist = tnear; return tnear <= tfar; } @@ -93,7 +95,8 @@ ccl_device_forceinline int bvh_unaligned_node_intersect(KernelGlobals kg, const float3 P, const float3 dir, const float3 idir, - const float t, + const float tmin, + const float tmax, const int node_addr, const uint visibility, float dist[2]) @@ -102,7 +105,7 @@ ccl_device_forceinline int bvh_unaligned_node_intersect(KernelGlobals kg, #ifdef __VISIBILITY_FLAG__ float4 cnodes = kernel_data_fetch(bvh_nodes, node_addr + 0); #endif - if (bvh_unaligned_node_intersect_child(kg, P, dir, t, node_addr, 0, &dist[0])) { + if (bvh_unaligned_node_intersect_child(kg, P, dir, tmin, tmax, node_addr, 0, &dist[0])) { #ifdef __VISIBILITY_FLAG__ if ((__float_as_uint(cnodes.x) & visibility)) #endif @@ -110,7 +113,7 @@ ccl_device_forceinline int bvh_unaligned_node_intersect(KernelGlobals kg, mask |= 1; } } - if (bvh_unaligned_node_intersect_child(kg, P, dir, t, node_addr, 1, &dist[1])) { + if (bvh_unaligned_node_intersect_child(kg, P, dir, tmin, tmax, node_addr, 1, &dist[1])) { #ifdef __VISIBILITY_FLAG__ if ((__float_as_uint(cnodes.y) & visibility)) #endif @@ -125,16 +128,17 @@ ccl_device_forceinline int bvh_node_intersect(KernelGlobals kg, const float3 P, const float3 dir, const float3 idir, - const float t, + const float tmin, + const float tmax, const int node_addr, const uint visibility, float dist[2]) { float4 node = kernel_data_fetch(bvh_nodes, node_addr); if (__float_as_uint(node.x) & PATH_RAY_NODE_UNALIGNED) { - return bvh_unaligned_node_intersect(kg, P, dir, idir, t, node_addr, visibility, dist); + return bvh_unaligned_node_intersect(kg, P, dir, idir, tmin, tmax, node_addr, visibility, dist); } else { - return bvh_aligned_node_intersect(kg, P, idir, t, node_addr, visibility, dist); + return bvh_aligned_node_intersect(kg, P, idir, tmin, tmax, node_addr, visibility, dist); } } diff --git a/intern/cycles/kernel/bvh/shadow_all.h b/intern/cycles/kernel/bvh/shadow_all.h index e86fe867eac..db3c91569aa 100644 --- a/intern/cycles/kernel/bvh/shadow_all.h +++ b/intern/cycles/kernel/bvh/shadow_all.h @@ -49,6 +49,7 @@ ccl_device_inline float3 P = ray->P; float3 dir = bvh_clamp_direction(ray->D); float3 idir = bvh_inverse_direction(dir); + float tmin = ray->tmin; int object = OBJECT_NONE; uint num_hits = 0; @@ -59,12 +60,12 @@ ccl_device_inline /* 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->t; + 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->t; + float t_max_current = ray->tmax; /* Conversion from world to local space for the current instance if any, 1.0 * otherwise. */ @@ -88,6 +89,7 @@ ccl_device_inline dir, #endif idir, + tmin, t_max_current, node_addr, visibility, @@ -156,8 +158,16 @@ ccl_device_inline switch (type & PRIMITIVE_ALL) { case PRIMITIVE_TRIANGLE: { - hit = triangle_intersect( - kg, &isect, P, dir, t_max_current, visibility, prim_object, prim, prim_addr); + hit = triangle_intersect(kg, + &isect, + P, + dir, + tmin, + t_max_current, + visibility, + prim_object, + prim, + prim_addr); break; } #if BVH_FEATURE(BVH_MOTION) @@ -166,6 +176,7 @@ ccl_device_inline &isect, P, dir, + tmin, t_max_current, ray->time, visibility, @@ -189,8 +200,16 @@ ccl_device_inline } const int curve_type = kernel_data_fetch(prim_type, prim_addr); - hit = curve_intersect( - kg, &isect, P, dir, t_max_current, prim_object, prim, ray->time, curve_type); + hit = curve_intersect(kg, + &isect, + P, + dir, + tmin, + t_max_current, + prim_object, + prim, + ray->time, + curve_type); break; } @@ -207,8 +226,16 @@ ccl_device_inline } const int point_type = kernel_data_fetch(prim_type, prim_addr); - hit = point_intersect( - kg, &isect, P, dir, t_max_current, prim_object, prim, ray->time, point_type); + hit = point_intersect(kg, + &isect, + P, + dir, + tmin, + t_max_current, + prim_object, + prim, + ray->time, + point_type); break; } #endif /* BVH_FEATURE(BVH_POINTCLOUD) */ @@ -302,6 +329,7 @@ ccl_device_inline /* 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); @@ -323,7 +351,8 @@ ccl_device_inline #endif /* Restore world space ray length. */ - t_max_current = ray->t; + tmin = ray->tmin; + t_max_current = ray->tmax; object = OBJECT_NONE; t_world_to_instance = 1.0f; diff --git a/intern/cycles/kernel/bvh/traversal.h b/intern/cycles/kernel/bvh/traversal.h index 784fbf4fd11..0ff38bf02de 100644 --- a/intern/cycles/kernel/bvh/traversal.h +++ b/intern/cycles/kernel/bvh/traversal.h @@ -43,13 +43,14 @@ 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; int object = OBJECT_NONE; #if BVH_FEATURE(BVH_MOTION) Transform ob_itfm; #endif - isect->t = ray->t; + isect->t = ray->tmax; isect->u = 0.0f; isect->v = 0.0f; isect->prim = PRIM_NONE; @@ -71,6 +72,7 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg, dir, #endif idir, + tmin, isect->t, node_addr, visibility, @@ -133,8 +135,16 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg, switch (type & PRIMITIVE_ALL) { case PRIMITIVE_TRIANGLE: { - if (triangle_intersect( - kg, isect, P, dir, isect->t, visibility, prim_object, prim, prim_addr)) { + if (triangle_intersect(kg, + isect, + P, + dir, + tmin, + isect->t, + visibility, + prim_object, + prim, + prim_addr)) { /* shadow ray early termination */ if (visibility & PATH_RAY_SHADOW_OPAQUE) return true; @@ -147,6 +157,7 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg, isect, P, dir, + tmin, isect->t, ray->time, visibility, @@ -174,7 +185,7 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg, const int curve_type = kernel_data_fetch(prim_type, prim_addr); const bool hit = curve_intersect( - kg, isect, P, dir, isect->t, prim_object, prim, ray->time, curve_type); + kg, isect, P, dir, tmin, isect->t, prim_object, prim, ray->time, curve_type); if (hit) { /* shadow ray early termination */ if (visibility & PATH_RAY_SHADOW_OPAQUE) @@ -195,7 +206,7 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg, const int point_type = kernel_data_fetch(prim_type, prim_addr); const bool hit = point_intersect( - kg, isect, P, dir, isect->t, prim_object, prim, ray->time, point_type); + kg, isect, P, dir, tmin, isect->t, prim_object, prim, ray->time, point_type); if (hit) { /* shadow ray early termination */ if (visibility & PATH_RAY_SHADOW_OPAQUE) @@ -212,11 +223,15 @@ 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) - isect->t *= bvh_instance_motion_push(kg, object, ray, &P, &dir, &idir, &ob_itfm); + const float t_world_to_instance = bvh_instance_motion_push( + kg, object, ray, &P, &dir, &idir, &ob_itfm); #else - isect->t *= bvh_instance_push(kg, object, ray, &P, &dir, &idir); + const float t_world_to_instance = 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; @@ -235,6 +250,7 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg, #else isect->t = bvh_instance_pop(kg, object, ray, &P, &dir, &idir, isect->t); #endif + tmin = ray->tmin; object = OBJECT_NONE; node_addr = traversal_stack[stack_ptr]; diff --git a/intern/cycles/kernel/bvh/util.h b/intern/cycles/kernel/bvh/util.h index 572e023db25..1795ae4c790 100644 --- a/intern/cycles/kernel/bvh/util.h +++ b/intern/cycles/kernel/bvh/util.h @@ -5,6 +5,19 @@ CCL_NAMESPACE_BEGIN +/* Offset intersection distance by the smallest possible amount, to skip + * intersections at this distance. This works in cases where the ray start + * position is unchanged and only tmin is updated, since for self + * intersection we'll be comparing against the exact same distances. */ +ccl_device_forceinline float intersection_t_offset(const float t) +{ + /* This is a simplified version of nextafterf(t, FLT_MAX), only dealing with + * non-negative and finite t. */ + kernel_assert(t >= 0.0f && isfinite_safe(t)); + const uint32_t bits = (t == 0.0f) ? 1 : __float_as_uint(t) + 1; + return __uint_as_float(bits); +} + #if defined(__KERNEL_CPU__) ccl_device int intersections_compare(const void *a, const void *b) { diff --git a/intern/cycles/kernel/bvh/volume.h b/intern/cycles/kernel/bvh/volume.h index 9715712a8f2..bd4e508ecac 100644 --- a/intern/cycles/kernel/bvh/volume.h +++ b/intern/cycles/kernel/bvh/volume.h @@ -46,13 +46,14 @@ ccl_device_inline float3 P = ray->P; float3 dir = bvh_clamp_direction(ray->D); float3 idir = bvh_inverse_direction(dir); + float tmin = ray->tmin; int object = OBJECT_NONE; #if BVH_FEATURE(BVH_MOTION) Transform ob_itfm; #endif - isect->t = ray->t; + isect->t = ray->tmax; isect->u = 0.0f; isect->v = 0.0f; isect->prim = PRIM_NONE; @@ -73,6 +74,7 @@ ccl_device_inline dir, #endif idir, + tmin, isect->t, node_addr, visibility, @@ -140,7 +142,7 @@ ccl_device_inline continue; } triangle_intersect( - kg, isect, P, dir, isect->t, visibility, prim_object, prim, prim_addr); + kg, isect, P, dir, tmin, isect->t, visibility, prim_object, prim, prim_addr); } break; } @@ -165,6 +167,7 @@ ccl_device_inline isect, P, dir, + tmin, isect->t, ray->time, visibility, @@ -186,11 +189,15 @@ ccl_device_inline int object_flag = kernel_data_fetch(object_flag, object); if (object_flag & SD_OBJECT_HAS_VOLUME) { #if BVH_FEATURE(BVH_MOTION) - isect->t *= bvh_instance_motion_push(kg, object, ray, &P, &dir, &idir, &ob_itfm); + const float t_world_to_instance = bvh_instance_motion_push( + kg, object, ray, &P, &dir, &idir, &ob_itfm); #else - isect->t *= bvh_instance_push(kg, object, ray, &P, &dir, &idir); + const float t_world_to_instance = 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; @@ -217,6 +224,8 @@ ccl_device_inline isect->t = bvh_instance_pop(kg, object, ray, &P, &dir, &idir, isect->t); #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 d06ea8fe557..c6eeb07a14d 100644 --- a/intern/cycles/kernel/bvh/volume_all.h +++ b/intern/cycles/kernel/bvh/volume_all.h @@ -44,12 +44,12 @@ ccl_device_inline int node_addr = kernel_data.bvh.root; /* ray parameters in registers */ - const float tmax = ray->t; float3 P = ray->P; float3 dir = bvh_clamp_direction(ray->D); float3 idir = bvh_inverse_direction(dir); + float tmin = ray->tmin; int object = OBJECT_NONE; - float isect_t = tmax; + float isect_t = ray->tmax; #if BVH_FEATURE(BVH_MOTION) Transform ob_itfm; @@ -58,7 +58,7 @@ ccl_device_inline int num_hits_in_instance = 0; uint num_hits = 0; - isect_array->t = tmax; + isect_array->t = ray->tmax; /* traversal loop */ do { @@ -75,6 +75,7 @@ ccl_device_inline dir, #endif idir, + tmin, isect_t, node_addr, visibility, @@ -141,8 +142,16 @@ ccl_device_inline if ((object_flag & SD_OBJECT_HAS_VOLUME) == 0) { continue; } - hit = triangle_intersect( - kg, isect_array, P, dir, isect_t, visibility, prim_object, prim, prim_addr); + hit = triangle_intersect(kg, + isect_array, + P, + dir, + tmin, + isect_t, + visibility, + prim_object, + prim, + prim_addr); if (hit) { /* Move on to next entry in intersections array. */ isect_array++; @@ -189,6 +198,7 @@ ccl_device_inline isect_array, P, dir, + tmin, isect_t, ray->time, visibility, @@ -232,11 +242,15 @@ ccl_device_inline int object_flag = kernel_data_fetch(object_flag, object); if (object_flag & SD_OBJECT_HAS_VOLUME) { #if BVH_FEATURE(BVH_MOTION) - isect_t *= bvh_instance_motion_push(kg, object, ray, &P, &dir, &idir, &ob_itfm); + const float t_world_to_instance = bvh_instance_motion_push( + kg, object, ray, &P, &dir, &idir, &ob_itfm); #else - isect_t *= bvh_instance_push(kg, object, ray, &P, &dir, &idir); + const float t_world_to_instance = 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; @@ -280,7 +294,8 @@ ccl_device_inline #endif } - isect_t = tmax; + tmin = ray->tmin; + isect_t = ray->tmax; isect_array->t = isect_t; object = OBJECT_NONE; diff --git a/intern/cycles/kernel/camera/camera.h b/intern/cycles/kernel/camera/camera.h index 7e1b1c037e9..926ccf7b86f 100644 --- a/intern/cycles/kernel/camera/camera.h +++ b/intern/cycles/kernel/camera/camera.h @@ -165,9 +165,11 @@ ccl_device void camera_sample_perspective(KernelGlobals kg, float nearclip = kernel_data.cam.nearclip * z_inv; ray->P += nearclip * ray->D; ray->dP += nearclip * ray->dD; - ray->t = kernel_data.cam.cliplength * z_inv; + ray->tmin = 0.0f; + ray->tmax = kernel_data.cam.cliplength * z_inv; #else - ray->t = FLT_MAX; + ray->tmin = 0.0f; + ray->tmax = FLT_MAX; #endif } @@ -231,9 +233,11 @@ ccl_device void camera_sample_orthographic(KernelGlobals kg, #ifdef __CAMERA_CLIPPING__ /* clipping */ - ray->t = kernel_data.cam.cliplength; + ray->tmin = 0.0f; + ray->tmax = kernel_data.cam.cliplength; #else - ray->t = FLT_MAX; + ray->tmin = 0.0f; + ray->tmax = FLT_MAX; #endif } @@ -258,7 +262,7 @@ ccl_device_inline void camera_sample_panorama(ccl_constant KernelCamera *cam, /* indicates ray should not receive any light, outside of the lens */ if (is_zero(D)) { - ray->t = 0.0f; + ray->tmax = 0.0f; return; } @@ -349,9 +353,11 @@ ccl_device_inline void camera_sample_panorama(ccl_constant KernelCamera *cam, float nearclip = cam->nearclip; ray->P += nearclip * ray->D; ray->dP += nearclip * ray->dD; - ray->t = cam->cliplength; + ray->tmin = 0.0f; + ray->tmax = cam->cliplength; #else - ray->t = FLT_MAX; + ray->tmin = 0.0f; + ray->tmax = FLT_MAX; #endif } diff --git a/intern/cycles/kernel/device/metal/kernel.metal b/intern/cycles/kernel/device/metal/kernel.metal index 74b4b079a32..764c26dbe8f 100644 --- a/intern/cycles/kernel/device/metal/kernel.metal +++ b/intern/cycles/kernel/device/metal/kernel.metal @@ -410,6 +410,7 @@ void metalrt_intersection_curve(constant KernelParamsMetal &launch_params_metal, const float3 ray_origin, const float3 ray_direction, float time, + const float ray_tmin, const float ray_tmax, thread BoundingBoxIntersectionResult &result) { @@ -434,7 +435,7 @@ void metalrt_intersection_curve(constant KernelParamsMetal &launch_params_metal, isect.t *= len; MetalKernelContext context(launch_params_metal); - if (context.curve_intersect(NULL, &isect, P, dir, isect.t, object, prim, time, type)) { + if (context.curve_intersect(NULL, &isect, P, dir, 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) { @@ -456,6 +457,7 @@ void metalrt_intersection_curve_shadow(constant KernelParamsMetal &launch_params const float3 ray_origin, const float3 ray_direction, float time, + const float ray_tmin, const float ray_tmax, thread BoundingBoxIntersectionResult &result) { @@ -475,7 +477,7 @@ void metalrt_intersection_curve_shadow(constant KernelParamsMetal &launch_params isect.t *= len; MetalKernelContext context(launch_params_metal); - if (context.curve_intersect(NULL, &isect, P, dir, isect.t, object, prim, time, type)) { + if (context.curve_intersect(NULL, &isect, P, dir, 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; @@ -494,6 +496,7 @@ __intersection__curve_ribbon(constant KernelParamsMetal &launch_params_metal [[b const uint primitive_id [[primitive_id]], const float3 ray_origin [[origin]], const float3 ray_direction [[direction]], + const float ray_tmin [[min_distance]], const float ray_tmax [[max_distance]]) { uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object); @@ -511,7 +514,7 @@ __intersection__curve_ribbon(constant KernelParamsMetal &launch_params_metal [[b # else 0.0f, # endif - ray_tmax, result); + ray_tmin, ray_tmax, result); } return result; @@ -525,6 +528,7 @@ __intersection__curve_ribbon_shadow(constant KernelParamsMetal &launch_params_me const uint primitive_id [[primitive_id]], const float3 ray_origin [[origin]], const float3 ray_direction [[direction]], + const float ray_tmin [[min_distance]], const float ray_tmax [[max_distance]]) { uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object); @@ -542,7 +546,7 @@ __intersection__curve_ribbon_shadow(constant KernelParamsMetal &launch_params_me # else 0.0f, # endif - ray_tmax, result); + ray_tmin, ray_tmax, result); } return result; @@ -556,6 +560,7 @@ __intersection__curve_all(constant KernelParamsMetal &launch_params_metal [[buff const uint primitive_id [[primitive_id]], const float3 ray_origin [[origin]], const float3 ray_direction [[direction]], + const float ray_tmin [[min_distance]], const float ray_tmax [[max_distance]]) { uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object); @@ -571,7 +576,7 @@ __intersection__curve_all(constant KernelParamsMetal &launch_params_metal [[buff # else 0.0f, # endif - ray_tmax, result); + ray_tmin, ray_tmax, result); return result; } @@ -584,6 +589,7 @@ __intersection__curve_all_shadow(constant KernelParamsMetal &launch_params_metal const uint primitive_id [[primitive_id]], const float3 ray_origin [[origin]], const float3 ray_direction [[direction]], + const float ray_tmin [[min_distance]], const float ray_tmax [[max_distance]]) { uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object); @@ -600,7 +606,7 @@ __intersection__curve_all_shadow(constant KernelParamsMetal &launch_params_metal # else 0.0f, # endif - ray_tmax, result); + ray_tmin, ray_tmax, result); return result; } @@ -616,6 +622,7 @@ void metalrt_intersection_point(constant KernelParamsMetal &launch_params_metal, const float3 ray_origin, const float3 ray_direction, float time, + const float ray_tmin, const float ray_tmax, thread BoundingBoxIntersectionResult &result) { @@ -640,7 +647,7 @@ void metalrt_intersection_point(constant KernelParamsMetal &launch_params_metal, isect.t *= len; MetalKernelContext context(launch_params_metal); - if (context.point_intersect(NULL, &isect, P, dir, isect.t, object, prim, time, type)) { + if (context.point_intersect(NULL, &isect, P, dir, 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) { @@ -662,6 +669,7 @@ void metalrt_intersection_point_shadow(constant KernelParamsMetal &launch_params const float3 ray_origin, const float3 ray_direction, float time, + const float ray_tmin, const float ray_tmax, thread BoundingBoxIntersectionResult &result) { @@ -681,7 +689,7 @@ void metalrt_intersection_point_shadow(constant KernelParamsMetal &launch_params isect.t *= len; MetalKernelContext context(launch_params_metal); - if (context.point_intersect(NULL, &isect, P, dir, isect.t, object, prim, time, type)) { + if (context.point_intersect(NULL, &isect, P, dir, 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; @@ -700,6 +708,7 @@ __intersection__point(constant KernelParamsMetal &launch_params_metal [[buffer(1 const uint primitive_id [[primitive_id]], const float3 ray_origin [[origin]], const float3 ray_direction [[direction]], + const float ray_tmin [[min_distance]], const float ray_tmax [[max_distance]]) { const uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object); @@ -716,7 +725,7 @@ __intersection__point(constant KernelParamsMetal &launch_params_metal [[buffer(1 # else 0.0f, # endif - ray_tmax, result); + ray_tmin, ray_tmax, result); return result; } @@ -729,6 +738,7 @@ __intersection__point_shadow(constant KernelParamsMetal &launch_params_metal [[b const uint primitive_id [[primitive_id]], const float3 ray_origin [[origin]], const float3 ray_direction [[direction]], + const float ray_tmin [[min_distance]], const float ray_tmax [[max_distance]]) { const uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object); @@ -745,7 +755,7 @@ __intersection__point_shadow(constant KernelParamsMetal &launch_params_metal [[b # else 0.0f, # endif - ray_tmax, result); + ray_tmin, ray_tmax, result); return result; } diff --git a/intern/cycles/kernel/device/optix/kernel.cu b/intern/cycles/kernel/device/optix/kernel.cu index 949bf41d171..510f7cca5d6 100644 --- a/intern/cycles/kernel/device/optix/kernel.cu +++ b/intern/cycles/kernel/device/optix/kernel.cu @@ -51,32 +51,36 @@ ccl_device_forceinline int get_object_id() extern "C" __global__ void __raygen__kernel_optix_integrator_intersect_closest() { const int global_index = optixGetLaunchIndex().x; - const int path_index = (kernel_params.path_index_array) ? kernel_params.path_index_array[global_index] : - global_index; + const int path_index = (kernel_params.path_index_array) ? + kernel_params.path_index_array[global_index] : + global_index; integrator_intersect_closest(nullptr, path_index, kernel_params.render_buffer); } extern "C" __global__ void __raygen__kernel_optix_integrator_intersect_shadow() { const int global_index = optixGetLaunchIndex().x; - const int path_index = (kernel_params.path_index_array) ? kernel_params.path_index_array[global_index] : - global_index; + const int path_index = (kernel_params.path_index_array) ? + kernel_params.path_index_array[global_index] : + global_index; integrator_intersect_shadow(nullptr, path_index); } extern "C" __global__ void __raygen__kernel_optix_integrator_intersect_subsurface() { const int global_index = optixGetLaunchIndex().x; - const int path_index = (kernel_params.path_index_array) ? kernel_params.path_index_array[global_index] : - global_index; + const int path_index = (kernel_params.path_index_array) ? + kernel_params.path_index_array[global_index] : + global_index; integrator_intersect_subsurface(nullptr, path_index); } extern "C" __global__ void __raygen__kernel_optix_integrator_intersect_volume_stack() { const int global_index = optixGetLaunchIndex().x; - const int path_index = (kernel_params.path_index_array) ? kernel_params.path_index_array[global_index] : - global_index; + const int path_index = (kernel_params.path_index_array) ? + kernel_params.path_index_array[global_index] : + global_index; integrator_intersect_volume_stack(nullptr, path_index); } @@ -408,6 +412,7 @@ ccl_device_inline void optix_intersection_curve(const int prim, const int type) 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; @@ -425,7 +430,7 @@ ccl_device_inline void optix_intersection_curve(const int prim, const int type) if (isect.t != FLT_MAX) isect.t *= len; - if (curve_intersect(NULL, &isect, P, dir, isect.t, object, prim, time, type)) { + if (curve_intersect(NULL, &isect, P, dir, 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, @@ -462,6 +467,7 @@ extern "C" __global__ void __intersection__point() float3 P = optixGetObjectRayOrigin(); float3 dir = optixGetObjectRayDirection(); + float tmin = optixGetRayTmin(); /* The direction is not normalized by default, the point intersection routine expects that. */ float len; @@ -480,7 +486,7 @@ extern "C" __global__ void __intersection__point() isect.t *= len; } - if (point_intersect(NULL, &isect, P, dir, isect.t, object, prim, time, type)) { + if (point_intersect(NULL, &isect, P, dir, 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); } diff --git a/intern/cycles/kernel/geom/curve_intersect.h b/intern/cycles/kernel/geom/curve_intersect.h index 001bec01749..9770105dd81 100644 --- a/intern/cycles/kernel/geom/curve_intersect.h +++ b/intern/cycles/kernel/geom/curve_intersect.h @@ -156,7 +156,8 @@ ccl_device_inline float2 half_plane_intersect(const float3 P, const float3 N, co } ccl_device bool curve_intersect_iterative(const float3 ray_dir, - ccl_private float *ray_tfar, + const float ray_tmin, + ccl_private float *ray_tmax, const float dt, const float4 curve[4], float u, @@ -220,7 +221,7 @@ ccl_device bool curve_intersect_iterative(const float3 ray_dir, if (fabsf(f) < f_err && fabsf(g) < g_err) { t += dt; - if (!(0.0f <= t && t <= *ray_tfar)) { + if (!(t >= ray_tmin && t <= *ray_tmax)) { return false; /* Rejects NaNs */ } if (!(u >= 0.0f && u <= 1.0f)) { @@ -237,7 +238,7 @@ ccl_device bool curve_intersect_iterative(const float3 ray_dir, } /* Record intersection. */ - *ray_tfar = t; + *ray_tmax = t; isect->t = t; isect->u = u; isect->v = 0.0f; @@ -250,7 +251,8 @@ ccl_device bool curve_intersect_iterative(const float3 ray_dir, ccl_device bool curve_intersect_recursive(const float3 ray_orig, const float3 ray_dir, - float ray_tfar, + const float ray_tmin, + float ray_tmax, float4 curve[4], ccl_private Intersection *isect) { @@ -331,7 +333,7 @@ ccl_device bool curve_intersect_recursive(const float3 ray_orig, } /* Intersect with cap-planes. */ - float2 tp = make_float2(-dt, ray_tfar - dt); + 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); @@ -394,19 +396,20 @@ ccl_device bool curve_intersect_recursive(const float3 ray_orig, CURVE_NUM_BEZIER_SUBDIVISIONS; if (depth >= termDepth) { found |= curve_intersect_iterative( - ray_dir, &ray_tfar, dt, curve, u_outer0, tp0.x, use_backfacing, isect); + ray_dir, ray_tmin, &ray_tmax, dt, curve, u_outer0, tp0.x, use_backfacing, isect); } else { recurse = true; } } - if (valid1 && (tp1.x + dt <= ray_tfar)) { + const float t1 = tp1.x + dt; + if (valid1 && (t1 >= ray_tmin && t1 <= ray_tmax)) { const int termDepth = unstable1 ? CURVE_NUM_BEZIER_SUBDIVISIONS_UNSTABLE : CURVE_NUM_BEZIER_SUBDIVISIONS; if (depth >= termDepth) { found |= curve_intersect_iterative( - ray_dir, &ray_tfar, dt, curve, u_outer1, tp1.y, use_backfacing, isect); + ray_dir, ray_tmin, &ray_tmax, dt, curve, u_outer1, tp1.y, use_backfacing, isect); } else { recurse = true; @@ -456,7 +459,8 @@ ccl_device_inline bool cylinder_culling_test(const float2 p1, const float2 p2, c * v0,v1,v3 and v2,v3,v1. The edge v1,v2 decides which of the two * triangles gets intersected. */ -ccl_device_inline bool ribbon_intersect_quad(const float ray_tfar, +ccl_device_inline bool ribbon_intersect_quad(const float ray_tmin, + const float ray_tmax, const float3 quad_v0, const float3 quad_v1, const float3 quad_v2, @@ -497,7 +501,7 @@ ccl_device_inline bool ribbon_intersect_quad(const float ray_tfar, /* Perform depth test? */ const float t = rcpDen * dot(v0, Ng); - if (!(0.0f <= t && t <= ray_tfar)) { + if (!(t >= ray_tmin && t <= ray_tmax)) { return false; } @@ -534,7 +538,8 @@ 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, - float ray_tfar, + const float ray_tmin, + float ray_tmax, const int N, float4 curve[4], ccl_private Intersection *isect) @@ -582,7 +587,7 @@ ccl_device_inline bool ribbon_intersect(const float3 ray_org, /* Intersect quad. */ float vu, vv, vt; - bool valid0 = ribbon_intersect_quad(ray_tfar, lp0, lp1, up1, up0, &vu, &vv, &vt); + bool valid0 = ribbon_intersect_quad(ray_tmin, ray_tmax, lp0, lp1, up1, up0, &vu, &vv, &vt); if (valid0) { /* ignore self intersections */ @@ -596,7 +601,7 @@ ccl_device_inline bool ribbon_intersect(const float3 ray_org, vv = 2.0f * vv - 1.0f; /* Record intersection. */ - ray_tfar = vt; + ray_tmax = vt; isect->t = vt; isect->u = u + vu * step_size; isect->v = vv; @@ -616,6 +621,7 @@ ccl_device_forceinline bool curve_intersect(KernelGlobals kg, ccl_private Intersection *isect, const float3 P, const float3 dir, + const float tmin, const float tmax, int object, int prim, @@ -645,7 +651,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, tmax, subdivisions, curve, isect)) { + if (ribbon_intersect(P, dir, tmin, tmax, subdivisions, curve, isect)) { isect->prim = prim; isect->object = object; isect->type = type; @@ -655,7 +661,7 @@ ccl_device_forceinline bool curve_intersect(KernelGlobals kg, return false; } else { - if (curve_intersect_recursive(P, dir, tmax, curve, isect)) { + if (curve_intersect_recursive(P, dir, tmin, tmax, curve, isect)) { isect->prim = prim; isect->object = object; isect->type = type; diff --git a/intern/cycles/kernel/geom/motion_triangle_intersect.h b/intern/cycles/kernel/geom/motion_triangle_intersect.h index 6eea5096567..b59c5c43c20 100644 --- a/intern/cycles/kernel/geom/motion_triangle_intersect.h +++ b/intern/cycles/kernel/geom/motion_triangle_intersect.h @@ -46,6 +46,7 @@ ccl_device_inline bool motion_triangle_intersect(KernelGlobals kg, ccl_private Intersection *isect, float3 P, float3 dir, + float tmin, float tmax, float time, uint visibility, @@ -58,7 +59,7 @@ ccl_device_inline bool motion_triangle_intersect(KernelGlobals kg, motion_triangle_vertices(kg, object, prim, time, verts); /* Ray-triangle intersection, unoptimized. */ float t, u, v; - if (ray_triangle_intersect(P, dir, tmax, verts[0], verts[1], verts[2], &u, &v, &t)) { + if (ray_triangle_intersect(P, dir, tmin, tmax, verts[0], verts[1], verts[2], &u, &v, &t)) { #ifdef __VISIBILITY_FLAG__ /* Visibility flag test. we do it here under the assumption * that most triangles are culled by node flags. @@ -92,6 +93,7 @@ ccl_device_inline bool motion_triangle_intersect_local(KernelGlobals kg, int object, int prim, int prim_addr, + float tmin, float tmax, ccl_private uint *lcg_state, int max_hits) @@ -101,7 +103,7 @@ ccl_device_inline bool motion_triangle_intersect_local(KernelGlobals kg, motion_triangle_vertices(kg, object, prim, time, verts); /* Ray-triangle intersection, unoptimized. */ float t, u, v; - if (!ray_triangle_intersect(P, dir, tmax, verts[0], verts[1], verts[2], &u, &v, &t)) { + if (!ray_triangle_intersect(P, dir, tmin, tmax, verts[0], verts[1], verts[2], &u, &v, &t)) { return false; } diff --git a/intern/cycles/kernel/geom/point_intersect.h b/intern/cycles/kernel/geom/point_intersect.h index dfd9d9a015b..ee5a564947b 100644 --- a/intern/cycles/kernel/geom/point_intersect.h +++ b/intern/cycles/kernel/geom/point_intersect.h @@ -9,8 +9,12 @@ CCL_NAMESPACE_BEGIN #ifdef __POINTCLOUD__ -ccl_device_forceinline bool point_intersect_test( - const float4 point, const float3 P, const float3 dir, const float tmax, ccl_private float *t) +ccl_device_forceinline bool point_intersect_test(const float4 point, + const float3 P, + const float3 dir, + const float tmin, + const float tmax, + ccl_private float *t) { const float3 center = float4_to_float3(point); const float radius = point.w; @@ -28,12 +32,12 @@ ccl_device_forceinline bool point_intersect_test( const float td = sqrt((r2 - l2) * rd2); const float t_front = projC0 - td; - const bool valid_front = (0.0f <= t_front) & (t_front <= tmax); + const bool valid_front = (tmin <= t_front) & (t_front <= tmax); /* Always back-face culling for now. */ # if 0 const float t_back = projC0 + td; - const bool valid_back = (0.0f <= t_back) & (t_back <= tmax); + const bool valid_back = (tmin <= t_back) & (t_back <= tmax); /* check if there is a first hit */ const bool valid_first = valid_front | valid_back; @@ -56,6 +60,7 @@ 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 int object, const int prim, @@ -65,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, tmax, &isect->t)) { + if (!point_intersect_test(point, P, dir, tmin, tmax, &isect->t)) { return false; } diff --git a/intern/cycles/kernel/geom/shader_data.h b/intern/cycles/kernel/geom/shader_data.h index e5dbeac5e66..99b9289cb4a 100644 --- a/intern/cycles/kernel/geom/shader_data.h +++ b/intern/cycles/kernel/geom/shader_data.h @@ -407,7 +407,7 @@ ccl_device_inline void shader_setup_from_volume(KernelGlobals kg, { /* vectors */ - sd->P = ray->P; + sd->P = ray->P + ray->D * ray->tmin; sd->N = -ray->D; sd->Ng = -ray->D; sd->I = -ray->D; @@ -441,7 +441,6 @@ ccl_device_inline void shader_setup_from_volume(KernelGlobals kg, /* for NDC coordinates */ sd->ray_P = ray->P; - sd->ray_dP = ray->dP; } #endif /* __VOLUME__ */ diff --git a/intern/cycles/kernel/geom/triangle_intersect.h b/intern/cycles/kernel/geom/triangle_intersect.h index 0c76de9ccc7..f968e537cfa 100644 --- a/intern/cycles/kernel/geom/triangle_intersect.h +++ b/intern/cycles/kernel/geom/triangle_intersect.h @@ -17,6 +17,7 @@ ccl_device_inline bool triangle_intersect(KernelGlobals kg, ccl_private Intersection *isect, float3 P, float3 dir, + float tmin, float tmax, uint visibility, int object, @@ -28,7 +29,7 @@ ccl_device_inline bool triangle_intersect(KernelGlobals kg, tri_b = kernel_data_fetch(tri_verts, tri_vindex + 1), tri_c = kernel_data_fetch(tri_verts, tri_vindex + 2); float t, u, v; - if (ray_triangle_intersect(P, dir, tmax, tri_a, tri_b, tri_c, &u, &v, &t)) { + if (ray_triangle_intersect(P, dir, tmin, tmax, tri_a, tri_b, tri_c, &u, &v, &t)) { #ifdef __VISIBILITY_FLAG__ /* Visibility flag test. we do it here under the assumption * that most triangles are culled by node flags. @@ -62,6 +63,7 @@ ccl_device_inline bool triangle_intersect_local(KernelGlobals kg, int object, int prim, int prim_addr, + float tmin, float tmax, ccl_private uint *lcg_state, int max_hits) @@ -71,7 +73,7 @@ ccl_device_inline bool triangle_intersect_local(KernelGlobals kg, tri_b = kernel_data_fetch(tri_verts, tri_vindex + 1), tri_c = kernel_data_fetch(tri_verts, tri_vindex + 2); float t, u, v; - if (!ray_triangle_intersect(P, dir, tmax, tri_a, tri_b, tri_c, &u, &v, &t)) { + if (!ray_triangle_intersect(P, dir, tmin, tmax, tri_a, tri_b, tri_c, &u, &v, &t)) { return false; } diff --git a/intern/cycles/kernel/integrator/init_from_bake.h b/intern/cycles/kernel/integrator/init_from_bake.h index afd174de9e8..bf3f41b52b9 100644 --- a/intern/cycles/kernel/integrator/init_from_bake.h +++ b/intern/cycles/kernel/integrator/init_from_bake.h @@ -174,7 +174,8 @@ ccl_device bool integrator_init_from_bake(KernelGlobals kg, Ray ray ccl_optional_struct_init; ray.P = zero_float3(); ray.D = normalize(P); - ray.t = FLT_MAX; + ray.tmin = 0.0f; + ray.tmax = FLT_MAX; ray.time = 0.5f; ray.dP = differential_zero_compact(); ray.dD = differential_zero_compact(); @@ -210,7 +211,8 @@ ccl_device bool integrator_init_from_bake(KernelGlobals kg, Ray ray ccl_optional_struct_init; ray.P = P + N; ray.D = -N; - ray.t = FLT_MAX; + ray.tmin = 0.0f; + ray.tmax = FLT_MAX; ray.time = 0.5f; /* Setup differentials. */ diff --git a/intern/cycles/kernel/integrator/init_from_camera.h b/intern/cycles/kernel/integrator/init_from_camera.h index 73db13be697..e89ab3991c7 100644 --- a/intern/cycles/kernel/integrator/init_from_camera.h +++ b/intern/cycles/kernel/integrator/init_from_camera.h @@ -86,7 +86,7 @@ ccl_device bool integrator_init_from_camera(KernelGlobals kg, /* Generate camera ray. */ Ray ray; integrate_camera_sample(kg, sample, x, y, rng_hash, &ray); - if (ray.t == 0.0f) { + if (ray.tmax == 0.0f) { return true; } diff --git a/intern/cycles/kernel/integrator/intersect_closest.h b/intern/cycles/kernel/integrator/intersect_closest.h index 923dab9591a..60299f2cb2f 100644 --- a/intern/cycles/kernel/integrator/intersect_closest.h +++ b/intern/cycles/kernel/integrator/intersect_closest.h @@ -324,7 +324,7 @@ ccl_device void integrator_intersect_closest(KernelGlobals kg, /* Read ray from integrator state into local memory. */ Ray ray ccl_optional_struct_init; integrator_state_read_ray(kg, state, &ray); - kernel_assert(ray.t != 0.0f); + kernel_assert(ray.tmax != 0.0f); const uint visibility = path_state_ray_visibility(state); const int last_isect_prim = INTEGRATOR_STATE(state, isect, prim); @@ -332,12 +332,12 @@ ccl_device void integrator_intersect_closest(KernelGlobals kg, /* Trick to use short AO rays to approximate indirect light at the end of the path. */ if (path_state_ao_bounce(kg, state)) { - ray.t = kernel_data.integrator.ao_bounces_distance; + ray.tmax = kernel_data.integrator.ao_bounces_distance; if (last_isect_object != OBJECT_NONE) { const float object_ao_distance = kernel_data_fetch(objects, last_isect_object).ao_distance; if (object_ao_distance != 0.0f) { - ray.t = object_ao_distance; + ray.tmax = object_ao_distance; } } } diff --git a/intern/cycles/kernel/integrator/intersect_volume_stack.h b/intern/cycles/kernel/integrator/intersect_volume_stack.h index 78f0b4d62aa..9ba4a0a3964 100644 --- a/intern/cycles/kernel/integrator/intersect_volume_stack.h +++ b/intern/cycles/kernel/integrator/intersect_volume_stack.h @@ -24,7 +24,8 @@ ccl_device void integrator_volume_stack_update_for_subsurface(KernelGlobals kg, Ray volume_ray ccl_optional_struct_init; volume_ray.P = from_P; - volume_ray.D = normalize_len(to_P - from_P, &volume_ray.t); + volume_ray.D = normalize_len(to_P - from_P, &volume_ray.tmax); + volume_ray.tmin = 0.0f; volume_ray.self.object = INTEGRATOR_STATE(state, isect, object); volume_ray.self.prim = INTEGRATOR_STATE(state, isect, prim); volume_ray.self.light_object = OBJECT_NONE; @@ -58,12 +59,9 @@ ccl_device void integrator_volume_stack_update_for_subsurface(KernelGlobals kg, volume_stack_enter_exit(kg, state, stack_sd); /* Move ray forward. */ - volume_ray.P = stack_sd->P; + volume_ray.tmin = intersection_t_offset(isect.t); volume_ray.self.object = isect.object; volume_ray.self.prim = isect.prim; - if (volume_ray.t != FLT_MAX) { - volume_ray.D = normalize_len(to_P - volume_ray.P, &volume_ray.t); - } ++step; } #endif @@ -82,7 +80,8 @@ ccl_device void integrator_volume_stack_init(KernelGlobals kg, IntegratorState s /* Trace ray in random direction. Any direction works, Z up is a guess to get the * fewest hits. */ volume_ray.D = make_float3(0.0f, 0.0f, 1.0f); - volume_ray.t = FLT_MAX; + volume_ray.tmin = 0.0f; + volume_ray.tmax = FLT_MAX; volume_ray.self.object = OBJECT_NONE; volume_ray.self.prim = PRIM_NONE; volume_ray.self.light_object = OBJECT_NONE; @@ -199,7 +198,7 @@ ccl_device void integrator_volume_stack_init(KernelGlobals kg, IntegratorState s } /* Move ray forward. */ - volume_ray.P = stack_sd->P; + volume_ray.tmin = intersection_t_offset(isect.t); volume_ray.self.object = isect.object; volume_ray.self.prim = isect.prim; ++step; diff --git a/intern/cycles/kernel/integrator/mnee.h b/intern/cycles/kernel/integrator/mnee.h index 70b009d3b48..7a6f866b1a0 100644 --- a/intern/cycles/kernel/integrator/mnee.h +++ b/intern/cycles/kernel/integrator/mnee.h @@ -442,6 +442,7 @@ ccl_device_forceinline bool mnee_newton_solver(KernelGlobals kg, projection_ray.self.light_prim = PRIM_NONE; projection_ray.dP = differential_make_compact(sd->dP); projection_ray.dD = differential_zero_compact(); + projection_ray.tmin = 0.0f; projection_ray.time = sd->time; Intersection projection_isect; @@ -505,8 +506,8 @@ ccl_device_forceinline bool mnee_newton_solver(KernelGlobals kg, projection_ray.self.prim = pv.prim; projection_ray.P = pv.p; } - projection_ray.D = normalize_len(tentative_p - projection_ray.P, &projection_ray.t); - projection_ray.t *= MNEE_PROJECTION_DISTANCE_MULTIPLIER; + projection_ray.D = normalize_len(tentative_p - projection_ray.P, &projection_ray.tmax); + projection_ray.tmax *= MNEE_PROJECTION_DISTANCE_MULTIPLIER; bool projection_success = false; for (int isect_count = 0; isect_count < MNEE_MAX_INTERSECTION_COUNT; isect_count++) { @@ -525,8 +526,7 @@ ccl_device_forceinline bool mnee_newton_solver(KernelGlobals kg, projection_ray.self.object = projection_isect.object; projection_ray.self.prim = projection_isect.prim; - projection_ray.P += projection_isect.t * projection_ray.D; - projection_ray.t -= projection_isect.t; + projection_ray.tmin = intersection_t_offset(projection_isect.t); } if (!projection_success) { reduce_stepsize = true; @@ -858,6 +858,7 @@ ccl_device_forceinline bool mnee_path_contribution(KernelGlobals kg, Ray probe_ray; probe_ray.self.light_object = ls->object; probe_ray.self.light_prim = ls->prim; + probe_ray.tmin = 0.0f; probe_ray.dP = differential_make_compact(sd->dP); probe_ray.dD = differential_zero_compact(); probe_ray.time = sd->time; @@ -873,13 +874,13 @@ ccl_device_forceinline bool mnee_path_contribution(KernelGlobals kg, ccl_private const ManifoldVertex &v = vertices[vi]; /* Check visibility. */ - probe_ray.D = normalize_len(v.p - probe_ray.P, &probe_ray.t); + probe_ray.D = normalize_len(v.p - probe_ray.P, &probe_ray.tmax); if (scene_intersect(kg, &probe_ray, PATH_RAY_TRANSMIT, &probe_isect)) { int hit_object = (probe_isect.object == OBJECT_NONE) ? kernel_data_fetch(prim_object, probe_isect.prim) : probe_isect.object; /* Test whether the ray hit the appropriate object at its intended location. */ - if (hit_object != v.object || fabsf(probe_ray.t - probe_isect.t) > MNEE_MIN_DISTANCE) + if (hit_object != v.object || fabsf(probe_ray.tmax - probe_isect.t) > MNEE_MIN_DISTANCE) return false; } probe_ray.self.object = v.object; @@ -958,15 +959,16 @@ ccl_device_forceinline int kernel_path_mnee_sample(KernelGlobals kg, probe_ray.self.light_object = ls->object; probe_ray.self.light_prim = ls->prim; probe_ray.P = sd->P; + probe_ray.tmin = 0.0f; if (ls->t == FLT_MAX) { /* Distant / env light. */ probe_ray.D = ls->D; - probe_ray.t = ls->t; + probe_ray.tmax = ls->t; } else { /* Other lights, avoid self-intersection. */ probe_ray.D = ls->P - probe_ray.P; - probe_ray.D = normalize_len(probe_ray.D, &probe_ray.t); + probe_ray.D = normalize_len(probe_ray.D, &probe_ray.tmax); } probe_ray.dP = differential_make_compact(sd->dP); probe_ray.dD = differential_zero_compact(); @@ -1048,9 +1050,7 @@ ccl_device_forceinline int kernel_path_mnee_sample(KernelGlobals kg, probe_ray.self.object = probe_isect.object; probe_ray.self.prim = probe_isect.prim; - probe_ray.P += probe_isect.t * probe_ray.D; - if (ls->t != FLT_MAX) - probe_ray.t -= probe_isect.t; + probe_ray.tmin = intersection_t_offset(probe_isect.t); }; /* Mark the manifold walk invalid to keep mollification on by default. */ diff --git a/intern/cycles/kernel/integrator/path_state.h b/intern/cycles/kernel/integrator/path_state.h index 1a085506a70..912c380cdb6 100644 --- a/intern/cycles/kernel/integrator/path_state.h +++ b/intern/cycles/kernel/integrator/path_state.h @@ -52,7 +52,6 @@ ccl_device_inline void path_state_init_integrator(KernelGlobals kg, INTEGRATOR_STATE_WRITE(state, path, flag) = PATH_RAY_CAMERA | PATH_RAY_MIS_SKIP | PATH_RAY_TRANSPARENT_BACKGROUND; INTEGRATOR_STATE_WRITE(state, path, mis_ray_pdf) = 0.0f; - INTEGRATOR_STATE_WRITE(state, path, mis_ray_t) = 0.0f; INTEGRATOR_STATE_WRITE(state, path, min_ray_pdf) = FLT_MAX; INTEGRATOR_STATE_WRITE(state, path, continuation_probability) = 1.0f; INTEGRATOR_STATE_WRITE(state, path, throughput) = make_float3(1.0f, 1.0f, 1.0f); diff --git a/intern/cycles/kernel/integrator/shade_background.h b/intern/cycles/kernel/integrator/shade_background.h index 47233634463..a7edfffd175 100644 --- a/intern/cycles/kernel/integrator/shade_background.h +++ b/intern/cycles/kernel/integrator/shade_background.h @@ -62,11 +62,10 @@ ccl_device float3 integrator_eval_background_shader(KernelGlobals kg, const float3 ray_P = INTEGRATOR_STATE(state, ray, P); const float3 ray_D = INTEGRATOR_STATE(state, ray, D); const float mis_ray_pdf = INTEGRATOR_STATE(state, path, mis_ray_pdf); - const float mis_ray_t = INTEGRATOR_STATE(state, path, mis_ray_t); /* multiple importance sampling, get background light pdf for ray * direction, and compute weight with respect to BSDF pdf */ - const float pdf = background_light_pdf(kg, ray_P - ray_D * mis_ray_t, ray_D); + const float pdf = background_light_pdf(kg, ray_P, ray_D); const float mis_weight = light_sample_mis_weight_forward(kg, mis_ray_pdf, pdf); L *= mis_weight; } diff --git a/intern/cycles/kernel/integrator/shade_light.h b/intern/cycles/kernel/integrator/shade_light.h index 1a5ac3637f6..910e3383f51 100644 --- a/intern/cycles/kernel/integrator/shade_light.h +++ b/intern/cycles/kernel/integrator/shade_light.h @@ -22,19 +22,8 @@ ccl_device_inline void integrate_light(KernelGlobals kg, const float3 ray_D = INTEGRATOR_STATE(state, ray, D); const float ray_time = INTEGRATOR_STATE(state, ray, time); - /* Advance ray beyond light. */ - /* TODO: can we make this more numerically robust to avoid reintersecting the - * same light in some cases? Ray should not intersect surface anymore as the - * object and prim ids will prevent self intersection. */ - const float3 new_ray_P = ray_P + ray_D * isect.t; - INTEGRATOR_STATE_WRITE(state, ray, P) = new_ray_P; - INTEGRATOR_STATE_WRITE(state, ray, t) -= isect.t; - - /* Set position to where the BSDF was sampled, for correct MIS PDF. */ - const float mis_ray_t = INTEGRATOR_STATE(state, path, mis_ray_t); - ray_P -= ray_D * mis_ray_t; - isect.t += mis_ray_t; - INTEGRATOR_STATE_WRITE(state, path, mis_ray_t) = isect.t; + /* Advance ray to new start distance. */ + INTEGRATOR_STATE_WRITE(state, ray, tmin) = intersection_t_offset(isect.t); LightSample ls ccl_optional_struct_init; const bool use_light_sample = light_sample_from_intersection(kg, &isect, ray_P, ray_D, &ls); diff --git a/intern/cycles/kernel/integrator/shade_shadow.h b/intern/cycles/kernel/integrator/shade_shadow.h index 830da0158cf..4b002a47bee 100644 --- a/intern/cycles/kernel/integrator/shade_shadow.h +++ b/intern/cycles/kernel/integrator/shade_shadow.h @@ -75,13 +75,9 @@ ccl_device_inline void integrate_transparent_volume_shadow(KernelGlobals kg, ray.self.light_object = OBJECT_NONE; ray.self.light_prim = PRIM_NONE; /* Modify ray position and length to match current segment. */ - const float start_t = (hit == 0) ? 0.0f : - INTEGRATOR_STATE_ARRAY(state, shadow_isect, hit - 1, t); - const float end_t = (hit < num_recorded_hits) ? - INTEGRATOR_STATE_ARRAY(state, shadow_isect, hit, t) : - ray.t; - ray.P += start_t * ray.D; - ray.t = end_t - start_t; + ray.tmin = (hit == 0) ? ray.tmin : INTEGRATOR_STATE_ARRAY(state, shadow_isect, hit - 1, t); + ray.tmax = (hit < num_recorded_hits) ? INTEGRATOR_STATE_ARRAY(state, shadow_isect, hit, t) : + ray.tmax; shader_setup_from_volume(kg, shadow_sd, &ray); @@ -137,10 +133,7 @@ ccl_device_inline bool integrate_transparent_shadow(KernelGlobals kg, /* There are more hits that we could not recorded due to memory usage, * adjust ray to intersect again from the last hit. */ const float last_hit_t = INTEGRATOR_STATE_ARRAY(state, shadow_isect, num_recorded_hits - 1, t); - const float3 ray_P = INTEGRATOR_STATE(state, shadow_ray, P); - const float3 ray_D = INTEGRATOR_STATE(state, shadow_ray, D); - INTEGRATOR_STATE_WRITE(state, shadow_ray, P) = ray_P + last_hit_t * ray_D; - INTEGRATOR_STATE_WRITE(state, shadow_ray, t) -= last_hit_t; + INTEGRATOR_STATE_WRITE(state, shadow_ray, tmin) = intersection_t_offset(last_hit_t); } return false; diff --git a/intern/cycles/kernel/integrator/shade_surface.h b/intern/cycles/kernel/integrator/shade_surface.h index 91e34968148..1514b3956ad 100644 --- a/intern/cycles/kernel/integrator/shade_surface.h +++ b/intern/cycles/kernel/integrator/shade_surface.h @@ -77,7 +77,7 @@ ccl_device_forceinline void integrate_surface_emission(KernelGlobals kg, # endif { const float bsdf_pdf = INTEGRATOR_STATE(state, path, mis_ray_pdf); - const float t = sd->ray_length + INTEGRATOR_STATE(state, path, mis_ray_t); + const float t = sd->ray_length; /* Multiple importance sampling, get triangle light pdf, * and compute weight with respect to BSDF pdf. */ @@ -323,16 +323,21 @@ ccl_device_forceinline int integrate_surface_bsdf_bssrdf_bounce( return LABEL_NONE; } - /* Setup ray. Note that clipping works through transparent bounces. */ - INTEGRATOR_STATE_WRITE(state, ray, P) = sd->P; - INTEGRATOR_STATE_WRITE(state, ray, D) = normalize(bsdf_omega_in); - INTEGRATOR_STATE_WRITE(state, ray, t) = (label & LABEL_TRANSPARENT) ? - INTEGRATOR_STATE(state, ray, t) - sd->ray_length : - FLT_MAX; + if (label & LABEL_TRANSPARENT) { + /* Only need to modify start distance for transparent. */ + INTEGRATOR_STATE_WRITE(state, ray, tmin) = intersection_t_offset(sd->ray_length); + } + else { + /* Setup ray with changed origin and direction. */ + INTEGRATOR_STATE_WRITE(state, ray, P) = sd->P; + INTEGRATOR_STATE_WRITE(state, ray, D) = normalize(bsdf_omega_in); + INTEGRATOR_STATE_WRITE(state, ray, tmin) = 0.0f; + INTEGRATOR_STATE_WRITE(state, ray, tmax) = FLT_MAX; #ifdef __RAY_DIFFERENTIALS__ - INTEGRATOR_STATE_WRITE(state, ray, dP) = differential_make_compact(sd->dP); - INTEGRATOR_STATE_WRITE(state, ray, dD) = differential_make_compact(bsdf_domega_in); + INTEGRATOR_STATE_WRITE(state, ray, dP) = differential_make_compact(sd->dP); + INTEGRATOR_STATE_WRITE(state, ray, dD) = differential_make_compact(bsdf_domega_in); #endif + } /* Update throughput. */ float3 throughput = INTEGRATOR_STATE(state, path, throughput); @@ -349,12 +354,8 @@ ccl_device_forceinline int integrate_surface_bsdf_bssrdf_bounce( } /* Update path state */ - if (label & LABEL_TRANSPARENT) { - INTEGRATOR_STATE_WRITE(state, path, mis_ray_t) += sd->ray_length; - } - else { + if (!(label & LABEL_TRANSPARENT)) { INTEGRATOR_STATE_WRITE(state, path, mis_ray_pdf) = bsdf_pdf; - INTEGRATOR_STATE_WRITE(state, path, mis_ray_t) = 0.0f; INTEGRATOR_STATE_WRITE(state, path, min_ray_pdf) = fminf( bsdf_pdf, INTEGRATOR_STATE(state, path, min_ray_pdf)); } @@ -371,17 +372,8 @@ ccl_device_forceinline int integrate_surface_volume_only_bounce(IntegratorState return LABEL_NONE; } - /* Setup ray position, direction stays unchanged. */ - INTEGRATOR_STATE_WRITE(state, ray, P) = sd->P; - - /* Clipping works through transparent. */ - INTEGRATOR_STATE_WRITE(state, ray, t) -= sd->ray_length; - -# ifdef __RAY_DIFFERENTIALS__ - INTEGRATOR_STATE_WRITE(state, ray, dP) = differential_make_compact(sd->dP); -# endif - - INTEGRATOR_STATE_WRITE(state, path, mis_ray_t) += sd->ray_length; + /* Only modify start distance. */ + INTEGRATOR_STATE_WRITE(state, ray, tmin) = intersection_t_offset(sd->ray_length); return LABEL_TRANSMIT | LABEL_TRANSPARENT; } @@ -432,7 +424,8 @@ ccl_device_forceinline void integrate_surface_ao(KernelGlobals kg, Ray ray ccl_optional_struct_init; ray.P = shadow_ray_offset(kg, sd, ao_D, &skip_self); ray.D = ao_D; - ray.t = kernel_data.integrator.ao_bounces_distance; + ray.tmin = 0.0f; + ray.tmax = kernel_data.integrator.ao_bounces_distance; ray.time = sd->time; ray.self.object = (skip_self) ? sd->object : OBJECT_NONE; ray.self.prim = (skip_self) ? sd->prim : PRIM_NONE; @@ -616,7 +609,7 @@ ccl_device_forceinline void integrator_shade_surface(KernelGlobals kg, kg, state, current_kernel, DEVICE_KERNEL_INTEGRATOR_INTERSECT_SUBSURFACE); } else { - kernel_assert(INTEGRATOR_STATE(state, ray, t) != 0.0f); + kernel_assert(INTEGRATOR_STATE(state, ray, tmax) != 0.0f); integrator_path_next(kg, state, current_kernel, DEVICE_KERNEL_INTEGRATOR_INTERSECT_CLOSEST); } } diff --git a/intern/cycles/kernel/integrator/shade_volume.h b/intern/cycles/kernel/integrator/shade_volume.h index 683c031f0d9..4aab097a7d8 100644 --- a/intern/cycles/kernel/integrator/shade_volume.h +++ b/intern/cycles/kernel/integrator/shade_volume.h @@ -114,7 +114,8 @@ ccl_device_inline bool volume_shader_sample(KernelGlobals kg, ccl_device_forceinline void volume_step_init(KernelGlobals kg, ccl_private const RNGState *rng_state, const float object_step_size, - float t, + const float tmin, + const float tmax, ccl_private float *step_size, ccl_private float *step_shade_offset, ccl_private float *steps_offset, @@ -122,7 +123,7 @@ ccl_device_forceinline void volume_step_init(KernelGlobals kg, { if (object_step_size == FLT_MAX) { /* Homogeneous volume. */ - *step_size = t; + *step_size = tmax - tmin; *step_shade_offset = 0.0f; *steps_offset = 1.0f; *max_steps = 1; @@ -130,6 +131,7 @@ ccl_device_forceinline void volume_step_init(KernelGlobals kg, else { /* Heterogeneous volume. */ *max_steps = kernel_data.integrator.volume_max_steps; + const float t = tmax - tmin; float step = min(object_step_size, t); /* compute exact steps in advance for malloc */ @@ -165,7 +167,7 @@ ccl_device void volume_shadow_homogeneous(KernelGlobals kg, IntegratorState stat float3 sigma_t = zero_float3(); if (shadow_volume_shader_sample(kg, state, sd, &sigma_t)) { - *throughput *= volume_color_transmittance(sigma_t, ray->t); + *throughput *= volume_color_transmittance(sigma_t, ray->tmax - ray->tmin); } } # endif @@ -194,7 +196,8 @@ ccl_device void volume_shadow_heterogeneous(KernelGlobals kg, volume_step_init(kg, &rng_state, object_step_size, - ray->t, + ray->tmin, + ray->tmax, &step_size, &step_shade_offset, &unused, @@ -202,13 +205,13 @@ ccl_device void volume_shadow_heterogeneous(KernelGlobals kg, const float steps_offset = 1.0f; /* compute extinction at the start */ - float t = 0.0f; + float t = ray->tmin; float3 sum = zero_float3(); for (int i = 0; i < max_steps; i++) { /* advance to new position */ - float new_t = min(ray->t, (i + steps_offset) * step_size); + float new_t = min(ray->tmax, ray->tmin + (i + steps_offset) * step_size); float dt = new_t - t; float3 new_P = ray->P + ray->D * (t + dt * step_shade_offset); @@ -233,7 +236,7 @@ ccl_device void volume_shadow_heterogeneous(KernelGlobals kg, /* stop if at the end of the volume */ t = new_t; - if (t == ray->t) { + if (t == ray->tmax) { /* Update throughput in case we haven't done it above */ tp = *throughput * exp(sum); break; @@ -257,15 +260,16 @@ ccl_device float volume_equiangular_sample(ccl_private const Ray *ccl_restrict r const float xi, ccl_private float *pdf) { - const float t = ray->t; + const float tmin = ray->tmin; + const float tmax = ray->tmax; const float delta = dot((light_P - ray->P), ray->D); const float D = safe_sqrtf(len_squared(light_P - ray->P) - delta * delta); if (UNLIKELY(D == 0.0f)) { *pdf = 0.0f; return 0.0f; } - const float theta_a = -atan2f(delta, D); - const float theta_b = atan2f(t - delta, D); + const float theta_a = atan2f(tmin - delta, D); + const float theta_b = atan2f(tmax - delta, D); const float t_ = D * tanf((xi * theta_b) + (1 - xi) * theta_a); if (UNLIKELY(theta_b == theta_a)) { *pdf = 0.0f; @@ -273,7 +277,7 @@ ccl_device float volume_equiangular_sample(ccl_private const Ray *ccl_restrict r } *pdf = D / ((theta_b - theta_a) * (D * D + t_ * t_)); - return min(t, delta + t_); /* min is only for float precision errors */ + return clamp(delta + t_, tmin, tmax); /* clamp is only for float precision errors */ } ccl_device float volume_equiangular_pdf(ccl_private const Ray *ccl_restrict ray, @@ -286,11 +290,12 @@ ccl_device float volume_equiangular_pdf(ccl_private const Ray *ccl_restrict ray, return 0.0f; } - const float t = ray->t; + const float tmin = ray->tmin; + const float tmax = ray->tmax; const float t_ = sample_t - delta; - const float theta_a = -atan2f(delta, D); - const float theta_b = atan2f(t - delta, D); + const float theta_a = atan2f(tmin - delta, D); + const float theta_b = atan2f(tmax - delta, D); if (UNLIKELY(theta_b == theta_a)) { return 0.0f; } @@ -310,11 +315,12 @@ ccl_device float volume_equiangular_cdf(ccl_private const Ray *ccl_restrict ray, return 0.0f; } - const float t = ray->t; + const float tmin = ray->tmin; + const float tmax = ray->tmax; const float t_ = sample_t - delta; - const float theta_a = -atan2f(delta, D); - const float theta_b = atan2f(t - delta, D); + const float theta_a = atan2f(tmin - delta, D); + const float theta_b = atan2f(tmax - delta, D); if (UNLIKELY(theta_b == theta_a)) { return 0.0f; } @@ -390,8 +396,8 @@ ccl_device float3 volume_emission_integrate(ccl_private VolumeShaderCoefficients typedef struct VolumeIntegrateState { /* Volume segment extents. */ - float start_t; - float end_t; + float tmin; + float tmax; /* If volume is absorption-only up to this point, and no probabilistic * scattering or termination has been used yet. */ @@ -426,9 +432,9 @@ ccl_device_forceinline void volume_integrate_step_scattering( /* Equiangular sampling for direct lighting. */ if (vstate.direct_sample_method == VOLUME_SAMPLE_EQUIANGULAR && !result.direct_scatter) { - if (result.direct_t >= vstate.start_t && result.direct_t <= vstate.end_t && + if (result.direct_t >= vstate.tmin && result.direct_t <= vstate.tmax && vstate.equiangular_pdf > VOLUME_SAMPLE_PDF_CUTOFF) { - const float new_dt = result.direct_t - vstate.start_t; + const float new_dt = result.direct_t - vstate.tmin; const float3 new_transmittance = volume_color_transmittance(coeff.sigma_t, new_dt); result.direct_scatter = true; @@ -458,7 +464,7 @@ ccl_device_forceinline void volume_integrate_step_scattering( /* compute sampling distance */ const float sample_sigma_t = volume_channel_get(coeff.sigma_t, channel); const float new_dt = -logf(1.0f - vstate.rscatter) / sample_sigma_t; - const float new_t = vstate.start_t + new_dt; + const float new_t = vstate.tmin + new_dt; /* transmittance and pdf */ const float3 new_transmittance = volume_color_transmittance(coeff.sigma_t, new_dt); @@ -528,7 +534,8 @@ ccl_device_forceinline void volume_integrate_heterogeneous( volume_step_init(kg, rng_state, object_step_size, - ray->t, + ray->tmin, + ray->tmax, &step_size, &step_shade_offset, &steps_offset, @@ -536,8 +543,8 @@ ccl_device_forceinline void volume_integrate_heterogeneous( /* Initialize volume integration state. */ VolumeIntegrateState vstate ccl_optional_struct_init; - vstate.start_t = 0.0f; - vstate.end_t = 0.0f; + vstate.tmin = ray->tmin; + vstate.tmax = ray->tmin; vstate.absorption_only = true; vstate.rscatter = path_state_rng_1D(kg, rng_state, PRNG_SCATTER_DISTANCE); vstate.rphase = path_state_rng_1D(kg, rng_state, PRNG_PHASE_CHANNEL); @@ -578,8 +585,8 @@ ccl_device_forceinline void volume_integrate_heterogeneous( for (int i = 0; i < max_steps; i++) { /* Advance to new position */ - vstate.end_t = min(ray->t, (i + steps_offset) * step_size); - const float shade_t = vstate.start_t + (vstate.end_t - vstate.start_t) * step_shade_offset; + vstate.tmax = min(ray->tmax, ray->tmin + (i + steps_offset) * step_size); + const float shade_t = vstate.tmin + (vstate.tmax - vstate.tmin) * step_shade_offset; sd->P = ray->P + ray->D * shade_t; /* compute segment */ @@ -588,7 +595,7 @@ ccl_device_forceinline void volume_integrate_heterogeneous( const int closure_flag = sd->flag; /* Evaluate transmittance over segment. */ - const float dt = (vstate.end_t - vstate.start_t); + const float dt = (vstate.tmax - vstate.tmin); const float3 transmittance = (closure_flag & SD_EXTINCTION) ? volume_color_transmittance(coeff.sigma_t, dt) : one_float3(); @@ -645,8 +652,8 @@ ccl_device_forceinline void volume_integrate_heterogeneous( } /* Stop if at the end of the volume. */ - vstate.start_t = vstate.end_t; - if (vstate.start_t == ray->t) { + vstate.tmin = vstate.tmax; + if (vstate.tmin == ray->tmax) { break; } } @@ -880,7 +887,8 @@ ccl_device_forceinline bool integrate_volume_phase_scatter( /* Setup ray. */ INTEGRATOR_STATE_WRITE(state, ray, P) = sd->P; INTEGRATOR_STATE_WRITE(state, ray, D) = normalize(phase_omega_in); - INTEGRATOR_STATE_WRITE(state, ray, t) = FLT_MAX; + INTEGRATOR_STATE_WRITE(state, ray, tmin) = 0.0f; + INTEGRATOR_STATE_WRITE(state, ray, tmax) = FLT_MAX; # ifdef __RAY_DIFFERENTIALS__ INTEGRATOR_STATE_WRITE(state, ray, dP) = differential_make_compact(sd->dP); INTEGRATOR_STATE_WRITE(state, ray, dD) = differential_make_compact(phase_domega_in); @@ -901,7 +909,6 @@ ccl_device_forceinline bool integrate_volume_phase_scatter( /* Update path state */ INTEGRATOR_STATE_WRITE(state, path, mis_ray_pdf) = phase_pdf; - INTEGRATOR_STATE_WRITE(state, path, mis_ray_t) = 0.0f; INTEGRATOR_STATE_WRITE(state, path, min_ray_pdf) = fminf( phase_pdf, INTEGRATOR_STATE(state, path, min_ray_pdf)); @@ -1021,7 +1028,7 @@ ccl_device void integrator_shade_volume(KernelGlobals kg, integrator_state_read_isect(kg, state, &isect); /* Set ray length to current segment. */ - ray.t = (isect.prim != PRIM_NONE) ? isect.t : FLT_MAX; + ray.tmax = (isect.prim != PRIM_NONE) ? isect.t : FLT_MAX; /* Clean volume stack for background rays. */ if (isect.prim == PRIM_NONE) { diff --git a/intern/cycles/kernel/integrator/shadow_state_template.h b/intern/cycles/kernel/integrator/shadow_state_template.h index eaee65ada40..c340467606d 100644 --- a/intern/cycles/kernel/integrator/shadow_state_template.h +++ b/intern/cycles/kernel/integrator/shadow_state_template.h @@ -47,7 +47,8 @@ KERNEL_STRUCT_END(shadow_path) KERNEL_STRUCT_BEGIN(shadow_ray) KERNEL_STRUCT_MEMBER(shadow_ray, packed_float3, P, KERNEL_FEATURE_PATH_TRACING) KERNEL_STRUCT_MEMBER(shadow_ray, packed_float3, D, KERNEL_FEATURE_PATH_TRACING) -KERNEL_STRUCT_MEMBER(shadow_ray, float, t, KERNEL_FEATURE_PATH_TRACING) +KERNEL_STRUCT_MEMBER(shadow_ray, float, tmin, KERNEL_FEATURE_PATH_TRACING) +KERNEL_STRUCT_MEMBER(shadow_ray, float, tmax, KERNEL_FEATURE_PATH_TRACING) KERNEL_STRUCT_MEMBER(shadow_ray, float, time, KERNEL_FEATURE_PATH_TRACING) KERNEL_STRUCT_MEMBER(shadow_ray, float, dP, KERNEL_FEATURE_PATH_TRACING) KERNEL_STRUCT_MEMBER(shadow_ray, int, object, KERNEL_FEATURE_PATH_TRACING) diff --git a/intern/cycles/kernel/integrator/state_template.h b/intern/cycles/kernel/integrator/state_template.h index e7e6db037b0..5c2af131945 100644 --- a/intern/cycles/kernel/integrator/state_template.h +++ b/intern/cycles/kernel/integrator/state_template.h @@ -37,11 +37,10 @@ KERNEL_STRUCT_MEMBER(path, uint32_t, flag, KERNEL_FEATURE_PATH_TRACING) /* enum PathRayMNEE */ KERNEL_STRUCT_MEMBER(path, uint8_t, mnee, KERNEL_FEATURE_PATH_TRACING) /* Multiple importance sampling - * The PDF of BSDF sampling at the last scatter point, and distance to the - * last scatter point minus the last ray segment. This distance lets us - * compute the complete distance through transparent surfaces and volumes. */ + * The PDF of BSDF sampling at the last scatter point, which is at ray distance + * zero and distance. Note that transparency and volume attenuation increase + * the ray tmin but keep P unmodified so that this works. */ KERNEL_STRUCT_MEMBER(path, float, mis_ray_pdf, KERNEL_FEATURE_PATH_TRACING) -KERNEL_STRUCT_MEMBER(path, float, mis_ray_t, KERNEL_FEATURE_PATH_TRACING) /* Filter glossy. */ KERNEL_STRUCT_MEMBER(path, float, min_ray_pdf, KERNEL_FEATURE_PATH_TRACING) /* Continuation probability for path termination. */ @@ -63,7 +62,8 @@ KERNEL_STRUCT_END(path) KERNEL_STRUCT_BEGIN(ray) KERNEL_STRUCT_MEMBER(ray, packed_float3, P, KERNEL_FEATURE_PATH_TRACING) KERNEL_STRUCT_MEMBER(ray, packed_float3, D, KERNEL_FEATURE_PATH_TRACING) -KERNEL_STRUCT_MEMBER(ray, float, t, KERNEL_FEATURE_PATH_TRACING) +KERNEL_STRUCT_MEMBER(ray, float, tmin, KERNEL_FEATURE_PATH_TRACING) +KERNEL_STRUCT_MEMBER(ray, float, tmax, KERNEL_FEATURE_PATH_TRACING) KERNEL_STRUCT_MEMBER(ray, float, time, KERNEL_FEATURE_PATH_TRACING) KERNEL_STRUCT_MEMBER(ray, float, dP, KERNEL_FEATURE_PATH_TRACING) KERNEL_STRUCT_MEMBER(ray, float, dD, KERNEL_FEATURE_PATH_TRACING) diff --git a/intern/cycles/kernel/integrator/state_util.h b/intern/cycles/kernel/integrator/state_util.h index 280db2d1aac..8dd58ad6bcd 100644 --- a/intern/cycles/kernel/integrator/state_util.h +++ b/intern/cycles/kernel/integrator/state_util.h @@ -17,7 +17,8 @@ ccl_device_forceinline void integrator_state_write_ray(KernelGlobals kg, { INTEGRATOR_STATE_WRITE(state, ray, P) = ray->P; INTEGRATOR_STATE_WRITE(state, ray, D) = ray->D; - INTEGRATOR_STATE_WRITE(state, ray, t) = ray->t; + INTEGRATOR_STATE_WRITE(state, ray, tmin) = ray->tmin; + INTEGRATOR_STATE_WRITE(state, ray, tmax) = ray->tmax; INTEGRATOR_STATE_WRITE(state, ray, time) = ray->time; INTEGRATOR_STATE_WRITE(state, ray, dP) = ray->dP; INTEGRATOR_STATE_WRITE(state, ray, dD) = ray->dD; @@ -29,7 +30,8 @@ ccl_device_forceinline void integrator_state_read_ray(KernelGlobals kg, { ray->P = INTEGRATOR_STATE(state, ray, P); ray->D = INTEGRATOR_STATE(state, ray, D); - ray->t = INTEGRATOR_STATE(state, ray, t); + ray->tmin = INTEGRATOR_STATE(state, ray, tmin); + ray->tmax = INTEGRATOR_STATE(state, ray, tmax); ray->time = INTEGRATOR_STATE(state, ray, time); ray->dP = INTEGRATOR_STATE(state, ray, dP); ray->dD = INTEGRATOR_STATE(state, ray, dD); @@ -42,7 +44,8 @@ ccl_device_forceinline void integrator_state_write_shadow_ray( { INTEGRATOR_STATE_WRITE(state, shadow_ray, P) = ray->P; INTEGRATOR_STATE_WRITE(state, shadow_ray, D) = ray->D; - INTEGRATOR_STATE_WRITE(state, shadow_ray, t) = ray->t; + INTEGRATOR_STATE_WRITE(state, shadow_ray, tmin) = ray->tmin; + INTEGRATOR_STATE_WRITE(state, shadow_ray, tmax) = ray->tmax; INTEGRATOR_STATE_WRITE(state, shadow_ray, time) = ray->time; INTEGRATOR_STATE_WRITE(state, shadow_ray, dP) = ray->dP; } @@ -53,7 +56,8 @@ ccl_device_forceinline void integrator_state_read_shadow_ray(KernelGlobals kg, { ray->P = INTEGRATOR_STATE(state, shadow_ray, P); ray->D = INTEGRATOR_STATE(state, shadow_ray, D); - ray->t = INTEGRATOR_STATE(state, shadow_ray, t); + ray->tmin = INTEGRATOR_STATE(state, shadow_ray, tmin); + ray->tmax = INTEGRATOR_STATE(state, shadow_ray, tmax); ray->time = INTEGRATOR_STATE(state, shadow_ray, time); ray->dP = INTEGRATOR_STATE(state, shadow_ray, dP); ray->dD = differential_zero_compact(); diff --git a/intern/cycles/kernel/integrator/subsurface.h b/intern/cycles/kernel/integrator/subsurface.h index ab26a2d93cc..2f96f215d8a 100644 --- a/intern/cycles/kernel/integrator/subsurface.h +++ b/intern/cycles/kernel/integrator/subsurface.h @@ -38,7 +38,8 @@ ccl_device int subsurface_bounce(KernelGlobals kg, /* Setup ray into surface. */ INTEGRATOR_STATE_WRITE(state, ray, P) = sd->P; INTEGRATOR_STATE_WRITE(state, ray, D) = bssrdf->N; - INTEGRATOR_STATE_WRITE(state, ray, t) = FLT_MAX; + INTEGRATOR_STATE_WRITE(state, ray, tmin) = 0.0f; + INTEGRATOR_STATE_WRITE(state, ray, tmax) = FLT_MAX; INTEGRATOR_STATE_WRITE(state, ray, dP) = differential_make_compact(sd->dP); INTEGRATOR_STATE_WRITE(state, ray, dD) = differential_zero_compact(); @@ -160,7 +161,7 @@ ccl_device_inline bool subsurface_scatter(KernelGlobals kg, IntegratorState stat /* Pretend ray is coming from the outside towards the exit point. This ensures * correct front/back facing normals. * TODO: find a more elegant solution? */ - ray.P += ray.D * ray.t * 2.0f; + ray.P += ray.D * ray.tmax * 2.0f; ray.D = -ray.D; integrator_state_write_isect(kg, state, &ss_isect.hits[0]); diff --git a/intern/cycles/kernel/integrator/subsurface_disk.h b/intern/cycles/kernel/integrator/subsurface_disk.h index ae857c50493..2836934f6dd 100644 --- a/intern/cycles/kernel/integrator/subsurface_disk.h +++ b/intern/cycles/kernel/integrator/subsurface_disk.h @@ -82,7 +82,8 @@ ccl_device_inline bool subsurface_disk(KernelGlobals kg, /* Create ray. */ ray.P = P + disk_N * disk_height + disk_P; ray.D = -disk_N; - ray.t = 2.0f * disk_height; + ray.tmin = 0.0f; + ray.tmax = 2.0f * disk_height; ray.dP = ray_dP; ray.dD = differential_zero_compact(); ray.time = time; @@ -188,7 +189,8 @@ ccl_device_inline bool subsurface_disk(KernelGlobals kg, ray.P = ray.P + ray.D * ss_isect.hits[hit].t; ray.D = ss_isect.Ng[hit]; - ray.t = 1.0f; + ray.tmin = 0.0f; + ray.tmax = 1.0f; return true; } diff --git a/intern/cycles/kernel/integrator/subsurface_random_walk.h b/intern/cycles/kernel/integrator/subsurface_random_walk.h index 8094bf7159e..c1691030817 100644 --- a/intern/cycles/kernel/integrator/subsurface_random_walk.h +++ b/intern/cycles/kernel/integrator/subsurface_random_walk.h @@ -195,7 +195,8 @@ ccl_device_inline bool subsurface_random_walk(KernelGlobals kg, /* Setup ray. */ ray.P = P; ray.D = D; - ray.t = FLT_MAX; + ray.tmin = 0.0f; + ray.tmax = FLT_MAX; ray.time = time; ray.dP = ray_dP; ray.dD = differential_zero_compact(); @@ -370,10 +371,10 @@ ccl_device_inline bool subsurface_random_walk(KernelGlobals kg, * chance of connecting to it. * TODO: Maybe use less than 10 times the mean free path? */ if (bounce == 0) { - ray.t = max(t, 10.0f / (reduce_min(sigma_t))); + ray.tmax = max(t, 10.0f / (reduce_min(sigma_t))); } else { - ray.t = t; + ray.tmax = t; /* After the first bounce the object can intersect the same surface again */ ray.self.object = OBJECT_NONE; ray.self.prim = PRIM_NONE; @@ -384,12 +385,12 @@ ccl_device_inline bool subsurface_random_walk(KernelGlobals kg, if (hit) { #ifdef __KERNEL_GPU_RAYTRACING__ /* t is always in world space with OptiX and MetalRT. */ - ray.t = ss_isect.hits[0].t; + ray.tmax = ss_isect.hits[0].t; #else /* Compute world space distance to surface hit. */ float3 D = transform_direction(&ob_itfm, ray.D); D = normalize(D) * ss_isect.hits[0].t; - ray.t = len(transform_direction(&ob_tfm, D)); + ray.tmax = len(transform_direction(&ob_tfm, D)); #endif } @@ -397,16 +398,16 @@ ccl_device_inline bool subsurface_random_walk(KernelGlobals kg, /* Check if we hit the opposite side. */ if (hit) { have_opposite_interface = true; - opposite_distance = dot(ray.P + ray.t * ray.D - P, -N); + opposite_distance = dot(ray.P + ray.tmax * ray.D - P, -N); } /* Apart from the opposite side check, we were supposed to only trace up to distance t, * so check if there would have been a hit in that case. */ - hit = ray.t < t; + hit = ray.tmax < t; } /* Use the distance to the exit point for the throughput update if we found one. */ if (hit) { - t = ray.t; + t = ray.tmax; } /* Advance to new scatter location. */ diff --git a/intern/cycles/kernel/light/light.h b/intern/cycles/kernel/light/light.h index 1e7a333d013..b939489bb18 100644 --- a/intern/cycles/kernel/light/light.h +++ b/intern/cycles/kernel/light/light.h @@ -270,31 +270,26 @@ ccl_device bool lights_intersect(KernelGlobals kg, if (type == LIGHT_SPOT) { /* Spot/Disk light. */ - const float mis_ray_t = INTEGRATOR_STATE(state, path, mis_ray_t); - const float3 ray_P = ray->P - ray->D * mis_ray_t; - const float3 lightP = make_float3(klight->co[0], klight->co[1], klight->co[2]); const float radius = klight->spot.radius; if (radius == 0.0f) { continue; } /* disk oriented normal */ - const float3 lightN = normalize(ray_P - lightP); + const float3 lightN = normalize(ray->P - lightP); /* One sided. */ if (dot(ray->D, lightN) >= 0.0f) { continue; } float3 P; - if (!ray_disk_intersect(ray->P, ray->D, ray->t, lightP, lightN, radius, &P, &t)) { + if (!ray_disk_intersect( + ray->P, ray->D, ray->tmin, ray->tmax, lightP, lightN, radius, &P, &t)) { continue; } } else if (type == LIGHT_POINT) { /* Sphere light (aka, aligned disk light). */ - const float mis_ray_t = INTEGRATOR_STATE(state, path, mis_ray_t); - const float3 ray_P = ray->P - ray->D * mis_ray_t; - const float3 lightP = make_float3(klight->co[0], klight->co[1], klight->co[2]); const float radius = klight->spot.radius; if (radius == 0.0f) { @@ -302,9 +297,10 @@ ccl_device bool lights_intersect(KernelGlobals kg, } /* disk oriented normal */ - const float3 lightN = normalize(ray_P - lightP); + const float3 lightN = normalize(ray->P - lightP); float3 P; - if (!ray_disk_intersect(ray->P, ray->D, ray->t, lightP, lightN, radius, &P, &t)) { + if (!ray_disk_intersect( + ray->P, ray->D, ray->tmin, ray->tmax, lightP, lightN, radius, &P, &t)) { continue; } } @@ -330,8 +326,19 @@ ccl_device bool lights_intersect(KernelGlobals kg, const float3 light_P = make_float3(klight->co[0], klight->co[1], klight->co[2]); float3 P; - if (!ray_quad_intersect( - ray->P, ray->D, 0.0f, ray->t, light_P, axisu, axisv, Ng, &P, &t, &u, &v, is_round)) { + if (!ray_quad_intersect(ray->P, + ray->D, + ray->tmin, + ray->tmax, + light_P, + axisu, + axisv, + Ng, + &P, + &t, + &u, + &v, + is_round)) { continue; } } @@ -775,7 +782,8 @@ ccl_device_forceinline void triangle_light_sample(KernelGlobals kg, ls->D = z * B + safe_sqrtf(1.0f - z * z) * safe_normalize(C_ - dot(C_, B) * B); /* calculate intersection with the planar triangle */ - if (!ray_triangle_intersect(P, ls->D, FLT_MAX, V[0], V[1], V[2], &ls->u, &ls->v, &ls->t)) { + if (!ray_triangle_intersect( + P, ls->D, 0.0f, FLT_MAX, V[0], V[1], V[2], &ls->u, &ls->v, &ls->t)) { ls->pdf = 0.0f; return; } diff --git a/intern/cycles/kernel/light/sample.h b/intern/cycles/kernel/light/sample.h index 5cf7dce683a..210bb1b35c2 100644 --- a/intern/cycles/kernel/light/sample.h +++ b/intern/cycles/kernel/light/sample.h @@ -227,23 +227,24 @@ ccl_device_inline void shadow_ray_setup(ccl_private const ShaderData *ccl_restri if (ls->shader & SHADER_CAST_SHADOW) { /* setup ray */ ray->P = P; + ray->tmin = 0.0f; if (ls->t == FLT_MAX) { /* distant light */ ray->D = ls->D; - ray->t = ls->t; + ray->tmax = ls->t; } else { /* other lights, avoid self-intersection */ ray->D = ls->P - P; - ray->D = normalize_len(ray->D, &ray->t); + ray->D = normalize_len(ray->D, &ray->tmax); } } else { /* signal to not cast shadow ray */ ray->P = zero_float3(); ray->D = zero_float3(); - ray->t = 0.0f; + ray->tmax = 0.0f; } ray->dP = differential_make_compact(sd->dP); diff --git a/intern/cycles/kernel/osl/services.cpp b/intern/cycles/kernel/osl/services.cpp index 78c23b858c4..6b7981b7f3a 100644 --- a/intern/cycles/kernel/osl/services.cpp +++ b/intern/cycles/kernel/osl/services.cpp @@ -1094,10 +1094,8 @@ bool OSLRenderServices::get_background_attribute(const KernelGlobalsCPU *kg, ndc[0] = camera_world_to_ndc(kg, sd, sd->ray_P); if (derivatives) { - ndc[1] = camera_world_to_ndc(kg, sd, sd->ray_P + make_float3(sd->ray_dP, 0.0f, 0.0f)) - - ndc[0]; - ndc[2] = camera_world_to_ndc(kg, sd, sd->ray_P + make_float3(0.0f, sd->ray_dP, 0.0f)) - - ndc[0]; + ndc[1] = zero_float3(); + ndc[2] = zero_float3(); } } else { @@ -1671,7 +1669,8 @@ bool OSLRenderServices::trace(TraceOpt &options, ray.P = TO_FLOAT3(P); ray.D = TO_FLOAT3(R); - ray.t = (options.maxdist == 1.0e30f) ? FLT_MAX : options.maxdist - options.mindist; + ray.tmin = 0.0f; + ray.tmax = (options.maxdist == 1.0e30f) ? FLT_MAX : options.maxdist - options.mindist; ray.time = sd->time; ray.self.object = OBJECT_NONE; ray.self.prim = PRIM_NONE; diff --git a/intern/cycles/kernel/svm/ao.h b/intern/cycles/kernel/svm/ao.h index e66c535824c..c57c68d6230 100644 --- a/intern/cycles/kernel/svm/ao.h +++ b/intern/cycles/kernel/svm/ao.h @@ -59,7 +59,8 @@ ccl_device float svm_ao( Ray ray; ray.P = sd->P; ray.D = D.x * T + D.y * B + D.z * N; - ray.t = max_dist; + ray.tmin = 0.0f; + ray.tmax = max_dist; ray.time = sd->time; ray.self.object = sd->object; ray.self.prim = sd->prim; diff --git a/intern/cycles/kernel/svm/bevel.h b/intern/cycles/kernel/svm/bevel.h index 790437d8e82..4617a056a52 100644 --- a/intern/cycles/kernel/svm/bevel.h +++ b/intern/cycles/kernel/svm/bevel.h @@ -179,7 +179,8 @@ ccl_device float3 svm_bevel( Ray ray ccl_optional_struct_init; ray.P = sd->P + disk_N * disk_height + disk_P; ray.D = -disk_N; - ray.t = 2.0f * disk_height; + ray.tmin = 0.0f; + ray.tmax = 2.0f * disk_height; ray.dP = differential_zero_compact(); ray.dD = differential_zero_compact(); ray.time = sd->time; diff --git a/intern/cycles/kernel/svm/tex_coord.h b/intern/cycles/kernel/svm/tex_coord.h index d9138796c45..2a0130e11d4 100644 --- a/intern/cycles/kernel/svm/tex_coord.h +++ b/intern/cycles/kernel/svm/tex_coord.h @@ -138,7 +138,7 @@ ccl_device_noinline int svm_node_tex_coord_bump_dx(KernelGlobals kg, case NODE_TEXCO_WINDOW: { if ((path_flag & PATH_RAY_CAMERA) && sd->object == OBJECT_NONE && kernel_data.cam.type == CAMERA_ORTHOGRAPHIC) - data = camera_world_to_ndc(kg, sd, sd->ray_P + make_float3(sd->ray_dP, 0.0f, 0.0f)); + data = camera_world_to_ndc(kg, sd, sd->ray_P); else data = camera_world_to_ndc(kg, sd, sd->P + sd->dP.dx); data.z = 0.0f; @@ -223,7 +223,7 @@ ccl_device_noinline int svm_node_tex_coord_bump_dy(KernelGlobals kg, case NODE_TEXCO_WINDOW: { if ((path_flag & PATH_RAY_CAMERA) && sd->object == OBJECT_NONE && kernel_data.cam.type == CAMERA_ORTHOGRAPHIC) - data = camera_world_to_ndc(kg, sd, sd->ray_P + make_float3(0.0f, sd->ray_dP, 0.0f)); + data = camera_world_to_ndc(kg, sd, sd->ray_P); else data = camera_world_to_ndc(kg, sd, sd->P + sd->dP.dy); data.z = 0.0f; diff --git a/intern/cycles/kernel/types.h b/intern/cycles/kernel/types.h index 62ac75e5e4d..05320deed19 100644 --- a/intern/cycles/kernel/types.h +++ b/intern/cycles/kernel/types.h @@ -535,7 +535,8 @@ typedef struct RaySelfPrimitives { typedef struct Ray { float3 P; /* origin */ float3 D; /* direction */ - float t; /* length of the ray */ + float tmin; /* start distance */ + float tmax; /* end distance */ float time; /* time (for motion blur) */ RaySelfPrimitives self; |