griefith/test2
1
0
Fork 0
Code Issues Pull Requests Actions 1 Packages Projects Releases Wiki Activity

Cycles: refactor rays to have start and end distance, fix precision issues

Browse Source 
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
This commit is contained in:
Brecht Van Lommel
2022-07-13 16:54:53 +02:00
parent bb376da6df
commit 5152c7c152
42 changed files with 427 additions and 294 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -175,8 +175,8 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg,
optixTrace(scene_intersect_valid(ray) ? kernel_data.device_bvh : 0, optixTrace(scene_intersect_valid(ray) ? kernel_data.device_bvh : 0,
ray->P, ray->P,
ray->D, ray->D,
0.0f, ray->tmin,
ray->t, ray->tmax,
ray->time, ray->time,
ray_mask, ray_mask,
ray_flags, ray_flags,
@@ -203,28 +203,28 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg,
#elif defined(__METALRT__) #elif defined(__METALRT__)
if (!scene_intersect_valid(ray)) { if (!scene_intersect_valid(ray)) {
isect->t = ray->t; isect->t = ray->tmax;
isect->type = PRIMITIVE_NONE; isect->type = PRIMITIVE_NONE;
return false; return false;
} }
# if defined(__KERNEL_DEBUG__) # if defined(__KERNEL_DEBUG__)
if (is_null_instance_acceleration_structure(metal_ancillaries->accel_struct)) { if (is_null_instance_acceleration_structure(metal_ancillaries->accel_struct)) {
isect->t = ray->t; isect->t = ray->tmax;
isect->type = PRIMITIVE_NONE; isect->type = PRIMITIVE_NONE;
kernel_assert(!"Invalid metal_ancillaries->accel_struct pointer"); kernel_assert(!"Invalid metal_ancillaries->accel_struct pointer");
return false; return false;
} }
if (is_null_intersection_function_table(metal_ancillaries->ift_default)) { if (is_null_intersection_function_table(metal_ancillaries->ift_default)) {
isect->t = ray->t; isect->t = ray->tmax;
isect->type = PRIMITIVE_NONE; isect->type = PRIMITIVE_NONE;
kernel_assert(!"Invalid ift_default"); kernel_assert(!"Invalid ift_default");
return false; return false;
} }
# endif # 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_intersector_type metalrt_intersect;
if (!kernel_data.bvh.have_curves) { if (!kernel_data.bvh.have_curves) {
@@ -263,7 +263,7 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg,
# endif # endif
if (intersection.type == intersection_type::none) { if (intersection.type == intersection_type::none) {
isect->t = ray->t; isect->t = ray->tmax;
isect->type = PRIMITIVE_NONE; isect->type = PRIMITIVE_NONE;
return false; return false;
@@ -296,7 +296,7 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg,
# ifdef __EMBREE__ # ifdef __EMBREE__
if (kernel_data.device_bvh) { if (kernel_data.device_bvh) {
isect->t = ray->t; isect->t = ray->tmax;
CCLIntersectContext ctx(kg, CCLIntersectContext::RAY_REGULAR); CCLIntersectContext ctx(kg, CCLIntersectContext::RAY_REGULAR);
IntersectContext rtc_ctx(&ctx); IntersectContext rtc_ctx(&ctx);
RTCRayHit ray_hit; 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, optixTrace(scene_intersect_valid(ray) ? kernel_data.device_bvh : 0,
ray->P, ray->P,
ray->D, ray->D,
0.0f, ray->tmin,
ray->t, ray->tmax,
ray->time, ray->time,
0xFF, 0xFF,
/* Need to always call into __anyhit__kernel_optix_local_hit. */ /* Need to always call into __anyhit__kernel_optix_local_hit. */
@@ -405,7 +405,7 @@ ccl_device_intersect bool scene_intersect_local(KernelGlobals kg,
} }
# endif # 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_intersector_type metalrt_intersect;
metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque); 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 dir = ray->D;
float3 idir = ray->D; float3 idir = ray->D;
Transform ob_itfm; 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(kg, local_object, ray, &P, &dir, &idir, &ob_itfm);
/* bvh_instance_motion_push() returns the inverse transform but /* bvh_instance_motion_push() returns the inverse transform but
* it's not needed here. */ * 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, optixTrace(scene_intersect_valid(ray) ? kernel_data.device_bvh : 0,
ray->P, ray->P,
ray->D, ray->D,
0.0f, ray->tmin,
ray->t, ray->tmax,
ray->time, ray->time,
ray_mask, ray_mask,
/* Need to always call into __anyhit__kernel_optix_shadow_all_hit. */ /* 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 # 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_intersector_type metalrt_intersect;
metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque); 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, optixTrace(scene_intersect_valid(ray) ? kernel_data.device_bvh : 0,
ray->P, ray->P,
ray->D, ray->D,
0.0f, ray->tmin,
ray->t, ray->tmax,
ray->time, ray->time,
ray_mask, ray_mask,
/* Need to always call into __anyhit__kernel_optix_volume_test. */ /* Need to always call into __anyhit__kernel_optix_volume_test. */
@@ -744,7 +744,7 @@ ccl_device_intersect bool scene_intersect_volume(KernelGlobals kg,
} }
# endif # 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_intersector_type metalrt_intersect;
metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque); metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque);

4
intern/cycles/kernel/bvh/embree.h
View File

@@ -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_x = ray.D.x;
rtc_ray.dir_y = ray.D.y; rtc_ray.dir_y = ray.D.y;
rtc_ray.dir_z = ray.D.z; rtc_ray.dir_z = ray.D.z;
rtc_ray.tnear = 0.0f; rtc_ray.tnear = ray.tmin;
rtc_ray.tfar = ray.t; rtc_ray.tfar = ray.tmax;
rtc_ray.time = ray.time; rtc_ray.time = ray.time;
rtc_ray.mask = visibility; rtc_ray.mask = visibility;
} }

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

@@ -47,8 +47,9 @@ ccl_device_inline
float3 P = ray->P; float3 P = ray->P;
float3 dir = bvh_clamp_direction(ray->D); float3 dir = bvh_clamp_direction(ray->D);
float3 idir = bvh_inverse_direction(dir); float3 idir = bvh_inverse_direction(dir);
float tmin = ray->tmin;
int object = OBJECT_NONE; int object = OBJECT_NONE;
float isect_t = ray->t; float isect_t = ray->tmax;
if (local_isect != NULL) { if (local_isect != NULL) {
local_isect->num_hits = 0; local_isect->num_hits = 0;
@@ -59,10 +60,13 @@ ccl_device_inline
if (!(object_flag & SD_OBJECT_TRANSFORM_APPLIED)) { if (!(object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
#if BVH_FEATURE(BVH_MOTION) #if BVH_FEATURE(BVH_MOTION)
Transform ob_itfm; 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 #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 #endif
isect_t *= t_world_to_instance;
tmin *= t_world_to_instance;
object = local_object; object = local_object;
} }
@@ -81,6 +85,7 @@ ccl_device_inline
dir, dir,
#endif #endif
idir, idir,
tmin,
isect_t, isect_t,
node_addr, node_addr,
PATH_RAY_ALL_VISIBILITY, PATH_RAY_ALL_VISIBILITY,
@@ -155,6 +160,7 @@ ccl_device_inline
local_object, local_object,
prim, prim,
prim_addr, prim_addr,
tmin,
isect_t, isect_t,
lcg_state, lcg_state,
max_hits)) { max_hits)) {
@@ -191,6 +197,7 @@ ccl_device_inline
local_object, local_object,
prim, prim,
prim_addr, prim_addr,
tmin,
isect_t, isect_t,
lcg_state, lcg_state,
max_hits)) { max_hits)) {

32
intern/cycles/kernel/bvh/nodes.h
View File

@@ -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, ccl_device_forceinline int bvh_aligned_node_intersect(KernelGlobals kg,
const float3 P, const float3 P,
const float3 idir, const float3 idir,
const float t, const float tmin,
const float tmax,
const int node_addr, const int node_addr,
const uint visibility, const uint visibility,
float dist[2]) 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 c0hiy = (node1.z - P.y) * idir.y;
float c0loz = (node2.x - P.z) * idir.z; float c0loz = (node2.x - P.z) * idir.z;
float c0hiz = (node2.z - 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 c0min = max4(tmin, min(c0lox, c0hix), min(c0loy, c0hiy), min(c0loz, c0hiz));
float c0max = min4(t, max(c0lox, c0hix), max(c0loy, c0hiy), max(c0loz, c0hiz)); float c0max = min4(tmax, max(c0lox, c0hix), max(c0loy, c0hiy), max(c0loz, c0hiz));
float c1lox = (node0.y - P.x) * idir.x; float c1lox = (node0.y - P.x) * idir.x;
float c1hix = (node0.w - 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 c1hiy = (node1.w - P.y) * idir.y;
float c1loz = (node2.y - P.z) * idir.z; float c1loz = (node2.y - P.z) * idir.z;
float c1hiz = (node2.w - 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 c1min = max4(tmin, min(c1lox, c1hix), min(c1loy, c1hiy), min(c1loz, c1hiz));
float c1max = min4(t, max(c1lox, c1hix), max(c1loy, c1hiy), max(c1loz, c1hiz)); float c1max = min4(tmax, max(c1lox, c1hix), max(c1loy, c1hiy), max(c1loz, c1hiz));
dist[0] = c0min; dist[0] = c0min;
dist[1] = c1min; 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, ccl_device_forceinline bool bvh_unaligned_node_intersect_child(KernelGlobals kg,
const float3 P, const float3 P,
const float3 dir, const float3 dir,
const float t, const float tmin,
const float tmax,
int node_addr, int node_addr,
int child, int child,
float dist[2]) 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_x = max(lower_xyz.x, upper_xyz.x);
const float far_y = max(lower_xyz.y, upper_xyz.y); const float far_y = max(lower_xyz.y, upper_xyz.y);
const float far_z = max(lower_xyz.z, upper_xyz.z); 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 tnear = max4(tmin, near_x, near_y, near_z);
const float tfar = min4(t, far_x, far_y, far_z); const float tfar = min4(tmax, far_x, far_y, far_z);
*dist = tnear; *dist = tnear;
return tnear <= tfar; return tnear <= tfar;
} }
@@ -93,7 +95,8 @@ ccl_device_forceinline int bvh_unaligned_node_intersect(KernelGlobals kg,
const float3 P, const float3 P,
const float3 dir, const float3 dir,
const float3 idir, const float3 idir,
const float t, const float tmin,
const float tmax,
const int node_addr, const int node_addr,
const uint visibility, const uint visibility,
float dist[2]) float dist[2])
@@ -102,7 +105,7 @@ ccl_device_forceinline int bvh_unaligned_node_intersect(KernelGlobals kg,
#ifdef __VISIBILITY_FLAG__ #ifdef __VISIBILITY_FLAG__
float4 cnodes = kernel_data_fetch(bvh_nodes, node_addr + 0); float4 cnodes = kernel_data_fetch(bvh_nodes, node_addr + 0);
#endif #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__ #ifdef __VISIBILITY_FLAG__
if ((__float_as_uint(cnodes.x) & visibility)) if ((__float_as_uint(cnodes.x) & visibility))
#endif #endif
@@ -110,7 +113,7 @@ ccl_device_forceinline int bvh_unaligned_node_intersect(KernelGlobals kg,
mask |= 1; 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__ #ifdef __VISIBILITY_FLAG__
if ((__float_as_uint(cnodes.y) & visibility)) if ((__float_as_uint(cnodes.y) & visibility))
#endif #endif
@@ -125,16 +128,17 @@ ccl_device_forceinline int bvh_node_intersect(KernelGlobals kg,
const float3 P, const float3 P,
const float3 dir, const float3 dir,
const float3 idir, const float3 idir,
const float t, const float tmin,
const float tmax,
const int node_addr, const int node_addr,
const uint visibility, const uint visibility,
float dist[2]) float dist[2])
{ {
float4 node = kernel_data_fetch(bvh_nodes, node_addr); float4 node = kernel_data_fetch(bvh_nodes, node_addr);
if (__float_as_uint(node.x) & PATH_RAY_NODE_UNALIGNED) { 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 { 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);
} }
} }

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

@@ -49,6 +49,7 @@ ccl_device_inline
float3 P = ray->P; float3 P = ray->P;
float3 dir = bvh_clamp_direction(ray->D); float3 dir = bvh_clamp_direction(ray->D);
float3 idir = bvh_inverse_direction(dir); float3 idir = bvh_inverse_direction(dir);
float tmin = ray->tmin;
int object = OBJECT_NONE; int object = OBJECT_NONE;
uint num_hits = 0; uint num_hits = 0;
@@ -59,12 +60,12 @@ ccl_device_inline
/* Max distance in world space. May be dynamically reduced when max number of /* 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 * recorded hits is exceeded and we no longer need to find hits beyond the max
* distance found. */ * distance found. */
float t_max_world = ray->t; float t_max_world = ray->tmax;
/* Current maximum distance to the intersection. /* Current maximum distance to the intersection.
* Is calculated as a ray length, transformed to an object space when entering * Is calculated as a ray length, transformed to an object space when entering
* instance node. */ * 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 /* Conversion from world to local space for the current instance if any, 1.0
* otherwise. */ * otherwise. */
@@ -88,6 +89,7 @@ ccl_device_inline
dir, dir,
#endif #endif
idir, idir,
tmin,
t_max_current, t_max_current,
node_addr, node_addr,
visibility, visibility,
@@ -156,8 +158,16 @@ ccl_device_inline
switch (type & PRIMITIVE_ALL) { switch (type & PRIMITIVE_ALL) {
case PRIMITIVE_TRIANGLE: { case PRIMITIVE_TRIANGLE: {
hit = triangle_intersect( hit = triangle_intersect(kg,
kg, &isect, P, dir, t_max_current, visibility, prim_object, prim, prim_addr); &isect,
P,
dir,
tmin,
t_max_current,
visibility,
prim_object,
prim,
prim_addr);
break; break;
} }
#if BVH_FEATURE(BVH_MOTION) #if BVH_FEATURE(BVH_MOTION)
@@ -166,6 +176,7 @@ ccl_device_inline
&isect, &isect,
P, P,
dir, dir,
tmin,
t_max_current, t_max_current,
ray->time, ray->time,
visibility, visibility,
@@ -189,8 +200,16 @@ ccl_device_inline
} }
const int curve_type = kernel_data_fetch(prim_type, prim_addr); const int curve_type = kernel_data_fetch(prim_type, prim_addr);
hit = curve_intersect( hit = curve_intersect(kg,
kg, &isect, P, dir, t_max_current, prim_object, prim, ray->time, curve_type); &isect,
P,
dir,
tmin,
t_max_current,
prim_object,
prim,
ray->time,
curve_type);
break; break;
} }
@@ -207,8 +226,16 @@ ccl_device_inline
} }
const int point_type = kernel_data_fetch(prim_type, prim_addr); const int point_type = kernel_data_fetch(prim_type, prim_addr);
hit = point_intersect( hit = point_intersect(kg,
kg, &isect, P, dir, t_max_current, prim_object, prim, ray->time, point_type); &isect,
P,
dir,
tmin,
t_max_current,
prim_object,
prim,
ray->time,
point_type);
break; break;
} }
#endif /* BVH_FEATURE(BVH_POINTCLOUD) */ #endif /* BVH_FEATURE(BVH_POINTCLOUD) */
@@ -302,6 +329,7 @@ ccl_device_inline
/* Convert intersection to object space. */ /* Convert intersection to object space. */
t_max_current *= t_world_to_instance; t_max_current *= t_world_to_instance;
tmin *= t_world_to_instance;
++stack_ptr; ++stack_ptr;
kernel_assert(stack_ptr < BVH_STACK_SIZE); kernel_assert(stack_ptr < BVH_STACK_SIZE);
@@ -323,7 +351,8 @@ ccl_device_inline
#endif #endif
/* Restore world space ray length. */ /* Restore world space ray length. */
t_max_current = ray->t; tmin = ray->tmin;
t_max_current = ray->tmax;
object = OBJECT_NONE; object = OBJECT_NONE;
t_world_to_instance = 1.0f; t_world_to_instance = 1.0f;

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

