无法在 CUDA 中找到 1 到 100 个数字的简单总和

Unable to find simple sum of 1 to 100 numbers in CUDA?

本文关键字：数字 简单 CUDA      更新时间：2023-10-16

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

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

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

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

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

在 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 示例代码来演示这一点，并随附一个介绍该方法的演示文稿。

相关文章：