Hello. I am trying to understand how to make a conditional atomicCAS if a previous atomicCAS has swapped a value.
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 do not 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;
do{
if( old <= value ) return old;
assumed = old;
old = atomicCAS((unsigned int*)addr, __float_as_int(assumed), __float_as_int(value));
}while( old!=assumed );
return old;
}
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, I only want to store the indexes of the two points that currently have the smaller distance if and only if its distance has just been successfully swapped in the global variable
typedef struct {float gbl_min_dist,
unsigned int point1,
unsigned int point2;} global_closest_points;
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. Is there any return value to check if atomicCAS has made the swap? Any ideas on how to implement this within the atomicMIN using a conditional? Thanks in advance