Small cleanup to sparse primitives. (rapidsai#203)

* Remove unused code.
* Fix comments.

Ported from rapidsai/cuml#3575 .  Rebased onto 0.20.

Authors:
  - Jiaming Yuan (https://github.com/trivialfis)

Approvers:
  - Corey J. Nolet (https://github.com/cjnolet)

URL: rapidsai#203

cpp/include/raft/sparse/linalg/degree.cuh

@@ -105,7 +105,7 @@ __global__ void coo_degree_scalar_kernel(const int *rows, const T *vals,
 }
 
 /**
- * @brief Count the number of values for each row matching a particular scalar
+ * @brief Count the number of values for each row that doesn't match a particular scalar
  * @tparam TPB_X: number of threads to use per block
  * @tparam T: the type name of the underlying value arrays
  * @param in: Input COO array
@@ -124,7 +124,7 @@ void coo_degree_scalar(COO<T> *in, T scalar, int *results,
 }
 
 /**
- * @brief Count the number of values for each row matching a particular scalar
+ * @brief Count the number of values for each row that doesn't match a particular scalar
  * @tparam TPB_X: number of threads to use per block
  * @tparam T: the type name of the underlying value arrays
  * @param rows: Input COO row array
@@ -181,4 +181,4 @@ void coo_degree_nz(COO<T> *in, int *results, cudaStream_t stream) {
 
 };  // end NAMESPACE linalg
 };  // end NAMESPACE sparse
-};  // end NAMESPACE raft
+};  // end NAMESPACE raft

cpp/include/raft/sparse/op/filter.cuh

@@ -28,9 +28,9 @@
 #include <thrust/scan.h>
 
 #include <cuda_runtime.h>
-#include <stdio.h>
 
 #include <algorithm>
+#include <cstdio>
 #include <iostream>
 
 #include <raft/sparse/utils.h>
@@ -41,29 +41,6 @@ namespace raft {
 namespace sparse {
 namespace op {
 
-template <int TPB_X, typename T>
-__global__ void coo_remove_zeros_kernel(const int *rows, const int *cols,
-                                        const T *vals, int nnz, int *crows,
-                                        int *ccols, T *cvals, int *ex_scan,
-                                        int *cur_ex_scan, int m) {
-  int row = (blockIdx.x * TPB_X) + threadIdx.x;
-
-  if (row < m) {
-    int start = cur_ex_scan[row];
-    int stop = get_stop_idx(row, m, nnz, cur_ex_scan);
-    int cur_out_idx = ex_scan[row];
-
-    for (int idx = start; idx < stop; idx++) {
-      if (vals[idx] != 0.0) {
-        crows[cur_out_idx] = rows[idx];
-        ccols[cur_out_idx] = cols[idx];
-        cvals[cur_out_idx] = vals[idx];
-        ++cur_out_idx;
-      }
-    }
-  }
-}
-
 template <int TPB_X, typename T>
 __global__ void coo_remove_scalar_kernel(const int *rows, const int *cols,
                                          const T *vals, int nnz, int *crows,
@@ -198,4 +175,4 @@ void coo_remove_zeros(COO<T> *in, COO<T> *out,
 
 };  // namespace op
 };  // end NAMESPACE sparse
-};  // end NAMESPACE raft
+};  // end NAMESPACE raft