2

我正在研究使用 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;
}
4

2 回答 2

3

您的所有线程同时写入相同的内存位置。

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

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

__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在线程开始读取数据之前完成加载数据的操作，并在第二次调用 to__syncthreads()以确保在从 复制最终结果之前完成所有求和操作sdata[0]。请注意，只有线程 0 将其结果写入per_block_results[blockIdx.x]，因此那里也没有竞争条件。

您可以在Google Code上找到上述的完整示例代码（我没有写这个）。该幻灯片对 CUDA 的减少进行了合理的总结。它包括真正有助于理解交错内存读取和写入如何不相互冲突的图表。

您可以找到许多其他关于在 GPU 上有效实现缩减的材料。确保您的实现最有效地利用内存是从内存绑定操作（如缩减）中获得最佳性能的关键。

于 2013-11-13T05:58:47.307 回答
1

在 GPU 代码中，我们有多个线程并行执行。如果所有这些线程都试图更新内存中的同一位置，我们就会有未定义的行为，除非我们使用特殊操作，调用atomics来执行更新。

在您的情况下，由于sum由所有线程更新，并且sum是一个double数量，我们可以使用编程指南中描述的特殊自定义原子函数来完成此操作。

如果我用以下代码替换您的内核代码：

__device__ double atomicAdd(double* address, double val)
{
    unsigned long long int* address_as_ull =
                              (unsigned long long int*)address;
    unsigned long long int old = *address_as_ull, assumed;
    do {
        assumed = old;
        old = atomicCAS(address_as_ull, assumed,
                        __double_as_longlong(val +
                               __longlong_as_double(assumed)));
    } while (assumed != old);
    return __longlong_as_double(old);
}

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

        int x = threadIdx.x;
        atomicAdd(sum, pixels[x]);
}

sum并在内核之前将值初始化为零：

double gpuSum = 0.0;
cudaMemcpy(dev_total, &gpuSum, sizeof(double), cudaMemcpyHostToDevice);

然后我想你会得到匹配的结果。

正如@AdeMiller 指出的那样，执行这样的并行求和的更快方法是通过经典的并行归约。

有一个 CUDA 示例代码演示了这一点，并附带了一个涵盖该方法的演示文稿。

于 2013-11-13T06:07:50.437 回答