c++ - 无法在 CUDA 中找到 1 到 100 数字的简单和？

Question

我正在研究使用 CUDA 的图像处理算法。在我的算法中，我想使用 CUDA 内核找到图像所有像素的总和。所以我在cuda中制作了内核方法​​来测量16位灰度图像的所有像素的总和，但我得到了错误的答案。 所以我在cuda中编写了一个简单的程序来查找1到100个数字的总和，我的代码如下。 在我的代码中，我没有使用 GPU 得到 1 到 100 个数字的精确总和，但使用 CPU 得到了 1 到 100 个数字的精确总和。那么我在该代码中做了什么？

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#include <conio.h>
#include <malloc.h>
#include <limits>
#include <math.h>

using namespace std;

__global__ void computeMeanValue1(double *pixels,double *sum){

        int x = threadIdx.x;

        sum[0] = sum[0] + (pixels[(x)]);
        __syncthreads();
}

int main(int argc, char **argv)
{
    double *data;
    double *dev_data;
    double *dev_total;
    double *total;

    data=new double[(100) * sizeof(double)];
    total=new double[(1) * sizeof(double)];

    double cpuSum=0.0;

    for(int i=0;i<100;i++){
        data[i]=i+1;
        cpuSum=cpuSum+data[i];
    }
    cout<<"CPU total = "<<cpuSum<<std::endl;

    cudaMalloc( (void**)&dev_data, 100 * sizeof(double));
    cudaMalloc( (void**)&dev_total, 1 * sizeof(double));

    cudaMemcpy(dev_data, data, 100 * sizeof(double), cudaMemcpyHostToDevice);

    computeMeanValue1<<<1,100>>>(dev_data,dev_total);
    cudaDeviceSynchronize();

    cudaMemcpy(total, dev_total, 1* sizeof(double), cudaMemcpyDeviceToHost);
    cout<<"GPU total = "<<total[0]<<std::endl;

    cudaFree(dev_data);
    cudaFree(dev_total);

    free(data);
    free(total);

    getch();
    return 0;
}

Answer 1

所有线程同时写入同一内​​存位置。

sum[0] = sum[0] + (pixels[(x)]);

您不能这样做并期望得到正确的结果。您的内核需要采取不同的方法来避免从不同的线程写入相同的内存。通常用于执行此操作的模式是减少。简而言之，每个线程负责对数组中的元素 block 求和，然后存储结果。通过采用一系列这些归约运算，可以对数组的整个内容进行求和。

__global__ void block_sum(const float *input,
                          float *per_block_results,
                          const size_t n)
{
  extern __shared__ float sdata[];
  unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;

  // load input into __shared__ memory
  float x = 0;
  if(i < n)
  {
    x = input[i];
  }
  sdata[threadIdx.x] = x;
  __syncthreads();

  // contiguous range pattern
  for(int offset = blockDim.x / 2;
      offset > 0;
      offset >>= 1)
  {
    if(threadIdx.x < offset)
    {
      // add a partial sum upstream to our own
      sdata[threadIdx.x] += sdata[threadIdx.x + offset];
    }

    // wait until all threads in the block have
    // updated their partial sums
    __syncthreads();
  }

  // thread 0 writes the final result
  if(threadIdx.x == 0)
  {
    per_block_results[blockIdx.x] = sdata[0];
  }
}

每个线程写入 sdata[threadIdx.x] 中的不同位置没有竞争条件。线程可以自由访问 sdata 中的其他元素因为他们只读取它们，所以不存在竞争条件。注意 __syncthreads() 的使用确保加载数据到sdata的操作在线程开始读取数据和第二次调用 __syncthreads() 之前完成确保在从 sdata[0] 复制最终结果之前已完成所有求和操作。 。请注意，只有线程 0 将其结果写入 per_block_results[blockIdx.x] ，所以那里也不存在竞争条件。

您可以在 Google Code 上找到上述完整的示例代码(这个不是我写的)。该幻灯片有一个合理的总结:reductions in CUDA 。它包含的图表确实有助于理解交错内存的读取和写入如何不会相互冲突。

您可以找到许多有关在 GPU 上高效实现缩减的其他 Material 。确保您的实现最有效地利用内存是从内存限制操作(如缩减)中获得最佳性能的关键。