cuda - calling Thrust device_vector from a device function -

i have struct cap inside have thrust::device_vector of structure. when compile code, error complains calling host function (thrust::device_vector<floatintpair>) device function sphericalfacemanager::makecaps. when add __host__ __device__ instead of __device__ member functions , constructors code compiles receive warning same aforementioned error , think copies data between host , device. question how can access device vectors in classes avoiding data transfer between cpu , gpu?

hereafter can find code:

struct particleid {  int solver; int ngb; int oldngb; llint no; llint masterno;  __device__ particleid() {     solver = -8;     ngb = 0;     oldngb = 0;     no = 0;     masterno = -1; } };   struct baseparticle {  float h; float3 pos; particleid id;  __device__ baseparticle(const float3& _pos, const float& _h, const particleid& _id) :      h(_h), pos(_pos), id(_id) { }  };   struct floatintpair{  float first; int second;  __device__ floatintpair(const float& _first, int _second) : first(_first), second(_second) { } __device__ floatintpair(const floatintpair& sample) : first(sample.first), second(sample.second) { }  static struct {     __device__ bool operator()(const floatintpair& a, const floatintpair& b) {  return a.first < b.first; } } lessop; };   struct cap {  float3 ex; float3 ey; float radius; float height;  float3 center; float3 normal;  baseparticle* ap; baseparticle* bp;  thrust::device_vector<floatintpair> vertices; // ordered list of vertices generated intersections other circles  __device__ inline float findangle(const float3& vertex) const {      float result;     float3 r = (vertex - center);     result = atan2(r|ey,r|ex);     return result += (result < 0.0) * (2.0 * _pi); }  __device__ void insertvertex(const float3& vertex, int id) {      float theta;     if (!vertices.empty())         theta = findangle(vertex);     else {         ex = normalvec(vertex - center);         ey = normal ^ ex;         theta = 0.0;     }     vertices.push_back(floatintpair(theta,id)); }  __device__ cap(baseparticle* _ap, baseparticle* _bp) : ap(_ap), bp(_bp) {      //compute normal, center, radius     float d = mag(bp->pos - ap->pos);     if(d == 0.0){         normal = vector1(0.0);         center = ap->pos;         radius = height = 0.0;     } else {         normal = (bp->pos - ap->pos) / d;         float x = (d * d - bp->h * bp->h + ap->h * ap->h) / (2.0 * d);         center = ap->pos + normal * x;         if (x >= ap->h) {             radius = height = 0.0;             return;         }         radius = sqrt(ap->h * ap->h - x * x);         height = min(2.0 * ap->h, ap->h - x);          float3 vec001 = vector(0.0,0.0,1.0);             float3 vec011 = vector(0.0,1.0,1.0);          ex = normalvec(vec001 ^ normal);         if (mag2(ex) < geoeps()) {             ex = ex = normalvec(vec011 ^ normal);         }          ey = normal ^ ex;     } } };  class sphericalfacemanager { baseparticle* particle; int basesigma; public: thrust::device_vector<cap> caps;   thrust::device_vector<float3> vertexpool;     __device__ void makecaps(); };   __device__ void sphericalfacemanager::makecaps() {  baseparticle* ap; baseparticle* bp; cap acap(ap,bp); } 

you cannot use thrust vectors (or std::vector) directly in device code. mentioned in various other questions such here

if want use data in thrust::device_vector in device code, should pass pointer data functor initializing parameter. various other questions give examples of this, such here

likewise, cannot use vector methods, e.g. .empty() or .push_back() in device code.

you need replace these ordinary c-style allocators , c-style indexed data access.

for multi-threaded implementation of push_back in device code, recommend this. worked example demonstrates how allocate space vector , how each thread can use insertvertex example.