cuda内核中的动态扩展数组

Question

我正在尝试在 GPU 上运行 Brandes 算法（基本上是带有一些额外操作和数据结构的 bfs），并且我正在为每个线程分配一个顶点来运行 Brandes。我面临的问题是在我的代码中

我需要存储在 bfs 期间访问的每个顶点的父节点

。在 CPU 实现中，只要我找到一个技术上是动态扩展数组的新父对象，就可以通过创建向量映射并调用 push_back 来轻松实现。我不知道如何在 CUDA 中做到这一点。

这是我需要的功能的示例代码：

    vector<int> distance;               //Initialized to 0
    vector<int> paths;                  //Initialized to 0
    vector<bool> visited;               //Initialized to false
    map <int, vector<int> > parents;    //Parent vector of each key is empty
    queue<int> q;


    // Running bfs from vertex
    q.push(vertex);                     
    while(!q.empty())
    {
        int source = q.front();
        q.pop();

        for(auto neighbour : adjacency_list[source])
        {
            if(!visited[neighbour])
            {
                visited[neighbour] = true;
                q.push(neighbour);
                distance[neighbour] = distance[source] + 1;
            }
            if(distance[neighbour] == distance[source] + 1)
            {
                paths[neighbour] += paths[source];
                parents[neighbour].push_back(source);
            }
        }
    }

    {
        // Use data accumulated above for calculations
        ....
    }

这是我在设备代码中难以实现的行（功能）

父母[邻居].push_back(source);

我的印象：

1. 我可以为每个顶点过度分配（最大程度的图形）父列表，但这会花费我很多未使用的内存

2. 将父关系存储为大小为 2*Edges 的数组中的边，但我需要一个顶点的所有父级一起（连续存储或存储在同一容器中），这在此实现中是不可能的

   李>

3. 我知道 gpu 堆内存，但想不出一种方法来利用它供我使用

4. 最坏的情况：我首先运行一个 bfs 来发现没有。每个顶点的父节点，然后为每个顶点分配适当的内存，然后再次运行品牌。

Answer 1

1. 我认为您的印象 1 可以通过here 所描述的大致实现（每个线程堆栈，预分配）。它有你提到的与过度分配有关的问题。在较新的 GPU 中，数 GB（或更多）的内存很常见，因此如果总内存不是问题，过度分配问题可能不是很严重。

2. 我认为您的印象 2 可以通过here （设备范围的线程安全向量 push_back）所描述的内容大致实现。它有你提到的问题，与结果向量中的结果缺乏排序有关。这些可能会在收集操作完成后通过排序操作来解决。

（4. 听起来您可能已经知道如何制作“最坏情况”印象 4。）

1. 留下印象 3。我们可以使用印象 1 和印象 2 组合的想法，即创建每个线程向量 push_back，但通过内核内 malloc 或 @987654324 使用按需分配@。像这样的内核内存分配非常慢，并且并非没有其自身的问题（例如，您可能必须保留额外的堆空间，内核分配的堆内存不能参与到主机的传输，小分配可能在内存中效率低下用法），但如果没有有关问题维度的更多信息，确实无法判断哪种方法可能是最好的。如果在遍历图时跟踪父节点是一个相对不常见的操作，那么动态分配方法可能不是问题。

这是一个如何创建简单向量（每线程）的示例：

$ cat t376.cu
#include <iostream>
#include <cstdio>

#include <assert.h>
template <typename T>
class cu_vec{  // simple implementation of per-thread "vector"
  const size_t alloc_block_size = 4096; // tuning parameter
  T *my_ptr;
  size_t n_items;
  size_t alloc_blocks;
  public:
    __host__ __device__
    cu_vec(){
      assert(sizeof(T) <= alloc_block_size);
      n_items = 0;
      my_ptr = (T *)new char[alloc_block_size];
      assert(my_ptr != NULL);
      alloc_blocks = 1;}

