archive/blender-archive
Archived
1
1
Fork 0
Code Pull Requests Packages Activity

Merge branch 'master' into retopo_transform

Browse Source
This commit is contained in:
jon denning
2022-07-16 06:53:04 -04:00
parent a2938c86ca 0a8d21e0c9
commit cc761cdae6
52 changed files with 527 additions and 362 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,
ray->P,
ray->D,
0.0f,
ray->t,
ray->tmin,
ray->tmax,
ray->time,
ray_mask,
ray_flags,
@@ -203,28 +203,28 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg,
#elif defined(__METALRT__)
if (!scene_intersect_valid(ray)) {
isect->t = ray->t;
isect->t = ray->tmax;
isect->type = PRIMITIVE_NONE;
return false;
}
# if defined(__KERNEL_DEBUG__)
if (is_null_instance_acceleration_structure(metal_ancillaries->accel_struct)) {
isect->t = ray->t;
isect->t = ray->tmax;
isect->type = PRIMITIVE_NONE;
kernel_assert(!"Invalid metal_ancillaries->accel_struct pointer");
return false;
}
if (is_null_intersection_function_table(metal_ancillaries->ift_default)) {
isect->t = ray->t;
isect->t = ray->tmax;
isect->type = PRIMITIVE_NONE;
kernel_assert(!"Invalid ift_default");
return false;
}
# endif
metal::raytracing::ray r(ray->P, ray->D, 0.0f, ray->t);
metal::raytracing::ray r(ray->P, ray->D, ray->tmin, ray->tmax);
metalrt_intersector_type metalrt_intersect;
if (!kernel_data.bvh.have_curves) {
@@ -263,7 +263,7 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg,
# endif
if (intersection.type == intersection_type::none) {
isect->t = ray->t;
isect->t = ray->tmax;
isect->type = PRIMITIVE_NONE;
return false;
@@ -296,7 +296,7 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg,
# ifdef __EMBREE__
if (kernel_data.device_bvh) {
isect->t = ray->t;
isect->t = ray->tmax;
CCLIntersectContext ctx(kg, CCLIntersectContext::RAY_REGULAR);
IntersectContext rtc_ctx(&ctx);
RTCRayHit ray_hit;
@@ -360,8 +360,8 @@ ccl_device_intersect bool scene_intersect_local(KernelGlobals kg,
optixTrace(scene_intersect_valid(ray) ? kernel_data.device_bvh : 0,
ray->P,
ray->D,
0.0f,
ray->t,
ray->tmin,
ray->tmax,
ray->time,
0xFF,
/* Need to always call into __anyhit__kernel_optix_local_hit. */
@@ -405,7 +405,7 @@ ccl_device_intersect bool scene_intersect_local(KernelGlobals kg,
}
# endif
metal::raytracing::ray r(ray->P, ray->D, 0.0f, ray->t);
metal::raytracing::ray r(ray->P, ray->D, ray->tmin, ray->tmax);
metalrt_intersector_type metalrt_intersect;
metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque);
@@ -476,7 +476,7 @@ ccl_device_intersect bool scene_intersect_local(KernelGlobals kg,
float3 dir = ray->D;
float3 idir = ray->D;
Transform ob_itfm;
rtc_ray.tfar = ray->t *
rtc_ray.tfar = ray->tmax *
bvh_instance_motion_push(kg, local_object, ray, &P, &dir, &idir, &ob_itfm);
/* bvh_instance_motion_push() returns the inverse transform but
* it's not needed here. */
@@ -542,8 +542,8 @@ ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg,
optixTrace(scene_intersect_valid(ray) ? kernel_data.device_bvh : 0,
ray->P,
ray->D,
0.0f,
ray->t,
ray->tmin,
ray->tmax,
ray->time,
ray_mask,
/* Need to always call into __anyhit__kernel_optix_shadow_all_hit. */
@@ -582,7 +582,7 @@ ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg,
}
# endif
metal::raytracing::ray r(ray->P, ray->D, 0.0f, ray->t);
metal::raytracing::ray r(ray->P, ray->D, ray->tmin, ray->tmax);
metalrt_intersector_type metalrt_intersect;
metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque);
@@ -701,8 +701,8 @@ ccl_device_intersect bool scene_intersect_volume(KernelGlobals kg,
optixTrace(scene_intersect_valid(ray) ? kernel_data.device_bvh : 0,
ray->P,
ray->D,
0.0f,
ray->t,
ray->tmin,
ray->tmax,
ray->time,
ray_mask,
/* Need to always call into __anyhit__kernel_optix_volume_test. */
@@ -744,7 +744,7 @@ ccl_device_intersect bool scene_intersect_volume(KernelGlobals kg,
}
# endif
metal::raytracing::ray r(ray->P, ray->D, 0.0f, ray->t);
metal::raytracing::ray r(ray->P, ray->D, ray->tmin, ray->tmax);
metalrt_intersector_type metalrt_intersect;
metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque);

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

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

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

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

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

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

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

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

@@ -5,6 +5,19 @@
CCL_NAMESPACE_BEGIN
/* Offset intersection distance by the smallest possible amount, to skip
* intersections at this distance. This works in cases where the ray start
* position is unchanged and only tmin is updated, since for self
* intersection we'll be comparing against the exact same distances. */
ccl_device_forceinline float intersection_t_offset(const float t)
{
/* This is a simplified version of nextafterf(t, FLT_MAX), only dealing with
* non-negative and finite t. */
kernel_assert(t >= 0.0f && isfinite_safe(t));
const uint32_t bits = (t == 0.0f) ? 1 : __float_as_uint(t) + 1;
return __uint_as_float(bits);
}
#if defined(__KERNEL_CPU__)
ccl_device int intersections_compare(const void *a, const void *b)
{

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

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

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

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

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

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_direction,
float time,
const float ray_tmin,
const float ray_tmax,
thread BoundingBoxIntersectionResult &result)
{
@@ -434,7 +435,7 @@ void metalrt_intersection_curve(constant KernelParamsMetal &launch_params_metal,
isect.t *= len;
MetalKernelContext context(launch_params_metal);
if (context.curve_intersect(NULL, &isect, P, dir, isect.t, object, prim, time, type)) {
if (context.curve_intersect(NULL, &isect, P, dir, ray_tmin, isect.t, object, prim, time, type)) {
result = metalrt_visibility_test<BoundingBoxIntersectionResult, METALRT_HIT_BOUNDING_BOX>(
launch_params_metal, payload, object, prim, isect.u);
if (result.accept) {
@@ -456,6 +457,7 @@ void metalrt_intersection_curve_shadow(constant KernelParamsMetal &launch_params
const float3 ray_origin,
const float3 ray_direction,
float time,
const float ray_tmin,
const float ray_tmax,
thread BoundingBoxIntersectionResult &result)
{
@@ -475,7 +477,7 @@ void metalrt_intersection_curve_shadow(constant KernelParamsMetal &launch_params
isect.t *= len;
MetalKernelContext context(launch_params_metal);
if (context.curve_intersect(NULL, &isect, P, dir, isect.t, object, prim, time, type)) {
if (context.curve_intersect(NULL, &isect, P, dir, ray_tmin, isect.t, object, prim, time, type)) {
result.continue_search = metalrt_shadow_all_hit<METALRT_HIT_BOUNDING_BOX>(
launch_params_metal, payload, object, prim, float2(isect.u, isect.v), ray_tmax);
result.accept = !result.continue_search;
@@ -494,6 +496,7 @@ __intersection__curve_ribbon(constant KernelParamsMetal &launch_params_metal [[b
const uint primitive_id [[primitive_id]],
const float3 ray_origin [[origin]],
const float3 ray_direction [[direction]],
const float ray_tmin [[min_distance]],
const float ray_tmax [[max_distance]])
{
uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object);
@@ -511,7 +514,7 @@ __intersection__curve_ribbon(constant KernelParamsMetal &launch_params_metal [[b
# else
0.0f,
# endif
ray_tmax, result);
ray_tmin, ray_tmax, result);
}
return result;
@@ -525,6 +528,7 @@ __intersection__curve_ribbon_shadow(constant KernelParamsMetal &launch_params_me
const uint primitive_id [[primitive_id]],
const float3 ray_origin [[origin]],
const float3 ray_direction [[direction]],
const float ray_tmin [[min_distance]],
const float ray_tmax [[max_distance]])
{
uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object);
@@ -542,7 +546,7 @@ __intersection__curve_ribbon_shadow(constant KernelParamsMetal &launch_params_me
# else
0.0f,
# endif
ray_tmax, result);
ray_tmin, ray_tmax, result);
}
return result;
@@ -556,6 +560,7 @@ __intersection__curve_all(constant KernelParamsMetal &launch_params_metal [[buff
const uint primitive_id [[primitive_id]],
const float3 ray_origin [[origin]],
const float3 ray_direction [[direction]],
const float ray_tmin [[min_distance]],
const float ray_tmax [[max_distance]])
{
uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object);
@@ -571,7 +576,7 @@ __intersection__curve_all(constant KernelParamsMetal &launch_params_metal [[buff
# else
0.0f,
# endif
ray_tmax, result);
ray_tmin, ray_tmax, result);
return result;
}
@@ -584,6 +589,7 @@ __intersection__curve_all_shadow(constant KernelParamsMetal &launch_params_metal
const uint primitive_id [[primitive_id]],
const float3 ray_origin [[origin]],
const float3 ray_direction [[direction]],
const float ray_tmin [[min_distance]],
const float ray_tmax [[max_distance]])
{
uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object);
@@ -600,7 +606,7 @@ __intersection__curve_all_shadow(constant KernelParamsMetal &launch_params_metal
# else
0.0f,
# endif
ray_tmax, result);
ray_tmin, ray_tmax, result);
return result;
}
@@ -616,6 +622,7 @@ void metalrt_intersection_point(constant KernelParamsMetal &launch_params_metal,
const float3 ray_origin,
const float3 ray_direction,
float time,
const float ray_tmin,
const float ray_tmax,
thread BoundingBoxIntersectionResult &result)
{
@@ -640,7 +647,7 @@ void metalrt_intersection_point(constant KernelParamsMetal &launch_params_metal,
isect.t *= len;
MetalKernelContext context(launch_params_metal);
if (context.point_intersect(NULL, &isect, P, dir, isect.t, object, prim, time, type)) {
if (context.point_intersect(NULL, &isect, P, dir, ray_tmin, isect.t, object, prim, time, type)) {
result = metalrt_visibility_test<BoundingBoxIntersectionResult, METALRT_HIT_BOUNDING_BOX>(
launch_params_metal, payload, object, prim, isect.u);
if (result.accept) {
@@ -662,6 +669,7 @@ void metalrt_intersection_point_shadow(constant KernelParamsMetal &launch_params
const float3 ray_origin,
const float3 ray_direction,
float time,
const float ray_tmin,
const float ray_tmax,
thread BoundingBoxIntersectionResult &result)
{
@@ -681,7 +689,7 @@ void metalrt_intersection_point_shadow(constant KernelParamsMetal &launch_params
isect.t *= len;
MetalKernelContext context(launch_params_metal);
if (context.point_intersect(NULL, &isect, P, dir, isect.t, object, prim, time, type)) {
if (context.point_intersect(NULL, &isect, P, dir, ray_tmin, isect.t, object, prim, time, type)) {
result.continue_search = metalrt_shadow_all_hit<METALRT_HIT_BOUNDING_BOX>(
launch_params_metal, payload, object, prim, float2(isect.u, isect.v), ray_tmax);
result.accept = !result.continue_search;
@@ -700,6 +708,7 @@ __intersection__point(constant KernelParamsMetal &launch_params_metal [[buffer(1
const uint primitive_id [[primitive_id]],
const float3 ray_origin [[origin]],
const float3 ray_direction [[direction]],
const float ray_tmin [[min_distance]],
const float ray_tmax [[max_distance]])
{
const uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object);
@@ -716,7 +725,7 @@ __intersection__point(constant KernelParamsMetal &launch_params_metal [[buffer(1
# else
0.0f,
# endif
ray_tmax, result);
ray_tmin, ray_tmax, result);
return result;
}
@@ -729,6 +738,7 @@ __intersection__point_shadow(constant KernelParamsMetal &launch_params_metal [[b
const uint primitive_id [[primitive_id]],
const float3 ray_origin [[origin]],
const float3 ray_direction [[direction]],
const float ray_tmin [[min_distance]],
const float ray_tmax [[max_distance]])
{
const uint prim = primitive_id + kernel_data_fetch(object_prim_offset, object);
@@ -745,7 +755,7 @@ __intersection__point_shadow(constant KernelParamsMetal &launch_params_metal [[b
# else
0.0f,
# endif
ray_tmax, result);
ray_tmin, ray_tmax, result);
return result;
}

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

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

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,
float3 P,
float3 dir,
float tmin,
float tmax,
float time,
uint visibility,
@@ -58,7 +59,7 @@ ccl_device_inline bool motion_triangle_intersect(KernelGlobals kg,
motion_triangle_vertices(kg, object, prim, time, verts);
/* Ray-triangle intersection, unoptimized. */
float t, u, v;
if (ray_triangle_intersect(P, dir, tmax, verts[0], verts[1], verts[2], &u, &v, &t)) {
if (ray_triangle_intersect(P, dir, tmin, tmax, verts[0], verts[1], verts[2], &u, &v, &t)) {
#ifdef __VISIBILITY_FLAG__
/* Visibility flag test. we do it here under the assumption
* that most triangles are culled by node flags.
@@ -92,6 +93,7 @@ ccl_device_inline bool motion_triangle_intersect_local(KernelGlobals kg,
int object,
int prim,
int prim_addr,
float tmin,
float tmax,
ccl_private uint *lcg_state,
int max_hits)
@@ -101,7 +103,7 @@ ccl_device_inline bool motion_triangle_intersect_local(KernelGlobals kg,
motion_triangle_vertices(kg, object, prim, time, verts);
/* Ray-triangle intersection, unoptimized. */
float t, u, v;
if (!ray_triangle_intersect(P, dir, tmax, verts[0], verts[1], verts[2], &u, &v, &t)) {
if (!ray_triangle_intersect(P, dir, tmin, tmax, verts[0], verts[1], verts[2], &u, &v, &t)) {
return false;
}

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

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

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

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

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

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

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;
volume_ray.P = from_P;
volume_ray.D = normalize_len(to_P - from_P, &volume_ray.t);
volume_ray.D = normalize_len(to_P - from_P, &volume_ray.tmax);
volume_ray.tmin = 0.0f;
volume_ray.self.object = INTEGRATOR_STATE(state, isect, object);
volume_ray.self.prim = INTEGRATOR_STATE(state, isect, prim);
volume_ray.self.light_object = OBJECT_NONE;
@@ -58,12 +59,9 @@ ccl_device void integrator_volume_stack_update_for_subsurface(KernelGlobals kg,
volume_stack_enter_exit(kg, state, stack_sd);
/* Move ray forward. */
volume_ray.P = stack_sd->P;
volume_ray.tmin = intersection_t_offset(isect.t);
volume_ray.self.object = isect.object;
volume_ray.self.prim = isect.prim;
if (volume_ray.t != FLT_MAX) {
volume_ray.D = normalize_len(to_P - volume_ray.P, &volume_ray.t);
}
++step;
}
#endif
@@ -82,7 +80,8 @@ ccl_device void integrator_volume_stack_init(KernelGlobals kg, IntegratorState s
/* Trace ray in random direction. Any direction works, Z up is a guess to get the
* fewest hits. */
volume_ray.D = make_float3(0.0f, 0.0f, 1.0f);
volume_ray.t = FLT_MAX;
volume_ray.tmin = 0.0f;
volume_ray.tmax = FLT_MAX;
volume_ray.self.object = OBJECT_NONE;
volume_ray.self.prim = PRIM_NONE;
volume_ray.self.light_object = OBJECT_NONE;
@@ -199,7 +198,7 @@ ccl_device void integrator_volume_stack_init(KernelGlobals kg, IntegratorState s
}
/* Move ray forward. */
volume_ray.P = stack_sd->P;
volume_ray.tmin = intersection_t_offset(isect.t);
volume_ray.self.object = isect.object;
volume_ray.self.prim = isect.prim;
++step;

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

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

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

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 float ray_time = INTEGRATOR_STATE(state, ray, time);
/* Advance ray beyond light. */
/* TODO: can we make this more numerically robust to avoid reintersecting the
* same light in some cases? Ray should not intersect surface anymore as the
* object and prim ids will prevent self intersection. */
const float3 new_ray_P = ray_P + ray_D * isect.t;
INTEGRATOR_STATE_WRITE(state, ray, P) = new_ray_P;
INTEGRATOR_STATE_WRITE(state, ray, t) -= isect.t;
/* Set position to where the BSDF was sampled, for correct MIS PDF. */
const float mis_ray_t = INTEGRATOR_STATE(state, path, mis_ray_t);
ray_P -= ray_D * mis_ray_t;
isect.t += mis_ray_t;
INTEGRATOR_STATE_WRITE(state, path, mis_ray_t) = isect.t;
/* Advance ray to new start distance. */
INTEGRATOR_STATE_WRITE(state, ray, tmin) = intersection_t_offset(isect.t);
LightSample ls ccl_optional_struct_init;
const bool use_light_sample = light_sample_from_intersection(kg, &isect, ray_P, ray_D, &ls);

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

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

@@ -77,7 +77,7 @@ ccl_device_forceinline void integrate_surface_emission(KernelGlobals kg,
# endif
{
const float bsdf_pdf = INTEGRATOR_STATE(state, path, mis_ray_pdf);
const float t = sd->ray_length + INTEGRATOR_STATE(state, path, mis_ray_t);
const float t = sd->ray_length;
/* Multiple importance sampling, get triangle light pdf,
* and compute weight with respect to BSDF pdf. */
@@ -323,16 +323,21 @@ ccl_device_forceinline int integrate_surface_bsdf_bssrdf_bounce(
return LABEL_NONE;
}
/* Setup ray. Note that clipping works through transparent bounces. */
INTEGRATOR_STATE_WRITE(state, ray, P) = sd->P;
INTEGRATOR_STATE_WRITE(state, ray, D) = normalize(bsdf_omega_in);
INTEGRATOR_STATE_WRITE(state, ray, t) = (label & LABEL_TRANSPARENT) ?
INTEGRATOR_STATE(state, ray, t) - sd->ray_length :
FLT_MAX;
if (label & LABEL_TRANSPARENT) {
/* Only need to modify start distance for transparent. */
INTEGRATOR_STATE_WRITE(state, ray, tmin) = intersection_t_offset(sd->ray_length);
}
else {
/* Setup ray with changed origin and direction. */
INTEGRATOR_STATE_WRITE(state, ray, P) = sd->P;
INTEGRATOR_STATE_WRITE(state, ray, D) = normalize(bsdf_omega_in);
INTEGRATOR_STATE_WRITE(state, ray, tmin) = 0.0f;
INTEGRATOR_STATE_WRITE(state, ray, tmax) = FLT_MAX;
#ifdef __RAY_DIFFERENTIALS__
INTEGRATOR_STATE_WRITE(state, ray, dP) = differential_make_compact(sd->dP);
INTEGRATOR_STATE_WRITE(state, ray, dD) = differential_make_compact(bsdf_domega_in);
INTEGRATOR_STATE_WRITE(state, ray, dP) = differential_make_compact(sd->dP);
INTEGRATOR_STATE_WRITE(state, ray, dD) = differential_make_compact(bsdf_domega_in);
#endif
}
/* Update throughput. */
float3 throughput = INTEGRATOR_STATE(state, path, throughput);
@@ -349,12 +354,8 @@ ccl_device_forceinline int integrate_surface_bsdf_bssrdf_bounce(
}
/* Update path state */
if (label & LABEL_TRANSPARENT) {
INTEGRATOR_STATE_WRITE(state, path, mis_ray_t) += sd->ray_length;
}
else {
if (!(label & LABEL_TRANSPARENT)) {
INTEGRATOR_STATE_WRITE(state, path, mis_ray_pdf) = bsdf_pdf;
INTEGRATOR_STATE_WRITE(state, path, mis_ray_t) = 0.0f;
INTEGRATOR_STATE_WRITE(state, path, min_ray_pdf) = fminf(
bsdf_pdf, INTEGRATOR_STATE(state, path, min_ray_pdf));
}
@@ -371,17 +372,8 @@ ccl_device_forceinline int integrate_surface_volume_only_bounce(IntegratorState
return LABEL_NONE;
}
/* Setup ray position, direction stays unchanged. */
INTEGRATOR_STATE_WRITE(state, ray, P) = sd->P;
/* Clipping works through transparent. */
INTEGRATOR_STATE_WRITE(state, ray, t) -= sd->ray_length;
# ifdef __RAY_DIFFERENTIALS__
INTEGRATOR_STATE_WRITE(state, ray, dP) = differential_make_compact(sd->dP);
# endif
INTEGRATOR_STATE_WRITE(state, path, mis_ray_t) += sd->ray_length;
/* Only modify start distance. */
INTEGRATOR_STATE_WRITE(state, ray, tmin) = intersection_t_offset(sd->ray_length);
return LABEL_TRANSMIT | LABEL_TRANSPARENT;
}
@@ -432,7 +424,8 @@ ccl_device_forceinline void integrate_surface_ao(KernelGlobals kg,
Ray ray ccl_optional_struct_init;
ray.P = shadow_ray_offset(kg, sd, ao_D, &skip_self);
ray.D = ao_D;
ray.t = kernel_data.integrator.ao_bounces_distance;
ray.tmin = 0.0f;
ray.tmax = kernel_data.integrator.ao_bounces_distance;
ray.time = sd->time;
ray.self.object = (skip_self) ? sd->object : OBJECT_NONE;
ray.self.prim = (skip_self) ? sd->prim : PRIM_NONE;
@@ -616,7 +609,7 @@ ccl_device_forceinline void integrator_shade_surface(KernelGlobals kg,
kg, state, current_kernel, DEVICE_KERNEL_INTEGRATOR_INTERSECT_SUBSURFACE);
}
else {
kernel_assert(INTEGRATOR_STATE(state, ray, t) != 0.0f);
kernel_assert(INTEGRATOR_STATE(state, ray, tmax) != 0.0f);
integrator_path_next(kg, state, current_kernel, DEVICE_KERNEL_INTEGRATOR_INTERSECT_CLOSEST);
}
}

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

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_MEMBER(shadow_ray, packed_float3, P, KERNEL_FEATURE_PATH_TRACING)
KERNEL_STRUCT_MEMBER(shadow_ray, packed_float3, D, KERNEL_FEATURE_PATH_TRACING)
KERNEL_STRUCT_MEMBER(shadow_ray, float, t, KERNEL_FEATURE_PATH_TRACING)
KERNEL_STRUCT_MEMBER(shadow_ray, float, tmin, KERNEL_FEATURE_PATH_TRACING)
KERNEL_STRUCT_MEMBER(shadow_ray, float, tmax, KERNEL_FEATURE_PATH_TRACING)
KERNEL_STRUCT_MEMBER(shadow_ray, float, time, KERNEL_FEATURE_PATH_TRACING)
KERNEL_STRUCT_MEMBER(shadow_ray, float, dP, KERNEL_FEATURE_PATH_TRACING)
KERNEL_STRUCT_MEMBER(shadow_ray, int, object, KERNEL_FEATURE_PATH_TRACING)

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

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, D) = ray->D;
INTEGRATOR_STATE_WRITE(state, ray, t) = ray->t;
INTEGRATOR_STATE_WRITE(state, ray, tmin) = ray->tmin;
INTEGRATOR_STATE_WRITE(state, ray, tmax) = ray->tmax;
INTEGRATOR_STATE_WRITE(state, ray, time) = ray->time;
INTEGRATOR_STATE_WRITE(state, ray, dP) = ray->dP;
INTEGRATOR_STATE_WRITE(state, ray, dD) = ray->dD;
@@ -29,7 +30,8 @@ ccl_device_forceinline void integrator_state_read_ray(KernelGlobals kg,
{
ray->P = INTEGRATOR_STATE(state, ray, P);
ray->D = INTEGRATOR_STATE(state, ray, D);
ray->t = INTEGRATOR_STATE(state, ray, t);
ray->tmin = INTEGRATOR_STATE(state, ray, tmin);
ray->tmax = INTEGRATOR_STATE(state, ray, tmax);
ray->time = INTEGRATOR_STATE(state, ray, time);
ray->dP = INTEGRATOR_STATE(state, ray, dP);
ray->dD = INTEGRATOR_STATE(state, ray, dD);
@@ -42,7 +44,8 @@ ccl_device_forceinline void integrator_state_write_shadow_ray(
{
INTEGRATOR_STATE_WRITE(state, shadow_ray, P) = ray->P;
INTEGRATOR_STATE_WRITE(state, shadow_ray, D) = ray->D;
INTEGRATOR_STATE_WRITE(state, shadow_ray, t) = ray->t;
INTEGRATOR_STATE_WRITE(state, shadow_ray, tmin) = ray->tmin;
INTEGRATOR_STATE_WRITE(state, shadow_ray, tmax) = ray->tmax;
INTEGRATOR_STATE_WRITE(state, shadow_ray, time) = ray->time;
INTEGRATOR_STATE_WRITE(state, shadow_ray, dP) = ray->dP;
}
@@ -53,7 +56,8 @@ ccl_device_forceinline void integrator_state_read_shadow_ray(KernelGlobals kg,
{
ray->P = INTEGRATOR_STATE(state, shadow_ray, P);
ray->D = INTEGRATOR_STATE(state, shadow_ray, D);
ray->t = INTEGRATOR_STATE(state, shadow_ray, t);
ray->tmin = INTEGRATOR_STATE(state, shadow_ray, tmin);
ray->tmax = INTEGRATOR_STATE(state, shadow_ray, tmax);
ray->time = INTEGRATOR_STATE(state, shadow_ray, time);
ray->dP = INTEGRATOR_STATE(state, shadow_ray, dP);
ray->dD = differential_zero_compact();

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

@@ -38,7 +38,8 @@ ccl_device int subsurface_bounce(KernelGlobals kg,
/* Setup ray into surface. */
INTEGRATOR_STATE_WRITE(state, ray, P) = sd->P;
INTEGRATOR_STATE_WRITE(state, ray, D) = bssrdf->N;
INTEGRATOR_STATE_WRITE(state, ray, t) = FLT_MAX;
INTEGRATOR_STATE_WRITE(state, ray, tmin) = 0.0f;
INTEGRATOR_STATE_WRITE(state, ray, tmax) = FLT_MAX;
INTEGRATOR_STATE_WRITE(state, ray, dP) = differential_make_compact(sd->dP);
INTEGRATOR_STATE_WRITE(state, ray, dD) = differential_zero_compact();
@@ -160,7 +161,7 @@ ccl_device_inline bool subsurface_scatter(KernelGlobals kg, IntegratorState stat
/* Pretend ray is coming from the outside towards the exit point. This ensures
* correct front/back facing normals.
* TODO: find a more elegant solution? */
ray.P += ray.D * ray.t * 2.0f;
ray.P += ray.D * ray.tmax * 2.0f;
ray.D = -ray.D;
integrator_state_write_isect(kg, state, &ss_isect.hits[0]);

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

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

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

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

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

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) {
/* setup ray */
ray->P = P;
ray->tmin = 0.0f;
if (ls->t == FLT_MAX) {
/* distant light */
ray->D = ls->D;
ray->t = ls->t;
ray->tmax = ls->t;
}
else {
/* other lights, avoid self-intersection */
ray->D = ls->P - P;
ray->D = normalize_len(ray->D, &ray->t);
ray->D = normalize_len(ray->D, &ray->tmax);
}
}
else {
/* signal to not cast shadow ray */
ray->P = zero_float3();
ray->D = zero_float3();
ray->t = 0.0f;
ray->tmax = 0.0f;
}
ray->dP = differential_make_compact(sd->dP);

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);
if (derivatives) {
ndc[1] = camera_world_to_ndc(kg, sd, sd->ray_P + make_float3(sd->ray_dP, 0.0f, 0.0f)) -
ndc[0];
ndc[2] = camera_world_to_ndc(kg, sd, sd->ray_P + make_float3(0.0f, sd->ray_dP, 0.0f)) -
ndc[0];
ndc[1] = zero_float3();
ndc[2] = zero_float3();
}
}
else {
@@ -1671,7 +1669,8 @@ bool OSLRenderServices::trace(TraceOpt &options,
ray.P = TO_FLOAT3(P);
ray.D = TO_FLOAT3(R);
ray.t = (options.maxdist == 1.0e30f) ? FLT_MAX : options.maxdist - options.mindist;
ray.tmin = 0.0f;
ray.tmax = (options.maxdist == 1.0e30f) ? FLT_MAX : options.maxdist - options.mindist;
ray.time = sd->time;
ray.self.object = OBJECT_NONE;
ray.self.prim = PRIM_NONE;

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

@@ -59,7 +59,8 @@ ccl_device float svm_ao(
Ray ray;
ray.P = sd->P;
ray.D = D.x * T + D.y * B + D.z * N;
ray.t = max_dist;
ray.tmin = 0.0f;
ray.tmax = max_dist;
ray.time = sd->time;
ray.self.object = sd->object;
ray.self.prim = sd->prim;

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.P = sd->P + disk_N * disk_height + disk_P;
ray.D = -disk_N;
ray.t = 2.0f * disk_height;
ray.tmin = 0.0f;
ray.tmax = 2.0f * disk_height;
ray.dP = differential_zero_compact();
ray.dD = differential_zero_compact();
ray.time = sd->time;

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

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

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

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

15
intern/guardedalloc/intern/mallocn_lockfree_impl.c
View File

@@ -44,6 +44,7 @@ enum {
#define PTR_FROM_MEMHEAD(memhead) (memhead + 1)
#define MEMHEAD_ALIGNED_FROM_PTR(ptr) (((MemHeadAligned *)ptr) - 1)
#define MEMHEAD_IS_ALIGNED(memhead) ((memhead)->len & (size_t)MEMHEAD_ALIGN_FLAG)
#define MEMHEAD_LEN(memhead) ((memhead)->len & ~((size_t)(MEMHEAD_ALIGN_FLAG)))
/* Uncomment this to have proper peak counter. */
#define USE_ATOMIC_MAX
@@ -78,8 +79,8 @@ print_error(const char *str, ...)
size_t MEM_lockfree_allocN_len(const void *vmemh)
{
if (vmemh) {
return MEMHEAD_FROM_PTR(vmemh)->len & ~((size_t)(MEMHEAD_ALIGN_FLAG));
if (LIKELY(vmemh)) {
return MEMHEAD_LEN(MEMHEAD_FROM_PTR(vmemh));
}
return 0;
@@ -87,14 +88,11 @@ size_t MEM_lockfree_allocN_len(const void *vmemh)
void MEM_lockfree_freeN(void *vmemh)
{
if (leak_detector_has_run) {
if (UNLIKELY(leak_detector_has_run)) {
print_error("%s\n", free_after_leak_detection_message);
}
MemHead *memh = MEMHEAD_FROM_PTR(vmemh);
size_t len = MEM_lockfree_allocN_len(vmemh);
if (vmemh == NULL) {
if (UNLIKELY(vmemh == NULL)) {
print_error("Attempt to free NULL pointer\n");
#ifdef WITH_ASSERT_ABORT
abort();
@@ -102,6 +100,9 @@ void MEM_lockfree_freeN(void *vmemh)
return;
}
MemHead *memh = MEMHEAD_FROM_PTR(vmemh);
size_t len = MEMHEAD_LEN(memh);
atomic_sub_and_fetch_u(&totblock, 1);
atomic_sub_and_fetch_z(&mem_in_use, len);

5
source/blender/blenlib/intern/filereader_zstd.c
View File

@@ -281,7 +281,10 @@ static void zstd_close(FileReader *reader)
if (zstd->reader.seek) {
MEM_freeN(zstd->seek.uncompressed_ofs);
MEM_freeN(zstd->seek.compressed_ofs);
MEM_freeN(zstd->seek.cached_content);
/* When an error has occurred this may be NULL, see: T99744. */
if (zstd->seek.cached_content) {
MEM_freeN(zstd->seek.cached_content);
}
}
else {
MEM_freeN((void *)zstd->in_buf.src);

12
source/blender/gpu/GPU_context.h
View File

@@ -17,14 +17,10 @@
extern "C" {
#endif
/* GPU backends abstract the differences between different APIs. These must be
* initialized before creating contexts, and deleted after the last context is
* discarded. GPU_context_create automatically initializes a backend if none
* exists yet. */
bool GPU_backend_init_once(void);
void GPU_backend_exit(void);
bool GPU_backend_supported(eGPUBackendType type);
/* GPU backends abstract the differences between different APIs. GPU_context_create
* automatically initializes the backend, and GPU_context_discard frees it when there
* are no more contexts. */
bool GPU_backend_supported(void);
eGPUBackendType GPU_backend_get_type(void);
/** Opaque type hiding blender::gpu::Context. */

51
source/blender/gpu/intern/gpu_context.cc
View File

@@ -44,6 +44,12 @@ using namespace blender::gpu;
static thread_local Context *active_ctx = nullptr;
static std::mutex backend_users_mutex;
static int num_backend_users = 0;
static void gpu_backend_create();
static void gpu_backend_discard();
/* -------------------------------------------------------------------- */
/** \name gpu::Context methods
* \{ */
@@ -86,7 +92,14 @@ Context *Context::get()
GPUContext *GPU_context_create(void *ghost_window)
{
GPU_backend_init_once();
{
std::scoped_lock lock(backend_users_mutex);
if (num_backend_users == 0) {
/* Automatically create backend when first context is created. */
gpu_backend_create();
}
num_backend_users++;
}
Context *ctx = GPUBackend::get()->context_alloc(ghost_window);
@@ -99,6 +112,16 @@ void GPU_context_discard(GPUContext *ctx_)
Context *ctx = unwrap(ctx_);
delete ctx;
active_ctx = nullptr;
{
std::scoped_lock lock(backend_users_mutex);
num_backend_users--;
BLI_assert(num_backend_users >= 0);
if (num_backend_users == 0) {
/* Discard backend when last context is discarded. */
gpu_backend_discard();
}
}
}
void GPU_context_active_set(GPUContext *ctx_)
@@ -189,11 +212,12 @@ void GPU_render_step()
/** \name Backend selection
* \{ */
static const eGPUBackendType g_backend_type = GPU_BACKEND_OPENGL;
static GPUBackend *g_backend = nullptr;
bool GPU_backend_supported(eGPUBackendType type)
bool GPU_backend_supported(void)
{
switch (type) {
switch (g_backend_type) {
case GPU_BACKEND_OPENGL:
#ifdef WITH_OPENGL_BACKEND
return true;
@@ -212,21 +236,12 @@ bool GPU_backend_supported(eGPUBackendType type)
}
}
bool GPU_backend_init_once()
static void gpu_backend_create()
{
if (GPUBackend::get() != nullptr) {
return true;
}
BLI_assert(g_backend == nullptr);
BLI_assert(GPU_backend_supported());
const eGPUBackendType backend_type = GPU_BACKEND_OPENGL;
if (!GPU_backend_supported(backend_type)) {
return false;
}
static std::mutex backend_init_mutex;
std::scoped_lock lock(backend_init_mutex);
switch (backend_type) {
switch (g_backend_type) {
#ifdef WITH_OPENGL_BACKEND
case GPU_BACKEND_OPENGL:
g_backend = new GLBackend;
@@ -241,8 +256,6 @@ bool GPU_backend_init_once()
BLI_assert(0);
break;
}
return true;
}
void gpu_backend_delete_resources()
@@ -251,7 +264,7 @@ void gpu_backend_delete_resources()
g_backend->delete_resources();
}
void GPU_backend_exit()
void gpu_backend_discard()
{
/* TODO: assert no resource left. */
delete g_backend;

2
source/blender/gpu/tests/gpu_testing.cc
View File

@@ -17,7 +17,6 @@ void GPUTest::SetUp()
GHOST_GLSettings glSettings = {0};
CLG_init();
ghost_system = GHOST_CreateSystem();
GPU_backend_init_once();
ghost_context = GHOST_CreateOpenGLContext(ghost_system, glSettings);
GHOST_ActivateOpenGLContext(ghost_context);
context = GPU_context_create(nullptr);
@@ -29,7 +28,6 @@ void GPUTest::TearDown()
GPU_exit();
GPU_context_discard(context);
GHOST_DisposeOpenGLContext(ghost_system, ghost_context);
GPU_backend_exit();
GHOST_DisposeSystem(ghost_system);
CLG_exit();
}

10
source/blender/python/generic/CMakeLists.txt
View File

@@ -17,6 +17,7 @@ set(INC_SYS
)
set(SRC
bgl.c
bl_math_py_api.c
blf_py_api.c
bpy_threads.c
@@ -26,6 +27,7 @@ set(SRC
py_capi_rna.c
py_capi_utils.c
bgl.h
bl_math_py_api.h
blf_py_api.h
idprop_py_api.h
@@ -38,14 +40,6 @@ set(SRC
python_utildefines.h
)
if(WITH_OPENGL)
list(APPEND SRC
bgl.c
bgl.h
)
endif()
set(LIB
${GLEW_LIBRARY}
${PYTHON_LINKFLAGS}

66
source/blender/python/generic/bgl.c
View File

@@ -1082,18 +1082,34 @@ static PyObject *Buffer_repr(Buffer *self)
/** \name OpenGL API Wrapping
* \{ */
#define BGL_Wrap(funcname, ret, arg_list) \
static PyObject *Method_##funcname(PyObject *UNUSED(self), PyObject *args) \
{ \
arg_def arg_list; \
ret_def_##ret; \
if (!PyArg_ParseTuple(args, arg_str arg_list, arg_ref arg_list)) { \
#ifdef WITH_OPENGL
# define BGL_Wrap(funcname, ret, arg_list) \
static PyObject *Method_##funcname(PyObject *UNUSED(self), PyObject *args) \
{ \
arg_def arg_list; \
ret_def_##ret; \
if (!PyArg_ParseTuple(args, arg_str arg_list, arg_ref arg_list)) { \
return NULL; \
} \
GPU_bgl_start(); \
ret_set_##ret gl##funcname(arg_var arg_list); \
ret_ret_##ret; \
}
#else
static void bgl_no_opengl_error(void)
{
PyErr_SetString(PyExc_RuntimeError, "Built without OpenGL support");
}
# define BGL_Wrap(funcname, ret, arg_list) \
static PyObject *Method_##funcname(PyObject *UNUSED(self), PyObject *args) \
{ \
(void)args; \
bgl_no_opengl_error(); \
return NULL; \
} \
GPU_bgl_start(); \
ret_set_##ret gl##funcname(arg_var arg_list); \
ret_ret_##ret; \
}
}
#endif
/* GL_VERSION_1_0 */
BGL_Wrap(BlendFunc, void, (GLenum, GLenum));
@@ -1421,12 +1437,22 @@ static void py_module_dict_add_method(PyObject *submodule,
#ifdef __GNUC__
# pragma GCC diagnostic ignored "-Waddress"
#endif
#define PY_MOD_ADD_METHOD(func) \
{ \
static PyMethodDef method_def = {"gl" #func, Method_##func, METH_VARARGS}; \
py_module_dict_add_method(submodule, dict, &method_def, (gl##func != NULL)); \
} \
((void)0)
#ifdef WITH_OPENGL
# define PY_MOD_ADD_METHOD(func) \
{ \
static PyMethodDef method_def = {"gl" #func, Method_##func, METH_VARARGS}; \
py_module_dict_add_method(submodule, dict, &method_def, (gl##func != NULL)); \
} \
((void)0)
#else
# define PY_MOD_ADD_METHOD(func) \
{ \
static PyMethodDef method_def = {"gl" #func, Method_##func, METH_VARARGS}; \
py_module_dict_add_method(submodule, dict, &method_def, false); \
} \
((void)0)
#endif
static void init_bgl_version_1_0_methods(PyObject *submodule, PyObject *dict)
{
@@ -2620,9 +2646,13 @@ static PyObject *Method_ShaderSource(PyObject *UNUSED(self), PyObject *args)
return NULL;
}
#ifdef WITH_OPENGL
glShaderSource(shader, 1, (const char **)&source, NULL);
Py_RETURN_NONE;
#else
bgl_no_opengl_error();
return NULL;
#endif
}
/** \} */

2
source/blender/python/intern/bpy_interface.c
View File

@@ -259,9 +259,7 @@ static struct _inittab bpy_internal_modules[] = {
{"mathutils.kdtree", PyInit_mathutils_kdtree},
#endif
{"_bpy_path", BPyInit__bpy_path},
#ifdef WITH_OPENGL
{"bgl", BPyInit_bgl},
#endif
{"blf", BPyInit_blf},
{"bl_math", BPyInit_bl_math},
{"imbuf", BPyInit_imbuf},

2
source/blender/render/intern/engine.c
View File

@@ -952,7 +952,7 @@ bool RE_engine_render(Render *re, bool do_all)
re->draw_lock(re->dlh, true);
}
if ((type->flag & RE_USE_GPU_CONTEXT) && (GPU_backend_get_type() == GPU_BACKEND_NONE)) {
if ((type->flag & RE_USE_GPU_CONTEXT) && !GPU_backend_supported()) {
/* Clear UI drawing locks. */
if (re->draw_lock) {
re->draw_lock(re->dlh, false);

3
source/blender/windowmanager/intern/wm_init_exit.c
View File

@@ -169,7 +169,7 @@ void WM_init_opengl(void)
wm_ghost_init(NULL);
}
if (!GPU_backend_init_once()) {
if (!GPU_backend_supported()) {
return;
}
@@ -613,7 +613,6 @@ void WM_exit_ex(bContext *C, const bool do_python)
else {
UI_exit();
}
GPU_backend_exit();
BKE_blender_userdef_data_free(&U, false);