如何在CUDA中将atomicCAS用于带有条件的多个变量 [英] How to use atomicCAS for multiple variables with conditionals in CUDA

问题描述

我最近在编程中遇到了一个简单的概念，但是当我尝试在cuda中实现它时，我陷入了困境. 假设我有成千上万个元素，而我想找到它们之间最接近的对.我在全局内存中使用atomicMIN(假设我们不想减少)，因此，如果每个线程计算的距离小于全局变量中存储的距离，atomicCAS将用较小的值替换它.例如我有全局变量float gbl_min_dist

I recently encountered a simple notion in programming but i stuck when i tried to implement it in cuda. Suppose that i have thousands of elements and i want to find the closest pair between them. I use atomicMIN in global memory (suppose that we dont want to reduce) so if the distance which is calculated by each thread is smaller than the distance stored in the global variable the atomicCAS will replace it with the smaller value. For example i have the global variable float gbl_min_dist

为此，我使用以下代码:

To do this I use the following code:

__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已成功地将旧的最小距离替换为由这两个点计算出的最小距离.我的意思是我只想存储当前距离较小的两个点的索引，当且仅当它的距离已成功在全局变量中成功交换时

Suppose now that we want to store the index of the two points that were closer together and for which the atomicMIN has successfully replaced the old minimum distance with the one calculated by those two points. What I mean is that I only want to store the indeces of the two points that currently have the smaller distance if and only if its distance has just been successfully swaped in the global variable

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

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

So here, when a thread executes the atomicMIN, if the value that is proposed by that tread to be compared is swapped in the gbl_min_dist then i also need to swap the p1, p2 with the values from the thread. If the gbl_min_dist is not swapped then I dont want to store the points cause this would give wrong points but correct minimum distance.

是否有返回值可检查atomicCAS是否进行了交换?

Is there any return value to check if atomicCAS has made the swap?

关于如何在atomicMIN内实现此目标的任何想法?

Any ideas on how to implement this within the atomicMIN?

预先感谢

推荐答案

1. 您可以使用关键部分以使每个线程都具有独占访问权限在数据更新时对其进行处理.
2. 由于您的gbl_min_dist是32位值，因此，如果您可以找到将p1和p2都压缩为单个32位值的方法，则可以使用自定义原子答案之类的方法我在此处给出了.

1. You could use a critical section to have each thread have exclusive access to the data while it is updating it.
2. Since your gbl_min_dist is a 32-bit value, if you can figure out a way to squeeze both p1 and p2 into a single 32-bit value, you could use an approach like the custom atomics answer I gave here.

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

If you simply use whether or not the atomicCAS made the first swap to condition additional code to update p1 and p2, I think it's still possible to have a race condition that allows your data to get out of sync between thread updates.

这篇关于如何在CUDA中将atomicCAS用于带有条件的多个变量的文章就介绍到这了，希望我们推荐的答案对大家有所帮助，也希望大家多多支持IT屋！