diff options
author | Brecht Van Lommel <brecht@blender.org> | 2022-07-13 17:54:53 +0300 |
---|---|---|
committer | Brecht Van Lommel <brecht@blender.org> | 2022-07-15 19:46:24 +0300 |
commit | 5152c7c152e52d563cbd3ba3c792de3af0c2c14f (patch) | |
tree | 45cad004066413ce309279ad52f2a363e5477150 /intern/cycles/kernel/bvh/bvh.h | |
parent | bb376da6dfdd2476fc3738ce1fc89dac27825cef (diff) |
Cycles: refactor rays to have start and end distance, fix precision issues
For transparency, volume and light intersection rays, adjust these distances
rather than the ray start position. This way we increment the start distance
by the smallest possible float increment to avoid self intersections, and be
sure it works as the distance compared to be will be exactly the same as
before, due to the ray start position and direction remaining the same.
Fix T98764, T96537, hair ray tracing precision issues.
Differential Revision: https://developer.blender.org/D15455
Diffstat (limited to 'intern/cycles/kernel/bvh/bvh.h')
-rw-r--r-- | intern/cycles/kernel/bvh/bvh.h | 36 |
1 files changed, 18 insertions, 18 deletions
diff --git a/intern/cycles/kernel/bvh/bvh.h b/intern/cycles/kernel/bvh/bvh.h index f375529a6f6..9972de86c47 100644 --- a/intern/cycles/kernel/bvh/bvh.h +++ b/intern/cycles/kernel/bvh/bvh.h @@ -175,8 +175,8 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg, optixTrace(scene_intersect_valid(ray) ? kernel_data.device_bvh : 0, ray->P, ray->D, - 0.0f, - ray->t, + ray->tmin, + ray->tmax, ray->time, ray_mask, ray_flags, @@ -203,28 +203,28 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg, #elif defined(__METALRT__) if (!scene_intersect_valid(ray)) { - isect->t = ray->t; + isect->t = ray->tmax; isect->type = PRIMITIVE_NONE; return false; } # if defined(__KERNEL_DEBUG__) if (is_null_instance_acceleration_structure(metal_ancillaries->accel_struct)) { - isect->t = ray->t; + isect->t = ray->tmax; isect->type = PRIMITIVE_NONE; kernel_assert(!"Invalid metal_ancillaries->accel_struct pointer"); return false; } if (is_null_intersection_function_table(metal_ancillaries->ift_default)) { - isect->t = ray->t; + isect->t = ray->tmax; isect->type = PRIMITIVE_NONE; kernel_assert(!"Invalid ift_default"); return false; } # endif - metal::raytracing::ray r(ray->P, ray->D, 0.0f, ray->t); + metal::raytracing::ray r(ray->P, ray->D, ray->tmin, ray->tmax); metalrt_intersector_type metalrt_intersect; if (!kernel_data.bvh.have_curves) { @@ -263,7 +263,7 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg, # endif if (intersection.type == intersection_type::none) { - isect->t = ray->t; + isect->t = ray->tmax; isect->type = PRIMITIVE_NONE; return false; @@ -296,7 +296,7 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg, # ifdef __EMBREE__ if (kernel_data.device_bvh) { - isect->t = ray->t; + isect->t = ray->tmax; CCLIntersectContext ctx(kg, CCLIntersectContext::RAY_REGULAR); IntersectContext rtc_ctx(&ctx); RTCRayHit ray_hit; @@ -360,8 +360,8 @@ ccl_device_intersect bool scene_intersect_local(KernelGlobals kg, optixTrace(scene_intersect_valid(ray) ? kernel_data.device_bvh : 0, ray->P, ray->D, - 0.0f, - ray->t, + ray->tmin, + ray->tmax, ray->time, 0xFF, /* Need to always call into __anyhit__kernel_optix_local_hit. */ @@ -405,7 +405,7 @@ ccl_device_intersect bool scene_intersect_local(KernelGlobals kg, } # endif - metal::raytracing::ray r(ray->P, ray->D, 0.0f, ray->t); + metal::raytracing::ray r(ray->P, ray->D, ray->tmin, ray->tmax); metalrt_intersector_type metalrt_intersect; metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque); @@ -476,7 +476,7 @@ ccl_device_intersect bool scene_intersect_local(KernelGlobals kg, float3 dir = ray->D; float3 idir = ray->D; Transform ob_itfm; - rtc_ray.tfar = ray->t * + rtc_ray.tfar = ray->tmax * bvh_instance_motion_push(kg, local_object, ray, &P, &dir, &idir, &ob_itfm); /* bvh_instance_motion_push() returns the inverse transform but * it's not needed here. */ @@ -542,8 +542,8 @@ ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg, optixTrace(scene_intersect_valid(ray) ? kernel_data.device_bvh : 0, ray->P, ray->D, - 0.0f, - ray->t, + ray->tmin, + ray->tmax, ray->time, ray_mask, /* Need to always call into __anyhit__kernel_optix_shadow_all_hit. */ @@ -582,7 +582,7 @@ ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg, } # endif - metal::raytracing::ray r(ray->P, ray->D, 0.0f, ray->t); + metal::raytracing::ray r(ray->P, ray->D, ray->tmin, ray->tmax); metalrt_intersector_type metalrt_intersect; metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque); @@ -701,8 +701,8 @@ ccl_device_intersect bool scene_intersect_volume(KernelGlobals kg, optixTrace(scene_intersect_valid(ray) ? kernel_data.device_bvh : 0, ray->P, ray->D, - 0.0f, - ray->t, + ray->tmin, + ray->tmax, ray->time, ray_mask, /* Need to always call into __anyhit__kernel_optix_volume_test. */ @@ -744,7 +744,7 @@ ccl_device_intersect bool scene_intersect_volume(KernelGlobals kg, } # endif - metal::raytracing::ray r(ray->P, ray->D, 0.0f, ray->t); + metal::raytracing::ray r(ray->P, ray->D, ray->tmin, ray->tmax); metalrt_intersector_type metalrt_intersect; metalrt_intersect.force_opacity(metal::raytracing::forced_opacity::non_opaque); |