[Bf-blender-cvs] [451ccf73961] master: Cycles: Cleanup, move curve intersection functions to own file
Sergey Sharybin
noreply at git.blender.org
Mon Aug 7 21:12:48 CEST 2017
Commit: 451ccf739611ebea82f4e6fb05c3640bd330f8b6
Author: Sergey Sharybin
Date: Mon Aug 7 19:53:12 2017 +0200
Branches: master
https://developer.blender.org/rB451ccf739611ebea82f4e6fb05c3640bd330f8b6
Cycles: Cleanup, move curve intersection functions to own file
This way curve file becomes much shorter and it's also easier to write a
benchmark application to check performance before/after future changes.
===================================================================
M intern/cycles/kernel/CMakeLists.txt
M intern/cycles/kernel/geom/geom.h
M intern/cycles/kernel/geom/geom_curve.h
A intern/cycles/kernel/geom/geom_curve_intersect.h
===================================================================
diff --git a/intern/cycles/kernel/CMakeLists.txt b/intern/cycles/kernel/CMakeLists.txt
index 88c4c4e3282..9fe61515570 100644
--- a/intern/cycles/kernel/CMakeLists.txt
+++ b/intern/cycles/kernel/CMakeLists.txt
@@ -202,6 +202,7 @@ set(SRC_GEOM_HEADERS
geom/geom.h
geom/geom_attribute.h
geom/geom_curve.h
+ geom/geom_curve_intersect.h
geom/geom_motion_curve.h
geom/geom_motion_triangle.h
geom/geom_motion_triangle_intersect.h
diff --git a/intern/cycles/kernel/geom/geom.h b/intern/cycles/kernel/geom/geom.h
index c623e3490fd..f34b77ebc07 100644
--- a/intern/cycles/kernel/geom/geom.h
+++ b/intern/cycles/kernel/geom/geom.h
@@ -27,6 +27,7 @@
#include "kernel/geom/geom_motion_triangle_shader.h"
#include "kernel/geom/geom_motion_curve.h"
#include "kernel/geom/geom_curve.h"
+#include "kernel/geom/geom_curve_intersect.h"
#include "kernel/geom/geom_volume.h"
#include "kernel/geom/geom_primitive.h"
diff --git a/intern/cycles/kernel/geom/geom_curve.h b/intern/cycles/kernel/geom/geom_curve.h
index b3ee0343a1e..e35267f02bf 100644
--- a/intern/cycles/kernel/geom/geom_curve.h
+++ b/intern/cycles/kernel/geom/geom_curve.h
@@ -16,18 +16,13 @@ CCL_NAMESPACE_BEGIN
/* Curve Primitive
*
- * Curve primitive for rendering hair and fur. These can be render as flat ribbons
- * or curves with actual thickness. The curve can also be rendered as line segments
- * rather than curves for better performance */
+ * Curve primitive for rendering hair and fur. These can be render as flat
+ * ribbons or curves with actual thickness. The curve can also be rendered as
+ * line segments rather than curves for better performance.
+ */
#ifdef __HAIR__
-#if defined(__KERNEL_CUDA__) && (__CUDA_ARCH__ < 300)
-# define ccl_device_curveintersect ccl_device
-#else
-# define ccl_device_curveintersect ccl_device_forceinline
-#endif
-
/* Reading attributes on various curve elements */
ccl_device float curve_attribute_float(KernelGlobals *kg, const ShaderData *sd, const AttributeDescriptor desc, float *dx, float *dy)
@@ -219,925 +214,6 @@ ccl_device_inline void curvebounds(float *lower, float *upper, float *extremta,
}
}
-#ifdef __KERNEL_SSE2__
-ccl_device_inline ssef transform_point_T3(const ssef t[3], const ssef &a)
-{
- return madd(shuffle<0>(a), t[0], madd(shuffle<1>(a), t[1], shuffle<2>(a) * t[2]));
-}
-#endif
-
-#ifdef __KERNEL_SSE2__
-/* Pass P and dir by reference to aligned vector */
-ccl_device_curveintersect bool cardinal_curve_intersect(KernelGlobals *kg,
- Intersection *isect,
- const float3 &P,
- const float3 &dir,
- uint visibility,
- int object,
- int curveAddr,
- float time,
- int type,
- uint *lcg_state,
- float difl,
- float extmax)
-#else
-ccl_device_curveintersect bool cardinal_curve_intersect(KernelGlobals *kg,
- Intersection *isect,
- float3 P,
- float3 dir,
- uint visibility,
- int object,
- int curveAddr,
- float time,
- int type,
- uint *lcg_state,
- float difl,
- float extmax)
-#endif
-{
- const bool is_curve_primitive = (type & PRIMITIVE_CURVE);
-
- if(!is_curve_primitive && kernel_data.bvh.use_bvh_steps) {
- const float2 prim_time = kernel_tex_fetch(__prim_time, curveAddr);
- if(time < prim_time.x || time > prim_time.y) {
- return false;
- }
- }
-
- int segment = PRIMITIVE_UNPACK_SEGMENT(type);
- float epsilon = 0.0f;
- float r_st, r_en;
-
- int depth = kernel_data.curve.subdivisions;
- int flags = kernel_data.curve.curveflags;
- int prim = kernel_tex_fetch(__prim_index, curveAddr);
-
-#ifdef __KERNEL_SSE2__
- ssef vdir = load4f(dir);
- ssef vcurve_coef[4];
- const float3 *curve_coef = (float3 *)vcurve_coef;
-
- {
- ssef dtmp = vdir * vdir;
- ssef d_ss = mm_sqrt(dtmp + shuffle<2>(dtmp));
- ssef rd_ss = load1f_first(1.0f) / d_ss;
-
- ssei v00vec = load4i((ssei *)&kg->__curves.data[prim]);
- int2 &v00 = (int2 &)v00vec;
-
- int k0 = v00.x + segment;
- int k1 = k0 + 1;
- int ka = max(k0 - 1, v00.x);
- int kb = min(k1 + 1, v00.x + v00.y - 1);
-
-#if defined(__KERNEL_AVX2__) && defined(__KERNEL_SSE__) && (!defined(_MSC_VER) || _MSC_VER > 1800)
- avxf P_curve_0_1, P_curve_2_3;
- if(is_curve_primitive) {
- P_curve_0_1 = _mm256_loadu2_m128(&kg->__curve_keys.data[k0].x, &kg->__curve_keys.data[ka].x);
- P_curve_2_3 = _mm256_loadu2_m128(&kg->__curve_keys.data[kb].x, &kg->__curve_keys.data[k1].x);
- }
- else {
- int fobject = (object == OBJECT_NONE) ? kernel_tex_fetch(__prim_object, curveAddr) : object;
- motion_cardinal_curve_keys_avx(kg, fobject, prim, time, ka, k0, k1, kb, &P_curve_0_1,&P_curve_2_3);
- }
-#else /* __KERNEL_AVX2__ */
- ssef P_curve[4];
-
- if(is_curve_primitive) {
- P_curve[0] = load4f(&kg->__curve_keys.data[ka].x);
- P_curve[1] = load4f(&kg->__curve_keys.data[k0].x);
- P_curve[2] = load4f(&kg->__curve_keys.data[k1].x);
- P_curve[3] = load4f(&kg->__curve_keys.data[kb].x);
- }
- else {
- int fobject = (object == OBJECT_NONE)? kernel_tex_fetch(__prim_object, curveAddr): object;
- motion_cardinal_curve_keys(kg, fobject, prim, time, ka, k0, k1, kb, (float4*)&P_curve);
- }
-#endif /* __KERNEL_AVX2__ */
-
- ssef rd_sgn = set_sign_bit<0, 1, 1, 1>(shuffle<0>(rd_ss));
- ssef mul_zxxy = shuffle<2, 0, 0, 1>(vdir) * rd_sgn;
- ssef mul_yz = shuffle<1, 2, 1, 2>(vdir) * mul_zxxy;
- ssef mul_shuf = shuffle<0, 1, 2, 3>(mul_zxxy, mul_yz);
- ssef vdir0 = vdir & cast(ssei(0xFFFFFFFF, 0xFFFFFFFF, 0xFFFFFFFF, 0));
-
- ssef htfm0 = shuffle<0, 2, 0, 3>(mul_shuf, vdir0);
- ssef htfm1 = shuffle<1, 0, 1, 3>(load1f_first(extract<0>(d_ss)), vdir0);
- ssef htfm2 = shuffle<1, 3, 2, 3>(mul_shuf, vdir0);
-
-#if defined(__KERNEL_AVX2__) && defined(__KERNEL_SSE__) && (!defined(_MSC_VER) || _MSC_VER > 1800)
- const avxf vPP = _mm256_broadcast_ps(&P.m128);
- const avxf htfm00 = avxf(htfm0.m128, htfm0.m128);
- const avxf htfm11 = avxf(htfm1.m128, htfm1.m128);
- const avxf htfm22 = avxf(htfm2.m128, htfm2.m128);
-
- const avxf p01 = madd(shuffle<0>(P_curve_0_1 - vPP),
- htfm00,
- madd(shuffle<1>(P_curve_0_1 - vPP),
- htfm11,
- shuffle<2>(P_curve_0_1 - vPP) * htfm22));
- const avxf p23 = madd(shuffle<0>(P_curve_2_3 - vPP),
- htfm00,
- madd(shuffle<1>(P_curve_2_3 - vPP),
- htfm11,
- shuffle<2>(P_curve_2_3 - vPP)*htfm22));
-
- const ssef p0 = _mm256_castps256_ps128(p01);
- const ssef p1 = _mm256_extractf128_ps(p01, 1);
- const ssef p2 = _mm256_castps256_ps128(p23);
- const ssef p3 = _mm256_extractf128_ps(p23, 1);
-
- const ssef P_curve_1 = _mm256_extractf128_ps(P_curve_0_1, 1);
- r_st = ((float4 &)P_curve_1).w;
- const ssef P_curve_2 = _mm256_castps256_ps128(P_curve_2_3);
- r_en = ((float4 &)P_curve_2).w;
-#else /* __KERNEL_AVX2__ */
- ssef htfm[] = { htfm0, htfm1, htfm2 };
- ssef vP = load4f(P);
- ssef p0 = transform_point_T3(htfm, P_curve[0] - vP);
- ssef p1 = transform_point_T3(htfm, P_curve[1] - vP);
- ssef p2 = transform_point_T3(htfm, P_curve[2] - vP);
- ssef p3 = transform_point_T3(htfm, P_curve[3] - vP);
-
- r_st = ((float4 &)P_curve[1]).w;
- r_en = ((float4 &)P_curve[2]).w;
-#endif /* __KERNEL_AVX2__ */
-
- float fc = 0.71f;
- ssef vfc = ssef(fc);
- ssef vfcxp3 = vfc * p3;
-
- vcurve_coef[0] = p1;
- vcurve_coef[1] = vfc * (p2 - p0);
- vcurve_coef[2] = madd(ssef(fc * 2.0f), p0, madd(ssef(fc - 3.0f), p1, msub(ssef(3.0f - 2.0f * fc), p2, vfcxp3)));
- vcurve_coef[3] = msub(ssef(fc - 2.0f), p2 - p1, msub(vfc, p0, vfcxp3));
-
- }
-#else
- float3 curve_coef[4];
-
- /* curve Intersection check */
- /* obtain curve parameters */
- {
- /* ray transform created - this should be created at beginning of intersection loop */
- Transform htfm;
- float d = sqrtf(dir.x * dir.x + dir.z * dir.z);
- htfm = make_transform(
- dir.z / d, 0, -dir.x /d, 0,
- -dir.x * dir.y /d, d, -dir.y * dir.z /d, 0,
- dir.x, dir.y, dir.z, 0,
- 0, 0, 0, 1);
-
- float4 v00 = kernel_tex_fetch(__curves, prim);
-
- int k0 = __float_as_int(v00.x) + segment;
- int k1 = k0 + 1;
-
- int ka = max(k0 - 1,__float_as_int(v00.x));
- int kb = min(k1 + 1,__float_as_int(v00.x) + __float_as_int(v00.y) - 1);
-
- float4 P_curve[4];
-
- if(is_curve_primitive) {
- P_curve[0] = kernel_tex_fetch(__curve_keys, ka);
- P_curve[1] = kernel_tex_fetch(__curve_keys, k0);
- P_curve[2] = kernel_tex_fetch(__curve_keys, k1);
- P_curve[3] = kernel_tex_fetch(__curve_keys, kb);
- }
- else {
- int fobject = (object == OBJECT_NONE)? kernel_tex_fetch(__prim_object, curveAddr): object;
- motion_cardinal_curve_keys(kg, fobject, prim, time, ka, k0, k1, kb, P_curve);
- }
-
- float3 p0 = transform_point(&htfm, float4_to_float3(P_curve[0]) - P);
- float3 p1 = transform_point(&htfm, float4_to_float3(P_curve[1]) - P);
- float3 p2 = transform_point(&htfm, float4_to_float3(P_curve[2]) - P);
- float3 p3 = transform_point(&htfm, float4_to_float3(P_curve[3]) - P);
-
- float fc = 0.71f;
- curve_coef[0] = p1;
- curve_coef[1] = -fc*p0 + fc*p2;
-
@@ Diff output truncated at 10240 characters. @@
More information about the Bf-blender-cvs
mailing list