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

【问题标题】：Allocating a graph using adjacency list in cuda?在cuda中使用邻接列表分配图？
【发布时间】：2015-03-02 13:18:23
【问题描述】：

我想使用邻接表分配一个图形，即“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] 但有更好的方法吗？

【问题讨论】：

对于在数据元素之间具有大量指针的复杂数据结构（例如链表），UVA（统一虚拟寻址）提供了代码简化，但可能会降低性能（后者实际上取决于您如何访问来自内核的数据）。除此之外，只需将所有内容从主机深度复制到设备，复制哪些点的结构。

标签： c++ c cuda

【解决方案1】：

如果你想坚持你想要的，这似乎是每个顶点分配一个 CUDA 内存，你的方法是正确的，但也是低效且耗时的。

这是低效的，因为每个 CUDA 分配都有对齐要求。 This post（加上 CUDA 文档本身）告诉任何 CUDA malloc 将消耗至少 256 字节的全局内存。结果，无论您的顶点的outdegree 多么小；使用您的方法保存指针将消耗每个顶点 256 个字节。随着图形大小的增加，这将导致内存很快耗尽。例如，考虑在您的图中，每个顶点的 outdegree 等于 4。虽然每个顶点的大小 required 是 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 );

【讨论】：

我实际上也在使用 CSR 表示，正如您的回答所建议的那样，但我想将此表示用于动态图，并比较使用 CSR 与纯邻接表用于动态图的性能。

【解决方案2】：

如果您的顶点的 outdegree() 很小，不会有太大的变化，并且您可以接受一些浪费的空间（请注意，您已经通过双重分配至少浪费了 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 代替，并且可能会将您的空间使用量减少大约一半。

【讨论】：