madrocketsci/amsculib3
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

codegenerated matrix logic

Browse Source
This commit is contained in:
madrocketsci
2026-04-13 11:42:40 -04:00
parent 978dad618f
commit 17a1c3f84a
30 changed files with 6468 additions and 1291 deletions
Download Patch File Download Diff File Expand all files Collapse all files

BIN
build_linux64/libamsculib3.linux64.a
View File

Binary file not shown.

BIN
build_linux64/objstore/cuvec2f.o
View File

Binary file not shown.

BIN
build_linux64/objstore/cuvec3f.o
View File

Binary file not shown.

BIN
build_linux64/objstore/cuvec4f.o
View File

Binary file not shown.

BIN
build_linux64/test
View File

Binary file not shown.

4
include/amsculib3/math/cuvec2f.hpp
View File

@ -45,8 +45,8 @@ namespace amscuda
__host__ __device__ cumat2f();
__host__ __device__ ~cumat2f();
__host__ __device__ cumat2f(
const float& _m00, const float& _m10,
const float& _m01, const float& _m11
const float& _m00, const float& _m01,
const float& _m10, const float& _m11
);
__host__ __device__ cumat2f(const float* data4);

85
include/amsculib3/math/cuvec3f.hpp
View File

@ -15,75 +15,72 @@ namespace amscuda
__host__ __device__ ~cuvec3f();
__host__ __device__ cuvec3f(const float &_x, const float &_y, const float &_z);
__host__ __device__ float& operator[](const int &I);
__host__ __device__ const float& operator[](const int &I) const;
__host__ __device__ cuvec3f operator+(const cuvec3f &rhs);
__host__ __device__ cuvec3f operator-(const cuvec3f &rhs);
__host__ __device__ cuvec3f operator+(const cuvec3f& rhs) const;
__host__ __device__ cuvec3f operator-(const cuvec3f& rhs) const;
__host__ __device__ cuvec3f operator*(const cuvec3f& rhs) const; //elementwise product
__host__ __device__ cuvec3f operator/(const cuvec3f& rhs) const; //elementwise division
__host__ __device__ friend cuvec3f operator*(const cuvec3f& lhs, const float &rhs);
__host__ __device__ friend cuvec3f operator/(const cuvec3f& lhs, const float &rhs);
__host__ __device__ friend cuvec3f operator*(const float& lhs, const cuvec3f &rhs);
__host__ __device__ friend cuvec3f operator/(const float& lhs, const cuvec3f &rhs);
__host__ __device__ friend cuvec3f operator-(const cuvec3f &rhs);
__host__ __device__ cuvec3f& operator+=(const cuvec3f &rhs);
__host__ __device__ cuvec3f& operator-=(const cuvec3f &rhs);
__host__ __device__ cuvec3f& operator/=(const float &rhs);
__host__ __device__ cuvec3f& operator*=(const float &rhs);
__host__ __device__ friend cuvec3f operator*(const cuvec3f& lhs, const float& rhs);
__host__ __device__ friend cuvec3f operator*(const float& lhs, const cuvec3f& rhs);
__host__ __device__ friend cuvec3f operator/(const cuvec3f& lhs, const float& rhs);
__host__ __device__ friend cuvec3f operator/(const float& lhs, const cuvec3f& rhs);
__host__ __device__ friend cuvec3f operator-(const cuvec3f& other);
__host__ __device__ cuvec3f& operator+=(const cuvec3f& rhs);
__host__ __device__ cuvec3f& operator-=(const cuvec3f& rhs);
__host__ __device__ cuvec3f& operator*=(const float& rhs);
__host__ __device__ cuvec3f& operator/=(const float& rhs);
};
class cumat3f
{
public:
float m00,m10,m20; //named references to force register use?
float m01,m11,m21; //switched to column-major-order to match GLSL/lapack
float m00,m10,m20;
float m01,m11,m21;
float m02,m12,m22;
__host__ __device__ cumat3f();
__host__ __device__ ~cumat3f();
__host__ __device__ cumat3f(
const float & _m00, const float & _m01, const float & _m02,
const float & _m10, const float & _m11, const float & _m12,
const float & _m20, const float & _m21, const float & _m22
const float& _m00, const float& _m01, const float& _m02,
const float& _m10, const float& _m11, const float& _m12,
const float& _m20, const float& _m21, const float& _m22
);
__host__ __device__ explicit cumat3f(const float *data9);
__host__ __device__ cumat3f(const float* data9);
__host__ __device__ float& operator[](const int &I);
__host__ __device__ float& operator()(const int &I, const int &J);
__host__ __device__ float& at(const int &I, const int &J);
__host__ __device__ const float& operator[](const int &I) const;
__host__ __device__ float& operator()(const int &I, const int &J);
__host__ __device__ const float& operator()(const int &I, const int &J) const;
__host__ __device__ float& at(const int &I, const int &J);
__host__ __device__ const float& at(const int &I, const int &J) const;
__host__ __device__ float* data(); //pointer to float9 representation of matrix
__host__ __device__ const float* data() const; //pointer to float9 representation of matrix
__host__ __device__ cumat3f operator+(const cumat3f &rhs);
__host__ __device__ cumat3f operator-(const cumat3f &rhs);
__host__ __device__ cumat3f operator*(const float &rhs);
__host__ __device__ cumat3f operator/(const float &rhs);
__host__ __device__ cuvec3f operator*(const cuvec3f &rhs);
__host__ __device__ cumat3f operator*(const cumat3f &rhs);
__host__ __device__ friend cumat3f operator-(const cumat3f &rhs);
__host__ __device__ cumat3f operator+(const cumat3f& rhs) const;
__host__ __device__ cumat3f operator-(const cumat3f& rhs) const;
__host__ __device__ cumat3f operator*(const cumat3f& rhs) const;
__host__ __device__ friend cumat3f operator*(const cumat3f& lhs, const float& rhs);
__host__ __device__ friend cumat3f operator/(const cumat3f& lhs, const float& rhs);
__host__ __device__ friend cumat3f operator*(const float& lhs, const cumat3f& rhs);
__host__ __device__ friend cuvec3f operator*(const cumat3f& lhs, const cuvec3f& rhs);
__host__ __device__ friend cuvec3f operator*(const cuvec3f& lhs, const cumat3f& rhs);
__host__ __device__ friend cumat3f operator-(const cumat3f& rhs);
__host__ __device__ cumat3f& operator+=(const cumat3f& rhs);
__host__ __device__ cumat3f& operator-=(const cumat3f& rhs);
__host__ __device__ cumat3f& operator*=(const float& rhs);
__host__ __device__ cumat3f& operator/=(const float& rhs);
__host__ __device__ cumat3f& operator*=(const cumat3f& rhs);
__host__ __device__ cumat3f transpose() const;
__host__ __device__ float det();
__host__ __device__ cumat3f transpose();
__host__ __device__ cumat3f inverse();
__host__ __device__ float* data(); //pointer to float[9] representation of matrix
__host__ __device__ const float* data() const; //pointer to float[9] representation of matrix
//In place operations (to save GPU register use)
__host__ __device__ cumat3f& operator+=(const cumat3f &rhs);
__host__ __device__ cumat3f& operator-=(const cumat3f &rhs);
__host__ __device__ cumat3f& operator/=(const float &rhs);
__host__ __device__ cumat3f& operator*=(const float &rhs);
__host__ __device__ cumat3f& operator*=(const cumat3f &rhs);
};
__host__ __device__ float cuvec3f_dot(const cuvec3f &a,const cuvec3f &b);

88
include/amsculib3/math/cuvec4f.hpp
View File

@ -16,77 +16,79 @@ namespace amscuda
__host__ __device__ ~cuvec4f();
__host__ __device__ cuvec4f(const float &_x, const float &_y, const float &_z, const float &_w);
__host__ __device__ float& operator[](const int &I);
__host__ __device__ const float& operator[](const int &I) const;
__host__ __device__ cuvec4f operator+(const cuvec4f &rhs);
__host__ __device__ cuvec4f operator-(const cuvec4f &rhs);
__host__ __device__ cuvec4f operator+(const cuvec4f& rhs) const;
__host__ __device__ cuvec4f operator-(const cuvec4f& rhs) const;
__host__ __device__ cuvec4f operator*(const cuvec4f& rhs) const; //elementwise product
__host__ __device__ cuvec4f operator/(const cuvec4f& rhs) const; //elementwise division
__host__ __device__ friend cuvec4f operator*(const cuvec4f& lhs, const float &rhs);
__host__ __device__ friend cuvec4f operator/(const cuvec4f& lhs, const float &rhs);
__host__ __device__ friend cuvec4f operator*(const float& lhs, const cuvec4f &rhs);
__host__ __device__ friend cuvec4f operator/(const float& lhs, const cuvec4f &rhs);
__host__ __device__ friend cuvec4f operator*(const cuvec4f& lhs, const float& rhs);
__host__ __device__ friend cuvec4f operator*(const float& lhs, const cuvec4f& rhs);
__host__ __device__ friend cuvec4f operator/(const cuvec4f& lhs, const float& rhs);
__host__ __device__ friend cuvec4f operator/(const float& lhs, const cuvec4f& rhs);
__host__ __device__ friend cuvec4f operator-(const cuvec4f& other);
__host__ __device__ friend cuvec4f operator-(const cuvec4f &rhs);
__host__ __device__ cuvec4f& operator+=(const cuvec4f &rhs);
__host__ __device__ cuvec4f& operator-=(const cuvec4f &rhs);
__host__ __device__ cuvec4f& operator/=(const float &rhs);
__host__ __device__ cuvec4f& operator*=(const float &rhs);
__host__ __device__ cuvec4f& operator+=(const cuvec4f& rhs);
__host__ __device__ cuvec4f& operator-=(const cuvec4f& rhs);
__host__ __device__ cuvec4f& operator*=(const float& rhs);
__host__ __device__ cuvec4f& operator/=(const float& rhs);
};
class cumat4f
{
public:
//float dat[16];
//__forceinline__
float m00,m10,m20,m30; //named references to force register use?
float m01,m11,m21,m31; //switched to column-major-order to match GLSL/lapack
float m00,m10,m20,m30;
float m01,m11,m21,m31;
float m02,m12,m22,m32;
float m03,m13,m23,m33;
__host__ __device__ cumat4f();
__host__ __device__ ~cumat4f();
__host__ __device__ cumat4f(
const float & _m00, const float & _m01, const float & _m02, const float & _m03,
const float & _m10, const float & _m11, const float & _m12, const float & _m13,
const float & _m20, const float & _m21, const float & _m22, const float & _m23,
const float & _m30, const float & _m31, const float & _m32, const float & _m33
const float& _m00, const float& _m01, const float& _m02, const float& _m03,
const float& _m10, const float& _m11, const float& _m12, const float& _m13,
const float& _m20, const float& _m21, const float& _m22, const float& _m23,
const float& _m30, const float& _m31, const float& _m32, const float& _m33
);
__host__ __device__ cumat4f(const float* data16);
__host__ __device__ explicit cumat4f(const float *data16);
//__forceinline__
__host__ __device__ float& operator[](const int &I);
__host__ __device__ float& operator()(const int &I, const int &J);
__host__ __device__ float& at(const int &I, const int &J);
__host__ __device__ const float& operator[](const int &I) const;
__host__ __device__ float& operator()(const int &I, const int &J);
__host__ __device__ const float& operator()(const int &I, const int &J) const;
__host__ __device__ float& at(const int &I, const int &J);
__host__ __device__ const float& at(const int &I, const int &J) const;
__host__ __device__ cumat4f operator+(const cumat4f &rhs);
__host__ __device__ cumat4f operator-(const cumat4f &rhs);
__host__ __device__ cumat4f operator*(const float &rhs);
__host__ __device__ cumat4f operator/(const float &rhs);
__host__ __device__ cuvec4f operator*(const cuvec4f &rhs);
__host__ __device__ cumat4f operator*(const cumat4f &rhs);
__host__ __device__ friend cumat4f operator-(const cumat4f &rhs);
__host__ __device__ float* data(); //pointer to float16 representation of matrix
__host__ __device__ const float* data() const; //pointer to float16 representation of matrix
__host__ __device__ cumat4f operator+(const cumat4f& rhs) const;
__host__ __device__ cumat4f operator-(const cumat4f& rhs) const;
__host__ __device__ cumat4f operator*(const cumat4f& rhs) const;
__host__ __device__ friend cumat4f operator*(const cumat4f& lhs, const float& rhs);
__host__ __device__ friend cumat4f operator/(const cumat4f& lhs, const float& rhs);
__host__ __device__ friend cumat4f operator*(const float& lhs, const cumat4f& rhs);
__host__ __device__ friend cuvec4f operator*(const cumat4f& lhs, const cuvec4f& rhs);
__host__ __device__ friend cuvec4f operator*(const cuvec4f& lhs, const cumat4f& rhs);
__host__ __device__ friend cumat4f operator-(const cumat4f& rhs);
__host__ __device__ cumat4f& operator+=(const cumat4f& rhs);
__host__ __device__ cumat4f& operator-=(const cumat4f& rhs);
__host__ __device__ cumat4f& operator*=(const float& rhs);
__host__ __device__ cumat4f& operator/=(const float& rhs);
__host__ __device__ cumat4f& operator*=(const cumat4f& rhs);
__host__ __device__ cumat4f transpose() const;
__host__ __device__ float det();
__host__ __device__ cumat4f transpose();
__host__ __device__ cumat4f inverse();
__host__ __device__ float* data(); //pointer to float[9] representation of matrix
__host__ __device__ const float* data() const; //pointer to float[9] representation of matrix
//In place operations (to save GPU register use)
__host__ __device__ cumat4f& operator+=(const cumat4f &rhs);
__host__ __device__ cumat4f& operator-=(const cumat4f &rhs);
__host__ __device__ cumat4f& operator/=(const float &rhs);
__host__ __device__ cumat4f& operator*=(const float &rhs);
__host__ __device__ cumat4f& operator*=(const cumat4f &rhs);
};
__host__ __device__ float cuvec4f_dot(cuvec4f &a, cuvec4f &b);

276
src/amsculib3/math/cuvec2f.cu
View File

@ -176,8 +176,8 @@ namespace amscuda
}
__host__ __device__ cumat2f::cumat2f(
const float& _m00, const float& _m10,
const float& _m01, const float& _m11
const float& _m00, const float& _m01,
const float& _m10, const float& _m11
)
{
m00 = _m00;
@ -256,165 +256,163 @@ namespace amscuda
return (*this)[I+2*J];
}
__host__ __device__ float* cumat2f::data()
{
return (float*)this;
}
__host__ __device__ float* cumat2f::data()
{
return (float*)this;
}
__host__ __device__ const float* cumat2f::data() const
{
return (float*)this;
}
__host__ __device__ const float* cumat2f::data() const
{
return (float*)this;
}
__host__ __device__ cumat2f cumat2f::operator+(const cumat2f& rhs) const
{
cumat2f ret;
ret.m00 = m00 + rhs.m00;
ret.m10 = m10 + rhs.m10;
__host__ __device__ cumat2f cumat2f::operator+(const cumat2f& rhs) const
{
cumat2f ret;
ret.m00 = m00 + rhs.m00;
ret.m10 = m10 + rhs.m10;
ret.m01 = m01 + rhs.m01;
ret.m11 = m11 + rhs.m11;
ret.m01 = m01 + rhs.m01;
ret.m11 = m11 + rhs.m11;
return ret;
}
return ret;
}
__host__ __device__ cumat2f cumat2f::operator-(const cumat2f& rhs) const
{
cumat2f ret;
ret.m00 = m00 - rhs.m00;
ret.m10 = m10 - rhs.m10;
__host__ __device__ cumat2f cumat2f::operator-(const cumat2f& rhs) const
{
cumat2f ret;
ret.m00 = m00 - rhs.m00;
ret.m10 = m10 - rhs.m10;
ret.m01 = m01 - rhs.m01;
ret.m11 = m11 - rhs.m11;
ret.m01 = m01 - rhs.m01;
ret.m11 = m11 - rhs.m11;
return ret;
}
return ret;
}
__host__ __device__ cumat2f cumat2f::operator*(const cumat2f& rhs) const
{
cumat2f ret; //should be zeroed in constructor
ret.m00 = m00*rhs.m00 + m01*rhs.m10;
ret.m01 = m00*rhs.m01 + m01*rhs.m11;
__host__ __device__ cumat2f cumat2f::operator*(const cumat2f& rhs) const
{
cumat2f ret; //should be zeroed in constructor
ret.m00 = m00*rhs.m00 + m01*rhs.m10;
ret.m01 = m00*rhs.m01 + m01*rhs.m11;
ret.m10 = m10*rhs.m00 + m11*rhs.m10;
ret.m11 = m10*rhs.m01 + m11*rhs.m11;
ret.m10 = m10*rhs.m00 + m11*rhs.m10;
ret.m11 = m10*rhs.m01 + m11*rhs.m11;
return ret;
}
return ret;
}
__host__ __device__ cumat2f operator*(const cumat2f& lhs, const float& rhs)
{
cumat2f ret;
ret.m00=lhs.m00*rhs;
ret.m10=lhs.m10*rhs;
ret.m01=lhs.m01*rhs;
ret.m11=lhs.m11*rhs;
return ret;
}
__host__ __device__ cumat2f operator*(const cumat2f& lhs, const float& rhs)
{
cumat2f ret;
ret.m00=lhs.m00*rhs;
ret.m10=lhs.m10*rhs;
ret.m01=lhs.m01*rhs;
ret.m11=lhs.m11*rhs;
return ret;
}
__host__ __device__ cumat2f operator/(const cumat2f& lhs, const float& rhs)
{
cumat2f ret;
ret.m00=lhs.m00/rhs;
ret.m10=lhs.m10/rhs;
ret.m01=lhs.m01/rhs;
ret.m11=lhs.m11/rhs;
return ret;
}
__host__ __device__ cumat2f operator/(const cumat2f& lhs, const float& rhs)
{
cumat2f ret;
ret.m00=lhs.m00/rhs;
ret.m10=lhs.m10/rhs;
ret.m01=lhs.m01/rhs;
ret.m11=lhs.m11/rhs;
return ret;
}
__host__ __device__ cumat2f operator*(const float& lhs, const cumat2f& rhs)
{
cumat2f ret;
ret.m00=lhs*rhs.m00;
ret.m10=lhs*rhs.m10;
ret.m01=lhs*rhs.m01;
ret.m11=lhs*rhs.m11;
return ret;
}
__host__ __device__ cumat2f operator*(const float& lhs, const cumat2f& rhs)
{
cumat2f ret;
ret.m00=lhs*rhs.m00;
ret.m10=lhs*rhs.m10;
ret.m01=lhs*rhs.m01;
ret.m11=lhs*rhs.m11;
return ret;
}
__host__ __device__ cuvec2f operator*(const cumat2f& lhs, const cuvec2f& rhs)
{
cuvec2f ret;
ret.x = lhs.m00*rhs.x + lhs.m01*rhs.y;
ret.y = lhs.m10*rhs.x + lhs.m11*rhs.y;
return ret;
}
__host__ __device__ cuvec2f operator*(const cumat2f& lhs, const cuvec2f& rhs)
{
cuvec2f ret;
ret.x = lhs.m00*rhs.x + lhs.m01*rhs.y;
ret.y = lhs.m10*rhs.x + lhs.m11*rhs.y;
return ret;
}
__host__ __device__ cuvec2f operator*(const cuvec2f& lhs, const cumat2f& rhs)
{
cuvec2f ret;
ret.x = lhs.x*rhs.m00 + lhs.y*rhs.m10;
ret.y = lhs.x*rhs.m01 + lhs.y*rhs.m11;
return ret;
}
__host__ __device__ cuvec2f operator*(const cuvec2f& lhs, const cumat2f& rhs)
{
cuvec2f ret;
ret.x = lhs.x*rhs.m00 + lhs.y*rhs.m10;
ret.y = lhs.x*rhs.m01 + lhs.y*rhs.m11;
return ret;
}
__host__ __device__ cumat2f operator-(const cumat2f& rhs)
{
cumat2f ret;
ret.m00 = -rhs.m00;
ret.m10 = -rhs.m10;
ret.m01 = -rhs.m01;
ret.m11 = -rhs.m11;
return ret;
}
__host__ __device__ cumat2f operator-(const cumat2f& rhs)
{
cumat2f ret;
ret.m00 = -rhs.m00;
ret.m10 = -rhs.m10;
ret.m01 = -rhs.m01;
ret.m11 = -rhs.m11;
return ret;
}
__host__ __device__ cumat2f& cumat2f::operator+=(const cumat2f& rhs)
{
m00 += rhs.m00;
m10 += rhs.m10;
m01 += rhs.m01;
m11 += rhs.m11;
return *this;
}
__host__ __device__ cumat2f& cumat2f::operator+=(const cumat2f& rhs)
{
m00 += rhs.m00;
m10 += rhs.m10;
m01 += rhs.m01;
m11 += rhs.m11;
return *this;
}
__host__ __device__ cumat2f& cumat2f::operator-=(const cumat2f& rhs)
{
m00 -= rhs.m00;
m10 -= rhs.m10;
m01 -= rhs.m01;
m11 -= rhs.m11;
return *this;
}
__host__ __device__ cumat2f& cumat2f::operator-=(const cumat2f& rhs)
{
m00 -= rhs.m00;
m10 -= rhs.m10;
m01 -= rhs.m01;
m11 -= rhs.m11;
return *this;
}
__host__ __device__ cumat2f& cumat2f::operator*=(const float& rhs)
{
m00 *= rhs;
m10 *= rhs;
m01 *= rhs;
m11 *= rhs;
return *this;
}
__host__ __device__ cumat2f& cumat2f::operator*=(const float& rhs)
{
m00 *= rhs;
m10 *= rhs;
m01 *= rhs;
m11 *= rhs;
return *this;
}
__host__ __device__ cumat2f& cumat2f::operator/=(const float& rhs)
{
m00 /= rhs;
m10 /= rhs;
m01 /= rhs;
m11 /= rhs;
return *this;
}
__host__ __device__ cumat2f& cumat2f::operator*=(const cumat2f& rhs)
{
cumat2f tmp = *this;
m00 = tmp.m00*rhs.m00 + tmp.m01*rhs.m10;
m01 = tmp.m00*rhs.m01 + tmp.m01*rhs.m11;
m10 = tmp.m10*rhs.m00 + tmp.m11*rhs.m10;
m11 = tmp.m10*rhs.m01 + tmp.m11*rhs.m11;
return *this;
}
__host__ __device__ cumat2f cumat2f::transpose() const
{
cumat2f ret;
ret.m00 = m00;
ret.m10 = m01;
ret.m01 = m10;
ret.m11 = m11;
return ret;
}
__host__ __device__ cumat2f& cumat2f::operator/=(const float& rhs)
{
m00 /= rhs;
m10 /= rhs;
m01 /= rhs;
m11 /= rhs;
return *this;
}
__host__ __device__ cumat2f& cumat2f::operator*=(const cumat2f& rhs)
{
cumat2f tmp = *this;
m00 = tmp.m00*rhs.m00 + tmp.m01*rhs.m10;
m01 = tmp.m00*rhs.m01 + tmp.m01*rhs.m11;
m10 = tmp.m10*rhs.m00 + tmp.m11*rhs.m10;
m11 = tmp.m10*rhs.m01 + tmp.m11*rhs.m11;
return *this;
}
__host__ __device__ cumat2f cumat2f::transpose() const
{
cumat2f ret;
ret.m00 = m00;
ret.m10 = m01;
ret.m01 = m10;
ret.m11 = m11;
return ret;
}
///////////////////
//Det and Inverse//

