Re: "const" qualifier vs. OpenACC data/OpenMP map clauses

[Thomas Schwinge <thomas at codesourcery dot com>]

Hi!

Ping.  Or should I open a PR for that?

On Fri, 15 Apr 2016 10:47:34 +0200, I wrote:
> On Thu, 14 Apr 2016 14:21:33 -0700, Cesar Philippidis <cesar@codesourcery.com> wrote:
> > This patch fixes a segfault in libgomp.oacc-fortran/non-scalar-data.f90.
> > The problem here is that 'n' is a parameter, and the kernels region
> > implicitly adds a copy clause to n. Naturally, the test segfaults when
> > it comes time to write the value back to the host as the kernels region
> > terminates.
> 
> ;-| Ha!  Glad that you found this rather quickly!
> 
> > This problem only occurs on nvptx targets.
> 
> It's a generic problem, I would say.  Just in a lot of cases, it seem
> that the writeback doesn't have a destructive effect (if the "const" data
> happens to live in a writeable memory location, for suppose).
> 
> So, in C (which I'm more comfortable with than Fortran), we're basically
> talking about the following case:
> 
>     const int b = 0;
>     #pragma acc kernels /* implicit copy(b) */
>     {
>       b = 1;
>     }
> 
> ..., which is invalid, as the copyout of "b" writes to "const" memory.
> GCC does not currently diagnose that (but should, I think?).
> 
> But, GCC does diagnose "assignment of read-only variable" inside the
> offloaded block, which I wonder whether that's correct: at least in the
> case of OpenACC, the specification says that given a data clause, "the
> compiler will allocate and manage a copy of the variable or array in the
> memory of the current device, creating a visible device copy of that
> variable or array, for non-shared memory devices".  (Have not looked up
> the corresponding for OpenMP.)  Does that "visible device copy of that
> variable" still have the "const" property?  Or only for shared memory
> devices?  (Uh...)  So, to maintain coherence between shared and
> non-shared memory devices, it probably makes sense to indeed emit this
> diagnostic.  (That is, a "visible device copy" keeps the "const" property
> of the original variable.)
> 
> With that settled, does it then follow that, for example, the create
> (OpenMP: map(alloc)) and "const" qualifiers are conflicting?  That is,
> should the following emit some kind of "conflicting data clause for
> read-only variable" diagnostic?
> 
>     const int b = 0;
>     #pragma acc kernels create(b)
>     {
>       b = 1;
>     }
> 
> (I hope the intention of the specification is not to allow for, for
> example, a create clause to override the original variable's "const"
> property.)
> 
> Here's what I got so far; OK to commit to trunk (as
> gcc/testsuite/c-c++-common/{goacc,gomp}/read-only.c?)?  I suppose more
> test cases to be added once resolving the XFAILs.
> 
>     void openacc(const double a)
>     {
>       const short b = 0;
>     #pragma acc kernels /* implicit copy(b) */ /* { dg-error "assignment of read-only variable" "TODO" { xfail *-*-* } } */
>       {
>         b = 1; /* { dg-error "assignment of read-only variable" } */
>       }
>     
>     #pragma acc kernels create(b) /* { dg-error "TODO conflicting data clause for read-only variable" "TODO" { xfail *-*-* } } */
>       {
>         b = 1; /* { dg-error "assignment of read-only variable" } */
>       }
>     
>       (void) b;
>     }
> 
>     void openmp(const int a)
>     {
>       a = 10; /* { dg-error "assignment of read-only parameter" } */
>     
>     #pragma omp target data map(from:a) /* { dg-error "assignment of read-only parameter" "TODO" { xfail *-*-* } } */
>       {
>         a = 10; /* { dg-error "assignment of read-only parameter" } */
>       }
>     
>     #pragma omp target map(from:a) /* { dg-error "assignment of read-only parameter" "TODO" { xfail *-*-* } } */
>       {
>         a = 10; /* { dg-error "assignment of read-only parameter" } */
>       }
>     
>     #pragma omp target map(alloc:a) /* { dg-error "TODO conflicting map clause for read-only parameter" "TODO" { xfail *-*-* } } */
>       {
>         a = 10; /* { dg-error "assignment of read-only parameter" } */
>       }
>     
>       const float b = 1;
>       b = 0.123;  /* { dg-error "assignment of read-only variable" } */
>     
>     #pragma omp target data map(from:b) /* { dg-error "assignment of read-only variable" "TODO" { xfail *-*-* } } */
>       {
>         b = 10; /* { dg-error "assignment of read-only variable" } */
>       }
>     
>     #pragma omp target map(from:b) /* { dg-error "assignment of read-only variable" "TODO" { xfail *-*-* } } */
>       {
>         b = 10; /* { dg-error "assignment of read-only variable" } */
>       }
>     
>     #pragma omp target map(alloc:b) /* { dg-error "TODO conflicting map clause for read-only variable" "TODO" { xfail *-*-* } } */
>       {
>         b = 10; /* { dg-error "assignment of read-only variable" } */
>       }
>     
>       (void) b;
>     }


GrÃÃe
 Thomas