Reduce matrix rows with CUDA

by

Since you mentioned you need general reduction algorithm other than sum only. I will try to give 3 approaches here. kernel approach may have the highest performance. thrust approach is easiest to implement. cuBLAS approach works only with sum and have good performance.

Kernel Approach

Here’s a very good doc introducing how to optimize standard parallel reduction. Standard reduction can be divide into 2 stages.

  1. Multiple thread blocks each reduces one part of the data;
  2. One thread block reduces from result of stage 1 to the final 1 element.

For your multi-reduction (reduce rows of mat) problem, only stage 1 is enough. The idea is to reduce 1 row per thread block. For further considerations like multi-row per thread block or 1 row per multiple thread blocks, you can refer to the paper provided by @Novak. This may improve the performance more, especially for matrices with bad shape.

Thrust Approach

General multi-reduction can be done by thrust::reduction_by_key in a few minutes. You can find some discussions here Determining the least element and its position in each matrix column with CUDA Thrust.

However thrust::reduction_by_key does not assume each row has the same length, so you will get performance penalty. Another post How to normalize matrix columns in CUDA with max performance? gives profiling comparison between thrust::reduction_by_key and cuBLAS approach on sum of rows. It may give you a basic understanding about the performance.

cuBLAS Approach

Sum of rows/cols of a matrix A can be seen as a matrix-vector multiplication where the elements of the vector are all ones. it can be represented by the following matlab code.

y = A * ones(size(A,2),1);

where y is the sum of rows of A.

cuBLAS libary provides a high performance matrix-vector multiplication function cublas<t>gemv() for this operation.

Timing result shows that this routine is only 10~50% slower than simply read all the elements of A once, which can be seen as the theoretical upper limit of the performance for this operation.

More Related Contents:

  1. double pointer to constant array
  2. Sort 2D matrix in C
  3. Dynamic allocation of an unknown matrix in C
  4. Cholesky decomposition with OpenMP
  5. Dynamic programming – Largest square block
  6. sending 3d array to CUDA kernel
  7. Function to dynamically allocate matrix
  8. C – allocating a matrix in a function
  9. Traverse Matrix in Diagonal strips
  10. allocate matrix in C
  11. How to work on a sub-matrix in a matrix by pointer?
  12. Use of cudamalloc(). Why the double pointer?
  13. Optimized matrix multiplication in C
  14. undefined reference to `std::ios_base::Init::Init()’
  15. MPI partition matrix into blocks
  16. How performing multiple matrix multiplications in CUDA?
  17. Cuda C – Linker error – undefined reference
  18. Passing a matrix in a function (C)
  19. Timing CUDA operations
  20. why do we need cudaDeviceSynchronize(); in kernels with device-printf?
  21. C – scanf() vs gets() vs fgets()
  22. how to control popen stdin, stdout, stderr redirection?
  23. reading a text file into an array in c
  24. fflush(stdin) function does not work
  25. Are floating point operations in C associative?
  26. An alternative for the deprecated __malloc_hook functionality of glibc
  27. What’s the difference between array and &array?
  28. Convert Eigen Matrix to C array
  29. C Function to Convert float to byte array
  30. Efficient way to find task_struct by pid

Leave a Comment