1050
src/amsculib3/math/cuvec3f.cu
View File

File diff suppressed because it is too large Load Diff

972
src/amsculib3/math/cuvec4f.cu
View File

File diff suppressed because it is too large Load Diff

22
test_scripts/cuvec_codegen.py
View File

@ -523,7 +523,7 @@ def mconstructor2(dim,dtype):
for I in range(0,dim):
lnsh+="\t\t"
for J in range(0,dim):
lnsh += f"const {dtype}& _m{J}{I}"
lnsh += f"const {dtype}& _m{I}{J}"
if(J+I*dim<dim*dim-1):
lnsh+=", "
lnsh+="\n"
@ -533,7 +533,7 @@ def mconstructor2(dim,dtype):
for I in range(0,dim):
lnss+="\t\t"
for J in range(0,dim):
lnss += f"const {dtype}& _m{J}{I}"
lnss += f"const {dtype}& _m{I}{J}"
if(J+I*dim<dim*dim-1):
lnss+=", "
lnss+="\n"
@ -1080,12 +1080,9 @@ def genmatrixstuff(dim,dtype):
return [lnsh,lnss]
if(__name__=="__main__"):
dim = 2
dtype = "float"
genheader = "./test_scripts/{}_codegen1.hpp".format(vtypename(dim,dtype))
gensource = "./test_scripts/{}_codegen1.cu".format(vtypename(dim,dtype))
def gencode(dim,dtype):
genheader = "./test_scripts/cuvec_codegen/{}_codegen1.hpp".format(vtypename(dim,dtype))
gensource = "./test_scripts/cuvec_codegen/{}_codegen1.cu".format(vtypename(dim,dtype))
lnsh = ""
lnss = ""
@ -1099,6 +1096,11 @@ if(__name__=="__main__"):
fwrite(genheader,lnsh)
fwrite(gensource,lnss)
if(__name__=="__main__"):
#lns = genvecheaderv1(dim,dtype)
#print(lns)
dim = 4
dtype = "int"
for I in range(2,5):
for d in ["float","double","int"]:
gencode(I,d)

412
test_scripts/cuvec_codegen/cuvec2_codegen1.cu Normal file
View File

@ -0,0 +1,412 @@
__host__ __device__ cuvec2::cuvec2()
{
x = 0; y = 0;
return;
}
__host__ __device__ cuvec2::~cuvec2()
{
x = 0; y = 0;
return;
}
__host__ __device__ cuvec2::cuvec2(const double &_x, const double &_y)
{
x = _x; y = _y;
return;
}
__host__ __device__ float& cuvec2::operator[](const int &I)
{
switch(I)
{
case 0:
return x;
case 1:
return y;
}
return x;
}
__host__ __device__ const float& cuvec2::operator[](const int &I) const
{
switch(I)
{
case 0:
return x;
case 1:
return y;
}
return x;
}
__host__ __device__ cuvec2 cuvec2::operator+(const cuvec2& rhs) const
{
cuvec2 ret;
ret.x = x + rhs.x;
ret.y = y + rhs.y;
return ret;
}
__host__ __device__ cuvec2 cuvec2::operator-(const cuvec2& rhs) const
{
cuvec2 ret;
ret.x = x - rhs.x;
ret.y = y - rhs.y;
return ret;
}
__host__ __device__ cuvec2 cuvec2::operator*(const cuvec2& rhs) const
{
//Elementwise product
cuvec2 ret;
ret.x = x * rhs.x;
ret.y = y * rhs.y;
return ret;
}
__host__ __device__ cuvec2 cuvec2::operator/(const cuvec2& rhs) const
{
//Elementwise division
cuvec2 ret;
ret.x = x / rhs.x;
ret.y = y / rhs.y;
return ret;
}
__host__ __device__ cuvec2 operator*(const cuvec2& lhs, const double& rhs)
{
cuvec2 ret;
ret.x = lhs.x*rhs;
ret.y = lhs.y*rhs;
return ret;
}
__host__ __device__ cuvec2 operator*(const double& lhs, const cuvec2& rhs)
{
cuvec2 ret;
ret.x = lhs*rhs.x;
ret.y = lhs*rhs.y;
return ret;
}
__host__ __device__ cuvec2 operator/(const cuvec2& lhs, const double& rhs)
{
cuvec2 ret;
ret.x = lhs.x/rhs;
ret.y = lhs.y/rhs;
return ret;
}
__host__ __device__ cuvec2 operator/(const double& lhs, const cuvec2& rhs)
{
cuvec2 ret;
ret.x = lhs/rhs.x;
ret.y = lhs/rhs.y;
return ret;
}
__host__ __device__ cuvec2 operator-(const cuvec2& other)
{
cuvec2 ret;
ret.x = -other.x;
ret.y = -other.y;
return ret;
}
__host__ __device__ cuvec2& cuvec2::operator+=(const cuvec2& rhs)
{
x += rhs.x;
y += rhs.y;
return *this;
}
__host__ __device__ cuvec2& cuvec2::operator-=(const cuvec2& rhs)
{
x -= rhs.x;
y -= rhs.y;
return *this;
}
__host__ __device__ cuvec2& cuvec2::operator*=(const double& rhs)
{
x *= rhs;
y *= rhs;
return *this;
}
__host__ __device__ cuvec2& cuvec2::operator/=(const double& rhs)
{
x /= rhs;
y /= rhs;
return *this;
}
//
//Matrix Header Stuff
//
__host__ __device__ cumat2::cumat2()
{
m00 = 0;
m01 = 0;
m10 = 0;
m11 = 0;
return;
}
__host__ __device__ cumat2::~cumat2()
{
//m00 = 0;
//m01 = 0;
//m10 = 0;
//m11 = 0;
return;
}
__host__ __device__ cumat2::cumat2(
const double& _m00, const double& _m01,
const double& _m10, const double& _m11
)
{
m00 = _m00;
m10 = _m10;
m01 = _m01;
m11 = _m11;
return;
}
__host__ __device__ cumat2::cumat2(const double* data4)
{
m00 = data4[0];
m10 = data4[1];
m01 = data4[2];
m11 = data4[3];
return;
}
__host__ __device__ float& cumat2::operator[](const int &I)
{
switch(I)
{
case 0:
return m00;
case 1:
return m10;
case 2:
return m01;
case 3:
return m11;
}
return m00;
}
__host__ __device__ const float& cumat2::operator[](const int &I) const
{
switch(I)
{
case 0:
return m00;
case 1:
return m10;
case 2:
return m01;
case 3:
return m11;
}
return m00;
}
__host__ __device__ float& cumat2::operator()(const int &I, const int &J)
{
return (*this)[I+2*J];
}
__host__ __device__ const float& cumat2::operator()(const int &I, const int &J) const
{
return (*this)[I+2*J];
}
__host__ __device__ float& cumat2::at(const int &I, const int &J)
{
return (*this)[I+2*J];
}
__host__ __device__ const float& cumat2::at(const int &I, const int &J) const
{
return (*this)[I+2*J];
}
__host__ __device__ double* cumat2::data()
{
return (double*)this;
}
__host__ __device__ const double* cumat2::data() const
{
return (double*)this;
}
__host__ __device__ cumat2 cumat2::operator+(const cumat2& rhs) const
{
cumat2 ret;
ret.m00 = m00 + rhs.m00;
ret.m10 = m10 + rhs.m10;
ret.m01 = m01 + rhs.m01;
ret.m11 = m11 + rhs.m11;
return ret;
}
__host__ __device__ cumat2 cumat2::operator-(const cumat2& rhs) const
{
cumat2 ret;
ret.m00 = m00 - rhs.m00;
ret.m10 = m10 - rhs.m10;
ret.m01 = m01 - rhs.m01;
ret.m11 = m11 - rhs.m11;
return ret;
}
__host__ __device__ cumat2 cumat2::operator*(const cumat2& rhs) const
{
cumat2 ret; //should be zeroed in constructor
ret.m00 = m00*rhs.m00 + m01*rhs.m10;
ret.m01 = m00*rhs.m01 + m01*rhs.m11;
ret.m10 = m10*rhs.m00 + m11*rhs.m10;
ret.m11 = m10*rhs.m01 + m11*rhs.m11;
return ret;
}
__host__ __device__ cumat2 operator*(const cumat2& lhs, const double& rhs)
{
cumat2 ret;
ret.m00=lhs.m00*rhs;
ret.m10=lhs.m10*rhs;
ret.m01=lhs.m01*rhs;
ret.m11=lhs.m11*rhs;
return ret;
}
__host__ __device__ cumat2 operator/(const cumat2& lhs, const double& rhs)
{
cumat2 ret;
ret.m00=lhs.m00/rhs;
ret.m10=lhs.m10/rhs;
ret.m01=lhs.m01/rhs;
ret.m11=lhs.m11/rhs;
return ret;
}
__host__ __device__ cumat2 operator*(const double& lhs, const cumat2& rhs)
{
cumat2 ret;
ret.m00=lhs*rhs.m00;
ret.m10=lhs*rhs.m10;
ret.m01=lhs*rhs.m01;
ret.m11=lhs*rhs.m11;
return ret;
}
__host__ __device__ cuvec2 operator*(const cumat2& lhs, const cuvec2& rhs)
{
cuvec2 ret;
ret.x = lhs.m00*rhs.x + lhs.m01*rhs.y;
ret.y = lhs.m10*rhs.x + lhs.m11*rhs.y;
return ret;
}
__host__ __device__ cuvec2 operator*(const cuvec2& lhs, const cumat2& rhs)
{
cuvec2 ret;
ret.x = lhs.x*rhs.m00 + lhs.y*rhs.m10;
ret.y = lhs.x*rhs.m01 + lhs.y*rhs.m11;
return ret;
}
__host__ __device__ cumat2 operator-(const cumat2& rhs)
{
cumat2 ret;
ret.m00 = -rhs.m00;
ret.m10 = -rhs.m10;
ret.m01 = -rhs.m01;
ret.m11 = -rhs.m11;
return ret;
}
__host__ __device__ cumat2& cumat2::operator+=(const cumat2& rhs)
{
m00 += rhs.m00;
m10 += rhs.m10;
m01 += rhs.m01;
m11 += rhs.m11;
return *this;
}
__host__ __device__ cumat2& cumat2::operator-=(const cumat2& rhs)
{
m00 -= rhs.m00;
m10 -= rhs.m10;
m01 -= rhs.m01;
m11 -= rhs.m11;
return *this;
}
__host__ __device__ cumat2& cumat2::operator*=(const double& rhs)
{
m00 *= rhs;
m10 *= rhs;
m01 *= rhs;
m11 *= rhs;
return *this;
}
__host__ __device__ cumat2& cumat2::operator/=(const double& rhs)
{
m00 /= rhs;
m10 /= rhs;
m01 /= rhs;
m11 /= rhs;
return *this;
}
__host__ __device__ cumat2& cumat2::operator*=(const cumat2& rhs)
{
cumat2 tmp = *this;
m00 = tmp.m00*rhs.m00 + tmp.m01*rhs.m10;
m01 = tmp.m00*rhs.m01 + tmp.m01*rhs.m11;
m10 = tmp.m10*rhs.m00 + tmp.m11*rhs.m10;
m11 = tmp.m10*rhs.m01 + tmp.m11*rhs.m11;
return *this;
}
__host__ __device__ cumat2 cumat2::transpose() const
{
cumat2 ret;
ret.m00 = m00;
ret.m10 = m01;
ret.m01 = m10;
ret.m11 = m11;
return ret;
}

54
test_scripts/cuvec_codegen/cuvec2_codegen1.hpp Normal file
View File

@ -0,0 +1,54 @@
__host__ __device__ cuvec2();
__host__ __device__ ~cuvec2();
__host__ __device__ cuvec2(const double &_x, const double &_y);
__host__ __device__ float& operator[](const int &I);
__host__ __device__ const float& operator[](const int &I) const;
__host__ __device__ cuvec2 operator+(const cuvec2& rhs) const;
__host__ __device__ cuvec2 operator-(const cuvec2& rhs) const;
__host__ __device__ cuvec2 operator*(const cuvec2& rhs) const; //elementwise product
__host__ __device__ cuvec2 operator/(const cuvec2& rhs) const; //elementwise division
__host__ __device__ friend cuvec2 operator*(const cuvec2& lhs, const double& rhs);
__host__ __device__ friend cuvec2 operator*(const double& lhs, const cuvec2& rhs);
__host__ __device__ friend cuvec2 operator/(const cuvec2& lhs, const double& rhs);
__host__ __device__ friend cuvec2 operator/(const double& lhs, const cuvec2& rhs);
__host__ __device__ friend cuvec2 operator-(const cuvec2& other);
__host__ __device__ cuvec2& operator+=(const cuvec2& rhs);
__host__ __device__ cuvec2& operator-=(const cuvec2& rhs);
__host__ __device__ cuvec2& operator*=(const double& rhs);
__host__ __device__ cuvec2& operator/=(const double& rhs);
//
//Matrix Header Stuff
//
double m00,m10;
double m01,m11;
__host__ __device__ cumat2();
__host__ __device__ ~cumat2();
__host__ __device__ cumat2(
const double& _m00, const double& _m01,
const double& _m10, const double& _m11
);
__host__ __device__ cumat2(const double* data4);
__host__ __device__ float& operator[](const int &I);
__host__ __device__ const float& operator[](const int &I) const;
__host__ __device__ float& operator()(const int &I, const int &J);
__host__ __device__ const float& operator()(const int &I, const int &J) const;
__host__ __device__ float& at(const int &I, const int &J);
__host__ __device__ const float& at(const int &I, const int &J) const;
__host__ __device__ double* data(); //pointer to float4 representation of matrix
__host__ __device__ const double* data() const; //pointer to float4 representation of matrix
__host__ __device__ cumat2 operator+(const cumat2& rhs) const;
__host__ __device__ cumat2 operator-(const cumat2& rhs) const;
__host__ __device__ cumat2 operator*(const cumat2& rhs) const;
__host__ __device__ friend cumat2 operator*(const cumat2& lhs, const double& rhs);
__host__ __device__ friend cumat2 operator/(const cumat2& lhs, const double& rhs);
__host__ __device__ friend cumat2 operator*(const double& lhs, const cumat2& rhs);
__host__ __device__ friend cuvec2 operator*(const cumat2& lhs, const cuvec2& rhs);
__host__ __device__ friend cuvec2 operator*(const cuvec2& lhs, const cumat2& rhs);
__host__ __device__ friend cumat2 operator-(const cumat2& rhs);
__host__ __device__ cumat2& operator+=(const cumat2& rhs);
__host__ __device__ cumat2& operator-=(const cumat2& rhs);
__host__ __device__ cumat2& operator*=(const double& rhs);
__host__ __device__ cumat2& operator/=(const double& rhs);
__host__ __device__ cumat2& operator*=(const cumat2& rhs);
__host__ __device__ cumat2 transpose() const;

4
test_scripts/cuvec2f_codegen1.cu → test_scripts/cuvec_codegen/cuvec2f_codegen1.cu
View File

