CUDA 编程

Question

我是 CUDA 的新手。我有一个关于一个简单程序的问题，希望有人能注意到我的错误。

__global__ void ADD(float* A, float* B, float* C)
{
   const int ix = blockDim.x * blockIdx.x + threadIdx.x;
   const int iy = blockDim.y * blockIdx.y + threadIdx.y;

   if(ix < 16 && iy < 16)
   {
      for(int i = 0; i<256; i++)
      C[i] = A[ix+iy*16] + B[ix+iy*16] + C[i]; // << I wish to store all in C
   }
}

extern "C" void cuda_p(float* A, float* B, float* C)
{
    float* dev_A;
    float* dev_B;
    float* dev_C;
    cudaMalloc((void**) &dev_A,  sizeof(float) * 256);
    cudaMalloc((void**) &dev_B,  sizeof(float) * 256);
    cudaMalloc((void**) &dev_C,  sizeof(float) * 256);
    cudaMemcpy(dev_A, A, sizeof(float) * 256, cudaMemcpyHostToDevice);
    cudaMemcpy(dev_B, B, sizeof(float) * 256, cudaMemcpyHostToDevice);
    cudaMemcpy(dev_C, C, sizeof(float) * 256, cudaMemcpyHostToDevice);
    ADDD<<<16,16>>>(dev_A,dev_B,dev_C);
    cudaMemcpy(A, dev_A, sizeof(float) * 256, cudaMemcpyDeviceToHost);
    cudaMemcpy(B, dev_B, sizeof(float) * 256, cudaMemcpyDeviceToHost);
    cudaMemcpy(C, dev_C, sizeof(float) * 256, cudaMemcpyDeviceToHost);
 cudaFree(dev_A);
 cudaFree(dev_B);
 cudaFree(dev_C);
}

Answer 1

1. 你确定内核启动配置吗？在您的代码中，您尝试启动一些未知函数ADDD。而你的执行配置是:gridDim = (16, 0, 0) 和 blockDim = (16, 0, 0)。所以在你的内核中 blockIdx.x = [0..16) 和 threadIdx.x = [0..16)。如果我没听错，那么

   ix = threadIdx.x;
iy = blockIdx.x;

   在 CUDA 编程指南(附录 B.15)中了解它。

2. 但这不仅仅是一个错误。当您在 C[i] 中累积值时，您会遇到竞争条件。 16 个线程(1 个 warp)同时读取 C[i]，添加一些值(A[ix+iy*16] + B[ix+iy*16])和将结果写回 C[i]。您应该使用原子添加操作(CUDA 编程指南，附录 B.11.1.1)或重新设计内核以最大化内存合并(CUDA C 最佳实践指南 3.2.1)，因为原子操作非常非常慢...