Обрабатывать массив массивов (различной длины) для ядра CUDA

Question

У меня есть массив массивов (с разной длиной) в C, которые я хотел бы обработать в «ядре CUDA».

const int N_ARRAYS = 1000;
int *arrayOfArrays[N_ARRAYS];
int arr1[3] = {1,2,3};
int arr2[2] = {1,4};
int arr3[4] = {1,5,3,6};
//....
int arr1000[5] = {9,9,9,10,10};

arrayOfArrays[0] = arr1;
arrayOfArrays[1] = arr2;
arrayOfArrays[2] = arr3;
//...
arrayOfArrays[1000] = arr1000;

Я нашел этот пост: CUDA выделяет массив массивов который дал хорошую идею о том, как это может работать. Но, честно говоря, я не получил его на работу.

Я снова подведу итоги:

1. Вы должны выделить указатели для памяти хоста,
2. затем выделите память устройства для каждого массива
3. и сохранить его указатель в памяти хоста.
4. Затем выделите память для хранения указателей на устройстве
5. , а затем скопируйте память хоста в память устройства.

Вот то, что я пробовал до сих пор, основываясь на данном ответе. В целях иллюстрации я покажу его с N_ARRAYS = 3, но на самом деле его путь выше (> 1000).

int main(){
    const int N_ARRAYS = 3;
    int *arrayOfArrays[N_ARRAYS];
    int arr1[1] = {1,2,3};
    int arr2[2] = {1,4};
    int arr3[3] = {1,5,3};

    arrayOfArrays[0] = arr1;
    arrayOfArrays[1] = arr2;
    arrayOfArrays[2] = arr3;

    // 1) You have to allocate the pointers to a host memory, 
    //void *h_array = malloc(sizeof(void*) * N_ARRAYS); // i use arrayOfArrays instead
    for(int i = 0; i < N_ARRAYS; i++){
        //2) then allocate device memory for each array
        cudaMalloc(&arrayOfArrays[i], i * sizeof(void*));
    }

    // 4) Allocate the memmory for storing the pointers into the device to *d_array
    void *d_array = cudaMalloc(sizeof(void*) * N_ARRAYS);

    // 5) Copy arrayOfArrays to d_array of size sizeof(void*) * N_ARRAYS from Host to device
    cudaMemcpy(d_array, arrayOfArrays, sizeof(void*) * N_ARRAYS, cudaMemcpyHostToDevice);

    // Call kernel
    multi_array_kernel<1,1>(N_ARRAYS, d_array);
    cudaThreadSynchronize();

    for(int i = 0; i < N_ARRAYS; i++){
        cudaFree(arrayOfArrays[i]); //host not device memory
        //TODO: check error
    }
    cudaFree(d_array);
    free(arrayOfArrays);
}

и ядро:

__global__ void multi_array_kernel( int N, void** arrays ){
    int nr;
    int sum = 0;
    for(nr = 0; nr < N; nr++){
        if(arrays[nr+0] == arrays[nr-1+0]) sum +=1; // some sample calc.
    }

}

Answer 1

1. Размеры вашего массива не имеют смысла:

   int arr1[1] = {1,2,3};
            ^
            array length of 1 integer storage
   

   Вы не можете инициализировать массив длиной 1 с 3 целыми числами. Я думаю, что это приведет к ошибке компиляции.

2. Вы уже используете arrayOfArrays[] для хранения указателей на массивы хостов, нам понадобится другая аналогичная переменная для хранения массива соответствующих указателей устройств. Теперь у вас есть просто перезаписать предыдущие значения там, когда вы используете его в cudaMalloc

3. Ваша операция cudaMalloc в цикле for настроена неправильно. Это должно выделять место для каждого массива целых чисел (например, arr1 и т. Д.). Поэтому нам нужно написать что-то вроде этого:

   cudaMalloc(&darrayOfArrays[i], arr_len[i] * sizeof(arr[0]));
   

