Hi!
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;
}
> I'll apply this patch to trunk as obvious.
> libgomp/
> * testsuite/libgomp.oacc-fortran/non-scalar-data.f90: Don't
> pass parameter variables to subroutines.
> --- a/libgomp/testsuite/libgomp.oacc-fortran/non-scalar-data.f90
> +++ b/libgomp/testsuite/libgomp.oacc-fortran/non-scalar-data.f90
> @@ -6,9 +6,11 @@
> program main
> implicit none
>
> - integer, parameter :: n = 100
> - integer :: array(n), i
> -
> + integer,parameter :: size = 100
> + integer :: array(size), i, n
> +
> + n = size
> +
> !$acc data copy(array)
> call kernels(array, n)
;-) To keep the code idiomatic, I'd rather have used a copyin(n) with the
kernels directive, to override the implicit copy(b). (Will do this
change together with committing my test cases above, if approved.)
Grüße
Thomas