    __host__ __device__
    cu_vec(size_t sz){
      assert(sizeof(T) <= alloc_block_size);
      n_items = sz;
      alloc_blocks = (n_items*sizeof(T)+alloc_block_size-1)/alloc_block_size;
      my_ptr = (T *)new char[alloc_blocks*alloc_block_size];
      assert(my_ptr != NULL);
      memset(my_ptr, 0, alloc_blocks*alloc_block_size);}

    __host__ __device__
    ~cu_vec(){
      if (my_ptr != NULL) delete[] my_ptr;
      }

    __host__ __device__
    void push_back(T const &item){ // first test if we can just store new item
      if ((n_items+1)*sizeof(T) > alloc_blocks*alloc_block_size){
        T *temp = (T *)new char[(alloc_blocks+1)*alloc_block_size];
        assert(temp != NULL);
        memcpy(temp, my_ptr, alloc_blocks*alloc_block_size);
        delete[] my_ptr;
        my_ptr = temp;
        alloc_blocks++;}
      my_ptr[n_items] = item;
      n_items++;}

    __host__ __device__
    size_t size(){
      return n_items;}

    __host__ __device__
    void clear(){
      n_items = 0;}

    __host__ __device__
    T& operator[](size_t idx){
      assert(idx < n_items);
      return my_ptr[idx];}

    __host__ __device__
    T& pop_back(){
      if (n_items > 0){
        n_items--;}
      return my_ptr[n_items];}

    __host__ __device__
    T* data(){
      return my_ptr;}

    __host__ __device__
    size_t storage_ratio(){
      return alloc_block_size/sizeof(T);}
};

struct ss
{
   unsigned x;
   float y;
};

__global__ void test(){

  cu_vec<ss> my_vec;
  ss temp = {threadIdx.x, 2.0f};
  my_vec.push_back(temp);
  assert(my_vec.size() == 1);
  assert(my_vec.storage_ratio() >= 1);
  ss temp2 = my_vec[0];
  printf("threadIdx.x: %u, ss.x: %u, ss.y: %f\n", threadIdx.x, temp2.x, temp2.y);
  temp.y = 3.0f;
  my_vec[0].x = temp.x;
  my_vec[0].y = temp.y;
  ss temp3 = my_vec.pop_back();
  printf("threadIdx.x: %u, ss.x: %u, ss.y: %f\n", threadIdx.x, temp3.x, temp3.y);
  my_vec.clear();
  temp.x = 0;
  for (int i = 0; i < 10000; i++){
    my_vec.push_back(temp);
    temp.x++;}
  temp.x--;
  for (int i = 0; i < 10000; i++) {
    assert(my_vec.pop_back().x == temp.x);
    temp.x--;}
  cu_vec<ss> my_vec2(2);
  assert(my_vec2[1].x == 0);
  assert(my_vec2[1].y == 0.0f);
}

int main(){

  //default heap space is 8MB, if needed reserve more with:
  cudaDeviceSetLimit(cudaLimitMallocHeapSize, (1048576*32));
  test<<<1, 4>>>();
  cudaDeviceSynchronize();
}
$ nvcc -std=c++11 -o t376 t376.cu
$ cuda-memcheck ./t376
========= CUDA-MEMCHECK
threadIdx.x: 0, ss.x: 0, ss.y: 2.000000
threadIdx.x: 1, ss.x: 1, ss.y: 2.000000
threadIdx.x: 2, ss.x: 2, ss.y: 2.000000
threadIdx.x: 3, ss.x: 3, ss.y: 2.000000
threadIdx.x: 0, ss.x: 0, ss.y: 3.000000
threadIdx.x: 1, ss.x: 1, ss.y: 3.000000
threadIdx.x: 2, ss.x: 2, ss.y: 3.000000
threadIdx.x: 3, ss.x: 3, ss.y: 3.000000
========= ERROR SUMMARY: 0 errors
$

与您在此处看到的相比，该代码尚未经过测试。