使用cuda中的邻接列表来分配图

Allocating a graph using adjacency list in cuda?

本文关键字:列表 分配 cuda 使用      更新时间:2023-10-16

我想使用邻接列表分配一个图,即"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],但有更好的方法吗?

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

它是低效的,因为每个CUDA分配都有一个对齐要求。这篇文章(加上CUDA文档本身)告诉任何CUDAmalloc都将消耗至少256字节的全局内存。因此,无论顶点的outdegree有多小;使用您的方法保存指针将消耗每个顶点256个字节。随着图形大小的增加,这将导致内存很快耗尽。例如,考虑图中每个顶点的outdegree等于4。虽然假设64位寻址,每个顶点所需的大小为4*8=32,但每个顶点将消耗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 );

如果你的顶点的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,并可能将你的空间使用量减少大约一半。