do thrust::min_element on arrays on the gpu

thrust::min_element returns an iterator.

From the documentation:

min_element finds the smallest element in the range [first, last). It returns the first iterator i in [first, last) such that no other iterator in [first, last) points to a value smaller than *i.

An iterator is something like a pointer. It indicates a position of an element in a container. Like pointers, iterators can be added to, subtracted from, etc.

So we could extract this iterator directly:

thrust::device_ptr<double> result_position = thrust::min_element(... 

or another approach would be to get a relative offset to that position, from the start of the container:

int result_offset = thrust::min_element(g_ptr, ...) - g_ptr; 

This works because iterators (or thrust::device_ptr) can be subtracted. The iterator returned from min_element minus the start of the container will give the offset to the position of the minimum element.

Here's a worked example based on your code:

$ cat t957.cu #include <stdio.h> #include <stdlib.h> /* for rand() */ #include <iostream> #include <thrust/device_ptr.h> #include <thrust/extrema.h> using namespace std; bool errorAsk(const char *s="n/a") { cudaError_t err=cudaGetLastError(); if(err==cudaSuccess) return false; printf("CUDA error [%s]: %s\n",s,cudaGetErrorString(err)); return true; }; double *fillArray(double *c_idata,int N,double constant) { int n; for (n = 0; n < N; n++) { c_idata[n] = constant*floor(drand48()*10.0); } return c_idata; } int main(int argc,char *argv[]) { int N; N = 100; double *c_data,*g_data; // result = new double[N]; c_data = new double[N]; c_data = fillArray(c_data,N,1.0); c_data[32] = -1.0; cudaMalloc(&g_data,N*sizeof(double)); cudaMemcpy(g_data,c_data,N*sizeof(double),cudaMemcpyHostToDevice); thrust::device_ptr<double> g_ptr = thrust::device_pointer_cast(g_data); int result_offset = thrust::min_element(g_ptr, g_ptr + N) - g_ptr; double min_value = *(g_ptr + result_offset); // we could also do this: // double min_value = c_data[result_offset]; std::cout<< "min value found at position: " << result_offset << " value: " << min_value << std::endl; } $ nvcc -o t957 t957.cu $ ./t957 min value found at position: 32 value: -1 $ 

The thrust quick start guide gives a brief description of iterators and their usage in thrust.