support for atomic update of unsigned char

When compiling this code (using PGI 15.1 compiler)

unsigned char counter = 0;

#ifdef _OPENACC
#pragma acc kernels
#pragma acc loop independent
   for( int i( 0 ); i < 100; ++i )
#ifdef _OPENACC
#pragma acc atomic update

I get the following error

PGCC-S-0155-Invalid atomic expression ( 21)
PGCC-S-0155-Invalid atomic region. ( 21)
16, Loop is parallelizable
Accelerator kernel generated
16, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
PGCC/x86 Linux 15.1-0: compilation completed with severe errors

However, when the type of variable counter is set to unsigned int, it compiles without problems.

Is this a known limitation? The standard doesn’t say anything about type restrictions for atomic.

Thank you.


Hi L,

This is a hardware restriction rather than a restriction in the OpenACC spec. You can find which NVIDIA devices support which data types here:

Hope this helps,

Hi Mat,

thanks for the useful link for NVIDIA devices. Is there a similar link for AMD devices?