https://gcc.gnu.org/bugzilla/show_bug.cgi?id=90742
Bug ID: 90742
Summary: OpenACC/OpenMP target offloading: Fortran
'allocatable' scalars in 'firstprivate' clauses
Product: gcc
Version: 10.0
Status: UNCONFIRMED
Keywords: openacc, openmp
Severity: normal
Priority: P3
Component: fortran
Assignee: unassigned at gcc dot gnu.org
Reporter: tschwinge at gcc dot gnu.org
CC: jakub at gcc dot gnu.org, jules at gcc dot gnu.org
Target Milestone: ---
Created attachment 46449
--> https://gcc.gnu.org/bugzilla/attachment.cgi?id=46449&action=edit
libgomp.oacc-fortran/allocatable-scalar-1.f90
For OpenACC/OpenMP target offloading, Fortran 'allocatable' scalars in
'firstprivate' clauses don't work, with nvptx offloading: "libgomp:
cuStreamSynchronize error: unspecified launch failure (perhaps abort was
called)"; see the first half of the attached test case. Should they not work
in the same way as they do for regular data clauses; see the second half of the
attached test case?
Submitted originally as part of
<http://mid.mail-archive.com/486f1f87-16b3-46b5-17f5-6857756d4115@codesourcery.com>
"[gomp4] add support for allocatable scalars in OpenACC declare constructs"
(which in turn depends on
<http://mid.mail-archive.com/86f51209-c59d-a4cf-297d-9a072823aa61@codesourcery.com>
"[gomp4] add support for fortran allocate support with declare create", and the
patches to enable 'GOMP_MAP_FIRSTPRIVATE_INT' for OpenACC 'firstprivate'), and
recently re-submitted as part of
<http://mid.mail-archive.com/20180920195908.04486d45@squid.athome> '[PATCH,
OpenACC] Fortran "declare create"/allocate support for OpenACC', Cesar also
once singled out the changes to make this work (but for OpenACC only):
<http://mid.mail-archive.com/1f88e441-d3da-5b59-4278-058ff1368a73@codesourcery.com>
"[PATCH][OpenACC] Add support for firstprivate Fortran allocatable scalars".
The latter patch still applies, and makes the OpenACC (but not OpenMP) case
work by means of the following '-fdump-tree-omplower' changes:
[...]
- .omp_data_arr.28.a = &a;
+ a.32 = a;
+ .omp_data_arr.28.a = a.32;
.omp_data_arr.28.b = &b;
#pragma omp target oacc_parallel firstprivate(a) map(from:b [len:
4]) [child fn: MAIN__._omp_fn.0 (.omp_data_arr.28, .omp_data_sizes.29,
.omp_data_kinds.30)]
{
.omp_data_i = (const struct .omp_data_t.25 & restrict)
&.omp_data_arr.28;
D.3974 = .omp_data_i->a;
a.27 = *D.3974;
- a = a.27;
+ a.28 = a.27;
+ a = &a.28;
a.5 = a;
b.6 = *a.5;
[...]
I'm currently working on figuring out if that's correct (OpenACC/OpenMP
standards, and then the GCC implementation), but will appreciate any insights.
(The old code version cited above shows what likely causes the 'SIGSEGV' on the
device: dereferencing a host pointer.)
Also need to look up how that relates to the changes done for PR77371
(<http://mid.mail-archive.com/21a94405-6cb7-e474-1e14-727de6939eaf@mentor.com>),
and PR85879
(<http://mid.mail-archive.com/34d2a115-c01c-26c1-57d9-4989a0a00095@mentor.com>),
where the latter's commit r261025 subsumed the former's changes.