@ -172,8 +172,8 @@
}
__host__ __device__ cumat2f::cumat2f(
const float& _m00, const float& _m10,
const float& _m01, const float& _m11
const float& _m00, const float& _m01,
const float& _m10, const float& _m11
)
{
m00 = _m00;

4
test_scripts/cuvec2f_codegen1.hpp → test_scripts/cuvec_codegen/cuvec2f_codegen1.hpp
View File

@ -25,8 +25,8 @@ float m01,m11;
__host__ __device__ cumat2f();
__host__ __device__ ~cumat2f();
__host__ __device__ cumat2f(
const float& _m00, const float& _m10,
const float& _m01, const float& _m11
const float& _m00, const float& _m01,
const float& _m10, const float& _m11
);
__host__ __device__ cumat2f(const float* data4);
__host__ __device__ float& operator[](const int &I);

412
test_scripts/cuvec_codegen/cuvec2i_codegen1.cu Normal file
View File

@ -0,0 +1,412 @@
__host__ __device__ cuvec2i::cuvec2i()
{
x = 0; y = 0;
return;
}
__host__ __device__ cuvec2i::~cuvec2i()
{
x = 0; y = 0;
return;
}
__host__ __device__ cuvec2i::cuvec2i(const int &_x, const int &_y)
{
x = _x; y = _y;
return;
}
__host__ __device__ float& cuvec2i::operator[](const int &I)
{
switch(I)
{
case 0:
return x;
case 1:
return y;
}
return x;
}
__host__ __device__ const float& cuvec2i::operator[](const int &I) const
{
switch(I)
{
case 0:
return x;
case 1:
return y;
}
return x;
}
__host__ __device__ cuvec2i cuvec2i::operator+(const cuvec2i& rhs) const
{
cuvec2i ret;
ret.x = x + rhs.x;
ret.y = y + rhs.y;
return ret;
}
__host__ __device__ cuvec2i cuvec2i::operator-(const cuvec2i& rhs) const
{
cuvec2i ret;
ret.x = x - rhs.x;
ret.y = y - rhs.y;
return ret;
}
__host__ __device__ cuvec2i cuvec2i::operator*(const cuvec2i& rhs) const
{
//Elementwise product
cuvec2i ret;
ret.x = x * rhs.x;
ret.y = y * rhs.y;
return ret;
}
__host__ __device__ cuvec2i cuvec2i::operator/(const cuvec2i& rhs) const
{
//Elementwise division
cuvec2i ret;
ret.x = x / rhs.x;
ret.y = y / rhs.y;
return ret;
}
__host__ __device__ cuvec2i operator*(const cuvec2i& lhs, const int& rhs)
{
cuvec2i ret;
ret.x = lhs.x*rhs;
ret.y = lhs.y*rhs;
return ret;
}
__host__ __device__ cuvec2i operator*(const int& lhs, const cuvec2i& rhs)
{
cuvec2i ret;
ret.x = lhs*rhs.x;
ret.y = lhs*rhs.y;
return ret;
}
__host__ __device__ cuvec2i operator/(const cuvec2i& lhs, const int& rhs)
{
cuvec2i ret;
ret.x = lhs.x/rhs;
ret.y = lhs.y/rhs;
return ret;
}
__host__ __device__ cuvec2i operator/(const int& lhs, const cuvec2i& rhs)
{
cuvec2i ret;
ret.x = lhs/rhs.x;
ret.y = lhs/rhs.y;
return ret;
}
__host__ __device__ cuvec2i operator-(const cuvec2i& other)
{
cuvec2i ret;
ret.x = -other.x;
ret.y = -other.y;
return ret;
}
__host__ __device__ cuvec2i& cuvec2i::operator+=(const cuvec2i& rhs)
{
x += rhs.x;
y += rhs.y;
return *this;
}
__host__ __device__ cuvec2i& cuvec2i::operator-=(const cuvec2i& rhs)
{
x -= rhs.x;
y -= rhs.y;
return *this;
}
__host__ __device__ cuvec2i& cuvec2i::operator*=(const int& rhs)
{
x *= rhs;
y *= rhs;
return *this;
}
__host__ __device__ cuvec2i& cuvec2i::operator/=(const int& rhs)
{
x /= rhs;
y /= rhs;
return *this;
}
//
//Matrix Header Stuff
//
__host__ __device__ cumat2i::cumat2i()
{
m00 = 0;
m01 = 0;
m10 = 0;
m11 = 0;
return;
}
__host__ __device__ cumat2i::~cumat2i()
{
//m00 = 0;
//m01 = 0;
//m10 = 0;
//m11 = 0;
return;
}
__host__ __device__ cumat2i::cumat2i(
const int& _m00, const int& _m01,
const int& _m10, const int& _m11
)
{
m00 = _m00;
m10 = _m10;
m01 = _m01;
m11 = _m11;
return;
}
__host__ __device__ cumat2i::cumat2i(const int* data4)
{
m00 = data4[0];
m10 = data4[1];
m01 = data4[2];
m11 = data4[3];
return;
}
__host__ __device__ float& cumat2i::operator[](const int &I)
{
switch(I)
{
case 0:
return m00;
case 1:
return m10;
case 2:
return m01;
case 3:
return m11;
}
return m00;
}
__host__ __device__ const float& cumat2i::operator[](const int &I) const
{
switch(I)
{
case 0:
return m00;
case 1:
return m10;
case 2:
return m01;
case 3:
return m11;
}
return m00;
}
__host__ __device__ float& cumat2i::operator()(const int &I, const int &J)
{
return (*this)[I+2*J];
}
__host__ __device__ const float& cumat2i::operator()(const int &I, const int &J) const
{
return (*this)[I+2*J];
}
__host__ __device__ float& cumat2i::at(const int &I, const int &J)
{
return (*this)[I+2*J];
}
__host__ __device__ const float& cumat2i::at(const int &I, const int &J) const
{
return (*this)[I+2*J];
}
__host__ __device__ int* cumat2i::data()
{
return (int*)this;
}
__host__ __device__ const int* cumat2i::data() const
{
return (int*)this;
}
__host__ __device__ cumat2i cumat2i::operator+(const cumat2i& rhs) const
{
cumat2i ret;
ret.m00 = m00 + rhs.m00;
ret.m10 = m10 + rhs.m10;
ret.m01 = m01 + rhs.m01;
ret.m11 = m11 + rhs.m11;
return ret;
}
__host__ __device__ cumat2i cumat2i::operator-(const cumat2i& rhs) const
{
cumat2i ret;
ret.m00 = m00 - rhs.m00;
ret.m10 = m10 - rhs.m10;
ret.m01 = m01 - rhs.m01;
ret.m11 = m11 - rhs.m11;
return ret;
}
__host__ __device__ cumat2i cumat2i::operator*(const cumat2i& rhs) const
{
cumat2i ret; //should be zeroed in constructor
ret.m00 = m00*rhs.m00 + m01*rhs.m10;
ret.m01 = m00*rhs.m01 + m01*rhs.m11;
ret.m10 = m10*rhs.m00 + m11*rhs.m10;
ret.m11 = m10*rhs.m01 + m11*rhs.m11;
return ret;
}
__host__ __device__ cumat2i operator*(const cumat2i& lhs, const int& rhs)
{
cumat2i ret;
ret.m00=lhs.m00*rhs;
ret.m10=lhs.m10*rhs;
ret.m01=lhs.m01*rhs;
ret.m11=lhs.m11*rhs;
return ret;
}
__host__ __device__ cumat2i operator/(const cumat2i& lhs, const int& rhs)
{
cumat2i ret;
ret.m00=lhs.m00/rhs;
ret.m10=lhs.m10/rhs;
ret.m01=lhs.m01/rhs;
ret.m11=lhs.m11/rhs;
return ret;
}
__host__ __device__ cumat2i operator*(const int& lhs, const cumat2i& rhs)
{
cumat2i ret;
ret.m00=lhs*rhs.m00;
ret.m10=lhs*rhs.m10;
ret.m01=lhs*rhs.m01;
ret.m11=lhs*rhs.m11;
return ret;
}
__host__ __device__ cuvec2i operator*(const cumat2i& lhs, const cuvec2i& rhs)
{
cuvec2i ret;
ret.x = lhs.m00*rhs.x + lhs.m01*rhs.y;
ret.y = lhs.m10*rhs.x + lhs.m11*rhs.y;
return ret;
}
__host__ __device__ cuvec2i operator*(const cuvec2i& lhs, const cumat2i& rhs)
{
cuvec2i ret;
ret.x = lhs.x*rhs.m00 + lhs.y*rhs.m10;
ret.y = lhs.x*rhs.m01 + lhs.y*rhs.m11;
return ret;
}
__host__ __device__ cumat2i operator-(const cumat2i& rhs)
{
cumat2i ret;
ret.m00 = -rhs.m00;
ret.m10 = -rhs.m10;
ret.m01 = -rhs.m01;
ret.m11 = -rhs.m11;
return ret;
}
__host__ __device__ cumat2i& cumat2i::operator+=(const cumat2i& rhs)
{
m00 += rhs.m00;
m10 += rhs.m10;
m01 += rhs.m01;
m11 += rhs.m11;
return *this;
}
__host__ __device__ cumat2i& cumat2i::operator-=(const cumat2i& rhs)
{
m00 -= rhs.m00;
m10 -= rhs.m10;
m01 -= rhs.m01;
m11 -= rhs.m11;
return *this;
}
__host__ __device__ cumat2i& cumat2i::operator*=(const int& rhs)
{
m00 *= rhs;
m10 *= rhs;
m01 *= rhs;
m11 *= rhs;
return *this;
}
__host__ __device__ cumat2i& cumat2i::operator/=(const int& rhs)
{
m00 /= rhs;
m10 /= rhs;
m01 /= rhs;
m11 /= rhs;
return *this;
}
__host__ __device__ cumat2i& cumat2i::operator*=(const cumat2i& rhs)
{
cumat2i tmp = *this;
m00 = tmp.m00*rhs.m00 + tmp.m01*rhs.m10;
m01 = tmp.m00*rhs.m01 + tmp.m01*rhs.m11;
m10 = tmp.m10*rhs.m00 + tmp.m11*rhs.m10;
m11 = tmp.m10*rhs.m01 + tmp.m11*rhs.m11;
return *this;
}
__host__ __device__ cumat2i cumat2i::transpose() const
{
cumat2i ret;
ret.m00 = m00;
ret.m10 = m01;
ret.m01 = m10;
ret.m11 = m11;
return ret;
}

54
test_scripts/cuvec_codegen/cuvec2i_codegen1.hpp Normal file
View File

@ -0,0 +1,54 @@
__host__ __device__ cuvec2i();
__host__ __device__ ~cuvec2i();
__host__ __device__ cuvec2i(const int &_x, const int &_y);
__host__ __device__ float& operator[](const int &I);
__host__ __device__ const float& operator[](const int &I) const;
__host__ __device__ cuvec2i operator+(const cuvec2i& rhs) const;
__host__ __device__ cuvec2i operator-(const cuvec2i& rhs) const;
__host__ __device__ cuvec2i operator*(const cuvec2i& rhs) const; //elementwise product
__host__ __device__ cuvec2i operator/(const cuvec2i& rhs) const; //elementwise division
__host__ __device__ friend cuvec2i operator*(const cuvec2i& lhs, const int& rhs);
__host__ __device__ friend cuvec2i operator*(const int& lhs, const cuvec2i& rhs);
__host__ __device__ friend cuvec2i operator/(const cuvec2i& lhs, const int& rhs);
__host__ __device__ friend cuvec2i operator/(const int& lhs, const cuvec2i& rhs);
__host__ __device__ friend cuvec2i operator-(const cuvec2i& other);
__host__ __device__ cuvec2i& operator+=(const cuvec2i& rhs);
__host__ __device__ cuvec2i& operator-=(const cuvec2i& rhs);
__host__ __device__ cuvec2i& operator*=(const int& rhs);
__host__ __device__ cuvec2i& operator/=(const int& rhs);
//
//Matrix Header Stuff
//
int m00,m10;
int m01,m11;
__host__ __device__ cumat2i();
__host__ __device__ ~cumat2i();
__host__ __device__ cumat2i(
const int& _m00, const int& _m01,
const int& _m10, const int& _m11
);
__host__ __device__ cumat2i(const int* data4);
__host__ __device__ float& operator[](const int &I);
__host__ __device__ const float& operator[](const int &I) const;
__host__ __device__ float& operator()(const int &I, const int &J);
__host__ __device__ const float& operator()(const int &I, const int &J) const;
__host__ __device__ float& at(const int &I, const int &J);
__host__ __device__ const float& at(const int &I, const int &J) const;
__host__ __device__ int* data(); //pointer to float4 representation of matrix
__host__ __device__ const int* data() const; //pointer to float4 representation of matrix
__host__ __device__ cumat2i operator+(const cumat2i& rhs) const;
__host__ __device__ cumat2i operator-(const cumat2i& rhs) const;
__host__ __device__ cumat2i operator*(const cumat2i& rhs) const;
__host__ __device__ friend cumat2i operator*(const cumat2i& lhs, const int& rhs);
__host__ __device__ friend cumat2i operator/(const cumat2i& lhs, const int& rhs);
__host__ __device__ friend cumat2i operator*(const int& lhs, const cumat2i& rhs);
__host__ __device__ friend cuvec2i operator*(const cumat2i& lhs, const cuvec2i& rhs);
__host__ __device__ friend cuvec2i operator*(const cuvec2i& lhs, const cumat2i& rhs);
__host__ __device__ friend cumat2i operator-(const cumat2i& rhs);
__host__ __device__ cumat2i& operator+=(const cumat2i& rhs);
__host__ __device__ cumat2i& operator-=(const cumat2i& rhs);
__host__ __device__ cumat2i& operator*=(const int& rhs);
__host__ __device__ cumat2i& operator/=(const int& rhs);
__host__ __device__ cumat2i& operator*=(const cumat2i& rhs);
__host__ __device__ cumat2i transpose() const;

544
test_scripts/cuvec_codegen/cuvec3_codegen1.cu Normal file
View File

@ -0,0 +1,544 @@
__host__ __device__ cuvec3::cuvec3()
{
x = 0; y = 0; z = 0;
return;
}
__host__ __device__ cuvec3::~cuvec3()
{
x = 0; y = 0; z = 0;
return;
}
__host__ __device__ cuvec3::cuvec3(const double &_x, const double &_y, const double &_z)
{
x = _x; y = _y; z = _z;
return;
}
__host__ __device__ float& cuvec3::operator[](const int &I)
{
switch(I)
{
case 0:
return x;
case 1:
return y;
case 2:
return z;
}
return x;
}
__host__ __device__ const float& cuvec3::operator[](const int &I) const
{
switch(I)
{
case 0:
return x;
case 1:
return y;
case 2:
return z;
}
return x;
}
__host__ __device__ cuvec3 cuvec3::operator+(const cuvec3& rhs) const
{
cuvec3 ret;
ret.x = x + rhs.x;
ret.y = y + rhs.y;
ret.z = z + rhs.z;
return ret;
}
__host__ __device__ cuvec3 cuvec3::operator-(const cuvec3& rhs) const
{
cuvec3 ret;
ret.x = x - rhs.x;
ret.y = y - rhs.y;
ret.z = z - rhs.z;
return ret;
}
__host__ __device__ cuvec3 cuvec3::operator*(const cuvec3& rhs) const
{
//Elementwise product
cuvec3 ret;
ret.x = x * rhs.x;
ret.y = y * rhs.y;
ret.z = z * rhs.z;
return ret;
}
__host__ __device__ cuvec3 cuvec3::operator/(const cuvec3& rhs) const
{
//Elementwise division
cuvec3 ret;
ret.x = x / rhs.x;
ret.y = y / rhs.y;
ret.z = z / rhs.z;
return ret;
}
__host__ __device__ cuvec3 operator*(const cuvec3& lhs, const double& rhs)
{
cuvec3 ret;
ret.x = lhs.x*rhs;
ret.y = lhs.y*rhs;
ret.z = lhs.z*rhs;
return ret;
}
__host__ __device__ cuvec3 operator*(const double& lhs, const cuvec3& rhs)
{
cuvec3 ret;
ret.x = lhs*rhs.x;
ret.y = lhs*rhs.y;
ret.z = lhs*rhs.z;
return ret;
}
__host__ __device__ cuvec3 operator/(const cuvec3& lhs, const double& rhs)
{
cuvec3 ret;
ret.x = lhs.x/rhs;
ret.y = lhs.y/rhs;
ret.z = lhs.z/rhs;
return ret;
}
__host__ __device__ cuvec3 operator/(const double& lhs, const cuvec3& rhs)
{
cuvec3 ret;
ret.x = lhs/rhs.x;
ret.y = lhs/rhs.y;
ret.z = lhs/rhs.z;
return ret;
}
__host__ __device__ cuvec3 operator-(const cuvec3& other)
{
cuvec3 ret;
ret.x = -other.x;
ret.y = -other.y;
ret.z = -other.z;
return ret;
}
__host__ __device__ cuvec3& cuvec3::operator+=(const cuvec3& rhs)
{
x += rhs.x;
y += rhs.y;
z += rhs.z;
return *this;
}
__host__ __device__ cuvec3& cuvec3::operator-=(const cuvec3& rhs)
{
x -= rhs.x;
y -= rhs.y;
z -= rhs.z;
return *this;
}
__host__ __device__ cuvec3& cuvec3::operator*=(const double& rhs)
{
x *= rhs;
y *= rhs;
z *= rhs;
return *this;
}
__host__ __device__ cuvec3& cuvec3::operator/=(const double& rhs)
{
x /= rhs;
y /= rhs;
z /= rhs;
return *this;
}
//
//Matrix Header Stuff
//
__host__ __device__ cumat3::cumat3()
{
m00 = 0;
m01 = 0;
m02 = 0;
m10 = 0;
m11 = 0;
m12 = 0;
m20 = 0;
m21 = 0;
m22 = 0;
return;
}
__host__ __device__ cumat3::~cumat3()
{
//m00 = 0;
//m01 = 0;
//m02 = 0;
//m10 = 0;
//m11 = 0;
//m12 = 0;
//m20 = 0;
//m21 = 0;
//m22 = 0;
return;
}
__host__ __device__ cumat3::cumat3(
const double& _m00, const double& _m01, const double& _m02,
const double& _m10, const double& _m11, const double& _m12,
const double& _m20, const double& _m21, const double& _m22
)
{
m00 = _m00;
m10 = _m10;
m20 = _m20;
m01 = _m01;
m11 = _m11;
m21 = _m21;
m02 = _m02;
m12 = _m12;
m22 = _m22;
return;
}
__host__ __device__ cumat3::cumat3(const double* data9)
{
m00 = data9[0];
m10 = data9[1];
m20 = data9[2];
m01 = data9[3];
m11 = data9[4];
m21 = data9[5];
m02 = data9[6];
m12 = data9[7];
m22 = data9[8];
return;
}
__host__ __device__ float& cumat3::operator[](const int &I)
{
switch(I)
{
case 0:
return m00;
case 1:
return m10;
case 2:
return m20;
case 3:
return m01;
case 4:
return m11;
case 5:
return m21;
case 6:
return m02;
case 7:
return m12;
case 8:
return m22;
}
return m00;
}
__host__ __device__ const float& cumat3::operator[](const int &I) const
{
switch(I)
{
case 0:
return m00;
case 1:
return m10;
case 2:
return m20;
case 3:
return m01;
case 4:
return m11;
case 5:
return m21;
case 6:
return m02;
case 7:
return m12;
case 8:
return m22;
}
return m00;
}
__host__ __device__ float& cumat3::operator()(const int &I, const int &J)
{
return (*this)[I+3*J];
}
__host__ __device__ const float& cumat3::operator()(const int &I, const int &J) const
{
return (*this)[I+3*J];
}
__host__ __device__ float& cumat3::at(const int &I, const int &J)
{
return (*this)[I+3*J];
}
__host__ __device__ const float& cumat3::at(const int &I, const int &J) const
{
return (*this)[I+3*J];
}
__host__ __device__ double* cumat3::data()
{
return (double*)this;
}
__host__ __device__ const double* cumat3::data() const
{
return (double*)this;
}
__host__ __device__ cumat3 cumat3::operator+(const cumat3& rhs) const
{
cumat3 ret;
ret.m00 = m00 + rhs.m00;
ret.m10 = m10 + rhs.m10;
ret.m20 = m20 + rhs.m20;
ret.m01 = m01 + rhs.m01;
ret.m11 = m11 + rhs.m11;
ret.m21 = m21 + rhs.m21;
ret.m02 = m02 + rhs.m02;
ret.m12 = m12 + rhs.m12;
ret.m22 = m22 + rhs.m22;
return ret;
}
__host__ __device__ cumat3 cumat3::operator-(const cumat3& rhs) const
{
cumat3 ret;
ret.m00 = m00 - rhs.m00;
ret.m10 = m10 - rhs.m10;
ret.m20 = m20 - rhs.m20;
ret.m01 = m01 - rhs.m01;
ret.m11 = m11 - rhs.m11;
ret.m21 = m21 - rhs.m21;
ret.m02 = m02 - rhs.m02;
ret.m12 = m12 - rhs.m12;
ret.m22 = m22 - rhs.m22;
return ret;
}
__host__ __device__ cumat3 cumat3::operator*(const cumat3& rhs) const
{
cumat3 ret; //should be zeroed in constructor
ret.m00 = m00*rhs.m00 + m01*rhs.m10 + m02*rhs.m20;
ret.m01 = m00*rhs.m01 + m01*rhs.m11 + m02*rhs.m21;
ret.m02 = m00*rhs.m02 + m01*rhs.m12 + m02*rhs.m22;
ret.m10 = m10*rhs.m00 + m11*rhs.m10 + m12*rhs.m20;
ret.m11 = m10*rhs.m01 + m11*rhs.m11 + m12*rhs.m21;
ret.m12 = m10*rhs.m02 + m11*rhs.m12 + m12*rhs.m22;
ret.m20 = m20*rhs.m00 + m21*rhs.m10 + m22*rhs.m20;
ret.m21 = m20*rhs.m01 + m21*rhs.m11 + m22*rhs.m21;
ret.m22 = m20*rhs.m02 + m21*rhs.m12 + m22*rhs.m22;
return ret;
}
__host__ __device__ cumat3 operator*(const cumat3& lhs, const double& rhs)
{
cumat3 ret;
ret.m00=lhs.m00*rhs;
ret.m10=lhs.m10*rhs;
ret.m20=lhs.m20*rhs;
ret.m01=lhs.m01*rhs;
ret.m11=lhs.m11*rhs;
ret.m21=lhs.m21*rhs;
ret.m02=lhs.m02*rhs;
ret.m12=lhs.m12*rhs;
ret.m22=lhs.m22*rhs;
return ret;
}
__host__ __device__ cumat3 operator/(const cumat3& lhs, const double& rhs)
{
cumat3 ret;
ret.m00=lhs.m00/rhs;
ret.m10=lhs.m10/rhs;
ret.m20=lhs.m20/rhs;
ret.m01=lhs.m01/rhs;
ret.m11=lhs.m11/rhs;
ret.m21=lhs.m21/rhs;
ret.m02=lhs.m02/rhs;
ret.m12=lhs.m12/rhs;
ret.m22=lhs.m22/rhs;
return ret;
}
__host__ __device__ cumat3 operator*(const double& lhs, const cumat3& rhs)
{
cumat3 ret;
ret.m00=lhs*rhs.m00;
ret.m10=lhs*rhs.m10;
ret.m20=lhs*rhs.m20;
ret.m01=lhs*rhs.m01;
ret.m11=lhs*rhs.m11;
ret.m21=lhs*rhs.m21;
ret.m02=lhs*rhs.m02;
ret.m12=lhs*rhs.m12;
ret.m22=lhs*rhs.m22;
return ret;
}
__host__ __device__ cuvec3 operator*(const cumat3& lhs, const cuvec3& rhs)
{
cuvec3 ret;
ret.x = lhs.m00*rhs.x + lhs.m01*rhs.y + lhs.m02*rhs.z;
ret.y = lhs.m10*rhs.x + lhs.m11*rhs.y + lhs.m12*rhs.z;
ret.z = lhs.m20*rhs.x + lhs.m21*rhs.y + lhs.m22*rhs.z;
return ret;
}
__host__ __device__ cuvec3 operator*(const cuvec3& lhs, const cumat3& rhs)
{
cuvec3 ret;
ret.x = lhs.x*rhs.m00 + lhs.y*rhs.m10 + lhs.z*rhs.m20;
ret.y = lhs.x*rhs.m01 + lhs.y*rhs.m11 + lhs.z*rhs.m21;
ret.z = lhs.x*rhs.m02 + lhs.y*rhs.m12 + lhs.z*rhs.m22;
return ret;
}
__host__ __device__ cumat3 operator-(const cumat3& rhs)
{
cumat3 ret;
ret.m00 = -rhs.m00;
ret.m10 = -rhs.m10;
ret.m20 = -rhs.m20;
ret.m01 = -rhs.m01;
ret.m11 = -rhs.m11;
ret.m21 = -rhs.m21;
ret.m02 = -rhs.m02;
ret.m12 = -rhs.m12;
ret.m22 = -rhs.m22;
return ret;
}
__host__ __device__ cumat3& cumat3::operator+=(const cumat3& rhs)
{
m00 += rhs.m00;
m10 += rhs.m10;
m20 += rhs.m20;
m01 += rhs.m01;
m11 += rhs.m11;
m21 += rhs.m21;
m02 += rhs.m02;
m12 += rhs.m12;
m22 += rhs.m22;
return *this;
}
__host__ __device__ cumat3& cumat3::operator-=(const cumat3& rhs)
{
m00 -= rhs.m00;
m10 -= rhs.m10;
m20 -= rhs.m20;
m01 -= rhs.m01;
m11 -= rhs.m11;
m21 -= rhs.m21;
m02 -= rhs.m02;
m12 -= rhs.m12;
m22 -= rhs.m22;
return *this;
}
__host__ __device__ cumat3& cumat3::operator*=(const double& rhs)
{
m00 *= rhs;
m10 *= rhs;
m20 *= rhs;
m01 *= rhs;
m11 *= rhs;
m21 *= rhs;
m02 *= rhs;
m12 *= rhs;
m22 *= rhs;
return *this;
}
__host__ __device__ cumat3& cumat3::operator/=(const double& rhs)
{
m00 /= rhs;
m10 /= rhs;
m20 /= rhs;
m01 /= rhs;
m11 /= rhs;
m21 /= rhs;
m02 /= rhs;
m12 /= rhs;
m22 /= rhs;
return *this;
}
__host__ __device__ cumat3& cumat3::operator*=(const cumat3& rhs)
{
cumat3 tmp = *this;
m00 = tmp.m00*rhs.m00 + tmp.m01*rhs.m10 + tmp.m02*rhs.m20;
m01 = tmp.m00*rhs.m01 + tmp.m01*rhs.m11 + tmp.m02*rhs.m21;
m02 = tmp.m00*rhs.m02 + tmp.m01*rhs.m12 + tmp.m02*rhs.m22;
m10 = tmp.m10*rhs.m00 + tmp.m11*rhs.m10 + tmp.m12*rhs.m20;
m11 = tmp.m10*rhs.m01 + tmp.m11*rhs.m11 + tmp.m12*rhs.m21;
m12 = tmp.m10*rhs.m02 + tmp.m11*rhs.m12 + tmp.m12*rhs.m22;
m20 = tmp.m20*rhs.m00 + tmp.m21*rhs.m10 + tmp.m22*rhs.m20;
m21 = tmp.m20*rhs.m01 + tmp.m21*rhs.m11 + tmp.m22*rhs.m21;
m22 = tmp.m20*rhs.m02 + tmp.m21*rhs.m12 + tmp.m22*rhs.m22;
return *this;
}
__host__ __device__ cumat3 cumat3::transpose() const
{
cumat3 ret;
ret.m00 = m00;
ret.m10 = m01;
ret.m20 = m02;
ret.m01 = m10;
ret.m11 = m11;
ret.m21 = m12;
ret.m02 = m20;
ret.m12 = m21;
ret.m22 = m22;
return ret;
}

56
test_scripts/cuvec_codegen/cuvec3_codegen1.hpp Normal file
View File

@ -0,0 +1,56 @@
__host__ __device__ cuvec3();
__host__ __device__ ~cuvec3();
__host__ __device__ cuvec3(const double &_x, const double &_y, const double &_z);
__host__ __device__ float& operator[](const int &I);
__host__ __device__ const float& operator[](const int &I) const;
__host__ __device__ cuvec3 operator+(const cuvec3& rhs) const;
__host__ __device__ cuvec3 operator-(const cuvec3& rhs) const;
__host__ __device__ cuvec3 operator*(const cuvec3& rhs) const; //elementwise product
__host__ __device__ cuvec3 operator/(const cuvec3& rhs) const; //elementwise division
__host__ __device__ friend cuvec3 operator*(const cuvec3& lhs, const double& rhs);
__host__ __device__ friend cuvec3 operator*(const double& lhs, const cuvec3& rhs);
__host__ __device__ friend cuvec3 operator/(const cuvec3& lhs, const double& rhs);
__host__ __device__ friend cuvec3 operator/(const double& lhs, const cuvec3& rhs);
__host__ __device__ friend cuvec3 operator-(const cuvec3& other);
__host__ __device__ cuvec3& operator+=(const cuvec3& rhs);
__host__ __device__ cuvec3& operator-=(const cuvec3& rhs);
__host__ __device__ cuvec3& operator*=(const double& rhs);
__host__ __device__ cuvec3& operator/=(const double& rhs);
//
//Matrix Header Stuff
//
double m00,m10,m20;
double m01,m11,m21;
double m02,m12,m22;
__host__ __device__ cumat3();
__host__ __device__ ~cumat3();
__host__ __device__ cumat3(
const double& _m00, const double& _m01, const double& _m02,
const double& _m10, const double& _m11, const double& _m12,
const double& _m20, const double& _m21, const double& _m22
);
__host__ __device__ cumat3(const double* data9);
__host__ __device__ float& operator[](const int &I);
__host__ __device__ const float& operator[](const int &I) const;
__host__ __device__ float& operator()(const int &I, const int &J);
__host__ __device__ const float& operator()(const int &I, const int &J) const;
__host__ __device__ float& at(const int &I, const int &J);
__host__ __device__ const float& at(const int &I, const int &J) const;
__host__ __device__ double* data(); //pointer to float9 representation of matrix
__host__ __device__ const double* data() const; //pointer to float9 representation of matrix
__host__ __device__ cumat3 operator+(const cumat3& rhs) const;
__host__ __device__ cumat3 operator-(const cumat3& rhs) const;
__host__ __device__ cumat3 operator*(const cumat3& rhs) const;
__host__ __device__ friend cumat3 operator*(const cumat3& lhs, const double& rhs);
__host__ __device__ friend cumat3 operator/(const cumat3& lhs, const double& rhs);
__host__ __device__ friend cumat3 operator*(const double& lhs, const cumat3& rhs);
__host__ __device__ friend cuvec3 operator*(const cumat3& lhs, const cuvec3& rhs);
__host__ __device__ friend cuvec3 operator*(const cuvec3& lhs, const cumat3& rhs);
__host__ __device__ friend cumat3 operator-(const cumat3& rhs);
__host__ __device__ cumat3& operator+=(const cumat3& rhs);
__host__ __device__ cumat3& operator-=(const cumat3& rhs);
__host__ __device__ cumat3& operator*=(const double& rhs);
__host__ __device__ cumat3& operator/=(const double& rhs);
__host__ __device__ cumat3& operator*=(const cumat3& rhs);
__host__ __device__ cumat3 transpose() const;

544
test_scripts/cuvec_codegen/cuvec3f_codegen1.cu Normal file
View File

@ -0,0 +1,544 @@
__host__ __device__ cuvec3f::cuvec3f()
{
x = 0; y = 0; z = 0;
return;
}
__host__ __device__ cuvec3f::~cuvec3f()
{
x = 0; y = 0; z = 0;
return;
}
__host__ __device__ cuvec3f::cuvec3f(const float &_x, const float &_y, const float &_z)
{
x = _x; y = _y; z = _z;
return;
}
__host__ __device__ float& cuvec3f::operator[](const int &I)
{
switch(I)
{
case 0:
return x;
case 1:
return y;
case 2:
return z;
}
return x;
}
__host__ __device__ const float& cuvec3f::operator[](const int &I) const
{
switch(I)
{
case 0:
return x;
case 1:
return y;
case 2:
return z;
}
return x;
}
__host__ __device__ cuvec3f cuvec3f::operator+(const cuvec3f& rhs) const
{
cuvec3f ret;
ret.x = x + rhs.x;
ret.y = y + rhs.y;
ret.z = z + rhs.z;
return ret;
}
__host__ __device__ cuvec3f cuvec3f::operator-(const cuvec3f& rhs) const
{
cuvec3f ret;
ret.x = x - rhs.x;
ret.y = y - rhs.y;
ret.z = z - rhs.z;
return ret;
}
__host__ __device__ cuvec3f cuvec3f::operator*(const cuvec3f& rhs) const
{
//Elementwise product
cuvec3f ret;
ret.x = x * rhs.x;
ret.y = y * rhs.y;
ret.z = z * rhs.z;
return ret;
}
__host__ __device__ cuvec3f cuvec3f::operator/(const cuvec3f& rhs) const
{
//Elementwise division
cuvec3f ret;
ret.x = x / rhs.x;
ret.y = y / rhs.y;
ret.z = z / rhs.z;
return ret;
}
__host__ __device__ cuvec3f operator*(const cuvec3f& lhs, const float& rhs)
{
cuvec3f ret;
ret.x = lhs.x*rhs;
ret.y = lhs.y*rhs;
ret.z = lhs.z*rhs;
return ret;
}
__host__ __device__ cuvec3f operator*(const float& lhs, const cuvec3f& rhs)
{
cuvec3f ret;
ret.x = lhs*rhs.x;
ret.y = lhs*rhs.y;
ret.z = lhs*rhs.z;
return ret;
}
__host__ __device__ cuvec3f operator/(const cuvec3f& lhs, const float& rhs)
{
cuvec3f ret;
ret.x = lhs.x/rhs;
ret.y = lhs.y/rhs;
ret.z = lhs.z/rhs;
return ret;
}
__host__ __device__ cuvec3f operator/(const float& lhs, const cuvec3f& rhs)
{
cuvec3f ret;
ret.x = lhs/rhs.x;
ret.y = lhs/rhs.y;
ret.z = lhs/rhs.z;
return ret;
}
__host__ __device__ cuvec3f operator-(const cuvec3f& other)
{
cuvec3f ret;
ret.x = -other.x;
ret.y = -other.y;
ret.z = -other.z;
return ret;
}
__host__ __device__ cuvec3f& cuvec3f::operator+=(const cuvec3f& rhs)
{
x += rhs.x;
y += rhs.y;
z += rhs.z;
return *this;
}
__host__ __device__ cuvec3f& cuvec3f::operator-=(const cuvec3f& rhs)
{
x -= rhs.x;
y -= rhs.y;
z -= rhs.z;
return *this;
}
__host__ __device__ cuvec3f& cuvec3f::operator*=(const float& rhs)
{
x *= rhs;
y *= rhs;
z *= rhs;
return *this;
}
__host__ __device__ cuvec3f& cuvec3f::operator/=(const float& rhs)
{
x /= rhs;
y /= rhs;
z /= rhs;
return *this;
}
//
//Matrix Header Stuff
//
__host__ __device__ cumat3f::cumat3f()
{
m00 = 0;
m01 = 0;
m02 = 0;
m10 = 0;
m11 = 0;
m12 = 0;
m20 = 0;
m21 = 0;
m22 = 0;
return;
}
__host__ __device__ cumat3f::~cumat3f()
{
//m00 = 0;
//m01 = 0;
//m02 = 0;
//m10 = 0;
//m11 = 0;
//m12 = 0;
//m20 = 0;
//m21 = 0;
//m22 = 0;
return;
}
__host__ __device__ cumat3f::cumat3f(
const float& _m00, const float& _m01, const float& _m02,
const float& _m10, const float& _m11, const float& _m12,
const float& _m20, const float& _m21, const float& _m22
)
{
m00 = _m00;
m10 = _m10;
m20 = _m20;
m01 = _m01;
m11 = _m11;
m21 = _m21;
m02 = _m02;
m12 = _m12;
m22 = _m22;
return;
}
__host__ __device__ cumat3f::cumat3f(const float* data9)
{
m00 = data9[0];
m10 = data9[1];
m20 = data9[2];
m01 = data9[3];
m11 = data9[4];
m21 = data9[5];
m02 = data9[6];
m12 = data9[7];
m22 = data9[8];
return;
}
__host__ __device__ float& cumat3f::operator[](const int &I)
{
switch(I)
{
case 0:
return m00;
case 1:
return m10;
case 2:
return m20;
case 3:
return m01;
case 4:
return m11;
case 5:
return m21;
case 6:
return m02;
case 7:
return m12;
case 8:
return m22;
}
return m00;
}
__host__ __device__ const float& cumat3f::operator[](const int &I) const
{
switch(I)
{
case 0:
return m00;
case 1:
return m10;
case 2:
return m20;
case 3:
return m01;
case 4:
return m11;
case 5:
return m21;
case 6:
return m02;
case 7:
return m12;
case 8:
return m22;
}
return m00;
}
__host__ __device__ float& cumat3f::operator()(const int &I, const int &J)
{
return (*this)[I+3*J];
}
__host__ __device__ const float& cumat3f::operator()(const int &I, const int &J) const
{
return (*this)[I+3*J];
}
__host__ __device__ float& cumat3f::at(const int &I, const int &J)
{
return (*this)[I+3*J];
}
__host__ __device__ const float& cumat3f::at(const int &I, const int &J) const
{
return (*this)[I+3*J];
}
__host__ __device__ float* cumat3f::data()
{
return (float*)this;
}
__host__ __device__ const float* cumat3f::data() const
{
return (float*)this;
}
__host__ __device__ cumat3f cumat3f::operator+(const cumat3f& rhs) const
{
cumat3f ret;
ret.m00 = m00 + rhs.m00;
ret.m10 = m10 + rhs.m10;
ret.m20 = m20 + rhs.m20;
ret.m01 = m01 + rhs.m01;
ret.m11 = m11 + rhs.m11;
ret.m21 = m21 + rhs.m21;
ret.m02 = m02 + rhs.m02;
ret.m12 = m12 + rhs.m12;
ret.m22 = m22 + rhs.m22;
return ret;
}
__host__ __device__ cumat3f cumat3f::operator-(const cumat3f& rhs) const
{
cumat3f ret;
ret.m00 = m00 - rhs.m00;
ret.m10 = m10 - rhs.m10;
ret.m20 = m20 - rhs.m20;
ret.m01 = m01 - rhs.m01;
ret.m11 = m11 - rhs.m11;
ret.m21 = m21 - rhs.m21;
ret.m02 = m02 - rhs.m02;
ret.m12 = m12 - rhs.m12;
ret.m22 = m22 - rhs.m22;
return ret;
}
__host__ __device__ cumat3f cumat3f::operator*(const cumat3f& rhs) const
{
cumat3f ret; //should be zeroed in constructor
ret.m00 = m00*rhs.m00 + m01*rhs.m10 + m02*rhs.m20;
ret.m01 = m00*rhs.m01 + m01*rhs.m11 + m02*rhs.m21;
ret.m02 = m00*rhs.m02 + m01*rhs.m12 + m02*rhs.m22;
ret.m10 = m10*rhs.m00 + m11*rhs.m10 + m12*rhs.m20;
ret.m11 = m10*rhs.m01 + m11*rhs.m11 + m12*rhs.m21;
ret.m12 = m10*rhs.m02 + m11*rhs.m12 + m12*rhs.m22;
ret.m20 = m20*rhs.m00 + m21*rhs.m10 + m22*rhs.m20;
ret.m21 = m20*rhs.m01 + m21*rhs.m11 + m22*rhs.m21;
ret.m22 = m20*rhs.m02 + m21*rhs.m12 + m22*rhs.m22;
return ret;
}
__host__ __device__ cumat3f operator*(const cumat3f& lhs, const float& rhs)
{
cumat3f ret;
ret.m00=lhs.m00*rhs;
ret.m10=lhs.m10*rhs;
ret.m20=lhs.m20*rhs;
ret.m01=lhs.m01*rhs;
ret.m11=lhs.m11*rhs;
ret.m21=lhs.m21*rhs;
ret.m02=lhs.m02*rhs;
ret.m12=lhs.m12*rhs;
ret.m22=lhs.m22*rhs;
return ret;
}
__host__ __device__ cumat3f operator/(const cumat3f& lhs, const float& rhs)
{
cumat3f ret;
ret.m00=lhs.m00/rhs;
ret.m10=lhs.m10/rhs;
ret.m20=lhs.m20/rhs;
ret.m01=lhs.m01/rhs;
ret.m11=lhs.m11/rhs;
ret.m21=lhs.m21/rhs;
ret.m02=lhs.m02/rhs;
ret.m12=lhs.m12/rhs;
ret.m22=lhs.m22/rhs;
return ret;
}
__host__ __device__ cumat3f operator*(const float& lhs, const cumat3f& rhs)
{
cumat3f ret;
ret.m00=lhs*rhs.m00;
ret.m10=lhs*rhs.m10;
ret.m20=lhs*rhs.m20;
ret.m01=lhs*rhs.m01;
ret.m11=lhs*rhs.m11;
ret.m21=lhs*rhs.m21;
ret.m02=lhs*rhs.m02;
ret.m12=lhs*rhs.m12;
ret.m22=lhs*rhs.m22;
return ret;
}
__host__ __device__ cuvec3f operator*(const cumat3f& lhs, const cuvec3f& rhs)
{
cuvec3f ret;
ret.x = lhs.m00*rhs.x + lhs.m01*rhs.y + lhs.m02*rhs.z;
ret.y = lhs.m10*rhs.x + lhs.m11*rhs.y + lhs.m12*rhs.z;
ret.z = lhs.m20*rhs.x + lhs.m21*rhs.y + lhs.m22*rhs.z;
return ret;
}
__host__ __device__ cuvec3f operator*(const cuvec3f& lhs, const cumat3f& rhs)
{
cuvec3f ret;
ret.x = lhs.x*rhs.m00 + lhs.y*rhs.m10 + lhs.z*rhs.m20;
ret.y = lhs.x*rhs.m01 + lhs.y*rhs.m11 + lhs.z*rhs.m21;
ret.z = lhs.x*rhs.m02 + lhs.y*rhs.m12 + lhs.z*rhs.m22;
return ret;
}
__host__ __device__ cumat3f operator-(const cumat3f& rhs)
{
cumat3f ret;
ret.m00 = -rhs.m00;
ret.m10 = -rhs.m10;
ret.m20 = -rhs.m20;
ret.m01 = -rhs.m01;
ret.m11 = -rhs.m11;
ret.m21 = -rhs.m21;
ret.m02 = -rhs.m02;
ret.m12 = -rhs.m12;
ret.m22 = -rhs.m22;
return ret;
}
__host__ __device__ cumat3f& cumat3f::operator+=(const cumat3f& rhs)
{
m00 += rhs.m00;
m10 += rhs.m10;
m20 += rhs.m20;
m01 += rhs.m01;
m11 += rhs.m11;
m21 += rhs.m21;
m02 += rhs.m02;
m12 += rhs.m12;
m22 += rhs.m22;
return *this;
}
__host__ __device__ cumat3f& cumat3f::operator-=(const cumat3f& rhs)
{
m00 -= rhs.m00;
m10 -= rhs.m10;
m20 -= rhs.m20;
m01 -= rhs.m01;
m11 -= rhs.m11;
m21 -= rhs.m21;
m02 -= rhs.m02;
m12 -= rhs.m12;
m22 -= rhs.m22;
return *this;
}
__host__ __device__ cumat3f& cumat3f::operator*=(const float& rhs)
{
m00 *= rhs;
m10 *= rhs;
m20 *= rhs;
m01 *= rhs;
m11 *= rhs;
m21 *= rhs;
m02 *= rhs;
m12 *= rhs;
m22 *= rhs;
return *this;
}
__host__ __device__ cumat3f& cumat3f::operator/=(const float& rhs)
{
m00 /= rhs;
m10 /= rhs;
m20 /= rhs;
m01 /= rhs;
m11 /= rhs;
m21 /= rhs;
m02 /= rhs;
m12 /= rhs;
m22 /= rhs;
return *this;
}
__host__ __device__ cumat3f& cumat3f::operator*=(const cumat3f& rhs)
{
cumat3f tmp = *this;
m00 = tmp.m00*rhs.m00 + tmp.m01*rhs.m10 + tmp.m02*rhs.m20;
m01 = tmp.m00*rhs.m01 + tmp.m01*rhs.m11 + tmp.m02*rhs.m21;
m02 = tmp.m00*rhs.m02 + tmp.m01*rhs.m12 + tmp.m02*rhs.m22;
m10 = tmp.m10*rhs.m00 + tmp.m11*rhs.m10 + tmp.m12*rhs.m20;
m11 = tmp.m10*rhs.m01 + tmp.m11*rhs.m11 + tmp.m12*rhs.m21;
m12 = tmp.m10*rhs.m02 + tmp.m11*rhs.m12 + tmp.m12*rhs.m22;
m20 = tmp.m20*rhs.m00 + tmp.m21*rhs.m10 + tmp.m22*rhs.m20;
m21 = tmp.m20*rhs.m01 + tmp.m21*rhs.m11 + tmp.m22*rhs.m21;
m22 = tmp.m20*rhs.m02 + tmp.m21*rhs.m12 + tmp.m22*rhs.m22;
return *this;
}
__host__ __device__ cumat3f cumat3f::transpose() const
{
cumat3f ret;
ret.m00 = m00;
ret.m10 = m01;
ret.m20 = m02;
ret.m01 = m10;
ret.m11 = m11;
ret.m21 = m12;
ret.m02 = m20;
ret.m12 = m21;
ret.m22 = m22;
return ret;
}

56
test_scripts/cuvec_codegen/cuvec3f_codegen1.hpp Normal file
View File

@ -0,0 +1,56 @@
__host__ __device__ cuvec3f();
__host__ __device__ ~cuvec3f();
__host__ __device__ cuvec3f(const float &_x, const float &_y, const float &_z);
__host__ __device__ float& operator[](const int &I);
__host__ __device__ const float& operator[](const int &I) const;
__host__ __device__ cuvec3f operator+(const cuvec3f& rhs) const;
__host__ __device__ cuvec3f operator-(const cuvec3f& rhs) const;
__host__ __device__ cuvec3f operator*(const cuvec3f& rhs) const; //elementwise product
__host__ __device__ cuvec3f operator/(const cuvec3f& rhs) const; //elementwise division
__host__ __device__ friend cuvec3f operator*(const cuvec3f& lhs, const float& rhs);
__host__ __device__ friend cuvec3f operator*(const float& lhs, const cuvec3f& rhs);
__host__ __device__ friend cuvec3f operator/(const cuvec3f& lhs, const float& rhs);
__host__ __device__ friend cuvec3f operator/(const float& lhs, const cuvec3f& rhs);
__host__ __device__ friend cuvec3f operator-(const cuvec3f& other);
__host__ __device__ cuvec3f& operator+=(const cuvec3f& rhs);
__host__ __device__ cuvec3f& operator-=(const cuvec3f& rhs);
__host__ __device__ cuvec3f& operator*=(const float& rhs);
__host__ __device__ cuvec3f& operator/=(const float& rhs);
//
//Matrix Header Stuff
//
float m00,m10,m20;
float m01,m11,m21;
float m02,m12,m22;
__host__ __device__ cumat3f();
__host__ __device__ ~cumat3f();
__host__ __device__ cumat3f(
const float& _m00, const float& _m01, const float& _m02,
const float& _m10, const float& _m11, const float& _m12,
const float& _m20, const float& _m21, const float& _m22
);
__host__ __device__ cumat3f(const float* data9);
__host__ __device__ float& operator[](const int &I);
__host__ __device__ const float& operator[](const int &I) const;
__host__ __device__ float& operator()(const int &I, const int &J);
__host__ __device__ const float& operator()(const int &I, const int &J) const;
__host__ __device__ float& at(const int &I, const int &J);
__host__ __device__ const float& at(const int &I, const int &J) const;
__host__ __device__ float* data(); //pointer to float9 representation of matrix
__host__ __device__ const float* data() const; //pointer to float9 representation of matrix
__host__ __device__ cumat3f operator+(const cumat3f& rhs) const;
__host__ __device__ cumat3f operator-(const cumat3f& rhs) const;
__host__ __device__ cumat3f operator*(const cumat3f& rhs) const;
__host__ __device__ friend cumat3f operator*(const cumat3f& lhs, const float& rhs);
__host__ __device__ friend cumat3f operator/(const cumat3f& lhs, const float& rhs);
__host__ __device__ friend cumat3f operator*(const float& lhs, const cumat3f& rhs);
__host__ __device__ friend cuvec3f operator*(const cumat3f& lhs, const cuvec3f& rhs);
__host__ __device__ friend cuvec3f operator*(const cuvec3f& lhs, const cumat3f& rhs);
__host__ __device__ friend cumat3f operator-(const cumat3f& rhs);
__host__ __device__ cumat3f& operator+=(const cumat3f& rhs);
__host__ __device__ cumat3f& operator-=(const cumat3f& rhs);
__host__ __device__ cumat3f& operator*=(const float& rhs);
__host__ __device__ cumat3f& operator/=(const float& rhs);
__host__ __device__ cumat3f& operator*=(const cumat3f& rhs);
__host__ __device__ cumat3f transpose() const;

544
test_scripts/cuvec_codegen/cuvec3i_codegen1.cu Normal file
View File

@ -0,0 +1,544 @@
__host__ __device__ cuvec3i::cuvec3i()
{
x = 0; y = 0; z = 0;
return;
}
__host__ __device__ cuvec3i::~cuvec3i()
{
x = 0; y = 0; z = 0;
return;
}
__host__ __device__ cuvec3i::cuvec3i(const int &_x, const int &_y, const int &_z)
{
x = _x; y = _y; z = _z;
return;
}
__host__ __device__ float& cuvec3i::operator[](const int &I)
{
switch(I)
{
case 0:
return x;
case 1:
return y;
case 2:
return z;
}
return x;
}
__host__ __device__ const float& cuvec3i::operator[](const int &I) const
{
switch(I)
{
case 0:
return x;
case 1:
return y;
case 2:
return z;
}
return x;
}
__host__ __device__ cuvec3i cuvec3i::operator+(const cuvec3i& rhs) const
{
cuvec3i ret;
ret.x = x + rhs.x;
ret.y = y + rhs.y;
ret.z = z + rhs.z;
return ret;
}
__host__ __device__ cuvec3i cuvec3i::operator-(const cuvec3i& rhs) const
{
cuvec3i ret;
ret.x = x - rhs.x;
ret.y = y - rhs.y;
ret.z = z - rhs.z;
return ret;
}
__host__ __device__ cuvec3i cuvec3i::operator*(const cuvec3i& rhs) const
{
//Elementwise product
cuvec3i ret;
ret.x = x * rhs.x;
ret.y = y * rhs.y;
ret.z = z * rhs.z;
return ret;
}
__host__ __device__ cuvec3i cuvec3i::operator/(const cuvec3i& rhs) const
{
//Elementwise division
cuvec3i ret;
ret.x = x / rhs.x;
ret.y = y / rhs.y;
ret.z = z / rhs.z;
return ret;
}
__host__ __device__ cuvec3i operator*(const cuvec3i& lhs, const int& rhs)
{
cuvec3i ret;
ret.x = lhs.x*rhs;
ret.y = lhs.y*rhs;
ret.z = lhs.z*rhs;
return ret;
}
__host__ __device__ cuvec3i operator*(const int& lhs, const cuvec3i& rhs)
{
cuvec3i ret;
ret.x = lhs*rhs.x;
ret.y = lhs*rhs.y;
ret.z = lhs*rhs.z;
return ret;
}
__host__ __device__ cuvec3i operator/(const cuvec3i& lhs, const int& rhs)
{
cuvec3i ret;
ret.x = lhs.x/rhs;
ret.y = lhs.y/rhs;
ret.z = lhs.z/rhs;
return ret;
}
__host__ __device__ cuvec3i operator/(const int& lhs, const cuvec3i& rhs)
{
cuvec3i ret;
ret.x = lhs/rhs.x;
ret.y = lhs/rhs.y;
ret.z = lhs/rhs.z;
return ret;
}
__host__ __device__ cuvec3i operator-(const cuvec3i& other)
{
cuvec3i ret;
ret.x = -other.x;
ret.y = -other.y;
ret.z = -other.z;
return ret;
}
__host__ __device__ cuvec3i& cuvec3i::operator+=(const cuvec3i& rhs)
{
x += rhs.x;
y += rhs.y;
z += rhs.z;
return *this;
}
__host__ __device__ cuvec3i& cuvec3i::operator-=(const cuvec3i& rhs)
{
x -= rhs.x;
y -= rhs.y;
z -= rhs.z;
return *this;
}
__host__ __device__ cuvec3i& cuvec3i::operator*=(const int& rhs)
{
x *= rhs;
y *= rhs;
z *= rhs;
return *this;
}
__host__ __device__ cuvec3i& cuvec3i::operator/=(const int& rhs)
{
x /= rhs;
y /= rhs;
z /= rhs;
return *this;
}
//
//Matrix Header Stuff
//
__host__ __device__ cumat3i::cumat3i()
{
m00 = 0;
m01 = 0;
m02 = 0;
m10 = 0;
m11 = 0;
m12 = 0;
m20 = 0;
m21 = 0;
m22 = 0;
return;
}
__host__ __device__ cumat3i::~cumat3i()
{
//m00 = 0;
//m01 = 0;
//m02 = 0;
//m10 = 0;
//m11 = 0;
//m12 = 0;
//m20 = 0;
//m21 = 0;
//m22 = 0;
return;
}
__host__ __device__ cumat3i::cumat3i(
const int& _m00, const int& _m01, const int& _m02,
const int& _m10, const int& _m11, const int& _m12,
const int& _m20, const int& _m21, const int& _m22
)
{
m00 = _m00;
m10 = _m10;
m20 = _m20;
m01 = _m01;
m11 = _m11;
m21 = _m21;
m02 = _m02;
m12 = _m12;
m22 = _m22;
return;
}
__host__ __device__ cumat3i::cumat3i(const int* data9)
{
m00 = data9[0];
m10 = data9[1];
m20 = data9[2];
m01 = data9[3];
m11 = data9[4];
m21 = data9[5];
m02 = data9[6];
m12 = data9[7];
m22 = data9[8];
return;
}
__host__ __device__ float& cumat3i::operator[](const int &I)
{
switch(I)
{
case 0:
return m00;
case 1:
return m10;
case 2:
return m20;
case 3:
return m01;
case 4:
return m11;
case 5:
return m21;
case 6:
return m02;
case 7:
return m12;
case 8:
return m22;
}
return m00;
}
__host__ __device__ const float& cumat3i::operator[](const int &I) const
{
switch(I)
{
case 0:
return m00;
case 1:
return m10;
case 2:
return m20;
case 3:
return m01;
case 4:
return m11;
case 5:
return m21;
case 6:
return m02;
case 7:
return m12;
case 8:
return m22;
}
return m00;
}
__host__ __device__ float& cumat3i::operator()(const int &I, const int &J)
{
return (*this)[I+3*J];
}
__host__ __device__ const float& cumat3i::operator()(const int &I, const int &J) const
{
return (*this)[I+3*J];
}
__host__ __device__ float& cumat3i::at(const int &I, const int &J)
{
return (*this)[I+3*J];
}
__host__ __device__ const float& cumat3i::at(const int &I, const int &J) const
{
return (*this)[I+3*J];
}
__host__ __device__ int* cumat3i::data()
{
return (int*)this;
}
__host__ __device__ const int* cumat3i::data() const
{
return (int*)this;
}
__host__ __device__ cumat3i cumat3i::operator+(const cumat3i& rhs) const
{
cumat3i ret;
ret.m00 = m00 + rhs.m00;
ret.m10 = m10 + rhs.m10;
ret.m20 = m20 + rhs.m20;
ret.m01 = m01 + rhs.m01;
ret.m11 = m11 + rhs.m11;
ret.m21 = m21 + rhs.m21;
ret.m02 = m02 + rhs.m02;
ret.m12 = m12 + rhs.m12;
ret.m22 = m22 + rhs.m22;
return ret;
}
__host__ __device__ cumat3i cumat3i::operator-(const cumat3i& rhs) const
{
cumat3i ret;
ret.m00 = m00 - rhs.m00;
ret.m10 = m10 - rhs.m10;
ret.m20 = m20 - rhs.m20;
ret.m01 = m01 - rhs.m01;
ret.m11 = m11 - rhs.m11;
ret.m21 = m21 - rhs.m21;
ret.m02 = m02 - rhs.m02;
ret.m12 = m12 - rhs.m12;
ret.m22 = m22 - rhs.m22;
return ret;
}
__host__ __device__ cumat3i cumat3i::operator*(const cumat3i& rhs) const
{
cumat3i ret; //should be zeroed in constructor
ret.m00 = m00*rhs.m00 + m01*rhs.m10 + m02*rhs.m20;
ret.m01 = m00*rhs.m01 + m01*rhs.m11 + m02*rhs.m21;
ret.m02 = m00*rhs.m02 + m01*rhs.m12 + m02*rhs.m22;
ret.m10 = m10*rhs.m00 + m11*rhs.m10 + m12*rhs.m20;
ret.m11 = m10*rhs.m01 + m11*rhs.m11 + m12*rhs.m21;
ret.m12 = m10*rhs.m02 + m11*rhs.m12 + m12*rhs.m22;
ret.m20 = m20*rhs.m00 + m21*rhs.m10 + m22*rhs.m20;
ret.m21 = m20*rhs.m01 + m21*rhs.m11 + m22*rhs.m21;
ret.m22 = m20*rhs.m02 + m21*rhs.m12 + m22*rhs.m22;
return ret;
}
__host__ __device__ cumat3i operator*(const cumat3i& lhs, const int& rhs)
{
cumat3i ret;
ret.m00=lhs.m00*rhs;
ret.m10=lhs.m10*rhs;
ret.m20=lhs.m20*rhs;
ret.m01=lhs.m01*rhs;
ret.m11=lhs.m11*rhs;
ret.m21=lhs.m21*rhs;
ret.m02=lhs.m02*rhs;
ret.m12=lhs.m12*rhs;
ret.m22=lhs.m22*rhs;
return ret;
}
__host__ __device__ cumat3i operator/(const cumat3i& lhs, const int& rhs)
{
cumat3i ret;
ret.m00=lhs.m00/rhs;
ret.m10=lhs.m10/rhs;
ret.m20=lhs.m20/rhs;
ret.m01=lhs.m01/rhs;
ret.m11=lhs.m11/rhs;
ret.m21=lhs.m21/rhs;
ret.m02=lhs.m02/rhs;
ret.m12=lhs.m12/rhs;
ret.m22=lhs.m22/rhs;
return ret;
}
__host__ __device__ cumat3i operator*(const int& lhs, const cumat3i& rhs)
{
cumat3i ret;
ret.m00=lhs*rhs.m00;
ret.m10=lhs*rhs.m10;
ret.m20=lhs*rhs.m20;
ret.m01=lhs*rhs.m01;
ret.m11=lhs*rhs.m11;
ret.m21=lhs*rhs.m21;
ret.m02=lhs*rhs.m02;
ret.m12=lhs*rhs.m12;
ret.m22=lhs*rhs.m22;
return ret;
}
__host__ __device__ cuvec3i operator*(const cumat3i& lhs, const cuvec3i& rhs)
{
cuvec3i ret;
ret.x = lhs.m00*rhs.x + lhs.m01*rhs.y + lhs.m02*rhs.z;
ret.y = lhs.m10*rhs.x + lhs.m11*rhs.y + lhs.m12*rhs.z;
ret.z = lhs.m20*rhs.x + lhs.m21*rhs.y + lhs.m22*rhs.z;
return ret;
}
__host__ __device__ cuvec3i operator*(const cuvec3i& lhs, const cumat3i& rhs)
{
cuvec3i ret;
ret.x = lhs.x*rhs.m00 + lhs.y*rhs.m10 + lhs.z*rhs.m20;
ret.y = lhs.x*rhs.m01 + lhs.y*rhs.m11 + lhs.z*rhs.m21;
ret.z = lhs.x*rhs.m02 + lhs.y*rhs.m12 + lhs.z*rhs.m22;
return ret;
}
__host__ __device__ cumat3i operator-(const cumat3i& rhs)
{
cumat3i ret;
ret.m00 = -rhs.m00;
ret.m10 = -rhs.m10;
ret.m20 = -rhs.m20;
ret.m01 = -rhs.m01;
ret.m11 = -rhs.m11;
ret.m21 = -rhs.m21;
ret.m02 = -rhs.m02;
ret.m12 = -rhs.m12;
ret.m22 = -rhs.m22;
return ret;
}
__host__ __device__ cumat3i& cumat3i::operator+=(const cumat3i& rhs)
{
m00 += rhs.m00;
m10 += rhs.m10;
m20 += rhs.m20;
m01 += rhs.m01;
m11 += rhs.m11;
m21 += rhs.m21;
m02 += rhs.m02;
m12 += rhs.m12;
m22 += rhs.m22;
return *this;
}
__host__ __device__ cumat3i& cumat3i::operator-=(const cumat3i& rhs)
{
m00 -= rhs.m00;
m10 -= rhs.m10;
m20 -= rhs.m20;
m01 -= rhs.m01;
m11 -= rhs.m11;
m21 -= rhs.m21;
m02 -= rhs.m02;
m12 -= rhs.m12;
m22 -= rhs.m22;
return *this;
}
__host__ __device__ cumat3i& cumat3i::operator*=(const int& rhs)
{
m00 *= rhs;
m10 *= rhs;
m20 *= rhs;
m01 *= rhs;
m11 *= rhs;
m21 *= rhs;
m02 *= rhs;
m12 *= rhs;
m22 *= rhs;
return *this;
}
__host__ __device__ cumat3i& cumat3i::operator/=(const int& rhs)
{
m00 /= rhs;
m10 /= rhs;
m20 /= rhs;
m01 /= rhs;
m11 /= rhs;
m21 /= rhs;
m02 /= rhs;
m12 /= rhs;
m22 /= rhs;
return *this;
}
__host__ __device__ cumat3i& cumat3i::operator*=(const cumat3i& rhs)
{
cumat3i tmp = *this;
m00 = tmp.m00*rhs.m00 + tmp.m01*rhs.m10 + tmp.m02*rhs.m20;
m01 = tmp.m00*rhs.m01 + tmp.m01*rhs.m11 + tmp.m02*rhs.m21;
m02 = tmp.m00*rhs.m02 + tmp.m01*rhs.m12 + tmp.m02*rhs.m22;
m10 = tmp.m10*rhs.m00 + tmp.m11*rhs.m10 + tmp.m12*rhs.m20;
m11 = tmp.m10*rhs.m01 + tmp.m11*rhs.m11 + tmp.m12*rhs.m21;
m12 = tmp.m10*rhs.m02 + tmp.m11*rhs.m12 + tmp.m12*rhs.m22;
m20 = tmp.m20*rhs.m00 + tmp.m21*rhs.m10 + tmp.m22*rhs.m20;
m21 = tmp.m20*rhs.m01 + tmp.m21*rhs.m11 + tmp.m22*rhs.m21;
m22 = tmp.m20*rhs.m02 + tmp.m21*rhs.m12 + tmp.m22*rhs.m22;
return *this;
}
__host__ __device__ cumat3i cumat3i::transpose() const
{
cumat3i ret;
ret.m00 = m00;
ret.m10 = m01;
ret.m20 = m02;
ret.m01 = m10;
ret.m11 = m11;
ret.m21 = m12;
ret.m02 = m20;
ret.m12 = m21;
ret.m22 = m22;
return ret;
}

56
test_scripts/cuvec_codegen/cuvec3i_codegen1.hpp Normal file
View File

@ -0,0 +1,56 @@
__host__ __device__ cuvec3i();
__host__ __device__ ~cuvec3i();
__host__ __device__ cuvec3i(const int &_x, const int &_y, const int &_z);
__host__ __device__ float& operator[](const int &I);
__host__ __device__ const float& operator[](const int &I) const;
__host__ __device__ cuvec3i operator+(const cuvec3i& rhs) const;
__host__ __device__ cuvec3i operator-(const cuvec3i& rhs) const;
__host__ __device__ cuvec3i operator*(const cuvec3i& rhs) const; //elementwise product
__host__ __device__ cuvec3i operator/(const cuvec3i& rhs) const; //elementwise division
__host__ __device__ friend cuvec3i operator*(const cuvec3i& lhs, const int& rhs);
__host__ __device__ friend cuvec3i operator*(const int& lhs, const cuvec3i& rhs);
__host__ __device__ friend cuvec3i operator/(const cuvec3i& lhs, const int& rhs);
__host__ __device__ friend cuvec3i operator/(const int& lhs, const cuvec3i& rhs);
__host__ __device__ friend cuvec3i operator-(const cuvec3i& other);
__host__ __device__ cuvec3i& operator+=(const cuvec3i& rhs);
__host__ __device__ cuvec3i& operator-=(const cuvec3i& rhs);
__host__ __device__ cuvec3i& operator*=(const int& rhs);
__host__ __device__ cuvec3i& operator/=(const int& rhs);
//
//Matrix Header Stuff
//
int m00,m10,m20;
int m01,m11,m21;
int m02,m12,m22;
__host__ __device__ cumat3i();
__host__ __device__ ~cumat3i();
__host__ __device__ cumat3i(
const int& _m00, const int& _m01, const int& _m02,
const int& _m10, const int& _m11, const int& _m12,
const int& _m20, const int& _m21, const int& _m22
);
__host__ __device__ cumat3i(const int* data9);
__host__ __device__ float& operator[](const int &I);
__host__ __device__ const float& operator[](const int &I) const;
__host__ __device__ float& operator()(const int &I, const int &J);
__host__ __device__ const float& operator()(const int &I, const int &J) const;
__host__ __device__ float& at(const int &I, const int &J);
__host__ __device__ const float& at(const int &I, const int &J) const;
__host__ __device__ int* data(); //pointer to float9 representation of matrix
__host__ __device__ const int* data() const; //pointer to float9 representation of matrix
__host__ __device__ cumat3i operator+(const cumat3i& rhs) const;
__host__ __device__ cumat3i operator-(const cumat3i& rhs) const;
__host__ __device__ cumat3i operator*(const cumat3i& rhs) const;
__host__ __device__ friend cumat3i operator*(const cumat3i& lhs, const int& rhs);
__host__ __device__ friend cumat3i operator/(const cumat3i& lhs, const int& rhs);
__host__ __device__ friend cumat3i operator*(const int& lhs, const cumat3i& rhs);
__host__ __device__ friend cuvec3i operator*(const cumat3i& lhs, const cuvec3i& rhs);
__host__ __device__ friend cuvec3i operator*(const cuvec3i& lhs, const cumat3i& rhs);
__host__ __device__ friend cumat3i operator-(const cumat3i& rhs);
__host__ __device__ cumat3i& operator+=(const cumat3i& rhs);
__host__ __device__ cumat3i& operator-=(const cumat3i& rhs);
__host__ __device__ cumat3i& operator*=(const int& rhs);
__host__ __device__ cumat3i& operator/=(const int& rhs);
__host__ __device__ cumat3i& operator*=(const cumat3i& rhs);
__host__ __device__ cumat3i transpose() const;

718
test_scripts/cuvec_codegen/cuvec4_codegen1.cu Normal file
View File

@ -0,0 +1,718 @@
__host__ __device__ cuvec4::cuvec4()
{
x = 0; y = 0; z = 0; w = 0;
return;
}
__host__ __device__ cuvec4::~cuvec4()
{
x = 0; y = 0; z = 0; w = 0;
return;
}
__host__ __device__ cuvec4::cuvec4(const double &_x, const double &_y, const double &_z, const double &_w)
{
x = _x; y = _y; z = _z; w = _w;
return;
}
__host__ __device__ float& cuvec4::operator[](const int &I)
{
switch(I)
{
case 0:
return x;
case 1:
return y;
case 2:
return z;
case 3:
return w;
}
return x;
}
__host__ __device__ const float& cuvec4::operator[](const int &I) const
{
switch(I)
{
case 0:
return x;
case 1:
return y;
case 2:
return z;
case 3:
return w;
}
return x;
}
__host__ __device__ cuvec4 cuvec4::operator+(const cuvec4& rhs) const
{
cuvec4 ret;
ret.x = x + rhs.x;
ret.y = y + rhs.y;
ret.z = z + rhs.z;
ret.w = w + rhs.w;
return ret;
}
__host__ __device__ cuvec4 cuvec4::operator-(const cuvec4& rhs) const
{
cuvec4 ret;
ret.x = x - rhs.x;
ret.y = y - rhs.y;
ret.z = z - rhs.z;
ret.w = w - rhs.w;
return ret;
}
__host__ __device__ cuvec4 cuvec4::operator*(const cuvec4& rhs) const
{
//Elementwise product
cuvec4 ret;
ret.x = x * rhs.x;
ret.y = y * rhs.y;
ret.z = z * rhs.z;
ret.w = w * rhs.w;
return ret;
}
__host__ __device__ cuvec4 cuvec4::operator/(const cuvec4& rhs) const
{
//Elementwise division
cuvec4 ret;
ret.x = x / rhs.x;
ret.y = y / rhs.y;
ret.z = z / rhs.z;
ret.w = w / rhs.w;
return ret;
}
__host__ __device__ cuvec4 operator*(const cuvec4& lhs, const double& rhs)
{
cuvec4 ret;
ret.x = lhs.x*rhs;
ret.y = lhs.y*rhs;
ret.z = lhs.z*rhs;
ret.w = lhs.w*rhs;
return ret;
}
__host__ __device__ cuvec4 operator*(const double& lhs, const cuvec4& rhs)
{
cuvec4 ret;
ret.x = lhs*rhs.x;
ret.y = lhs*rhs.y;
ret.z = lhs*rhs.z;
ret.w = lhs*rhs.w;
return ret;
}
__host__ __device__ cuvec4 operator/(const cuvec4& lhs, const double& rhs)
{
cuvec4 ret;
ret.x = lhs.x/rhs;
ret.y = lhs.y/rhs;
ret.z = lhs.z/rhs;
ret.w = lhs.w/rhs;
return ret;
}
__host__ __device__ cuvec4 operator/(const double& lhs, const cuvec4& rhs)
{
cuvec4 ret;
ret.x = lhs/rhs.x;
ret.y = lhs/rhs.y;
ret.z = lhs/rhs.z;
ret.w = lhs/rhs.w;
return ret;
}
__host__ __device__ cuvec4 operator-(const cuvec4& other)
{
cuvec4 ret;
ret.x = -other.x;
ret.y = -other.y;
ret.z = -other.z;
ret.w = -other.w;
return ret;
}
__host__ __device__ cuvec4& cuvec4::operator+=(const cuvec4& rhs)
{
x += rhs.x;
y += rhs.y;
z += rhs.z;
w += rhs.w;
return *this;
}
__host__ __device__ cuvec4& cuvec4::operator-=(const cuvec4& rhs)
{
x -= rhs.x;
y -= rhs.y;
z -= rhs.z;
w -= rhs.w;
return *this;
}
__host__ __device__ cuvec4& cuvec4::operator*=(const double& rhs)
{
x *= rhs;
y *= rhs;
z *= rhs;
w *= rhs;
return *this;
}
__host__ __device__ cuvec4& cuvec4::operator/=(const double& rhs)
{
x /= rhs;
y /= rhs;
z /= rhs;
w /= rhs;
return *this;
}
//
//Matrix Header Stuff
//
__host__ __device__ cumat4::cumat4()
{
m00 = 0;
m01 = 0;
m02 = 0;
m03 = 0;
m10 = 0;
m11 = 0;
m12 = 0;
m13 = 0;
m20 = 0;
m21 = 0;
m22 = 0;
m23 = 0;
m30 = 0;
m31 = 0;
m32 = 0;
m33 = 0;
return;
}
__host__ __device__ cumat4::~cumat4()
{
//m00 = 0;
//m01 = 0;
//m02 = 0;
//m03 = 0;
//m10 = 0;
//m11 = 0;
//m12 = 0;
//m13 = 0;
//m20 = 0;
//m21 = 0;
//m22 = 0;
//m23 = 0;
//m30 = 0;
//m31 = 0;
//m32 = 0;
//m33 = 0;
return;
}
__host__ __device__ cumat4::cumat4(
const double& _m00, const double& _m01, const double& _m02, const double& _m03,
const double& _m10, const double& _m11, const double& _m12, const double& _m13,
const double& _m20, const double& _m21, const double& _m22, const double& _m23,
const double& _m30, const double& _m31, const double& _m32, const double& _m33
)
{
m00 = _m00;
m10 = _m10;
m20 = _m20;
m30 = _m30;
m01 = _m01;
m11 = _m11;
m21 = _m21;
m31 = _m31;
m02 = _m02;
m12 = _m12;
m22 = _m22;
m32 = _m32;
m03 = _m03;
m13 = _m13;
m23 = _m23;
m33 = _m33;
return;
}
__host__ __device__ cumat4::cumat4(const double* data16)
{
m00 = data16[0];
m10 = data16[1];
m20 = data16[2];
m30 = data16[3];
m01 = data16[4];
m11 = data16[5];
m21 = data16[6];
m31 = data16[7];
m02 = data16[8];
m12 = data16[9];
m22 = data16[10];
m32 = data16[11];
m03 = data16[12];
m13 = data16[13];
m23 = data16[14];
m33 = data16[15];
return;
}
__host__ __device__ float& cumat4::operator[](const int &I)
{
switch(I)
{
case 0:
return m00;
case 1:
return m10;
case 2:
return m20;
case 3:
return m30;
case 4:
return m01;
case 5:
return m11;
case 6:
return m21;
case 7:
return m31;
case 8:
return m02;
case 9:
return m12;
case 10:
return m22;
case 11:
return m32;
case 12:
return m03;
case 13:
return m13;
case 14:
return m23;
case 15:
return m33;
}
return m00;
}
__host__ __device__ const float& cumat4::operator[](const int &I) const
{
switch(I)
{
case 0:
return m00;
case 1:
return m10;
case 2:
return m20;
case 3:
return m30;
case 4:
return m01;
case 5:
return m11;
case 6:
return m21;
case 7:
return m31;
case 8:
return m02;
case 9:
return m12;
case 10:
return m22;
case 11:
return m32;
case 12:
return m03;
case 13:
return m13;
case 14:
return m23;
case 15:
return m33;
}
return m00;
}
__host__ __device__ float& cumat4::operator()(const int &I, const int &J)
{
return (*this)[I+4*J];
}
__host__ __device__ const float& cumat4::operator()(const int &I, const int &J) const
{
return (*this)[I+4*J];
}
__host__ __device__ float& cumat4::at(const int &I, const int &J)
{
return (*this)[I+4*J];
}
__host__ __device__ const float& cumat4::at(const int &I, const int &J) const
{
return (*this)[I+4*J];
}
__host__ __device__ double* cumat4::data()
{
return (double*)this;
}
__host__ __device__ const double* cumat4::data() const
{
return (double*)this;
}
__host__ __device__ cumat4 cumat4::operator+(const cumat4& rhs) const
{
cumat4 ret;
ret.m00 = m00 + rhs.m00;
ret.m10 = m10 + rhs.m10;
ret.m20 = m20 + rhs.m20;
ret.m30 = m30 + rhs.m30;
ret.m01 = m01 + rhs.m01;
ret.m11 = m11 + rhs.m11;
ret.m21 = m21 + rhs.m21;
ret.m31 = m31 + rhs.m31;
ret.m02 = m02 + rhs.m02;
ret.m12 = m12 + rhs.m12;
ret.m22 = m22 + rhs.m22;
ret.m32 = m32 + rhs.m32;
ret.m03 = m03 + rhs.m03;
ret.m13 = m13 + rhs.m13;
ret.m23 = m23 + rhs.m23;
ret.m33 = m33 + rhs.m33;
return ret;
}
__host__ __device__ cumat4 cumat4::operator-(const cumat4& rhs) const
{
cumat4 ret;
ret.m00 = m00 - rhs.m00;
ret.m10 = m10 - rhs.m10;
ret.m20 = m20 - rhs.m20;
ret.m30 = m30 - rhs.m30;
ret.m01 = m01 - rhs.m01;
ret.m11 = m11 - rhs.m11;
ret.m21 = m21 - rhs.m21;
ret.m31 = m31 - rhs.m31;
ret.m02 = m02 - rhs.m02;
ret.m12 = m12 - rhs.m12;
ret.m22 = m22 - rhs.m22;
ret.m32 = m32 - rhs.m32;
ret.m03 = m03 - rhs.m03;
ret.m13 = m13 - rhs.m13;
ret.m23 = m23 - rhs.m23;
ret.m33 = m33 - rhs.m33;
return ret;
}
__host__ __device__ cumat4 cumat4::operator*(const cumat4& rhs) const
{
cumat4 ret; //should be zeroed in constructor
ret.m00 = m00*rhs.m00 + m01*rhs.m10 + m02*rhs.m20 + m03*rhs.m30;
ret.m01 = m00*rhs.m01 + m01*rhs.m11 + m02*rhs.m21 + m03*rhs.m31;
ret.m02 = m00*rhs.m02 + m01*rhs.m12 + m02*rhs.m22 + m03*rhs.m32;
ret.m03 = m00*rhs.m03 + m01*rhs.m13 + m02*rhs.m23 + m03*rhs.m33;
ret.m10 = m10*rhs.m00 + m11*rhs.m10 + m12*rhs.m20 + m13*rhs.m30;
ret.m11 = m10*rhs.m01 + m11*rhs.m11 + m12*rhs.m21 + m13*rhs.m31;
ret.m12 = m10*rhs.m02 + m11*rhs.m12 + m12*rhs.m22 + m13*rhs.m32;
ret.m13 = m10*rhs.m03 + m11*rhs.m13 + m12*rhs.m23 + m13*rhs.m33;
ret.m20 = m20*rhs.m00 + m21*rhs.m10 + m22*rhs.m20 + m23*rhs.m30;
ret.m21 = m20*rhs.m01 + m21*rhs.m11 + m22*rhs.m21 + m23*rhs.m31;
ret.m22 = m20*rhs.m02 + m21*rhs.m12 + m22*rhs.m22 + m23*rhs.m32;
ret.m23 = m20*rhs.m03 + m21*rhs.m13 + m22*rhs.m23 + m23*rhs.m33;
ret.m30 = m30*rhs.m00 + m31*rhs.m10 + m32*rhs.m20 + m33*rhs.m30;
ret.m31 = m30*rhs.m01 + m31*rhs.m11 + m32*rhs.m21 + m33*rhs.m31;
ret.m32 = m30*rhs.m02 + m31*rhs.m12 + m32*rhs.m22 + m33*rhs.m32;
ret.m33 = m30*rhs.m03 + m31*rhs.m13 + m32*rhs.m23 + m33*rhs.m33;
return ret;
}
__host__ __device__ cumat4 operator*(const cumat4& lhs, const double& rhs)
{
cumat4 ret;
ret.m00=lhs.m00*rhs;
ret.m10=lhs.m10*rhs;
ret.m20=lhs.m20*rhs;
ret.m30=lhs.m30*rhs;
ret.m01=lhs.m01*rhs;
ret.m11=lhs.m11*rhs;
ret.m21=lhs.m21*rhs;
ret.m31=lhs.m31*rhs;
ret.m02=lhs.m02*rhs;
ret.m12=lhs.m12*rhs;
ret.m22=lhs.m22*rhs;
ret.m32=lhs.m32*rhs;
ret.m03=lhs.m03*rhs;
ret.m13=lhs.m13*rhs;
ret.m23=lhs.m23*rhs;
ret.m33=lhs.m33*rhs;
return ret;
}
__host__ __device__ cumat4 operator/(const cumat4& lhs, const double& rhs)
{
cumat4 ret;
ret.m00=lhs.m00/rhs;
ret.m10=lhs.m10/rhs;
ret.m20=lhs.m20/rhs;
ret.m30=lhs.m30/rhs;
ret.m01=lhs.m01/rhs;
ret.m11=lhs.m11/rhs;
ret.m21=lhs.m21/rhs;
ret.m31=lhs.m31/rhs;
ret.m02=lhs.m02/rhs;
ret.m12=lhs.m12/rhs;
ret.m22=lhs.m22/rhs;
ret.m32=lhs.m32/rhs;
ret.m03=lhs.m03/rhs;
ret.m13=lhs.m13/rhs;
ret.m23=lhs.m23/rhs;
ret.m33=lhs.m33/rhs;
return ret;
}
__host__ __device__ cumat4 operator*(const double& lhs, const cumat4& rhs)
{
cumat4 ret;
ret.m00=lhs*rhs.m00;
ret.m10=lhs*rhs.m10;
ret.m20=lhs*rhs.m20;
ret.m30=lhs*rhs.m30;
ret.m01=lhs*rhs.m01;
ret.m11=lhs*rhs.m11;
ret.m21=lhs*rhs.m21;
ret.m31=lhs*rhs.m31;
ret.m02=lhs*rhs.m02;
ret.m12=lhs*rhs.m12;
ret.m22=lhs*rhs.m22;
ret.m32=lhs*rhs.m32;
ret.m03=lhs*rhs.m03;
ret.m13=lhs*rhs.m13;
ret.m23=lhs*rhs.m23;
ret.m33=lhs*rhs.m33;
return ret;
}
__host__ __device__ cuvec4 operator*(const cumat4& lhs, const cuvec4& rhs)
{
cuvec4 ret;
ret.x = lhs.m00*rhs.x + lhs.m01*rhs.y + lhs.m02*rhs.z + lhs.m03*rhs.w;
ret.y = lhs.m10*rhs.x + lhs.m11*rhs.y + lhs.m12*rhs.z + lhs.m13*rhs.w;
ret.z = lhs.m20*rhs.x + lhs.m21*rhs.y + lhs.m22*rhs.z + lhs.m23*rhs.w;
ret.w = lhs.m30*rhs.x + lhs.m31*rhs.y + lhs.m32*rhs.z + lhs.m33*rhs.w;
return ret;
}
__host__ __device__ cuvec4 operator*(const cuvec4& lhs, const cumat4& rhs)
{
cuvec4 ret;
ret.x = lhs.x*rhs.m00 + lhs.y*rhs.m10 + lhs.z*rhs.m20 + lhs.w*rhs.m30;
ret.y = lhs.x*rhs.m01 + lhs.y*rhs.m11 + lhs.z*rhs.m21 + lhs.w*rhs.m31;
ret.z = lhs.x*rhs.m02 + lhs.y*rhs.m12 + lhs.z*rhs.m22 + lhs.w*rhs.m32;
ret.w = lhs.x*rhs.m03 + lhs.y*rhs.m13 + lhs.z*rhs.m23 + lhs.w*rhs.m33;
return ret;
}
__host__ __device__ cumat4 operator-(const cumat4& rhs)
{
cumat4 ret;
ret.m00 = -rhs.m00;
ret.m10 = -rhs.m10;
ret.m20 = -rhs.m20;
ret.m30 = -rhs.m30;
ret.m01 = -rhs.m01;
ret.m11 = -rhs.m11;
ret.m21 = -rhs.m21;
ret.m31 = -rhs.m31;
ret.m02 = -rhs.m02;
ret.m12 = -rhs.m12;
ret.m22 = -rhs.m22;
ret.m32 = -rhs.m32;
ret.m03 = -rhs.m03;
ret.m13 = -rhs.m13;
ret.m23 = -rhs.m23;
ret.m33 = -rhs.m33;
return ret;
}
__host__ __device__ cumat4& cumat4::operator+=(const cumat4& rhs)
{
m00 += rhs.m00;
m10 += rhs.m10;
m20 += rhs.m20;
m30 += rhs.m30;
m01 += rhs.m01;
m11 += rhs.m11;
m21 += rhs.m21;
m31 += rhs.m31;
m02 += rhs.m02;
m12 += rhs.m12;
m22 += rhs.m22;
m32 += rhs.m32;
m03 += rhs.m03;
m13 += rhs.m13;
m23 += rhs.m23;
m33 += rhs.m33;
return *this;
}
__host__ __device__ cumat4& cumat4::operator-=(const cumat4& rhs)
{
m00 -= rhs.m00;
m10 -= rhs.m10;
m20 -= rhs.m20;
m30 -= rhs.m30;
m01 -= rhs.m01;
m11 -= rhs.m11;
m21 -= rhs.m21;
m31 -= rhs.m31;
m02 -= rhs.m02;
m12 -= rhs.m12;
m22 -= rhs.m22;
m32 -= rhs.m32;
m03 -= rhs.m03;
m13 -= rhs.m13;
m23 -= rhs.m23;
m33 -= rhs.m33;
return *this;
}
__host__ __device__ cumat4& cumat4::operator*=(const double& rhs)
{
m00 *= rhs;
m10 *= rhs;
m20 *= rhs;
m30 *= rhs;
m01 *= rhs;
m11 *= rhs;
m21 *= rhs;
m31 *= rhs;
m02 *= rhs;
m12 *= rhs;
m22 *= rhs;
m32 *= rhs;
m03 *= rhs;
m13 *= rhs;
m23 *= rhs;
m33 *= rhs;
return *this;
}
__host__ __device__ cumat4& cumat4::operator/=(const double& rhs)
{
m00 /= rhs;
m10 /= rhs;
m20 /= rhs;
m30 /= rhs;
m01 /= rhs;
m11 /= rhs;
m21 /= rhs;
m31 /= rhs;
m02 /= rhs;
m12 /= rhs;
m22 /= rhs;
m32 /= rhs;
m03 /= rhs;
m13 /= rhs;
m23 /= rhs;
m33 /= rhs;
return *this;
}
__host__ __device__ cumat4& cumat4::operator*=(const cumat4& rhs)
{
cumat4 tmp = *this;
m00 = tmp.m00*rhs.m00 + tmp.m01*rhs.m10 + tmp.m02*rhs.m20 + tmp.m03*rhs.m30;
m01 = tmp.m00*rhs.m01 + tmp.m01*rhs.m11 + tmp.m02*rhs.m21 + tmp.m03*rhs.m31;
m02 = tmp.m00*rhs.m02 + tmp.m01*rhs.m12 + tmp.m02*rhs.m22 + tmp.m03*rhs.m32;
m03 = tmp.m00*rhs.m03 + tmp.m01*rhs.m13 + tmp.m02*rhs.m23 + tmp.m03*rhs.m33;
m10 = tmp.m10*rhs.m00 + tmp.m11*rhs.m10 + tmp.m12*rhs.m20 + tmp.m13*rhs.m30;
m11 = tmp.m10*rhs.m01 + tmp.m11*rhs.m11 + tmp.m12*rhs.m21 + tmp.m13*rhs.m31;
m12 = tmp.m10*rhs.m02 + tmp.m11*rhs.m12 + tmp.m12*rhs.m22 + tmp.m13*rhs.m32;
m13 = tmp.m10*rhs.m03 + tmp.m11*rhs.m13 + tmp.m12*rhs.m23 + tmp.m13*rhs.m33;
m20 = tmp.m20*rhs.m00 + tmp.m21*rhs.m10 + tmp.m22*rhs.m20 + tmp.m23*rhs.m30;
m21 = tmp.m20*rhs.m01 + tmp.m21*rhs.m11 + tmp.m22*rhs.m21 + tmp.m23*rhs.m31;
m22 = tmp.m20*rhs.m02 + tmp.m21*rhs.m12 + tmp.m22*rhs.m22 + tmp.m23*rhs.m32;
m23 = tmp.m20*rhs.m03 + tmp.m21*rhs.m13 + tmp.m22*rhs.m23 + tmp.m23*rhs.m33;
m30 = tmp.m30*rhs.m00 + tmp.m31*rhs.m10 + tmp.m32*rhs.m20 + tmp.m33*rhs.m30;
m31 = tmp.m30*rhs.m01 + tmp.m31*rhs.m11 + tmp.m32*rhs.m21 + tmp.m33*rhs.m31;
m32 = tmp.m30*rhs.m02 + tmp.m31*rhs.m12 + tmp.m32*rhs.m22 + tmp.m33*rhs.m32;
m33 = tmp.m30*rhs.m03 + tmp.m31*rhs.m13 + tmp.m32*rhs.m23 + tmp.m33*rhs.m33;
return *this;
}
__host__ __device__ cumat4 cumat4::transpose() const
{
cumat4 ret;
ret.m00 = m00;
ret.m10 = m01;
ret.m20 = m02;
ret.m30 = m03;
ret.m01 = m10;
ret.m11 = m11;
ret.m21 = m12;
ret.m31 = m13;
ret.m02 = m20;
ret.m12 = m21;
ret.m22 = m22;
ret.m32 = m23;
ret.m03 = m30;
ret.m13 = m31;
ret.m23 = m32;
ret.m33 = m33;
return ret;
}

58
test_scripts/cuvec_codegen/cuvec4_codegen1.hpp Normal file
View File

@ -0,0 +1,58 @@
__host__ __device__ cuvec4();
__host__ __device__ ~cuvec4();
__host__ __device__ cuvec4(const double &_x, const double &_y, const double &_z, const double &_w);
__host__ __device__ float& operator[](const int &I);
__host__ __device__ const float& operator[](const int &I) const;
__host__ __device__ cuvec4 operator+(const cuvec4& rhs) const;
__host__ __device__ cuvec4 operator-(const cuvec4& rhs) const;
__host__ __device__ cuvec4 operator*(const cuvec4& rhs) const; //elementwise product
__host__ __device__ cuvec4 operator/(const cuvec4& rhs) const; //elementwise division
__host__ __device__ friend cuvec4 operator*(const cuvec4& lhs, const double& rhs);
__host__ __device__ friend cuvec4 operator*(const double& lhs, const cuvec4& rhs);
__host__ __device__ friend cuvec4 operator/(const cuvec4& lhs, const double& rhs);
__host__ __device__ friend cuvec4 operator/(const double& lhs, const cuvec4& rhs);
__host__ __device__ friend cuvec4 operator-(const cuvec4& other);
__host__ __device__ cuvec4& operator+=(const cuvec4& rhs);
__host__ __device__ cuvec4& operator-=(const cuvec4& rhs);
__host__ __device__ cuvec4& operator*=(const double& rhs);
__host__ __device__ cuvec4& operator/=(const double& rhs);
//
//Matrix Header Stuff
//
double m00,m10,m20,m30;
double m01,m11,m21,m31;
double m02,m12,m22,m32;
double m03,m13,m23,m33;
__host__ __device__ cumat4();
__host__ __device__ ~cumat4();
__host__ __device__ cumat4(
const double& _m00, const double& _m01, const double& _m02, const double& _m03,
const double& _m10, const double& _m11, const double& _m12, const double& _m13,
const double& _m20, const double& _m21, const double& _m22, const double& _m23,
const double& _m30, const double& _m31, const double& _m32, const double& _m33
);
__host__ __device__ cumat4(const double* data16);
__host__ __device__ float& operator[](const int &I);
__host__ __device__ const float& operator[](const int &I) const;
__host__ __device__ float& operator()(const int &I, const int &J);
__host__ __device__ const float& operator()(const int &I, const int &J) const;
__host__ __device__ float& at(const int &I, const int &J);
__host__ __device__ const float& at(const int &I, const int &J) const;
__host__ __device__ double* data(); //pointer to float16 representation of matrix
__host__ __device__ const double* data() const; //pointer to float16 representation of matrix
__host__ __device__ cumat4 operator+(const cumat4& rhs) const;
__host__ __device__ cumat4 operator-(const cumat4& rhs) const;
__host__ __device__ cumat4 operator*(const cumat4& rhs) const;
__host__ __device__ friend cumat4 operator*(const cumat4& lhs, const double& rhs);
__host__ __device__ friend cumat4 operator/(const cumat4& lhs, const double& rhs);
__host__ __device__ friend cumat4 operator*(const double& lhs, const cumat4& rhs);
__host__ __device__ friend cuvec4 operator*(const cumat4& lhs, const cuvec4& rhs);
__host__ __device__ friend cuvec4 operator*(const cuvec4& lhs, const cumat4& rhs);
__host__ __device__ friend cumat4 operator-(const cumat4& rhs);
__host__ __device__ cumat4& operator+=(const cumat4& rhs);
__host__ __device__ cumat4& operator-=(const cumat4& rhs);
__host__ __device__ cumat4& operator*=(const double& rhs);
__host__ __device__ cumat4& operator/=(const double& rhs);
__host__ __device__ cumat4& operator*=(const cumat4& rhs);
__host__ __device__ cumat4 transpose() const;

718
test_scripts/cuvec_codegen/cuvec4f_codegen1.cu Normal file
View File

@ -0,0 +1,718 @@
__host__ __device__ cuvec4f::cuvec4f()
{
x = 0; y = 0; z = 0; w = 0;
return;
}
__host__ __device__ cuvec4f::~cuvec4f()
{
x = 0; y = 0; z = 0; w = 0;
return;
}
__host__ __device__ cuvec4f::cuvec4f(const float &_x, const float &_y, const float &_z, const float &_w)
{
x = _x; y = _y; z = _z; w = _w;
return;
}
__host__ __device__ float& cuvec4f::operator[](const int &I)
{
switch(I)
{
case 0:
return x;
case 1:
return y;
case 2:
return z;
case 3:
return w;
}
return x;
}
__host__ __device__ const float& cuvec4f::operator[](const int &I) const
{
switch(I)
{
case 0:
return x;
case 1:
return y;
case 2:
return z;
case 3:
return w;
}
return x;
}
__host__ __device__ cuvec4f cuvec4f::operator+(const cuvec4f& rhs) const
{
cuvec4f ret;
ret.x = x + rhs.x;
ret.y = y + rhs.y;
ret.z = z + rhs.z;
ret.w = w + rhs.w;
return ret;
}
__host__ __device__ cuvec4f cuvec4f::operator-(const cuvec4f& rhs) const
{
cuvec4f ret;
ret.x = x - rhs.x;
ret.y = y - rhs.y;
ret.z = z - rhs.z;
ret.w = w - rhs.w;
return ret;
}
__host__ __device__ cuvec4f cuvec4f::operator*(const cuvec4f& rhs) const
{
//Elementwise product
cuvec4f ret;
ret.x = x * rhs.x;
ret.y = y * rhs.y;
ret.z = z * rhs.z;
ret.w = w * rhs.w;
return ret;
}
__host__ __device__ cuvec4f cuvec4f::operator/(const cuvec4f& rhs) const
{
//Elementwise division
cuvec4f ret;
ret.x = x / rhs.x;
ret.y = y / rhs.y;
ret.z = z / rhs.z;
ret.w = w / rhs.w;
return ret;
}
__host__ __device__ cuvec4f operator*(const cuvec4f& lhs, const float& rhs)
{
cuvec4f ret;
ret.x = lhs.x*rhs;
ret.y = lhs.y*rhs;
ret.z = lhs.z*rhs;
ret.w = lhs.w*rhs;
return ret;
}
__host__ __device__ cuvec4f operator*(const float& lhs, const cuvec4f& rhs)
{
cuvec4f ret;
ret.x = lhs*rhs.x;
ret.y = lhs*rhs.y;
ret.z = lhs*rhs.z;
ret.w = lhs*rhs.w;
return ret;
}
__host__ __device__ cuvec4f operator/(const cuvec4f& lhs, const float& rhs)
{
cuvec4f ret;
ret.x = lhs.x/rhs;
ret.y = lhs.y/rhs;
ret.z = lhs.z/rhs;
ret.w = lhs.w/rhs;
return ret;
}
__host__ __device__ cuvec4f operator/(const float& lhs, const cuvec4f& rhs)
{
cuvec4f ret;
ret.x = lhs/rhs.x;
ret.y = lhs/rhs.y;
ret.z = lhs/rhs.z;
ret.w = lhs/rhs.w;
return ret;
}
__host__ __device__ cuvec4f operator-(const cuvec4f& other)
{
cuvec4f ret;
ret.x = -other.x;
ret.y = -other.y;
ret.z = -other.z;
ret.w = -other.w;
return ret;
}
__host__ __device__ cuvec4f& cuvec4f::operator+=(const cuvec4f& rhs)
{
x += rhs.x;
y += rhs.y;
z += rhs.z;
w += rhs.w;
return *this;
}
__host__ __device__ cuvec4f& cuvec4f::operator-=(const cuvec4f& rhs)
{
x -= rhs.x;
y -= rhs.y;
z -= rhs.z;
w -= rhs.w;
return *this;
}
__host__ __device__ cuvec4f& cuvec4f::operator*=(const float& rhs)
{
x *= rhs;
y *= rhs;
z *= rhs;
w *= rhs;
return *this;
}
__host__ __device__ cuvec4f& cuvec4f::operator/=(const float& rhs)
{
x /= rhs;
y /= rhs;
z /= rhs;
w /= rhs;
return *this;
}
//
//Matrix Header Stuff
//
__host__ __device__ cumat4f::cumat4f()
{
m00 = 0;
m01 = 0;
m02 = 0;
m03 = 0;
m10 = 0;
m11 = 0;
m12 = 0;
m13 = 0;
m20 = 0;
m21 = 0;
m22 = 0;
m23 = 0;
m30 = 0;
m31 = 0;
m32 = 0;
m33 = 0;
return;
}
__host__ __device__ cumat4f::~cumat4f()
{
//m00 = 0;
//m01 = 0;
//m02 = 0;
//m03 = 0;
//m10 = 0;
//m11 = 0;
//m12 = 0;
//m13 = 0;
//m20 = 0;
//m21 = 0;
//m22 = 0;
//m23 = 0;
//m30 = 0;
//m31 = 0;
//m32 = 0;
//m33 = 0;
return;
}
__host__ __device__ cumat4f::cumat4f(
const float& _m00, const float& _m01, const float& _m02, const float& _m03,
const float& _m10, const float& _m11, const float& _m12, const float& _m13,
const float& _m20, const float& _m21, const float& _m22, const float& _m23,
const float& _m30, const float& _m31, const float& _m32, const float& _m33
)
{
m00 = _m00;
m10 = _m10;
m20 = _m20;
m30 = _m30;
m01 = _m01;
m11 = _m11;
m21 = _m21;
m31 = _m31;
m02 = _m02;
m12 = _m12;
m22 = _m22;
m32 = _m32;
m03 = _m03;
m13 = _m13;
m23 = _m23;
m33 = _m33;
return;
}
__host__ __device__ cumat4f::cumat4f(const float* data16)
{
m00 = data16[0];
m10 = data16[1];
m20 = data16[2];
m30 = data16[3];
m01 = data16[4];
m11 = data16[5];
m21 = data16[6];
m31 = data16[7];
m02 = data16[8];
m12 = data16[9];
m22 = data16[10];
m32 = data16[11];
m03 = data16[12];
m13 = data16[13];
m23 = data16[14];
m33 = data16[15];
return;
}
__host__ __device__ float& cumat4f::operator[](const int &I)
{
switch(I)
{
case 0:
return m00;
case 1:
return m10;
case 2:
return m20;
case 3:
return m30;
case 4:
return m01;
case 5:
return m11;
case 6:
return m21;
case 7:
return m31;
case 8:
return m02;
case 9:
return m12;
case 10:
return m22;
case 11:
return m32;
case 12:
return m03;
case 13:
return m13;
case 14:
return m23;
case 15:
return m33;
}
return m00;
}
__host__ __device__ const float& cumat4f::operator[](const int &I) const
{
switch(I)
{
case 0:
return m00;
case 1:
return m10;
case 2:
return m20;
case 3:
return m30;
case 4:
return m01;
case 5:
return m11;
case 6:
return m21;
case 7:
return m31;
case 8:
return m02;
case 9:
return m12;
case 10:
return m22;
case 11:
return m32;
case 12:
return m03;
case 13:
return m13;
case 14:
return m23;
case 15:
return m33;
}
return m00;
}
__host__ __device__ float& cumat4f::operator()(const int &I, const int &J)
{
return (*this)[I+4*J];
}
__host__ __device__ const float& cumat4f::operator()(const int &I, const int &J) const
{
return (*this)[I+4*J];
}
__host__ __device__ float& cumat4f::at(const int &I, const int &J)
{
return (*this)[I+4*J];
}
__host__ __device__ const float& cumat4f::at(const int &I, const int &J) const
{
return (*this)[I+4*J];
}
__host__ __device__ float* cumat4f::data()
{
return (float*)this;
}
__host__ __device__ const float* cumat4f::data() const
{
return (float*)this;
}
__host__ __device__ cumat4f cumat4f::operator+(const cumat4f& rhs) const
{
cumat4f ret;
ret.m00 = m00 + rhs.m00;
ret.m10 = m10 + rhs.m10;
ret.m20 = m20 + rhs.m20;
ret.m30 = m30 + rhs.m30;
ret.m01 = m01 + rhs.m01;
ret.m11 = m11 + rhs.m11;
ret.m21 = m21 + rhs.m21;
ret.m31 = m31 + rhs.m31;
ret.m02 = m02 + rhs.m02;
ret.m12 = m12 + rhs.m12;
ret.m22 = m22 + rhs.m22;
ret.m32 = m32 + rhs.m32;
ret.m03 = m03 + rhs.m03;
ret.m13 = m13 + rhs.m13;
ret.m23 = m23 + rhs.m23;
ret.m33 = m33 + rhs.m33;
return ret;
}
__host__ __device__ cumat4f cumat4f::operator-(const cumat4f& rhs) const
{
cumat4f ret;
ret.m00 = m00 - rhs.m00;
ret.m10 = m10 - rhs.m10;
ret.m20 = m20 - rhs.m20;
ret.m30 = m30 - rhs.m30;
ret.m01 = m01 - rhs.m01;
ret.m11 = m11 - rhs.m11;
ret.m21 = m21 - rhs.m21;
ret.m31 = m31 - rhs.m31;
ret.m02 = m02 - rhs.m02;
ret.m12 = m12 - rhs.m12;
ret.m22 = m22 - rhs.m22;
ret.m32 = m32 - rhs.m32;
ret.m03 = m03 - rhs.m03;
ret.m13 = m13 - rhs.m13;
ret.m23 = m23 - rhs.m23;
ret.m33 = m33 - rhs.m33;
return ret;
}
__host__ __device__ cumat4f cumat4f::operator*(const cumat4f& rhs) const
{
cumat4f ret; //should be zeroed in constructor
ret.m00 = m00*rhs.m00 + m01*rhs.m10 + m02*rhs.m20 + m03*rhs.m30;
ret.m01 = m00*rhs.m01 + m01*rhs.m11 + m02*rhs.m21 + m03*rhs.m31;
ret.m02 = m00*rhs.m02 + m01*rhs.m12 + m02*rhs.m22 + m03*rhs.m32;
ret.m03 = m00*rhs.m03 + m01*rhs.m13 + m02*rhs.m23 + m03*rhs.m33;
ret.m10 = m10*rhs.m00 + m11*rhs.m10 + m12*rhs.m20 + m13*rhs.m30;
ret.m11 = m10*rhs.m01 + m11*rhs.m11 + m12*rhs.m21 + m13*rhs.m31;
ret.m12 = m10*rhs.m02 + m11*rhs.m12 + m12*rhs.m22 + m13*rhs.m32;
ret.m13 = m10*rhs.m03 + m11*rhs.m13 + m12*rhs.m23 + m13*rhs.m33;
ret.m20 = m20*rhs.m00 + m21*rhs.m10 + m22*rhs.m20 + m23*rhs.m30;
ret.m21 = m20*rhs.m01 + m21*rhs.m11 + m22*rhs.m21 + m23*rhs.m31;
ret.m22 = m20*rhs.m02 + m21*rhs.m12 + m22*rhs.m22 + m23*rhs.m32;
ret.m23 = m20*rhs.m03 + m21*rhs.m13 + m22*rhs.m23 + m23*rhs.m33;
ret.m30 = m30*rhs.m00 + m31*rhs.m10 + m32*rhs.m20 + m33*rhs.m30;
ret.m31 = m30*rhs.m01 + m31*rhs.m11 + m32*rhs.m21 + m33*rhs.m31;
ret.m32 = m30*rhs.m02 + m31*rhs.m12 + m32*rhs.m22 + m33*rhs.m32;
ret.m33 = m30*rhs.m03 + m31*rhs.m13 + m32*rhs.m23 + m33*rhs.m33;
return ret;
}
__host__ __device__ cumat4f operator*(const cumat4f& lhs, const float& rhs)
{
cumat4f ret;
ret.m00=lhs.m00*rhs;
ret.m10=lhs.m10*rhs;
ret.m20=lhs.m20*rhs;
ret.m30=lhs.m30*rhs;
ret.m01=lhs.m01*rhs;
ret.m11=lhs.m11*rhs;
ret.m21=lhs.m21*rhs;
ret.m31=lhs.m31*rhs;
ret.m02=lhs.m02*rhs;
ret.m12=lhs.m12*rhs;
ret.m22=lhs.m22*rhs;
ret.m32=lhs.m32*rhs;
ret.m03=lhs.m03*rhs;
ret.m13=lhs.m13*rhs;
ret.m23=lhs.m23*rhs;
ret.m33=lhs.m33*rhs;
return ret;
}
__host__ __device__ cumat4f operator/(const cumat4f& lhs, const float& rhs)
{
cumat4f ret;
ret.m00=lhs.m00/rhs;
ret.m10=lhs.m10/rhs;
ret.m20=lhs.m20/rhs;
ret.m30=lhs.m30/rhs;
ret.m01=lhs.m01/rhs;
ret.m11=lhs.m11/rhs;
ret.m21=lhs.m21/rhs;
ret.m31=lhs.m31/rhs;
ret.m02=lhs.m02/rhs;
ret.m12=lhs.m12/rhs;
ret.m22=lhs.m22/rhs;
ret.m32=lhs.m32/rhs;
ret.m03=lhs.m03/rhs;
ret.m13=lhs.m13/rhs;
ret.m23=lhs.m23/rhs;
ret.m33=lhs.m33/rhs;
return ret;
}
__host__ __device__ cumat4f operator*(const float& lhs, const cumat4f& rhs)
{
cumat4f ret;
ret.m00=lhs*rhs.m00;
ret.m10=lhs*rhs.m10;
ret.m20=lhs*rhs.m20;
ret.m30=lhs*rhs.m30;
ret.m01=lhs*rhs.m01;
ret.m11=lhs*rhs.m11;
ret.m21=lhs*rhs.m21;
ret.m31=lhs*rhs.m31;
ret.m02=lhs*rhs.m02;
ret.m12=lhs*rhs.m12;
ret.m22=lhs*rhs.m22;
ret.m32=lhs*rhs.m32;
ret.m03=lhs*rhs.m03;
ret.m13=lhs*rhs.m13;
ret.m23=lhs*rhs.m23;
ret.m33=lhs*rhs.m33;
return ret;
}
__host__ __device__ cuvec4f operator*(const cumat4f& lhs, const cuvec4f& rhs)
{
cuvec4f ret;
ret.x = lhs.m00*rhs.x + lhs.m01*rhs.y + lhs.m02*rhs.z + lhs.m03*rhs.w;
ret.y = lhs.m10*rhs.x + lhs.m11*rhs.y + lhs.m12*rhs.z + lhs.m13*rhs.w;
ret.z = lhs.m20*rhs.x + lhs.m21*rhs.y + lhs.m22*rhs.z + lhs.m23*rhs.w;
ret.w = lhs.m30*rhs.x + lhs.m31*rhs.y + lhs.m32*rhs.z + lhs.m33*rhs.w;
return ret;
}
__host__ __device__ cuvec4f operator*(const cuvec4f& lhs, const cumat4f& rhs)
{
cuvec4f ret;
ret.x = lhs.x*rhs.m00 + lhs.y*rhs.m10 + lhs.z*rhs.m20 + lhs.w*rhs.m30;
ret.y = lhs.x*rhs.m01 + lhs.y*rhs.m11 + lhs.z*rhs.m21 + lhs.w*rhs.m31;
ret.z = lhs.x*rhs.m02 + lhs.y*rhs.m12 + lhs.z*rhs.m22 + lhs.w*rhs.m32;
ret.w = lhs.x*rhs.m03 + lhs.y*rhs.m13 + lhs.z*rhs.m23 + lhs.w*rhs.m33;
return ret;
}
__host__ __device__ cumat4f operator-(const cumat4f& rhs)
{
cumat4f ret;
ret.m00 = -rhs.m00;
ret.m10 = -rhs.m10;
ret.m20 = -rhs.m20;
ret.m30 = -rhs.m30;
ret.m01 = -rhs.m01;
ret.m11 = -rhs.m11;
ret.m21 = -rhs.m21;
ret.m31 = -rhs.m31;
ret.m02 = -rhs.m02;
ret.m12 = -rhs.m12;
ret.m22 = -rhs.m22;
ret.m32 = -rhs.m32;
ret.m03 = -rhs.m03;
ret.m13 = -rhs.m13;
ret.m23 = -rhs.m23;
ret.m33 = -rhs.m33;
return ret;
}
__host__ __device__ cumat4f& cumat4f::operator+=(const cumat4f& rhs)
{
m00 += rhs.m00;
m10 += rhs.m10;
m20 += rhs.m20;
m30 += rhs.m30;
m01 += rhs.m01;
m11 += rhs.m11;
m21 += rhs.m21;
m31 += rhs.m31;
m02 += rhs.m02;
m12 += rhs.m12;
m22 += rhs.m22;
m32 += rhs.m32;
m03 += rhs.m03;
m13 += rhs.m13;
m23 += rhs.m23;
m33 += rhs.m33;
return *this;
}
__host__ __device__ cumat4f& cumat4f::operator-=(const cumat4f& rhs)
{
m00 -= rhs.m00;
m10 -= rhs.m10;
m20 -= rhs.m20;
m30 -= rhs.m30;
m01 -= rhs.m01;
m11 -= rhs.m11;
m21 -= rhs.m21;
m31 -= rhs.m31;
m02 -= rhs.m02;
m12 -= rhs.m12;
m22 -= rhs.m22;
m32 -= rhs.m32;
m03 -= rhs.m03;
m13 -= rhs.m13;
m23 -= rhs.m23;
m33 -= rhs.m33;
return *this;
}
__host__ __device__ cumat4f& cumat4f::operator*=(const float& rhs)
{
m00 *= rhs;
m10 *= rhs;
m20 *= rhs;
m30 *= rhs;
m01 *= rhs;
m11 *= rhs;
m21 *= rhs;
m31 *= rhs;
m02 *= rhs;
m12 *= rhs;
m22 *= rhs;
m32 *= rhs;
m03 *= rhs;
m13 *= rhs;
m23 *= rhs;
m33 *= rhs;
return *this;
}
__host__ __device__ cumat4f& cumat4f::operator/=(const float& rhs)
{
m00 /= rhs;
m10 /= rhs;
m20 /= rhs;
m30 /= rhs;
m01 /= rhs;
m11 /= rhs;
m21 /= rhs;
m31 /= rhs;
m02 /= rhs;
m12 /= rhs;
m22 /= rhs;
m32 /= rhs;
m03 /= rhs;
m13 /= rhs;
m23 /= rhs;
m33 /= rhs;
return *this;
}
__host__ __device__ cumat4f& cumat4f::operator*=(const cumat4f& rhs)
{
cumat4f tmp = *this;
m00 = tmp.m00*rhs.m00 + tmp.m01*rhs.m10 + tmp.m02*rhs.m20 + tmp.m03*rhs.m30;
m01 = tmp.m00*rhs.m01 + tmp.m01*rhs.m11 + tmp.m02*rhs.m21 + tmp.m03*rhs.m31;
m02 = tmp.m00*rhs.m02 + tmp.m01*rhs.m12 + tmp.m02*rhs.m22 + tmp.m03*rhs.m32;
m03 = tmp.m00*rhs.m03 + tmp.m01*rhs.m13 + tmp.m02*rhs.m23 + tmp.m03*rhs.m33;
m10 = tmp.m10*rhs.m00 + tmp.m11*rhs.m10 + tmp.m12*rhs.m20 + tmp.m13*rhs.m30;
m11 = tmp.m10*rhs.m01 + tmp.m11*rhs.m11 + tmp.m12*rhs.m21 + tmp.m13*rhs.m31;
m12 = tmp.m10*rhs.m02 + tmp.m11*rhs.m12 + tmp.m12*rhs.m22 + tmp.m13*rhs.m32;
m13 = tmp.m10*rhs.m03 + tmp.m11*rhs.m13 + tmp.m12*rhs.m23 + tmp.m13*rhs.m33;
m20 = tmp.m20*rhs.m00 + tmp.m21*rhs.m10 + tmp.m22*rhs.m20 + tmp.m23*rhs.m30;
m21 = tmp.m20*rhs.m01 + tmp.m21*rhs.m11 + tmp.m22*rhs.m21 + tmp.m23*rhs.m31;
m22 = tmp.m20*rhs.m02 + tmp.m21*rhs.m12 + tmp.m22*rhs.m22 + tmp.m23*rhs.m32;
m23 = tmp.m20*rhs.m03 + tmp.m21*rhs.m13 + tmp.m22*rhs.m23 + tmp.m23*rhs.m33;
m30 = tmp.m30*rhs.m00 + tmp.m31*rhs.m10 + tmp.m32*rhs.m20 + tmp.m33*rhs.m30;
m31 = tmp.m30*rhs.m01 + tmp.m31*rhs.m11 + tmp.m32*rhs.m21 + tmp.m33*rhs.m31;
m32 = tmp.m30*rhs.m02 + tmp.m31*rhs.m12 + tmp.m32*rhs.m22 + tmp.m33*rhs.m32;
m33 = tmp.m30*rhs.m03 + tmp.m31*rhs.m13 + tmp.m32*rhs.m23 + tmp.m33*rhs.m33;
return *this;
}
__host__ __device__ cumat4f cumat4f::transpose() const
{
cumat4f ret;
ret.m00 = m00;
ret.m10 = m01;
ret.m20 = m02;
ret.m30 = m03;
ret.m01 = m10;
ret.m11 = m11;
ret.m21 = m12;
ret.m31 = m13;
ret.m02 = m20;
ret.m12 = m21;
ret.m22 = m22;
ret.m32 = m23;
ret.m03 = m30;
ret.m13 = m31;
ret.m23 = m32;
ret.m33 = m33;
return ret;
}

58
test_scripts/cuvec_codegen/cuvec4f_codegen1.hpp Normal file
View File

@ -0,0 +1,58 @@
__host__ __device__ cuvec4f();
__host__ __device__ ~cuvec4f();
__host__ __device__ cuvec4f(const float &_x, const float &_y, const float &_z, const float &_w);
__host__ __device__ float& operator[](const int &I);
__host__ __device__ const float& operator[](const int &I) const;
__host__ __device__ cuvec4f operator+(const cuvec4f& rhs) const;
__host__ __device__ cuvec4f operator-(const cuvec4f& rhs) const;
__host__ __device__ cuvec4f operator*(const cuvec4f& rhs) const; //elementwise product
__host__ __device__ cuvec4f operator/(const cuvec4f& rhs) const; //elementwise division
__host__ __device__ friend cuvec4f operator*(const cuvec4f& lhs, const float& rhs);
__host__ __device__ friend cuvec4f operator*(const float& lhs, const cuvec4f& rhs);
__host__ __device__ friend cuvec4f operator/(const cuvec4f& lhs, const float& rhs);
__host__ __device__ friend cuvec4f operator/(const float& lhs, const cuvec4f& rhs);
__host__ __device__ friend cuvec4f operator-(const cuvec4f& other);
__host__ __device__ cuvec4f& operator+=(const cuvec4f& rhs);
__host__ __device__ cuvec4f& operator-=(const cuvec4f& rhs);
__host__ __device__ cuvec4f& operator*=(const float& rhs);
__host__ __device__ cuvec4f& operator/=(const float& rhs);
//
//Matrix Header Stuff
//
float m00,m10,m20,m30;
float m01,m11,m21,m31;
float m02,m12,m22,m32;
float m03,m13,m23,m33;
__host__ __device__ cumat4f();
__host__ __device__ ~cumat4f();
__host__ __device__ cumat4f(
const float& _m00, const float& _m01, const float& _m02, const float& _m03,
const float& _m10, const float& _m11, const float& _m12, const float& _m13,
const float& _m20, const float& _m21, const float& _m22, const float& _m23,
const float& _m30, const float& _m31, const float& _m32, const float& _m33
);
__host__ __device__ cumat4f(const float* data16);
__host__ __device__ float& operator[](const int &I);
__host__ __device__ const float& operator[](const int &I) const;
__host__ __device__ float& operator()(const int &I, const int &J);
__host__ __device__ const float& operator()(const int &I, const int &J) const;
__host__ __device__ float& at(const int &I, const int &J);
__host__ __device__ const float& at(const int &I, const int &J) const;
__host__ __device__ float* data(); //pointer to float16 representation of matrix
__host__ __device__ const float* data() const; //pointer to float16 representation of matrix
__host__ __device__ cumat4f operator+(const cumat4f& rhs) const;
__host__ __device__ cumat4f operator-(const cumat4f& rhs) const;
__host__ __device__ cumat4f operator*(const cumat4f& rhs) const;
__host__ __device__ friend cumat4f operator*(const cumat4f& lhs, const float& rhs);
__host__ __device__ friend cumat4f operator/(const cumat4f& lhs, const float& rhs);
__host__ __device__ friend cumat4f operator*(const float& lhs, const cumat4f& rhs);
__host__ __device__ friend cuvec4f operator*(const cumat4f& lhs, const cuvec4f& rhs);
__host__ __device__ friend cuvec4f operator*(const cuvec4f& lhs, const cumat4f& rhs);
__host__ __device__ friend cumat4f operator-(const cumat4f& rhs);
__host__ __device__ cumat4f& operator+=(const cumat4f& rhs);
__host__ __device__ cumat4f& operator-=(const cumat4f& rhs);
__host__ __device__ cumat4f& operator*=(const float& rhs);
__host__ __device__ cumat4f& operator/=(const float& rhs);
__host__ __device__ cumat4f& operator*=(const cumat4f& rhs);
__host__ __device__ cumat4f transpose() const;

718
test_scripts/cuvec_codegen/cuvec4i_codegen1.cu Normal file
View File

@ -0,0 +1,718 @@
__host__ __device__ cuvec4i::cuvec4i()
{
x = 0; y = 0; z = 0; w = 0;
return;
}
__host__ __device__ cuvec4i::~cuvec4i()
{
x = 0; y = 0; z = 0; w = 0;
return;
}
__host__ __device__ cuvec4i::cuvec4i(const int &_x, const int &_y, const int &_z, const int &_w)
{
x = _x; y = _y; z = _z; w = _w;
return;
}
__host__ __device__ float& cuvec4i::operator[](const int &I)
{
switch(I)
{
case 0:
return x;
case 1:
return y;
case 2:
return z;
case 3:
return w;
}
return x;
}
__host__ __device__ const float& cuvec4i::operator[](const int &I) const
{
switch(I)
{
case 0:
return x;
case 1:
return y;
case 2:
return z;
case 3:
return w;
}
return x;
}
__host__ __device__ cuvec4i cuvec4i::operator+(const cuvec4i& rhs) const
{
cuvec4i ret;
ret.x = x + rhs.x;
ret.y = y + rhs.y;
ret.z = z + rhs.z;
ret.w = w + rhs.w;
return ret;
}
__host__ __device__ cuvec4i cuvec4i::operator-(const cuvec4i& rhs) const
{
cuvec4i ret;
ret.x = x - rhs.x;
ret.y = y - rhs.y;
ret.z = z - rhs.z;
ret.w = w - rhs.w;
return ret;
}
__host__ __device__ cuvec4i cuvec4i::operator*(const cuvec4i& rhs) const
{
//Elementwise product
cuvec4i ret;
ret.x = x * rhs.x;
ret.y = y * rhs.y;
ret.z = z * rhs.z;
ret.w = w * rhs.w;
return ret;
}
__host__ __device__ cuvec4i cuvec4i::operator/(const cuvec4i& rhs) const
{
//Elementwise division
cuvec4i ret;
ret.x = x / rhs.x;
ret.y = y / rhs.y;
ret.z = z / rhs.z;
ret.w = w / rhs.w;
return ret;
}
__host__ __device__ cuvec4i operator*(const cuvec4i& lhs, const int& rhs)
{
cuvec4i ret;
ret.x = lhs.x*rhs;
ret.y = lhs.y*rhs;
ret.z = lhs.z*rhs;
ret.w = lhs.w*rhs;
return ret;
}
__host__ __device__ cuvec4i operator*(const int& lhs, const cuvec4i& rhs)
{
cuvec4i ret;
ret.x = lhs*rhs.x;
ret.y = lhs*rhs.y;
ret.z = lhs*rhs.z;
ret.w = lhs*rhs.w;
return ret;
}
__host__ __device__ cuvec4i operator/(const cuvec4i& lhs, const int& rhs)
{
cuvec4i ret;
ret.x = lhs.x/rhs;
ret.y = lhs.y/rhs;
ret.z = lhs.z/rhs;
ret.w = lhs.w/rhs;
return ret;
}
__host__ __device__ cuvec4i operator/(const int& lhs, const cuvec4i& rhs)
{
cuvec4i ret;
ret.x = lhs/rhs.x;
ret.y = lhs/rhs.y;
ret.z = lhs/rhs.z;
ret.w = lhs/rhs.w;
return ret;
}
__host__ __device__ cuvec4i operator-(const cuvec4i& other)
{
cuvec4i ret;
ret.x = -other.x;
ret.y = -other.y;
ret.z = -other.z;
ret.w = -other.w;
return ret;
}
__host__ __device__ cuvec4i& cuvec4i::operator+=(const cuvec4i& rhs)
{
x += rhs.x;
y += rhs.y;
z += rhs.z;
w += rhs.w;
return *this;
}
__host__ __device__ cuvec4i& cuvec4i::operator-=(const cuvec4i& rhs)
{
x -= rhs.x;
y -= rhs.y;
z -= rhs.z;
w -= rhs.w;
return *this;
}
__host__ __device__ cuvec4i& cuvec4i::operator*=(const int& rhs)
{
x *= rhs;
y *= rhs;
z *= rhs;
w *= rhs;
return *this;
}
__host__ __device__ cuvec4i& cuvec4i::operator/=(const int& rhs)
{
x /= rhs;
y /= rhs;
z /= rhs;
w /= rhs;
return *this;
}
//
//Matrix Header Stuff
//
__host__ __device__ cumat4i::cumat4i()
{
m00 = 0;
m01 = 0;
m02 = 0;
m03 = 0;
m10 = 0;
m11 = 0;
m12 = 0;
m13 = 0;
m20 = 0;
m21 = 0;
m22 = 0;
m23 = 0;
m30 = 0;
m31 = 0;
m32 = 0;
m33 = 0;
return;
}
__host__ __device__ cumat4i::~cumat4i()
{
//m00 = 0;
//m01 = 0;
//m02 = 0;
//m03 = 0;
//m10 = 0;
//m11 = 0;
//m12 = 0;
//m13 = 0;
//m20 = 0;
//m21 = 0;
//m22 = 0;
//m23 = 0;
//m30 = 0;
//m31 = 0;
//m32 = 0;
//m33 = 0;
return;
}
__host__ __device__ cumat4i::cumat4i(
const int& _m00, const int& _m01, const int& _m02, const int& _m03,
const int& _m10, const int& _m11, const int& _m12, const int& _m13,
const int& _m20, const int& _m21, const int& _m22, const int& _m23,
const int& _m30, const int& _m31, const int& _m32, const int& _m33
)
{
m00 = _m00;
m10 = _m10;
m20 = _m20;
m30 = _m30;
m01 = _m01;
m11 = _m11;
m21 = _m21;
m31 = _m31;
m02 = _m02;
m12 = _m12;
m22 = _m22;
m32 = _m32;
m03 = _m03;
m13 = _m13;
m23 = _m23;
m33 = _m33;
return;
}
__host__ __device__ cumat4i::cumat4i(const int* data16)
{
m00 = data16[0];
m10 = data16[1];
m20 = data16[2];
m30 = data16[3];
m01 = data16[4];
m11 = data16[5];
m21 = data16[6];
m31 = data16[7];
m02 = data16[8];
m12 = data16[9];
m22 = data16[10];
m32 = data16[11];
m03 = data16[12];
m13 = data16[13];
m23 = data16[14];
m33 = data16[15];
return;
}
__host__ __device__ float& cumat4i::operator[](const int &I)
{
switch(I)
{
case 0:
return m00;
case 1:
return m10;
case 2:
return m20;
case 3:
return m30;
case 4:
return m01;
case 5:
return m11;
case 6:
return m21;
case 7:
return m31;
case 8:
return m02;
case 9:
return m12;
case 10:
return m22;
case 11:
return m32;
case 12:
return m03;
case 13:
return m13;
case 14:
return m23;
case 15:
return m33;
}
return m00;
}
__host__ __device__ const float& cumat4i::operator[](const int &I) const
{
switch(I)
{
case 0:
return m00;
case 1:
return m10;
case 2:
return m20;
case 3:
return m30;
case 4:
return m01;
case 5:
return m11;
case 6:
return m21;
case 7:
return m31;
case 8:
return m02;
case 9:
return m12;
case 10:
return m22;
case 11:
return m32;
case 12:
return m03;
case 13:
return m13;
case 14:
return m23;
case 15:
return m33;
}
return m00;
}
__host__ __device__ float& cumat4i::operator()(const int &I, const int &J)
{
return (*this)[I+4*J];
}
__host__ __device__ const float& cumat4i::operator()(const int &I, const int &J) const
{
return (*this)[I+4*J];
}
__host__ __device__ float& cumat4i::at(const int &I, const int &J)
{
return (*this)[I+4*J];
}
__host__ __device__ const float& cumat4i::at(const int &I, const int &J) const
{
return (*this)[I+4*J];
}
__host__ __device__ int* cumat4i::data()
{
return (int*)this;
}
__host__ __device__ const int* cumat4i::data() const
{
return (int*)this;
}
__host__ __device__ cumat4i cumat4i::operator+(const cumat4i& rhs) const
{
cumat4i ret;
ret.m00 = m00 + rhs.m00;
ret.m10 = m10 + rhs.m10;
ret.m20 = m20 + rhs.m20;
ret.m30 = m30 + rhs.m30;
ret.m01 = m01 + rhs.m01;
ret.m11 = m11 + rhs.m11;
ret.m21 = m21 + rhs.m21;
ret.m31 = m31 + rhs.m31;
ret.m02 = m02 + rhs.m02;
ret.m12 = m12 + rhs.m12;
ret.m22 = m22 + rhs.m22;
ret.m32 = m32 + rhs.m32;
ret.m03 = m03 + rhs.m03;
ret.m13 = m13 + rhs.m13;
ret.m23 = m23 + rhs.m23;
ret.m33 = m33 + rhs.m33;
return ret;
}
__host__ __device__ cumat4i cumat4i::operator-(const cumat4i& rhs) const
{
cumat4i ret;
ret.m00 = m00 - rhs.m00;
ret.m10 = m10 - rhs.m10;
ret.m20 = m20 - rhs.m20;
ret.m30 = m30 - rhs.m30;
ret.m01 = m01 - rhs.m01;
ret.m11 = m11 - rhs.m11;
ret.m21 = m21 - rhs.m21;
ret.m31 = m31 - rhs.m31;
ret.m02 = m02 - rhs.m02;
ret.m12 = m12 - rhs.m12;
ret.m22 = m22 - rhs.m22;
ret.m32 = m32 - rhs.m32;
ret.m03 = m03 - rhs.m03;
ret.m13 = m13 - rhs.m13;
ret.m23 = m23 - rhs.m23;
ret.m33 = m33 - rhs.m33;
return ret;
}
__host__ __device__ cumat4i cumat4i::operator*(const cumat4i& rhs) const
{
cumat4i ret; //should be zeroed in constructor
ret.m00 = m00*rhs.m00 + m01*rhs.m10 + m02*rhs.m20 + m03*rhs.m30;
ret.m01 = m00*rhs.m01 + m01*rhs.m11 + m02*rhs.m21 + m03*rhs.m31;
ret.m02 = m00*rhs.m02 + m01*rhs.m12 + m02*rhs.m22 + m03*rhs.m32;
ret.m03 = m00*rhs.m03 + m01*rhs.m13 + m02*rhs.m23 + m03*rhs.m33;
ret.m10 = m10*rhs.m00 + m11*rhs.m10 + m12*rhs.m20 + m13*rhs.m30;
ret.m11 = m10*rhs.m01 + m11*rhs.m11 + m12*rhs.m21 + m13*rhs.m31;
ret.m12 = m10*rhs.m02 + m11*rhs.m12 + m12*rhs.m22 + m13*rhs.m32;
ret.m13 = m10*rhs.m03 + m11*rhs.m13 + m12*rhs.m23 + m13*rhs.m33;
ret.m20 = m20*rhs.m00 + m21*rhs.m10 + m22*rhs.m20 + m23*rhs.m30;
ret.m21 = m20*rhs.m01 + m21*rhs.m11 + m22*rhs.m21 + m23*rhs.m31;
ret.m22 = m20*rhs.m02 + m21*rhs.m12 + m22*rhs.m22 + m23*rhs.m32;
ret.m23 = m20*rhs.m03 + m21*rhs.m13 + m22*rhs.m23 + m23*rhs.m33;
ret.m30 = m30*rhs.m00 + m31*rhs.m10 + m32*rhs.m20 + m33*rhs.m30;
ret.m31 = m30*rhs.m01 + m31*rhs.m11 + m32*rhs.m21 + m33*rhs.m31;
ret.m32 = m30*rhs.m02 + m31*rhs.m12 + m32*rhs.m22 + m33*rhs.m32;
ret.m33 = m30*rhs.m03 + m31*rhs.m13 + m32*rhs.m23 + m33*rhs.m33;
return ret;
}
__host__ __device__ cumat4i operator*(const cumat4i& lhs, const int& rhs)
{
cumat4i ret;
ret.m00=lhs.m00*rhs;
ret.m10=lhs.m10*rhs;
ret.m20=lhs.m20*rhs;
ret.m30=lhs.m30*rhs;
ret.m01=lhs.m01*rhs;
ret.m11=lhs.m11*rhs;
ret.m21=lhs.m21*rhs;
ret.m31=lhs.m31*rhs;
ret.m02=lhs.m02*rhs;
ret.m12=lhs.m12*rhs;
ret.m22=lhs.m22*rhs;
ret.m32=lhs.m32*rhs;
ret.m03=lhs.m03*rhs;
ret.m13=lhs.m13*rhs;
ret.m23=lhs.m23*rhs;
ret.m33=lhs.m33*rhs;
return ret;
}
__host__ __device__ cumat4i operator/(const cumat4i& lhs, const int& rhs)
{
cumat4i ret;
ret.m00=lhs.m00/rhs;
ret.m10=lhs.m10/rhs;
ret.m20=lhs.m20/rhs;
ret.m30=lhs.m30/rhs;
ret.m01=lhs.m01/rhs;
ret.m11=lhs.m11/rhs;
ret.m21=lhs.m21/rhs;
ret.m31=lhs.m31/rhs;
ret.m02=lhs.m02/rhs;
ret.m12=lhs.m12/rhs;
ret.m22=lhs.m22/rhs;
ret.m32=lhs.m32/rhs;
ret.m03=lhs.m03/rhs;
ret.m13=lhs.m13/rhs;
ret.m23=lhs.m23/rhs;
ret.m33=lhs.m33/rhs;
return ret;
}
__host__ __device__ cumat4i operator*(const int& lhs, const cumat4i& rhs)
{
cumat4i ret;
ret.m00=lhs*rhs.m00;
ret.m10=lhs*rhs.m10;
ret.m20=lhs*rhs.m20;
ret.m30=lhs*rhs.m30;
ret.m01=lhs*rhs.m01;
ret.m11=lhs*rhs.m11;
ret.m21=lhs*rhs.m21;
ret.m31=lhs*rhs.m31;
ret.m02=lhs*rhs.m02;
ret.m12=lhs*rhs.m12;
ret.m22=lhs*rhs.m22;
ret.m32=lhs*rhs.m32;
ret.m03=lhs*rhs.m03;
ret.m13=lhs*rhs.m13;
ret.m23=lhs*rhs.m23;
ret.m33=lhs*rhs.m33;
return ret;
}
__host__ __device__ cuvec4i operator*(const cumat4i& lhs, const cuvec4i& rhs)
{
cuvec4i ret;
ret.x = lhs.m00*rhs.x + lhs.m01*rhs.y + lhs.m02*rhs.z + lhs.m03*rhs.w;
ret.y = lhs.m10*rhs.x + lhs.m11*rhs.y + lhs.m12*rhs.z + lhs.m13*rhs.w;
ret.z = lhs.m20*rhs.x + lhs.m21*rhs.y + lhs.m22*rhs.z + lhs.m23*rhs.w;
ret.w = lhs.m30*rhs.x + lhs.m31*rhs.y + lhs.m32*rhs.z + lhs.m33*rhs.w;
return ret;
}
__host__ __device__ cuvec4i operator*(const cuvec4i& lhs, const cumat4i& rhs)
{
cuvec4i ret;
ret.x = lhs.x*rhs.m00 + lhs.y*rhs.m10 + lhs.z*rhs.m20 + lhs.w*rhs.m30;
ret.y = lhs.x*rhs.m01 + lhs.y*rhs.m11 + lhs.z*rhs.m21 + lhs.w*rhs.m31;
ret.z = lhs.x*rhs.m02 + lhs.y*rhs.m12 + lhs.z*rhs.m22 + lhs.w*rhs.m32;
ret.w = lhs.x*rhs.m03 + lhs.y*rhs.m13 + lhs.z*rhs.m23 + lhs.w*rhs.m33;
return ret;
}
__host__ __device__ cumat4i operator-(const cumat4i& rhs)
{
cumat4i ret;
ret.m00 = -rhs.m00;
ret.m10 = -rhs.m10;
ret.m20 = -rhs.m20;
ret.m30 = -rhs.m30;
ret.m01 = -rhs.m01;
ret.m11 = -rhs.m11;
ret.m21 = -rhs.m21;
ret.m31 = -rhs.m31;
ret.m02 = -rhs.m02;
ret.m12 = -rhs.m12;
ret.m22 = -rhs.m22;
ret.m32 = -rhs.m32;
ret.m03 = -rhs.m03;
ret.m13 = -rhs.m13;
ret.m23 = -rhs.m23;
ret.m33 = -rhs.m33;
return ret;
}
__host__ __device__ cumat4i& cumat4i::operator+=(const cumat4i& rhs)
{
m00 += rhs.m00;
m10 += rhs.m10;
m20 += rhs.m20;
m30 += rhs.m30;
m01 += rhs.m01;
m11 += rhs.m11;
m21 += rhs.m21;
m31 += rhs.m31;
m02 += rhs.m02;
m12 += rhs.m12;
m22 += rhs.m22;
m32 += rhs.m32;
m03 += rhs.m03;
m13 += rhs.m13;
m23 += rhs.m23;
m33 += rhs.m33;
return *this;
}
__host__ __device__ cumat4i& cumat4i::operator-=(const cumat4i& rhs)
{
m00 -= rhs.m00;
m10 -= rhs.m10;
m20 -= rhs.m20;
m30 -= rhs.m30;
m01 -= rhs.m01;
m11 -= rhs.m11;
m21 -= rhs.m21;
m31 -= rhs.m31;
m02 -= rhs.m02;
m12 -= rhs.m12;
m22 -= rhs.m22;
m32 -= rhs.m32;
m03 -= rhs.m03;
m13 -= rhs.m13;
m23 -= rhs.m23;
m33 -= rhs.m33;
return *this;
}
__host__ __device__ cumat4i& cumat4i::operator*=(const int& rhs)
{
m00 *= rhs;
m10 *= rhs;
m20 *= rhs;
m30 *= rhs;
m01 *= rhs;
m11 *= rhs;
m21 *= rhs;
m31 *= rhs;
m02 *= rhs;
m12 *= rhs;
m22 *= rhs;
m32 *= rhs;
m03 *= rhs;
m13 *= rhs;
m23 *= rhs;
m33 *= rhs;
return *this;
}
__host__ __device__ cumat4i& cumat4i::operator/=(const int& rhs)
{
m00 /= rhs;
m10 /= rhs;
m20 /= rhs;
m30 /= rhs;
m01 /= rhs;
m11 /= rhs;
m21 /= rhs;
m31 /= rhs;
m02 /= rhs;
m12 /= rhs;
m22 /= rhs;
m32 /= rhs;
m03 /= rhs;
m13 /= rhs;
m23 /= rhs;
m33 /= rhs;
return *this;
}
__host__ __device__ cumat4i& cumat4i::operator*=(const cumat4i& rhs)
{
cumat4i tmp = *this;
m00 = tmp.m00*rhs.m00 + tmp.m01*rhs.m10 + tmp.m02*rhs.m20 + tmp.m03*rhs.m30;
m01 = tmp.m00*rhs.m01 + tmp.m01*rhs.m11 + tmp.m02*rhs.m21 + tmp.m03*rhs.m31;
m02 = tmp.m00*rhs.m02 + tmp.m01*rhs.m12 + tmp.m02*rhs.m22 + tmp.m03*rhs.m32;
m03 = tmp.m00*rhs.m03 + tmp.m01*rhs.m13 + tmp.m02*rhs.m23 + tmp.m03*rhs.m33;
m10 = tmp.m10*rhs.m00 + tmp.m11*rhs.m10 + tmp.m12*rhs.m20 + tmp.m13*rhs.m30;
m11 = tmp.m10*rhs.m01 + tmp.m11*rhs.m11 + tmp.m12*rhs.m21 + tmp.m13*rhs.m31;
m12 = tmp.m10*rhs.m02 + tmp.m11*rhs.m12 + tmp.m12*rhs.m22 + tmp.m13*rhs.m32;
m13 = tmp.m10*rhs.m03 + tmp.m11*rhs.m13 + tmp.m12*rhs.m23 + tmp.m13*rhs.m33;
m20 = tmp.m20*rhs.m00 + tmp.m21*rhs.m10 + tmp.m22*rhs.m20 + tmp.m23*rhs.m30;
m21 = tmp.m20*rhs.m01 + tmp.m21*rhs.m11 + tmp.m22*rhs.m21 + tmp.m23*rhs.m31;
m22 = tmp.m20*rhs.m02 + tmp.m21*rhs.m12 + tmp.m22*rhs.m22 + tmp.m23*rhs.m32;
m23 = tmp.m20*rhs.m03 + tmp.m21*rhs.m13 + tmp.m22*rhs.m23 + tmp.m23*rhs.m33;
m30 = tmp.m30*rhs.m00 + tmp.m31*rhs.m10 + tmp.m32*rhs.m20 + tmp.m33*rhs.m30;
m31 = tmp.m30*rhs.m01 + tmp.m31*rhs.m11 + tmp.m32*rhs.m21 + tmp.m33*rhs.m31;
m32 = tmp.m30*rhs.m02 + tmp.m31*rhs.m12 + tmp.m32*rhs.m22 + tmp.m33*rhs.m32;
m33 = tmp.m30*rhs.m03 + tmp.m31*rhs.m13 + tmp.m32*rhs.m23 + tmp.m33*rhs.m33;
return *this;
}
__host__ __device__ cumat4i cumat4i::transpose() const
{
cumat4i ret;
ret.m00 = m00;
ret.m10 = m01;
ret.m20 = m02;
ret.m30 = m03;
ret.m01 = m10;
ret.m11 = m11;
ret.m21 = m12;
ret.m31 = m13;
ret.m02 = m20;
ret.m12 = m21;
ret.m22 = m22;
ret.m32 = m23;
ret.m03 = m30;
ret.m13 = m31;
ret.m23 = m32;
ret.m33 = m33;
return ret;
}

58
test_scripts/cuvec_codegen/cuvec4i_codegen1.hpp Normal file
View File

@ -0,0 +1,58 @@
__host__ __device__ cuvec4i();
__host__ __device__ ~cuvec4i();
__host__ __device__ cuvec4i(const int &_x, const int &_y, const int &_z, const int &_w);
__host__ __device__ float& operator[](const int &I);
__host__ __device__ const float& operator[](const int &I) const;
__host__ __device__ cuvec4i operator+(const cuvec4i& rhs) const;
__host__ __device__ cuvec4i operator-(const cuvec4i& rhs) const;
__host__ __device__ cuvec4i operator*(const cuvec4i& rhs) const; //elementwise product
__host__ __device__ cuvec4i operator/(const cuvec4i& rhs) const; //elementwise division
__host__ __device__ friend cuvec4i operator*(const cuvec4i& lhs, const int& rhs);
__host__ __device__ friend cuvec4i operator*(const int& lhs, const cuvec4i& rhs);
__host__ __device__ friend cuvec4i operator/(const cuvec4i& lhs, const int& rhs);
__host__ __device__ friend cuvec4i operator/(const int& lhs, const cuvec4i& rhs);
__host__ __device__ friend cuvec4i operator-(const cuvec4i& other);
__host__ __device__ cuvec4i& operator+=(const cuvec4i& rhs);
__host__ __device__ cuvec4i& operator-=(const cuvec4i& rhs);
__host__ __device__ cuvec4i& operator*=(const int& rhs);
__host__ __device__ cuvec4i& operator/=(const int& rhs);
//
//Matrix Header Stuff
//
int m00,m10,m20,m30;
int m01,m11,m21,m31;
int m02,m12,m22,m32;
int m03,m13,m23,m33;
__host__ __device__ cumat4i();
__host__ __device__ ~cumat4i();
__host__ __device__ cumat4i(
const int& _m00, const int& _m01, const int& _m02, const int& _m03,
const int& _m10, const int& _m11, const int& _m12, const int& _m13,
const int& _m20, const int& _m21, const int& _m22, const int& _m23,
const int& _m30, const int& _m31, const int& _m32, const int& _m33
);
__host__ __device__ cumat4i(const int* data16);
__host__ __device__ float& operator[](const int &I);
__host__ __device__ const float& operator[](const int &I) const;
__host__ __device__ float& operator()(const int &I, const int &J);
__host__ __device__ const float& operator()(const int &I, const int &J) const;
__host__ __device__ float& at(const int &I, const int &J);
__host__ __device__ const float& at(const int &I, const int &J) const;
__host__ __device__ int* data(); //pointer to float16 representation of matrix
__host__ __device__ const int* data() const; //pointer to float16 representation of matrix
__host__ __device__ cumat4i operator+(const cumat4i& rhs) const;
__host__ __device__ cumat4i operator-(const cumat4i& rhs) const;
__host__ __device__ cumat4i operator*(const cumat4i& rhs) const;
__host__ __device__ friend cumat4i operator*(const cumat4i& lhs, const int& rhs);
__host__ __device__ friend cumat4i operator/(const cumat4i& lhs, const int& rhs);
__host__ __device__ friend cumat4i operator*(const int& lhs, const cumat4i& rhs);
__host__ __device__ friend cuvec4i operator*(const cumat4i& lhs, const cuvec4i& rhs);
__host__ __device__ friend cuvec4i operator*(const cuvec4i& lhs, const cumat4i& rhs);
__host__ __device__ friend cumat4i operator-(const cumat4i& rhs);
__host__ __device__ cumat4i& operator+=(const cumat4i& rhs);
__host__ __device__ cumat4i& operator-=(const cumat4i& rhs);
__host__ __device__ cumat4i& operator*=(const int& rhs);
__host__ __device__ cumat4i& operator/=(const int& rhs);
__host__ __device__ cumat4i& operator*=(const cumat4i& rhs);
__host__ __device__ cumat4i transpose() const;