[SYSTEMDS-3352] CUDA code gen support for connected components

General cleanup and bug fixing to make components.dml run.
Also contains improvements to handle single precision execution.

src/main/cuda/headers/Matrix.h

@@ -18,8 +18,6 @@
  */
 
 #pragma once
-#ifndef SYSTEMDS_MATRIX_H
-#define SYSTEMDS_MATRIX_H
 
 using uint32_t = unsigned int;
 using int32_t = int;
@@ -43,22 +41,22 @@ struct Matrix {
 
 #ifdef __CUDACC__
 
-template<typename T>
-uint32_t bin_search(T* values, uint32_t lower, uint32_t upper, T val) {
-	upper -= 1;
-	while(lower <= (upper-1)) {
-		uint32_t idx = (lower + upper) >> 1;
-		uint32_t vi = values[idx];
-		if (vi < val)
-			lower = idx + 1;
-		else {
-			if (vi <= val)
-				return idx;
-			upper = idx - 1;
-		}
-	}
-	return upper + 1;
-}
+//template<typename T>
+//uint32_t bin_search(T* values, uint32_t lower, uint32_t upper, T val) {
+//	upper -= 1;
+//	while(lower <= (upper-1)) {
+//		uint32_t idx = (lower + upper) >> 1;
+//		uint32_t vi = values[idx];
+//		if (vi < val)
+//			lower = idx + 1;
+//		else {
+//			if (vi <= val)
+//				return idx;
+//			upper = idx - 1;
+//		}
+//	}
+//	return upper + 1;
+//}
 
 template<typename T>
 class MatrixAccessor {
@@ -68,11 +66,11 @@ class MatrixAccessor {
 public:
 	MatrixAccessor() = default;
 
-	__device__ MatrixAccessor(Matrix<T>* mat) : _mat(mat) {}
+	__device__ explicit MatrixAccessor(Matrix<T>* mat) : _mat(mat) {}
 
 	__device__ void init(Matrix<T>* mat) { _mat = mat; }
 
-	__device__ uint32_t& nnz() { return return _mat->row_ptr == nullptr ? _mat->rows * _mat->cols : _mat->nnz; }
+//	__device__ uint32_t& nnz() { return _mat->row_ptr == nullptr ? _mat->rows * _mat->cols : _mat->nnz; }
 	__device__ uint32_t cols() { return _mat->cols; }
 	__device__ uint32_t rows() { return _mat->rows; }
 
@@ -96,14 +94,14 @@ class MatrixAccessor {
 	}
 
 	__device__ uint32_t row_len(uint32_t rix) {
-		return _mat->row_ptr == nullptr ? row_len_dense(rix) : row_len_sparse(rix);
+		return _mat->row_ptr == nullptr ? _mat->rows : row_len_sparse(rix);
 	}
 
 	__device__ uint32_t* col_idxs(uint32_t rix) { return cols_sparse(rix); }
 
 	__device__ void set(uint32_t r, uint32_t c, T v) { set_sparse(r,c,v); }
 
-	__device__ uint32_t* indexes() {  return _mat->row_ptr;	}
+//	__device__ uint32_t* indexes() {  return _mat->row_ptr;	}
 
 	__device__ bool hasData() { return _mat->data != nullptr; }
 private:
@@ -127,10 +125,6 @@ class MatrixAccessor {
 		return &(_mat->data[rix]);
 	}
 
-	__device__ uint32_t row_len_dense(uint32_t rix) {
-		return _mat->rows;
-	}
-
 	//ToDo sparse accessors
 	__device__ uint32_t len_sparse() {
 		return _mat->row_ptr[_mat->rows];
@@ -145,8 +139,8 @@ class MatrixAccessor {
 	}
 
 	__device__ T& val_sparse_rc(uint32_t r, uint32_t c) {
-//		printf("TBI: val_sparse_rc\n");
-//		asm("trap;");
+		printf("TBI: val_sparse_rc(%d, %d)\n", r, c);
+		asm("trap;");
 
 		return _mat->data[0];
 	}
@@ -228,5 +222,3 @@ class RingBuffer {
 };
 
 #endif // __CUDACC_RTC__
-
-#endif //SYSTEMDS_MATRIX_H