Welcome to mirror list, hosted at ThFree Co, Russian Federation.

git.blender.org/blender.git - Unnamed repository; edit this file 'description' to name the repository.
summaryrefslogtreecommitdiff
path: root/intern/cycles/kernel/geom/geom_triangle_intersect.h
diff options
context:
space:
mode:
Diffstat (limited to 'intern/cycles/kernel/geom/geom_triangle_intersect.h')
-rw-r--r--intern/cycles/kernel/geom/geom_triangle_intersect.h396
1 files changed, 56 insertions, 340 deletions
diff --git a/intern/cycles/kernel/geom/geom_triangle_intersect.h b/intern/cycles/kernel/geom/geom_triangle_intersect.h
index 4db121d94f4..804e74d7e37 100644
--- a/intern/cycles/kernel/geom/geom_triangle_intersect.h
+++ b/intern/cycles/kernel/geom/geom_triangle_intersect.h
@@ -22,232 +22,50 @@
CCL_NAMESPACE_BEGIN
-/* Workaround stupidness of CUDA/OpenCL which doesn't allow to access indexed
- * component of float3 value.
- */
-#ifndef __KERNEL_CPU__
-# define IDX(vec, idx) \
- ((idx == 0) ? ((vec).x) : ( (idx == 1) ? ((vec).y) : ((vec).z) ))
-#else
-# define IDX(vec, idx) ((vec)[idx])
-#endif
-
-/* Ray-Triangle intersection for BVH traversal
- *
- * Sven Woop
- * Watertight Ray/Triangle Intersection
- *
- * http://jcgt.org/published/0002/01/05/paper.pdf
- */
-
-/* Precalculated data for the ray->tri intersection. */
-typedef struct IsectPrecalc {
- /* Maximal dimension kz, and orthogonal dimensions. */
- int kx, ky, kz;
-
- /* Shear constants. */
- float Sx, Sy, Sz;
-} IsectPrecalc;
-
-#if (defined(__KERNEL_OPENCL_APPLE__)) || \
- (defined(__KERNEL_CUDA__) && (defined(i386) || defined(_M_IX86)))
-ccl_device_noinline
-#else
-ccl_device_inline
-#endif
-void triangle_intersect_precalc(float3 dir,
- IsectPrecalc *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;
-}
-
-/* TODO(sergey): Make it general utility function. */
-ccl_device_inline float xor_signmask(float x, int y)
-{
- return __int_as_float(__float_as_int(x) ^ y);
-}
-
ccl_device_inline bool triangle_intersect(KernelGlobals *kg,
- const IsectPrecalc *isect_precalc,
Intersection *isect,
float3 P,
+ float3 dir,
uint visibility,
int object,
int prim_addr)
{
- 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;
-
- /* Calculate vertices relative to ray origin. */
const uint tri_vindex = kernel_tex_fetch(__prim_tri_index, prim_addr);
-
-#if defined(__KERNEL_AVX2__) && defined(__KERNEL_SSE__)
- const avxf avxf_P(P.m128, P.m128);
-
- const avxf tri_ab = kernel_tex_fetch_avxf(__prim_tri_verts, tri_vindex + 0);
- const avxf tri_bc = kernel_tex_fetch_avxf(__prim_tri_verts, tri_vindex + 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__)
+ 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),
tri_b = kernel_tex_fetch(__prim_tri_verts, tri_vindex+1),
tri_c = kernel_tex_fetch(__prim_tri_verts, tri_vindex+2);
- const float3 A = make_float3(tri_a.x - P.x, tri_a.y - P.y, tri_a.z - P.z);
- const float3 B = make_float3(tri_b.x - P.x, tri_b.y - P.y, tri_b.z - P.z);
- const float3 C = make_float3(tri_c.x - P.x, tri_c.y - P.y, tri_c.z - P.z);
-
- 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);
-
- /* 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;
-
- /* 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)) {
- 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);
- if((sign_T < 0.0f) ||
- (sign_T > isect->t * xor_signmask(det, sign_det)))
+ float t, u, v;
+ if(ray_triangle_intersect(P,
+ dir,
+ isect->t,
+#if defined(__KERNEL_SSE2__) && defined(__KERNEL_SSE__)
+ ssef_verts,
+#else
+ float4_to_float3(tri_a),
+ float4_to_float3(tri_b),
+ float4_to_float3(tri_c),
+#endif
+ &u, &v, &t))
{
- return false;
- }
-
#ifdef __VISIBILITY_FLAG__
- /* visibility flag test. we do it here under the assumption
- * that most triangles are culled by node flags */
- if(kernel_tex_fetch(__prim_visibility, prim_addr) & visibility)
+ /* Visibility flag test. we do it here under the assumption
+ * that most triangles are culled by node flags.
+ */
+ if(kernel_tex_fetch(__prim_visibility, prim_addr) & visibility)
#endif
- {
-#ifdef __KERNEL_CUDA__
- if(A == B && B == C) {
- return false;
+ {
+ isect->prim = prim_addr;
+ isect->object = object;
+ isect->type = PRIMITIVE_TRIANGLE;
+ isect->u = u;
+ isect->v = v;
+ isect->t = t;
+ return true;
}
-#endif
- /* Normalize U, V, W, and T. */
- const float inv_det = 1.0f / det;
- isect->prim = prim_addr;
- isect->object = object;
- isect->type = PRIMITIVE_TRIANGLE;
- isect->u = U * inv_det;
- isect->v = V * inv_det;
- isect->t = T * inv_det;
- return true;
}
return false;
}
@@ -260,138 +78,37 @@ ccl_device_inline bool triangle_intersect(KernelGlobals *kg,
#ifdef __SUBSURFACE__
ccl_device_inline void triangle_intersect_subsurface(
KernelGlobals *kg,
- const IsectPrecalc *isect_precalc,
SubsurfaceIntersection *ss_isect,
float3 P,
+ float3 dir,
int object,
int prim_addr,
float tmax,
uint *lcg_state,
int max_hits)
{
- 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;
-
- /* Calculate vertices relative to ray origin. */
const uint tri_vindex = kernel_tex_fetch(__prim_tri_index, prim_addr);
- const float4 tri_a = kernel_tex_fetch(__prim_tri_verts, tri_vindex+0),
- tri_b = kernel_tex_fetch(__prim_tri_verts, tri_vindex+1),
- tri_c = kernel_tex_fetch(__prim_tri_verts, tri_vindex+2);
-
-#if defined(__KERNEL_AVX2__) && defined(__KERNEL_SSE__)
- const avxf avxf_P(P.m128, P.m128);
-
- const avxf tri_ab = kernel_tex_fetch_avxf(__prim_tri_verts, tri_vindex + 0);
- const avxf tri_bc = kernel_tex_fetch_avxf(__prim_tri_verts, tri_vindex + 1);
-
- const avxf AB = tri_ab - avxf_P;
- const avxf BC = tri_bc - avxf_P;
-
- const __m256i permuteMask = _mm256_set_epi32(0x3, kz, ky, kx, 0x3, kz, ky, kx);
-
- const avxf AB_k = shuffle(AB, permuteMask);
- const avxf BC_k = shuffle(BC, permuteMask);
-
- /* 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 negMask(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, negMask /* Dont care */);
-
- const avxf WUVWnegWUVW = permute<0,1,5,0,0,1,5,0>(WUxxxxVxx_neg) ^ negMask;
-
- /* 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;
- }
+#if defined(__KERNEL_SSE2__) && defined(__KERNEL_SSE__)
+ const ssef *ssef_verts = (ssef*)&kg->__prim_tri_verts.data[tri_vindex];
#else
- const float3 A = make_float3(tri_a.x - P.x, tri_a.y - P.y, tri_a.z - P.z);
- const float3 B = make_float3(tri_b.x - P.x, tri_b.y - P.y, tri_b.z - P.z);
- const float3 C = make_float3(tri_c.x - P.x, tri_c.y - P.y, tri_c.z - P.z);
-
- 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);
-
- /* 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;
-
- /* 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;
- }
+ 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
-
- /* Calculate determinant. */
- float det = U + V + W;
- if(UNLIKELY(det == 0.0f)) {
- return;
- }
-
- /* Calculate scaled z−coordinates of vertices and use them to calculate
- * the hit distance.
- */
- const int sign_det = (__float_as_int(det) & 0x80000000);
- const float T = (U * A_kz + V * B_kz + W * C_kz) * Sz;
- const float sign_T = xor_signmask(T, sign_det);
- if((sign_T < 0.0f) ||
- (sign_T > tmax * xor_signmask(det, sign_det)))
+ float t, u, v;
+ if(!ray_triangle_intersect(P,
+ dir,
+ tmax,
+#if defined(__KERNEL_SSE2__) && defined(__KERNEL_SSE__)
+ ssef_verts,
+#else
+ tri_a, tri_b, tri_c,
+#endif
+ &u, &v, &t))
{
return;
}
- /* Normalize U, V, W, and T. */
- const float inv_det = 1.0f / det;
-
- const float t = T * inv_det;
for(int i = min(max_hits, ss_isect->num_hits) - 1; i >= 0; --i) {
if(ss_isect->hits[i].t == t) {
return;
@@ -418,18 +135,19 @@ ccl_device_inline void triangle_intersect_subsurface(
isect->prim = prim_addr;
isect->object = object;
isect->type = PRIMITIVE_TRIANGLE;
- isect->u = U * inv_det;
- isect->v = V * inv_det;
+ isect->u = u;
+ isect->v = v;
isect->t = t;
/* Record geometric normal. */
- /* TODO(sergey): Use float4_to_float3() on just an edges. */
- const float3 v0 = float4_to_float3(tri_a);
- const float3 v1 = float4_to_float3(tri_b);
- const float3 v2 = float4_to_float3(tri_c);
- ss_isect->Ng[hit] = normalize(cross(v1 - v0, v2 - v0));
-}
+#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 /* __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
@@ -457,7 +175,7 @@ ccl_device_inline float3 triangle_refine(KernelGlobals *kg,
return P;
}
# ifdef __OBJECT_MOTION__
- Transform tfm = ccl_fetch(sd, ob_itfm);
+ Transform tfm = sd->ob_itfm;
# else
Transform tfm = object_fetch_transform(kg, isect->object, OBJECT_INVERSE_TRANSFORM);
# endif
@@ -491,7 +209,7 @@ ccl_device_inline float3 triangle_refine(KernelGlobals *kg,
if(isect->object != OBJECT_NONE) {
# ifdef __OBJECT_MOTION__
- Transform tfm = ccl_fetch(sd, ob_tfm);
+ Transform tfm = sd->ob_tfm;
# else
Transform tfm = object_fetch_transform(kg, isect->object, OBJECT_TRANSFORM);
# endif
@@ -519,7 +237,7 @@ ccl_device_inline float3 triangle_refine_subsurface(KernelGlobals *kg,
if(isect->object != OBJECT_NONE) {
#ifdef __OBJECT_MOTION__
- Transform tfm = ccl_fetch(sd, ob_itfm);
+ Transform tfm = sd->ob_itfm;
#else
Transform tfm = object_fetch_transform(kg,
isect->object,
@@ -557,7 +275,7 @@ ccl_device_inline float3 triangle_refine_subsurface(KernelGlobals *kg,
if(isect->object != OBJECT_NONE) {
#ifdef __OBJECT_MOTION__
- Transform tfm = ccl_fetch(sd, ob_tfm);
+ Transform tfm = sd->ob_tfm;
#else
Transform tfm = object_fetch_transform(kg,
isect->object,
@@ -570,6 +288,4 @@ ccl_device_inline float3 triangle_refine_subsurface(KernelGlobals *kg,
return P;
}
-#undef IDX
-
CCL_NAMESPACE_END
generated by cgit v1.2.3 (git 2.25.1) at 2024-05-22 09:13:02 +0300