使用 CUDA C 进行大型矩阵加法 [关闭]答案

【问题标题】：Large Matrix Addition using CUDA C [closed]使用 CUDA C 进行大型矩阵加法 [关闭]
【发布时间】：2016-03-23 01:21:47
【问题描述】：

我想添加两个大型矩阵 NxN（N 是 2 的倍数）并使用 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]

【问题讨论】：

标签： c matrix cuda parallel-processing

【解决方案1】：

如果您在使用 CUDA 代码时遇到问题，建议您使用 cuda-memcheck 运行您的代码并添加 proper cuda error checking。
这个：
```
 dim3 dimBlock(BLOCK_DIM, BLOCK_DIM);
```

BLOCK_DIM 大于 22 或 32（取决于 GPU 和 CUDA 版本）在 CUDA 中是不合法的。 CUDA 内核限制为每个块 512 或 1024 个线程，这是各个线程块维度的乘积。因此，在您的情况下，BLOCK_DIM*BLOCK_DIM 必须小于 512 或 1024，具体取决于 GPU 和 CUDA 版本。在任何情况下，将BLOCK_DIM 设置为 512 都不起作用。

如果你在堆栈上创建这样的变量：
```
 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((double)N/dimBlock.x),(int)ceil((double)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
$

【讨论】：