diff options
author | Brecht Van Lommel <brecht@blender.org> | 2022-07-26 17:07:50 +0300 |
---|---|---|
committer | Brecht Van Lommel <brecht@blender.org> | 2022-07-27 22:03:33 +0300 |
commit | 38af5b0501005d8ee84a59f027417bf6a31fcc5e (patch) | |
tree | 0856dd31787e59abe9300a8d4b3c6b37be3cfbc8 /intern/cycles/kernel/device | |
parent | 69f2732a1391680d252c86365b2df62b084ceeb8 (diff) |
Cycles: switch Cycles triangle barycentric convention to match Embree/OptiX
Simplifies intersection code a little and slightly improves precision regarding
self intersection.
The parametric texture coordinate in shader nodes is still the same as before
for compatibility.
Diffstat (limited to 'intern/cycles/kernel/device')
-rw-r--r-- | intern/cycles/kernel/device/cpu/bvh.h | 8 | ||||
-rw-r--r-- | intern/cycles/kernel/device/metal/bvh.h | 10 | ||||
-rw-r--r-- | intern/cycles/kernel/device/metal/kernel.metal | 12 | ||||
-rw-r--r-- | intern/cycles/kernel/device/optix/bvh.h | 12 |
4 files changed, 18 insertions, 24 deletions
diff --git a/intern/cycles/kernel/device/cpu/bvh.h b/intern/cycles/kernel/device/cpu/bvh.h index 8ebeee99c47..6c06232a692 100644 --- a/intern/cycles/kernel/device/cpu/bvh.h +++ b/intern/cycles/kernel/device/cpu/bvh.h @@ -166,16 +166,16 @@ ccl_device_inline void kernel_embree_convert_hit(KernelGlobals kg, } else { isect->type = kernel_data_fetch(objects, isect->object).primitive_type; - isect->u = 1.0f - hit->v - hit->u; - isect->v = hit->u; + isect->u = hit->u; + isect->v = hit->v; } } ccl_device_inline void kernel_embree_convert_sss_hit( KernelGlobals kg, const RTCRay *ray, const RTCHit *hit, Intersection *isect, int object) { - isect->u = 1.0f - hit->v - hit->u; - isect->v = hit->u; + isect->u = hit->u; + isect->v = hit->v; isect->t = ray->tfar; RTCScene inst_scene = (RTCScene)rtcGetGeometryUserData( rtcGetGeometry(kernel_data.device_bvh, object * 2)); diff --git a/intern/cycles/kernel/device/metal/bvh.h b/intern/cycles/kernel/device/metal/bvh.h index f30b21abaf9..03faa3f020f 100644 --- a/intern/cycles/kernel/device/metal/bvh.h +++ b/intern/cycles/kernel/device/metal/bvh.h @@ -129,9 +129,8 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg, isect->t = intersection.distance; if (intersection.type == intersection_type::triangle) { - isect->u = 1.0f - intersection.triangle_barycentric_coord.y - - intersection.triangle_barycentric_coord.x; - isect->v = intersection.triangle_barycentric_coord.x; + isect->u = intersection.triangle_barycentric_coord.x; + isect->v = intersection.triangle_barycentric_coord.y; } else { isect->u = payload.u; @@ -346,9 +345,8 @@ ccl_device_intersect bool scene_intersect_volume(KernelGlobals kg, isect->t = intersection.distance; if (intersection.type == intersection_type::triangle) { - isect->u = 1.0f - intersection.triangle_barycentric_coord.y - - intersection.triangle_barycentric_coord.x; - isect->v = intersection.triangle_barycentric_coord.x; + isect->u = intersection.triangle_barycentric_coord.x; + isect->v = intersection.triangle_barycentric_coord.y; } else { isect->u = payload.u; diff --git a/intern/cycles/kernel/device/metal/kernel.metal b/intern/cycles/kernel/device/metal/kernel.metal index b295e081f3f..3d173b0d601 100644 --- a/intern/cycles/kernel/device/metal/kernel.metal +++ b/intern/cycles/kernel/device/metal/kernel.metal @@ -122,8 +122,8 @@ TReturn metalrt_local_hit(constant KernelParamsMetal &launch_params_metal, isect->object = object; isect->type = kernel_data_fetch(objects, object).primitive_type; - isect->u = 1.0f - barycentrics.y - barycentrics.x; - isect->v = barycentrics.x; + isect->u = barycentrics.x; + isect->v = barycentrics.y; /* Record geometric normal */ const uint tri_vindex = kernel_data_fetch(tri_vindex, isect->prim).w; @@ -185,18 +185,14 @@ bool metalrt_shadow_all_hit(constant KernelParamsMetal &launch_params_metal, return true; } - float u = 0.0f, v = 0.0f; + const float u = barycentrics.x; + const float v = barycentrics.y; int type = 0; if (intersection_type == METALRT_HIT_TRIANGLE) { - u = 1.0f - barycentrics.y - barycentrics.x; - v = barycentrics.x; type = kernel_data_fetch(objects, object).primitive_type; } # ifdef __HAIR__ else { - u = barycentrics.x; - v = barycentrics.y; - const KernelCurveSegment segment = kernel_data_fetch(curve_segments, prim); type = segment.type; prim = segment.prim; diff --git a/intern/cycles/kernel/device/optix/bvh.h b/intern/cycles/kernel/device/optix/bvh.h index 0fb8156c27d..d1d342f6034 100644 --- a/intern/cycles/kernel/device/optix/bvh.h +++ b/intern/cycles/kernel/device/optix/bvh.h @@ -116,8 +116,8 @@ extern "C" __global__ void __anyhit__kernel_optix_local_hit() isect->type = kernel_data_fetch(objects, isect->object).primitive_type; const float2 barycentrics = optixGetTriangleBarycentrics(); - isect->u = 1.0f - barycentrics.y - barycentrics.x; - isect->v = barycentrics.x; + isect->u = barycentrics.x; + isect->v = barycentrics.y; /* Record geometric normal. */ const uint tri_vindex = kernel_data_fetch(tri_vindex, prim).w; @@ -152,8 +152,8 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit() int type = 0; if (optixIsTriangleHit()) { const float2 barycentrics = optixGetTriangleBarycentrics(); - u = 1.0f - barycentrics.y - barycentrics.x; - v = barycentrics.x; + u = barycentrics.x; + v = barycentrics.y; type = kernel_data_fetch(objects, object).primitive_type; } # ifdef __HAIR__ @@ -336,8 +336,8 @@ extern "C" __global__ void __closesthit__kernel_optix_hit() if (optixIsTriangleHit()) { const float2 barycentrics = optixGetTriangleBarycentrics(); - optixSetPayload_1(__float_as_uint(1.0f - barycentrics.y - barycentrics.x)); - optixSetPayload_2(__float_as_uint(barycentrics.x)); + optixSetPayload_1(__float_as_uint(barycentrics.x)); + optixSetPayload_2(__float_as_uint(barycentrics.y)); optixSetPayload_3(prim); optixSetPayload_5(kernel_data_fetch(objects, object).primitive_type); } |