diff options
-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 | ||||
-rw-r--r-- | intern/cycles/kernel/geom/motion_triangle_intersect.h | 4 | ||||
-rw-r--r-- | intern/cycles/kernel/geom/subd_triangle.h | 72 | ||||
-rw-r--r-- | intern/cycles/kernel/geom/triangle.h | 36 | ||||
-rw-r--r-- | intern/cycles/kernel/geom/triangle_intersect.h | 4 | ||||
-rw-r--r-- | intern/cycles/kernel/integrator/init_from_bake.h | 5 | ||||
-rw-r--r-- | intern/cycles/kernel/integrator/mnee.h | 14 | ||||
-rw-r--r-- | intern/cycles/kernel/light/sample.h | 5 | ||||
-rw-r--r-- | intern/cycles/kernel/osl/shaders/node_geometry.osl | 2 | ||||
-rw-r--r-- | intern/cycles/kernel/svm/geometry.h | 6 | ||||
-rw-r--r-- | intern/cycles/scene/mesh_displace.cpp | 8 | ||||
-rw-r--r-- | intern/cycles/util/math_intersect.h | 20 |
15 files changed, 110 insertions, 108 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); } diff --git a/intern/cycles/kernel/geom/motion_triangle_intersect.h b/intern/cycles/kernel/geom/motion_triangle_intersect.h index b59c5c43c20..b30ee7258dc 100644 --- a/intern/cycles/kernel/geom/motion_triangle_intersect.h +++ b/intern/cycles/kernel/geom/motion_triangle_intersect.h @@ -27,8 +27,8 @@ ccl_device_inline float3 motion_triangle_point_from_uv(KernelGlobals kg, const float v, float3 verts[3]) { - float w = 1.0f - u - v; - float3 P = u * verts[0] + v * verts[1] + w * verts[2]; + /* This appears to give slightly better precision than interpolating with w = (1 - u - v). */ + float3 P = verts[0] + u * (verts[1] - verts[0]) + v * (verts[2] - verts[0]); if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) { const Transform tfm = object_get_transform(kg, sd); diff --git a/intern/cycles/kernel/geom/subd_triangle.h b/intern/cycles/kernel/geom/subd_triangle.h index 8b73b342e16..c6f883461bd 100644 --- a/intern/cycles/kernel/geom/subd_triangle.h +++ b/intern/cycles/kernel/geom/subd_triangle.h @@ -94,11 +94,11 @@ ccl_device_noinline float subd_triangle_attribute_float(KernelGlobals kg, float2 uv[3]; subd_triangle_patch_uv(kg, sd, uv); - float2 dpdu = uv[0] - uv[2]; - float2 dpdv = uv[1] - uv[2]; + float2 dpdu = uv[1] - uv[0]; + float2 dpdv = uv[2] - uv[0]; /* p is [s, t] */ - float2 p = dpdu * sd->u + dpdv * sd->v + uv[2]; + float2 p = dpdu * sd->u + dpdv * sd->v + uv[0]; float a, dads, dadt; a = patch_eval_float(kg, sd, desc.offset, patch, p.x, p.y, 0, &dads, &dadt); @@ -165,12 +165,12 @@ ccl_device_noinline float subd_triangle_attribute_float(KernelGlobals kg, #ifdef __RAY_DIFFERENTIALS__ if (dx) - *dx = sd->du.dx * a + sd->dv.dx * b - (sd->du.dx + sd->dv.dx) * c; + *dx = sd->du.dx * b + sd->dv.dx * c - (sd->du.dx + sd->dv.dx) * a; if (dy) - *dy = sd->du.dy * a + sd->dv.dy * b - (sd->du.dy + sd->dv.dy) * c; + *dy = sd->du.dy * b + sd->dv.dy * c - (sd->du.dy + sd->dv.dy) * a; #endif - return sd->u * a + sd->v * b + (1.0f - sd->u - sd->v) * c; + return sd->u * b + sd->v * c + (1.0f - sd->u - sd->v) * a; } else if (desc.element == ATTR_ELEMENT_CORNER) { float2 uv[3]; @@ -195,12 +195,12 @@ ccl_device_noinline float subd_triangle_attribute_float(KernelGlobals kg, #ifdef __RAY_DIFFERENTIALS__ if (dx) - *dx = sd->du.dx * a + sd->dv.dx * b - (sd->du.dx + sd->dv.dx) * c; + *dx = sd->du.dx * b + sd->dv.dx * c - (sd->du.dx + sd->dv.dx) * a; if (dy) - *dy = sd->du.dy * a + sd->dv.dy * b - (sd->du.dy + sd->dv.dy) * c; + *dy = sd->du.dy * b + sd->dv.dy * c - (sd->du.dy + sd->dv.dy) * a; #endif - return sd->u * a + sd->v * b + (1.0f - sd->u - sd->v) * c; + return sd->u * b + sd->v * c + (1.0f - sd->u - sd->v) * a; } else if (desc.element == ATTR_ELEMENT_OBJECT || desc.element == ATTR_ELEMENT_MESH) { if (dx) @@ -233,11 +233,11 @@ ccl_device_noinline float2 subd_triangle_attribute_float2(KernelGlobals kg, float2 uv[3]; subd_triangle_patch_uv(kg, sd, uv); - float2 dpdu = uv[0] - uv[2]; - float2 dpdv = uv[1] - uv[2]; + float2 dpdu = uv[1] - uv[0]; + float2 dpdv = uv[2] - uv[0]; /* p is [s, t] */ - float2 p = dpdu * sd->u + dpdv * sd->v + uv[2]; + float2 p = dpdu * sd->u + dpdv * sd->v + uv[0]; float2 a, dads, dadt; @@ -305,12 +305,12 @@ ccl_device_noinline float2 subd_triangle_attribute_float2(KernelGlobals kg, #ifdef __RAY_DIFFERENTIALS__ if (dx) - *dx = sd->du.dx * a + sd->dv.dx * b - (sd->du.dx + sd->dv.dx) * c; + *dx = sd->du.dx * b + sd->dv.dx * c - (sd->du.dx + sd->dv.dx) * a; if (dy) - *dy = sd->du.dy * a + sd->dv.dy * b - (sd->du.dy + sd->dv.dy) * c; + *dy = sd->du.dy * b + sd->dv.dy * c - (sd->du.dy + sd->dv.dy) * a; #endif - return sd->u * a + sd->v * b + (1.0f - sd->u - sd->v) * c; + return sd->u * b + sd->v * c + (1.0f - sd->u - sd->v) * a; } else if (desc.element == ATTR_ELEMENT_CORNER) { float2 uv[3]; @@ -337,12 +337,12 @@ ccl_device_noinline float2 subd_triangle_attribute_float2(KernelGlobals kg, #ifdef __RAY_DIFFERENTIALS__ if (dx) - *dx = sd->du.dx * a + sd->dv.dx * b - (sd->du.dx + sd->dv.dx) * c; + *dx = sd->du.dx * b + sd->dv.dx * c - (sd->du.dx + sd->dv.dx) * a; if (dy) - *dy = sd->du.dy * a + sd->dv.dy * b - (sd->du.dy + sd->dv.dy) * c; + *dy = sd->du.dy * b + sd->dv.dy * c - (sd->du.dy + sd->dv.dy) * a; #endif - return sd->u * a + sd->v * b + (1.0f - sd->u - sd->v) * c; + return sd->u * b + sd->v * c + (1.0f - sd->u - sd->v) * a; } else if (desc.element == ATTR_ELEMENT_OBJECT || desc.element == ATTR_ELEMENT_MESH) { if (dx) @@ -375,11 +375,11 @@ ccl_device_noinline float3 subd_triangle_attribute_float3(KernelGlobals kg, float2 uv[3]; subd_triangle_patch_uv(kg, sd, uv); - float2 dpdu = uv[0] - uv[2]; - float2 dpdv = uv[1] - uv[2]; + float2 dpdu = uv[1] - uv[0]; + float2 dpdv = uv[2] - uv[0]; /* p is [s, t] */ - float2 p = dpdu * sd->u + dpdv * sd->v + uv[2]; + float2 p = dpdu * sd->u + dpdv * sd->v + uv[0]; float3 a, dads, dadt; a = patch_eval_float3(kg, sd, desc.offset, patch, p.x, p.y, 0, &dads, &dadt); @@ -446,12 +446,12 @@ ccl_device_noinline float3 subd_triangle_attribute_float3(KernelGlobals kg, #ifdef __RAY_DIFFERENTIALS__ if (dx) - *dx = sd->du.dx * a + sd->dv.dx * b - (sd->du.dx + sd->dv.dx) * c; + *dx = sd->du.dx * b + sd->dv.dx * c - (sd->du.dx + sd->dv.dx) * a; if (dy) - *dy = sd->du.dy * a + sd->dv.dy * b - (sd->du.dy + sd->dv.dy) * c; + *dy = sd->du.dy * b + sd->dv.dy * c - (sd->du.dy + sd->dv.dy) * a; #endif - return sd->u * a + sd->v * b + (1.0f - sd->u - sd->v) * c; + return sd->u * b + sd->v * c + (1.0f - sd->u - sd->v) * a; } else if (desc.element == ATTR_ELEMENT_CORNER) { float2 uv[3]; @@ -478,12 +478,12 @@ ccl_device_noinline float3 subd_triangle_attribute_float3(KernelGlobals kg, #ifdef __RAY_DIFFERENTIALS__ if (dx) - *dx = sd->du.dx * a + sd->dv.dx * b - (sd->du.dx + sd->dv.dx) * c; + *dx = sd->du.dx * b + sd->dv.dx * c - (sd->du.dx + sd->dv.dx) * a; if (dy) - *dy = sd->du.dy * a + sd->dv.dy * b - (sd->du.dy + sd->dv.dy) * c; + *dy = sd->du.dy * b + sd->dv.dy * c - (sd->du.dy + sd->dv.dy) * a; #endif - return sd->u * a + sd->v * b + (1.0f - sd->u - sd->v) * c; + return sd->u * b + sd->v * c + (1.0f - sd->u - sd->v) * a; } else if (desc.element == ATTR_ELEMENT_OBJECT || desc.element == ATTR_ELEMENT_MESH) { if (dx) @@ -516,11 +516,11 @@ ccl_device_noinline float4 subd_triangle_attribute_float4(KernelGlobals kg, float2 uv[3]; subd_triangle_patch_uv(kg, sd, uv); - float2 dpdu = uv[0] - uv[2]; - float2 dpdv = uv[1] - uv[2]; + float2 dpdu = uv[1] - uv[0]; + float2 dpdv = uv[2] - uv[0]; /* p is [s, t] */ - float2 p = dpdu * sd->u + dpdv * sd->v + uv[2]; + float2 p = dpdu * sd->u + dpdv * sd->v + uv[0]; float4 a, dads, dadt; if (desc.type == NODE_ATTR_RGBA) { @@ -592,12 +592,12 @@ ccl_device_noinline float4 subd_triangle_attribute_float4(KernelGlobals kg, #ifdef __RAY_DIFFERENTIALS__ if (dx) - *dx = sd->du.dx * a + sd->dv.dx * b - (sd->du.dx + sd->dv.dx) * c; + *dx = sd->du.dx * b + sd->dv.dx * c - (sd->du.dx + sd->dv.dx) * a; if (dy) - *dy = sd->du.dy * a + sd->dv.dy * b - (sd->du.dy + sd->dv.dy) * c; + *dy = sd->du.dy * b + sd->dv.dy * c - (sd->du.dy + sd->dv.dy) * a; #endif - return sd->u * a + sd->v * b + (1.0f - sd->u - sd->v) * c; + return sd->u * b + sd->v * c + (1.0f - sd->u - sd->v) * a; } else if (desc.element == ATTR_ELEMENT_CORNER || desc.element == ATTR_ELEMENT_CORNER_BYTE) { float2 uv[3]; @@ -636,12 +636,12 @@ ccl_device_noinline float4 subd_triangle_attribute_float4(KernelGlobals kg, #ifdef __RAY_DIFFERENTIALS__ if (dx) - *dx = sd->du.dx * a + sd->dv.dx * b - (sd->du.dx + sd->dv.dx) * c; + *dx = sd->du.dx * b + sd->dv.dx * c - (sd->du.dx + sd->dv.dx) * a; if (dy) - *dy = sd->du.dy * a + sd->dv.dy * b - (sd->du.dy + sd->dv.dy) * c; + *dy = sd->du.dy * b + sd->dv.dy * c - (sd->du.dy + sd->dv.dy) * a; #endif - return sd->u * a + sd->v * b + (1.0f - sd->u - sd->v) * c; + return sd->u * b + sd->v * c + (1.0f - sd->u - sd->v) * a; } else if (desc.element == ATTR_ELEMENT_OBJECT || desc.element == ATTR_ELEMENT_MESH) { if (dx) diff --git a/intern/cycles/kernel/geom/triangle.h b/intern/cycles/kernel/geom/triangle.h index 788bfaca7cf..6b9450d59ef 100644 --- a/intern/cycles/kernel/geom/triangle.h +++ b/intern/cycles/kernel/geom/triangle.h @@ -45,8 +45,8 @@ ccl_device_inline void triangle_point_normal(KernelGlobals kg, float3 v1 = kernel_data_fetch(tri_verts, tri_vindex.w + 1); float3 v2 = kernel_data_fetch(tri_verts, tri_vindex.w + 2); /* compute point */ - float t = 1.0f - u - v; - *P = (u * v0 + v * v1 + t * v2); + float w = 1.0f - u - v; + *P = (w * v0 + u * v1 + v * v2); /* get object flags */ int object_flag = kernel_data_fetch(object_flag, object); /* compute normal */ @@ -97,7 +97,7 @@ triangle_smooth_normal(KernelGlobals kg, float3 Ng, int prim, float u, float v) float3 n1 = kernel_data_fetch(tri_vnormal, tri_vindex.y); float3 n2 = kernel_data_fetch(tri_vnormal, tri_vindex.z); - float3 N = safe_normalize((1.0f - u - v) * n2 + u * n0 + v * n1); + float3 N = safe_normalize((1.0f - u - v) * n0 + u * n1 + v * n2); return is_zero(N) ? Ng : N; } @@ -118,7 +118,7 @@ ccl_device_inline float3 triangle_smooth_normal_unnormalized( object_inverse_normal_transform(kg, sd, &n2); } - float3 N = (1.0f - u - v) * n2 + u * n0 + v * n1; + float3 N = (1.0f - u - v) * n0 + u * n1 + v * n2; return is_zero(N) ? Ng : N; } @@ -137,8 +137,8 @@ ccl_device_inline void triangle_dPdudv(KernelGlobals kg, const float3 p2 = kernel_data_fetch(tri_verts, tri_vindex.w + 2); /* compute derivatives of P w.r.t. uv */ - *dPdu = (p0 - p2); - *dPdv = (p1 - p2); + *dPdu = (p1 - p0); + *dPdv = (p2 - p0); } /* Reading attributes on various triangle elements */ @@ -167,12 +167,12 @@ ccl_device float triangle_attribute_float(KernelGlobals kg, #ifdef __RAY_DIFFERENTIALS__ if (dx) - *dx = sd->du.dx * f0 + sd->dv.dx * f1 - (sd->du.dx + sd->dv.dx) * f2; + *dx = sd->du.dx * f1 + sd->dv.dx * f2 - (sd->du.dx + sd->dv.dx) * f0; if (dy) - *dy = sd->du.dy * f0 + sd->dv.dy * f1 - (sd->du.dy + sd->dv.dy) * f2; + *dy = sd->du.dy * f1 + sd->dv.dy * f2 - (sd->du.dy + sd->dv.dy) * f0; #endif - return sd->u * f0 + sd->v * f1 + (1.0f - sd->u - sd->v) * f2; + return sd->u * f1 + sd->v * f2 + (1.0f - sd->u - sd->v) * f0; } else { #ifdef __RAY_DIFFERENTIALS__ @@ -217,12 +217,12 @@ ccl_device float2 triangle_attribute_float2(KernelGlobals kg, #ifdef __RAY_DIFFERENTIALS__ if (dx) - *dx = sd->du.dx * f0 + sd->dv.dx * f1 - (sd->du.dx + sd->dv.dx) * f2; + *dx = sd->du.dx * f1 + sd->dv.dx * f2 - (sd->du.dx + sd->dv.dx) * f0; if (dy) - *dy = sd->du.dy * f0 + sd->dv.dy * f1 - (sd->du.dy + sd->dv.dy) * f2; + *dy = sd->du.dy * f1 + sd->dv.dy * f2 - (sd->du.dy + sd->dv.dy) * f0; #endif - return sd->u * f0 + sd->v * f1 + (1.0f - sd->u - sd->v) * f2; + return sd->u * f1 + sd->v * f2 + (1.0f - sd->u - sd->v) * f0; } else { #ifdef __RAY_DIFFERENTIALS__ @@ -267,12 +267,12 @@ ccl_device float3 triangle_attribute_float3(KernelGlobals kg, #ifdef __RAY_DIFFERENTIALS__ if (dx) - *dx = sd->du.dx * f0 + sd->dv.dx * f1 - (sd->du.dx + sd->dv.dx) * f2; + *dx = sd->du.dx * f1 + sd->dv.dx * f2 - (sd->du.dx + sd->dv.dx) * f0; if (dy) - *dy = sd->du.dy * f0 + sd->dv.dy * f1 - (sd->du.dy + sd->dv.dy) * f2; + *dy = sd->du.dy * f1 + sd->dv.dy * f2 - (sd->du.dy + sd->dv.dy) * f0; #endif - return sd->u * f0 + sd->v * f1 + (1.0f - sd->u - sd->v) * f2; + return sd->u * f1 + sd->v * f2 + (1.0f - sd->u - sd->v) * f0; } else { #ifdef __RAY_DIFFERENTIALS__ @@ -328,12 +328,12 @@ ccl_device float4 triangle_attribute_float4(KernelGlobals kg, #ifdef __RAY_DIFFERENTIALS__ if (dx) - *dx = sd->du.dx * f0 + sd->dv.dx * f1 - (sd->du.dx + sd->dv.dx) * f2; + *dx = sd->du.dx * f1 + sd->dv.dx * f2 - (sd->du.dx + sd->dv.dx) * f0; if (dy) - *dy = sd->du.dy * f0 + sd->dv.dy * f1 - (sd->du.dy + sd->dv.dy) * f2; + *dy = sd->du.dy * f1 + sd->dv.dy * f2 - (sd->du.dy + sd->dv.dy) * f0; #endif - return sd->u * f0 + sd->v * f1 + (1.0f - sd->u - sd->v) * f2; + return sd->u * f1 + sd->v * f2 + (1.0f - sd->u - sd->v) * f0; } else { #ifdef __RAY_DIFFERENTIALS__ diff --git a/intern/cycles/kernel/geom/triangle_intersect.h b/intern/cycles/kernel/geom/triangle_intersect.h index f968e537cfa..847ed22fddd 100644 --- a/intern/cycles/kernel/geom/triangle_intersect.h +++ b/intern/cycles/kernel/geom/triangle_intersect.h @@ -145,9 +145,9 @@ ccl_device_inline float3 triangle_point_from_uv(KernelGlobals kg, const packed_float3 tri_a = kernel_data_fetch(tri_verts, tri_vindex + 0), tri_b = kernel_data_fetch(tri_verts, tri_vindex + 1), tri_c = kernel_data_fetch(tri_verts, tri_vindex + 2); - float w = 1.0f - u - v; - float3 P = u * tri_a + v * tri_b + w * tri_c; + /* This appears to give slightly better precision than interpolating with w = (1 - u - v). */ + float3 P = tri_a + u * (tri_b - tri_a) + v * (tri_c - tri_a); if (!(sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) { const Transform tfm = object_get_transform(kg, sd); diff --git a/intern/cycles/kernel/integrator/init_from_bake.h b/intern/cycles/kernel/integrator/init_from_bake.h index bf3f41b52b9..dd26215bcd2 100644 --- a/intern/cycles/kernel/integrator/init_from_bake.h +++ b/intern/cycles/kernel/integrator/init_from_bake.h @@ -155,6 +155,11 @@ ccl_device bool integrator_init_from_bake(KernelGlobals kg, 1.0f - u); } + /* Convert from Blender to Cycles/Embree/OptiX barycentric convention. */ + const float tmp = u; + u = v; + v = 1.0f - tmp - v; + /* Position and normal on triangle. */ const int object = kernel_data.bake.object_index; float3 P, Ng; diff --git a/intern/cycles/kernel/integrator/mnee.h b/intern/cycles/kernel/integrator/mnee.h index 7a6f866b1a0..aa72b93c9d2 100644 --- a/intern/cycles/kernel/integrator/mnee.h +++ b/intern/cycles/kernel/integrator/mnee.h @@ -186,7 +186,7 @@ ccl_device_forceinline void mnee_setup_manifold_vertex(KernelGlobals kg, triangle_vertices_and_normals(kg, sd_vtx->prim, verts, normals); /* Compute refined position (same code as in triangle_point_from_uv). */ - sd_vtx->P = isect->u * verts[0] + isect->v * verts[1] + (1.f - isect->u - isect->v) * verts[2]; + sd_vtx->P = (1.f - isect->u - isect->v) * verts[0] + isect->u * verts[1] + isect->v * verts[2]; if (!(sd_vtx->object_flag & SD_OBJECT_TRANSFORM_APPLIED)) { const Transform tfm = object_get_transform(kg, sd_vtx); sd_vtx->P = transform_point(&tfm, sd_vtx->P); @@ -213,8 +213,8 @@ ccl_device_forceinline void mnee_setup_manifold_vertex(KernelGlobals kg, } /* Tangent space (position derivatives) WRT barycentric (u, v). */ - float3 dp_du = verts[0] - verts[2]; - float3 dp_dv = verts[1] - verts[2]; + float3 dp_du = verts[1] - verts[0]; + float3 dp_dv = verts[2] - verts[0]; /* Geometric normal. */ vtx->ng = normalize(cross(dp_du, dp_dv)); @@ -223,16 +223,16 @@ ccl_device_forceinline void mnee_setup_manifold_vertex(KernelGlobals kg, /* Shading normals: Interpolate normals between vertices. */ float n_len; - vtx->n = normalize_len(normals[0] * sd_vtx->u + normals[1] * sd_vtx->v + - normals[2] * (1.0f - sd_vtx->u - sd_vtx->v), + vtx->n = normalize_len(normals[0] * (1.0f - sd_vtx->u - sd_vtx->v) + normals[1] * sd_vtx->u + + normals[2] * sd_vtx->v, &n_len); /* Shading normal derivatives WRT barycentric (u, v) * we calculate the derivative of n = |u*n0 + v*n1 + (1-u-v)*n2| using: * d/du [f(u)/|f(u)|] = [d/du f(u)]/|f(u)| - f(u)/|f(u)|^3 <f(u), d/du f(u)>. */ const float inv_n_len = 1.f / n_len; - float3 dn_du = inv_n_len * (normals[0] - normals[2]); - float3 dn_dv = inv_n_len * (normals[1] - normals[2]); + float3 dn_du = inv_n_len * (normals[1] - normals[0]); + float3 dn_dv = inv_n_len * (normals[2] - normals[0]); dn_du -= vtx->n * dot(vtx->n, dn_du); dn_dv -= vtx->n * dot(vtx->n, dn_dv); diff --git a/intern/cycles/kernel/light/sample.h b/intern/cycles/kernel/light/sample.h index 210bb1b35c2..2e309cee43f 100644 --- a/intern/cycles/kernel/light/sample.h +++ b/intern/cycles/kernel/light/sample.h @@ -137,8 +137,9 @@ ccl_device_inline float3 shadow_ray_smooth_surface_offset( triangle_vertices_and_normals(kg, sd->prim, V, N); } - const float u = sd->u, v = sd->v; - const float w = 1 - u - v; + const float u = 1.0f - sd->u - sd->v; + const float v = sd->u; + const float w = sd->v; float3 P = V[0] * u + V[1] * v + V[2] * w; /* Local space */ float3 n = N[0] * u + N[1] * v + N[2] * w; /* We get away without normalization */ diff --git a/intern/cycles/kernel/osl/shaders/node_geometry.osl b/intern/cycles/kernel/osl/shaders/node_geometry.osl index 23d4c2ee66f..cc891abd6e3 100644 --- a/intern/cycles/kernel/osl/shaders/node_geometry.osl +++ b/intern/cycles/kernel/osl/shaders/node_geometry.osl @@ -20,7 +20,7 @@ shader node_geometry(normal NormalIn = N, Normal = NormalIn; TrueNormal = Ng; Incoming = I; - Parametric = point(u, v, 0.0); + Parametric = point(1.0 - u - v, u, 0.0); Backfacing = backfacing(); if (bump_offset == "dx") { diff --git a/intern/cycles/kernel/svm/geometry.h b/intern/cycles/kernel/svm/geometry.h index 4b5368dd765..bbefdcfa755 100644 --- a/intern/cycles/kernel/svm/geometry.h +++ b/intern/cycles/kernel/svm/geometry.h @@ -34,7 +34,7 @@ ccl_device_noinline void svm_node_geometry(KernelGlobals kg, data = sd->Ng; break; case NODE_GEOM_uv: - data = make_float3(sd->u, sd->v, 0.0f); + data = make_float3(1.0f - sd->u - sd->v, sd->u, 0.0f); break; default: data = make_float3(0.0f, 0.0f, 0.0f); @@ -57,7 +57,7 @@ ccl_device_noinline void svm_node_geometry_bump_dx(KernelGlobals kg, data = sd->P + sd->dP.dx; break; case NODE_GEOM_uv: - data = make_float3(sd->u + sd->du.dx, sd->v + sd->dv.dx, 0.0f); + data = make_float3(1.0f - sd->u - sd->du.dx - sd->v - sd->dv.dx, sd->u + sd->du.dx, 0.0f); break; default: svm_node_geometry(kg, sd, stack, type, out_offset); @@ -84,7 +84,7 @@ ccl_device_noinline void svm_node_geometry_bump_dy(KernelGlobals kg, data = sd->P + sd->dP.dy; break; case NODE_GEOM_uv: - data = make_float3(sd->u + sd->du.dy, sd->v + sd->dv.dy, 0.0f); + data = make_float3(1.0f - sd->u - sd->du.dy - sd->v - sd->dv.dy, sd->u + sd->du.dy, 0.0f); break; default: svm_node_geometry(kg, sd, stack, type, out_offset); diff --git a/intern/cycles/scene/mesh_displace.cpp b/intern/cycles/scene/mesh_displace.cpp index e180145daac..cdd5a793530 100644 --- a/intern/cycles/scene/mesh_displace.cpp +++ b/intern/cycles/scene/mesh_displace.cpp @@ -73,16 +73,16 @@ static int fill_shader_input(const Scene *scene, switch (j) { case 0: - u = 1.0f; + u = 0.0f; v = 0.0f; break; case 1: - u = 0.0f; - v = 1.0f; + u = 1.0f; + v = 0.0f; break; default: u = 0.0f; - v = 0.0f; + v = 1.0f; break; } diff --git a/intern/cycles/util/math_intersect.h b/intern/cycles/util/math_intersect.h index 3e5891b2507..37bdc5f1ccf 100644 --- a/intern/cycles/util/math_intersect.h +++ b/intern/cycles/util/math_intersect.h @@ -120,9 +120,9 @@ ccl_device_forceinline bool ray_triangle_intersect(const float3 ray_P, * in Embree. */ /* Calculate vertices relative to ray origin. */ - const float3 v0 = tri_c - ray_P; - const float3 v1 = tri_a - ray_P; - const float3 v2 = tri_b - ray_P; + const float3 v0 = tri_a - ray_P; + const float3 v1 = tri_b - ray_P; + const float3 v2 = tri_c - ray_P; /* Calculate triangle edges. */ const float3 e0 = v2 - v0; @@ -130,11 +130,12 @@ ccl_device_forceinline bool ray_triangle_intersect(const float3 ray_P, const float3 e2 = v1 - v2; /* Perform edge tests. */ - const float U = dot(cross(v2 + v0, e0), ray_D); - const float V = dot(cross(v0 + v1, e1), ray_D); - const float W = dot(cross(v1 + v2, e2), ray_D); + const float U = dot(cross(e0, v2 + v0), ray_D); + const float V = dot(cross(e1, v0 + v1), ray_D); + const float W = dot(cross(e2, v1 + v2), ray_D); - const float eps = FLT_EPSILON * fabsf(U + V + W); + const float UVW = U + V + W; + const float eps = FLT_EPSILON * fabsf(UVW); const float minUVW = min(U, min(V, W)); const float maxUVW = max(U, max(V, W)); @@ -158,8 +159,9 @@ ccl_device_forceinline bool ray_triangle_intersect(const float3 ray_P, return false; } - *isect_u = U / den; - *isect_v = V / den; + const float rcp_UVW = (fabsf(UVW) < 1e-18f) ? 0.0f : 1.0f / UVW; + *isect_u = min(U * rcp_UVW, 1.0f); + *isect_v = min(V * rcp_UVW, 1.0f); *isect_t = t; return true; } |