Как распределить память и скопировать 2D-массивы между CPU / GPU в CUDA, не выравнивая их?

Question

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

Обратите внимание, что высота!= Ширина для массивов, это то, что еще больше смущает меня, если это возможно, так как я всегда борюсь с выбором размера сетки.

Я подумал об их сведении, но я действительно хочу, чтобы это работало таким образом.

Это то, как далеко я продвинулся в своем собственном исследовании.

__global__ void myKernel(int *firstArray, int *secondArray, int rows, int columns) {
    int row = blockIdx.x * blockDim.x + threadIdx.x;
    int column = blockIdx.y * blockDim.y + threadIdx.y;

    if (row >= rows || column >= columns)
        return;

    // Do something with the arrays like you would on a CPU, like:
    firstArray[row][column] = row * 2;
    secondArray[row[column] = row * 3;  
}


int main() {
    int rows = 300, columns = 200;
    int h_firstArray[rows][columns], h_secondArray[rows][columns];
    int *d_firstArray[rows][columns], *d_secondArray[rows][columns];

    // populate h_ arrays (Can do this bit myself)

    // Allocate memory on device, no idea how to do for 2D arrays.
    // Do memcopies to GPU, no idea how to do for 2D arrays.

    dim3 block(rows,columns);
    dim3 grid (1,1);
    myKernel<<<grid,block>>>(d_firstArray, d_secondArray, rows, columns);

    // Do memcopies back to host, no idea how to do for 2D arrays.

    cudaFree(d_firstArray);
    cudaFree(d_secondArray);

    return 0;
}

РЕДАКТИРОВАТЬ: меня спросили, будет ли ширина массива будет известна во время компиляции в задачах, которые я бы попытался решить. Вы можете предположить, что это так, как я интересуюсь прежде всего этой конкретной ситуацией на данный момент.

Answer 1

В общем случае (размеры массива неизвестны до времени выполнения), обработка двукратно подписанного доступа в коде устройства CUDA требует массива указателей, как это происходит в коде хоста. C и C++ обрабатывают каждый нижний индекс как разыменование указателя, чтобы достичь конечного местоположения в "2D массиве".

Двойной указатель / двукратный доступ к коду устройства в общем случае уже описан в каноническом ответе, связанном со страницей информации тега cuda. В этом есть несколько недостатков, которые освещены в этом ответе, поэтому я не буду их здесь повторять.

Однако, если ширина массива известна во время компиляции (высота массива может быть динамической - то есть определяется во время выполнения), тогда мы можем использовать компилятор и механизмы типизации языка, чтобы позволить нам обойти большинство недостатков. Ваш код демонстрирует несколько других неправильных шаблонов для использования CUDA и / или C/C++:

1. Передача элемента для двукратно подписанного доступа к функции C или C++ не может быть выполнена с простым типом одного указателя, таким как int *firstarray

2. Выделение больших массивов хостов с помощью стековых механизмов:

   int h_firstArray[rows][columns], h_secondArray[rows][columns];
   

   часто проблематично в C и C++. Это переменные, основанные на стеке, и они часто выходят за пределы стека, если они достаточно велики.

3. Блоки потоков CUDA ограничены до 1024 потоков. Поэтому такой размер нити блока:

   dim3 block(rows,columns);
   

   не будет работать, за исключением очень маленьких размеров rows а также columns (произведение должно быть меньше или равно 1024).

4. При объявлении переменных-указателей для массива устройств в CUDA почти никогда не правильно создавать массивы указателей:

   int *d_firstArray[rows][columns], *d_secondArray[rows][columns];
   

   мы также не выделяем место на хосте, а затем "перераспределяем" эти указатели для использования устройством.

Ниже приведен работающий пример с указанными выше элементами, демонстрирующий вышеупомянутый метод, в котором ширина массива известна во время выполнения:

$ cat t50.cu
#include <stdio.h>

const int array_width = 200;

typedef int my_arr[array_width];

__global__ void myKernel(my_arr *firstArray, my_arr *secondArray, int rows, int columns) {
    int column = blockIdx.x * blockDim.x + threadIdx.x;
    int row = blockIdx.y * blockDim.y + threadIdx.y;

    if (row >= rows || column >= columns)
        return;

    // Do something with the arrays like you would on a CPU, like:
    firstArray[row][column] = row * 2;
    secondArray[row][column] = row * 3;
}


int main() {
    int rows = 300, columns = array_width;
    my_arr *h_firstArray, *h_secondArray;
    my_arr *d_firstArray, *d_secondArray;
    size_t dsize = rows*columns*sizeof(int);
    h_firstArray = (my_arr *)malloc(dsize);
    h_secondArray = (my_arr *)malloc(dsize);
    // populate h_ arrays
    memset(h_firstArray, 0, dsize);
    memset(h_secondArray, 0, dsize);

    // Allocate memory on device
    cudaMalloc(&d_firstArray, dsize);
    cudaMalloc(&d_secondArray, dsize);
    // Do memcopies to GPU
    cudaMemcpy(d_firstArray, h_firstArray, dsize, cudaMemcpyHostToDevice);
    cudaMemcpy(d_secondArray, h_secondArray, dsize, cudaMemcpyHostToDevice);

    dim3 block(32,32);
    dim3 grid ((columns+block.x-1)/block.x,(rows+block.y-1)/block.y);
    myKernel<<<grid,block>>>(d_firstArray, d_secondArray, rows, columns);

    // Do memcopies back to host
    cudaMemcpy(h_firstArray, d_firstArray, dsize, cudaMemcpyDeviceToHost);
    cudaMemcpy(h_secondArray, d_secondArray, dsize, cudaMemcpyDeviceToHost);
    // validate
    if (cudaGetLastError() != cudaSuccess) {printf("cuda error\n"); return -1;}
    for (int i = 0; i < rows; i++)
      for (int j = 0; j < columns; j++){
        if (h_firstArray[i][j] != i*2) {printf("first mismatch at %d,%d, was: %d, should be: %d\n", i,j,h_firstArray[i][j], i*2); return -1;}
        if (h_secondArray[i][j] != i*3) {printf("second mismatch at %d,%d, was: %d, should be: %d\n", i,j,h_secondArray[i][j], i*3); return -1;}}

    printf("success!\n");


    cudaFree(d_firstArray);
    cudaFree(d_secondArray);

    return 0;
}
$ nvcc -arch=sm_61 -o t50 t50.cu
$ cuda-memcheck ./t50
========= CUDA-MEMCHECK
success!
========= ERROR SUMMARY: 0 errors
$

Я изменил смысл индексации вашего ядра (x,y), чтобы помочь объединить глобальный доступ к памяти. Мы видим, что с таким типом создания типов мы можем использовать функции компилятора и языка, чтобы в итоге получить код, который обеспечивает двукратно подписанный доступ как к коду хоста, так и к коду устройства, в то же время позволяя выполнять операции CUDA (например, cudaMemcpy), как будто мы имеем дело с массивами с одним указателем (например, "плоскими").