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.h412
1 files changed, 169 insertions, 243 deletions
diff --git a/intern/cycles/kernel/geom/geom_triangle_intersect.h b/intern/cycles/kernel/geom/geom_triangle_intersect.h
index 6b4fbcbbb7c..57f4c86d403 100644
--- a/intern/cycles/kernel/geom/geom_triangle_intersect.h
+++ b/intern/cycles/kernel/geom/geom_triangle_intersect.h
@@ -71,28 +71,23 @@ ccl_device_inline bool triangle_intersect(KernelGlobals *kg,
}
#ifdef __KERNEL_AVX2__
-
#define cross256(A,B, C,D) _mm256_fmsub_ps(A,B, _mm256_mul_ps(C,D))
-#if defined(__KERNEL_CUDA__) && __CUDA_ARCH__ < 300
-ccl_device_inline
-#else
-ccl_device_forceinline
-#endif
-int ray_triangle_intersect8(KernelGlobals *kg,
- float3 ray_P,
- float3 ray_dir,
- Intersection **isect,
- uint visibility,
- int object,
- __m256 *triA,
- __m256 *triB,
- __m256 *triC,
- int prim_addr,
- int prim_num,
- uint *num_hits,
- uint max_hits,
- int *num_hits_in_instance,
- float isec_t)
+ccl_device_inline int ray_triangle_intersect8(
+ KernelGlobals *kg,
+ float3 ray_P,
+ float3 ray_dir,
+ Intersection **isect,
+ uint visibility,
+ int object,
+ __m256 *triA,
+ __m256 *triB,
+ __m256 *triC,
+ int prim_addr,
+ int prim_num,
+ uint *num_hits,
+ uint max_hits,
+ int *num_hits_in_instance,
+ float isect_t)
{
const unsigned char prim_num_mask = (1 << prim_num) - 1;
@@ -108,10 +103,6 @@ int ray_triangle_intersect8(KernelGlobals *kg,
const __m256 dirz256 = _mm256_set1_ps(ray_dir.z);
/* Calculate vertices relative to ray origin. */
- /* const float3 v0 = tri_c - P;
- const float3 v1 = tri_a - P;
- const float3 v2 = tri_b - P; */
-
__m256 v0_x_256 = _mm256_sub_ps(triC[0], Px256);
__m256 v0_y_256 = _mm256_sub_ps(triC[1], Py256);
__m256 v0_z_256 = _mm256_sub_ps(triC[2], Pz256);
@@ -136,11 +127,7 @@ int ray_triangle_intersect8(KernelGlobals *kg,
__m256 v1_v2_y_256 = _mm256_add_ps(v1_y_256, v2_y_256);
__m256 v1_v2_z_256 = _mm256_add_ps(v1_z_256, v2_z_256);
- /* Calculate triangle edges.
- const float3 e0 = v2 - v0;
- const float3 e1 = v0 - v1;
- const float3 e2 = v1 - v2;*/
-
+ /* Calculate triangle edges. */
__m256 e0_x_256 = _mm256_sub_ps(v2_x_256, v0_x_256);
__m256 e0_y_256 = _mm256_sub_ps(v2_y_256, v0_y_256);
__m256 e0_z_256 = _mm256_sub_ps(v2_z_256, v0_z_256);
@@ -153,48 +140,32 @@ int ray_triangle_intersect8(KernelGlobals *kg,
__m256 e2_y_256 = _mm256_sub_ps(v1_y_256, v2_y_256);
__m256 e2_z_256 = _mm256_sub_ps(v1_z_256, v2_z_256);
- /* Perform edge tests.
- 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);*/
-
- //cross (AyBz - AzBy, AzBx -AxBz, AxBy - AyBx)
+ /* Perform edge tests. */
+ /* cross (AyBz - AzBy, AzBx -AxBz, AxBy - AyBx) */
__m256 U_x_256 = cross256(v0_v2_y_256, e0_z_256, v0_v2_z_256, e0_y_256);
__m256 U_y_256 = cross256(v0_v2_z_256, e0_x_256, v0_v2_x_256, e0_z_256);
__m256 U_z_256 = cross256(v0_v2_x_256, e0_y_256, v0_v2_y_256, e0_x_256);
- //vertical dot
+ /* vertical dot */
__m256 U_256 = _mm256_mul_ps(U_x_256, dirx256);
- U_256 = _mm256_fmadd_ps(U_y_256, diry256, U_256); //_mm256_add_ps(U_256, _mm256_mul_ps(U_y_256, diry256));
- U_256 = _mm256_fmadd_ps(U_z_256, dirz256, U_256); //_mm256_add_ps(U_256, _mm256_mul_ps(U_z_256, dirz256));
+ U_256 = _mm256_fmadd_ps(U_y_256, diry256, U_256);
+ U_256 = _mm256_fmadd_ps(U_z_256, dirz256, U_256);
__m256 V_x_256 = cross256(v0_v1_y_256, e1_z_256, v0_v1_z_256, e1_y_256);
__m256 V_y_256 = cross256(v0_v1_z_256, e1_x_256, v0_v1_x_256, e1_z_256);
__m256 V_z_256 = cross256(v0_v1_x_256, e1_y_256, v0_v1_y_256, e1_x_256);
- //vertical dot
+ /* vertical dot */
__m256 V_256 = _mm256_mul_ps(V_x_256, dirx256);
- V_256 = _mm256_fmadd_ps(V_y_256, diry256, V_256);// _mm256_add_ps(V_256, _mm256_mul_ps(V_y_256, diry256));
- V_256 = _mm256_fmadd_ps(V_z_256, dirz256, V_256);// _mm256_add_ps(V_256, _mm256_mul_ps(V_z_256, dirz256));
+ V_256 = _mm256_fmadd_ps(V_y_256, diry256, V_256);
+ V_256 = _mm256_fmadd_ps(V_z_256, dirz256, V_256);
__m256 W_x_256 = cross256(v1_v2_y_256, e2_z_256, v1_v2_z_256, e2_y_256);
__m256 W_y_256 = cross256(v1_v2_z_256, e2_x_256, v1_v2_x_256, e2_z_256);
__m256 W_z_256 = cross256(v1_v2_x_256, e2_y_256, v1_v2_y_256, e2_x_256);
- //vertical dot
+ /* vertical dot */
__m256 W_256 = _mm256_mul_ps(W_x_256, dirx256);
- W_256 = _mm256_fmadd_ps(W_y_256, diry256,W_256);//_mm256_add_ps(W_256, _mm256_mul_ps(W_y_256, diry256));
- W_256 = _mm256_fmadd_ps(W_z_256, dirz256,W_256);//_mm256_add_ps(W_256, _mm256_mul_ps(W_z_256, dirz256));
-
- //const float minUVW = min(U, min(V, W));
- //const float maxUVW = max(U, max(V, W));
-#if 0
- __m256 minUVW_256 = _mm256_min_ps(U_256, _mm256_min_ps(V_256, W_256));
- __m256 maxUVW_256 = _mm256_max_ps(U_256, _mm256_max_ps(V_256, W_256));
-
- //if(minUVW < 0.0f && maxUVW > 0.0f)
- __m256i mask_minmaxUVW_256 = _mm256_and_si256(
- _mm256_cmpgt_epi32(zero256, _mm256_castps_si256(minUVW_256)),
- //_mm256_castps_si256(minUVW_256),
- _mm256_cmpgt_epi32(_mm256_castps_si256(maxUVW_256), zero256));
-#else
+ W_256 = _mm256_fmadd_ps(W_y_256, diry256,W_256);
+ W_256 = _mm256_fmadd_ps(W_z_256, dirz256,W_256);
+
__m256i U_256_1 = _mm256_srli_epi32(_mm256_castps_si256(U_256), 31);
__m256i V_256_1 = _mm256_srli_epi32(_mm256_castps_si256(V_256), 31);
__m256i W_256_1 = _mm256_srli_epi32(_mm256_castps_si256(W_256), 31);
@@ -204,9 +175,8 @@ int ray_triangle_intersect8(KernelGlobals *kg,
const __m256i two256 = _mm256_set1_epi32(2);
__m256i mask_minmaxUVW_256 = _mm256_or_si256(
- _mm256_cmpeq_epi32(one256, UVW_256_1),
- _mm256_cmpeq_epi32(two256, UVW_256_1) );
-#endif
+ _mm256_cmpeq_epi32(one256, UVW_256_1),
+ _mm256_cmpeq_epi32(two256, UVW_256_1));
unsigned char mask_minmaxUVW_pos = _mm256_movemask_ps(_mm256_castsi256_ps(mask_minmaxUVW_256));
if((mask_minmaxUVW_pos & prim_num_mask) == prim_num_mask) { //all bits set
@@ -214,231 +184,187 @@ int ray_triangle_intersect8(KernelGlobals *kg,
}
/* Calculate geometry normal and denominator. */
- // const float3 Ng1 = cross(e1, e0);
- //const Vec3vfM Ng1 = stable_triangle_normal(e2,e1,e0);
-
__m256 Ng1_x_256 = cross256(e1_y_256, e0_z_256, e1_z_256, e0_y_256);
__m256 Ng1_y_256 = cross256(e1_z_256, e0_x_256, e1_x_256, e0_z_256);
__m256 Ng1_z_256 = cross256(e1_x_256, e0_y_256, e1_y_256, e0_x_256);
- //const float3 Ng = Ng1 + Ng1;
Ng1_x_256 = _mm256_add_ps(Ng1_x_256, Ng1_x_256);
Ng1_y_256 = _mm256_add_ps(Ng1_y_256, Ng1_y_256);
Ng1_z_256 = _mm256_add_ps(Ng1_z_256, Ng1_z_256);
- //const float den = dot3(Ng, dir);
- //vertical dot
+ /* vertical dot */
__m256 den_256 = _mm256_mul_ps(Ng1_x_256, dirx256);
- den_256 = _mm256_fmadd_ps(Ng1_y_256, diry256,den_256);//_mm256_add_ps(den_256, _mm256_mul_ps(Ng1_y_256, diry256));
- den_256 = _mm256_fmadd_ps(Ng1_z_256, dirz256,den_256);//_mm256_add_ps(den_256, _mm256_mul_ps(Ng1_z_256, dirz256));
-
- // __m256i maskden256 = _mm256_cmpeq_epi32(_mm256_castps_si256(den_256), zero256);
+ den_256 = _mm256_fmadd_ps(Ng1_y_256, diry256,den_256);
+ den_256 = _mm256_fmadd_ps(Ng1_z_256, dirz256,den_256);
/* Perform depth test. */
- //const float T = dot3(v0, Ng);
__m256 T_256 = _mm256_mul_ps(Ng1_x_256, v0_x_256);
- T_256 = _mm256_fmadd_ps(Ng1_y_256, v0_y_256,T_256);//_mm256_add_ps(T_256, _mm256_mul_ps(Ng1_y_256, v0_y_256));
- T_256 = _mm256_fmadd_ps(Ng1_z_256, v0_z_256,T_256);//_mm256_add_ps(T_256, _mm256_mul_ps(Ng1_z_256, v0_z_256));
+ T_256 = _mm256_fmadd_ps(Ng1_y_256, v0_y_256,T_256);
+ T_256 = _mm256_fmadd_ps(Ng1_z_256, v0_z_256,T_256);
- //const int sign_den = (__float_as_int(den) & 0x80000000);
const __m256i c0x80000000 = _mm256_set1_epi32(0x80000000);
__m256i sign_den_256 = _mm256_and_si256(_mm256_castps_si256(den_256), c0x80000000);
- //const float sign_T = xor_signmask(T, sign_den);
__m256 sign_T_256 = _mm256_castsi256_ps(_mm256_xor_si256(_mm256_castps_si256(T_256), sign_den_256));
- /*if((sign_T < 0.0f) || mask_minmaxUVW_pos { return false;} */
unsigned char mask_sign_T = _mm256_movemask_ps(sign_T_256);
if(((mask_minmaxUVW_pos | mask_sign_T) & prim_num_mask) == prim_num_mask) {
return false;
- } /**/
+ }
__m256 xor_signmask_256 = _mm256_castsi256_ps(_mm256_xor_si256(_mm256_castps_si256(den_256), sign_den_256));
-
ccl_align(32) float den8[8], U8[8], V8[8], T8[8], sign_T8[8], xor_signmask8[8];
ccl_align(32) unsigned int mask_minmaxUVW8[8];
- if(visibility == PATH_RAY_SHADOW_OPAQUE){
- __m256i mask_final_256 = _mm256_cmpeq_epi32(mask_minmaxUVW_256, zero256);//~mask_minmaxUVW_256
-
- __m256i maskden256 = _mm256_cmpeq_epi32(_mm256_castps_si256(den_256), zero256);
-
- __m256i mask0 = _mm256_cmpgt_epi32(zero256, _mm256_castps_si256(sign_T_256));
- __m256 rayt_256 = _mm256_set1_ps((*isect)->t);
-
- __m256i mask1 = _mm256_cmpgt_epi32(_mm256_castps_si256(sign_T_256),
- _mm256_castps_si256(
- _mm256_mul_ps(_mm256_castsi256_ps(_mm256_xor_si256(_mm256_castps_si256(den_256), sign_den_256)), rayt_256)
- )
- );
- /* __m256i mask1 = _mm256_castps_si256(_mm256_cmp_ps(sign_T_256,
- _mm256_mul_ps(_mm256_castsi256_ps(_mm256_xor_si256(_mm256_castps_si256(den_256), sign_den_256)), rayt_256),
- _CMP_GT_OS
- ) );*/
-
- mask0 = _mm256_or_si256(mask1, mask0);
- //unsigned char mask = _mm256_movemask_ps(_mm256_castsi256_ps(mask0));
- //unsigned char maskden = _mm256_movemask_ps(_mm256_castsi256_ps(maskden256));
- //unsigned char mask_final = ((~mask) & (~maskden) & (~mask_minmaxUVW_pos));
- mask_final_256 = _mm256_andnot_si256(mask0, mask_final_256); //(~mask_minmaxUVW_pos) &(~mask)
- mask_final_256 = _mm256_andnot_si256(maskden256, mask_final_256); //(~mask_minmaxUVW_pos) &(~mask) & (~maskden)
-
- unsigned char mask_final = _mm256_movemask_ps(_mm256_castsi256_ps(mask_final_256));
- if((mask_final & prim_num_mask) == 0) { //all bits NOT set
- return false;
- } /**/
-
- unsigned long i = 0;
-#if defined(_MSC_VER)
- unsigned char res = _BitScanForward(&i, (unsigned long)mask_final);
-#else
- i = __builtin_ffs(mask_final)-1;
-#endif
-
- den_256 = _mm256_rcp_ps(den_256); //inv_den
- U_256 = _mm256_mul_ps(U_256, den_256); //*inv_den
- V_256 = _mm256_mul_ps(V_256, den_256); //*inv_den
- T_256 = _mm256_mul_ps(T_256, den_256); //*inv_den
-
- _mm256_store_ps(U8, U_256);
- _mm256_store_ps(V8, V_256);
- _mm256_store_ps(T8, T_256);
-
-
- //here we assume (kernel_tex_fetch(__prim_visibility, (prim_addr +i)) & visibility) is always true
-
- (*isect)->u = U8[i];
- (*isect)->v = V8[i];
- (*isect)->t = T8[i];
-
- (*isect)->prim = (prim_addr + i);
- (*isect)->object = object;
- (*isect)->type = PRIMITIVE_TRIANGLE;
-
- return true;
+ if(visibility == PATH_RAY_SHADOW_OPAQUE) {
+ __m256i mask_final_256 = _mm256_cmpeq_epi32(mask_minmaxUVW_256, zero256);
+ __m256i maskden256 = _mm256_cmpeq_epi32(_mm256_castps_si256(den_256), zero256);
+ __m256i mask0 = _mm256_cmpgt_epi32(zero256, _mm256_castps_si256(sign_T_256));
+ __m256 rayt_256 = _mm256_set1_ps((*isect)->t);
+ __m256i mask1 = _mm256_cmpgt_epi32(_mm256_castps_si256(sign_T_256),
+ _mm256_castps_si256(
+ _mm256_mul_ps(_mm256_castsi256_ps(_mm256_xor_si256(_mm256_castps_si256(den_256), sign_den_256)), rayt_256)
+ )
+ );
+ mask0 = _mm256_or_si256(mask1, mask0);
+ mask_final_256 = _mm256_andnot_si256(mask0, mask_final_256); //(~mask_minmaxUVW_pos) &(~mask)
+ mask_final_256 = _mm256_andnot_si256(maskden256, mask_final_256); //(~mask_minmaxUVW_pos) &(~mask) & (~maskden)
+ unsigned char mask_final = _mm256_movemask_ps(_mm256_castsi256_ps(mask_final_256));
+ if((mask_final & prim_num_mask) == 0) {
+ return false;
}
+ const int i = __bsf(mask_final);
+ __m256 inv_den_256 = _mm256_rcp_ps(den_256);
+ U_256 = _mm256_mul_ps(U_256, inv_den_256);
+ V_256 = _mm256_mul_ps(V_256, inv_den_256);
+ T_256 = _mm256_mul_ps(T_256, inv_den_256);
+ _mm256_store_ps(U8, U_256);
+ _mm256_store_ps(V8, V_256);
+ _mm256_store_ps(T8, T_256);
+ /* NOTE: Here we assume visibility for all triangles in the node is
+ * the same. */
+ (*isect)->u = U8[i];
+ (*isect)->v = V8[i];
+ (*isect)->t = T8[i];
+ (*isect)->prim = (prim_addr + i);
+ (*isect)->object = object;
+ (*isect)->type = PRIMITIVE_TRIANGLE;
+ return true;
+ }
else {
- _mm256_store_ps(den8, den_256);
- _mm256_store_ps(U8, U_256);
- _mm256_store_ps(V8, V_256);
- _mm256_store_ps(T8, T_256);
+ _mm256_store_ps(den8, den_256);
+ _mm256_store_ps(U8, U_256);
+ _mm256_store_ps(V8, V_256);
+ _mm256_store_ps(T8, T_256);
- _mm256_store_ps(sign_T8, sign_T_256);
- _mm256_store_ps(xor_signmask8, xor_signmask_256);
- _mm256_store_si256((__m256i*)mask_minmaxUVW8, mask_minmaxUVW_256);
+ _mm256_store_ps(sign_T8, sign_T_256);
+ _mm256_store_ps(xor_signmask8, xor_signmask_256);
+ _mm256_store_si256((__m256i*)mask_minmaxUVW8, mask_minmaxUVW_256);
- int ret = false;
+ int ret = false;
- if(visibility == PATH_RAY_SHADOW) {
- for(int i = 0; i < prim_num; i++) {
- if(!mask_minmaxUVW8[i]) {
+ if(visibility == PATH_RAY_SHADOW) {
+ for(int i = 0; i < prim_num; i++) {
+ if(mask_minmaxUVW8[i]) {
+ continue;
+ }
#ifdef __VISIBILITY_FLAG__
- if(kernel_tex_fetch(__prim_visibility, (prim_addr + i)) & visibility)
+ if((kernel_tex_fetch(__prim_visibility, (prim_addr + i)) & visibility) == 0) {
+ continue;
+ }
#endif
- {
- if((sign_T8[i] >= 0.0f) &&
- (sign_T8[i] <= (*isect)->t * xor_signmask8[i]))
- {
- if(den8[i]) {
- const float inv_den = 1.0f / den8[i];
-
- (*isect)->u = U8[i] * inv_den;
- (*isect)->v = V8[i] * inv_den;
- (*isect)->t = T8[i] * inv_den;
-
- (*isect)->prim = (prim_addr + i);
- (*isect)->object = object;
- (*isect)->type = PRIMITIVE_TRIANGLE;
-
- int prim = kernel_tex_fetch(__prim_index, (*isect)->prim);
- int shader = 0;
-
+ if((sign_T8[i] < 0.0f) ||
+ (sign_T8[i] > (*isect)->t * xor_signmask8[i]))
+ {
+ continue;
+ }
+ if(!den8[i]) {
+ continue;
+ }
+ const float inv_den = 1.0f / den8[i];
+ (*isect)->u = U8[i] * inv_den;
+ (*isect)->v = V8[i] * inv_den;
+ (*isect)->t = T8[i] * inv_den;
+ (*isect)->prim = (prim_addr + i);
+ (*isect)->object = object;
+ (*isect)->type = PRIMITIVE_TRIANGLE;
+ const int prim = kernel_tex_fetch(__prim_index, (*isect)->prim);
+ int shader = 0;
#ifdef __HAIR__
- if(kernel_tex_fetch(__prim_type, (*isect)->prim) & PRIMITIVE_ALL_TRIANGLE)
+ if(kernel_tex_fetch(__prim_type, (*isect)->prim) & PRIMITIVE_ALL_TRIANGLE)
#endif
- {
- shader = kernel_tex_fetch(__tri_shader, prim);
- }
+ {
+ shader = kernel_tex_fetch(__tri_shader, prim);
+ }
#ifdef __HAIR__
- else {
- float4 str = kernel_tex_fetch(__curves, prim);
- shader = __float_as_int(str.z);
- }
+ else {
+ float4 str = kernel_tex_fetch(__curves, prim);
+ shader = __float_as_int(str.z);
+ }
#endif
- int flag = kernel_tex_fetch(__shaders, (shader & SHADER_MASK)).flags;
-
- /* if no transparent shadows, all light is blocked */
- if(!(flag & SD_HAS_TRANSPARENT_SHADOW)) {
- return 2;
- }
- /* if maximum number of hits reached, block all light */
- else if(*num_hits == max_hits) {
- return 2;
- }
- /* move on to next entry in intersections array */
- ret = true;
-
- (*isect)++;
- (*num_hits)++;
-
- (*num_hits_in_instance)++;
-
- (*isect)->t = isec_t;
-
- } //den
- } //if sign
- } //vis
- }//if mask
- } //for
+ const int flag = kernel_tex_fetch(__shaders, (shader & SHADER_MASK)).flags;
+ /* If no transparent shadows, all light is blocked. */
+ if(!(flag & SD_HAS_TRANSPARENT_SHADOW)) {
+ return 2;
+ }
+ /* If maximum number of hits reached, block all light. */
+ else if(num_hits == NULL || *num_hits == max_hits) {
+ return 2;
+ }
+ /* Move on to next entry in intersections array. */
+ ret = true;
+ (*isect)++;
+ (*num_hits)++;
+ (*num_hits_in_instance)++;
+ (*isect)->t = isect_t;
+ }
}
- else { //default case
+ else {
for(int i = 0; i < prim_num; i++) {
- if(!mask_minmaxUVW8[i]) {
+ if(mask_minmaxUVW8[i]) {
+ continue;
+ }
#ifdef __VISIBILITY_FLAG__
- if(kernel_tex_fetch(__prim_visibility, (prim_addr + i)) & visibility)
+ if((kernel_tex_fetch(__prim_visibility, (prim_addr + i)) & visibility) == 0) {
+ continue;
+ }
#endif
- {
- if((sign_T8[i] >= 0.0f) &&
- (sign_T8[i] <= (*isect)->t * xor_signmask8[i]))
- {
- if(den8[i]) {
- const float inv_den = 1.0f / den8[i];
-
- (*isect)->u = U8[i] * inv_den;
- (*isect)->v = V8[i] * inv_den;
- (*isect)->t = T8[i] * inv_den;
-
- (*isect)->prim = (prim_addr + i);
- (*isect)->object = object;
- (*isect)->type = PRIMITIVE_TRIANGLE;
-
- ret = true;
- } //den
- } //if sign
- } //vis
- }//if mask
- } //for
- } //default
- return ret;
-}// else PATH_RAY_SHADOW_OPAQUE
-
+ if((sign_T8[i] < 0.0f) ||
+ (sign_T8[i] > (*isect)->t * xor_signmask8[i]))
+ {
+ continue;
+ }
+ if(!den8[i]) {
+ continue;
+ }
+ const float inv_den = 1.0f / den8[i];
+ (*isect)->u = U8[i] * inv_den;
+ (*isect)->v = V8[i] * inv_den;
+ (*isect)->t = T8[i] * inv_den;
+ (*isect)->prim = (prim_addr + i);
+ (*isect)->object = object;
+ (*isect)->type = PRIMITIVE_TRIANGLE;
+ ret = true;
+ }
+ }
+ return ret;
+ }
}
-//vz static
-ccl_device_inline
-int triangle_intersect8(KernelGlobals *kg,
- Intersection **isect,
- float3 P,
- float3 dir,
- uint visibility,
- int object,
- int prim_addr,
- int prim_num,
- uint *num_hits,
- uint max_hits,
- int *num_hits_in_instance,
- float isec_t)
+ccl_device_inline int triangle_intersect8(
+ KernelGlobals *kg,
+ Intersection **isect,
+ float3 P,
+ float3 dir,
+ uint visibility,
+ int object,
+ int prim_addr,
+ int prim_num,
+ uint *num_hits,
+ uint max_hits,
+ int *num_hits_in_instance,
+ float isect_t)
{
__m128 tri_a[8], tri_b[8], tri_c[8];
__m256 tritmp[12], tri[12];
@@ -540,7 +466,7 @@ int triangle_intersect8(KernelGlobals *kg,
num_hits,
max_hits,
num_hits_in_instance,
- isec_t);
+ isect_t);
return result;
}
generated by cgit v1.2.3 (git 2.25.1) at 2024-06-21 05:33:00 +0300