调用从设备函数调用device_vector

Question

我有一个struct Cap其中我有一个thrust::device_vector另一种结构。当我编译代码时，出现一个错误，提示从设备函数SphericalFaceManager::makeCaps调用主机函数（thrust::device_vector<FloatIntPair>）。当我将__host__ __device__而不是__device__添加到成员函数和构造函数中时，代码会进行编译，但会收到与上述错误相同的警告，我认为它会在主机和设备之间复制数据。我的问题是，如何访问我的类中的设备向量，避免CPU和GPU之间的任何数据传输？

此后，你可以找到的代码：

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; // The ordered list of vertices generated from intersections by 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); 
} 

Answer 1

不能使用推力矢量（或std::vector）直接在设备代码。这在其他各种SO问题中都有提及，如here

如果要在设备代码中使用thrust::device_vector中的数据，则应该将指针传递给函数作为函数初始化参数。各种其他的SO问题给出了这样的例子，例如here

同样，你不能使用矢量方法，例如，设备代码中的.empty()或.push_back()。

您将需要用普通的C风格的分配器和C风格的索引数据访问来替换它们。

对于设备代码中push_back的多线程实现，我会推荐类似this的东西。这是一个完整的工作示例，演示了如何为矢量分配空间以及每个线程如何使用该空间以用于例如insertVertex。