在 cuda 内核中创建向量

Creating a vector in cuda kernel

本文关键字：创建 向量 内核 cuda      更新时间：2023-10-16

我需要创建某种堆栈来爬行 cuda 内核中的树。 我以为我可以使用thrust：:d evice_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 设备代码中不可用。

我不知道有任何类似内核容器的 API 是 CUDA 发行版本身的一部分。 但是，如果您四处搜索，您可能会发现数十种可能有用/有趣的实现。 像trove这样的较低级别的库可能会为这种用例提供更高的性能。

在您的示例中，似乎每个线程都会维护自己的"堆栈"或"向量"来跟踪树遍历。 (我将在这里提供的方法取决于没有线程同时访问同一堆栈。 如果需要从多个线程进行并发访问，此处的方法可能会作为起点。

如果您知道这种堆栈的最大可能大小是多少，我建议您提前为其分配，要么是内核中每个线程的静态(局部(变量定义，要么是动态分配，例如通过 cudaMalloc . (出于性能原因，我不会建议内核内malloc，并且出于性能原因，我绝对不会建议即时分配/解除分配。 选择哪种分配方法将提供最大的性能可能取决于您的实际测试用例。 对于访问全局指针和访问本地指针，合并规则(即底层存储方法(略有不同。 如果您的线程倾向于在扭曲中均匀地推动或弹出，并且随着代码的进行，那么任何一种分配方法都可以提供良好的性能。 您可以尝试任一方法。

下面是您在示例中概述的"堆栈"方法的一个相当简单的部分工作示例，假设每个线程的最大堆栈大小是先验已知的。 它绝不经过全面测试;我的目的是给你一些想法或起点。 但是，如果您发现错误，请随时指出它们，我将尝试解决它们。

$ 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
$

笔记：

  1. 除了您在此处看到的内容之外，此代码未经过测试。 我建议在按原样使用它之前进行更多验证。

  2. 正如您在代码中看到的，基本上没有错误检查。

  3. 这种随机访问通常往往很慢，可能无论您选择哪种分配方法。 如果可能，请尽量减少使用此类"堆栈"。 如果您知道每个线程的堆栈大小非常小，您还可以尝试将此构造与__shared__内存分配一起使用。

    4. 我在这里

  4. 没有演示的另一种分配方法是为每个线程提供全局分配，但让线程连续推送和弹出，而不是以我在这里展示的跨步方式(算法上是我在这里概述的两种方法的组合(。 这种方法肯定会降低"统一"情况下的性能，但在某些"随机"访问模式中可能会提供更好的性能。

相关文章：