[Bf-blender-cvs] [38af5b05010] blender-v3.3-release: Cycles: switch Cycles triangle barycentric convention to match Embree/OptiX
Brecht Van Lommel
noreply at git.blender.org
Wed Jul 27 21:25:39 CEST 2022
Commit: 38af5b0501005d8ee84a59f027417bf6a31fcc5e
Author: Brecht Van Lommel
Date: Tue Jul 26 16:07:50 2022 +0200
Branches: blender-v3.3-release
https://developer.blender.org/rB38af5b0501005d8ee84a59f027417bf6a31fcc5e
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.
===================================================================
M intern/cycles/kernel/device/cpu/bvh.h
M intern/cycles/kernel/device/metal/bvh.h
M intern/cycles/kernel/device/metal/kernel.metal
M intern/cycles/kernel/device/optix/bvh.h
M intern/cycles/kernel/geom/motion_triangle_intersect.h
M intern/cycles/kernel/geom/subd_triangle.h
M intern/cycles/kernel/geom/triangle.h
M intern/cycles/kernel/geom/triangle_intersect.h
M intern/cycles/kernel/integrator/init_from_bake.h
M intern/cycles/kernel/integrator/mnee.h
M intern/cycles/kernel/light/sample.h
M intern/cycles/kernel/osl/shaders/node_geometry.osl
M intern/cycles/kernel/svm/geometry.h
M intern/cycles/scene/mesh_displace.cpp
M intern/cycles/util/math_intersect.h
===================================================================
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 =
@@ Diff output truncated at 10240 characters. @@
More information about the Bf-blender-cvs
mailing list