blender
/
blender
134
401
Fork 584
Code Issues 6.3k Pull Requests 760 Projects 19 Wiki Activity

Merge branch 'blender-v3.3-release'

This commit is contained in:
Brecht Van Lommel 2022-07-27 21:26:18 +02:00
parent 55fb2abc81 38af5b0501
commit 165fa9e2a1
23 changed files with 227 additions and 251 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

26
intern/cycles/kernel/bvh/bvh.h
View File

@ -12,6 +12,7 @@
#if defined(__EMBREE__)
# include "kernel/device/cpu/bvh.h"
# define __BVH2__
#elif defined(__METALRT__)
# include "kernel/device/metal/bvh.h"
#elif defined(__KERNEL_OPTIX__)
@ -72,6 +73,12 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg,
return false;
}
# ifdef __EMBREE__
if (kernel_data.device_bvh) {
return kernel_embree_intersect(kg, ray, visibility, isect);
}
# endif
# ifdef __OBJECT_MOTION__
if (kernel_data.bvh.have_motion) {
# ifdef __HAIR__
@ -121,6 +128,12 @@ ccl_device_intersect bool scene_intersect_local(KernelGlobals kg,
return false;
}
# ifdef __EMBREE__
if (kernel_data.device_bvh) {
return kernel_embree_intersect_local(kg, ray, local_isect, local_object, lcg_state, max_hits);
}
# endif
# ifdef __OBJECT_MOTION__
if (kernel_data.bvh.have_motion) {
return bvh_intersect_local_motion(kg, ray, local_isect, local_object, lcg_state, max_hits);
@ -170,6 +183,13 @@ ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg,
return false;
}
# ifdef __EMBREE__
if (kernel_data.device_bvh) {
return kernel_embree_intersect_shadow_all(
kg, state, ray, visibility, max_hits, num_recorded_hits, throughput);
}
# endif
# ifdef __OBJECT_MOTION__
if (kernel_data.bvh.have_motion) {
# ifdef __HAIR__
@ -254,6 +274,12 @@ ccl_device_intersect uint scene_intersect_volume(KernelGlobals kg,
return false;
}
# ifdef __EMBREE__
if (kernel_data.device_bvh) {
return kernel_embree_intersect_volume(kg, ray, isect, max_hits, visibility);
}
# endif
# ifdef __OBJECT_MOTION__
if (kernel_data.bvh.have_motion) {
return bvh_intersect_volume_all_motion(kg, ray, isect, max_hits, visibility);

6
intern/cycles/kernel/bvh/shadow_all.h
View File

@ -300,11 +300,7 @@ ccl_device_inline
kernel_assert(object != OBJECT_NONE);
/* Instance pop. */
#if BVH_FEATURE(BVH_MOTION)
bvh_instance_motion_pop(kg, object, ray, &P, &dir, &idir);
#else
bvh_instance_pop(kg, object, ray, &P, &dir, &idir);
#endif
bvh_instance_pop(ray, &P, &dir, &idir);
object = OBJECT_NONE;
node_addr = traversal_stack[stack_ptr];

6
intern/cycles/kernel/bvh/traversal.h
View File

@ -237,11 +237,7 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals kg,
kernel_assert(object != OBJECT_NONE);
/* instance pop */
#if BVH_FEATURE(BVH_MOTION)
bvh_instance_motion_pop(kg, object, ray, &P, &dir, &idir);
#else
bvh_instance_pop(kg, object, ray, &P, &dir, &idir);
#endif
bvh_instance_pop(ray, &P, &dir, &idir);
object = OBJECT_NONE;
node_addr = traversal_stack[stack_ptr];

6
intern/cycles/kernel/bvh/volume.h
View File

@ -210,11 +210,7 @@ ccl_device_inline
kernel_assert(object != OBJECT_NONE);
/* instance pop */
#if BVH_FEATURE(BVH_MOTION)
bvh_instance_motion_pop(kg, object, ray, &P, &dir, &idir);
#else
bvh_instance_pop(kg, object, ray, &P, &dir, &idir);
#endif
bvh_instance_pop(ray, &P, &dir, &idir);
object = OBJECT_NONE;
node_addr = traversal_stack[stack_ptr];

6
intern/cycles/kernel/bvh/volume_all.h
View File

@ -242,11 +242,7 @@ ccl_device_inline
kernel_assert(object != OBJECT_NONE);
/* Instance pop. */
#if BVH_FEATURE(BVH_MOTION)
bvh_instance_motion_pop(kg, object, ray, &P, &dir, &idir);
#else
bvh_instance_pop(kg, object, ray, &P, &dir, &idir);
#endif
bvh_instance_pop(ray, &P, &dir, &idir);
object = OBJECT_NONE;
node_addr = traversal_stack[stack_ptr];

89
intern/cycles/kernel/device/cpu/bvh.h
View File

@ -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));
@ -446,19 +446,11 @@ ccl_device void kernel_embree_filter_occluded_func_backface_cull(
/* Scene intersection. */
ccl_device_intersect bool scene_intersect(KernelGlobals kg,
ccl_private const Ray *ray,
const uint visibility,
ccl_private Intersection *isect)
ccl_device_intersect bool kernel_embree_intersect(KernelGlobals kg,
ccl_private const Ray *ray,
const uint visibility,
ccl_private Intersection *isect)
{
if (!intersection_ray_valid(ray)) {
return false;
}
if (!kernel_data.device_bvh) {
return false;
}
isect->t = ray->tmax;
CCLIntersectContext ctx(kg, CCLIntersectContext::RAY_REGULAR);
IntersectContext rtc_ctx(&ctx);
@ -476,24 +468,13 @@ ccl_device_intersect bool scene_intersect(KernelGlobals kg,
}
#ifdef __BVH_LOCAL__
ccl_device_intersect bool scene_intersect_local(KernelGlobals kg,
ccl_private const Ray *ray,
ccl_private LocalIntersection *local_isect,
int local_object,
ccl_private uint *lcg_state,
int max_hits)
ccl_device_intersect bool kernel_embree_intersect_local(KernelGlobals kg,
ccl_private const Ray *ray,
ccl_private LocalIntersection *local_isect,
int local_object,
ccl_private uint *lcg_state,
int max_hits)
{
if (!intersection_ray_valid(ray)) {
if (local_isect) {
local_isect->num_hits = 0;
}
return false;
}
if (!kernel_data.device_bvh) {
return false;
}
const bool has_bvh = !(kernel_data_fetch(object_flag, local_object) &
SD_OBJECT_TRANSFORM_APPLIED);
CCLIntersectContext ctx(kg,
@ -544,24 +525,14 @@ ccl_device_intersect bool scene_intersect_local(KernelGlobals kg,
#endif
#ifdef __SHADOW_RECORD_ALL__
ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg,
IntegratorShadowStateCPU *state,
ccl_private const Ray *ray,
uint visibility,
uint max_hits,
ccl_private uint *num_recorded_hits,
ccl_private float *throughput)
ccl_device_intersect bool kernel_embree_intersect_shadow_all(KernelGlobals kg,
IntegratorShadowStateCPU *state,
ccl_private const Ray *ray,
uint visibility,
uint max_hits,
ccl_private uint *num_recorded_hits,
ccl_private float *throughput)
{
if (!intersection_ray_valid(ray)) {
*num_recorded_hits = 0;
*throughput = 1.0f;
return false;
}
if (!kernel_data.device_bvh) {
return false;
}
CCLIntersectContext ctx(kg, CCLIntersectContext::RAY_SHADOW_ALL);
Intersection *isect_array = (Intersection *)state->shadow_isect;
ctx.isect_s = isect_array;
@ -579,20 +550,12 @@ ccl_device_intersect bool scene_intersect_shadow_all(KernelGlobals kg,
#endif
#ifdef __VOLUME__
ccl_device_intersect uint scene_intersect_volume(KernelGlobals kg,
ccl_private const Ray *ray,
ccl_private Intersection *isect,
const uint max_hits,
const uint visibility)
ccl_device_intersect uint kernel_embree_intersect_volume(KernelGlobals kg,
ccl_private const Ray *ray,
ccl_private Intersection *isect,
const uint max_hits,
const uint visibility)
{
if (!intersection_ray_valid(ray)) {
return false;
}
if (!kernel_data.device_bvh) {
return false;
}
CCLIntersectContext ctx(kg, CCLIntersectContext::RAY_VOLUME_ALL);
ctx.isect_s = isect;
ctx.max_hits = max_hits;

10
intern/cycles/kernel/device/metal/bvh.h
View File

@ -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;

12
intern/cycles/kernel/device/metal/kernel.metal
View File

@ -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;

12
intern/cycles/kernel/device/optix/bvh.h
View File

@ -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);
}

4
intern/cycles/kernel/geom/motion_triangle_intersect.h
View File

@ -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);

30
intern/cycles/kernel/geom/object.h
View File

@ -503,20 +503,6 @@ ccl_device_inline void bvh_instance_push(KernelGlobals kg,
*idir = bvh_inverse_direction(*dir);
}
/* Transform ray to exit static object in BVH. */
ccl_device_inline void bvh_instance_pop(KernelGlobals kg,
int object,
ccl_private const Ray *ray,
ccl_private float3 *P,
ccl_private float3 *dir,
ccl_private float3 *idir)
{
*P = ray->P;
*dir = bvh_clamp_direction(ray->D);
*idir = bvh_inverse_direction(*dir);
}
#ifdef __OBJECT_MOTION__
/* Transform ray into object space to enter motion blurred object in BVH */
@ -536,22 +522,20 @@ ccl_device_inline void bvh_instance_motion_push(KernelGlobals kg,
*idir = bvh_inverse_direction(*dir);
}
/* Transform ray to exit motion blurred object in BVH. */
#endif
ccl_device_inline void bvh_instance_motion_pop(KernelGlobals kg,
int object,
ccl_private const Ray *ray,
ccl_private float3 *P,
ccl_private float3 *dir,
ccl_private float3 *idir)
/* Transform ray to exit static object in BVH. */
ccl_device_inline void bvh_instance_pop(ccl_private const Ray *ray,
ccl_private float3 *P,
ccl_private float3 *dir,
ccl_private float3 *idir)
{
*P = ray->P;
*dir = bvh_clamp_direction(ray->D);
*idir = bvh_inverse_direction(*dir);
}
#endif
/* TODO: This can be removed when we know if no devices will require explicit
* address space qualifiers for this case. */

72
intern/cycles/kernel/geom/subd_triangle.h
View File

@ -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)

36
intern/cycles/kernel/geom/triangle.h
View File

@ -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__

4
intern/cycles/kernel/geom/triangle_intersect.h
View File

@ -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);

5
intern/cycles/kernel/integrator/init_from_bake.h
View File

@ -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;

14
intern/cycles/kernel/integrator/mnee.h
View File

@ -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);

5
intern/cycles/kernel/light/sample.h
View File

@ -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 */

2
intern/cycles/kernel/osl/shaders/node_geometry.osl
View File

@ -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") {

6
intern/cycles/kernel/svm/geometry.h
View File

@ -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);

8
intern/cycles/scene/mesh_displace.cpp
View File

@ -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;
}

