diff --git a/build_linux64/libamsculib3.linux64.a b/build_linux64/libamsculib3.linux64.a index 1b5956f..534f417 100644 Binary files a/build_linux64/libamsculib3.linux64.a and b/build_linux64/libamsculib3.linux64.a differ diff --git a/build_linux64/objstore/cuvec2i.o b/build_linux64/objstore/cuvec2i.o new file mode 100644 index 0000000..89cffc5 Binary files /dev/null and b/build_linux64/objstore/cuvec2i.o differ diff --git a/build_linux64/objstore/cuvec3i.o b/build_linux64/objstore/cuvec3i.o new file mode 100644 index 0000000..6913e2b Binary files /dev/null and b/build_linux64/objstore/cuvec3i.o differ diff --git a/build_linux64/objstore/cuvec4i.o b/build_linux64/objstore/cuvec4i.o new file mode 100644 index 0000000..56697de Binary files /dev/null and b/build_linux64/objstore/cuvec4i.o differ diff --git a/include/amsculib3/math/cuvec2i.hpp b/include/amsculib3/math/cuvec2i.hpp index 172d7e5..a1a7d80 100644 --- a/include/amsculib3/math/cuvec2i.hpp +++ b/include/amsculib3/math/cuvec2i.hpp @@ -4,6 +4,32 @@ namespace amscuda { + class cuvec2i + { + public: + int x; + int y; + + __host__ __device__ cuvec2i(); + __host__ __device__ ~cuvec2i(); + __host__ __device__ cuvec2i(const int &_x, const int &_y); + __host__ __device__ int& operator[](const int &I); + __host__ __device__ const int& 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); + + }; }; //end namespace amscuda diff --git a/include/amsculib3/math/cuvec3i.hpp b/include/amsculib3/math/cuvec3i.hpp index 90f6c1c..e98a341 100644 --- a/include/amsculib3/math/cuvec3i.hpp +++ b/include/amsculib3/math/cuvec3i.hpp @@ -4,6 +4,33 @@ namespace amscuda { + class cuvec3i + { + public: + int x; + int y; + int z; + + __host__ __device__ cuvec3i(); + __host__ __device__ ~cuvec3i(); + __host__ __device__ cuvec3i(const int &_x, const int &_y, const int &_z); + __host__ __device__ int& operator[](const int &I); + __host__ __device__ const int& 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); + + }; }; //end namespace amscuda diff --git a/include/amsculib3/math/cuvec4i.hpp b/include/amsculib3/math/cuvec4i.hpp index ae44b21..658ebd8 100644 --- a/include/amsculib3/math/cuvec4i.hpp +++ b/include/amsculib3/math/cuvec4i.hpp @@ -4,6 +4,34 @@ namespace amscuda { + class cuvec4i + { + public: + int x; + int y; + int z; + int w; + + __host__ __device__ cuvec4i(); + __host__ __device__ ~cuvec4i(); + __host__ __device__ cuvec4i(const int &_x, const int &_y, const int &_z, const int &_w); + __host__ __device__ int& operator[](const int &I); + __host__ __device__ const int& 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); + + }; }; //end namespace amscuda diff --git a/src/amsculib3/math/cuvec2i.cu b/src/amsculib3/math/cuvec2i.cu new file mode 100644 index 0000000..829f578 --- /dev/null +++ b/src/amsculib3/math/cuvec2i.cu @@ -0,0 +1,167 @@ +#include + +namespace amscuda +{ + +////////////////// +// Vector Class // +////////////////// + +__host__ __device__ cuvec2i::cuvec2i() +{ + x = 0; y = 0; + return; +} + +__host__ __device__ cuvec2i::~cuvec2i() +{ + x = 0; y = 0; + return; +} + +__host__ __device__ cuvec2i::cuvec2i(const int32_t &_x, const int32_t &_y) +{ + x = _x; y = _y; + return; +} + +__host__ __device__ int32_t& cuvec2i::operator[](const int &I) +{ + switch(I) + { + case 0: + return x; + case 1: + return y; + } + + return x; +} + +__host__ __device__ const int32_t& 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 int32_t& rhs) +{ + cuvec2i ret; + ret.x = lhs.x*rhs; + ret.y = lhs.y*rhs; + return ret; +} + +__host__ __device__ cuvec2i operator*(const int32_t& 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 int32_t& rhs) +{ + cuvec2i ret; + ret.x = lhs.x/rhs; + ret.y = lhs.y/rhs; + return ret; +} + +__host__ __device__ cuvec2i operator/(const int32_t& 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 int32_t& rhs) +{ + x *= rhs; + y *= rhs; + return *this; +} + +__host__ __device__ cuvec2i& cuvec2i::operator/=(const int32_t& rhs) +{ + x /= rhs; + y /= rhs; + return *this; +} + + + +////////////////////// +// Standalone Funcs // +////////////////////// + + +/////////// +// Tests // +/////////// + +}; \ No newline at end of file diff --git a/src/amsculib3/math/cuvec3i.cu b/src/amsculib3/math/cuvec3i.cu new file mode 100644 index 0000000..6e2d7e7 --- /dev/null +++ b/src/amsculib3/math/cuvec3i.cu @@ -0,0 +1,183 @@ +#include + +namespace amscuda +{ + +////////////////// +// Vector Class // +////////////////// + +__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__ int& cuvec3i::operator[](const int &I) +{ + switch(I) + { + case 0: + return x; + case 1: + return y; + case 2: + return z; + } + + return x; +} + +__host__ __device__ const int& 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; +} + + +////////////////////// +// Standalone Funcs // +////////////////////// + + +/////////// +// Tests // +/////////// + +}; \ No newline at end of file diff --git a/src/amsculib3/math/cuvec4i.cu b/src/amsculib3/math/cuvec4i.cu new file mode 100644 index 0000000..7de0e66 --- /dev/null +++ b/src/amsculib3/math/cuvec4i.cu @@ -0,0 +1,200 @@ +#include + +namespace amscuda +{ + +////////////////// +// Vector Class // +////////////////// + +__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__ int& 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 int& 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; +} + + +////////////////////// +// Standalone Funcs // +////////////////////// + + +/////////// +// Tests // +/////////// + +}; \ No newline at end of file diff --git a/test_scripts/cuvec_codegen.py b/test_scripts/cuvec_codegen.py index 134d892..23efc97 100644 --- a/test_scripts/cuvec_codegen.py +++ b/test_scripts/cuvec_codegen.py @@ -155,9 +155,9 @@ def v_operatorbk1(dim,dtype): lnss = "" vname = vtypename(dim,dtype) - lnsh += "\t{} float& operator[](const int &I);\n".format(cudadec) + lnsh += f"\t{cudadec} {dtype}& operator[](const int &I);\n" - lnss += "\t{} float& {}::operator[](const int &I)\n".format(cudadec,vname) + lnss += f"\t{cudadec} {dtype}& {vname}::operator[](const int &I)\n" lnss += "\t{\n" lnss += "\t\tswitch(I)\n\t\t{\n" for I in range(0,dim): @@ -174,9 +174,9 @@ def v_operatorbk2(dim,dtype): lnss = "" vname = vtypename(dim,dtype) - lnsh += "\t{} const float& operator[](const int &I) const;\n".format(cudadec) + lnsh += f"\t{cudadec} const {dtype}& operator[](const int &I) const;\n" - lnss += "\t{} const float& {}::operator[](const int &I) const\n".format(cudadec,vname) + lnss += f"\t{cudadec} const {dtype}& {vname}::operator[](const int &I) const\n" lnss += "\t{\n" lnss += "\t\tswitch(I)\n\t\t{\n" for I in range(0,dim): @@ -580,9 +580,9 @@ def moperatorbk1(dim,dtype): lnss = "" mname = mtypename(dim,dtype) - lnsh += f"\t{cudadec} float& operator[](const int &I);\n" + lnsh += f"\t{cudadec} {dtype}& operator[](const int &I);\n" - lnss += f"\t{cudadec} float& {mname}::operator[](const int &I)\n" + lnss += f"\t{cudadec} {dtype}& {mname}::operator[](const int &I)\n" lnss += "\t{\n" lnss += "\t\tswitch(I)\n\t\t{\n" for I in range(0,dim*dim): @@ -601,9 +601,9 @@ def moperatorbk2(dim,dtype): lnss = "" mname = mtypename(dim,dtype) - lnsh += f"\t{cudadec} const float& operator[](const int &I) const;\n" + lnsh += f"\t{cudadec} const {dtype}& operator[](const int &I) const;\n" - lnss += f"\t{cudadec} const float& {mname}::operator[](const int &I) const\n" + lnss += f"\t{cudadec} const {dtype}& {mname}::operator[](const int &I) const\n" lnss += "\t{\n" lnss += "\t\tswitch(I)\n\t\t{\n" for I in range(0,dim*dim): @@ -622,9 +622,9 @@ def moperatorbk3(dim,dtype): lnss = "" mname = mtypename(dim,dtype) - lnsh += f"\t{cudadec} float& operator()(const int &I, const int &J);\n" + lnsh += f"\t{cudadec} {dtype}& operator()(const int &I, const int &J);\n" - lnss += f"\t{cudadec} float& {mname}::operator()(const int &I, const int &J)\n" + lnss += f"\t{cudadec} {dtype}& {mname}::operator()(const int &I, const int &J)\n" lnss += "\t{\n" lnss += f"\t\treturn (*this)[I+{dim}*J];\n" lnss += "\t}\n\n" @@ -636,9 +636,9 @@ def moperatorbk4(dim,dtype): lnss = "" mname = mtypename(dim,dtype) - lnsh += f"\t{cudadec} const float& operator()(const int &I, const int &J) const;\n" + lnsh += f"\t{cudadec} const {dtype}& operator()(const int &I, const int &J) const;\n" - lnss += f"\t{cudadec} const float& {mname}::operator()(const int &I, const int &J) const\n" + lnss += f"\t{cudadec} const {dtype}& {mname}::operator()(const int &I, const int &J) const\n" lnss += "\t{\n" lnss += f"\t\treturn (*this)[I+{dim}*J];\n" lnss += "\t}\n\n" @@ -650,9 +650,9 @@ def moperatorbk5(dim,dtype): lnss = "" mname = mtypename(dim,dtype) - lnsh += f"\t{cudadec} float& at(const int &I, const int &J);\n" + lnsh += f"\t{cudadec} {dtype}& at(const int &I, const int &J);\n" - lnss += f"\t{cudadec} float& {mname}::at(const int &I, const int &J)\n" + lnss += f"\t{cudadec} {dtype}& {mname}::at(const int &I, const int &J)\n" lnss += "\t{\n" lnss += f"\t\treturn (*this)[I+{dim}*J];\n" lnss += "\t}\n\n" @@ -664,9 +664,9 @@ def moperatorbk6(dim,dtype): lnss = "" mname = mtypename(dim,dtype) - lnsh += f"\t{cudadec} const float& at(const int &I, const int &J) const;\n" + lnsh += f"\t{cudadec} const {dtype}& at(const int &I, const int &J) const;\n" - lnss += f"\t{cudadec} const float& {mname}::at(const int &I, const int &J) const\n" + lnss += f"\t{cudadec} const {dtype}& {mname}::at(const int &I, const int &J) const\n" lnss += "\t{\n" lnss += f"\t\treturn (*this)[I+{dim}*J];\n" lnss += "\t}\n\n" @@ -680,7 +680,7 @@ def mdataoperator1(dim,dtype): lnss = "" mname = mtypename(dim,dtype) - lnsh += f"{cudadec} {dtype}* data(); //pointer to float{dim*dim} representation of matrix\n" + lnsh += f"{cudadec} {dtype}* data(); //pointer to {dtype}{dim*dim} representation of matrix\n" lnss += f"{cudadec} {dtype}* {mname}::data()\n" lnss += "{\n" @@ -694,7 +694,7 @@ def mdataoperator2(dim,dtype): lnss = "" mname = mtypename(dim,dtype) - lnsh += f"{cudadec} const {dtype}* data() const; //pointer to float{dim*dim} representation of matrix\n" + lnsh += f"{cudadec} const {dtype}* data() const; //pointer to {dtype}{dim*dim} representation of matrix\n" lnss += f"{cudadec} const {dtype}* {mname}::data() const\n" lnss += "{\n" diff --git a/test_scripts/cuvec_codegen/cuvec2_codegen1.cu b/test_scripts/cuvec_codegen/cuvec2_codegen1.cu index cf7349f..e7f8dda 100644 --- a/test_scripts/cuvec_codegen/cuvec2_codegen1.cu +++ b/test_scripts/cuvec_codegen/cuvec2_codegen1.cu @@ -16,7 +16,7 @@ return; } - __host__ __device__ float& cuvec2::operator[](const int &I) + __host__ __device__ double& cuvec2::operator[](const int &I) { switch(I) { @@ -29,7 +29,7 @@ return x; } - __host__ __device__ const float& cuvec2::operator[](const int &I) const + __host__ __device__ const double& cuvec2::operator[](const int &I) const { switch(I) { @@ -198,7 +198,7 @@ return; } - __host__ __device__ float& cumat2::operator[](const int &I) + __host__ __device__ double& cumat2::operator[](const int &I) { switch(I) { @@ -215,7 +215,7 @@ return m00; } - __host__ __device__ const float& cumat2::operator[](const int &I) const + __host__ __device__ const double& cumat2::operator[](const int &I) const { switch(I) { @@ -232,22 +232,22 @@ return m00; } - __host__ __device__ float& cumat2::operator()(const int &I, const int &J) + __host__ __device__ double& 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 + __host__ __device__ const double& 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) + __host__ __device__ double& 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 + __host__ __device__ const double& cumat2::at(const int &I, const int &J) const { return (*this)[I+2*J]; } diff --git a/test_scripts/cuvec_codegen/cuvec2_codegen1.hpp b/test_scripts/cuvec_codegen/cuvec2_codegen1.hpp index 74f2e44..4cb56e4 100644 --- a/test_scripts/cuvec_codegen/cuvec2_codegen1.hpp +++ b/test_scripts/cuvec_codegen/cuvec2_codegen1.hpp @@ -1,8 +1,8 @@ __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__ double& operator[](const int &I); + __host__ __device__ const double& 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 @@ -29,14 +29,14 @@ double m01,m11; 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__ double& operator[](const int &I); + __host__ __device__ const double& operator[](const int &I) const; + __host__ __device__ double& operator()(const int &I, const int &J); + __host__ __device__ const double& operator()(const int &I, const int &J) const; + __host__ __device__ double& at(const int &I, const int &J); + __host__ __device__ const double& at(const int &I, const int &J) const; +__host__ __device__ double* data(); //pointer to double4 representation of matrix +__host__ __device__ const double* data() const; //pointer to double4 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; diff --git a/test_scripts/cuvec_codegen/cuvec2i_codegen1.cu b/test_scripts/cuvec_codegen/cuvec2i_codegen1.cu index 57ea2c3..ad58ebe 100644 --- a/test_scripts/cuvec_codegen/cuvec2i_codegen1.cu +++ b/test_scripts/cuvec_codegen/cuvec2i_codegen1.cu @@ -16,7 +16,7 @@ return; } - __host__ __device__ float& cuvec2i::operator[](const int &I) + __host__ __device__ int& cuvec2i::operator[](const int &I) { switch(I) { @@ -29,7 +29,7 @@ return x; } - __host__ __device__ const float& cuvec2i::operator[](const int &I) const + __host__ __device__ const int& cuvec2i::operator[](const int &I) const { switch(I) { @@ -198,7 +198,7 @@ return; } - __host__ __device__ float& cumat2i::operator[](const int &I) + __host__ __device__ int& cumat2i::operator[](const int &I) { switch(I) { @@ -215,7 +215,7 @@ return m00; } - __host__ __device__ const float& cumat2i::operator[](const int &I) const + __host__ __device__ const int& cumat2i::operator[](const int &I) const { switch(I) { @@ -232,22 +232,22 @@ return m00; } - __host__ __device__ float& cumat2i::operator()(const int &I, const int &J) + __host__ __device__ int& 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 + __host__ __device__ const int& 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) + __host__ __device__ int& 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 + __host__ __device__ const int& cumat2i::at(const int &I, const int &J) const { return (*this)[I+2*J]; } diff --git a/test_scripts/cuvec_codegen/cuvec2i_codegen1.hpp b/test_scripts/cuvec_codegen/cuvec2i_codegen1.hpp index 1ba6b12..1b84e78 100644 --- a/test_scripts/cuvec_codegen/cuvec2i_codegen1.hpp +++ b/test_scripts/cuvec_codegen/cuvec2i_codegen1.hpp @@ -1,8 +1,8 @@ __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__ int& operator[](const int &I); + __host__ __device__ const int& 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 @@ -29,14 +29,14 @@ int m01,m11; 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__ int& operator[](const int &I); + __host__ __device__ const int& operator[](const int &I) const; + __host__ __device__ int& operator()(const int &I, const int &J); + __host__ __device__ const int& operator()(const int &I, const int &J) const; + __host__ __device__ int& at(const int &I, const int &J); + __host__ __device__ const int& at(const int &I, const int &J) const; +__host__ __device__ int* data(); //pointer to int4 representation of matrix +__host__ __device__ const int* data() const; //pointer to int4 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; diff --git a/test_scripts/cuvec_codegen/cuvec3_codegen1.cu b/test_scripts/cuvec_codegen/cuvec3_codegen1.cu index 6b2faff..d9e5711 100644 --- a/test_scripts/cuvec_codegen/cuvec3_codegen1.cu +++ b/test_scripts/cuvec_codegen/cuvec3_codegen1.cu @@ -16,7 +16,7 @@ return; } - __host__ __device__ float& cuvec3::operator[](const int &I) + __host__ __device__ double& cuvec3::operator[](const int &I) { switch(I) { @@ -31,7 +31,7 @@ return x; } - __host__ __device__ const float& cuvec3::operator[](const int &I) const + __host__ __device__ const double& cuvec3::operator[](const int &I) const { switch(I) { @@ -240,7 +240,7 @@ return; } - __host__ __device__ float& cumat3::operator[](const int &I) + __host__ __device__ double& cumat3::operator[](const int &I) { switch(I) { @@ -267,7 +267,7 @@ return m00; } - __host__ __device__ const float& cumat3::operator[](const int &I) const + __host__ __device__ const double& cumat3::operator[](const int &I) const { switch(I) { @@ -294,22 +294,22 @@ return m00; } - __host__ __device__ float& cumat3::operator()(const int &I, const int &J) + __host__ __device__ double& 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 + __host__ __device__ const double& 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) + __host__ __device__ double& 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 + __host__ __device__ const double& cumat3::at(const int &I, const int &J) const { return (*this)[I+3*J]; } diff --git a/test_scripts/cuvec_codegen/cuvec3_codegen1.hpp b/test_scripts/cuvec_codegen/cuvec3_codegen1.hpp index d83bd0e..420855e 100644 --- a/test_scripts/cuvec_codegen/cuvec3_codegen1.hpp +++ b/test_scripts/cuvec_codegen/cuvec3_codegen1.hpp @@ -1,8 +1,8 @@ __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__ double& operator[](const int &I); + __host__ __device__ const double& 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 @@ -31,14 +31,14 @@ double m02,m12,m22; 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__ double& operator[](const int &I); + __host__ __device__ const double& operator[](const int &I) const; + __host__ __device__ double& operator()(const int &I, const int &J); + __host__ __device__ const double& operator()(const int &I, const int &J) const; + __host__ __device__ double& at(const int &I, const int &J); + __host__ __device__ const double& at(const int &I, const int &J) const; +__host__ __device__ double* data(); //pointer to double9 representation of matrix +__host__ __device__ const double* data() const; //pointer to double9 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; diff --git a/test_scripts/cuvec_codegen/cuvec3i_codegen1.cu b/test_scripts/cuvec_codegen/cuvec3i_codegen1.cu index 8417071..48b11c4 100644 --- a/test_scripts/cuvec_codegen/cuvec3i_codegen1.cu +++ b/test_scripts/cuvec_codegen/cuvec3i_codegen1.cu @@ -16,7 +16,7 @@ return; } - __host__ __device__ float& cuvec3i::operator[](const int &I) + __host__ __device__ int& cuvec3i::operator[](const int &I) { switch(I) { @@ -31,7 +31,7 @@ return x; } - __host__ __device__ const float& cuvec3i::operator[](const int &I) const + __host__ __device__ const int& cuvec3i::operator[](const int &I) const { switch(I) { @@ -240,7 +240,7 @@ return; } - __host__ __device__ float& cumat3i::operator[](const int &I) + __host__ __device__ int& cumat3i::operator[](const int &I) { switch(I) { @@ -267,7 +267,7 @@ return m00; } - __host__ __device__ const float& cumat3i::operator[](const int &I) const + __host__ __device__ const int& cumat3i::operator[](const int &I) const { switch(I) { @@ -294,22 +294,22 @@ return m00; } - __host__ __device__ float& cumat3i::operator()(const int &I, const int &J) + __host__ __device__ int& 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 + __host__ __device__ const int& 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) + __host__ __device__ int& 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 + __host__ __device__ const int& cumat3i::at(const int &I, const int &J) const { return (*this)[I+3*J]; } diff --git a/test_scripts/cuvec_codegen/cuvec3i_codegen1.hpp b/test_scripts/cuvec_codegen/cuvec3i_codegen1.hpp index 53a1936..5e6db5a 100644 --- a/test_scripts/cuvec_codegen/cuvec3i_codegen1.hpp +++ b/test_scripts/cuvec_codegen/cuvec3i_codegen1.hpp @@ -1,8 +1,8 @@ __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__ int& operator[](const int &I); + __host__ __device__ const int& 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 @@ -31,14 +31,14 @@ int m02,m12,m22; 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__ int& operator[](const int &I); + __host__ __device__ const int& operator[](const int &I) const; + __host__ __device__ int& operator()(const int &I, const int &J); + __host__ __device__ const int& operator()(const int &I, const int &J) const; + __host__ __device__ int& at(const int &I, const int &J); + __host__ __device__ const int& at(const int &I, const int &J) const; +__host__ __device__ int* data(); //pointer to int9 representation of matrix +__host__ __device__ const int* data() const; //pointer to int9 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; diff --git a/test_scripts/cuvec_codegen/cuvec4_codegen1.cu b/test_scripts/cuvec_codegen/cuvec4_codegen1.cu index 7884c7e..ed3f59a 100644 --- a/test_scripts/cuvec_codegen/cuvec4_codegen1.cu +++ b/test_scripts/cuvec_codegen/cuvec4_codegen1.cu @@ -16,7 +16,7 @@ return; } - __host__ __device__ float& cuvec4::operator[](const int &I) + __host__ __device__ double& cuvec4::operator[](const int &I) { switch(I) { @@ -33,7 +33,7 @@ return x; } - __host__ __device__ const float& cuvec4::operator[](const int &I) const + __host__ __device__ const double& cuvec4::operator[](const int &I) const { switch(I) { @@ -290,7 +290,7 @@ return; } - __host__ __device__ float& cumat4::operator[](const int &I) + __host__ __device__ double& cumat4::operator[](const int &I) { switch(I) { @@ -331,7 +331,7 @@ return m00; } - __host__ __device__ const float& cumat4::operator[](const int &I) const + __host__ __device__ const double& cumat4::operator[](const int &I) const { switch(I) { @@ -372,22 +372,22 @@ return m00; } - __host__ __device__ float& cumat4::operator()(const int &I, const int &J) + __host__ __device__ double& 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 + __host__ __device__ const double& 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) + __host__ __device__ double& 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 + __host__ __device__ const double& cumat4::at(const int &I, const int &J) const { return (*this)[I+4*J]; } diff --git a/test_scripts/cuvec_codegen/cuvec4_codegen1.hpp b/test_scripts/cuvec_codegen/cuvec4_codegen1.hpp index 1c61be8..b94c463 100644 --- a/test_scripts/cuvec_codegen/cuvec4_codegen1.hpp +++ b/test_scripts/cuvec_codegen/cuvec4_codegen1.hpp @@ -1,8 +1,8 @@ __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__ double& operator[](const int &I); + __host__ __device__ const double& 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 @@ -33,14 +33,14 @@ double m03,m13,m23,m33; 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__ double& operator[](const int &I); + __host__ __device__ const double& operator[](const int &I) const; + __host__ __device__ double& operator()(const int &I, const int &J); + __host__ __device__ const double& operator()(const int &I, const int &J) const; + __host__ __device__ double& at(const int &I, const int &J); + __host__ __device__ const double& at(const int &I, const int &J) const; +__host__ __device__ double* data(); //pointer to double16 representation of matrix +__host__ __device__ const double* data() const; //pointer to double16 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; diff --git a/test_scripts/cuvec_codegen/cuvec4i_codegen1.cu b/test_scripts/cuvec_codegen/cuvec4i_codegen1.cu index 3d38904..736d65d 100644 --- a/test_scripts/cuvec_codegen/cuvec4i_codegen1.cu +++ b/test_scripts/cuvec_codegen/cuvec4i_codegen1.cu @@ -16,7 +16,7 @@ return; } - __host__ __device__ float& cuvec4i::operator[](const int &I) + __host__ __device__ int& cuvec4i::operator[](const int &I) { switch(I) { @@ -33,7 +33,7 @@ return x; } - __host__ __device__ const float& cuvec4i::operator[](const int &I) const + __host__ __device__ const int& cuvec4i::operator[](const int &I) const { switch(I) { @@ -290,7 +290,7 @@ return; } - __host__ __device__ float& cumat4i::operator[](const int &I) + __host__ __device__ int& cumat4i::operator[](const int &I) { switch(I) { @@ -331,7 +331,7 @@ return m00; } - __host__ __device__ const float& cumat4i::operator[](const int &I) const + __host__ __device__ const int& cumat4i::operator[](const int &I) const { switch(I) { @@ -372,22 +372,22 @@ return m00; } - __host__ __device__ float& cumat4i::operator()(const int &I, const int &J) + __host__ __device__ int& 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 + __host__ __device__ const int& 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) + __host__ __device__ int& 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 + __host__ __device__ const int& cumat4i::at(const int &I, const int &J) const { return (*this)[I+4*J]; } diff --git a/test_scripts/cuvec_codegen/cuvec4i_codegen1.hpp b/test_scripts/cuvec_codegen/cuvec4i_codegen1.hpp index c66849e..91f183b 100644 --- a/test_scripts/cuvec_codegen/cuvec4i_codegen1.hpp +++ b/test_scripts/cuvec_codegen/cuvec4i_codegen1.hpp @@ -1,8 +1,8 @@ __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__ int& operator[](const int &I); + __host__ __device__ const int& 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 @@ -33,14 +33,14 @@ int m03,m13,m23,m33; 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__ int& operator[](const int &I); + __host__ __device__ const int& operator[](const int &I) const; + __host__ __device__ int& operator()(const int &I, const int &J); + __host__ __device__ const int& operator()(const int &I, const int &J) const; + __host__ __device__ int& at(const int &I, const int &J); + __host__ __device__ const int& at(const int &I, const int &J) const; +__host__ __device__ int* data(); //pointer to int16 representation of matrix +__host__ __device__ const int* data() const; //pointer to int16 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;