Большая матричная добавка с использованием CUDA C

Question

Я хочу добавить две большие матрицы NxN (N кратно двум) и распараллелить программу с помощью Cuda C. Мне удалось запустить программу с матрицами размером 512x512. Но если я выйду за пределы этого (например, 1024x1024), то это не сработает. Я считаю, что проблема может заключаться в том, что CUDA может запускать максимум 512 потоков на блок (?). Поэтому мой вопрос заключается в том, как изменить программу, чтобы я мог создавать матрицы любого размера.

CUDA ядра

__global__ void parMatrixAdd_kernel(float *a, float *b, float *c, int N) { 

    int col = threadIdx.x + blockIdx.x * blockDim.x; 
    int row = threadIdx.y + blockIdx.y * blockDim.y; 
    int index = col + row * N; 
    if (col < N && row < N) { 
     c[index] = a[index] + b[index]; 
    } 
} 

блок и сетки размеры:

//BLOCK_DIM is 512, N works upto 512, but not beyond 
    dim3 dimBlock(BLOCK_DIM, BLOCK_DIM); 
    dim3 dimGrid((int)ceil(N/dimBlock.x),(int)ceil(N/dimBlock.y)); 

массивы: matrix1[N][N]matrix2[N][N]

Answer 1

1. Каждый раз, когда у вас возникли проблемы с кодом CUDA, то целесообразно для запуска вашего кода с cuda-memcheck, а также добавить proper cuda error checking.

2. Это:

   dim3 dimBlock(BLOCK_DIM, BLOCK_DIM); 
   

   не является легальным в CUDA для BLOCK_DIM больше, чем 22 или 32 (в зависимости от GPU и CUDA версии). Ядра CUDA ограничены 512 или 1024 потоками на блок, что составляет продукт отдельных размеров резьбовых блоков. Поэтому в вашем случае BLOCK_DIM * BLOCK_DIM должен быть меньше 512 или 1024 в зависимости от версии графического процессора и CUDA. Установка BLOCK_DIM на 512 не может работать, в любом случае.

3. Если вы создаете переменную, как это на стеке:

   float matrix1[N][N]; 
   

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

Следующий код (построенный вокруг частей, которые вы показали), кажется, работает на меня и реализует вышеуказанные изменения. Я пропустил правильную проверку ошибок cuda для краткости презентации, но я рекомендую использовать ее, если у вас возникли проблемы с кодом CUDA. Вместо этого я запускаю его с помощью cuda-memcheck:

$ cat t1002.cu 
#include <stdio.h> 
#include <math.h> 

const size_t BLOCK_DIM = 16; 
const size_t MY_N = 2048; 

const float tval1 = 1.0f; 
const float tval2 = 2.0f; 

__global__ void parMatrixAdd_kernel(float *a, float *b, float *c, int N) { 

    int col = threadIdx.x + blockIdx.x * blockDim.x; 
    int row = threadIdx.y + blockIdx.y * blockDim.y; 
    int index = col + row * N; 
    if (col < N && row < N) { 
     c[index] = a[index] + b[index]; 
    } 
} 

typedef float my_mat[MY_N]; 

int main(){ 

    my_mat *A, *B, *C; 
    const size_t N = MY_N; 
    size_t dsize = N*N*sizeof(float); 
    A = (my_mat *)malloc(dsize); 
    B = (my_mat *)malloc(dsize); 
    C = (my_mat *)malloc(dsize); 

    for (int i = 0; i < N; i++) 
    for (int j = 0; j < N; j++) { 
     A[i][j] = tval1; 
     B[i][j] = tval2; 
     C[i][j] = 0.0f;} 

    float *d_A, *d_B, *d_C; 
    cudaMalloc(&d_A, dsize); 
    cudaMalloc(&d_B, dsize); 
    cudaMalloc(&d_C, dsize); 

    cudaMemcpy(d_A, A, dsize, cudaMemcpyHostToDevice); 
    cudaMemcpy(d_B, B, dsize, cudaMemcpyHostToDevice); 

    dim3 dimBlock(BLOCK_DIM, BLOCK_DIM); 
    dim3 dimGrid((int)ceil(N/dimBlock.x),(int)ceil(N/dimBlock.y)); 

    parMatrixAdd_kernel<<<dimGrid, dimBlock>>>(d_A, d_B, d_C, N); 
    cudaMemcpy(C, d_C, dsize, cudaMemcpyDeviceToHost); 
    for (int i = 0; i < N; i++) 
    for (int j = 0; j < N; j++) 
     if (C[i][j] != (tval1+tval2)) {printf("mismatch at %d,%d was: %f, should be %f\n", i, j, C[i][j], (tval1+tval2)); return 1;} 
    printf("Success!\n"); 
    return 0; 
} 
$ nvcc -o t1002 t1002.cu 
$ cuda-memcheck ./t1002 
========= CUDA-MEMCHECK 
Success! 
========= ERROR SUMMARY: 0 errors 
$