0

假设您有一个函数，该函数将接收一个向量、一组向量，并找出向量集中的哪个向量最接近原始向量。如果我包含一些代码可能会很有用：

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 个元素的块对应一个向量。inputVector是被比较的向量，它也有 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 一起使用？

4

1 回答 1

1

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

第一种方法是最简单的。我的猜测是它也会为您提供最佳性能，尽管它确实为免费的 Thrust 库添加了依赖项。您将使用推力::min_element()。

使用的方法使用具有 64 位模式atomicCAS()的事实atomicCAS()，在该模式中，您可以将任何您希望的语义分配给 64 位值。在您的情况下，您将使用 32 位来存储leastDistance和 32 位来存储winner. 要使用此方法，请在实现双精度浮点的 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);
}
于 2014-03-08T15:09:25.270 回答