我需要创建某种堆栈来爬取 cuda 内核中的树。我以为我可以使用 thrust::device_vector 但显然不行。有这方面的 API 还是我必须自己编写代码。
__global__
void step_objects_kernel(ContainerNode* root, ObjectNode** objs, ObjectNode* new_objs, size_t n, real dt, real g)
{
int idx = blockIdx.x * gridDim.x + threadIdx.x;
if(idx >= n) return;
thrust::device_vector<Node*> to_visit;
to_visit.push_back(root);
vec3 a = {0};
while(!to_visit.empty())
{
Node* n = to_visit.back();
to_visit.pop_back();
}
}
error: calling a __host__ function("thrust::device_vector<Node *, thrust::device_malloc_allocator<Node *> > ::device_vector") from a __global__ function("step_objects_kernel") is not allowed
最佳答案
thrust::device_vector
在 CUDA 设备代码中不可用是正确的。
我不知道任何属于 CUDA 发行版本身的内核内容器式 API。但是,如果您四处搜索,您可能会发现许多可能有用/有趣的实现。较低级别的库,如 trove可能会为这种用例提供改进的性能。
在您的示例中,似乎每个线程都将维护自己的“堆栈”或“vector ”以跟踪树遍历。 (我将在此处提供的方法取决于没有线程并发访问同一堆栈。如果您需要从多个线程并发访问,方法 here 可能会作为起点。)
如果您知道此类堆栈的最大可能大小是多少,我会建议提前为它分配,无论是内核中每个线程的静态(本地)变量定义,还是动态分配,例如通过 cudaMalloc
。 (出于性能原因,我不建议在内核中使用 malloc
,而且我绝对不建议即时分配/取消分配。)选择哪种分配方法将提供最佳性能可能取决于您的实际测试用例。合并规则(即底层存储方法)对于访问全局指针与访问本地指针有些不同。如果您的线程倾向于在 warp 上均匀地推送或弹出,并且随着您的代码的进行,那么任何一种分配方法都可能提供良好的性能。您可以尝试使用任何一种方法。
这是您在示例中概述的“堆栈”方法的一个相当简单的部分工作示例,假设每个线程的最大堆栈大小先验已知。它绝不是经过全面测试的;我的目的是给你一些想法或一个起点。但是,如果您发现错误,请随时指出,我会尽力解决。
$ cat t1082.cu
const size_t max_items = 256;
template <typename T>
class cu_st{ // simple implementation of "stack" function
T *my_ptr;
size_t n_items;
size_t my_width;
public:
__host__ __device__
cu_st(T *base, size_t id, size_t width=0){
if (width == 0){ // "local" stack allocated
my_ptr = base;
my_width = 1;}
else{ // "global" stack allocated
my_ptr = base + id;
my_width = width;}
n_items = 0;}
__host__ __device__
int push_back(T &item){
if (n_items < max_items){
*my_ptr = item;
my_ptr += my_width;
n_items++;
return 0;}
return -1;}
__host__ __device__
T pop_back(){
if (n_items > 0){
n_items--;
my_ptr -= my_width;}
return *my_ptr;}
__host__ __device__
T back(){
if (n_items > 0){
return *(my_ptr-my_width);}
return *my_ptr;}
__host__ __device__
bool empty(){
return (n_items == 0);}
__host__ __device__
size_t size(){
return n_items;}
__host__ __device__
size_t max_size(){
return max_items;}
};
const size_t nTPB = 256;
const size_t nBLK = 256;
typedef int Node;
__global__
void kernel(Node **g_stack, size_t n)
{
int idx = blockIdx.x * gridDim.x + threadIdx.x;
if(idx >= n) return;
Node *root = NULL;
//method 1 - global stack
cu_st<Node*> to_visit(g_stack, idx, gridDim.x*blockDim.x);
to_visit.push_back(root);
while(!to_visit.empty())
{
Node* n = to_visit.back();
to_visit.pop_back();
}
//method 2 - local stack
Node *l_stack[max_items];
cu_st<Node*> l_to_visit(l_stack, idx);
l_to_visit.push_back(root);
while(!l_to_visit.empty())
{
Node* n = l_to_visit.back();
l_to_visit.pop_back();
}
}
int main(){
Node **d_stack;
cudaMalloc(&d_stack, nTPB*nBLK*max_items*sizeof(Node *));
kernel<<<nBLK, nTPB>>>(d_stack, nTPB*nBLK);
cudaDeviceSynchronize();
}
$ nvcc -o t1082 t1082.cu
$ cuda-memcheck ./t1082
========= CUDA-MEMCHECK
========= ERROR SUMMARY: 0 errors
$
注意事项:
除您在此处看到的内容外,此代码未经过测试。我建议在按原样使用它之前进行更多验证。
正如您在代码中所见,基本上没有错误检查。
这种随机访问通常会很慢,可能与您选择哪种分配方法无关。如果可能,请尽量减少对此类“堆栈”的使用。如果您知道每个线程的堆栈大小非常小,您也可以尝试将此构造与
__shared__
内存分配一起使用。我在这里没有展示的另一种分配方法是给每个线程一个全局分配,但让线程连续推送和弹出,而不是我在这里展示的跨步方式(在算法上是这两种方法的组合我在这里概述了)。这种方法在“统一”情况下肯定会降低性能,但在某些“随机”访问模式下可能会提供更好的性能。
关于c++ - 在cuda内核中创建一个 vector ,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/35385595/