Parallelizing nested loops with OpenACC

Question

I've written a serial method that involves four nested for loops - I'd like to parallelize this method using OpenACC (this is the first time I've tried using it and I'm not very familiar with all the directives).

I tried the following but see the following error: call to cuStreamSynchronize returned error 700: Illegal address during kernel execution

I've pasted a simplified pseudocode version of my method below, I'd really appreciate help figuring out the best way to parallelize this four nested loop structure.

// a, b, and c are input arguments to this method
#pragma acc parallel
for(int j = 0; j < a; j++){
    for(int i = 0; i < b; i++){

        // computing mins and maxs based on formulas with i, j, a, b, and c
        int minX = ...
        int maxX = ...
        int minY = ...
        int maxY = ...

        double count = (maxX - minX + 1)*(maxY - minY + 1);
        int sum1 = 0;
        int sum2 = 0;
        int sum3 = 0;

        #pragma acc loop
        for (int y = minY; y < maxY; y++) {
          for (int x = minX; x < maxX; x++) {
            #pragma acc routine(function_call_name) seq
            sum1 += // some function call;
            sum2 += // some function call;
            sum3 += // some function call;
          }
        }
        int result1 = (int)(sum1/count);
        int result2 = (int)(sum2/count);
        int result3 = (int)(sum3/count);

        #pragma acc routine(function_call_name) seq
        // calling some function call to store result1, result2, result3 in the output
    }
}

Mat Colgrove Mat Colgrove · Accepted Answer · 2018-04-03T21:39:04

An "illegal address" means that your program is accessing a bad address on the GPU. Typically this is caused by an out-of-bounds access, accessing a host address on the device, using an aggregate data structure with dynamic data members and not "attaching" the members (i.e. setting the device pointers in the parent structure). Less common cases are heap or stack overflows.

How are you managing your data? Data regions elsewhere in your code?

If using PGI, try first targeting a multicore CPU (-ta=multicore) so you don't need to worry about data movement. Once you have the parallel regions working, you can then go back to using the GPU and work on the data movement. I'd recommend you start by using CUDA Unified Memory (-ta=tesla:managed) so the CUDA driver handles the data movement for you (dynamic data only). Then once this is working, try adding data regions to manually manage the data.

Other things I see:

The parallel construct needs a loop directive on the outer loops.

#pragma acc parallel loop
for(int j = 0; j < a; j++){
    for(int i = 0; i < b; i++){

You may consider collapsing the loops depending on the loop trip count:

#pragma acc parallel loop collapse(2)
for(int j = 0; j < a; j++){
    for(int i = 0; i < b; i++){

Also, the "routine" directive should decorate the routine's prototype or definition but shouldn't be used in a compute regions.

If you are using any global variables in your device routines, be sure to put then into "declare" directives so a global copy of the data is created on the device.

If you are using PGI, add the "-Minfo=accel" option to the compiler. This will give you the compiler feedback messages on how the compiler is parallelizing your code.

If you aren't using data directives, the compiler will need to implicitly copy the data. The messages will tell you what arrays are being copied along with the size being copied.

If you have trouble understanding the feedback messages, post the output from your compilation and I'll help you walk through them.

Parallelizing nested loops with OpenACC

1 Answers