1

我最近在编程中遇到了一个简单的概念，但是当我尝试在 cuda 中实现它时我卡住了。假设我有数千个元素，我想找到它们之间最近的一对。我atomicMIN在全局内存中使用（假设我们不想减少），所以如果每个线程计算的距离小于存储在全局变量中的距离，atomicCAS 将用较小的值替换它。例如我有全局变量float gbl_min_dist

为此，我使用以下代码：

__device__ inline float atomicMin(float *addr, float value){
    float old = *addr, assumed;
    if( old <= value ) return old;
    do{
        assumed = old;
        old = atomicCAS((unsigned int*)addr, __float_as_int(assumed), __float_as_int(value));
    }while( old!=assumed );
    return old;
}

现在假设我们想要存储靠得更近的两个点的索引，并且atomicMIN已经成功地将旧的最小距离替换为由这两个点计算出的最小距离。我的意思是，当且仅当它的距离刚刚在全局变量中成功交换时，我只想存储当前距离较小的两个点的索引 

typedef struct {float gbl_min_dist, 
                unsigned int point1,
                unsigned int point2;} global_closest_points;

因此，在这里，当一个线程执行时atomicMIN，如果该线程建议的要比较的值被交换，gbl_min_dist那么我还需要将 p1、p2 与线程中的值交换。如果gbl_min_dist没有交换，那么我不想存储这些点，因为这会给出错误的点但正确的最小距离。

是否有任何返回值来检查是否atomicCAS进行了交换？

关于如何在 中实现这一点的任何想法atomicMIN？

提前致谢

4

3 回答 3

1

您可以构造一个临界区来自动更新最小值和对应的点索引。以下链接提供了一个关于如何使用atomicCAS()和构建 CS 的示例atomicExch()。

https://github.com/ArchaeaSoftware/cudahandbook/blob/master/memory/spinlockReduction.cu

另一方面，我建议用并行缩减算法替换原子最小操作。这可能会提高性能。

于 2013-10-14T15:03:29.760 回答
1
  1. 您可以使用关键部分让每个线程在更新数据时拥有对数据的独占访问权限。
  2. 由于您gbl_min_dist是一个 32 位值，如果您能找到一种方法将两者都压缩p1到p2一个 32 位值中，您可以使用类似于我在此处给出的自定义原子答案的方法。

如果您只是使用是否atomicCAS进行了第一次交换来调节要更新的附加代码p1和p2，我认为仍然可能存在竞争条件，使您的数据在线程更新之间不同步。

于 2013-10-14T15:05:22.090 回答
0

我建议的方法是，不要依赖于存储的距离，而是在存储点可能已经改变的关键时刻重新计算它：

typedef struct {
    unsigned int point1, 
    unsigned int point2;
}

global_closest_points, local_closest_points, temp_c_p;

local_dist = distance(local_closest_points.point1, local_closest_points.point2);
temp_c_p = global_closest_points;
while (local_dist < distance(temp_c_p.point1, temp_c_p.point2)
    temp_c_p = atomicCAS(&global_closest_points, temp_c_p, local_closest_points);

旧习惯是，保存而不是重新计算。但是对于现代处理器，这通常不是最优的。在 CUDA 上，对全局内存的原子更新比计算数百个双精度距离需要更多时间。

于 2015-05-03T21:40:03.537 回答