integer vector types

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
 

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
 

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
 

src/amsculib3/math/cuvec2i.cu

@@ -0,0 +1,167 @@
+#include <amsculib3/amsculib3.hpp>
+
+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 //
+///////////
+
+};

src/amsculib3/math/cuvec3i.cu

@@ -0,0 +1,183 @@
+#include <amsculib3/amsculib3.hpp>
+
+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 //
+///////////
+
+};

src/amsculib3/math/cuvec4i.cu

@@ -0,0 +1,200 @@
+#include <amsculib3/amsculib3.hpp>
+
+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 //
+///////////
+
+};

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"

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];
 	}

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__ double& operator[](const int &I);
-	__host__ __device__ const float& operator[](const int &I) const;
+	__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__ double& operator[](const int &I);
-	__host__ __device__ const float& operator[](const int &I) const;
+	__host__ __device__ const double& operator[](const int &I) const;
-	__host__ __device__ float& operator()(const int &I, const int &J);
+	__host__ __device__ double& operator()(const int &I, const int &J);
-	__host__ __device__ const float& operator()(const int &I, const int &J) const;
+	__host__ __device__ const double& operator()(const int &I, const int &J) const;
-	__host__ __device__ float& at(const int &I, const int &J);
+	__host__ __device__ double& at(const int &I, const int &J);
-	__host__ __device__ const float& at(const int &I, const int &J) const;
+	__host__ __device__ const double& at(const int &I, const int &J) const;
-__host__ __device__ double* data(); //pointer to float4 representation of matrix
+__host__ __device__ double* data(); //pointer to double4 representation of matrix
-__host__ __device__ const double* data() const; //pointer to float4 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;

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];
 	}

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__ int& operator[](const int &I);
-	__host__ __device__ const float& operator[](const int &I) const;
+	__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__ int& operator[](const int &I);
-	__host__ __device__ const float& operator[](const int &I) const;
+	__host__ __device__ const int& operator[](const int &I) const;
-	__host__ __device__ float& operator()(const int &I, const int &J);
+	__host__ __device__ int& operator()(const int &I, const int &J);
-	__host__ __device__ const float& operator()(const int &I, const int &J) const;
+	__host__ __device__ const int& operator()(const int &I, const int &J) const;
-	__host__ __device__ float& at(const int &I, const int &J);
+	__host__ __device__ int& at(const int &I, const int &J);
-	__host__ __device__ const float& at(const int &I, const int &J) const;
+	__host__ __device__ const int& at(const int &I, const int &J) const;
-__host__ __device__ int* data(); //pointer to float4 representation of matrix
+__host__ __device__ int* data(); //pointer to int4 representation of matrix
-__host__ __device__ const int* data() const; //pointer to float4 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;

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];
 	}

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__ double& operator[](const int &I);
-	__host__ __device__ const float& operator[](const int &I) const;
+	__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__ double& operator[](const int &I);
-	__host__ __device__ const float& operator[](const int &I) const;
+	__host__ __device__ const double& operator[](const int &I) const;
-	__host__ __device__ float& operator()(const int &I, const int &J);
+	__host__ __device__ double& operator()(const int &I, const int &J);
-	__host__ __device__ const float& operator()(const int &I, const int &J) const;
+	__host__ __device__ const double& operator()(const int &I, const int &J) const;
-	__host__ __device__ float& at(const int &I, const int &J);
+	__host__ __device__ double& at(const int &I, const int &J);
-	__host__ __device__ const float& at(const int &I, const int &J) const;
+	__host__ __device__ const double& at(const int &I, const int &J) const;
-__host__ __device__ double* data(); //pointer to float9 representation of matrix
+__host__ __device__ double* data(); //pointer to double9 representation of matrix
-__host__ __device__ const double* data() const; //pointer to float9 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;

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];
 	}

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__ int& operator[](const int &I);
-	__host__ __device__ const float& operator[](const int &I) const;
+	__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__ int& operator[](const int &I);
-	__host__ __device__ const float& operator[](const int &I) const;
+	__host__ __device__ const int& operator[](const int &I) const;
-	__host__ __device__ float& operator()(const int &I, const int &J);
+	__host__ __device__ int& operator()(const int &I, const int &J);
-	__host__ __device__ const float& operator()(const int &I, const int &J) const;
+	__host__ __device__ const int& operator()(const int &I, const int &J) const;
-	__host__ __device__ float& at(const int &I, const int &J);
+	__host__ __device__ int& at(const int &I, const int &J);
-	__host__ __device__ const float& at(const int &I, const int &J) const;
+	__host__ __device__ const int& at(const int &I, const int &J) const;
-__host__ __device__ int* data(); //pointer to float9 representation of matrix
+__host__ __device__ int* data(); //pointer to int9 representation of matrix
-__host__ __device__ const int* data() const; //pointer to float9 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;

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];
 	}

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__ double& operator[](const int &I);
-	__host__ __device__ const float& operator[](const int &I) const;
+	__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__ double& operator[](const int &I);
-	__host__ __device__ const float& operator[](const int &I) const;
+	__host__ __device__ const double& operator[](const int &I) const;
-	__host__ __device__ float& operator()(const int &I, const int &J);
+	__host__ __device__ double& operator()(const int &I, const int &J);
-	__host__ __device__ const float& operator()(const int &I, const int &J) const;
+	__host__ __device__ const double& operator()(const int &I, const int &J) const;
-	__host__ __device__ float& at(const int &I, const int &J);
+	__host__ __device__ double& at(const int &I, const int &J);
-	__host__ __device__ const float& at(const int &I, const int &J) const;
+	__host__ __device__ const double& at(const int &I, const int &J) const;
-__host__ __device__ double* data(); //pointer to float16 representation of matrix
+__host__ __device__ double* data(); //pointer to double16 representation of matrix
-__host__ __device__ const double* data() const; //pointer to float16 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;

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];
 	}

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__ int& operator[](const int &I);
-	__host__ __device__ const float& operator[](const int &I) const;
+	__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__ int& operator[](const int &I);
-	__host__ __device__ const float& operator[](const int &I) const;
+	__host__ __device__ const int& operator[](const int &I) const;
-	__host__ __device__ float& operator()(const int &I, const int &J);
+	__host__ __device__ int& operator()(const int &I, const int &J);
-	__host__ __device__ const float& operator()(const int &I, const int &J) const;
+	__host__ __device__ const int& operator()(const int &I, const int &J) const;
-	__host__ __device__ float& at(const int &I, const int &J);
+	__host__ __device__ int& at(const int &I, const int &J);
-	__host__ __device__ const float& at(const int &I, const int &J) const;
+	__host__ __device__ const int& at(const int &I, const int &J) const;
-__host__ __device__ int* data(); //pointer to float16 representation of matrix
+__host__ __device__ int* data(); //pointer to int16 representation of matrix
-__host__ __device__ const int* data() const; //pointer to float16 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;