从设备函数调用 Thrust device_vector
calling Thrust device_vector from a device function
我有一个 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);
}
您不能在设备代码中直接使用推力矢量(或std::vector
)。这在其他各种 SO 问题中都有提到,例如 here
如果你想在设备代码中使用 thrust::device_vector
中的数据,你应该传递一个指向数据的指针作为函子初始化参数。各种其他 SO 问题给出了这方面的例子,例如
同样,您不能使用矢量方法,例如.empty()
或 .push_back()
设备代码。
您需要将它们替换为普通的 C 风格分配器和 C 风格索引数据访问。
对于 push_back 在设备代码中的多线程实现,我会推荐类似 this 的东西。这是一个完整的示例,演示了如何为向量分配 space 以及每个线程如何将它用于 insertVertex
例如。
我有一个 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);
}
您不能在设备代码中直接使用推力矢量(或std::vector
)。这在其他各种 SO 问题中都有提到,例如 here
如果你想在设备代码中使用 thrust::device_vector
中的数据,你应该传递一个指向数据的指针作为函子初始化参数。各种其他 SO 问题给出了这方面的例子,例如
同样,您不能使用矢量方法,例如.empty()
或 .push_back()
设备代码。
您需要将它们替换为普通的 C 风格分配器和 C 风格索引数据访问。
对于 push_back 在设备代码中的多线程实现,我会推荐类似 this 的东西。这是一个完整的示例,演示了如何为向量分配 space 以及每个线程如何将它用于 insertVertex
例如。