diff options
author | Sergey Sharybin <sergey.vfx@gmail.com> | 2018-11-21 13:28:49 +0300 |
---|---|---|
committer | Sergey Sharybin <sergey.vfx@gmail.com> | 2018-11-21 13:28:49 +0300 |
commit | 65d01def80a4fec855360b5e52d5d9724464ba17 (patch) | |
tree | 59f05258c972e8d56104435e6a978ae9063527c6 /intern/cycles/kernel/geom | |
parent | ec851efda93ed19c9f7063c6167ae9437ed6803c (diff) |
Cycles: Cleanup, CUDA code path is not possible inside AVX2
Diffstat (limited to 'intern/cycles/kernel/geom')
-rw-r--r-- | intern/cycles/kernel/geom/geom_triangle_intersect.h | 64 |
1 files changed, 29 insertions, 35 deletions
diff --git a/intern/cycles/kernel/geom/geom_triangle_intersect.h b/intern/cycles/kernel/geom/geom_triangle_intersect.h index 6b4fbcbbb7c..368ea25d941 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 isec_t) { const unsigned char prim_num_mask = (1 << prim_num) - 1; @@ -425,20 +420,19 @@ int ray_triangle_intersect8(KernelGlobals *kg, } -//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 isec_t) { __m128 tri_a[8], tri_b[8], tri_c[8]; __m256 tritmp[12], tri[12]; |