CUDA atomicAdd 跨块

Question

我无法让atomicAdd 函数作用于所有块。事实证明，下面的内核代码给了我一个块中的线程总数（例如< 5000）：

__global __ void kernelCode(float *result)
{
    int index = threadIdx.x+blockIdx.x*blockDim.x;
    if (index < 5000)
    {
        atomicAdd(result, 1.0f);
    }
}

您能告诉我如何在不分配1.0f 的整个数组的情况下为值添加一些内容吗？这是因为我在资源非常有限的系统上使用此代码 - 每一点都很重要。

Answer 1

此代码可以跨多个块工作，而无需分配 1.0f 数组。 if (index < 5000) 语句并非旨在将您限制为单个线程块。旨在确保只有整个网格中的合法线程才能参与操作。

试试这样的：

#include <iostream>
#define TOTAL_SIZE 100000
#define nTPB 256

#define cudaCheckErrors(msg) \
    do { \
        cudaError_t __err = cudaGetLastError(); \
        if (__err != cudaSuccess) { \
            fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \
                msg, cudaGetErrorString(__err), \
                __FILE__, __LINE__); \
            fprintf(stderr, "*** FAILED - ABORTING\n"); \
            exit(1); \
        } \
    } while (0)

__global__ void kernelCode(float *result)
{
    int index = threadIdx.x+blockIdx.x*blockDim.x;
    if (index < TOTAL_SIZE)
    {
        atomicAdd(result, 1.0f);
    }
}

int main(){

  float h_result, *d_result;
  cudaMalloc((void **)&d_result, sizeof(float));
  cudaCheckErrors("cuda malloc fail");
  h_result = 0.0f;
  cudaMemcpy(d_result, &h_result, sizeof(float), cudaMemcpyHostToDevice);
  cudaCheckErrors("cudaMemcpy 1 fail");
  kernelCode<<<(TOTAL_SIZE+nTPB-1)/nTPB, nTPB>>>(d_result);
  cudaDeviceSynchronize();
  cudaCheckErrors("kernel fail");
  cudaMemcpy(&h_result, d_result, sizeof(float), cudaMemcpyDeviceToHost);
  cudaCheckErrors("cudaMemcpy 2 fail");
  std::cout<< "result = " << h_result << std::endl;
  return 0;
}

您可以将TOTAL_SIZE 更改为任何适合float 的数字

请注意，我在浏览器中输入了这段代码，可能有印刷错误。