TAPsCUDA_UsefulDeviceFns.cu

Go to the documentation of this file.

/******************************************************************************
TAPsCUDA_UsefulDeviceFns.cu
******************************************************************************/
/******************************************************************************
SUKITTI PUNAK   (09/03/2009)
******************************************************************************/
#ifndef TAPs_CUDA_USEFUL_DEVICE_FNS_HPP
#define TAPs_CUDA_USEFUL_DEVICE_FNS_HPP

#include "../Core/TAPsListOfNamespaces.hpp"

BEGIN_NAMESPACE_TAPs__CUDA
//=============================================================================
// START: CUDA Constants
//-----------------------------------------------------------------------------
 //__constant__ float CUDA_K_PI = 4.0*atanf(1.0);
 //__constant__ float TAPsCUDA_K_PI = 3.1415926535897932384626433832795f;
//-----------------------------------------------------------------------------
// END: CUDA Constants
//=============================================================================


//=============================================================================
// START: Operators for build-in data-types
//-----------------------------------------------------------------------------

// float3
__device__
float3 operator+ ( const float3 & a, const float3 & b )
{
    float3 c;
    c.x = a.x + b.x;
    c.y = a.y + b.y;
    c.z = a.z + b.z;
    return c;
}

__device__
float3 operator- ( const float3 & a, const float3 & b )
{
    float3 c;
    c.x = a.x - b.x;
    c.y = a.y - b.y;
    c.z = a.z - b.z;
    return c;
}

__device__
float3 operator* ( const float3 & a, float s )
{
    float3 c;
    c.x = a.x * s;
    c.y = a.y * s;
    c.z = a.z * s;
    return c;
}

__device__
float3 operator* ( float s, const float3 & a )
{
    return a * s;
}

__device__
float3 operator/ ( const float3 & a, float s )
{
    float3 c;
    c.x = a.x / s;
    c.y = a.y / s;
    c.z = a.z / s;
    return c;
}


// float4
__device__
float4 operator+ ( const float4 & a, const float4 & b )
{
    float4 c;
    c.x = a.x + b.x;
    c.y = a.y + b.y;
    c.z = a.z + b.z;
    c.w = a.w + b.w;
    return c;
}

__device__
float4 operator- ( const float4 & a, const float4 & b )
{
    float4 c;
    c.x = a.x - b.x;
    c.y = a.y - b.y;
    c.z = a.z - b.z;
    c.w = a.w - b.w;
    return c;
}

__device__
float4 operator* ( const float4 & a, float s )
{
    float4 c;
    c.x = a.x * s;
    c.y = a.y * s;
    c.z = a.z * s;
    c.w = a.w * s;
    return c;
}

__device__
float4 operator* ( float s, const float4 & a )
{
    return a * s;
}

__device__
float4 operator/ ( const float4 & a, float s )
{
    float4 c;
    c.x = a.x / s;
    c.y = a.y / s;
    c.z = a.z / s;
    c.w = a.w / s;
    return c;
}
//-----------------------------------------------------------------------------
// END: Operators for build-in data-types
//=============================================================================


//=============================================================================
// START: 3-by-3 Matrix
//-----------------------------------------------------------------------------
struct matrix3x3 {
    float3 col[3];  
};
__device__
struct matrix3x3 make_matrix3x3( 
    float e00, float e01, float e02,
    float e10, float e11, float e12,
    float e20, float e21, float e22 
)
{
    struct matrix3x3 A;
    A.col[0].x = e00;  A.col[0].y = e10;  A.col[0].z = e20;
    A.col[1].x = e01;  A.col[1].y = e11;  A.col[1].z = e21;
    A.col[2].x = e02;  A.col[2].y = e12;  A.col[2].z = e22;
    return A;
}

__device__
struct matrix3x3 operator+ ( const struct matrix3x3 & A, const struct matrix3x3 & B )
{
    struct matrix3x3 C;
    #pragma unroll 3
    for ( int i = 0; i < 3; ++i ) {
        C.col[i] = A.col[i] + B.col[i];
    }
    return C;
}

__device__
struct matrix3x3 operator- ( const struct matrix3x3 & A, const struct matrix3x3 & B )
{
    struct matrix3x3 C;
    #pragma unroll 3
    for ( int i = 0; i < 3; ++i ) {
        C.col[i] = A.col[i] - B.col[i];
    }
    return C;
}

__device__
struct matrix3x3 operator* ( const struct matrix3x3 & A, float s )
{
    struct matrix3x3 C;
    #pragma unroll 3
    for ( int i = 0; i < 3; ++i ) {
        C.col[i] = A.col[i] * s;
    }
    return C;
}

