c++ - 使用 CUDA 在集合中查找最少的数字

Question

假设您有一个函数，它接受一个 vector 、一组 vector ，并找出这组 vector 中的哪个 vector 最接近原始 vector 。如果我包含一些代码，它可能会有用:

int findBMU(float * inputVector, float * weights){


    int count = 0;
    float currentDistance = 0;
    int winner = 0;
    float leastDistance = 99999;

    for(int i = 0; i<10; i++){
        for(int j = 0;j<10; j++){
            for(int k = 0; k<10; k++){

                int offset = (i*100+j*10+k)*644;
                for(int i = offset; i<offset+644; i++){
                    currentDistance += abs((inputVector[count]-weights[i]))*abs((inputVector[count]-weights[i]));
                    count++;
                }
                currentDistance = sqrt(currentDistance);

                count = 0;
                if(currentDistance<leastDistance){
                    winner = offset;

                    leastDistance = currentDistance;

                }
                currentDistance = 0;
            }
        }
    }
    return winner;
}

在这个例子中，weights是一个一维数组，有 644 个元素的 block 对应一个 vector 。 inputVector是被比较的 vector ，它也有 644 个元素。

为了加快程序速度，我决定看一看 NVIDIA 提供的 CUDA 框架。这是我更改代码以适应 CUDA 规范后的样子。

__global__ void findBMU(float * inputVector, float * weights, int * winner, float * leastDistance){




    int i = threadIdx.x+(blockIdx.x*blockDim.x);

    if(i<1000){

        int offset = i*644;
        int count = 0;
        float currentDistance = 0;
        for(int w = offset; w<offset+644; w++){
            currentDistance += abs((inputVector[count]-weights[w]))*abs((inputVector[count]-weights[w]));

            count++;
        }


        currentDistance = sqrt(currentDistance);

        count = 0;
        if(currentDistance<*leastDistance){
            *winner = offset;

            *leastDistance = currentDistance;

        }
        currentDistance = 0;
    }

}

为了调用该函数，我使用了:findBMU<<<20, 50>>>(d_data, d_weights, d_winner, d_least);

但是，当我调用该函数时，有时它会给我正确的答案，有时则不会。在做了一些研究之后，我发现 CUDA 有一些像这样的减少问题，但我找不到解决方法。如何修改我的程序以使其与 CUDA 一起工作？

Answer 1

问题是并发运行的线程会看到相同的 leastDistance并覆盖彼此的结果。线程之间共享两个值； leastDistance和 winner .你有两个基本的选择。您可以写出所有线程的结果，然后通过并行缩减对数据进行第二次传递以确定哪个 vector 具有最佳匹配，或者您可以使用 atomicCAS() 通过自定义原子操作来实现此目的。 .

第一种方法最简单。我的猜测是它也会给你最好的性能，尽管它确实增加了对免费 Thrust 库的依赖。你会使用 thrust::min_element() .

使用atomicCAS()的方法使用 atomicCAS() 的事实有一个 64 位模式，您可以在其中将您想要的任何语义分配给 64 位值。在您的情况下，您将使用 32 位来存储 leastDistance和 32 位存储 winner .要使用此方法，请在实现 double 浮点的 CUDA C 编程指南中改编此示例 atomicAdd() .

__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);
}