thrust::min_element Access violation while reading from location

Question

I want to find the minimum value of an array using the gpu therfore i want to use thrust::min_element my data is in the device this is why i have to use thrust::device but i have an "Access violation while reading from location 0x0000000701240000" and am taking into tuple.inl function:

inline __host__ __device__
  cons( T1& t1, T2& t2, T3& t3, T4& t4, T5& t5,
        T6& t6, T7& t7, T8& t8, T9& t9, T10& t10 )

but if i use thrust::host it works !!! this is my code. If there is something wrong plz tell me.

#include <thrust/extrema.h>
#include <thrust/execution_policy.h>
#include <time.h>
int main()
{
    int nx=200;
    int ny=200;
    float cpt=0;
    clock_t start,end;
    double time;
    float *in,*d_in;
    float moy,*d_moy;
    in=(float*)malloc(nx*ny*sizeof(float));
    moy=0.0f;
    cudaMalloc((void**)&d_in,nx*ny*sizeof(float));
    cudaMalloc((void**)&d_moy,sizeof(float));
    for(int i=0;i<nx*ny;i++){in[i]=i+0.07;cpt+=in[i];}
    cudaMemcpy(d_in,in,nx*ny*sizeof(float),cudaMemcpyHostToDevice);
     start=clock(); 
     //float result= thrust::reduce(thrust::device, d_in, d_in + nx*ny);
     float *result=thrust::min_element(thrust::device, d_in , d_in+ nx*ny);
     end=clock();
     time=((double)(end-start))/CLOCKS_PER_SEC;
     printf("result= %f and correct result is %f time= %lf  \n",*result,in[0],time);
     system("pause");
}

Answer 1

It is a bug of thrust::min_element. The program crashed when raw pointers are used. This bug only exists in CUDA7.5 Thrust1.8.2 or earlier.

You could use thrust::device_vector or thrust::device_ptr instead. This is a better way to use thrust.

#include <iostream>
#include <thrust/sequence.h>
#include <thrust/extrema.h>
#include <thrust/device_vector.h>
#include <thrust/execution_policy.h>

int main() {
  int n = 1000;
  thrust::device_vector<float> in(n);
  thrust::sequence(in.begin(), in.end(), 123);

  std::cerr << "by iterator:" << std::endl;
  thrust::device_vector<float>::iterator it_result =
      thrust::min_element(in.begin(), in.end());
  std::cerr << *it_result << std::endl;

  std::cerr << "by device_ptr:" << std::endl;
  thrust::device_ptr<float> ptr_in = in.data();
  thrust::device_ptr<float> ptr_result =
      thrust::min_element(ptr_in, ptr_in + in.size());
  std::cerr << *ptr_result << std::endl;

  std::cerr << "by pointer:" << std::endl;
  float* raw_in = thrust::raw_pointer_cast(in.data());
  std::cerr << "before min_element" << std::endl;
  float* result = thrust::min_element(thrust::device, raw_in, raw_in + in.size());
  std::cerr << "after min_element" << std::endl;
  std::cerr << in[result - raw_in] << std::endl;

  return 0;
}

The result demonstrates that thrust::min_element crashes with raw pointers.

$ nvcc -o test test.cu && ./test
by iterator:
123
by device_ptr:
123
by pointer:
before min_element
Segmentation fault (core dumped)

On the other hand, if this bug does not exist, your original code still have problem. As @talonmies said, result is device pointer. You need to copy the data pointed by it from device to host, before you can print it out.