Выделение графика с помощью списка смежности в 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], но есть ли лучший способ сделать это?

Answer 1

Если длина ваших вершин мала, она не будет сильно отличаться, и вы в порядке с каким-то потраченным впустую пространством (обратите внимание, что вы уже теряете по крайней мере 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 бита сосчитать/ID торговый центр? Если у вас меньше 2^16, вы можете вместо этого использовать uint16_t и, скорее всего, сократить использование вашего пространства примерно наполовину.

Answer 2

Если вы хотите придерживаться того, что вы хотите, то, что кажется одним распределением памяти CUDA на вершину, ваш подход правильный, но также неэффективен и занимает много времени.

Это неэффективно, поскольку каждое распределение CUDA имеет требование к выравниванию. This post (плюс сама документация CUDA) сообщает, что любой CUDA malloc будет потреблять не менее 256 байт глобальной памяти. В результате, независимо от того, насколько мал outdegree вашей вершины; сберегающие указатели с использованием вашего подхода будут потреблять 256 байт на вершину. Это приведет к быстрому исчерпанию памяти при увеличении размера графика. Например, учтите, что в вашем графике каждая вершина имеет outdegree, равную 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);