2D массив на CUDA

Question

Я хочу динамически распределить глобальный массив 2D в CUDA. Как мне этого добиться?

В моем основном я называю мой Kernel в цикле. Но прежде чем вызывать ядро, мне нужно выделить немного памяти на GPU. После вызова ядра одно целое число отправляется из GPU в CPU, чтобы сообщить, решена ли проблема или нет.
Если проблема не будет решена, я не освобожу старую память, так как в этом есть дальнейшая необходимость, и я должен выделить новую память для GPU и снова вызвать ядро.

показывается судокод:

int n=0,i=0;
while(n==0)
{
    //allocate 2d memory for MEM[i++] 
    //call kernel(MEM,i)
    // get n from kernel       
}


__global__ void kernerl(Mem,int i)
{
    Mem[0][5]=1;
    Mem[1][0]=Mem[0][5]+23;//can use this when MEM[1] is allocated before kernel call
}

Есть предложения? Спасибо.

Answer 1

Два вступительных комментария - использование динамически размещенного 2D-массива - плохая идея в CUDA, а повторное выделение памяти в цикле также не является хорошей идеей. Оба несут ненужные штрафы за производительность.

Для кода хоста, что-то вроде этого:

size_t allocsize = 16000 * sizeof(float);
int n_allocations = 16;
float * dpointer
cudaMalloc((void **)&dpointer, n_allocations * size_t(allocsize));

float * dcurrent = dpointer;
int n = 0;
for(int i=0; ((n==0) && (i<n_allocations)); i++, dcurrent+=allocsize) {

    // whatever you do before the kernel

    kernel <<< gridsize,blocksize >>> (dcurrent,.....);

    // whatever you do after the kernel

}

является предпочтительным. Здесь вы вызываете cudaMalloc только один раз и передаете смещения в выделение, что делает выделение памяти и управление свободным внутри цикла. Структура цикла также означает, что вы не можете работать бесконечно и исчерпать всю память графического процессора.

В самом вопросе о двумерном массиве есть две причины, по которым это плохая идея. Во-первых, для выделения из двумерного массива с N строками требуется (N + 1) вызовов cudaMalloc и копирование памяти хост-устройства, что является медленным и уродливым. Во-вторых, внутри кода ядра, чтобы получить ваши данные, графический процессор должен выполнить два глобальных чтения памяти, одно для перенаправления указателя, чтобы получить адрес строки, а затем одно для извлечения данных из строки. Это намного медленнее, чем эта альтернатива:

#define idx(i,j,lda) ( (j) + ((i)*(lda)) )
__global__ void kernerl(float * Mem, int lda, ....)
{
    Mem[idx(0,5,lda)]=1; // MemMem[0][5]=1;
}

, который использует индексацию в одномерном распределении. В GPU транзакции с памятью очень дороги, но FLOPS и IOPS дешевы. Единственное целочисленное умножение - самый эффективный способ сделать это. Если вам нужен доступ к результатам предыдущего вызова ядра, просто передайте смещение к предыдущим результатам и используйте два указателя внутри ядра, примерно так:

__global__ void kernel(float *Mem, int lda, int this, int previous)
{
   float * Mem0 = Mem + this;
   float * Mem1 = Mem + previous;

}

Эффективные программы с распределенной памятью (а CUDA на самом деле является типом программирования с распределенной памятью) через некоторое время начинают выглядеть как Fortran, но это цена, которую вы платите за мобильность, прозрачность и эффективность.

Надеюсь, это помогло.

Answer 2

Ну, вы можете сделать это так же, как это было бы на процессоре.

unsigned xSize = 666, ySize = 666;
int **h_ptr = (int**)malloc(sizeof(int*) * xSize);
int **d_ptr = NULL;
cudaMalloc( &d_ptr, xSize );
for(unsigned i = 0; i < xSize; ++i)
{
    cudaMalloc( &h_ptr[i], ySize );
}
cudaMemcpy( &d_ptr, &h_ptr, sizeof(int*) * xSize, cudaMemcpyHostToDevice );
free( h_ptr );

... и бесплатно примерно

int **h_ptr = (int**)malloc(sizeof(int*) * xSize);
cudaMemcpy( &h_ptr, &d_ptr, sizeof(int*) * xSize, cudaMemcpyDeviceToHost );
for(unsigned i = 0; i < xSize; ++i )
{
    cudaFree( h_ptr[i] );
}
cudaFree( d_ptr );
free( h_ptr );

Но вы должны помнить, чтокаждый доступ к ячейке этого массива будет включать двойной доступ к глобальной памяти графического процессора.Из-за этого доступ к памяти будет в два раза медленнее, чем с массивом 1d.

Answer 3

РЕДАКТИРОВАНИЕ:
Я пытался помочь вам привести пример, в котором, сгладив массив, вы можете достичь того же результата, но товарищи сказали, что это не то, что вы просите.
Итак, есть еще один пост здесь , рассказывающий о том, как вы можете выделить 2d массивы в CUDA.