[Bf-blender-cvs] [5152c7c152e] master: Cycles: refactor rays to have start and end distance, fix precision issues
Brecht Van Lommel
noreply at git.blender.org
Fri Jul 15 18:56:03 CEST 2022
Commit: 5152c7c152e52d563cbd3ba3c792de3af0c2c14f
Author: Brecht Van Lommel
Date: Wed Jul 13 16:54:53 2022 +0200
Branches: master
https://developer.blender.org/rB5152c7c152e52d563cbd3ba3c792de3af0c2c14f
Cycles: refactor rays to have start and end distance, fix precision issues
For transparency, volume and light intersection rays, adjust these distances
rather than the ray start position. This way we increment the start distance
by the smallest possible float increment to avoid self intersections, and be
sure it works as the distance compared to be will be exactly the same as
before, due to the ray start position and direction remaining the same.
Fix T98764, T96537, hair ray tracing precision issues.
Differential Revision: https://developer.blender.org/D15455
===================================================================
M intern/cycles/kernel/bvh/bvh.h
M intern/cycles/kernel/bvh/embree.h
M intern/cycles/kernel/bvh/local.h
M intern/cycles/kernel/bvh/nodes.h
M intern/cycles/kernel/bvh/shadow_all.h
M intern/cycles/kernel/bvh/traversal.h
M intern/cycles/kernel/bvh/util.h
M intern/cycles/kernel/bvh/volume.h
M intern/cycles/kernel/bvh/volume_all.h
M intern/cycles/kernel/camera/camera.h
M intern/cycles/kernel/device/metal/kernel.metal
M intern/cycles/kernel/device/optix/kernel.cu
M intern/cycles/kernel/geom/curve_intersect.h
M intern/cycles/kernel/geom/motion_triangle_intersect.h
M intern/cycles/kernel/geom/point_intersect.h
M intern/cycles/kernel/geom/shader_data.h
M intern/cycles/kernel/geom/triangle_intersect.h
M intern/cycles/kernel/integrator/init_from_bake.h
M intern/cycles/kernel/integrator/init_from_camera.h
M intern/cycles/kernel/integrator/intersect_closest.h
M intern/cycles/kernel/integrator/intersect_volume_stack.h
M intern/cycles/kernel/integrator/mnee.h
M intern/cycles/kernel/integrator/path_state.h
M intern/cycles/kernel/integrator/shade_background.h
M intern/cycles/kernel/integrator/shade_light.h
M intern/cycles/kernel/integrator/shade_shadow.h
M intern/cycles/kernel/integrator/shade_surface.h
M intern/cycles/kernel/integrator/shade_volume.h
M intern/cycles/kernel/integrator/shadow_state_template.h
M intern/cycles/kernel/integrator/state_template.h
M intern/cycles/kernel/integrator/state_util.h
M intern/cycles/kernel/integrator/subsurface.h
M intern/cycles/kernel/integrator/subsurface_disk.h
M intern/cycles/kernel/integrator/subsurface_random_walk.h
M intern/cycles/kernel/light/light.h
M intern/cycles/kernel/light/sample.h
M intern/cycles/kernel/osl/services.cpp
M intern/cycles/kernel/svm/ao.h
M intern/cycles/kernel/svm/bevel.h
M intern/cycles/kernel/svm/tex_coord.h
M intern/cycles/kernel/types.h
M intern/cycles/util/math_intersect.h
===================================================================
diff --git a/intern/cycles/kernel/bvh/bvh.h b/intern/cycles/kernel/bvh/bvh.h
index f375529a6f6..9972de86c47 100644
--- a/intern/cycles/kernel/bvh/bvh.h
+++ b/intern/cycles/kernel/bvh/bvh.h
@@ -175,8 +175,8 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg,
optixTrace(scene_intersect_valid(ray) ? kernel_data.device_bvh : 0,
ray->P,
ray->D,
- 0.0f,
- ray->t,
+ ray->tmin,
+ ray->tmax,
ray->time,
ray_mask,
ray_flags,
@@ -203,28 +203,28 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg,
#elif defined(__METALRT__)
if (!scene_intersect_valid(ray)) {
- isect->t = ray->t;
+ isect->t = ray->tmax;
isect->type = PRIMITIVE_NONE;
return false;
}
# if defined(__KERNEL_DEBUG__)
if (is_null_instance_acceleration_structure(metal_ancillaries->accel_struct)) {
- isect->t = ray->t;
+ isect->t = ray->tmax;
isect->type = PRIMITIVE_NONE;
kernel_assert(!"Invalid metal_ancillaries->accel_struct pointer");
return false;
}
if (is_null_intersection_function_table(metal_ancillaries->ift_default)) {
- isect->t = ray->t;
+ isect->t = ray->tmax;
isect->type = PRIMITIVE_NONE;
kernel_assert(!"Invalid ift_default");
return false;
}
# endif
- metal::raytracing::ray r(ray->P, ray->D, 0.0f, ray->t);
+ metal::raytracing::ray r(ray->P, ray->D, ray->tmin, ray->tmax);
metalrt_intersector_type metalrt_intersect;
if (!kernel_data.bvh.have_curves) {
@@ -263,7 +263,7 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg,
# endif
if (intersection.type == intersection_type::none) {
- isect->t = ray->t;
+ isect->t = ray->tmax;
isect->type = PRIMITIVE_NONE;
return false;
@@ -296,7 +296,7 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg,
# ifdef __EMBREE__
if (kernel_data.device_bvh) {
- isect->t = ray->t;
+ isect->t = ray->tmax;
CCLIntersectContext ctx(kg, CCLIntersectContext::RAY_REGULAR);
IntersectContext rtc_ctx(&ctx);
RTCRayHit ray_hit;
@@ -360,8 +360,8 @@ ccl_device_intersect bool scene_intersect_local(KernelGlobals kg,
optixTrace(scene_intersect_valid(ray) ? kernel_data.device_bvh : 0,
ray->P,
ray->D,
- 0.0f,
- ray->t,
+ ray->tmin,
+ ray->tmax,
ray->time,
0xFF,
/* Need to always call into __anyhit__kernel_optix_local_hit. */
@@ -405,7 +405,7 @@ ccl_device_intersect bool scene_intersect_local(KernelGlobals kg,
}
# endif
- metal::raytracing::ray r(ray->P, ray->D, 0.0f, ray->t);
+ metal::raytracing::ray r(ray->P, ray->D, ray->tmin, ray->tmax);
metalrt_intersector_type metalrt_intersect;
metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque);
@@ -476,7 +476,7 @@ ccl_device_intersect bool scene_intersect_local(KernelGlobals kg,
float3 dir = ray->D;
float3 idir = ray->D;
Transform ob_itfm;
- rtc_ray.tfar = ray->t *
+ rtc_ray.tfar = ray->tmax *
bvh_instance_motion_push(kg, local_object, ray, &P, &dir, &idir, &ob_itfm);
/* bvh_instance_motion_push() returns the inverse transform but
* it's not needed here. */
@@ -542,8 +542,8 @@ ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg,
optixTrace(scene_intersect_valid(ray) ? kernel_data.device_bvh : 0,
ray->P,
ray->D,
- 0.0f,
- ray->t,
+ ray->tmin,
+ ray->tmax,
ray->time,
ray_mask,
/* Need to always call into __anyhit__kernel_optix_shadow_all_hit. */
@@ -582,7 +582,7 @@ ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg,
}
# endif
- metal::raytracing::ray r(ray->P, ray->D, 0.0f, ray->t);
+ metal::raytracing::ray r(ray->P, ray->D, ray->tmin, ray->tmax);
metalrt_intersector_type metalrt_intersect;
metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque);
@@ -701,8 +701,8 @@ ccl_device_intersect bool scene_intersect_volume(KernelGlobals kg,
optixTrace(scene_intersect_valid(ray) ? kernel_data.device_bvh : 0,
ray->P,
ray->D,
- 0.0f,
- ray->t,
+ ray->tmin,
+ ray->tmax,
ray->time,
ray_mask,
/* Need to always call into __anyhit__kernel_optix_volume_test. */
@@ -744,7 +744,7 @@ ccl_device_intersect bool scene_intersect_volume(KernelGlobals kg,
}
# endif
- metal::raytracing::ray r(ray->P, ray->D, 0.0f, ray->t);
+ metal::raytracing::ray r(ray->P, ray->D, ray->tmin, ray->tmax);
metalrt_intersector_type metalrt_intersect;
metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque);
diff --git a/intern/cycles/kernel/bvh/embree.h b/intern/cycles/kernel/bvh/embree.h
index 77eec2468f4..fecbccac2f8 100644
--- a/intern/cycles/kernel/bvh/embree.h
+++ b/intern/cycles/kernel/bvh/embree.h
@@ -83,8 +83,8 @@ ccl_device_inline void kernel_embree_setup_ray(const Ray &ray,
rtc_ray.dir_x = ray.D.x;
rtc_ray.dir_y = ray.D.y;
rtc_ray.dir_z = ray.D.z;
- rtc_ray.tnear = 0.0f;
- rtc_ray.tfar = ray.t;
+ rtc_ray.tnear = ray.tmin;
+ rtc_ray.tfar = ray.tmax;
rtc_ray.time = ray.time;
rtc_ray.mask = visibility;
}
diff --git a/intern/cycles/kernel/bvh/local.h b/intern/cycles/kernel/bvh/local.h
index 3b6b30ea93d..017a241ef4a 100644
--- a/intern/cycles/kernel/bvh/local.h
+++ b/intern/cycles/kernel/bvh/local.h
@@ -47,8 +47,9 @@ ccl_device_inline
float3 P = ray->P;
float3 dir = bvh_clamp_direction(ray->D);
float3 idir = bvh_inverse_direction(dir);
+ float tmin = ray->tmin;
int object = OBJECT_NONE;
- float isect_t = ray->t;
+ float isect_t = ray->tmax;
if (local_isect != NULL) {
local_isect->num_hits = 0;
@@ -59,10 +60,13 @@ ccl_device_inline
if (!(object_flag & SD_OBJECT_TRANSFORM_APPLIED)) {
#if BVH_FEATURE(BVH_MOTION)
Transform ob_itfm;
- isect_t *= bvh_instance_motion_push(kg, local_object, ray, &P, &dir, &idir, &ob_itfm);
+ const float t_world_to_instance = bvh_instance_motion_push(
+ kg, local_object, ray, &P, &dir, &idir, &ob_itfm);
#else
- isect_t *= bvh_instance_push(kg, local_object, ray, &P, &dir, &idir);
+ const float t_world_to_instance = bvh_instance_push(kg, local_object, ray, &P, &dir, &idir);
#endif
+ isect_t *= t_world_to_instance;
+ tmin *= t_world_to_instance;
object = local_object;
}
@@ -81,6 +85,7 @@ ccl_device_inline
dir,
#endif
idir,
+ tmin,
isect_t,
node_addr,
PATH_RAY_ALL_VISIBILITY,
@@ -155,6 +160,7 @@ ccl_device_inline
local_object,
prim,
prim_addr,
+ tmin,
isect_t,
lcg_state,
max_hits)) {
@@ -191,6 +197,7 @@ ccl_device_inline
local_object,
prim,
prim_addr,
+ tmin,
isect_t,
lcg_state,
max_hits)) {
diff --git a/intern/cycles/kernel/bvh/nodes.h b/intern/cycles/kernel/bvh/nodes.h
index c19dea9223b..e02841fad16 100644
--- a/intern/cycles/kernel/bvh/nodes.h
+++ b/intern/cycles/kernel/bvh/nodes.h
@@ -18,7 +18,8 @@ ccl_device_forceinline Transform bvh_unaligned_node_fetch_space(KernelGlobals kg
ccl_device_forceinline int bvh_aligned_node_intersect(KernelGlobals kg,
const float3 P,
const float3 idir,
- const float t,
+ const float tmin,
+ const float tmax,
const int node_addr,
const uint visibility,
float dist[2])
@@ -39,8 +40,8 @@ ccl_device_forceinline int bvh_aligned_node_intersect(KernelGlobals kg,
float c0hiy = (node1.z - P.y) * idir.y;
float c0loz = (node2.x - P.z) * idir.z;
float c0hiz = (node2.z - P.z) * idir.z;
- float c0min = max4(0.0f, min(c0lox, c0hix), min(c0loy, c0hiy), min(c0loz, c0hiz));
- float c0max = min4(t, max(c0lox, c0hix), max(c0loy, c0hiy), max(c0loz, c0hiz));
+ float c0min = max4(tmin, min(c0lox, c0hix), min(c0loy, c0hiy), min(c0loz, c0hiz));
+ float c0max = min4(tmax, max(c0lox, c0hix), max(c0loy, c0hiy), max(c0loz, c0hiz));
float c1lox = (node0.y - P.x) * idir.x;
float c1hix = (node0.w - P.x) * idir.x;
@@ -48,8 +49,8 @@ ccl_device_forceinline int bvh_aligned_node_intersect(KernelGlobals kg,
float c1hiy = (node1.w - P.y) * idir.y;
float c1loz = (node2.y - P.z) * idir.z;
float c1hiz = (node2.w - P.z) * idir.z;
- float c1min = max4(0.0f, min(c1lox, c1hix), min(c1loy, c1hiy), min(c1loz, c1hiz));
- float c1max = min4(t, max(c1lox, c1hix), max(c1loy, c1hiy), max(c1loz, c1hiz));
+ float c1min = max4(tmin, min(c1lox, c1hix), min(c1loy, c1hiy), min(c1loz, c1hiz));
+ float c1max = min4(tmax, max(c1lox, c1hix), max(c1loy, c1hiy), max(c1loz, c1hiz));
dist[0] = c0min;
dist[1] = c1min;
@@ -66,7 +67,8 @@ ccl_device_forceinline int bvh_aligned_node_intersect(KernelGlobals kg,
ccl_device_forceinline bool bvh_unaligned_no
@@ Diff output truncated at 10240 characters. @@
More information about the Bf-blender-cvs
mailing list