Merge branch 'master' into blender2.8

This commit is contained in:
Julian Eisel 2017-03-28 23:11:10 +02:00
commit ff3e1fa760
17 changed files with 142 additions and 324 deletions

View File

@ -183,11 +183,6 @@ class CyclesRender_PT_sampling(CyclesButtonsPanel, Panel):
else:
sub.label(text="AA Samples:")
sub.prop(cscene, "aa_samples", text="Render")
sub.prop(cscene, "preview_aa_samples", text="Preview")
sub.separator()
sub.prop(cscene, "sample_all_lights_direct")
sub.prop(cscene, "sample_all_lights_indirect")
col = split.column()
sub = col.column(align=True)
sub.label(text="Samples:")
@ -203,6 +198,10 @@ class CyclesRender_PT_sampling(CyclesButtonsPanel, Panel):
sub.prop(cscene, "subsurface_samples", text="Subsurface")
sub.prop(cscene, "volume_samples", text="Volume")
col = layout.column(align=True)
col.prop(cscene, "sample_all_lights_direct")
col.prop(cscene, "sample_all_lights_indirect")
if not (use_opencl(context) and cscene.feature_set != 'EXPERIMENTAL'):
layout.row().prop(cscene, "sampling_pattern", text="Pattern")

View File

@ -101,9 +101,6 @@ bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals *kg,
gen_idirsplat_swap(pn, shuf_identity, shuf_swap, idir, idirsplat, shufflexyz);
#endif /* __KERNEL_SSE2__ */
TriangleIsectPrecalc isect_precalc;
ray_triangle_intersect_precalc(dir, &isect_precalc);
/* traversal loop */
do {
do {
@ -209,9 +206,9 @@ bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals *kg,
switch(p_type) {
case PRIMITIVE_TRIANGLE: {
hit = triangle_intersect(kg,
&isect_precalc,
isect_array,
P,
dir,
PATH_RAY_SHADOW,
object,
prim_addr);
@ -220,9 +217,9 @@ bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals *kg,
#if BVH_FEATURE(BVH_MOTION)
case PRIMITIVE_MOTION_TRIANGLE: {
hit = motion_triangle_intersect(kg,
&isect_precalc,
isect_array,
P,
dir,
ray->time,
PATH_RAY_SHADOW,
object,
@ -325,7 +322,6 @@ bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals *kg,
isect_t = bvh_instance_push(kg, object, ray, &P, &dir, &idir, isect_t);
# endif
ray_triangle_intersect_precalc(dir, &isect_precalc);
num_hits_in_instance = 0;
isect_array->t = isect_t;
@ -365,8 +361,6 @@ bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals *kg,
bvh_instance_pop_factor(kg, object, ray, &P, &dir, &idir, &t_fac);
# endif
ray_triangle_intersect_precalc(dir, &isect_precalc);
/* scale isect->t to adjust for instancing */
for(int i = 0; i < num_hits_in_instance; i++) {
(isect_array-i-1)->t *= t_fac;
@ -378,7 +372,6 @@ bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals *kg,
# else
bvh_instance_pop(kg, object, ray, &P, &dir, &idir, FLT_MAX);
# endif
ray_triangle_intersect_precalc(dir, &isect_precalc);
}
isect_t = tmax;

View File

@ -109,9 +109,6 @@ void BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals *kg,
gen_idirsplat_swap(pn, shuf_identity, shuf_swap, idir, idirsplat, shufflexyz);
#endif
TriangleIsectPrecalc isect_precalc;
ray_triangle_intersect_precalc(dir, &isect_precalc);
/* traversal loop */
do {
do {
@ -197,9 +194,9 @@ void BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals *kg,
for(; prim_addr < prim_addr2; prim_addr++) {
kernel_assert(kernel_tex_fetch(__prim_type, prim_addr) == type);
triangle_intersect_subsurface(kg,
&isect_precalc,
ss_isect,
P,
dir,
object,
prim_addr,
isect_t,
@ -214,9 +211,9 @@ void BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals *kg,
for(; prim_addr < prim_addr2; prim_addr++) {
kernel_assert(kernel_tex_fetch(__prim_type, prim_addr) == type);
motion_triangle_intersect_subsurface(kg,
&isect_precalc,
ss_isect,
P,
dir,
ray->time,
object,
prim_addr,

View File

@ -104,9 +104,6 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals *kg,
gen_idirsplat_swap(pn, shuf_identity, shuf_swap, idir, idirsplat, shufflexyz);
#endif
TriangleIsectPrecalc isect_precalc;
ray_triangle_intersect_precalc(dir, &isect_precalc);
/* traversal loop */
do {
do {
@ -238,9 +235,9 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals *kg,
BVH_DEBUG_NEXT_INTERSECTION();
kernel_assert(kernel_tex_fetch(__prim_type, prim_addr) == type);
if(triangle_intersect(kg,
&isect_precalc,
isect,
P,
dir,
visibility,
object,
prim_addr))
@ -267,9 +264,9 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals *kg,
BVH_DEBUG_NEXT_INTERSECTION();
kernel_assert(kernel_tex_fetch(__prim_type, prim_addr) == type);
if(motion_triangle_intersect(kg,
&isect_precalc,
isect,
P,
dir,
ray->time,
visibility,
object,
@ -358,7 +355,6 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals *kg,
# else
isect->t = bvh_instance_push(kg, object, ray, &P, &dir, &idir, isect->t);
# endif
ray_triangle_intersect_precalc(dir, &isect_precalc);
# if defined(__KERNEL_SSE2__)
Psplat[0] = ssef(P.x);
@ -395,7 +391,6 @@ ccl_device_noinline bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals *kg,
# else
isect->t = bvh_instance_pop(kg, object, ray, &P, &dir, &idir, isect->t);
# endif
ray_triangle_intersect_precalc(dir, &isect_precalc);
# if defined(__KERNEL_SSE2__)
Psplat[0] = ssef(P.x);

View File

@ -97,9 +97,6 @@ bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals *kg,
gen_idirsplat_swap(pn, shuf_identity, shuf_swap, idir, idirsplat, shufflexyz);
#endif
TriangleIsectPrecalc isect_precalc;
ray_triangle_intersect_precalc(dir, &isect_precalc);
/* traversal loop */
do {
do {
@ -194,9 +191,9 @@ bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals *kg,
continue;
}
triangle_intersect(kg,
&isect_precalc,
isect,
P,
dir,
visibility,
object,
prim_addr);
@ -215,9 +212,9 @@ bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals *kg,
continue;
}
motion_triangle_intersect(kg,
&isect_precalc,
isect,
P,
dir,
ray->time,
visibility,
object,
@ -243,8 +240,6 @@ bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals *kg,
isect->t = bvh_instance_push(kg, object, ray, &P, &dir, &idir, isect->t);
# endif
ray_triangle_intersect_precalc(dir, &isect_precalc);
# if defined(__KERNEL_SSE2__)
Psplat[0] = ssef(P.x);
Psplat[1] = ssef(P.y);
@ -286,8 +281,6 @@ bool BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals *kg,
isect->t = bvh_instance_pop(kg, object, ray, &P, &dir, &idir, isect->t);
# endif
ray_triangle_intersect_precalc(dir, &isect_precalc);
# if defined(__KERNEL_SSE2__)
Psplat[0] = ssef(P.x);
Psplat[1] = ssef(P.y);

View File

@ -101,9 +101,6 @@ uint BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals *kg,
gen_idirsplat_swap(pn, shuf_identity, shuf_swap, idir, idirsplat, shufflexyz);
#endif /* __KERNEL_SSE2__ */
TriangleIsectPrecalc isect_precalc;
ray_triangle_intersect_precalc(dir, &isect_precalc);
/* traversal loop */
do {
do {
@ -199,9 +196,9 @@ uint BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals *kg,
continue;
}
hit = triangle_intersect(kg,
&isect_precalc,
isect_array,
P,
dir,
visibility,
object,
prim_addr);
@ -243,9 +240,9 @@ uint BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals *kg,
continue;
}
hit = motion_triangle_intersect(kg,
&isect_precalc,
isect_array,
P,
dir,
ray->time,
visibility,
object,
@ -294,7 +291,6 @@ uint BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals *kg,
isect_t = bvh_instance_push(kg, object, ray, &P, &dir, &idir, isect_t);
# endif
ray_triangle_intersect_precalc(dir, &isect_precalc);
num_hits_in_instance = 0;
isect_array->t = isect_t;
@ -340,7 +336,6 @@ uint BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals *kg,
# else
bvh_instance_pop_factor(kg, object, ray, &P, &dir, &idir, &t_fac);
# endif
ray_triangle_intersect_precalc(dir, &isect_precalc);
/* Scale isect->t to adjust for instancing. */
for(int i = 0; i < num_hits_in_instance; i++) {
(isect_array-i-1)->t *= t_fac;
@ -352,7 +347,6 @@ uint BVH_FUNCTION_FULL_NAME(BVH)(KernelGlobals *kg,
# else
bvh_instance_pop(kg, object, ray, &P, &dir, &idir, FLT_MAX);
# endif
ray_triangle_intersect_precalc(dir, &isect_precalc);
}
isect_t = tmax;

View File

@ -97,9 +97,6 @@ ccl_device bool BVH_FUNCTION_FULL_NAME(QBVH)(KernelGlobals *kg,
&near_x, &near_y, &near_z,
&far_x, &far_y, &far_z);
TriangleIsectPrecalc isect_precalc;
ray_triangle_intersect_precalc(dir, &isect_precalc);
/* Traversal loop. */
do {
do {
@ -290,9 +287,9 @@ ccl_device bool BVH_FUNCTION_FULL_NAME(QBVH)(KernelGlobals *kg,
switch(p_type) {
case PRIMITIVE_TRIANGLE: {
hit = triangle_intersect(kg,
&isect_precalc,
isect_array,
P,
dir,
PATH_RAY_SHADOW,
object,
prim_addr);
@ -301,9 +298,9 @@ ccl_device bool BVH_FUNCTION_FULL_NAME(QBVH)(KernelGlobals *kg,
#if BVH_FEATURE(BVH_MOTION)
case PRIMITIVE_MOTION_TRIANGLE: {
hit = motion_triangle_intersect(kg,
&isect_precalc,
isect_array,
P,
dir,
ray->time,
PATH_RAY_SHADOW,
object,
@ -425,8 +422,6 @@ ccl_device bool BVH_FUNCTION_FULL_NAME(QBVH)(KernelGlobals *kg,
org4 = sse3f(ssef(P.x), ssef(P.y), ssef(P.z));
# endif
ray_triangle_intersect_precalc(dir, &isect_precalc);
++stack_ptr;
kernel_assert(stack_ptr < BVH_QSTACK_SIZE);
traversal_stack[stack_ptr].addr = ENTRYPOINT_SENTINEL;
@ -482,8 +477,6 @@ ccl_device bool BVH_FUNCTION_FULL_NAME(QBVH)(KernelGlobals *kg,
org4 = sse3f(ssef(P.x), ssef(P.y), ssef(P.z));
# endif
ray_triangle_intersect_precalc(dir, &isect_precalc);
object = OBJECT_NONE;
node_addr = traversal_stack[stack_ptr].addr;
--stack_ptr;

View File

@ -105,9 +105,6 @@ ccl_device void BVH_FUNCTION_FULL_NAME(QBVH)(KernelGlobals *kg,
&near_x, &near_y, &near_z,
&far_x, &far_y, &far_z);
TriangleIsectPrecalc isect_precalc;
ray_triangle_intersect_precalc(dir, &isect_precalc);
/* Traversal loop. */
do {
do {
@ -253,9 +250,9 @@ ccl_device void BVH_FUNCTION_FULL_NAME(QBVH)(KernelGlobals *kg,
for(; prim_addr < prim_addr2; prim_addr++) {
kernel_assert(kernel_tex_fetch(__prim_type, prim_addr) == type);
triangle_intersect_subsurface(kg,
&isect_precalc,
ss_isect,
P,
dir,
object,
prim_addr,
isect_t,
@ -270,9 +267,9 @@ ccl_device void BVH_FUNCTION_FULL_NAME(QBVH)(KernelGlobals *kg,
for(; prim_addr < prim_addr2; prim_addr++) {
kernel_assert(kernel_tex_fetch(__prim_type, prim_addr) == type);
motion_triangle_intersect_subsurface(kg,
&isect_precalc,
ss_isect,
P,
dir,
ray->time,
object,
prim_addr,

View File

@ -106,9 +106,6 @@ ccl_device bool BVH_FUNCTION_FULL_NAME(QBVH)(KernelGlobals *kg,
&near_x, &near_y, &near_z,
&far_x, &far_y, &far_z);
TriangleIsectPrecalc isect_precalc;
ray_triangle_intersect_precalc(dir, &isect_precalc);
/* Traversal loop. */
do {
do {
@ -333,9 +330,9 @@ ccl_device bool BVH_FUNCTION_FULL_NAME(QBVH)(KernelGlobals *kg,
BVH_DEBUG_NEXT_INTERSECTION();
kernel_assert(kernel_tex_fetch(__prim_type, prim_addr) == type);
if(triangle_intersect(kg,
&isect_precalc,
isect,
P,
dir,
visibility,
object,
prim_addr)) {
@ -354,9 +351,9 @@ ccl_device bool BVH_FUNCTION_FULL_NAME(QBVH)(KernelGlobals *kg,
BVH_DEBUG_NEXT_INTERSECTION();
kernel_assert(kernel_tex_fetch(__prim_type, prim_addr) == type);
if(motion_triangle_intersect(kg,
&isect_precalc,
isect,
P,
dir,
ray->time,
visibility,
object,
@ -447,8 +444,6 @@ ccl_device bool BVH_FUNCTION_FULL_NAME(QBVH)(KernelGlobals *kg,
org4 = sse3f(ssef(P.x), ssef(P.y), ssef(P.z));
# endif
ray_triangle_intersect_precalc(dir, &isect_precalc);
++stack_ptr;
kernel_assert(stack_ptr < BVH_QSTACK_SIZE);
traversal_stack[stack_ptr].addr = ENTRYPOINT_SENTINEL;
@ -489,8 +484,6 @@ ccl_device bool BVH_FUNCTION_FULL_NAME(QBVH)(KernelGlobals *kg,
org4 = sse3f(ssef(P.x), ssef(P.y), ssef(P.z));
# endif
ray_triangle_intersect_precalc(dir, &isect_precalc);
object = OBJECT_NONE;
node_addr = traversal_stack[stack_ptr].addr;
node_dist = traversal_stack[stack_ptr].dist;

View File

@ -91,9 +91,6 @@ ccl_device bool BVH_FUNCTION_FULL_NAME(QBVH)(KernelGlobals *kg,
&near_x, &near_y, &near_z,
&far_x, &far_y, &far_z);
TriangleIsectPrecalc isect_precalc;
ray_triangle_intersect_precalc(dir, &isect_precalc);
/* Traversal loop. */
do {
do {
@ -266,7 +263,7 @@ ccl_device bool BVH_FUNCTION_FULL_NAME(QBVH)(KernelGlobals *kg,
continue;
}
/* Intersect ray against primitive. */
triangle_intersect(kg, &isect_precalc, isect, P, visibility, object, prim_addr);
triangle_intersect(kg, isect, P, dir, visibility, object, prim_addr);
}
break;
}
@ -281,7 +278,7 @@ ccl_device bool BVH_FUNCTION_FULL_NAME(QBVH)(KernelGlobals *kg,
continue;
}
/* Intersect ray against primitive. */
motion_triangle_intersect(kg, &isect_precalc, isect, P, ray->time, visibility, object, prim_addr);
motion_triangle_intersect(kg, isect, P, dir, ray->time, visibility, object, prim_addr);
}
break;
}
@ -316,8 +313,6 @@ ccl_device bool BVH_FUNCTION_FULL_NAME(QBVH)(KernelGlobals *kg,
org4 = sse3f(ssef(P.x), ssef(P.y), ssef(P.z));
# endif
ray_triangle_intersect_precalc(dir, &isect_precalc);
++stack_ptr;
kernel_assert(stack_ptr < BVH_QSTACK_SIZE);
traversal_stack[stack_ptr].addr = ENTRYPOINT_SENTINEL;
@ -362,8 +357,6 @@ ccl_device bool BVH_FUNCTION_FULL_NAME(QBVH)(KernelGlobals *kg,
org4 = sse3f(ssef(P.x), ssef(P.y), ssef(P.z));
# endif
ray_triangle_intersect_precalc(dir, &isect_precalc);
object = OBJECT_NONE;
node_addr = traversal_stack[stack_ptr].addr;
--stack_ptr;

View File

@ -95,9 +95,6 @@ ccl_device uint BVH_FUNCTION_FULL_NAME(QBVH)(KernelGlobals *kg,
&near_x, &near_y, &near_z,
&far_x, &far_y, &far_z);
TriangleIsectPrecalc isect_precalc;
ray_triangle_intersect_precalc(dir, &isect_precalc);
/* Traversal loop. */
do {
do {
@ -271,7 +268,7 @@ ccl_device uint BVH_FUNCTION_FULL_NAME(QBVH)(KernelGlobals *kg,
continue;
}
/* Intersect ray against primitive. */
hit = triangle_intersect(kg, &isect_precalc, isect_array, P, visibility, object, prim_addr);
hit = triangle_intersect(kg, isect_array, P, dir, visibility, object, prim_addr);
if(hit) {
/* Move on to next entry in intersections array. */
isect_array++;
@ -309,7 +306,7 @@ ccl_device uint BVH_FUNCTION_FULL_NAME(QBVH)(KernelGlobals *kg,
continue;
}
/* Intersect ray against primitive. */
hit = motion_triangle_intersect(kg, &isect_precalc, isect_array, P, ray->time, visibility, object, prim_addr);
hit = motion_triangle_intersect(kg, isect_array, P, dir, ray->time, visibility, object, prim_addr);
if(hit) {
/* Move on to next entry in intersections array. */
isect_array++;
@ -367,7 +364,6 @@ ccl_device uint BVH_FUNCTION_FULL_NAME(QBVH)(KernelGlobals *kg,
org4 = sse3f(ssef(P.x), ssef(P.y), ssef(P.z));
# endif
ray_triangle_intersect_precalc(dir, &isect_precalc);
num_hits_in_instance = 0;
isect_array->t = isect_t;
@ -432,8 +428,6 @@ ccl_device uint BVH_FUNCTION_FULL_NAME(QBVH)(KernelGlobals *kg,
org4 = sse3f(ssef(P.x), ssef(P.y), ssef(P.z));
# endif
ray_triangle_intersect_precalc(dir, &isect_precalc);
object = OBJECT_NONE;
node_addr = traversal_stack[stack_ptr].addr;
--stack_ptr;

View File

@ -168,9 +168,9 @@ float3 motion_triangle_refine_subsurface(KernelGlobals *kg,
ccl_device_inline bool motion_triangle_intersect(
KernelGlobals *kg,
const TriangleIsectPrecalc *isect_precalc,
Intersection *isect,
float3 P,
float3 dir,
float time,
uint visibility,
int object,
@ -186,10 +186,10 @@ ccl_device_inline bool motion_triangle_intersect(
motion_triangle_vertices(kg, fobject, prim, time, verts);
/* Ray-triangle intersection, unoptimized. */
float t, u, v;
if(ray_triangle_intersect(isect_precalc,
P,
if(ray_triangle_intersect(P,
dir,
isect->t,
#if defined(__KERNEL_AVX2__) && defined(__KERNEL_SSE__)
#if defined(__KERNEL_SSE2__) && defined(__KERNEL_SSE__)
(ssef*)verts,
#else
verts[0], verts[1], verts[2],
@ -222,9 +222,9 @@ ccl_device_inline bool motion_triangle_intersect(
#ifdef __SUBSURFACE__
ccl_device_inline void motion_triangle_intersect_subsurface(
KernelGlobals *kg,
const TriangleIsectPrecalc *isect_precalc,
SubsurfaceIntersection *ss_isect,
float3 P,
float3 dir,
float time,
int object,
int prim_addr,
@ -242,10 +242,10 @@ ccl_device_inline void motion_triangle_intersect_subsurface(
motion_triangle_vertices(kg, fobject, prim, time, verts);
/* Ray-triangle intersection, unoptimized. */
float t, u, v;
if(ray_triangle_intersect(isect_precalc,
P,
if(ray_triangle_intersect(P,
dir,
tmax,
#if defined(__KERNEL_AVX2__) && defined(__KERNEL_SSE__)
#if defined(__KERNEL_SSE2__) && defined(__KERNEL_SSE__)
(ssef*)verts,
#else
verts[0], verts[1], verts[2],

View File

@ -22,25 +22,16 @@
CCL_NAMESPACE_BEGIN
/* Ray-Triangle intersection for BVH traversal
*
* Sven Woop
* Watertight Ray/Triangle Intersection
*
* http://jcgt.org/published/0002/01/05/paper.pdf
*/
ccl_device_inline bool triangle_intersect(KernelGlobals *kg,
const TriangleIsectPrecalc *isect_precalc,
Intersection *isect,
float3 P,
float3 dir,
uint visibility,
int object,
int prim_addr)
{
const uint tri_vindex = kernel_tex_fetch(__prim_tri_index, prim_addr);
#if defined(__KERNEL_AVX2__) && defined(__KERNEL_SSE__)
#if defined(__KERNEL_SSE2__) && defined(__KERNEL_SSE__)
const ssef *ssef_verts = (ssef*)&kg->__prim_tri_verts.data[tri_vindex];
#else
const float4 tri_a = kernel_tex_fetch(__prim_tri_verts, tri_vindex+0),
@ -48,9 +39,10 @@ ccl_device_inline bool triangle_intersect(KernelGlobals *kg,
tri_c = kernel_tex_fetch(__prim_tri_verts, tri_vindex+2);
#endif
float t, u, v;
if(ray_triangle_intersect(isect_precalc,
P, isect->t,
#if defined(__KERNEL_AVX2__) && defined(__KERNEL_SSE__)
if(ray_triangle_intersect(P,
dir,
isect->t,
#if defined(__KERNEL_SSE2__) && defined(__KERNEL_SSE__)
ssef_verts,
#else
float4_to_float3(tri_a),
@ -86,9 +78,9 @@ ccl_device_inline bool triangle_intersect(KernelGlobals *kg,
#ifdef __SUBSURFACE__
ccl_device_inline void triangle_intersect_subsurface(
KernelGlobals *kg,
const TriangleIsectPrecalc *isect_precalc,
SubsurfaceIntersection *ss_isect,
float3 P,
float3 dir,
int object,
int prim_addr,
float tmax,
@ -96,8 +88,7 @@ ccl_device_inline void triangle_intersect_subsurface(
int max_hits)
{
const uint tri_vindex = kernel_tex_fetch(__prim_tri_index, prim_addr);
#if defined(__KERNEL_AVX2__) && defined(__KERNEL_SSE__)
#if defined(__KERNEL_SSE2__) && defined(__KERNEL_SSE__)
const ssef *ssef_verts = (ssef*)&kg->__prim_tri_verts.data[tri_vindex];
#else
const float3 tri_a = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex+0)),
@ -105,14 +96,13 @@ ccl_device_inline void triangle_intersect_subsurface(
tri_c = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex+2));
#endif
float t, u, v;
if(!ray_triangle_intersect(isect_precalc,
P, tmax,
#if defined(__KERNEL_AVX2__) && defined(__KERNEL_SSE__)
if(!ray_triangle_intersect(P,
dir,
tmax,
#if defined(__KERNEL_SSE2__) && defined(__KERNEL_SSE__)
ssef_verts,
#else
tri_a,
tri_b,
tri_c,
tri_a, tri_b, tri_c,
#endif
&u, &v, &t))
{
@ -150,15 +140,14 @@ ccl_device_inline void triangle_intersect_subsurface(
isect->t = t;
/* Record geometric normal. */
/* TODO(sergey): Check whether it's faster to re-use ssef verts. */
#if defined(__KERNEL_AVX2__) && defined(__KERNEL_SSE__)
#if defined(__KERNEL_SSE2__) && defined(__KERNEL_SSE__)
const float3 tri_a = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex+0)),
tri_b = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex+1)),
tri_c = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex+2));
#endif
ss_isect->Ng[hit] = normalize(cross(tri_b - tri_a, tri_c - tri_a));
}
#endif
#endif /* __SUBSURFACE__ */
/* Refine triangle intersection to more precise hit point. For rays that travel
* far the precision is often not so good, this reintersects the primitive from

View File

@ -79,216 +79,106 @@ ccl_device bool ray_aligned_disk_intersect(
return true;
}
/* Optimized watertight ray-triangle intersection.
*
* Sven Woop
* Watertight Ray/Triangle Intersection
*
* http://jcgt.org/published/0002/01/05/paper.pdf
*/
/* Precalculated data for the ray->tri intersection. */
typedef struct TriangleIsectPrecalc {
/* Maximal dimension kz, and orthogonal dimensions. */
int kx, ky, kz;
/* Shear constants. */
float Sx, Sy, Sz;
} TriangleIsectPrecalc;
/* Workaround stupidness of CUDA/OpenCL which doesn't allow to access indexed
* component of float3 value.
*/
#ifdef __KERNEL_GPU__
# define IDX(vec, idx) \
((idx == 0) ? ((vec).x) : ( (idx == 1) ? ((vec).y) : ((vec).z) ))
#else
# define IDX(vec, idx) ((vec)[idx])
#endif
#if (defined(__KERNEL_OPENCL_APPLE__)) || \
(defined(__KERNEL_CUDA__) && (defined(i386) || defined(_M_IX86)))
ccl_device_noinline
#else
ccl_device_inline
#endif
void ray_triangle_intersect_precalc(float3 dir,
TriangleIsectPrecalc *isect_precalc)
{
/* Calculate dimension where the ray direction is maximal. */
#ifndef __KERNEL_SSE__
int kz = util_max_axis(make_float3(fabsf(dir.x),
fabsf(dir.y),
fabsf(dir.z)));
int kx = kz + 1; if(kx == 3) kx = 0;
int ky = kx + 1; if(ky == 3) ky = 0;
#else
int kx, ky, kz;
/* Avoiding mispredicted branch on direction. */
kz = util_max_axis(fabs(dir));
static const char inc_xaxis[] = {1, 2, 0, 55};
static const char inc_yaxis[] = {2, 0, 1, 55};
kx = inc_xaxis[kz];
ky = inc_yaxis[kz];
#endif
float dir_kz = IDX(dir, kz);
/* Swap kx and ky dimensions to preserve winding direction of triangles. */
if(dir_kz < 0.0f) {
int tmp = kx;
kx = ky;
ky = tmp;
}
/* Calculate the shear constants. */
float inv_dir_z = 1.0f / dir_kz;
isect_precalc->Sx = IDX(dir, kx) * inv_dir_z;
isect_precalc->Sy = IDX(dir, ky) * inv_dir_z;
isect_precalc->Sz = inv_dir_z;
/* Store the dimensions. */
isect_precalc->kx = kx;
isect_precalc->ky = ky;
isect_precalc->kz = kz;
}
ccl_device_forceinline bool ray_triangle_intersect(
const TriangleIsectPrecalc *isect_precalc,
float3 ray_P, float ray_t,
#if defined(__KERNEL_AVX2__) && defined(__KERNEL_SSE__)
float3 ray_P, float3 ray_dir, float ray_t,
#if defined(__KERNEL_SSE2__) && defined(__KERNEL_SSE__)
const ssef *ssef_verts,
#else
const float3 tri_a, const float3 tri_b, const float3 tri_c,
#endif
float *isect_u, float *isect_v, float *isect_t)
{
const int kx = isect_precalc->kx;
const int ky = isect_precalc->ky;
const int kz = isect_precalc->kz;
const float Sx = isect_precalc->Sx;
const float Sy = isect_precalc->Sy;
const float Sz = isect_precalc->Sz;
#if defined(__KERNEL_AVX2__) && defined(__KERNEL_SSE__)
const avxf avxf_P(ray_P.m128, ray_P.m128);
const avxf tri_ab(_mm256_loadu_ps((float *)(ssef_verts)));
const avxf tri_bc(_mm256_loadu_ps((float *)(ssef_verts + 1)));
const avxf AB = tri_ab - avxf_P;
const avxf BC = tri_bc - avxf_P;
const __m256i permute_mask = _mm256_set_epi32(0x3, kz, ky, kx, 0x3, kz, ky, kx);
const avxf AB_k = shuffle(AB, permute_mask);
const avxf BC_k = shuffle(BC, permute_mask);
/* Akz, Akz, Bkz, Bkz, Bkz, Bkz, Ckz, Ckz */
const avxf ABBC_kz = shuffle<2>(AB_k, BC_k);
/* Akx, Aky, Bkx, Bky, Bkx,Bky, Ckx, Cky */
const avxf ABBC_kxy = shuffle<0,1,0,1>(AB_k, BC_k);
const avxf Sxy(Sy, Sx, Sy, Sx);
/* Ax, Ay, Bx, By, Bx, By, Cx, Cy */
const avxf ABBC_xy = nmadd(ABBC_kz, Sxy, ABBC_kxy);
float ABBC_kz_array[8];
_mm256_storeu_ps((float*)&ABBC_kz_array, ABBC_kz);
const float A_kz = ABBC_kz_array[0];
const float B_kz = ABBC_kz_array[2];
const float C_kz = ABBC_kz_array[6];
/* By, Bx, Cy, Cx, By, Bx, Ay, Ax */
const avxf BCBA_yx = permute<3,2,7,6,3,2,1,0>(ABBC_xy);
const avxf neg_mask(0,0,0,0,0x80000000, 0x80000000, 0x80000000, 0x80000000);
/* W U V
* (AxBy-AyBx) (BxCy-ByCx) XX XX (BxBy-ByBx) (CxAy-CyAx) XX XX
*/
const avxf WUxxxxVxx_neg = _mm256_hsub_ps(ABBC_xy * BCBA_yx, neg_mask /* Dont care */);
const avxf WUVWnegWUVW = permute<0,1,5,0,0,1,5,0>(WUxxxxVxx_neg) ^ neg_mask;
/* Calculate scaled barycentric coordinates. */
float WUVW_array[4];
_mm_storeu_ps((float*)&WUVW_array, _mm256_castps256_ps128 (WUVWnegWUVW));
const float W = WUVW_array[0];
const float U = WUVW_array[1];
const float V = WUVW_array[2];
const int WUVW_mask = 0x7 & _mm256_movemask_ps(WUVWnegWUVW);
const int WUVW_zero = 0x7 & _mm256_movemask_ps(_mm256_cmp_ps(WUVWnegWUVW,
_mm256_setzero_ps(), 0));
if(!((WUVW_mask == 7) || (WUVW_mask == 0)) && ((WUVW_mask | WUVW_zero) != 7)) {
return false;
}
#if defined(__KERNEL_SSE2__) && defined(__KERNEL_SSE__)
typedef ssef float3;
const float3 tri_a(ssef_verts[0]);
const float3 tri_b(ssef_verts[1]);
const float3 tri_c(ssef_verts[2]);
const float3 P(ray_P);
const float3 dir(ray_dir);
#else
# define dot3(a, b) dot(a, b)
const float3 P = ray_P;
const float3 dir = ray_dir;
#endif
/* Calculate vertices relative to ray origin. */
const float3 A = make_float3(tri_a.x - ray_P.x, tri_a.y - ray_P.y, tri_a.z - ray_P.z);
const float3 B = make_float3(tri_b.x - ray_P.x, tri_b.y - ray_P.y, tri_b.z - ray_P.z);
const float3 C = make_float3(tri_c.x - ray_P.x, tri_c.y - ray_P.y, tri_c.z - ray_P.z);
const float3 v0 = tri_c - P;
const float3 v1 = tri_a - P;
const float3 v2 = tri_b - P;
const float A_kx = IDX(A, kx), A_ky = IDX(A, ky), A_kz = IDX(A, kz);
const float B_kx = IDX(B, kx), B_ky = IDX(B, ky), B_kz = IDX(B, kz);
const float C_kx = IDX(C, kx), C_ky = IDX(C, ky), C_kz = IDX(C, kz);
/* Calculate triangle edges. */
const float3 e0 = v2 - v0;
const float3 e1 = v0 - v1;
const float3 e2 = v1 - v2;
/* Perform shear and scale of vertices. */
const float Ax = A_kx - Sx * A_kz;
const float Ay = A_ky - Sy * A_kz;
const float Bx = B_kx - Sx * B_kz;
const float By = B_ky - Sy * B_kz;
const float Cx = C_kx - Sx * C_kz;
const float Cy = C_ky - Sy * C_kz;
/* Perform edge tests. */
#ifdef __KERNEL_SSE2__
const float3 crossU = cross(v2 + v0, e0);
const float3 crossV = cross(v0 + v1, e1);
const float3 crossW = cross(v1 + v2, e2);
# ifndef __KERNEL_SSE__
const ssef crossX(crossU.x, crossV.x, crossW.x, crossW.x);
const ssef crossY(crossU.y, crossV.y, crossW.y, crossW.y);
const ssef crossZ(crossU.z, crossV.z, crossW.z, crossW.z);
# else
ssef crossX(crossU);
ssef crossY(crossV);
ssef crossZ(crossW);
ssef zero = _mm_setzero_ps();
_MM_TRANSPOSE4_PS(crossX, crossY, crossZ, zero);
# endif
const ssef dirX(ray_dir.x);
const ssef dirY(ray_dir.y);
const ssef dirZ(ray_dir.z);
/*const*/ ssef UVWW = crossX*dirX + crossY*dirY + crossZ*dirZ;
const float minUVW = reduce_min(UVWW);
const float maxUVW = reduce_max(UVWW);
#else /* __KERNEL_SSE2__ */
const float U = dot(cross(v2 + v0, e0), ray_dir);
const float V = dot(cross(v0 + v1, e1), ray_dir);
const float W = dot(cross(v1 + v2, e2), ray_dir);
const float minUVW = min(U, min(V, W));
const float maxUVW = max(U, max(V, W));
#endif /* __KERNEL_SSE2__ */
/* Calculate scaled barycentric coordinates. */
float U = Cx * By - Cy * Bx;
float V = Ax * Cy - Ay * Cx;
float W = Bx * Ay - By * Ax;
if((U < 0.0f || V < 0.0f || W < 0.0f) &&
(U > 0.0f || V > 0.0f || W > 0.0f))
{
return false;
}
#endif
/* Calculate determinant. */
float det = U + V + W;
if(UNLIKELY(det == 0.0f)) {
if(minUVW < 0.0f && maxUVW > 0.0f) {
return false;
}
/* Calculate scaled z-coordinates of vertices and use them to calculate
* the hit distance.
*/
const float T = (U * A_kz + V * B_kz + W * C_kz) * Sz;
const int sign_det = (__float_as_int(det) & 0x80000000);
const float sign_T = xor_signmask(T, sign_det);
/* Calculate geometry normal and denominator. */
const float3 Ng1 = cross(e1, e0);
//const Vec3vfM Ng1 = stable_triangle_normal(e2,e1,e0);
const float3 Ng = Ng1 + Ng1;
const float den = dot3(Ng, dir);
/* Avoid division by 0. */
if(UNLIKELY(den == 0.0f)) {
return false;
}
/* Perform depth test. */
const float T = dot3(v0, Ng);
const int sign_den = (__float_as_int(den) & 0x80000000);
const float sign_T = xor_signmask(T, sign_den);
if((sign_T < 0.0f) ||
(sign_T > ray_t * xor_signmask(det, sign_det)))
(sign_T > ray_t * xor_signmask(den, sign_den)))
{
return false;
}
/* Workaround precision error on CUDA. */
#ifdef __KERNEL_CUDA__
if(A == B && B == C) {
return false;
}
const float inv_den = 1.0f / den;
#ifdef __KERNEL_SSE2__
UVWW *= inv_den;
_mm_store_ss(isect_u, UVWW);
_mm_store_ss(isect_v, shuffle<1,1,3,3>(UVWW));
#else
*isect_u = U * inv_den;
*isect_v = V * inv_den;
#endif
const float inv_det = 1.0f / det;
*isect_u = U * inv_det;
*isect_v = V * inv_det;
*isect_t = T * inv_det;
*isect_t = T * inv_den;
return true;
}
#undef IDX
#undef dot3
}
ccl_device bool ray_quad_intersect(float3 ray_P, float3 ray_D,
float ray_mint, float ray_maxt,

View File

@ -180,7 +180,7 @@ struct ccl_try_align(16) int3 {
};
__forceinline int3() {}
__forceinline int3(const __m128i a) : m128(a) {}
__forceinline int3(const __m128i& a) : m128(a) {}
__forceinline operator const __m128i&(void) const { return m128; }
__forceinline operator __m128i&(void) { return m128; }
@ -202,7 +202,7 @@ struct ccl_try_align(16) int4 {
};
__forceinline int4() {}
__forceinline int4(const __m128i a) : m128(a) {}
__forceinline int4(const __m128i& a) : m128(a) {}
__forceinline operator const __m128i&(void) const { return m128; }
__forceinline operator __m128i&(void) { return m128; }
@ -274,7 +274,7 @@ struct ccl_try_align(16) float4 {
};
__forceinline float4() {}
__forceinline float4(const __m128 a) : m128(a) {}
__forceinline float4(const __m128& a) : m128(a) {}
__forceinline operator const __m128&(void) const { return m128; }
__forceinline operator __m128&(void) { return m128; }

View File

@ -2771,12 +2771,16 @@ static void lib_link_ntree(FileData *fd, ID *id, bNodeTree *ntree)
* of library blocks that implement this.*/
IDP_LibLinkProperty(node->prop, (fd->flags & FD_FLAGS_SWITCH_ENDIAN), fd);
node->id= newlibadr_us(fd, id->lib, node->id);
node->id = newlibadr_us(fd, id->lib, node->id);
for (sock = node->inputs.first; sock; sock = sock->next)
for (sock = node->inputs.first; sock; sock = sock->next) {
IDP_LibLinkProperty(sock->prop, (fd->flags & FD_FLAGS_SWITCH_ENDIAN), fd);
lib_link_node_socket(fd, id, sock);
for (sock = node->outputs.first; sock; sock = sock->next)
}
for (sock = node->outputs.first; sock; sock = sock->next) {
IDP_LibLinkProperty(sock->prop, (fd->flags & FD_FLAGS_SWITCH_ENDIAN), fd);
lib_link_node_socket(fd, id, sock);
}
}
for (sock = ntree->inputs.first; sock; sock = sock->next) {

View File

@ -895,24 +895,18 @@ static void widget_draw_icon(
float ofs = 1.0f / aspect;
if (but->drawflag & UI_BUT_ICON_LEFT) {
if (but->block->flag & UI_BLOCK_LOOP) {
if (but->type == UI_BTYPE_SEARCH_MENU)
xs = rect->xmin + 4.0f * ofs;
else
xs = rect->xmin + ofs;
}
else {
if (but->dt == UI_EMBOSS_NONE || but->type == UI_BTYPE_LABEL)
xs = rect->xmin + 2.0f * ofs;
else
xs = rect->xmin + 4.0f * ofs;
}
ys = (rect->ymin + rect->ymax - height) / 2.0f;
/* special case - icon_only pie buttons */
if (ui_block_is_pie_menu(but->block) && but->type != UI_BTYPE_MENU && but->str && but->str[0] == '\0')
xs = rect->xmin + 2.0f * ofs;
else if (but->dt == UI_EMBOSS_NONE || but->type == UI_BTYPE_LABEL)
xs = rect->xmin + 2.0f * ofs;
else
xs = rect->xmin + 4.0f * ofs;
}
else {
xs = (rect->xmin + rect->xmax - height) / 2.0f;
ys = (rect->ymin + rect->ymax - height) / 2.0f;
}
ys = (rect->ymin + rect->ymax - height) / 2.0f;
/* force positions to integers, for zoom levels near 1. draws icons crisp. */
if (aspect > 0.95f && aspect < 1.05f) {