20
intern/cycles/util/math_intersect.h
View File

@ -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;
}

98
source/blender/blenkernel/intern/attribute_access.cc
View File

@ -214,36 +214,33 @@ static bool add_custom_data_layer_from_attribute_init(const AttributeIDRef &attr
const int domain_num,
const AttributeInit &initializer)
{
const int old_layer_num = custom_data.totlayer;
switch (initializer.type) {
case AttributeInit::Type::Default: {
void *data = add_generic_custom_data_layer(
add_generic_custom_data_layer(
custom_data, data_type, CD_DEFAULT, nullptr, domain_num, attribute_id);
return data != nullptr;
break;
}
case AttributeInit::Type::VArray: {
void *data = add_generic_custom_data_layer(
custom_data, data_type, CD_DEFAULT, nullptr, domain_num, attribute_id);
if (data == nullptr) {
return false;
if (data != nullptr) {
const GVArray &varray = static_cast<const AttributeInitVArray &>(initializer).varray;
varray.materialize_to_uninitialized(varray.index_range(), data);
}
const GVArray &varray = static_cast<const AttributeInitVArray &>(initializer).varray;
varray.materialize_to_uninitialized(varray.index_range(), data);
return true;
break;
}
case AttributeInit::Type::MoveArray: {
void *source_data = static_cast<const AttributeInitMove &>(initializer).data;
void *data = add_generic_custom_data_layer(
custom_data, data_type, CD_ASSIGN, source_data, domain_num, attribute_id);
if (data == nullptr) {
if (source_data != nullptr && data == nullptr) {
MEM_freeN(source_data);
return false;
}
return true;
break;
}
}
BLI_assert_unreachable();
return false;
return old_layer_num < custom_data.totlayer;
}
static bool custom_data_layer_matches_attribute_id(const CustomDataLayer &layer,
@ -265,17 +262,25 @@ GVArray BuiltinCustomDataLayerProvider::try_get_for_read(const void *owner) cons
return {};
}
const void *data;
if (stored_as_named_attribute_) {
data = CustomData_get_layer_named(custom_data, stored_type_, name_.c_str());
const void *data = nullptr;
bool found_attribute = false;
for (const CustomDataLayer &layer : Span(custom_data->layers, custom_data->totlayer)) {
if (stored_as_named_attribute_) {
if (layer.name == name_) {
data = layer.data;
found_attribute = true;
break;
}
}
else if (layer.type == stored_type_) {
data = layer.data;
found_attribute = true;
break;
}
}
else {
data = CustomData_get_layer(custom_data, stored_type_);
}
if (data == nullptr) {
if (!found_attribute) {
return {};
}
const int element_num = custom_data_access_.get_element_num(owner);
return as_read_attribute_(data, element_num);
}
@ -291,31 +296,42 @@ GAttributeWriter BuiltinCustomDataLayerProvider::try_get_for_write(void *owner)
}
const int element_num = custom_data_access_.get_element_num(owner);
void *data;
if (stored_as_named_attribute_) {
data = CustomData_get_layer_named(custom_data, stored_type_, name_.c_str());
void *data = nullptr;
bool found_attribute = false;
for (const CustomDataLayer &layer : Span(custom_data->layers, custom_data->totlayer)) {
if (stored_as_named_attribute_) {
if (layer.name == name_) {
data = layer.data;
found_attribute = true;
break;
}
}
else if (layer.type == stored_type_) {
data = layer.data;
found_attribute = true;
break;
}
}
else {
data = CustomData_get_layer(custom_data, stored_type_);
}
if (data == nullptr) {
if (!found_attribute) {
return {};
}
void *new_data;
if (stored_as_named_attribute_) {
new_data = CustomData_duplicate_referenced_layer_named(
custom_data, stored_type_, name_.c_str(), element_num);
}
else {
new_data = CustomData_duplicate_referenced_layer(custom_data, stored_type_, element_num);
}
if (data != new_data) {
if (custom_data_access_.update_custom_data_pointers) {
custom_data_access_.update_custom_data_pointers(owner);
if (data != nullptr) {
void *new_data;
if (stored_as_named_attribute_) {
new_data = CustomData_duplicate_referenced_layer_named(
custom_data, stored_type_, name_.c_str(), element_num);
}
else {
new_data = CustomData_duplicate_referenced_layer(custom_data, stored_type_, element_num);
}
if (data != new_data) {
if (custom_data_access_.update_custom_data_pointers) {
custom_data_access_.update_custom_data_pointers(owner);
}
data = new_data;
}
data = new_data;
}
std::function<void()> tag_modified_fn;

1
source/blender/nodes/geometry/nodes/node_geo_join_geometry.cc
View File

@ -185,6 +185,7 @@ static void node_geo_exec(GeoNodeExecParams params)
join_component_type<InstancesComponent>(geometry_sets, geometry_set_result);
join_component_type<VolumeComponent>(geometry_sets, geometry_set_result);
join_component_type<CurveComponent>(geometry_sets, geometry_set_result);
join_component_type<GeometryComponentEditData>(geometry_sets, geometry_set_result);
params.set_output("Geometry", std::move(geometry_set_result));
}