Databases
StackOverflow2013
Postsdbo
POHow to use atomicCAS for multiple variables with conditionals in CUDA
 

Note that there are some explanatory texts on larger screens.

plurals
  1. POHow to use atomicCAS for multiple variables with conditionals in CUDA
    primarykey
    data
    text
    <p>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 <code>atomicMIN</code> 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 <code>float gbl_min_dist</code></p> <p>To do this I use the following code:</p> <pre><code>__device__ inline float atomicMin(float *addr, float value){ float old = *addr, assumed; if( old &lt;= value ) return old; do{ assumed = old; old = atomicCAS((unsigned int*)addr, __float_as_int(assumed), __float_as_int(value)); }while( old!=assumed ); return old; } </code></pre> <p>Suppose now that <strong>we want to store the index of the two points</strong> that were closer together and for which the <code>atomicMIN</code> has successfully replaced the old minimum distance with the one calculated by those two points. What I mean is that <strong>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</strong> </p> <pre><code>typedef struct {float gbl_min_dist, unsigned int point1, unsigned int point2;} global_closest_points; </code></pre> <p>So here, when a thread executes the <code>atomicMIN</code>, if the value that is proposed by that tread to be compared is swapped in the <code>gbl_min_dist</code> then i also need to swap the p1, p2 with the values from the thread. If the <code>gbl_min_dist</code> is not swapped then I dont want to store the points cause this would give wrong points but correct minimum distance. </p> <p>Is there any return value to check if <code>atomicCAS</code> has made the swap? </p> <p>Any ideas on how to implement this within the <code>atomicMIN</code>?</p> <p>Thanks in advance</p>
    singulars
    1. PO
      singulars
      1. PTAnswer
    1. This table or related slice is empty.
    1. PTQuestion
    1. USBugShotGG
    1. USBugShotGG
    plurals
    1. PL
      singulars
      1. LTLinked
    2. PL
      singulars
      1. LTLinked
    1. This table or related slice is empty.
    1. This table or related slice is empty.
    1. PO
      singulars
      1. PTAnswer
    2. PO
      singulars
      1. PTAnswer
    1. This table or related slice is empty.
    1. COBTW: Your atomicMin is wrong...you have to put the if( old <= value ) return old; inside the loop (after do {)! Otherwise you would overwrite a value that is possibly smaller when it is written after your first read and the CAS!
      singulars
      1. POHow to use atomicCAS for multiple variables with conditionals in CUDA
      1. USmmmmmmmm
 

Querying!

 
Guidance
A row detail

Detail views are divided into sections. All the information in the data section comes from columns in the selected row. The other sections display data from other, related rows.

Related data can be related in a to-one or a to-many fashion. Captions of data related in a to-many fashion link to a list view showing a filtered view of the table.

Try moving around until you find a non-empty to-many entry and click on the label to get to one. You can move back to the root by clicking on the database name in the header.

SQuiL has stopped working due to an internal error.

If you are curious you may find further information in the browser console, which is accessible through the devtools (F12).

Reload