0

我无法使该atomicAdd功能适用​​于所有块。事实证明，以下内核代码为我提供了一个块中的线程总数（< 5000例如）：

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

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

4

1 回答 1

2

此代码可以跨多个块工作，而无需分配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

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

于 2013-09-26T12:34:51.733 回答