__device__
struct matrix3x3 operator/ ( const struct matrix3x3 & A, float s )
{
    struct matrix3x3 C;
    #pragma unroll 3
    for ( int i = 0; i < 3; ++i ) {
        C.col[i] = A.col[i] / s;
    }
    return C;
}
//-----------------------------------------------------------------------------
// END: 3-by-3 Matrix
//=============================================================================


//=============================================================================
// START: 4-by-4 Matrix
//-----------------------------------------------------------------------------
struct matrix4x4 {
    float4 col[4];  
};

__device__
struct matrix4x4 make_matrix4x4( 
    float e00, float e01, float e02, float e03, 
    float e10, float e11, float e12, float e13, 
    float e20, float e21, float e22, float e23, 
    float e30, float e31, float e32, float e33
)
{
    struct matrix4x4 A;
    A.col[0].x = e00;  A.col[0].y = e10;  A.col[0].z = e20;  A.col[0].w = e30;
    A.col[1].x = e01;  A.col[1].y = e11;  A.col[1].z = e21;  A.col[1].w = e31;
    A.col[2].x = e02;  A.col[2].y = e12;  A.col[2].z = e22;  A.col[2].w = e32;
    A.col[3].x = e03;  A.col[3].y = e13;  A.col[3].z = e23;  A.col[3].w = e33;
    return A;
}

__device__
struct matrix4x4 operator+ ( const struct matrix4x4 & A, const struct matrix4x4 & B )
{
    struct matrix4x4 C;
    #pragma unroll 4
    for ( int i = 0; i < 4; ++i ) {
        C.col[i] = A.col[i] + B.col[i];
    }
    return C;
}

__device__
struct matrix4x4 operator- ( const struct matrix4x4 & A, const struct matrix4x4 & B )
{
    struct matrix4x4 C;
    #pragma unroll 4
    for ( int i = 0; i < 4; ++i ) {
        C.col[i] = A.col[i] - B.col[i];
    }
    return C;
}

__device__
struct matrix4x4 operator* ( const struct matrix4x4 & A, float s )
{
    struct matrix4x4 C;
    #pragma unroll 4
    for ( int i = 0; i < 4; ++i ) {
        C.col[i] = A.col[i] * s;
    }
    return C;
}

__device__
struct matrix4x4 operator/ ( const struct matrix4x4 & A, float s )
{
    struct matrix4x4 C;
    #pragma unroll 4
    for ( int i = 0; i < 4; ++i ) {
        C.col[i] = A.col[i] / s;
    }
    return C;
}

__device__
float4 operator* ( const struct matrix4x4 & A, const float4 & v )
{
    float4 r = make_float4(
        A.col[0].x*v.x + A.col[1].x*v.y + A.col[2].x*v.z + A.col[3].x*v.w,
        A.col[0].y*v.x + A.col[1].y*v.y + A.col[2].y*v.z + A.col[3].y*v.w,
        A.col[0].z*v.x + A.col[1].z*v.y + A.col[2].z*v.z + A.col[3].z*v.w,
        A.col[0].w*v.x + A.col[1].w*v.y + A.col[2].w*v.z + A.col[3].w*v.w
    );
    return r;
}
//-----------------------------------------------------------------------------
// END: 4-by-4 Matrix
//=============================================================================




//=============================================================================
//-----------------------------------------------------------------------------
__device__
float3 XYZ ( const float4 & V )
{
    return make_float3( V.x, V.y, V.z );
}

__device__
float4 XYZW ( const float4 & V, float w )
{
    return make_float4( V.x, V.y, V.z, w );
}

__device__
float3 Add ( const float3 & V, const float3 & W )
{
    return make_float3( V.x+W.x, V.y+W.y, V.z+W.z );
}

__device__
float4 Add ( const float4 & V, const float4 & W )
{
    return make_float4( V.x+W.x, V.y+W.y, V.z+W.z, V.w+W.w );
}

__device__
float3 Sub ( const float3 & V, const float3 & W )
{
    return make_float3( V.x-W.x, V.y-W.y, V.z-W.z );
}

__device__
float4 Sub ( const float4 & V, const float4 & W )
{
    return make_float4( V.x-W.x, V.y-W.y, V.z-W.z, V.w-W.w );
}

__device__
float InnerProduct ( const float3 & V, const float3 & W )
{
    return V.x*W.x + V.y*W.y + V.z*W.z;
}

__device__
float InnerProduct ( const float4 & V, const float4 & W )
{
    return V.x*W.x + V.y*W.y + V.z*W.z + V.w*W.w;
}

__device__
float3 Mul ( const float3 & V, float s )
{
    return make_float3( V.x*s, V.y*s, V.z*s );
}

