从设备函数调用 Thrust 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。