Skip to content

Use atomic functions in OpenACC

An answer to this question on Stack Overflow.

Question

I want to use atomic functions with OpenACC directives. What are the compile options of pgc++ that can help? Should I use a particular header file?

Answer

One thing to check: are you using the right data? Per this forum not all accelerators accept all data types and, possibly, not all data types can be accessed atomically. Table 13 at this link indicates that atomics on Nvidia seem to only be available for 32-bit integer and floating-point datatypes. 64-bit is available for compute capacities 6.x+.

You say your code looks like:

#pragma acc atomic
{
  res[i][i]=res[i][i]+x;
  res[j][j]=res[j][j]+y;
}

However, I don't think you can nest multiple atomics together like that. Try, instead:

#pragma acc atomic update
res[i][i]=res[i][i]+x;
#pragma acc atomic update
res[j][j]=res[j][j]+y;