Calling functions within OpenACC kernels

Parallel Programming using GPUs has become very easy because of OpenACC. Being directive based programming approach, unknowingly we end up getting a lot of error messages. It is very common that we see errors like these while compiling the code using PGI compiler –

PGC-S-0155-Procedures called in a compute region must have acc routine information:
PGC-S-0155-Accelerator region ignored; see -Minfo messages
PGC/x86-64 Linux 19.10-0: compilation completed with severe errors

As the error message suggests, we probably missed to add “acc routine” clause to one of the function being called from one of the OpenACC Kernels. So, our compiler is looking for the statement like following –

#pragma acc routine seq

Here, seq denotes that we would like to run this function sequentially by the calling thread. Instead we can also use other clauses such as gang, worker, vector etc.

So, for reference I have provided an example of pi value calculation (by calculating area of the unit circle – area of a circle with radius = 1 units)

#include<stdio.h>
#include<stdlib.h>
#include<math.h>

#define N 999999999

#pragma acc routine seq
double calc_area(double x, double y)
{
double area = 0.0;
area = x*y;
return area;
}

int main()
{
int i, j;
double area, pi, temp_area;
double dx, y, x;

dx = 1.0/N;
x = 0.0;
area = 0.0;

//Calculate area for one quadrant
#pragma acc parallel loop reduction(+:area)
for(i=0;i<N;i++)
{
x = i*dx;
y = sqrt(1 – x*x);
temp_area = calc_area(dx,y);
area += temp_area;
}

pi = 4.0*area; //as the area was caluculated only for one quadrant
printf(“\n Value of pi is = %.16lf\n”, pi);
}

To compile this code, you can use following command –

pgcc -Minfo=accel -ta=tesla:managed pi_calc.c

You should get following compilation output for PGI 19.10 Community Edition

calc_area:
9, Generating acc routine seq
Generating Tesla code
main:
25, Generating Tesla code
26, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
Generating reduction(+:area)
25, Generating implicit copy(area) [if not already present]

After executing the code, you will get following output –

Value of pi is = 3.1415926555897564

Thanks for visiting our website!

Leave a Reply