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
> <ce...@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
