support for atomic update of unsigned char

Accelerated Computing HPC Compilers Legacy PGI Compilers
#1

When compiling this code (using PGI 15.1 compiler)

unsigned char counter = 0;

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

I get the following error

PGCC-S-0155-Invalid atomic expression (uchar.cc: 21)
PGCC-S-0155-Invalid atomic region. (uchar.cc: 21)
main:
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.

L

#2

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: http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#features-and-technical-specifications__feature-support-per-compute-capability

Hope this helps,
Mat

#3

Hi Mat,

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

Thanks,
Lutz