__device__ float3 operator+(const float3 &a, const float3 &b) { return make_float3(a.x+b.x, a.y+b.y, a.z+b.z); } __device__ void operator+=(float3 &a, const float3 &b) { a.x += b.x; a.y += b.y; a.z += b.z; } __device__ float3 operator-(const float3 &a, const float3 &b) { return make_float3(a.x-b.x, a.y-b.y, a.z-b.z); } __device__ void operator-=(float3 &a, const float3 &b) { a.x -= b.x; a.y -= b.y; a.z -= b.z; } __device__ float3 operator*(const float3 &v, float s) { #ifdef USE_FAST_MATH return make_float3(v.x * s, v.y * s, v.z * s); #else return make_float3(__fmul_ru(v.x, s), __fmul_ru(v.y, s), __fmul_ru(v.z, s)); #endif } __device__ void operator*=(float3 &v, float s) { v.x *= s; v.y *= s; v.z *= s; } __device__ float3 operator/(const float3 &v, float s) { #ifdef USE_FAST_MATH const float coef = 1.0f / s; // __frcp_rn(s); return v * coef; #else return make_float3( __fdiv_rn(v.x, s), __fdiv_rn(v.y, s), __fdiv_rn(v.z, s) ); #endif } __device__ void operator/=(float3 &v, float s) { #ifdef USE_FAST_MATH const float coef = 1.0f / s; // __frcp_rn(s); v.x *= coef; v.y *= coef; v.z *= coef; #else v = v / s; #endif } __device__ float innerProduct(const float3 &a, const float3 &b) { #ifdef USE_FAST_MATH // Fast, inaccurate path return a.x*b.x + a.y*b.y + a.z*b.z; #else // Slow, accurate path return __fmaf_rn(a.x, b.x, __fmaf_rn(a.y, b.y, __fmul_ru(a.z, b.z))); #endif } __device__ float3 crossProduct(const float3 &a, const float3 &b) { #ifdef USE_FAST_MATH return make_float3( a.y*b.z - a.z*b.y, a.z*b.z - a.x*b.z, a.x*b.y - a.y*b.z ); #else return make_float3( __fmaf_rn(a.y, b.z, - __fmul_ru(a.z, b.y)), __fmaf_rn(a.z, b.z, - __fmul_ru(a.x, b.z)), __fmaf_rn(a.x, b.y, - __fmul_ru(a.y, b.z)) ); #endif } __device__ float magnitudeSquared(const float3 &v) { return innerProduct(v, v); } __device__ float magnitude(const float3 &v) { #ifdef USE_FAST_MATH return sqrt(innerProduct(v, v)); #else return __fsqrt_rn(innerProduct(v, v)); #endif } __device__ float3 normalize(const float3 &v) { return v / magnitude(v); } __device__ void unify(float3 &v) { v /= magnitude(v); }