c++ - CUDA 数组 vector 乘法

Question

您好，我正在迈出 CUDA 技术的第一步，但我认为我没有做对。

我正在尝试将二维数组乘以 vector ，但有些东西不起作用

这是我试图找出的代码:

#include <stdio.h>
#include <stdlib.h>

#define N 2

__global__ void Multiply(int A[N][N], int B[N], int C[N]){
           int i = threadIdx.x;
           int j = threadIdx.y;

           int sum = A[i][j] * B[j];
           C[i]= sum;
           printf("%d,%d ", sum, C[i]);
}

int main(){

int A[N][N] ={  {1,1},
                {1,1}
            };
int B[N] = {4,6};
int C[N] = {0,0};    
int (*aA)[N], (*aB), (*aC);

cudaMalloc((void**)&aA, (N*N)*sizeof(int));
cudaMalloc((void**)&aB, (N)*sizeof(int));
cudaMalloc((void**)&aC, (N)*sizeof(int));

cudaMemcpy(aA, A, (N*N)*sizeof(int), cudaMemcpyHostToDevice);
cudaMemcpy(aB, B, (N)*sizeof(int), cudaMemcpyHostToDevice);
cudaMemcpy(aC, C, (N)*sizeof(int), cudaMemcpyHostToDevice);

int numBlocks = 1;
dim3 threadsPerBlock(N,N);
Multiply<<<numBlocks,threadsPerBlock>>>(aA,aB,aC);

cudaMemcpy(C, aC, (N)*sizeof(int), cudaMemcpyDeviceToHost);


cudaFree(aA); 
cudaFree(aB); 
cudaFree(aC);

printf("\n");
system("pause");

}

在这种情况下，输出是:4,6 4,6 6,6 6,6 所以基本上总和 i​​ 给出了正确的值，但 C[i] 总是返回 6，尽管有分配给它的总和值。

我做错了什么？

Answer 1

1. 任何时候当您在使用 CUDA 代码时遇到问题，最好使用 proper cuda error checking并使用 cuda-memcheck 运行您的代码。这只是我所做的样板声明。它实际上不会出现您在本例中显示的代码的问题。

2. 正如已在现已删除的答案中指出的那样，您实际上并没有将任何东西加在一起。即使您有一个名为 sum 的变量，它实际上并不是任何内容的总和，并且您的内核代码中没有 + 或求和操作。您不是在编写一个将任何东西加在一起的内核。

3. 为了产生正确的结果，您的内核依赖于让多个线程协作更新单个位置 (C[i])。但是，这需要线程之间进行一些协调。如果没有任何协调，线程之间将处于竞争状态，结果将不可预测。我们可以使用 a parallel reduction 解决这个问题，将来自每个单独线程的部分产品加在一起，或者为了简单起见，我们可以使用 atomicAdd operation ，这将强制线程逐个更新(添加到)C[i]，因此它们不会相互踩踏。因此，使用 atomicAdd 还提供了您的内核所缺少的必要加法 (+) 操作。

这是解决了第 2 项和第 3 项的工作代码。您可以使用 cuda-memcheck 运行它来验证行为的正确性，即使它没有明确的错误检查:

 $ cat t1037.cu
#include <stdio.h>
#include <stdlib.h>

#define N 2

__global__ void Multiply(int A[N][N], int B[N], int C[N]){
           int i = threadIdx.x;
           int j = threadIdx.y;

           int product = A[i][j] * B[j];
           atomicAdd(C+i, product);
  //         printf("%d,%d ", product, C[i]);
}

int main(){

int A[N][N] ={  {1,1},
                {1,1}
            };
int B[N] = {4,6};
int C[N] = {0,0};
int (*aA)[N], (*aB), (*aC), i;

cudaMalloc((void**)&aA, (N*N)*sizeof(int));
cudaMalloc((void**)&aB, (N)*sizeof(int));
cudaMalloc((void**)&aC, (N)*sizeof(int));

cudaMemcpy(aA, A, (N*N)*sizeof(int), cudaMemcpyHostToDevice);
cudaMemcpy(aB, B, (N)*sizeof(int), cudaMemcpyHostToDevice);
cudaMemcpy(aC, C, (N)*sizeof(int), cudaMemcpyHostToDevice);

int numBlocks = 1;
dim3 threadsPerBlock(N,N);
Multiply<<<numBlocks,threadsPerBlock>>>(aA,aB,aC);

cudaMemcpy(C, aC, (N)*sizeof(int), cudaMemcpyDeviceToHost);

for (i=0; i<N; i++){
  printf("C[%d] = %d\n",i,C[i]);
  }
cudaFree(aA);
cudaFree(aB);
cudaFree(aC);

printf("\n");

}
$ nvcc -o t1037 t1037.cu
$ cuda-memcheck ./t1037
========= CUDA-MEMCHECK
C[0] = 10
C[1] = 10

========= ERROR SUMMARY: 0 errors
$