out of bounds array access on CsrMV

[samuelpmishLLNL]

Hi,

The example problem on the cub::DeviceSpmv::CsrMV documentation page fails cuda-memcheck with an out of bounds array access:

#include <iostream>

#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
#include <cub/device/device_spmv.cuh>

#define CUDACHECK(cmd)                                              \
  do {                                                              \
    cudaError_t e = cmd;                                            \
    if (e != cudaSuccess) {                                         \
      printf("Failed: Cuda error %s:%d '%s'\n", __FILE__, __LINE__, \
             cudaGetErrorString(e));                                \
      exit(EXIT_FAILURE);                                           \
    }                                                               \
  } while (0)

template < typename T >
auto raw(thrust::device_vector< T > & p) {
  return thrust::raw_pointer_cast(p.data());
}

int main() {

  int nnz = 24;
  int num_rows = 9;
  int num_cols = 9;
  std::vector< int > row_ptr = {0, 2, 5, 7, 10, 14, 17, 19, 22, 24};
  std::vector< int > col_ind = {
    1, 3, 0, 2, 4, 1, 5, 0,
    4, 6, 1, 3, 5, 7, 2, 4,
    8, 3, 7, 4, 6, 8, 5, 7
  };

  std::vector< double > values(nnz, 1);
  std::vector< double > x(num_cols, 1);
  #endif

  thrust::device_vector< int > row_ptr_d(row_ptr);
  thrust::device_vector< int > col_ind_d(col_ind);
  thrust::device_vector< double > values_d(values);
  thrust::device_vector< double > x_d = x;
  thrust::device_vector< double > b_d(num_rows);

  void * buf = NULL;
  size_t size = 0;
  cub::DeviceSpmv::CsrMV(
    buf, size,
    raw(values_d), raw(row_ptr_d), raw(col_ind_d),
    raw(x_d), raw(b_d),
    num_rows, num_cols, nnz);

  CUDACHECK(cudaMalloc(&buf, size));

  cub::DeviceSpmv::CsrMV(
    buf, size,
    raw(values_d), raw(row_ptr_d), raw(col_ind_d),
    raw(x_d), raw(b_d),
    num_rows, num_cols, nnz);

  // $ cuda-memcheck.exe csr_matrix_test.exe
  //========= CUDA-MEMCHECK
  //========= Invalid __global__ read of size 4
  //=========     at 0x00000510 in void cub::DeviceSpmvKernel<
  //                 cub::DispatchSpmv<double, int>::PtxSpmvPolicyT, 
  //                 cub::ReduceByKeyScanTileState<double, int, bool=1>, 
  //                 double, int, int2, bool=0, bool=0
  //                 >(
  //                 cub::SpmvParams<cub::DispatchSpmv<double, int>::PtxSpmvPolicyT, double>,
  //                 int*, 
  //                 cub::KeyValuePair<cub::DispatchSpmv<double, int>::PtxSpmvPolicyT, 
  //                 cub::SpmvParams>*, int, int, int)
  //=========     by thread (9,0,0) in block (0,0,0)
  //=========     Address 0xb02000028 is out of bounds

  cudaFree(buf);

  thrust::host_vector< double > b = b_d;

  for (auto value : b) {
    std::cout << value << std::endl;
  }

}

The address of the offending read is one past the end of the row_ptr array, and adding an additional entry to that list seems to stop the error. I'm using the version of cub that comes bundled with the cuda toolkit 11.0 (#define CUB_VERSION 100909 in cub/version.cuh). Using a GTX 1080 ti on Windows 10, with MSVC 2019 host compiler.

[samuelpmishLLNL]

oops, this is probably a duplicate of https://github.com/NVlabs/cub/issues/139, but it looks like the PR they submitted to fix it has been ignored for about a year.

[lilohuang]

@allisonvacanti I think this issue has been fixed with #352 Could you help me to check? Thanks.

[alliepiper]

Good catch, yes, this was fixed with the recent DeviceSpmv updates.