__device__
float4 Mul ( const float4 & V, float s )
{
    return make_float4( V.x*s, V.y*s, V.z*s, V.w*s );
}

__device__
float3 Div ( const float3 & V, float s )
{
    return make_float3( V.x/s, V.y/s, V.z/s );
}

__device__
float4 Div ( const float4 & V, float s )
{
    return make_float4( V.x/s, V.y/s, V.z/s, V.w/s );
}

__device__
float LengthSquare ( const float3 & V )
{
    return InnerProduct( V, V );
}

__device__
float LengthSquare ( const float4 & V )
{
    return InnerProduct( V, V );
}

__device__
float Length ( const float3 & V )
{
    return sqrtf( LengthSquare( V ) );
}

__device__
float Length ( const float4 & V )
{
    return sqrtf( LengthSquare( V ) );
}

__device__
float3 Unit ( const float3 & V )
{
    return Div( V, Length(V) );
}

__device__
float4 Unit ( const float4 & V )
{
    return Div( V, Length(V) );
}
//-----------------------------------------------------------------------------
//=============================================================================


//=============================================================================
// START: Fns for Quaternion
//-----------------------------------------------------------------------------
__device__
float4 QuaternionConjugate ( const float4 & Q )
{
    return make_float4( Q.x, -Q.y, -Q.z, -Q.w );
}

__device__
float QuaternionNormSquare ( const float4 & Q )
{
    return Q.x*Q.x + Q.y*Q.y + Q.z*Q.z + Q.w*Q.w;
}

__device__
float QuaternionNorm ( const float4 & Q )
{
    return sqrtf( QuaternionNormSquare( Q ) );
}

__device__
float4 QuaternionUnit ( const float4 & Q )
{
    return Div( Q, QuaternionNorm(Q) );
}

__device__
float4 QuaternionMul ( const float4 & Q1, const float4 & Q2 )
{
    return make_float4(
        Q1.x*Q2.x - Q1.y*Q2.y - Q1.z*Q2.z - Q1.w*Q2.w,
        Q1.x*Q2.y + Q1.y*Q2.x + Q1.z*Q2.w - Q1.w*Q2.z,
        Q1.x*Q2.z + Q1.z*Q2.x + Q1.w*Q2.y - Q1.y*Q2.w,
        Q1.x*Q2.w + Q1.w*Q2.x + Q1.y*Q2.z - Q1.z*Q2.y
    );
}

__device__
struct matrix3x3 QuatenionToRotationMatrix3x3 ( const float4 & Q )
{
    float rr = Q.x*Q.x;
    float ii = Q.y*Q.y;
    float jj = Q.z*Q.z;
    float kk = Q.w*Q.w;
    float ri = Q.x*Q.y;
    float rj = Q.x*Q.z;
    float rk = Q.x*Q.w;
    float ij = Q.y*Q.z;
    float jk = Q.z*Q.w;
    float ik = Q.y*Q.w;

    struct matrix3x3 R = make_matrix3x3(
        rr+ii-jj-kk,    2*(ij-rk),    2*(ik+rj),
          2*(ij+rk),  rr-ii+jj-kk,    2*(jk-ri),
          2*(ik-rj),    2*(jk+ri),  rr-ii-jj+kk
    );
    return R;
}

__device__
struct matrix4x4 QuatenionToRotationMatrix4x4 ( const float4 & Q )
{
    float rr = Q.x*Q.x;
    float ii = Q.y*Q.y;
    float jj = Q.z*Q.z;
    float kk = Q.w*Q.w;
    float ri = Q.x*Q.y;
    float rj = Q.x*Q.z;
    float rk = Q.x*Q.w;
    float ij = Q.y*Q.z;
    float jk = Q.z*Q.w;
    float ik = Q.y*Q.w;

    struct matrix4x4 R = make_matrix4x4(
        rr+ii-jj-kk,    2*(ij-rk),    2*(ik+rj),  0,
          2*(ij+rk),  rr-ii+jj-kk,    2*(jk-ri),  0,
          2*(ik-rj),    2*(jk+ri),  rr-ii-jj+kk,  0,
             0,            0,            0,       1
    );
    return R;
}

//-----------------------------------------------------------------------------
// END: Fns for Quaternion
//=============================================================================
END_NAMESPACE_TAPs__CUDA
//-----------------------------------------------------------------------------
//#include "TAPsCUDA_UsefulDeviceFns_Def.cu"
#endif
//34567890123456789012345678901234567890123456789012345678901234567890123456789
//--+----1----+----2----+----3----+----4----+----5----+----6----+----7----+----