optimize by reviewing comments

Author: rhdong

cpp/include/raft/core/bitmap.hpp

@@ -125,8 +125,6 @@ struct bitmap_view : public bitset_view<bitmap_t, index_t> {
    * The bitmap is interpreted as a row-major matrix, with rows and columns defined by
    * the dimensions of the bitmap.
    *
-   * @tparam bitmap_t The data type of the elements in the bitmap matrix.
-   * @tparam index_t The data type used for indexing the elements in the matrices.
    * @tparam csr_matrix_t Specifies the CSR matrix type, constrained to raft::device_csr_matrix.
    *
    * @param[in] res RAFT resources for managing CUDA streams and execution policies.

cpp/include/raft/core/bitset.hpp

@@ -231,8 +231,6 @@ struct bitset_view {
    * //                1, 1, 1, 1];
    * @endcode
    *
-   * @tparam bitset_t The data type of the elements in the bitset matrix.
-   * @tparam index_t The data type used for indexing the elements in the matrices.
    * @tparam csr_matrix_t Specifies the CSR matrix type, constrained to raft::device_csr_matrix.
    *
    * @param[in] res RAFT resources for managing CUDA streams and execution policies.

cpp/include/raft/sparse/convert/detail/bitmap_to_csr.cuh

@@ -330,9 +330,9 @@ void bitmap_to_csr(raft::resources const& handle,
     thrust_policy, sub_nnz.data(), sub_nnz.data() + sub_nnz_size + 1, sub_nnz.data());
 
   if constexpr (is_device_csr_sparsity_owning_v<csr_matrix_t>) {
-    index_t nnz = 0;
+    nnz_t nnz = 0;
     RAFT_CUDA_TRY(cudaMemcpyAsync(
-      &nnz, sub_nnz.data() + sub_nnz_size, sizeof(index_t), cudaMemcpyDeviceToHost, stream));
+      &nnz, sub_nnz.data() + sub_nnz_size, sizeof(nnz_t), cudaMemcpyDeviceToHost, stream));
     resource::sync_stream(handle);
     csr.initialize_sparsity(nnz);
   }

cpp/include/raft/sparse/convert/detail/bitset_to_csr.cuh

@@ -58,11 +58,8 @@ RAFT_KERNEL repeat_csr_kernel(const index_t* indptr,
 
   __syncthreads();
 
-  int block_offset = blockIdx.x * blockDim.x;
-
   index_t item;
-  int idx = block_offset + threadIdx.x;
-  item    = (idx < nnz) ? indices[idx] : -1;
+  item = (global_id < nnz) ? indices[global_id] : -1;
 
   __syncthreads();
 
@@ -144,10 +141,10 @@ void bitset_to_csr(raft::resources const& handle,
   thrust::exclusive_scan(
     thrust_policy, sub_nnz.data(), sub_nnz.data() + sub_nnz_size + 1, sub_nnz.data());
 
-  index_t bitset_nnz = 0;
+  nnz_t bitset_nnz = 0;
   if constexpr (is_device_csr_sparsity_owning_v<csr_matrix_t>) {
     RAFT_CUDA_TRY(cudaMemcpyAsync(
-      &bitset_nnz, sub_nnz.data() + sub_nnz_size, sizeof(index_t), cudaMemcpyDeviceToHost, stream));
+      &bitset_nnz, sub_nnz.data() + sub_nnz_size, sizeof(nnz_t), cudaMemcpyDeviceToHost, stream));
     resource::sync_stream(handle);
     csr.initialize_sparsity(bitset_nnz * csr_view.get_n_rows());
   } else {