@@ -43,13 +43,14 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg,
float3 P = ray->P; float3 P = ray->P;
float3 dir = bvh_clamp_direction(ray->D); float3 dir = bvh_clamp_direction(ray->D);
float3 idir = bvh_inverse_direction(dir); float3 idir = bvh_inverse_direction(dir);
float tmin = ray->tmin;
int object = OBJECT_NONE; int object = OBJECT_NONE;
#if BVH_FEATURE(BVH_MOTION) #if BVH_FEATURE(BVH_MOTION)
Transform ob_itfm; Transform ob_itfm;
#endif #endif
isect->t = ray->t; isect->t = ray->tmax;
isect->u = 0.0f; isect->u = 0.0f;
isect->v = 0.0f; isect->v = 0.0f;
isect->prim = PRIM_NONE; isect->prim = PRIM_NONE;
@@ -71,6 +72,7 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg,
dir, dir,
#endif #endif
idir, idir,
tmin,
isect->t, isect->t,
node_addr, node_addr,
visibility, visibility,
@@ -133,8 +135,16 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg,
switch (type & PRIMITIVE_ALL) { switch (type & PRIMITIVE_ALL) {
case PRIMITIVE_TRIANGLE: { case PRIMITIVE_TRIANGLE: {
if (triangle_intersect( if (triangle_intersect(kg,
kg, isect, P, dir, isect->t, visibility, prim_object, prim, prim_addr)) { isect,
P,
dir,
tmin,
isect->t,
visibility,
prim_object,
prim,
prim_addr)) {
/* shadow ray early termination */ /* shadow ray early termination */
if (visibility & PATH_RAY_SHADOW_OPAQUE) if (visibility & PATH_RAY_SHADOW_OPAQUE)
return true; return true;
@@ -147,6 +157,7 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg,
isect, isect,
P, P,
dir, dir,
tmin,
isect->t, isect->t,
ray->time, ray->time,
visibility, 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 int curve_type = kernel_data_fetch(prim_type, prim_addr);
const bool hit = curve_intersect( 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) { if (hit) {
/* shadow ray early termination */ /* shadow ray early termination */
if (visibility & PATH_RAY_SHADOW_OPAQUE) 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 int point_type = kernel_data_fetch(prim_type, prim_addr);
const bool hit = point_intersect( 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) { if (hit) {
/* shadow ray early termination */ /* shadow ray early termination */
if (visibility & PATH_RAY_SHADOW_OPAQUE) 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); object = kernel_data_fetch(prim_object, -prim_addr - 1);
#if BVH_FEATURE(BVH_MOTION) #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 #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 #endif
isect->t *= t_world_to_instance;
tmin *= t_world_to_instance;
++stack_ptr; ++stack_ptr;
kernel_assert(stack_ptr < BVH_STACK_SIZE); kernel_assert(stack_ptr < BVH_STACK_SIZE);
traversal_stack[stack_ptr] = ENTRYPOINT_SENTINEL; traversal_stack[stack_ptr] = ENTRYPOINT_SENTINEL;
@@ -235,6 +250,7 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg,
#else #else
isect->t = bvh_instance_pop(kg, object, ray, &P, &dir, &idir, isect->t); isect->t = bvh_instance_pop(kg, object, ray, &P, &dir, &idir, isect->t);
#endif #endif
tmin = ray->tmin;
object = OBJECT_NONE; object = OBJECT_NONE;
node_addr = traversal_stack[stack_ptr]; node_addr = traversal_stack[stack_ptr];

13
intern/cycles/kernel/bvh/util.h
View File

@@ -5,6 +5,19 @@
CCL_NAMESPACE_BEGIN 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__) #if defined(__KERNEL_CPU__)
ccl_device int intersections_compare(const void *a, const void *b) ccl_device int intersections_compare(const void *a, const void *b)
{ {

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

@@ -46,13 +46,14 @@ ccl_device_inline
float3 P = ray->P; float3 P = ray->P;
float3 dir = bvh_clamp_direction(ray->D); float3 dir = bvh_clamp_direction(ray->D);
float3 idir = bvh_inverse_direction(dir); float3 idir = bvh_inverse_direction(dir);
float tmin = ray->tmin;
int object = OBJECT_NONE; int object = OBJECT_NONE;
#if BVH_FEATURE(BVH_MOTION) #if BVH_FEATURE(BVH_MOTION)
Transform ob_itfm; Transform ob_itfm;
#endif #endif
isect->t = ray->t; isect->t = ray->tmax;
isect->u = 0.0f; isect->u = 0.0f;
isect->v = 0.0f; isect->v = 0.0f;
isect->prim = PRIM_NONE; isect->prim = PRIM_NONE;
@@ -73,6 +74,7 @@ ccl_device_inline
dir, dir,
#endif #endif
idir, idir,
tmin,
isect->t, isect->t,
node_addr, node_addr,
visibility, visibility,
@@ -140,7 +142,7 @@ ccl_device_inline
continue; continue;
} }
triangle_intersect( 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; break;
} }
@@ -165,6 +167,7 @@ ccl_device_inline
isect, isect,
P, P,
dir, dir,
tmin,
isect->t, isect->t,
ray->time, ray->time,
visibility, visibility,
@@ -186,11 +189,15 @@ ccl_device_inline
int object_flag = kernel_data_fetch(object_flag, object); int object_flag = kernel_data_fetch(object_flag, object);
if (object_flag & SD_OBJECT_HAS_VOLUME) { if (object_flag & SD_OBJECT_HAS_VOLUME) {
#if BVH_FEATURE(BVH_MOTION) #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 #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 #endif
isect->t *= t_world_to_instance;
tmin *= t_world_to_instance;
++stack_ptr; ++stack_ptr;
kernel_assert(stack_ptr < BVH_STACK_SIZE); kernel_assert(stack_ptr < BVH_STACK_SIZE);
traversal_stack[stack_ptr] = ENTRYPOINT_SENTINEL; 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); isect->t = bvh_instance_pop(kg, object, ray, &P, &dir, &idir, isect->t);
#endif #endif
tmin = ray->tmin;
object = OBJECT_NONE; object = OBJECT_NONE;
node_addr = traversal_stack[stack_ptr]; node_addr = traversal_stack[stack_ptr];
--stack_ptr; --stack_ptr;

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

@@ -44,12 +44,12 @@ ccl_device_inline
int node_addr = kernel_data.bvh.root; int node_addr = kernel_data.bvh.root;
/* ray parameters in registers */ /* ray parameters in registers */
const float tmax = ray->t;
float3 P = ray->P; float3 P = ray->P;
float3 dir = bvh_clamp_direction(ray->D); float3 dir = bvh_clamp_direction(ray->D);
float3 idir = bvh_inverse_direction(dir); float3 idir = bvh_inverse_direction(dir);
float tmin = ray->tmin;
int object = OBJECT_NONE; int object = OBJECT_NONE;
float isect_t = tmax; float isect_t = ray->tmax;
#if BVH_FEATURE(BVH_MOTION) #if BVH_FEATURE(BVH_MOTION)
Transform ob_itfm; Transform ob_itfm;
@@ -58,7 +58,7 @@ ccl_device_inline
int num_hits_in_instance = 0; int num_hits_in_instance = 0;
uint num_hits = 0; uint num_hits = 0;
isect_array->t = tmax; isect_array->t = ray->tmax;
/* traversal loop */ /* traversal loop */
do { do {
@@ -75,6 +75,7 @@ ccl_device_inline
dir, dir,
#endif #endif
idir, idir,
tmin,
isect_t, isect_t,
node_addr, node_addr,
visibility, visibility,
@@ -141,8 +142,16 @@ ccl_device_inline
if ((object_flag & SD_OBJECT_HAS_VOLUME) == 0) { if ((object_flag & SD_OBJECT_HAS_VOLUME) == 0) {
continue; continue;
} }
hit = triangle_intersect( hit = triangle_intersect(kg,
kg, isect_array, P, dir, isect_t, visibility, prim_object, prim, prim_addr); isect_array,
P,
dir,
tmin,
isect_t,
visibility,
prim_object,
prim,
prim_addr);
if (hit) { if (hit) {
/* Move on to next entry in intersections array. */ /* Move on to next entry in intersections array. */
isect_array++; isect_array++;
@@ -189,6 +198,7 @@ ccl_device_inline
isect_array, isect_array,
P, P,
dir, dir,
tmin,
isect_t, isect_t,
ray->time, ray->time,
visibility, visibility,
@@ -232,11 +242,15 @@ ccl_device_inline
int object_flag = kernel_data_fetch(object_flag, object); int object_flag = kernel_data_fetch(object_flag, object);
if (object_flag & SD_OBJECT_HAS_VOLUME) { if (object_flag & SD_OBJECT_HAS_VOLUME) {
#if BVH_FEATURE(BVH_MOTION) #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 #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 #endif
isect_t *= t_world_to_instance;
tmin *= t_world_to_instance;
num_hits_in_instance = 0; num_hits_in_instance = 0;
isect_array->t = isect_t; isect_array->t = isect_t;
@@ -280,7 +294,8 @@ ccl_device_inline
#endif #endif
} }
isect_t = tmax; tmin = ray->tmin;
isect_t = ray->tmax;
isect_array->t = isect_t; isect_array->t = isect_t;
object = OBJECT_NONE; object = OBJECT_NONE;

20
intern/cycles/kernel/camera/camera.h
View File

@@ -165,9 +165,11 @@ ccl_device void camera_sample_perspective(KernelGlobals kg,
float nearclip = kernel_data.cam.nearclip * z_inv; float nearclip = kernel_data.cam.nearclip * z_inv;
ray->P += nearclip * ray->D; ray->P += nearclip * ray->D;
ray->dP += nearclip * ray->dD; 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 #else
ray->t = FLT_MAX; ray->tmin = 0.0f;
ray->tmax = FLT_MAX;
#endif #endif
} }
@@ -231,9 +233,11 @@ ccl_device void camera_sample_orthographic(KernelGlobals kg,
#ifdef __CAMERA_CLIPPING__ #ifdef __CAMERA_CLIPPING__
/* clipping */ /* clipping */
ray->t = kernel_data.cam.cliplength; ray->tmin = 0.0f;
ray->tmax = kernel_data.cam.cliplength;
#else #else
ray->t = FLT_MAX; ray->tmin = 0.0f;
ray->tmax = FLT_MAX;
#endif #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 */ /* indicates ray should not receive any light, outside of the lens */
if (is_zero(D)) { if (is_zero(D)) {
ray->t = 0.0f; ray->tmax = 0.0f;
return; return;
} }
@@ -349,9 +353,11 @@ ccl_device_inline void camera_sample_panorama(ccl_constant KernelCamera *cam,
float nearclip = cam->nearclip; float nearclip = cam->nearclip;
ray->P += nearclip * ray->D; ray->P += nearclip * ray->D;
ray->dP += nearclip * ray->dD; ray->dP += nearclip * ray->dD;
ray->t = cam->cliplength; ray->tmin = 0.0f;
ray->tmax = cam->cliplength;
#else #else
ray->t = FLT_MAX; ray->tmin = 0.0f;
ray->tmax = FLT_MAX;
#endif #endif
} }

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

@@ -410,6 +410,7 @@ void metalrt_intersection_curve(constant KernelParamsMetal &launch_params_metal,
const float3 ray_origin, const float3 ray_origin,
const float3 ray_direction, const float3 ray_direction,
float time, float time,
const float ray_tmin,
const float ray_tmax, const float ray_tmax,
thread BoundingBoxIntersectionResult &result) thread BoundingBoxIntersectionResult &result)
{ {
@@ -434,7 +435,7 @@ void metalrt_intersection_curve(constant KernelParamsMetal &launch_params_metal,
isect.t *= len; isect.t *= len;
MetalKernelContext context(launch_params_metal); 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>( result = metalrt_visibility_test<BoundingBoxIntersectionResult, METALRT_HIT_BOUNDING_BOX>(
launch_params_metal, payload, object, prim, isect.u); launch_params_metal, payload, object, prim, isect.u);
if (result.accept) { if (result.accept) {
@@ -456,6 +457,7 @@ void metalrt_intersection_curve_shadow(constant KernelParamsMetal &launch_params
const float3 ray_origin, const float3 ray_origin,
const float3 ray_direction, const float3 ray_direction,
float time, float time,
const float ray_tmin,
const float ray_tmax, const float ray_tmax,
thread BoundingBoxIntersectionResult &result) thread BoundingBoxIntersectionResult &result)
{ {
@@ -475,7 +477,7 @@ void metalrt_intersection_curve_shadow(constant KernelParamsMetal &launch_params
isect.t *= len; isect.t *= len;
MetalKernelContext context(launch_params_metal); 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>( result.continue_search = metalrt_shadow_all_hit<METALRT_HIT_BOUNDING_BOX>(
launch_params_metal, payload, object, prim, float2(isect.u, isect.v), ray_tmax); launch_params_metal, payload, object, prim, float2(isect.u, isect.v), ray_tmax);
result.accept = !result.continue_search; 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 uint primitive_id [[primitive_id]],
const float3 ray_origin [[origin]], const float3 ray_origin [[origin]],
const float3 ray_direction [[direction]], const float3 ray_direction [[direction]],
const float ray_tmin [[min_distance]],
const float ray_tmax [[max_distance]]) const float ray_tmax [[max_distance]])
{ {
uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object); 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 # else
0.0f, 0.0f,
# endif # endif
ray_tmax, result); ray_tmin, ray_tmax, result);
} }
return result; return result;
@@ -525,6 +528,7 @@ __intersection__curve_ribbon_shadow(constant KernelParamsMetal &launch_params_me
const uint primitive_id [[primitive_id]], const uint primitive_id [[primitive_id]],
const float3 ray_origin [[origin]], const float3 ray_origin [[origin]],
const float3 ray_direction [[direction]], const float3 ray_direction [[direction]],
const float ray_tmin [[min_distance]],
const float ray_tmax [[max_distance]]) const float ray_tmax [[max_distance]])
{ {
uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object); 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 # else
0.0f, 0.0f,
# endif # endif
ray_tmax, result); ray_tmin, ray_tmax, result);
} }
return result; return result;
@@ -556,6 +560,7 @@ __intersection__curve_all(constant KernelParamsMetal &launch_params_metal [[buff
const uint primitive_id [[primitive_id]], const uint primitive_id [[primitive_id]],
const float3 ray_origin [[origin]], const float3 ray_origin [[origin]],
const float3 ray_direction [[direction]], const float3 ray_direction [[direction]],
const float ray_tmin [[min_distance]],
const float ray_tmax [[max_distance]]) const float ray_tmax [[max_distance]])
{ {
uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object); 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 # else
0.0f, 0.0f,
# endif # endif
ray_tmax, result); ray_tmin, ray_tmax, result);
return result; return result;
} }
@@ -584,6 +589,7 @@ __intersection__curve_all_shadow(constant KernelParamsMetal &launch_params_metal
const uint primitive_id [[primitive_id]], const uint primitive_id [[primitive_id]],
const float3 ray_origin [[origin]], const float3 ray_origin [[origin]],
const float3 ray_direction [[direction]], const float3 ray_direction [[direction]],
const float ray_tmin [[min_distance]],
const float ray_tmax [[max_distance]]) const float ray_tmax [[max_distance]])
{ {
uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object); 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 # else
0.0f, 0.0f,
# endif # endif
ray_tmax, result); ray_tmin, ray_tmax, result);
return result; return result;
} }
@@ -616,6 +622,7 @@ void metalrt_intersection_point(constant KernelParamsMetal &launch_params_metal,
const float3 ray_origin, const float3 ray_origin,
const float3 ray_direction, const float3 ray_direction,
float time, float time,
const float ray_tmin,
const float ray_tmax, const float ray_tmax,
thread BoundingBoxIntersectionResult &result) thread BoundingBoxIntersectionResult &result)
{ {
@@ -640,7 +647,7 @@ void metalrt_intersection_point(constant KernelParamsMetal &launch_params_metal,
isect.t *= len; isect.t *= len;
MetalKernelContext context(launch_params_metal); 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>( result = metalrt_visibility_test<BoundingBoxIntersectionResult, METALRT_HIT_BOUNDING_BOX>(
launch_params_metal, payload, object, prim, isect.u); launch_params_metal, payload, object, prim, isect.u);
if (result.accept) { if (result.accept) {
@@ -662,6 +669,7 @@ void metalrt_intersection_point_shadow(constant KernelParamsMetal &launch_params
const float3 ray_origin, const float3 ray_origin,
const float3 ray_direction, const float3 ray_direction,
float time, float time,
const float ray_tmin,
const float ray_tmax, const float ray_tmax,
thread BoundingBoxIntersectionResult &result) thread BoundingBoxIntersectionResult &result)
{ {
@@ -681,7 +689,7 @@ void metalrt_intersection_point_shadow(constant KernelParamsMetal &launch_params
isect.t *= len; isect.t *= len;
MetalKernelContext context(launch_params_metal); 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>( result.continue_search = metalrt_shadow_all_hit<METALRT_HIT_BOUNDING_BOX>(
launch_params_metal, payload, object, prim, float2(isect.u, isect.v), ray_tmax); launch_params_metal, payload, object, prim, float2(isect.u, isect.v), ray_tmax);
result.accept = !result.continue_search; 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 uint primitive_id [[primitive_id]],
const float3 ray_origin [[origin]], const float3 ray_origin [[origin]],
const float3 ray_direction [[direction]], const float3 ray_direction [[direction]],
const float ray_tmin [[min_distance]],
const float ray_tmax [[max_distance]]) const float ray_tmax [[max_distance]])
{ {
const uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object); 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 # else
0.0f, 0.0f,
# endif # endif
ray_tmax, result); ray_tmin, ray_tmax, result);
return result; return result;
} }
@@ -729,6 +738,7 @@ __intersection__point_shadow(constant KernelParamsMetal &launch_params_metal [[b
const uint primitive_id [[primitive_id]], const uint primitive_id [[primitive_id]],
const float3 ray_origin [[origin]], const float3 ray_origin [[origin]],
const float3 ray_direction [[direction]], const float3 ray_direction [[direction]],
const float ray_tmin [[min_distance]],
const float ray_tmax [[max_distance]]) const float ray_tmax [[max_distance]])
{ {
const uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object); 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 # else
0.0f, 0.0f,
# endif # endif
ray_tmax, result); ray_tmin, ray_tmax, result);
return result; return result;
} }

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

@@ -51,32 +51,36 @@ ccl_device_forceinline int get_object_id()
extern "C" __global__ void __raygen__kernel_optix_integrator_intersect_closest() extern "C" __global__ void __raygen__kernel_optix_integrator_intersect_closest()
{ {
const int global_index = optixGetLaunchIndex().x; const int global_index = optixGetLaunchIndex().x;
const int path_index = (kernel_params.path_index_array) ? kernel_params.path_index_array[global_index] : const int path_index = (kernel_params.path_index_array) ?
global_index; kernel_params.path_index_array[global_index] :
global_index;
integrator_intersect_closest(nullptr, path_index, kernel_params.render_buffer); integrator_intersect_closest(nullptr, path_index, kernel_params.render_buffer);
} }
extern "C" __global__ void __raygen__kernel_optix_integrator_intersect_shadow() extern "C" __global__ void __raygen__kernel_optix_integrator_intersect_shadow()
{ {
const int global_index = optixGetLaunchIndex().x; const int global_index = optixGetLaunchIndex().x;
const int path_index = (kernel_params.path_index_array) ? kernel_params.path_index_array[global_index] : const int path_index = (kernel_params.path_index_array) ?
global_index; kernel_params.path_index_array[global_index] :
global_index;
integrator_intersect_shadow(nullptr, path_index); integrator_intersect_shadow(nullptr, path_index);
} }
extern "C" __global__ void __raygen__kernel_optix_integrator_intersect_subsurface() extern "C" __global__ void __raygen__kernel_optix_integrator_intersect_subsurface()
{ {
const int global_index = optixGetLaunchIndex().x; const int global_index = optixGetLaunchIndex().x;
const int path_index = (kernel_params.path_index_array) ? kernel_params.path_index_array[global_index] : const int path_index = (kernel_params.path_index_array) ?
global_index; kernel_params.path_index_array[global_index] :
global_index;
integrator_intersect_subsurface(nullptr, path_index); integrator_intersect_subsurface(nullptr, path_index);
} }
extern "C" __global__ void __raygen__kernel_optix_integrator_intersect_volume_stack() extern "C" __global__ void __raygen__kernel_optix_integrator_intersect_volume_stack()
{ {
const int global_index = optixGetLaunchIndex().x; const int global_index = optixGetLaunchIndex().x;
const int path_index = (kernel_params.path_index_array) ? kernel_params.path_index_array[global_index] : const int path_index = (kernel_params.path_index_array) ?
global_index; kernel_params.path_index_array[global_index] :
global_index;
integrator_intersect_volume_stack(nullptr, path_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 P = optixGetObjectRayOrigin();
float3 dir = optixGetObjectRayDirection(); float3 dir = optixGetObjectRayDirection();
float tmin = optixGetRayTmin();
/* The direction is not normalized by default, but the curve intersection routine expects that */ /* The direction is not normalized by default, but the curve intersection routine expects that */
float len; float len;
@@ -425,7 +430,7 @@ ccl_device_inline void optix_intersection_curve(const int prim, const int type)
if (isect.t != FLT_MAX) if (isect.t != FLT_MAX)
isect.t *= len; 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"); static_assert(PRIMITIVE_ALL < 128, "Values >= 128 are reserved for OptiX internal use");
optixReportIntersection(isect.t / len, optixReportIntersection(isect.t / len,
type & PRIMITIVE_ALL, type & PRIMITIVE_ALL,
@@ -462,6 +467,7 @@ extern "C" __global__ void __intersection__point()
float3 P = optixGetObjectRayOrigin(); float3 P = optixGetObjectRayOrigin();
float3 dir = optixGetObjectRayDirection(); float3 dir = optixGetObjectRayDirection();
float tmin = optixGetRayTmin();
/* The direction is not normalized by default, the point intersection routine expects that. */ /* The direction is not normalized by default, the point intersection routine expects that. */
float len; float len;
@@ -480,7 +486,7 @@ extern "C" __global__ void __intersection__point()
isect.t *= len; 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"); static_assert(PRIMITIVE_ALL < 128, "Values >= 128 are reserved for OptiX internal use");
optixReportIntersection(isect.t / len, type & PRIMITIVE_ALL); optixReportIntersection(isect.t / len, type & PRIMITIVE_ALL);
} }

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

@@ -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_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 float dt,
const float4 curve[4], const float4 curve[4],
float u, 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) { if (fabsf(f) < f_err && fabsf(g) < g_err) {
t += dt; t += dt;
if (!(0.0f <= t && t <= *ray_tfar)) { if (!(t >= ray_tmin && t <= *ray_tmax)) {
return false; /* Rejects NaNs */ return false; /* Rejects NaNs */
} }
if (!(u >= 0.0f && u <= 1.0f)) { if (!(u >= 0.0f && u <= 1.0f)) {
@@ -237,7 +238,7 @@ ccl_device bool curve_intersect_iterative(const float3 ray_dir,
} }
/* Record intersection. */ /* Record intersection. */
*ray_tfar = t; *ray_tmax = t;
isect->t = t; isect->t = t;
isect->u = u; isect->u = u;
isect->v = 0.0f; 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, ccl_device bool curve_intersect_recursive(const float3 ray_orig,
const float3 ray_dir, const float3 ray_dir,
float ray_tfar, const float ray_tmin,
float ray_tmax,
float4 curve[4], float4 curve[4],
ccl_private Intersection *isect) ccl_private Intersection *isect)
{ {
@@ -331,7 +333,7 @@ ccl_device bool curve_intersect_recursive(const float3 ray_orig,
} }
/* Intersect with cap-planes. */ /* 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)); tp = make_float2(max(tp.x, tc_outer.x), min(tp.y, tc_outer.y));
const float2 h0 = half_plane_intersect( const float2 h0 = half_plane_intersect(
float4_to_float3(P0), float4_to_float3(dP0du), ray_dir); 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; CURVE_NUM_BEZIER_SUBDIVISIONS;
if (depth >= termDepth) { if (depth >= termDepth) {
found |= curve_intersect_iterative( 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 { else {
recurse = true; 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 : const int termDepth = unstable1 ? CURVE_NUM_BEZIER_SUBDIVISIONS_UNSTABLE :
CURVE_NUM_BEZIER_SUBDIVISIONS; CURVE_NUM_BEZIER_SUBDIVISIONS;
if (depth >= termDepth) { if (depth >= termDepth) {
found |= curve_intersect_iterative( 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 { else {
recurse = true; 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 * v0,v1,v3 and v2,v3,v1. The edge v1,v2 decides which of the two
* triangles gets intersected. * 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_v0,
const float3 quad_v1, const float3 quad_v1,
const float3 quad_v2, const float3 quad_v2,
@@ -497,7 +501,7 @@ ccl_device_inline bool ribbon_intersect_quad(const float ray_tfar,
/* Perform depth test? */ /* Perform depth test? */
const float t = rcpDen * dot(v0, Ng); const float t = rcpDen * dot(v0, Ng);
if (!(0.0f <= t && t <= ray_tfar)) { if (!(t >= ray_tmin && t <= ray_tmax)) {
return false; 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, ccl_device_inline bool ribbon_intersect(const float3 ray_org,
const float3 ray_dir, const float3 ray_dir,
float ray_tfar, const float ray_tmin,
float ray_tmax,
const int N, const int N,
float4 curve[4], float4 curve[4],
ccl_private Intersection *isect) ccl_private Intersection *isect)
@@ -582,7 +587,7 @@ ccl_device_inline bool ribbon_intersect(const float3 ray_org,
/* Intersect quad. */ /* Intersect quad. */
float vu, vv, vt; 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) { if (valid0) {
/* ignore self intersections */ /* ignore self intersections */
@@ -596,7 +601,7 @@ ccl_device_inline bool ribbon_intersect(const float3 ray_org,
vv = 2.0f * vv - 1.0f; vv = 2.0f * vv - 1.0f;
/* Record intersection. */ /* Record intersection. */
ray_tfar = vt; ray_tmax = vt;
isect->t = vt; isect->t = vt;
isect->u = u + vu * step_size; isect->u = u + vu * step_size;
isect->v = vv; isect->v = vv;
@@ -616,6 +621,7 @@ ccl_device_forceinline bool curve_intersect(KernelGlobals kg,
ccl_private Intersection *isect, ccl_private Intersection *isect,
const float3 P, const float3 P,
const float3 dir, const float3 dir,
const float tmin,
const float tmax, const float tmax,
int object, int object,
int prim, int prim,
@@ -645,7 +651,7 @@ ccl_device_forceinline bool curve_intersect(KernelGlobals kg,
if (type & PRIMITIVE_CURVE_RIBBON) { if (type & PRIMITIVE_CURVE_RIBBON) {
/* todo: adaptive number of subdivisions could help performance here. */ /* todo: adaptive number of subdivisions could help performance here. */
const int subdivisions = kernel_data.bvh.curve_subdivisions; 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->prim = prim;
isect->object = object; isect->object = object;
isect->type = type; isect->type = type;
@@ -655,7 +661,7 @@ ccl_device_forceinline bool curve_intersect(KernelGlobals kg,
return false; return false;
} }
else { else {
if (curve_intersect_recursive(P, dir, tmax, curve, isect)) { if (curve_intersect_recursive(P, dir, tmin, tmax, curve, isect)) {
isect->prim = prim; isect->prim = prim;
isect->object = object; isect->object = object;
isect->type = type; isect->type = type;

6
intern/cycles/kernel/geom/motion_triangle_intersect.h
View File

@@ -46,6 +46,7 @@ ccl_device_inline bool motion_triangle_intersect(KernelGlobals kg,
ccl_private Intersection *isect, ccl_private Intersection *isect,
float3 P, float3 P,
float3 dir, float3 dir,
float tmin,
float tmax, float tmax,
float time, float time,
uint visibility, uint visibility,
@@ -58,7 +59,7 @@ ccl_device_inline bool motion_triangle_intersect(KernelGlobals kg,
motion_triangle_vertices(kg, object, prim, time, verts); motion_triangle_vertices(kg, object, prim, time, verts);
/* Ray-triangle intersection, unoptimized. */ /* Ray-triangle intersection, unoptimized. */
float t, u, v; 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__ #ifdef __VISIBILITY_FLAG__
/* Visibility flag test. we do it here under the assumption /* Visibility flag test. we do it here under the assumption
* that most triangles are culled by node flags. * 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 object,
int prim, int prim,
int prim_addr, int prim_addr,
float tmin,
float tmax, float tmax,
ccl_private uint *lcg_state, ccl_private uint *lcg_state,
int max_hits) 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); motion_triangle_vertices(kg, object, prim, time, verts);
/* Ray-triangle intersection, unoptimized. */ /* Ray-triangle intersection, unoptimized. */
float t, u, v; 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; return false;
} }

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

@@ -9,8 +9,12 @@ CCL_NAMESPACE_BEGIN
#ifdef __POINTCLOUD__ #ifdef __POINTCLOUD__
ccl_device_forceinline bool point_intersect_test( ccl_device_forceinline bool point_intersect_test(const float4 point,
const float4 point, const float3 P, const float3 dir, const float tmax, ccl_private float *t) const float3 P,
const float3 dir,
const float tmin,
const float tmax,
ccl_private float *t)
{ {
const float3 center = float4_to_float3(point); const float3 center = float4_to_float3(point);
const float radius = point.w; 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 td = sqrt((r2 - l2) * rd2);
const float t_front = projC0 - td; 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. */ /* Always back-face culling for now. */
# if 0 # if 0
const float t_back = projC0 + td; 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 */ /* check if there is a first hit */
const bool valid_first = valid_front | valid_back; const bool valid_first = valid_front | valid_back;
@@ -56,6 +60,7 @@ ccl_device_forceinline bool point_intersect(KernelGlobals kg,
ccl_private Intersection *isect, ccl_private Intersection *isect,
const float3 P, const float3 P,
const float3 dir, const float3 dir,
const float tmin,
const float tmax, const float tmax,
const int object, const int object,
const int prim, 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) : const float4 point = (type & PRIMITIVE_MOTION) ? motion_point(kg, object, prim, time) :
kernel_data_fetch(points, prim); 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; return false;
} }

3
intern/cycles/kernel/geom/shader_data.h
View File

@@ -407,7 +407,7 @@ ccl_device_inline void shader_setup_from_volume(KernelGlobals kg,
{ {
/* vectors */ /* vectors */
sd->P = ray->P; sd->P = ray->P + ray->D * ray->tmin;
sd->N = -ray->D; sd->N = -ray->D;
sd->Ng = -ray->D; sd->Ng = -ray->D;
sd->I = -ray->D; sd->I = -ray->D;
@@ -441,7 +441,6 @@ ccl_device_inline void shader_setup_from_volume(KernelGlobals kg,
/* for NDC coordinates */ /* for NDC coordinates */
sd->ray_P = ray->P; sd->ray_P = ray->P;
sd->ray_dP = ray->dP;
} }
#endif /* __VOLUME__ */ #endif /* __VOLUME__ */

6
intern/cycles/kernel/geom/triangle_intersect.h
View File

@@ -17,6 +17,7 @@ ccl_device_inline bool triangle_intersect(KernelGlobals kg,
ccl_private Intersection *isect, ccl_private Intersection *isect,
float3 P, float3 P,
float3 dir, float3 dir,
float tmin,
float tmax, float tmax,
uint visibility, uint visibility,
int object, 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_b = kernel_data_fetch(tri_verts, tri_vindex + 1),
tri_c = kernel_data_fetch(tri_verts, tri_vindex + 2); tri_c = kernel_data_fetch(tri_verts, tri_vindex + 2);
float t, u, v; 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__ #ifdef __VISIBILITY_FLAG__
/* Visibility flag test. we do it here under the assumption /* Visibility flag test. we do it here under the assumption
* that most triangles are culled by node flags. * that most triangles are culled by node flags.
@@ -62,6 +63,7 @@ ccl_device_inline bool triangle_intersect_local(KernelGlobals kg,
int object, int object,
int prim, int prim,
int prim_addr, int prim_addr,
float tmin,
float tmax, float tmax,
ccl_private uint *lcg_state, ccl_private uint *lcg_state,
int max_hits) 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_b = kernel_data_fetch(tri_verts, tri_vindex + 1),
tri_c = kernel_data_fetch(tri_verts, tri_vindex + 2); tri_c = kernel_data_fetch(tri_verts, tri_vindex + 2);
float t, u, v; 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; return false;
} }

6
intern/cycles/kernel/integrator/init_from_bake.h
View File

@@ -174,7 +174,8 @@ ccl_device bool integrator_init_from_bake(KernelGlobals kg,
Ray ray ccl_optional_struct_init; Ray ray ccl_optional_struct_init;
ray.P = zero_float3(); ray.P = zero_float3();
ray.D = normalize(P); ray.D = normalize(P);
ray.t = FLT_MAX; ray.tmin = 0.0f;
ray.tmax = FLT_MAX;
ray.time = 0.5f; ray.time = 0.5f;
ray.dP = differential_zero_compact(); ray.dP = differential_zero_compact();
ray.dD = 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 ray ccl_optional_struct_init;
ray.P = P + N; ray.P = P + N;
ray.D = -N; ray.D = -N;
ray.t = FLT_MAX; ray.tmin = 0.0f;
ray.tmax = FLT_MAX;
ray.time = 0.5f; ray.time = 0.5f;
/* Setup differentials. */ /* Setup differentials. */

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

@@ -86,7 +86,7 @@ ccl_device bool integrator_init_from_camera(KernelGlobals kg,
/* Generate camera ray. */ /* Generate camera ray. */
Ray ray; Ray ray;
integrate_camera_sample(kg, sample, x, y, rng_hash, &ray); integrate_camera_sample(kg, sample, x, y, rng_hash, &ray);
if (ray.t == 0.0f) { if (ray.tmax == 0.0f) {
return true; return true;
} }

6
intern/cycles/kernel/integrator/intersect_closest.h
View File

@@ -324,7 +324,7 @@ ccl_device void integrator_intersect_closest(KernelGlobals kg,
/* Read ray from integrator state into local memory. */ /* Read ray from integrator state into local memory. */
Ray ray ccl_optional_struct_init; Ray ray ccl_optional_struct_init;
integrator_state_read_ray(kg, state, &ray); 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 uint visibility = path_state_ray_visibility(state);
const int last_isect_prim = INTEGRATOR_STATE(state, isect, prim); 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. */ /* Trick to use short AO rays to approximate indirect light at the end of the path. */
if (path_state_ao_bounce(kg, state)) { 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) { if (last_isect_object != OBJECT_NONE) {
const float object_ao_distance = kernel_data_fetch(objects, last_isect_object).ao_distance; const float object_ao_distance = kernel_data_fetch(objects, last_isect_object).ao_distance;
if (object_ao_distance != 0.0f) { if (object_ao_distance != 0.0f) {
ray.t = object_ao_distance; ray.tmax = object_ao_distance;
} }
} }
} }

13
intern/cycles/kernel/integrator/intersect_volume_stack.h
View File

@@ -24,7 +24,8 @@ ccl_device void integrator_volume_stack_update_for_subsurface(KernelGlobals kg,
Ray volume_ray ccl_optional_struct_init; Ray volume_ray ccl_optional_struct_init;
volume_ray.P = from_P; 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.object = INTEGRATOR_STATE(state, isect, object);
volume_ray.self.prim = INTEGRATOR_STATE(state, isect, prim); volume_ray.self.prim = INTEGRATOR_STATE(state, isect, prim);
volume_ray.self.light_object = OBJECT_NONE; 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); volume_stack_enter_exit(kg, state, stack_sd);
/* Move ray forward. */ /* 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.object = isect.object;
volume_ray.self.prim = isect.prim; 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; ++step;
} }
#endif #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 /* Trace ray in random direction. Any direction works, Z up is a guess to get the
* fewest hits. */ * fewest hits. */
volume_ray.D = make_float3(0.0f, 0.0f, 1.0f); 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.object = OBJECT_NONE;
volume_ray.self.prim = PRIM_NONE; volume_ray.self.prim = PRIM_NONE;
volume_ray.self.light_object = OBJECT_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. */ /* 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.object = isect.object;
volume_ray.self.prim = isect.prim; volume_ray.self.prim = isect.prim;
++step; ++step;

22
intern/cycles/kernel/integrator/mnee.h
View File

@@ -442,6 +442,7 @@ ccl_device_forceinline bool mnee_newton_solver(KernelGlobals kg,
projection_ray.self.light_prim = PRIM_NONE; projection_ray.self.light_prim = PRIM_NONE;
projection_ray.dP = differential_make_compact(sd->dP); projection_ray.dP = differential_make_compact(sd->dP);
projection_ray.dD = differential_zero_compact(); projection_ray.dD = differential_zero_compact();
projection_ray.tmin = 0.0f;
projection_ray.time = sd->time; projection_ray.time = sd->time;
Intersection projection_isect; Intersection projection_isect;
@@ -505,8 +506,8 @@ ccl_device_forceinline bool mnee_newton_solver(KernelGlobals kg,
projection_ray.self.prim = pv.prim; projection_ray.self.prim = pv.prim;
projection_ray.P = pv.p; projection_ray.P = pv.p;
} }
projection_ray.D = normalize_len(tentative_p - projection_ray.P, &projection_ray.t); projection_ray.D = normalize_len(tentative_p - projection_ray.P, &projection_ray.tmax);
projection_ray.t *= MNEE_PROJECTION_DISTANCE_MULTIPLIER; projection_ray.tmax *= MNEE_PROJECTION_DISTANCE_MULTIPLIER;
bool projection_success = false; bool projection_success = false;
for (int isect_count = 0; isect_count < MNEE_MAX_INTERSECTION_COUNT; isect_count++) { 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.object = projection_isect.object;
projection_ray.self.prim = projection_isect.prim; projection_ray.self.prim = projection_isect.prim;
projection_ray.P += projection_isect.t * projection_ray.D; projection_ray.tmin = intersection_t_offset(projection_isect.t);
projection_ray.t -= projection_isect.t;
} }
if (!projection_success) { if (!projection_success) {
reduce_stepsize = true; reduce_stepsize = true;
@@ -858,6 +858,7 @@ ccl_device_forceinline bool mnee_path_contribution(KernelGlobals kg,
Ray probe_ray; Ray probe_ray;
probe_ray.self.light_object = ls->object; probe_ray.self.light_object = ls->object;
probe_ray.self.light_prim = ls->prim; probe_ray.self.light_prim = ls->prim;
probe_ray.tmin = 0.0f;
probe_ray.dP = differential_make_compact(sd->dP); probe_ray.dP = differential_make_compact(sd->dP);
probe_ray.dD = differential_zero_compact(); probe_ray.dD = differential_zero_compact();
probe_ray.time = sd->time; 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]; ccl_private const ManifoldVertex &v = vertices[vi];
/* Check visibility. */ /* 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)) { if (scene_intersect(kg, &probe_ray, PATH_RAY_TRANSMIT, &probe_isect)) {
int hit_object = (probe_isect.object == OBJECT_NONE) ? int hit_object = (probe_isect.object == OBJECT_NONE) ?
kernel_data_fetch(prim_object, probe_isect.prim) : kernel_data_fetch(prim_object, probe_isect.prim) :
probe_isect.object; probe_isect.object;
/* Test whether the ray hit the appropriate object at its intended location. */ /* 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; return false;
} }
probe_ray.self.object = v.object; 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_object = ls->object;
probe_ray.self.light_prim = ls->prim; probe_ray.self.light_prim = ls->prim;
probe_ray.P = sd->P; probe_ray.P = sd->P;
probe_ray.tmin = 0.0f;
if (ls->t == FLT_MAX) { if (ls->t == FLT_MAX) {
/* Distant / env light. */ /* Distant / env light. */
probe_ray.D = ls->D; probe_ray.D = ls->D;
probe_ray.t = ls->t; probe_ray.tmax = ls->t;
} }
else { else {
/* Other lights, avoid self-intersection. */ /* Other lights, avoid self-intersection. */
probe_ray.D = ls->P - probe_ray.P; 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.dP = differential_make_compact(sd->dP);
probe_ray.dD = differential_zero_compact(); 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.object = probe_isect.object;
probe_ray.self.prim = probe_isect.prim; probe_ray.self.prim = probe_isect.prim;
probe_ray.P += probe_isect.t * probe_ray.D; probe_ray.tmin = intersection_t_offset(probe_isect.t);
if (ls->t != FLT_MAX)
probe_ray.t -= probe_isect.t;
}; };
/* Mark the manifold walk invalid to keep mollification on by default. */ /* Mark the manifold walk invalid to keep mollification on by default. */

1
intern/cycles/kernel/integrator/path_state.h
View File

@@ -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 | INTEGRATOR_STATE_WRITE(state, path, flag) = PATH_RAY_CAMERA | PATH_RAY_MIS_SKIP |
PATH_RAY_TRANSPARENT_BACKGROUND; PATH_RAY_TRANSPARENT_BACKGROUND;
INTEGRATOR_STATE_WRITE(state, path, mis_ray_pdf) = 0.0f; 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, min_ray_pdf) = FLT_MAX;
INTEGRATOR_STATE_WRITE(state, path, continuation_probability) = 1.0f; INTEGRATOR_STATE_WRITE(state, path, continuation_probability) = 1.0f;
INTEGRATOR_STATE_WRITE(state, path, throughput) = make_float3(1.0f, 1.0f, 1.0f); INTEGRATOR_STATE_WRITE(state, path, throughput) = make_float3(1.0f, 1.0f, 1.0f);

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

@@ -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_P = INTEGRATOR_STATE(state, ray, P);
const float3 ray_D = INTEGRATOR_STATE(state, ray, D); 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_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 /* multiple importance sampling, get background light pdf for ray
* direction, and compute weight with respect to BSDF pdf */ * 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); const float mis_weight = light_sample_mis_weight_forward(kg, mis_ray_pdf, pdf);
L *= mis_weight; L *= mis_weight;
} }

15
intern/cycles/kernel/integrator/shade_light.h
View File

@@ -22,19 +22,8 @@ ccl_device_inline void integrate_light(KernelGlobals kg,
const float3 ray_D = INTEGRATOR_STATE(state, ray, D); const float3 ray_D = INTEGRATOR_STATE(state, ray, D);
const float ray_time = INTEGRATOR_STATE(state, ray, time); const float ray_time = INTEGRATOR_STATE(state, ray, time);
/* Advance ray beyond light. */ /* Advance ray to new start distance. */
/* TODO: can we make this more numerically robust to avoid reintersecting the INTEGRATOR_STATE_WRITE(state, ray, tmin) = intersection_t_offset(isect.t);
* 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;
LightSample ls ccl_optional_struct_init; LightSample ls ccl_optional_struct_init;
const bool use_light_sample = light_sample_from_intersection(kg, &isect, ray_P, ray_D, &ls); const bool use_light_sample = light_sample_from_intersection(kg, &isect, ray_P, ray_D, &ls);

15
intern/cycles/kernel/integrator/shade_shadow.h
View File

@@ -75,13 +75,9 @@ ccl_device_inline void integrate_transparent_volume_shadow(KernelGlobals kg,
ray.self.light_object = OBJECT_NONE; ray.self.light_object = OBJECT_NONE;
ray.self.light_prim = PRIM_NONE; ray.self.light_prim = PRIM_NONE;
/* Modify ray position and length to match current segment. */ /* Modify ray position and length to match current segment. */
const float start_t = (hit == 0) ? 0.0f : ray.tmin = (hit == 0) ? ray.tmin : INTEGRATOR_STATE_ARRAY(state, shadow_isect, hit - 1, t);
INTEGRATOR_STATE_ARRAY(state, shadow_isect, hit - 1, t); ray.tmax = (hit < num_recorded_hits) ? INTEGRATOR_STATE_ARRAY(state, shadow_isect, hit, t) :
const float end_t = (hit < num_recorded_hits) ? ray.tmax;
INTEGRATOR_STATE_ARRAY(state, shadow_isect, hit, t) :
ray.t;
ray.P += start_t * ray.D;
ray.t = end_t - start_t;
shader_setup_from_volume(kg, shadow_sd, &ray); 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, /* There are more hits that we could not recorded due to memory usage,
* adjust ray to intersect again from the last hit. */ * 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 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); INTEGRATOR_STATE_WRITE(state, shadow_ray, tmin) = intersection_t_offset(last_hit_t);
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;
} }
return false; return false;

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

@@ -77,7 +77,7 @@ ccl_device_forceinline void integrate_surface_emission(KernelGlobals kg,
# endif # endif
{ {
const float bsdf_pdf = INTEGRATOR_STATE(state, path, mis_ray_pdf); 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, /* Multiple importance sampling, get triangle light pdf,
* and compute weight with respect to BSDF 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; return LABEL_NONE;
} }
/* Setup ray. Note that clipping works through transparent bounces. */ if (label & LABEL_TRANSPARENT) {
INTEGRATOR_STATE_WRITE(state, ray, P) = sd->P; /* Only need to modify start distance for transparent. */
INTEGRATOR_STATE_WRITE(state, ray, D) = normalize(bsdf_omega_in); INTEGRATOR_STATE_WRITE(state, ray, tmin) = intersection_t_offset(sd->ray_length);
INTEGRATOR_STATE_WRITE(state, ray, t) = (label & LABEL_TRANSPARENT) ? }
INTEGRATOR_STATE(state, ray, t) - sd->ray_length : else {
FLT_MAX; /* 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__ #ifdef __RAY_DIFFERENTIALS__
INTEGRATOR_STATE_WRITE(state, ray, dP) = differential_make_compact(sd->dP); 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, dD) = differential_make_compact(bsdf_domega_in);
#endif #endif
}
/* Update throughput. */ /* Update throughput. */
float3 throughput = INTEGRATOR_STATE(state, path, throughput); float3 throughput = INTEGRATOR_STATE(state, path, throughput);
@@ -349,12 +354,8 @@ ccl_device_forceinline int integrate_surface_bsdf_bssrdf_bounce(
} }
/* Update path state */ /* Update path state */
if (label & LABEL_TRANSPARENT) { if (!(label & LABEL_TRANSPARENT)) {
INTEGRATOR_STATE_WRITE(state, path, mis_ray_t) += sd->ray_length;
}
else {
INTEGRATOR_STATE_WRITE(state, path, mis_ray_pdf) = bsdf_pdf; 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( INTEGRATOR_STATE_WRITE(state, path, min_ray_pdf) = fminf(
bsdf_pdf, INTEGRATOR_STATE(state, path, min_ray_pdf)); 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; return LABEL_NONE;
} }
/* Setup ray position, direction stays unchanged. */ /* Only modify start distance. */
INTEGRATOR_STATE_WRITE(state, ray, P) = sd->P; INTEGRATOR_STATE_WRITE(state, ray, tmin) = intersection_t_offset(sd->ray_length);
/* 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;
return LABEL_TRANSMIT | LABEL_TRANSPARENT; 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 ray ccl_optional_struct_init;
ray.P = shadow_ray_offset(kg, sd, ao_D, &skip_self); ray.P = shadow_ray_offset(kg, sd, ao_D, &skip_self);
ray.D = ao_D; 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.time = sd->time;
ray.self.object = (skip_self) ? sd->object : OBJECT_NONE; ray.self.object = (skip_self) ? sd->object : OBJECT_NONE;
ray.self.prim = (skip_self) ? sd->prim : PRIM_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); kg, state, current_kernel, DEVICE_KERNEL_INTEGRATOR_INTERSECT_SUBSURFACE);
} }
else { 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); integrator_path_next(kg, state, current_kernel, DEVICE_KERNEL_INTEGRATOR_INTERSECT_CLOSEST);
} }
} }

73
intern/cycles/kernel/integrator/shade_volume.h
View File

@@ -114,7 +114,8 @@ ccl_device_inline bool volume_shader_sample(KernelGlobals kg,
ccl_device_forceinline void volume_step_init(KernelGlobals kg, ccl_device_forceinline void volume_step_init(KernelGlobals kg,
ccl_private const RNGState *rng_state, ccl_private const RNGState *rng_state,
const float object_step_size, const float object_step_size,
float t, const float tmin,
const float tmax,
ccl_private float *step_size, ccl_private float *step_size,
ccl_private float *step_shade_offset, ccl_private float *step_shade_offset,
ccl_private float *steps_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) { if (object_step_size == FLT_MAX) {
/* Homogeneous volume. */ /* Homogeneous volume. */
*step_size = t; *step_size = tmax - tmin;
*step_shade_offset = 0.0f; *step_shade_offset = 0.0f;
*steps_offset = 1.0f; *steps_offset = 1.0f;
*max_steps = 1; *max_steps = 1;
@@ -130,6 +131,7 @@ ccl_device_forceinline void volume_step_init(KernelGlobals kg,
else { else {
/* Heterogeneous volume. */ /* Heterogeneous volume. */
*max_steps = kernel_data.integrator.volume_max_steps; *max_steps = kernel_data.integrator.volume_max_steps;
const float t = tmax - tmin;
float step = min(object_step_size, t); float step = min(object_step_size, t);
/* compute exact steps in advance for malloc */ /* 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(); float3 sigma_t = zero_float3();
if (shadow_volume_shader_sample(kg, state, sd, &sigma_t)) { 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 # endif
@@ -194,7 +196,8 @@ ccl_device void volume_shadow_heterogeneous(KernelGlobals kg,
volume_step_init(kg, volume_step_init(kg,
&rng_state, &rng_state,
object_step_size, object_step_size,
ray->t, ray->tmin,
ray->tmax,
&step_size, &step_size,
&step_shade_offset, &step_shade_offset,
&unused, &unused,
@@ -202,13 +205,13 @@ ccl_device void volume_shadow_heterogeneous(KernelGlobals kg,
const float steps_offset = 1.0f; const float steps_offset = 1.0f;
/* compute extinction at the start */ /* compute extinction at the start */
float t = 0.0f; float t = ray->tmin;
float3 sum = zero_float3(); float3 sum = zero_float3();
for (int i = 0; i < max_steps; i++) { for (int i = 0; i < max_steps; i++) {
/* advance to new position */ /* 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; float dt = new_t - t;
float3 new_P = ray->P + ray->D * (t + dt * step_shade_offset); 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 */ /* stop if at the end of the volume */
t = new_t; t = new_t;
if (t == ray->t) { if (t == ray->tmax) {
/* Update throughput in case we haven't done it above */ /* Update throughput in case we haven't done it above */
tp = *throughput * exp(sum); tp = *throughput * exp(sum);
break; break;
@@ -257,15 +260,16 @@ ccl_device float volume_equiangular_sample(ccl_private const Ray *ccl_restrict r
const float xi, const float xi,
ccl_private float *pdf) 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 delta = dot((light_P - ray->P), ray->D);
const float D = safe_sqrtf(len_squared(light_P - ray->P) - delta * delta); const float D = safe_sqrtf(len_squared(light_P - ray->P) - delta * delta);
if (UNLIKELY(D == 0.0f)) { if (UNLIKELY(D == 0.0f)) {
*pdf = 0.0f; *pdf = 0.0f;
return 0.0f; return 0.0f;
} }
const float theta_a = -atan2f(delta, D); const float theta_a = atan2f(tmin - delta, D);
const float theta_b = atan2f(t - delta, D); const float theta_b = atan2f(tmax - delta, D);
const float t_ = D * tanf((xi * theta_b) + (1 - xi) * theta_a); const float t_ = D * tanf((xi * theta_b) + (1 - xi) * theta_a);
if (UNLIKELY(theta_b == theta_a)) { if (UNLIKELY(theta_b == theta_a)) {
*pdf = 0.0f; *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_)); *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, 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; 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 t_ = sample_t - delta;
const float theta_a = -atan2f(delta, D); const float theta_a = atan2f(tmin - delta, D);
const float theta_b = atan2f(t - delta, D); const float theta_b = atan2f(tmax - delta, D);
if (UNLIKELY(theta_b == theta_a)) { if (UNLIKELY(theta_b == theta_a)) {
return 0.0f; return 0.0f;
} }
@@ -310,11 +315,12 @@ ccl_device float volume_equiangular_cdf(ccl_private const Ray *ccl_restrict ray,
return 0.0f; 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 t_ = sample_t - delta;
const float theta_a = -atan2f(delta, D); const float theta_a = atan2f(tmin - delta, D);
const float theta_b = atan2f(t - delta, D); const float theta_b = atan2f(tmax - delta, D);
if (UNLIKELY(theta_b == theta_a)) { if (UNLIKELY(theta_b == theta_a)) {
return 0.0f; return 0.0f;
} }
@@ -390,8 +396,8 @@ ccl_device float3 volume_emission_integrate(ccl_private VolumeShaderCoefficients
typedef struct VolumeIntegrateState { typedef struct VolumeIntegrateState {
/* Volume segment extents. */ /* Volume segment extents. */
float start_t; float tmin;
float end_t; float tmax;
/* If volume is absorption-only up to this point, and no probabilistic /* If volume is absorption-only up to this point, and no probabilistic
* scattering or termination has been used yet. */ * scattering or termination has been used yet. */
@@ -426,9 +432,9 @@ ccl_device_forceinline void volume_integrate_step_scattering(
/* Equiangular sampling for direct lighting. */ /* Equiangular sampling for direct lighting. */
if (vstate.direct_sample_method == VOLUME_SAMPLE_EQUIANGULAR && !result.direct_scatter) { 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) { 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); const float3 new_transmittance = volume_color_transmittance(coeff.sigma_t, new_dt);
result.direct_scatter = true; result.direct_scatter = true;
@@ -458,7 +464,7 @@ ccl_device_forceinline void volume_integrate_step_scattering(
/* compute sampling distance */ /* compute sampling distance */
const float sample_sigma_t = volume_channel_get(coeff.sigma_t, channel); 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_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 */ /* transmittance and pdf */
const float3 new_transmittance = volume_color_transmittance(coeff.sigma_t, new_dt); 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, volume_step_init(kg,
rng_state, rng_state,
object_step_size, object_step_size,
ray->t, ray->tmin,
ray->tmax,
&step_size, &step_size,
&step_shade_offset, &step_shade_offset,
&steps_offset, &steps_offset,
@@ -536,8 +543,8 @@ ccl_device_forceinline void volume_integrate_heterogeneous(
/* Initialize volume integration state. */ /* Initialize volume integration state. */
VolumeIntegrateState vstate ccl_optional_struct_init; VolumeIntegrateState vstate ccl_optional_struct_init;
vstate.start_t = 0.0f; vstate.tmin = ray->tmin;
vstate.end_t = 0.0f; vstate.tmax = ray->tmin;
vstate.absorption_only = true; vstate.absorption_only = true;
vstate.rscatter = path_state_rng_1D(kg, rng_state, PRNG_SCATTER_DISTANCE); vstate.rscatter = path_state_rng_1D(kg, rng_state, PRNG_SCATTER_DISTANCE);
vstate.rphase = path_state_rng_1D(kg, rng_state, PRNG_PHASE_CHANNEL); 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++) { for (int i = 0; i < max_steps; i++) {
/* Advance to new position */ /* Advance to new position */
vstate.end_t = min(ray->t, (i + steps_offset) * step_size); vstate.tmax = min(ray->tmax, ray->tmin + (i + steps_offset) * step_size);
const float shade_t = vstate.start_t + (vstate.end_t - vstate.start_t) * step_shade_offset; const float shade_t = vstate.tmin + (vstate.tmax - vstate.tmin) * step_shade_offset;
sd->P = ray->P + ray->D * shade_t; sd->P = ray->P + ray->D * shade_t;
/* compute segment */ /* compute segment */
@@ -588,7 +595,7 @@ ccl_device_forceinline void volume_integrate_heterogeneous(
const int closure_flag = sd->flag; const int closure_flag = sd->flag;
/* Evaluate transmittance over segment. */ /* 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) ? const float3 transmittance = (closure_flag & SD_EXTINCTION) ?
volume_color_transmittance(coeff.sigma_t, dt) : volume_color_transmittance(coeff.sigma_t, dt) :
one_float3(); one_float3();
@@ -645,8 +652,8 @@ ccl_device_forceinline void volume_integrate_heterogeneous(
} }
/* Stop if at the end of the volume. */ /* Stop if at the end of the volume. */
vstate.start_t = vstate.end_t; vstate.tmin = vstate.tmax;
if (vstate.start_t == ray->t) { if (vstate.tmin == ray->tmax) {
break; break;
} }
} }
@@ -880,7 +887,8 @@ ccl_device_forceinline bool integrate_volume_phase_scatter(
/* Setup ray. */ /* Setup ray. */
INTEGRATOR_STATE_WRITE(state, ray, P) = sd->P; INTEGRATOR_STATE_WRITE(state, ray, P) = sd->P;
INTEGRATOR_STATE_WRITE(state, ray, D) = normalize(phase_omega_in); 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__ # ifdef __RAY_DIFFERENTIALS__
INTEGRATOR_STATE_WRITE(state, ray, dP) = differential_make_compact(sd->dP); INTEGRATOR_STATE_WRITE(state, ray, dP) = differential_make_compact(sd->dP);
INTEGRATOR_STATE_WRITE(state, ray, dD) = differential_make_compact(phase_domega_in); 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 */ /* Update path state */
INTEGRATOR_STATE_WRITE(state, path, mis_ray_pdf) = phase_pdf; 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( INTEGRATOR_STATE_WRITE(state, path, min_ray_pdf) = fminf(
phase_pdf, INTEGRATOR_STATE(state, path, min_ray_pdf)); 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); integrator_state_read_isect(kg, state, &isect);
/* Set ray length to current segment. */ /* 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. */ /* Clean volume stack for background rays. */
if (isect.prim == PRIM_NONE) { if (isect.prim == PRIM_NONE) {

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

@@ -47,7 +47,8 @@ KERNEL_STRUCT_END(shadow_path)
KERNEL_STRUCT_BEGIN(shadow_ray) KERNEL_STRUCT_BEGIN(shadow_ray)
KERNEL_STRUCT_MEMBER(shadow_ray, packed_float3, P, KERNEL_FEATURE_PATH_TRACING) 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, 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, time, KERNEL_FEATURE_PATH_TRACING)
KERNEL_STRUCT_MEMBER(shadow_ray, float, dP, 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) KERNEL_STRUCT_MEMBER(shadow_ray, int, object, KERNEL_FEATURE_PATH_TRACING)

10
intern/cycles/kernel/integrator/state_template.h
View File

@@ -37,11 +37,10 @@ KERNEL_STRUCT_MEMBER(path, uint32_t, flag, KERNEL_FEATURE_PATH_TRACING)
/* enum PathRayMNEE */ /* enum PathRayMNEE */
KERNEL_STRUCT_MEMBER(path, uint8_t, mnee, KERNEL_FEATURE_PATH_TRACING) KERNEL_STRUCT_MEMBER(path, uint8_t, mnee, KERNEL_FEATURE_PATH_TRACING)
/* Multiple importance sampling /* Multiple importance sampling
* The PDF of BSDF sampling at the last scatter point, and distance to the * The PDF of BSDF sampling at the last scatter point, which is at ray distance
* last scatter point minus the last ray segment. This distance lets us * zero and distance. Note that transparency and volume attenuation increase
* compute the complete distance through transparent surfaces and volumes. */ * 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_pdf, KERNEL_FEATURE_PATH_TRACING)
KERNEL_STRUCT_MEMBER(path, float, mis_ray_t, KERNEL_FEATURE_PATH_TRACING)
/* Filter glossy. */ /* Filter glossy. */
KERNEL_STRUCT_MEMBER(path, float, min_ray_pdf, KERNEL_FEATURE_PATH_TRACING) KERNEL_STRUCT_MEMBER(path, float, min_ray_pdf, KERNEL_FEATURE_PATH_TRACING)
/* Continuation probability for path termination. */ /* Continuation probability for path termination. */
@@ -63,7 +62,8 @@ KERNEL_STRUCT_END(path)
KERNEL_STRUCT_BEGIN(ray) KERNEL_STRUCT_BEGIN(ray)
KERNEL_STRUCT_MEMBER(ray, packed_float3, P, KERNEL_FEATURE_PATH_TRACING) 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, 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, time, KERNEL_FEATURE_PATH_TRACING)
KERNEL_STRUCT_MEMBER(ray, float, dP, KERNEL_FEATURE_PATH_TRACING) KERNEL_STRUCT_MEMBER(ray, float, dP, KERNEL_FEATURE_PATH_TRACING)
KERNEL_STRUCT_MEMBER(ray, float, dD, KERNEL_FEATURE_PATH_TRACING) KERNEL_STRUCT_MEMBER(ray, float, dD, KERNEL_FEATURE_PATH_TRACING)

12
intern/cycles/kernel/integrator/state_util.h
View File

@@ -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, P) = ray->P;
INTEGRATOR_STATE_WRITE(state, ray, D) = ray->D; 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, time) = ray->time;
INTEGRATOR_STATE_WRITE(state, ray, dP) = ray->dP; INTEGRATOR_STATE_WRITE(state, ray, dP) = ray->dP;
INTEGRATOR_STATE_WRITE(state, ray, dD) = ray->dD; 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->P = INTEGRATOR_STATE(state, ray, P);
ray->D = INTEGRATOR_STATE(state, ray, D); 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->time = INTEGRATOR_STATE(state, ray, time);
ray->dP = INTEGRATOR_STATE(state, ray, dP); ray->dP = INTEGRATOR_STATE(state, ray, dP);
ray->dD = INTEGRATOR_STATE(state, ray, dD); 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, P) = ray->P;
INTEGRATOR_STATE_WRITE(state, shadow_ray, D) = ray->D; 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, time) = ray->time;
INTEGRATOR_STATE_WRITE(state, shadow_ray, dP) = ray->dP; 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->P = INTEGRATOR_STATE(state, shadow_ray, P);
ray->D = INTEGRATOR_STATE(state, shadow_ray, D); 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->time = INTEGRATOR_STATE(state, shadow_ray, time);
ray->dP = INTEGRATOR_STATE(state, shadow_ray, dP); ray->dP = INTEGRATOR_STATE(state, shadow_ray, dP);
ray->dD = differential_zero_compact(); ray->dD = differential_zero_compact();

5
intern/cycles/kernel/integrator/subsurface.h
View File

@@ -38,7 +38,8 @@ ccl_device int subsurface_bounce(KernelGlobals kg,
/* Setup ray into surface. */ /* Setup ray into surface. */
INTEGRATOR_STATE_WRITE(state, ray, P) = sd->P; INTEGRATOR_STATE_WRITE(state, ray, P) = sd->P;
INTEGRATOR_STATE_WRITE(state, ray, D) = bssrdf->N; 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, dP) = differential_make_compact(sd->dP);
INTEGRATOR_STATE_WRITE(state, ray, dD) = differential_zero_compact(); 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 /* Pretend ray is coming from the outside towards the exit point. This ensures
* correct front/back facing normals. * correct front/back facing normals.
* TODO: find a more elegant solution? */ * 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; ray.D = -ray.D;
integrator_state_write_isect(kg, state, &ss_isect.hits[0]); integrator_state_write_isect(kg, state, &ss_isect.hits[0]);

6
intern/cycles/kernel/integrator/subsurface_disk.h
View File

@@ -82,7 +82,8 @@ ccl_device_inline bool subsurface_disk(KernelGlobals kg,
/* Create ray. */ /* Create ray. */
ray.P = P + disk_N * disk_height + disk_P; ray.P = P + disk_N * disk_height + disk_P;
ray.D = -disk_N; 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.dP = ray_dP;
ray.dD = differential_zero_compact(); ray.dD = differential_zero_compact();
ray.time = time; 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.P = ray.P + ray.D * ss_isect.hits[hit].t;
ray.D = ss_isect.Ng[hit]; ray.D = ss_isect.Ng[hit];
ray.t = 1.0f; ray.tmin = 0.0f;
ray.tmax = 1.0f;
return true; return true;
} }

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

@@ -195,7 +195,8 @@ ccl_device_inline bool subsurface_random_walk(KernelGlobals kg,
/* Setup ray. */ /* Setup ray. */
ray.P = P; ray.P = P;
ray.D = D; ray.D = D;
ray.t = FLT_MAX; ray.tmin = 0.0f;
ray.tmax = FLT_MAX;
ray.time = time; ray.time = time;
ray.dP = ray_dP; ray.dP = ray_dP;
ray.dD = differential_zero_compact(); ray.dD = differential_zero_compact();
@@ -370,10 +371,10 @@ ccl_device_inline bool subsurface_random_walk(KernelGlobals kg,
* chance of connecting to it. * chance of connecting to it.
* TODO: Maybe use less than 10 times the mean free path? */ * TODO: Maybe use less than 10 times the mean free path? */
if (bounce == 0) { if (bounce == 0) {
ray.t = max(t, 10.0f / (reduce_min(sigma_t))); ray.tmax = max(t, 10.0f / (reduce_min(sigma_t)));
} }
else { else {
ray.t = t; ray.tmax = t;
/* After the first bounce the object can intersect the same surface again */ /* After the first bounce the object can intersect the same surface again */
ray.self.object = OBJECT_NONE; ray.self.object = OBJECT_NONE;
ray.self.prim = PRIM_NONE; ray.self.prim = PRIM_NONE;
@@ -384,12 +385,12 @@ ccl_device_inline bool subsurface_random_walk(KernelGlobals kg,
if (hit) { if (hit) {
#ifdef __KERNEL_GPU_RAYTRACING__ #ifdef __KERNEL_GPU_RAYTRACING__
/* t is always in world space with OptiX and MetalRT. */ /* 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 #else
/* Compute world space distance to surface hit. */ /* Compute world space distance to surface hit. */
float3 D = transform_direction(&ob_itfm, ray.D); float3 D = transform_direction(&ob_itfm, ray.D);
D = normalize(D) * ss_isect.hits[0].t; 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 #endif
} }
@@ -397,16 +398,16 @@ ccl_device_inline bool subsurface_random_walk(KernelGlobals kg,
/* Check if we hit the opposite side. */ /* Check if we hit the opposite side. */
if (hit) { if (hit) {
have_opposite_interface = true; 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, /* 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. */ * 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. */ /* Use the distance to the exit point for the throughput update if we found one. */
if (hit) { if (hit) {
t = ray.t; t = ray.tmax;
} }
/* Advance to new scatter location. */ /* Advance to new scatter location. */

34
intern/cycles/kernel/light/light.h
View File

@@ -270,31 +270,26 @@ ccl_device bool lights_intersect(KernelGlobals kg,
if (type == LIGHT_SPOT) { if (type == LIGHT_SPOT) {
/* Spot/Disk light. */ /* 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 float3 lightP = make_float3(klight->co[0], klight->co[1], klight->co[2]);
const float radius = klight->spot.radius; const float radius = klight->spot.radius;
if (radius == 0.0f) { if (radius == 0.0f) {
continue; continue;
} }
/* disk oriented normal */ /* disk oriented normal */
const float3 lightN = normalize(ray_P - lightP); const float3 lightN = normalize(ray->P - lightP);
/* One sided. */ /* One sided. */
if (dot(ray->D, lightN) >= 0.0f) { if (dot(ray->D, lightN) >= 0.0f) {
continue; continue;
} }
float3 P; 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; continue;
} }
} }
else if (type == LIGHT_POINT) { else if (type == LIGHT_POINT) {
/* Sphere light (aka, aligned disk light). */ /* 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 float3 lightP = make_float3(klight->co[0], klight->co[1], klight->co[2]);
const float radius = klight->spot.radius; const float radius = klight->spot.radius;
if (radius == 0.0f) { if (radius == 0.0f) {
@@ -302,9 +297,10 @@ ccl_device bool lights_intersect(KernelGlobals kg,
} }
/* disk oriented normal */ /* disk oriented normal */
const float3 lightN = normalize(ray_P - lightP); const float3 lightN = normalize(ray->P - lightP);
float3 P; 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; 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]); const float3 light_P = make_float3(klight->co[0], klight->co[1], klight->co[2]);
float3 P; float3 P;
if (!ray_quad_intersect( if (!ray_quad_intersect(ray->P,
ray->P, ray->D, 0.0f, ray->t, light_P, axisu, axisv, Ng, &P, &t, &u, &v, is_round)) { ray->D,
ray->tmin,
ray->tmax,
light_P,
axisu,
axisv,
Ng,
&P,
&t,
&u,
&v,
is_round)) {
continue; 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); ls->D = z * B + safe_sqrtf(1.0f - z * z) * safe_normalize(C_ - dot(C_, B) * B);
/* calculate intersection with the planar triangle */ /* 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; ls->pdf = 0.0f;
return; return;
} }

7
intern/cycles/kernel/light/sample.h
View File

@@ -227,23 +227,24 @@ ccl_device_inline void shadow_ray_setup(ccl_private const ShaderData *ccl_restri
if (ls->shader & SHADER_CAST_SHADOW) { if (ls->shader & SHADER_CAST_SHADOW) {
/* setup ray */ /* setup ray */
ray->P = P; ray->P = P;
ray->tmin = 0.0f;
if (ls->t == FLT_MAX) { if (ls->t == FLT_MAX) {
/* distant light */ /* distant light */
ray->D = ls->D; ray->D = ls->D;
ray->t = ls->t; ray->tmax = ls->t;
} }
else { else {
/* other lights, avoid self-intersection */ /* other lights, avoid self-intersection */
ray->D = ls->P - P; ray->D = ls->P - P;
ray->D = normalize_len(ray->D, &ray->t); ray->D = normalize_len(ray->D, &ray->tmax);
} }
} }
else { else {
/* signal to not cast shadow ray */ /* signal to not cast shadow ray */
ray->P = zero_float3(); ray->P = zero_float3();
ray->D = zero_float3(); ray->D = zero_float3();
ray->t = 0.0f; ray->tmax = 0.0f;
} }
ray->dP = differential_make_compact(sd->dP); ray->dP = differential_make_compact(sd->dP);

9
intern/cycles/kernel/osl/services.cpp
View File

@@ -1094,10 +1094,8 @@ bool OSLRenderServices::get_background_attribute(const KernelGlobalsCPU *kg,
ndc[0] = camera_world_to_ndc(kg, sd, sd->ray_P); ndc[0] = camera_world_to_ndc(kg, sd, sd->ray_P);
if (derivatives) { if (derivatives) {
ndc[1] = camera_world_to_ndc(kg, sd, sd->ray_P + make_float3(sd->ray_dP, 0.0f, 0.0f)) - ndc[1] = zero_float3();
ndc[0]; ndc[2] = zero_float3();
ndc[2] = camera_world_to_ndc(kg, sd, sd->ray_P + make_float3(0.0f, sd->ray_dP, 0.0f)) -
ndc[0];
} }
} }
else { else {
@@ -1671,7 +1669,8 @@ bool OSLRenderServices::trace(TraceOpt &options,
ray.P = TO_FLOAT3(P); ray.P = TO_FLOAT3(P);
ray.D = TO_FLOAT3(R); 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.time = sd->time;
ray.self.object = OBJECT_NONE; ray.self.object = OBJECT_NONE;
ray.self.prim = PRIM_NONE; ray.self.prim = PRIM_NONE;

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

@@ -59,7 +59,8 @@ ccl_device float svm_ao(
Ray ray; Ray ray;
ray.P = sd->P; ray.P = sd->P;
ray.D = D.x * T + D.y * B + D.z * N; 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.time = sd->time;
ray.self.object = sd->object; ray.self.object = sd->object;
ray.self.prim = sd->prim; ray.self.prim = sd->prim;

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

@@ -179,7 +179,8 @@ ccl_device float3 svm_bevel(
Ray ray ccl_optional_struct_init; Ray ray ccl_optional_struct_init;
ray.P = sd->P + disk_N * disk_height + disk_P; ray.P = sd->P + disk_N * disk_height + disk_P;
ray.D = -disk_N; 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.dP = differential_zero_compact();
ray.dD = differential_zero_compact(); ray.dD = differential_zero_compact();
ray.time = sd->time; ray.time = sd->time;

4
intern/cycles/kernel/svm/tex_coord.h
View File

@@ -138,7 +138,7 @@ ccl_device_noinline int svm_node_tex_coord_bump_dx(KernelGlobals kg,
case NODE_TEXCO_WINDOW: { case NODE_TEXCO_WINDOW: {
if ((path_flag & PATH_RAY_CAMERA) && sd->object == OBJECT_NONE && if ((path_flag & PATH_RAY_CAMERA) && sd->object == OBJECT_NONE &&
kernel_data.cam.type == CAMERA_ORTHOGRAPHIC) 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 else
data = camera_world_to_ndc(kg, sd, sd->P + sd->dP.dx); data = camera_world_to_ndc(kg, sd, sd->P + sd->dP.dx);
data.z = 0.0f; data.z = 0.0f;
@@ -223,7 +223,7 @@ ccl_device_noinline int svm_node_tex_coord_bump_dy(KernelGlobals kg,
case NODE_TEXCO_WINDOW: { case NODE_TEXCO_WINDOW: {
if ((path_flag & PATH_RAY_CAMERA) && sd->object == OBJECT_NONE && if ((path_flag & PATH_RAY_CAMERA) && sd->object == OBJECT_NONE &&
kernel_data.cam.type == CAMERA_ORTHOGRAPHIC) 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 else
data = camera_world_to_ndc(kg, sd, sd->P + sd->dP.dy); data = camera_world_to_ndc(kg, sd, sd->P + sd->dP.dy);
data.z = 0.0f; data.z = 0.0f;

3
intern/cycles/kernel/types.h
View File

@@ -535,7 +535,8 @@ typedef struct RaySelfPrimitives {
typedef struct Ray { typedef struct Ray {
float3 P; /* origin */ float3 P; /* origin */
float3 D; /* direction */ float3 D; /* direction */
float t; /* length of the ray */ float tmin; /* start distance */
float tmax; /* end distance */
float time; /* time (for motion blur) */ float time; /* time (for motion blur) */
RaySelfPrimitives self; RaySelfPrimitives self;

37
intern/cycles/util/math_intersect.h
View File

@@ -10,7 +10,8 @@ CCL_NAMESPACE_BEGIN
ccl_device bool ray_sphere_intersect(float3 ray_P, ccl_device bool ray_sphere_intersect(float3 ray_P,
float3 ray_D, float3 ray_D,
float ray_t, float ray_tmin,
float ray_tmax,
float3 sphere_P, float3 sphere_P,
float sphere_radius, float sphere_radius,
ccl_private float3 *isect_P, ccl_private float3 *isect_P,
@@ -33,7 +34,7 @@ ccl_device bool ray_sphere_intersect(float3 ray_P,
return false; return false;
} }
const float t = tp - sqrtf(radiussq - dsq); /* pythagoras */ const float t = tp - sqrtf(radiussq - dsq); /* pythagoras */
if (t < ray_t) { if (t > ray_tmin && t < ray_tmax) {
*isect_t = t; *isect_t = t;
*isect_P = ray_P + ray_D * t; *isect_P = ray_P + ray_D * t;
return true; return true;
@@ -44,7 +45,8 @@ ccl_device bool ray_sphere_intersect(float3 ray_P,
ccl_device bool ray_aligned_disk_intersect(float3 ray_P, ccl_device bool ray_aligned_disk_intersect(float3 ray_P,
float3 ray_D, float3 ray_D,
float ray_t, float ray_tmin,
float ray_tmax,
float3 disk_P, float3 disk_P,
float disk_radius, float disk_radius,
ccl_private float3 *isect_P, ccl_private float3 *isect_P,
@@ -59,7 +61,7 @@ ccl_device bool ray_aligned_disk_intersect(float3 ray_P,
} }
/* Compute t to intersection point. */ /* Compute t to intersection point. */
const float t = -disk_t / div; const float t = -disk_t / div;
if (t < 0.0f || t > ray_t) { if (!(t > ray_tmin && t < ray_tmax)) {
return false; return false;
} }
/* Test if within radius. */ /* Test if within radius. */
@@ -74,7 +76,8 @@ ccl_device bool ray_aligned_disk_intersect(float3 ray_P,
ccl_device bool ray_disk_intersect(float3 ray_P, ccl_device bool ray_disk_intersect(float3 ray_P,
float3 ray_D, float3 ray_D,
float ray_t, float ray_tmin,
float ray_tmax,
float3 disk_P, float3 disk_P,
float3 disk_N, float3 disk_N,
float disk_radius, float disk_radius,
@@ -92,7 +95,8 @@ ccl_device bool ray_disk_intersect(float3 ray_P,
} }
float3 P = ray_P + t * ray_D; float3 P = ray_P + t * ray_D;
float3 T = P - disk_P; float3 T = P - disk_P;
if (dot(T, T) < sqr(disk_radius) /*&& t > 0.f*/ && t <= ray_t) {
if (dot(T, T) < sqr(disk_radius) && (t > ray_tmin && t < ray_tmax)) {
*isect_P = ray_P + t * ray_D; *isect_P = ray_P + t * ray_D;
*isect_t = t; *isect_t = t;
return true; return true;
@@ -103,7 +107,8 @@ ccl_device bool ray_disk_intersect(float3 ray_P,
ccl_device_forceinline bool ray_triangle_intersect(float3 ray_P, ccl_device_forceinline bool ray_triangle_intersect(float3 ray_P,
float3 ray_dir, float3 ray_dir,
float ray_t, float ray_tmin,
float ray_tmax,
const float3 tri_a, const float3 tri_a,
const float3 tri_b, const float3 tri_b,
const float3 tri_c, const float3 tri_c,
@@ -149,16 +154,14 @@ ccl_device_forceinline bool ray_triangle_intersect(float3 ray_P,
/* Perform depth test. */ /* Perform depth test. */
const float T = dot3(v0, Ng); const float T = dot3(v0, Ng);
const int sign_den = (__float_as_int(den) & 0x80000000); const float t = T / den;
const float sign_T = xor_signmask(T, sign_den); if (!(t >= ray_tmin && t <= ray_tmax)) {
if ((sign_T < 0.0f) || (sign_T > ray_t * xor_signmask(den, sign_den))) {
return false; return false;
} }
const float inv_den = 1.0f / den; *isect_u = U / den;
*isect_u = U * inv_den; *isect_v = V / den;
*isect_v = V * inv_den; *isect_t = t;
*isect_t = T * inv_den;
return true; return true;
#undef dot3 #undef dot3
@@ -171,8 +174,8 @@ ccl_device_forceinline bool ray_triangle_intersect(float3 ray_P,
*/ */
ccl_device bool ray_quad_intersect(float3 ray_P, ccl_device bool ray_quad_intersect(float3 ray_P,
float3 ray_D, float3 ray_D,
float ray_mint, float ray_tmin,
float ray_maxt, float ray_tmax,
float3 quad_P, float3 quad_P,
float3 quad_u, float3 quad_u,
float3 quad_v, float3 quad_v,
@@ -185,7 +188,7 @@ ccl_device bool ray_quad_intersect(float3 ray_P,
{ {
/* Perform intersection test. */ /* Perform intersection test. */
float t = -(dot(ray_P, quad_n) - dot(quad_P, quad_n)) / dot(ray_D, quad_n); float t = -(dot(ray_P, quad_n) - dot(quad_P, quad_n)) / dot(ray_D, quad_n);
if (t < ray_mint || t > ray_maxt) { if (!(t > ray_tmin && t < ray_tmax)) {
return false; return false;
} }
const float3 hit = ray_P + t * ray_D; const float3 hit = ray_P + t * ray_D;