OpenACC nvc++ default behavior changed without warning when copying array to GPU using [size] vs [:size]

When copying arrays to the device using acc enter data copyin(D[dim]), the default behavior changed between nvc++ 21.5 and 22.3:

#include <vector>

int main(int argc, char* argv[]){
    uint64_t dim = (1<<27);
    std::vector<double> D_v(dim, 0.0);
    double* D = &D_v[0];
    #pragma acc init
    #pragma acc enter data copyin(D[dim])
    #pragma acc parallel loop present(D[dim])
    for(uint64_t j = 0; j < dim; ++j){
        D[j] += 1;
    }
}

With nvc++ 21.5, D[dim] was automatically interpreted as D[:dim]. Obviously, this was a deliberate design choice (since when would you want to copy the single element behind the last element anyway):

$ nvc++ -std=c++17 -acc -fast -gopt -Minfo=accel -Mcuda -Wall -Wextra -pedantic -o minimal minimal.cpp 
main:
      7, Generating enter data copyin(D[:dim])
         Generating present(D[:dim])
         Generating Tesla code
         11, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */

With 22.3, this behavior changed:

$ nvc++ -std=c++17 -acc -fast -gopt -Minfo=accel -Mcuda -Wall -Wextra -pedantic -o minimal minimal.cpp 
main:
      7, Generating enter data copyin(D[dim])
         Generating present(D[dim])
         Generating NVIDIA GPU code
         11, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */

So all the code that has previously compiled and worked just fine now crashes during execution with

Failing in Thread:1
call to cuStreamSynchronize returned error 700: Illegal address during kernel execution

I agree that technically, D[:dim] would be “more correct”, but this is a breaking change that may affect many existing codes.

Why was this change made now?
Shouldn’t nvc++ at least print a warning?

Best wishes,
Dennis

Hi Dennis,

This was an intentional change in behavior to be inline with the OpenACC standard. We did have warning in the compiler for about 6 months and notice in the Release Notes of this change starting with the 21.7 release. You just happened to jump over these when moving directly from 21.5 to 22.3.

Obviously, this was a deliberate design choice (since when would you want to copy the single element behind the last element anyway):

Yes, it was a deliberate choice, though as the OpenACC standard was refined, the behavior was not inline with the standard and our implementation behaved differently than other compiler’s implementations. Plus it turns out there are cases where folks did want to copy single elements. Hence we decided to change the behavior to require the triplet notation.

Note you can revert to the old behavior via the flag “-gpu=implicitsections”, however, if you decide to use a different compiler such as g++, you’ll run into issues.

It’s never easy to do these type of changes and we tried our best to notify folks well in advance, but apologies that it’s caused you issues.

-Mat

1 Like

Thanks Mat!

For some reason, I cannot reproduce this warning. I tried compiling with 21.7, 21.9, and 22.1. And with 22.1, it still interprets [dim] as [:dim]. Is there a special switch necessary to trigger such warnings? Would be good to know for the future.

Also, in the release notes of 22.3 it still says that it “will change in a future release” but not that it has already been changed.

My error. I thought the warning was in both C++ and Fortran, but looks like it was only added to Fortran codes.

% nvfortran test.F90 -acc -V21.7
NVFORTRAN-W-0155-Upcoming behavior change: array reference in data/update clause with no triplet notation will be treated as an array element in a future release: arr (test.F90: 10)
NVFORTRAN-W-0155-Upcoming behavior change: array reference in data/update clause with no triplet notation will be treated as an array element in a future release: arr (test.F90: 15)
  0 inform,   2 warnings,   0 severes, 0 fatal for foo

I’ll let engineering know about the Release Notes. That should have been updated.

1 Like

This topic was automatically closed 14 days after the last reply. New replies are no longer allowed.