[Code/Python] Vector fixes && Cuda matrix-to-vec reduce fixes

Author: EgorOrachyov

cubool/sources/cuda/kernels/spreduce.cuh

@@ -143,7 +143,6 @@ namespace cubool {
             using VectorType = details::SpVector<IndexType, AllocType>;
 
             VectorType operator()(const MatrixType& m) {
-                auto nrows = m.m_rows;
                 auto ncols = m.m_cols;
                 auto nvals = m.m_vals;
 
@@ -156,33 +155,34 @@ namespace cubool {
                          __device__ (IndexType i) {
                          auto j = colIndex[i];
 
-                         atomicOr((mask + j).get(), (IndexType) 1);
+                         mask[j] = 0x1u;
                      }
                 );
 
                 // Count nnz of the vector
                 auto resultSize = thrust::reduce(mask.begin(), mask.end(), (IndexType) 0);
 
                 ContainerType<index> result(resultSize);
-                ContainerType<index> offsets(mask.size());
-
-                // Evaluate write offsets of the result
-                thrust::exclusive_scan(mask.begin(), mask.end(), offsets.begin(), (IndexType) 0, thrust::plus<IndexType>());
+                ContainerType<index> order(1);
+                thrust::fill(order.begin(), order.end(), (IndexType) 0);
 
                 // For each value of the mask write in the result buffer if it not zero
                 thrust::for_each(thrust::counting_iterator<IndexType>(0), thrust::counting_iterator<IndexType>(ncols),
-                     [mask = mask.data(), result = result.data(), offsets = offsets.data()]
+                     [mask = mask.data(), result = result.data(), order = order.data()]
                          __device__ (IndexType i) {
                         if (mask[i] > 0) {
-                            auto offset = offsets[i];
+                            auto offset = atomicAdd(order.get(), (IndexType) 1);
                             auto colId = i;
 
                             result[offset] = colId;
                         }
                      }
                 );
 
-                return VectorType(std::move(result), nrows, resultSize);
+                // Sort values within vector
+                thrust::sort(result.begin(), result.end());
+
+                return VectorType(std::move(result), ncols, resultSize);
             }
 
         };

python/features.py

@@ -111,5 +111,20 @@ def gen_matrix(size, seed):
 print(vals)
 print(a.equals(a))
 
-matrix = pycubool.Matrix.generate(shape=(4, 4), density=0.5)
-print(matrix)
+matrix = pycubool.Matrix.generate(shape=(5, 4), density=0.5)
+print(matrix, matrix.extract_row(1), matrix.extract_col(2), sep="\n")
+
+matrix = pycubool.Matrix.generate(shape=(5, 4), density=0.5)
+vector = pycubool.Vector.generate(nrows=4, density=0.6)
+print(matrix, vector, matrix.mxv(vector), sep="\n")
+
+matrix = pycubool.Matrix.generate(shape=(5, 4), density=0.5)
+vector = pycubool.Vector.generate(nrows=5, density=0.6)
+print(matrix, vector, vector.vxm(matrix), sep="\n")
+
+matrix = pycubool.Matrix.generate(shape=(10, 6), density=0.2)
+print(matrix, matrix.reduce_vector(transpose=False), matrix.reduce_vector(transpose=True), sep="\n")
+
+
+
+

python/pycubool/__init__.py

@@ -21,6 +21,7 @@
 from .wrapper import *
 from .utils import *
 from .matrix import *
+from .vector import *
 from .io import *
 from .gviz import *
 

python/pycubool/__main__.py

@@ -493,7 +493,7 @@ def extract_matrix(self, i, j, shape, out=None, time_check=False):
 
     def extract_row(self, i, out=None):
         if out is None:
-            out = Vector.empty(self.ncols)
+            out = vector.Vector.empty(self.ncols)
 
         status = wrapper.loaded_dll.cuBool_Matrix_ExtractRow(
             out.hnd,
@@ -507,7 +507,7 @@ def extract_row(self, i, out=None):
 
     def extract_col(self, j, out=None):
         if out is None:
-            out = Vector.empty(self.nrows)
+            out = vector.Vector.empty(self.nrows)
 
         status = wrapper.loaded_dll.cuBool_Matrix_ExtractCol(
             out.hnd,
@@ -563,7 +563,7 @@ def mxm(self, other, out=None, accumulate=False, time_check=False):
 
     def mxv(self, other, out=None, time_check=False):
         if out is None:
-            out = Vector.empty(self.nrows)
+            out = vector.Vector.empty(self.nrows)
 
         status = wrapper.loaded_dll.cuBool_MxV(
             out.hnd,
@@ -689,7 +689,7 @@ def reduce(self, time_check=False):
     def reduce_vector(self, out=None, transpose=False, time_check=False):
         if out is None:
             nrows = self.ncols if transpose else self.nrows
-            out = Vector.empty(nrows)
+            out = vector.Vector.empty(nrows)
 
         status = wrapper.loaded_dll.cuBool_Matrix_Reduce(
             out.hnd,

python/pycubool/matrix.py

@@ -229,6 +229,9 @@ def ewiseadd(self, other, out=None, time_check=False):
             ctypes.c_uint(bridge.get_ewiseadd_hints(time_check=time_check))
         )
 
+        bridge.check(status)
+        return out
+
     def reduce(self, time_check=False):
         value = ctypes.c_uint(0)