Tricks and Tips: Finding the minimum element position of a cudaMalloc'ed array by thrust::min element

In this post, we are considering the very common case when one has a cudaMalloc’ed array and wants to find the position and the value of its minimum element by thrust::min_element. Below, we are providing a full example of that.

Basically, the solution below exploits the idea of wrapping a thrust::device_ptr around the cudaMalloc’ed linear memory and the position is found by thrust::distance.
Here is the full code:

#include <thrust/device_vector.h>
#include <thrust/extrema.h>

/********************/
/* CUDA ERROR CHECK */
/********************/
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, char *file, int line, bool abort=true)
{
   if (code != cudaSuccess)
   {
      fprintf(stderr,"GPUassert: %s %s %dn", cudaGetErrorString(code), file, line);
      if (abort) exit(code);
   }
}

/********/
/* MAIN */
/********/
int main() {

    const int N = 16;

    srand(time(NULL));

    // --- Host side memory allocation and initialization
    float *h_A = (float*)malloc(N * sizeof(float));
    for (int i=0; i<N; i++) h_A[i] = rand();

    // --- Device side memory allocation and initialization
    float *d_A; gpuErrchk(cudaMalloc((void**)&d_A, N * sizeof(float)));
    gpuErrchk(cudaMemcpy(d_A, h_A, N * sizeof(float), cudaMemcpyHostToDevice));

    thrust::device_ptr<float> dp = thrust::device_pointer_cast(d_A);
    thrust::device_ptr<float> pos = thrust::min_element(dp, dp + N);

    unsigned int pos_index = thrust::distance(dp, pos);
    float min_val;
    gpuErrchk(cudaMemcpy(&min_val, &d_A[pos_index], sizeof(float), cudaMemcpyDeviceToHost));

    for (int i=0; i<N; i++) printf("d_A[%i] = %fn", i, h_A[i]);
    printf("n");
    printf("Position of the minimum element = %i; Value of the minimum element = %fn", thrust::distance(dp, pos), min_val);

    cudaDeviceReset();

    return 0;
}

Leave a Reply

Your email address will not be published. Required fields are marked *