4. Нигде вы не копируете содержимое arr1 на устройство (и аналогично для arr2, arr3). Итак, вы пропустили шаг. Мы можем легко сделать это в вашем первом цикле for, например,

   cudaMemcpy(darrayOfArrays[i], arrayOfArrays[i], arr_len[i]*sizeof(int), cudaMemcpyHostToDevice); // copy contents of each array to device
   

5. Это не так, как cudaMalloc работает:

     void *d_array = cudaMalloc(sizeof(void*) * N_ARRAYS);
   

   Похоже, вы уже знаете, как это работает, исходя из вашего другого использования.

6. Ваш код ядра будет индексироваться за пределами, когда nr равен 0:

   if(arrays[nr+0] == arrays[nr-1+0]) sum +=1;
                             ^
                             out of bounds when nr = 0
   

   Более того, я предполагаю, что весь смысл этого упражнения в том, чтобы иметь возможность выполнять индексирование массивов с двойной подпиской (в противном случае код вашего ядра не имеет особого смысла, если только вы не собираетесь сравнивать указатели).

7. Это не синтаксис запуска ядра:

   multi_array_kernel<1,1>(N_ARRAYS, d_array);
   

8. Если вы хотите обрабатывать массивы различной длины, нам нужно где-то отслеживать длины массивов и использовать их во время выделения / копирования.

Вот фиксированный пример с указанными выше пунктами:

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

__global__ void multi_array_kernel( int N, int **arrays ){
    int nr;
    for(nr = 1; nr < N; nr++){
        if(arrays[nr][0] == arrays[nr-1][0]) printf("match at index: %d  to index: %d\n", nr, nr-1); // some sample calc.
    }

}

int main(){
    const int N_ARRAYS = 3;
    int *arrayOfArrays[N_ARRAYS];
    int *darrayOfArrays[N_ARRAYS];
    int arr1[3] = {1,2,3};
    int arr2[2] = {1,4};
    int arr3[3] = {1,5,3};
    int **d_array;
    int arr_len[N_ARRAYS] = {3, 2, 3};
    arrayOfArrays[0] = arr1;
    arrayOfArrays[1] = arr2;
    arrayOfArrays[2] = arr3;

    // 1) You have to allocate the pointers to a host memory,
    //void *h_array = malloc(sizeof(void*) * N_ARRAYS); // i use arrayOfArrays instead
    for(int i = 0; i < N_ARRAYS; i++){
        //2) then allocate device memory for each array
        cudaMalloc(&(darrayOfArrays[i]), arr_len[i] * sizeof(int));
        cudaMemcpy(darrayOfArrays[i], arrayOfArrays[i], arr_len[i]*sizeof(int), cudaMemcpyHostToDevice); // copy contents of each array to device
    }

    // 4) Allocate the memmory for storing the pointers into the device to *d_array
    cudaMalloc(&d_array, sizeof(int*) * N_ARRAYS);

    // 5) Copy arrayOfArrays to d_array of size sizeof(void*) * N_ARRAYS from Host to device
    cudaMemcpy(d_array, darrayOfArrays, sizeof(int*) * N_ARRAYS, cudaMemcpyHostToDevice);

    // Call kernel
    multi_array_kernel<<<1,1>>>(N_ARRAYS, d_array);
    cudaDeviceSynchronize();

    for(int i = 0; i < N_ARRAYS; i++){
        cudaFree(darrayOfArrays[i]); //host not device memory
        //TODO: check error
    }
    cudaFree(d_array);
    // free(arrayOfArrays);
    printf("%s\n", cudaGetErrorString(cudaGetLastError()));
}
$ nvcc -o t103 t103.cu
$ ./t103
match at index: 1  to index: 0
match at index: 2  to index: 1
no error
$

Обратите внимание, что из-за "сложности" вышесказанного, обычный совет - просто "сплющить" хранилище. Если вы выполните поиск по тегу cuda в строке поиска в верхней части этой страницы с помощью [cuda] flatten, вы найдете множество определений и примеров.