How can I find the id of a gang in OpenACC?

2018-01-19

An answer to this question on Stack Overflow.

Question

In OpenMP I can use omp_get_thread_num() to get the numerical id of the thread executing the code.

Is there a similar function I can use in OpenACC to get id of the gang executing a piece of code?

Answer

The OpenACC standard does not yet include such a function, but, with the PGI compiler, you can use the compiler extension function __pgi_gangidx() as follows:

//pgc++ -fast -acc -ta=tesla,cc60 -Minfo=accel test.cpp
#include <iostream>
#include "openacc.h"
int main(){
  int gangs = 100;
  int *ids  = new int[gangs];
  //Ensure everything is zeroed
  for(int i=0;i<gangs;i++)
    ids[i] = 0;
  #pragma acc parallel num_gangs(gangs) copyout(ids[0:gangs])
  {
    ids[__pgi_gangidx()] = __pgi_gangidx();
  }
  for(int i=0;i<gangs;i++)
    std::cout<<ids[i]<<" ";
  std::cout<<std::endl;
}

Compile with:

pgc++ -fast -acc -ta=tesla,cc60 -Minfo=accel test.cpp

This gives as output:

0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35 36 37 38 39 40 41 42 43 44 45 46 47 48 49 50 51 52 53 54 55 56 57 58 59 60 61 62 63 64 65 66 67 68 69 70 71 72 73 74 75 76 77 78 79 80 81 82 83 84 85 86 87 88 89 90 91 92 93 94 95 96 97 98 99

as expected.

A suite of additional functions are available:

extern int __pgi_gangidx(void);
extern int __pgi_workeridx(void);
extern int __pgi_vectoridx(void);
extern int __pgi_blockidx(int);
extern int __pgi_threadidx(int);

Note that omp_get_thread_num() does not (yet?) work for GPU-targeted code.