在cuda中使用邻接表分配图？

Question

我想使用邻接表分配一个图，即一个包含 'V' 指针的数组，每个指针都指向一个数组，该数组将具有相邻的顶点（这些顶点将不相等）。所以...

unsigned **d_ptr; 
cudaMalloc(&d_ptr, sizeof(unsigned *)*Count);
  for(int i=0;i<Count;i++)
  {   
     cudaMalloc(&temp, sizeof(unsigned)*outdegree(i));  
  }

我可以将临时指针复制到d_ptr [i]中，但有更好的方法吗？

- hitish

---

对于具有许多数据元素之间指针的复杂数据结构（例如链表），UVA（统一虚拟寻址）提供了代码简化的可能，但可能会牺牲性能（后者实际上取决于您如何从内核访问数据）。除此之外，只需将所有内容从主机深度复制到设备，复制指向何处的结构即可。 - void_ptr

Answer 1

如果你想坚持你所期望的，也就是每个顶点分配一个CUDA内存，那么你的方法是正确的，但同时也是低效和耗时的。这是因为每个CUDA内存分配都有对齐要求。根据这篇文章（以及CUDA文档本身），任何CUDA malloc至少会消耗256字节的全局内存。因此，无论你的顶点的出度多小，使用你的方法保存指针将每个顶点消耗256字节。随着图形规模的增加，这将导致内存迅速耗尽。例如，假设在你的图中每个顶点的出度都等于4。虽然每个顶点所需的大小为4*8=32（假设64位寻址），但每个顶点将消耗256字节，比实际需要的多8倍。请注意，对齐要求可能更高。因此，你提出的方法没有充分利用可用的全局内存。

您的方法也很耗时。在主机或设备代码中进行内存分配和释放都是耗时的操作。每个顶点都需要分配一个内存区域，还需要将temp复制到设备上每个顶点一次。因此，与只分配一次内存区域相比，预计会花费更多时间。

如果您想要用指向设备上顶点的指针填充d_ptr，而不是为每个顶点分配一个缓冲区，可以在主机端先计算所有顶点的outdegree总数，并使用它来分配一个设备缓冲区。

// Allocate one device buffer for all vertices.
unsigned nOutEdges = 0;
for( int i=0; i < Count; i++ )
    nOutEdges += outdegree(i); // outdegree[ i ]??
unsigned* d_out_nbrs;
cudaMalloc( (void**)&d_out_nbrs, sizeof(unsigned) * nOutEdges );

// Collect pointers to the device buffer inside a host buffer.
unsigned** host_array = (unsigned**) malloc( sizeof(unsigned*) * Count );
host_array[ 0 ] = d_out_nbrs;
for( int i=1; i < Count; i++ )
    host_array[ i ] = host_array[ i - 1 ] + outdegree[ i - 1 ];

// Allocate a device buffer and copy collected host buffer into the device buffer.
unsigned **d_ptr; 
cudaMalloc( &d_ptr, sizeof(unsigned *) * Count );
cudaMemcpy( d_ptr, host_array , sizeof(unsigned*) * Count, cudaMemcpyHostToDevice );

- Farzad

---

我其实也使用了你的答案所建议的CSR表示法，但我想在动态图中使用这种表示法，并比较使用CSR和纯邻接表在动态图上的性能。 - hitish

Answer 2

如果你的顶点的出度很小，不会变化太多，并且可以容忍一些浪费空间（请注意，你已经通过双重分配浪费了至少Count指针大小的空间），那么你可能更喜欢像这样进行单一分配：

// compute max outdegree; here I'm assuming you don't already know it
int max_outdegree = -1;

for (i = 0; i < Count; ++i)
  if (outdegree(i) > max_outdegree)
    max_outdegree = outdegree(i);

// alloc d_ptr in one flat array
unsigned *d_ptr;

cudaMalloc(&d_ptr, sizeof(unsigned)*Count*max_outdegree);

...

然后像这样索引到d_ptr：

unsigned ith_adj_vert_of_v = d_ptr[v * max_outdegree + i];  // use and check for -1 to indicate no such vertex

另外，你是否有如此之多的顶点需要使用32位来计数/标识它们？如果你的顶点数量小于2^16，你可以使用uint16_t代替，这样可以将空间利用率减少大约一半。

- jschultz410