cuda - Create an object with fields on the device directly -

i'm trying create class allocated on device. want constructor run on device whole object including fields inside automatically allocated on device instead of having create host object copy manually device.

i'm using thrust device_new

here code:

using namespace thrust;  class particle {     public:     int* data;      __device__  particle()     {         data = new int[10];         (int i=0; i<10; i++)         {             data[i] = i*2;         }     } };   __global__ void test(particle* p) {     (int i=0; i<10; i++)         printf("%d\n", p->data[i]); }  int main() {      device_ptr<particle> p = device_new<particle>();      test<<<1,1>>>(thrust::raw_pointer_cast(p));       cudadevicesynchronize();      printf("done!\n");  } 

i annotated constructor __device__ , used device_new (thrust), doesn't work, can explain me why?

cheers help

i believe answer lies in description given here. knows thrust under hood come along , indicate whether true or not.

although thrust has changed lot since 2009, believe device_new may still using form of operation object temporarily instantiated on host, copied device. believe size limitation described in above reference no longer applicable, however.

i able work:

#include <stdio.h> #include <thrust/device_ptr.h> #include <thrust/device_new.h>  #define n 512  using namespace thrust;  class particle {     public:     int data[n];      __device__ __host__  particle()     { //        data = new int[10];         (int i=0; i<n; i++)         {             data[i] = i*2;         }     } };   __global__ void test(particle* p) {     (int i=0; i<n; i++)         printf("%d\n", p->data[i]); }  int main() {      device_ptr<particle> p = device_new<particle>();      test<<<1,1>>>(thrust::raw_pointer_cast(p));       cudadevicesynchronize();      printf("done!\n");  } 

interestingly, gives bogus results if omit __host__ decorator on constructor, suggesting me temporary object copy mechanism still in place. gives bogus results (and cuda-memcheck reports out-of-bounds access errors) if switch using dynamic allocation data instead of static, suggesting me device_new using temporary object creation on host followed copy device.