gpgpu - CUDA Thrust reduction with double2 arrays -

i have following (compilable , executable) code using cuda thrust perform reductions of float2 arrays. works correctly

using namespace std;  // includes, system  #include <stdlib.h> #include <stdio.h> #include <string.h> #include <math.h> #include <conio.h>  #include <typeinfo>   #include <iostream>  // includes cuda #include <cuda.h> #include <cuda_runtime.h>  // includes thrust #include <thrust/host_vector.h> #include <thrust/device_vector.h> #include <thrust/reduce.h>  // float2 + struct struct add_float2 {     __device__ float2 operator()(const float2& a, const float2& b) const {         float2 r;         r.x = a.x + b.x;         r.y = a.y + b.y;         return r;     }  };  // double2 + struct struct add_double2 {     __device__ double2 operator()(const double2& a, const double2& b) const {         double2 r;         r.x = a.x + b.x;         r.y = a.y + b.y;         return r;     }  };  void main( int argc, char** argv)  {     int n = 20;      // --- host     float2* ha; ha = (float2*) malloc(n*sizeof(float2));     (unsigned i=0; i<n; ++i) {         ha[i].x = 1;         ha[i].y = 2;     }      // --- device     float2* da; cudamalloc((void**)&da,n*sizeof(float2));     cudamemcpy(da,ha,n*sizeof(float2),cudamemcpyhosttodevice);      thrust::device_ptr<float2> dev_ptr_1(da);     thrust::device_ptr<float2> dev_ptr_2(da+n);      float2 init; init.x = init.y = 0.0f;      float2 sum = thrust::reduce(dev_ptr_1,dev_ptr_2,init,add_float2());      cout << " real part = " << sum.x << "; imaginary part = " << sum.y << endl;      getch();   } 

however, when change float2 double2 in main program, namely

void main( int argc, char** argv)  {     int n = 20;      // --- host     double2* ha; ha = (double2*) malloc(n*sizeof(double2));     (unsigned i=0; i<n; ++i) {         ha[i].x = 1;         ha[i].y = 2;     }      // --- device     double2* da; cudamalloc((void**)&da,n*sizeof(double2));     cudamemcpy(da,ha,n*sizeof(double2),cudamemcpyhosttodevice);      thrust::device_ptr<double2> dev_ptr_1(da);     thrust::device_ptr<double2> dev_ptr_2(da+n);      double2 init; init.x = init.y = 0.0;      double2 sum = thrust::reduce(dev_ptr_1,dev_ptr_2,init,add_double2());      cout << " real part = " << sum.x << "; imaginary part = " << sum.y << endl;      getch();  } 

i receive exception @ reduce line. how can use cuda thrust reduction double2 arrays? doing wrong? in advance.

working solution following talonmies' answer

using namespace std;

// includes, system #include <stdlib.h> #include <stdio.h> #include <string.h> #include <math.h> #include <conio.h>  #include <typeinfo> #include <iostream>  // includes cuda #include <cuda.h> #include <cuda_runtime.h>  // includes thrust #include <thrust/host_vector.h> #include <thrust/device_vector.h> #include <thrust/reduce.h>  struct my_double2 {     double x, y; };  // double2 + struct struct add_my_double2 {     __device__ my_double2 operator()(const my_double2& a, const my_double2& b) const {         my_double2 r;         r.x = a.x + b.x;         r.y = a.y + b.y;         return r;     } };  void main( int argc, char** argv)  {     int n = 20;      // --- host     my_double2* ha; ha = (my_double2*) malloc(n*sizeof(my_double2));     (unsigned i=0; i<n; ++i) {         ha[i].x = 1;         ha[i].y = 2;     }      // --- device     my_double2* da; cudamalloc((void**)&da,n*sizeof(my_double2));     cudamemcpy(da,ha,n*sizeof(my_double2),cudamemcpyhosttodevice);      thrust::device_ptr<my_double2> dev_ptr_1(da);     thrust::device_ptr<my_double2> dev_ptr_2(da+n);      my_double2 init; init.x = init.y = 0.0;      cout << "here3\n";     my_double2 sum = thrust::reduce(dev_ptr_1,dev_ptr_2,init,add_my_double2());      cout << " real part = " << sum.x << "; imaginary part = " << sum.y << endl;      getch();  } 

this known incompatibility msvc , nvcc. see here example. solution define own version of double2 , use instead.

just reference, can compile , run code correctly on linux 64 bit box cuda 5.5.