text: /*
   * Copyright 1993-2009 NVIDIA Corporation.  All rights reserved.
   *
   * NVIDIA Corporation and its licensors retain all intellectual property and 
   * proprietary rights in and to this software and related documentation and 
   * any modifications thereto.  Any use, reproduction, disclosure, or distribution 
   * of this software and related documentation without an express license 
   * agreement from NVIDIA Corporation is strictly prohibited.
   * 
   */
  
  // TODO: Nick, check that license above ^^^^
  
  /*
      This file implements common mathematical operations on vector types
      (float3, float4 etc.) since these are not provided as standard by CUDA.
  
      The syntax is modelled on the Cg standard library.
  */
  
  #ifndef _FTL_CUDA_OPERATORS_HPP_
  #define _FTL_CUDA_OPERATORS_HPP_
  
  //#include <cuda_runtime.h>
  #include <ftl/cuda_util.hpp>
  
  
  ////////////////////////////////////////////////////////////////////////////////
  typedef unsigned int uint;
  typedef unsigned short ushort;
  
  #ifndef __CUDACC__
  #include <math.h>
  
  inline float fminf(float a, float b)
  {
    return a < b ? a : b;
  }
  
  inline float fmaxf(float a, float b)
  {
    return a > b ? a : b;
  }
  
  inline int max(int a, int b)
  {
    return a > b ? a : b;
  }
  
  inline int min(int a, int b)
  {
    return a < b ? a : b;
  }
  
  inline float rsqrtf(float x)
  {
      return 1.0f / sqrtf(x);
  }
  #endif
  
  // float functions
  ////////////////////////////////////////////////////////////////////////////////
  
  // lerp
  inline __device__ __host__ float lerp(float a, float b, float t)
  {
      return a + t*(b-a);
  }
  
  // clamp
  inline __device__ __host__ float clamp(float f, float a, float b)
  {
      return fmaxf(a, fminf(f, b));
  }
  
  inline __device__ __host__ int sign(float x) { 
  	int t = x<0 ? -1 : 0;
  	return x > 0 ? 1 : t;
  }
  
  // int2 functions
  ////////////////////////////////////////////////////////////////////////////////
  
  inline __host__ __device__ int2 make_int2(float2 f)
  {
    int2 t; t.x = static_cast<int>(f.x); t.y = static_cast<int>(f.y); return t;
  }
  
  inline __host__ __device__ uint2 make_uint2(int2 i)
  {
    uint2 t; t.x = static_cast<uint>(i.x); t.y = static_cast<uint>(i.y); return t;
  }
  
  // negate
  inline __host__ __device__ int2 operator-(int2 &a)
  {
      return make_int2(-a.x, -a.y);
  }
  
  // addition
  inline __host__ __device__ int2 operator+(int2 a, int2 b)
  {
      return make_int2(a.x + b.x, a.y + b.y);
  }
  inline __host__ __device__ void operator+=(int2 &a, int2 b)
  {
      a.x += b.x; a.y += b.y;
  }
  
  // subtract
  inline __host__ __device__ int2 operator-(int2 a, int2 b)
  {
      return make_int2(a.x - b.x, a.y - b.y);
  }
  inline __host__ __device__ void operator-=(int2 &a, int2 b)
  {
      a.x -= b.x; a.y -= b.y;
  }
  
  // multiply
  inline __host__ __device__ int2 operator*(int2 a, int2 b)
  {
      return make_int2(a.x * b.x, a.y * b.y);
  }
  inline __host__ __device__ int2 operator*(int2 a, int s)
  {
      return make_int2(a.x * s, a.y * s);
  }
  inline __host__ __device__ int2 operator*(int s, int2 a)
  {
      return make_int2(a.x * s, a.y * s);
  }
  inline __host__ __device__ void operator*=(int2 &a, int s)
  {
      a.x *= s; a.y *= s;
  }
  
  // float2 functions
  ////////////////////////////////////////////////////////////////////////////////
  
  // negate
  inline __host__ __device__ float2 operator-(float2 &a)
  {
      return make_float2(-a.x, -a.y);
  }
  
  // addition
  inline __host__ __device__ float2 operator+(float2 a, float2 b)
  {
      return make_float2(a.x + b.x, a.y + b.y);
  }
  inline __host__ __device__ void operator+=(float2 &a, float2 b)
  {
      a.x += b.x; a.y += b.y;
  }
  
  // subtract
  inline __host__ __device__ float2 operator-(float2 a, float2 b)
  {
      return make_float2(a.x - b.x, a.y - b.y);
  }
  inline __host__ __device__ void operator-=(float2 &a, float2 b)
  {
      a.x -= b.x; a.y -= b.y;
  }
  
  // multiply
  inline __host__ __device__ float2 operator*(float2 a, float2 b)
  {
      return make_float2(a.x * b.x, a.y * b.y);
  }
  inline __host__ __device__ float2 operator*(float2 a, float s)
  {
      return make_float2(a.x * s, a.y * s);
  }
  inline __host__ __device__ float2 operator*(float s, float2 a)
  {
      return make_float2(a.x * s, a.y * s);
  }
  inline __host__ __device__ void operator*=(float2 &a, float s)
  {
      a.x *= s; a.y *= s;
  }
  
  // divide
  inline __host__ __device__ float2 operator/(float2 a, float2 b)
  {
      return make_float2(a.x / b.x, a.y / b.y);
  }
  inline __host__ __device__ float2 operator/(float2 a, float s)
  {
      float inv = 1.0f / s;
      return a * inv;
  }
  inline __host__ __device__ float2 operator/(float s, float2 a)
  {
      float inv = 1.0f / s;
      return a * inv;
  }
  inline __host__ __device__ void operator/=(float2 &a, float s)
  {
      float inv = 1.0f / s;
      a *= inv;
  }
  
  // lerp
  inline __device__ __host__ float2 lerp(float2 a, float2 b, float t)
  {
      return a + t*(b-a);
  }
  
  // clamp
  inline __device__ __host__ float2 clamp(float2 v, float a, float b)
  {
      return make_float2(clamp(v.x, a, b), clamp(v.y, a, b));
  }
  
  inline __device__ __host__ float2 clamp(float2 v, float2 a, float2 b)
  {
      return make_float2(clamp(v.x, a.x, b.x), clamp(v.y, a.y, b.y));
  }
  
  // dot product
  inline __host__ __device__ float dot(float2 a, float2 b)
  { 
      return a.x * b.x + a.y * b.y;
  }
  
  // length
  inline __host__ __device__ float length(float2 v)
  {
      return sqrtf(dot(v, v));
  }
  
  // normalize
  inline __host__ __device__ float2 normalize(float2 v)
  {
      float invLen = rsqrtf(dot(v, v));
      return v * invLen;
  }
  
  // floor
  inline __host__ __device__ float2 floor(const float2 v)
  {
      return make_float2(floorf(v.x), floorf(v.y));
  }
  
  // reflect
  inline __host__ __device__ float2 reflect(float2 i, float2 n)
  {
  	return i - 2.0f * n * dot(n,i);
  }
  
  // absolute value
  inline __host__ __device__ float2 fabs(float2 v)
  {
  	return make_float2(fabsf(v.x), fabsf(v.y));
  }
  
  inline __device__ __host__ int2 sign(float2 f) { 
  	return make_int2(sign(f.x), sign(f.y));
  }
  
  // float3 functions
  ////////////////////////////////////////////////////////////////////////////////
  
  inline __host__ __device__ float3 make_float3(int3 i)
  {
    float3 t; t.x = static_cast<float>(i.x); t.y = static_cast<float>(i.y); t.z = static_cast<float>(i.z); return t;
  }
  
  inline __host__ __device__ float3 make_float3(float4 f)
  {
      return make_float3(f.x,f.y,f.z);
  }
  
  inline __host__ __device__ float3 make_float3(uchar3 c)
  {
      return make_float3(static_cast<float>(c.x), static_cast<float>(c.y), static_cast<float>(c.z));
  }
  
  inline __host__ __device__ uchar3 make_uchar3(float3 f)
  {
      return make_uchar3(static_cast<unsigned char>(f.x), static_cast<unsigned char>(f.y), static_cast<unsigned char>(f.z));
  }
  
  inline __host__ __device__ float3 make_float3(float f)
  {
      return make_float3(f,f,f);
  }
  
  // negate
  inline __host__ __device__ float3 operator-(const float3 &a)
  {
      return make_float3(-a.x, -a.y, -a.z);
  }
  
  // min
  static __inline__ __host__ __device__ float3 fminf(float3 a, float3 b)
  {
  	return make_float3(fminf(a.x,b.x), fminf(a.y,b.y), fminf(a.z,b.z));
  }
  
  // max
  static __inline__ __host__ __device__ float3 fmaxf(float3 a, float3 b)
  {
  	return make_float3(fmaxf(a.x,b.x), fmaxf(a.y,b.y), fmaxf(a.z,b.z));
  }
  
  // addition
  inline __host__ __device__ float3 operator+(float3 a, float3 b)
  {
      return make_float3(a.x + b.x, a.y + b.y, a.z + b.z);
  }
  inline __host__ __device__ float3 operator+(float3 a, float b)
  {
      return make_float3(a.x + b, a.y + b, a.z + b);
  }
  inline __host__ __device__ void operator+=(float3 &a, float3 b)
  {
      a.x += b.x; a.y += b.y; a.z += b.z;
  }
  
  // subtract
  inline __host__ __device__ float3 operator-(float3 a, float3 b)
  {
      return make_float3(a.x - b.x, a.y - b.y, a.z - b.z);
  }
  inline __host__ __device__ float3 operator-(float3 a, float b)
  {
      return make_float3(a.x - b, a.y - b, a.z - b);
  }
  inline __host__ __device__ void operator-=(float3 &a, float3 b)
  {
      a.x -= b.x; a.y -= b.y; a.z -= b.z;
  }
  
  // multiply
  inline __host__ __device__ float3 operator*(float3 a, float3 b)
  {
      return make_float3(a.x * b.x, a.y * b.y, a.z * b.z);
  }
  inline __host__ __device__ float3 operator*(float3 a, float s)
  {
      return make_float3(a.x * s, a.y * s, a.z * s);
  }
  inline __host__ __device__ float3 operator*(float s, float3 a)
  {
      return make_float3(a.x * s, a.y * s, a.z * s);
  }
  inline __host__ __device__ void operator*=(float3 &a, float s)
  {
      a.x *= s; a.y *= s; a.z *= s;
  }
  
  // divide
  inline __host__ __device__ float3 operator/(float3 a, float3 b)
  {
      return make_float3(a.x / b.x, a.y / b.y, a.z / b.z);
  }
  inline __host__ __device__ float3 operator/(float3 a, float s)
  {
      float inv = 1.0f / s;
      return a * inv;
  }
  inline __host__ __device__ float3 operator/(float s, float3 a)
  {
      float inv = 1.0f / s;
      return a * inv;
  }
  inline __host__ __device__ void operator/=(float3 &a, float s)
  {
      float inv = 1.0f / s;
      a *= inv;
  }
  
  // lerp
  inline __device__ __host__ float3 lerp(float3 a, float3 b, float t)
  {
      return a + t*(b-a);
  }
  
  // clamp
  inline __device__ __host__ float3 clamp(float3 v, float a, float b)
  {
      return make_float3(clamp(v.x, a, b), clamp(v.y, a, b), clamp(v.z, a, b));
  }
  
  inline __device__ __host__ float3 clamp(float3 v, float3 a, float3 b)
  {
      return make_float3(clamp(v.x, a.x, b.x), clamp(v.y, a.y, b.y), clamp(v.z, a.z, b.z));
  }
  
  // dot product
  inline __host__ __device__ float dot(float3 a, float3 b)
  { 
      return a.x * b.x + a.y * b.y + a.z * b.z;
  }
  
  // cross product
  inline __host__ __device__ float3 cross(float3 a, float3 b)
  { 
      return make_float3(a.y*b.z - a.z*b.y, a.z*b.x - a.x*b.z, a.x*b.y - a.y*b.x); 
  }
  
  // length
  inline __host__ __device__ float length(float3 v)
  {
      return sqrtf(dot(v, v));
  }
  
  // length squared
  inline __host__ __device__ float length2(const float3 &v)
  {
      return dot(v, v);
  }
  
  // normalize
  inline __host__ __device__ float3 normalize(float3 v)
  {
      float invLen = rsqrtf(dot(v, v));
      return v * invLen;
  }
  
  // floor
  inline __host__ __device__ float3 floor(const float3 v)
  {
      return make_float3(floorf(v.x), floorf(v.y), floorf(v.z));
  }
  
  // reflect
  inline __host__ __device__ float3 reflect(float3 i, float3 n)
  {
  	return i - 2.0f * n * dot(n,i);
  }
  
  // absolute value
  inline __host__ __device__ float3 fabs(float3 v)
  {
  	return make_float3(fabsf(v.x), fabsf(v.y), fabsf(v.z));
  }
  
  inline __device__ __host__ int3 sign(float3 f) { 
  	return make_int3(sign(f.x), sign(f.y), sign(f.z));
  }
  
  // float4 functions
  ////////////////////////////////////////////////////////////////////////////////
  
  inline __host__ __device__ float4 make_float4(float a)
  {
      return make_float4(a,a,a,a);
  }
  
  inline __host__ __device__ float4 make_float4(float3 f, float a)
  {
      return make_float4(f.x,f.y,f.z,a);
  }
  
  // negate
  inline __host__ __device__ float4 operator-(float4 &a)
  {
      return make_float4(-a.x, -a.y, -a.z, -a.w);
  }
  
  // min
  static __inline__ __host__ __device__ float4 fminf(float4 a, float4 b)
  {
  	return make_float4(fminf(a.x,b.x), fminf(a.y,b.y), fminf(a.z,b.z), fminf(a.w,b.w));
  }
  
  // max
  static __inline__ __host__ __device__ float4 fmaxf(float4 a, float4 b)
  {
  	return make_float4(fmaxf(a.x,b.x), fmaxf(a.y,b.y), fmaxf(a.z,b.z), fmaxf(a.w,b.w));
  }
  
  // addition
  inline __host__ __device__ float4 operator+(float4 a, float4 b)
  {
      return make_float4(a.x + b.x, a.y + b.y, a.z + b.z,  a.w + b.w);
  }
  inline __host__ __device__ void operator+=(float4 &a, float4 b)
  {
      a.x += b.x; a.y += b.y; a.z += b.z; a.w += b.w;
  }
  
  // subtract
  inline __host__ __device__ float4 operator-(float4 a, float4 b)
  {
      return make_float4(a.x - b.x, a.y - b.y, a.z - b.z,  a.w - b.w);
  }
  inline __host__ __device__ void operator-=(float4 &a, float4 b)
  {
      a.x -= b.x; a.y -= b.y; a.z -= b.z; a.w -= b.w;
  }
  
  // multiply
  inline __host__ __device__ float4 operator*(float4 a, float s)
  {
      return make_float4(a.x * s, a.y * s, a.z * s, a.w * s);
  }
  inline __host__ __device__ float4 operator*(float s, float4 a)
  {
      return make_float4(a.x * s, a.y * s, a.z * s, a.w * s);
  }
  inline __host__ __device__ void operator*=(float4 &a, float s)
  {
      a.x *= s; a.y *= s; a.z *= s; a.w *= s;
  }
  
  // divide
  inline __host__ __device__ float4 operator/(float4 a, float4 b)
  {
      return make_float4(a.x / b.x, a.y / b.y, a.z / b.z, a.w / b.w);
  }
  inline __host__ __device__ float4 operator/(float4 a, float s)
  {
      float inv = 1.0f / s;
      return a * inv;
  }
  inline __host__ __device__ float4 operator/(float s, float4 a)
  {
      float inv = 1.0f / s;
      return a * inv;
  }
  inline __host__ __device__ void operator/=(float4 &a, float s)
  {
      float inv = 1.0f / s;
      a *= inv;
  }
  
  // lerp
  inline __device__ __host__ float4 lerp(float4 a, float4 b, float t)
  {
      return a + t*(b-a);
  }
  
  // clamp
  inline __device__ __host__ float4 clamp(float4 v, float a, float b)
  {
      return make_float4(clamp(v.x, a, b), clamp(v.y, a, b), clamp(v.z, a, b), clamp(v.w, a, b));
  }
  
  inline __device__ __host__ float4 clamp(float4 v, float4 a, float4 b)
  {
      return make_float4(clamp(v.x, a.x, b.x), clamp(v.y, a.y, b.y), clamp(v.z, a.z, b.z), clamp(v.w, a.w, b.w));
  }
  
  // dot product
  inline __host__ __device__ float dot(float4 a, float4 b)
  { 
      return a.x * b.x + a.y * b.y + a.z * b.z + a.w * b.w;
  }
  
  // length
  inline __host__ __device__ float length(float4 r)
  {
      return sqrtf(dot(r, r));
  }
  
  // normalize
  inline __host__ __device__ float4 normalize(float4 v)
  {
      float invLen = rsqrtf(dot(v, v));
      return v * invLen;
  }
  
  // floor
  inline __host__ __device__ float4 floor(const float4 v)
  {
      return make_float4(floorf(v.x), floorf(v.y), floorf(v.z), floorf(v.w));
  }
  
  // absolute value
  inline __host__ __device__ float4 fabs(float4 v)
  {
  	return make_float4(fabsf(v.x), fabsf(v.y), fabsf(v.z), fabsf(v.w));
  }
  
  // int3 functions
  ////////////////////////////////////////////////////////////////////////////////
  
  inline __host__ __device__ int3 make_int3(float3 f)
  {
    int3 t; t.x = static_cast<int>(f.x); t.y = static_cast<int>(f.y); t.z = static_cast<int>(f.z); return t;
  }
  
  inline __host__ __device__ int3 make_int3(uint3 i)
  {
    int3 t; t.x = static_cast<int>(i.x); t.y = static_cast<int>(i.y); t.z = static_cast<int>(i.z); return t;
  }
  
  inline __host__ __device__ int3 make_int3(int i)
  {
    int3 t; t.x = i; t.y = i; t.z = i; return t;
  }
  
  __device__ inline int3 make_int3(short3 s) {
  	return make_int3(s.x,s.y,s.z);
  }
  
  __device__ inline int3 make_int3(short4 s) {
  	return make_int3(s.x,s.y,s.z);
  }
  
  // negate
  inline __host__ __device__ int3 operator-(int3 &a)
  {
      return make_int3(-a.x, -a.y, -a.z);
  }
  
  // min
  inline __host__ __device__ int3 min(int3 a, int3 b)
  {
      return make_int3(min(a.x,b.x), min(a.y,b.y), min(a.z,b.z));
  }
  
  // max
  inline __host__ __device__ int3 max(int3 a, int3 b)
  {
      return make_int3(max(a.x,b.x), max(a.y,b.y), max(a.z,b.z));
  }
  
  // addition
  inline __host__ __device__ int3 operator+(int3 a, int3 b)
  {
      return make_int3(a.x + b.x, a.y + b.y, a.z + b.z);
  }
  inline __host__ __device__ void operator+=(int3 &a, int3 b)
  {
      a.x += b.x; a.y += b.y; a.z += b.z;
  }
  
  // subtract
  inline __host__ __device__ int3 operator-(int3 a, int3 b)
  {
      return make_int3(a.x - b.x, a.y - b.y, a.z - b.z);
  }
  
  inline __host__ __device__ void operator-=(int3 &a, int3 b)
  {
      a.x -= b.x; a.y -= b.y; a.z -= b.z;
  }
  
  // multiply
  inline __host__ __device__ int3 operator*(int3 a, int3 b)
  {
      return make_int3(a.x * b.x, a.y * b.y, a.z * b.z);
  }
  inline __host__ __device__ int3 operator*(int3 a, int s)
  {
      return make_int3(a.x * s, a.y * s, a.z * s);
  }
  inline __host__ __device__ int3 operator*(int s, int3 a)
  {
      return make_int3(a.x * s, a.y * s, a.z * s);
  }
  inline __host__ __device__ void operator*=(int3 &a, int s)
  {
      a.x *= s; a.y *= s; a.z *= s;
  }
  
  // divide
  inline __host__ __device__ int3 operator/(int3 a, int3 b)
  {
      return make_int3(a.x / b.x, a.y / b.y, a.z / b.z);
  }
  inline __host__ __device__ int3 operator/(int3 a, int s)
  {
      return make_int3(a.x / s, a.y / s, a.z / s);
  }
  inline __host__ __device__ int3 operator/(int s, int3 a)
  {
      return make_int3(a.x / s, a.y / s, a.z / s);
  }
  inline __host__ __device__ void operator/=(int3 &a, int s)
  {
      a.x /= s; a.y /= s; a.z /= s;
  }
  
  // Comparison
  inline __host__ __device__ bool operator==(const int3 &a, const int3 &b) {
  	return a.x == b.x && a.y == b.y && a.z == b.z;
  }
  
  // clamp
  inline __device__ __host__ int clamp(int f, int a, int b)
  {
      return max(a, min(f, b));
  }
  
  inline __device__ __host__ int3 clamp(int3 v, int a, int b)
  {
      return make_int3(clamp(v.x, a, b), clamp(v.y, a, b), clamp(v.z, a, b));
  }
  
  inline __device__ __host__ int3 clamp(int3 v, int3 a, int3 b)
  {
      return make_int3(clamp(v.x, a.x, b.x), clamp(v.y, a.y, b.y), clamp(v.z, a.z, b.z));
  }
  
  
  // uint3 functions
  ////////////////////////////////////////////////////////////////////////////////
  
  // min
  inline __host__ __device__ uint3 min(uint3 a, uint3 b)
  {
      return make_uint3(min(a.x,b.x), min(a.y,b.y), min(a.z,b.z));
  }
  
  // max
  inline __host__ __device__ uint3 max(uint3 a, uint3 b)
  {
      return make_uint3(max(a.x,b.x), max(a.y,b.y), max(a.z,b.z));
  }
  
  // addition
  inline __host__ __device__ uint3 operator+(uint3 a, uint3 b)
  {
      return make_uint3(a.x + b.x, a.y + b.y, a.z + b.z);
  }
  inline __host__ __device__ void operator+=(uint3 &a, uint3 b)
  {
      a.x += b.x; a.y += b.y; a.z += b.z;
  }
  
  // subtract
  inline __host__ __device__ uint3 operator-(uint3 a, uint3 b)
  {
      return make_uint3(a.x - b.x, a.y - b.y, a.z - b.z);
  }
  
  inline __host__ __device__ void operator-=(uint3 &a, uint3 b)
  {
      a.x -= b.x; a.y -= b.y; a.z -= b.z;
  }
  
  // multiply
  inline __host__ __device__ uint3 operator*(uint3 a, uint3 b)
  {
      return make_uint3(a.x * b.x, a.y * b.y, a.z * b.z);
  }
  inline __host__ __device__ uint3 operator*(uint3 a, uint s)
  {
      return make_uint3(a.x * s, a.y * s, a.z * s);
  }
  inline __host__ __device__ uint3 operator*(uint s, uint3 a)
  {
      return make_uint3(a.x * s, a.y * s, a.z * s);
  }
  inline __host__ __device__ void operator*=(uint3 &a, uint s)
  {
      a.x *= s; a.y *= s; a.z *= s;
  }
  
  // divide
  inline __host__ __device__ uint3 operator/(uint3 a, uint3 b)
  {
      return make_uint3(a.x / b.x, a.y / b.y, a.z / b.z);
  }
  inline __host__ __device__ uint3 operator/(uint3 a, uint s)
  {
      return make_uint3(a.x / s, a.y / s, a.z / s);
  }
  inline __host__ __device__ uint3 operator/(uint s, uint3 a)
  {
      return make_uint3(a.x / s, a.y / s, a.z / s);
  }
  inline __host__ __device__ void operator/=(uint3 &a, uint s)
  {
      a.x /= s; a.y /= s; a.z /= s;
  }
  
  // clamp
  inline __device__ __host__ uint clamp(uint f, uint a, uint b)
  {
      return max(a, min(f, b));
  }
  
  inline __device__ __host__ uint3 clamp(uint3 v, uint a, uint b)
  {
      return make_uint3(clamp(v.x, a, b), clamp(v.y, a, b), clamp(v.z, a, b));
  }
  
  inline __device__ __host__ uint3 clamp(uint3 v, uint3 a, uint3 b)
  {
      return make_uint3(clamp(v.x, a.x, b.x), clamp(v.y, a.y, b.y), clamp(v.z, a.z, b.z));
  }
  
  #endif  // _FTL_CUDA_OPERATORS_HPP_
  