* [PATCH] Fix component mappings with derived types for OpenACC @ 2020-01-10 1:51 Julian Brown 2020-01-24 10:53 ` Tobias Burnus 0 siblings, 1 reply; 11+ messages in thread From: Julian Brown @ 2020-01-10 1:51 UTC (permalink / raw) To: gcc-patches, Tobias Burnus, fortran, Thomas Schwinge [-- Attachment #1: Type: text/plain, Size: 1317 bytes --] Hi, This patch fixes a bug with mapping Fortran components (i.e. with the manual deep-copy support) which themselves have derived types. I've also added a couple of new tests to make sure such mappings are lowered correctly, and to check for the case that Tobias found in the message: https://gcc.gnu.org/ml/gcc-patches/2020-01/msg00215.html The previous incorrect mapping was causing the error condition mentioned in that message to fail to trigger, and I think (my!) code in libgomp (goacc_exit_data_internal) to handle GOMP_MAP_STRUCT specially was papering over the bad mapping also. Oops! I haven't attempted to implement the (harder) sub-copy detection, if that is indeed supposed to be forbidden by the spec. This patch should get us to the same behaviour in Fortran as in C & C++ though. Tested with offloading to nvptx, also with the (uncommitted) reference-count self-checking patch enabled. OK? Thanks, Julian ChangeLog gcc/fortran/ * trans-openmp.c (gfc_trans_omp_clauses): Use inner not decl for components with derived types. gcc/testsuite/ * gfortran.dg/goacc/mapping-tests-3.f90: New test. * gfortran.dg/goacc/mapping-tests-4.f90: New test. libgomp/ * oacc-mem.c (goacc_exit_data_internal): Remove special (no-copyback) behaviour for GOMP_MAP_STRUCT. [-- Attachment #2: component-mappings-derived-types-1.diff --] [-- Type: text/x-patch, Size: 3656 bytes --] commit 5e9d8846bbaa33a9bdb08adcf1ee9f224a8e8fc0 Author: Julian Brown <julian@codesourcery.com> Date: Wed Jan 8 15:57:46 2020 -0800 Fix component mappings with derived types for OpenACC gcc/fortran/ * trans-openmp.c (gfc_trans_omp_clauses): Use inner not decl for components with derived types. gcc/testsuite/ * gfortran.dg/goacc/mapping-tests-3.f90: New test. * gfortran.dg/goacc/mapping-tests-4.f90: New test. libgomp/ * oacc-mem.c (goacc_exit_data_internal): Remove special (no-copyback) behaviour for GOMP_MAP_STRUCT. diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c index 9835d2aae3c..efc392d7fa6 100644 --- a/gcc/fortran/trans-openmp.c +++ b/gcc/fortran/trans-openmp.c @@ -2783,9 +2783,9 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, } else { - OMP_CLAUSE_DECL (node) = decl; + OMP_CLAUSE_DECL (node) = inner; OMP_CLAUSE_SIZE (node) - = TYPE_SIZE_UNIT (TREE_TYPE (decl)); + = TYPE_SIZE_UNIT (TREE_TYPE (inner)); } } else if (lastcomp->next diff --git a/gcc/testsuite/gfortran.dg/goacc/mapping-tests-3.f90 b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-3.f90 new file mode 100644 index 00000000000..312f596461e --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-3.f90 @@ -0,0 +1,15 @@ +! { dg-options "-fopenacc -fdump-tree-omplower" } + +subroutine foo + type one + integer i, j + end type + type two + type(one) A, B + end type + + type(two) x + + !$acc enter data copyin(x%A) +! { dg-final { scan-tree-dump-times "omp target oacc_enter_exit_data map\\(struct:x \\\[len: 1\\\]\\) map\\(to:x.a \\\[len: \[0-9\]+\\\]\\)" 1 "omplower" } } +end diff --git a/gcc/testsuite/gfortran.dg/goacc/mapping-tests-4.f90 b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-4.f90 new file mode 100644 index 00000000000..6257af942df --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-4.f90 @@ -0,0 +1,17 @@ +subroutine foo + type one + integer i, j + end type + type two + type(one) A, B + end type + + type(two) x + +! This is accepted at present, although it represents a probably-unintentional +! overlapping subcopy. + !$acc enter data copyin(x%A, x%A%i) +! But this raises an error. + !$acc enter data copyin(x%A, x%A%i, x%A%i) +! { dg-error ".x.a.i. appears more than once in map clauses" "" { target "*-*-*" } 15 } +end diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index 2d4bba78efd..232683a85f0 100644 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.c @@ -1136,38 +1136,6 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum, break; case GOMP_MAP_STRUCT: - { - int elems = sizes[i]; - for (int j = 1; j <= elems; j++) - { - struct splay_tree_key_s k; - k.host_start = (uintptr_t) hostaddrs[i + j]; - k.host_end = k.host_start + sizes[i + j]; - splay_tree_key str; - str = splay_tree_lookup (&acc_dev->mem_map, &k); - if (str) - { - if (finalize) - { - if (str->refcount != REFCOUNT_INFINITY) - str->refcount -= str->virtual_refcount; - str->virtual_refcount = 0; - } - if (str->virtual_refcount > 0) - { - if (str->refcount != REFCOUNT_INFINITY) - str->refcount--; - str->virtual_refcount--; - } - else if (str->refcount > 0 - && str->refcount != REFCOUNT_INFINITY) - str->refcount--; - if (str->refcount == 0) - gomp_remove_var_async (acc_dev, str, aq); - } - } - i += elems; - } break; default: ^ permalink raw reply [flat|nested] 11+ messages in thread
* Re: [PATCH] Fix component mappings with derived types for OpenACC 2020-01-10 1:51 [PATCH] Fix component mappings with derived types for OpenACC Julian Brown @ 2020-01-24 10:53 ` Tobias Burnus 2020-01-28 14:37 ` Julian Brown 0 siblings, 1 reply; 11+ messages in thread From: Tobias Burnus @ 2020-01-24 10:53 UTC (permalink / raw) To: Julian Brown, gcc-patches, Tobias Burnus, fortran, Thomas Schwinge Hi Julian, the gfortran part is rather obvious and it and the test case look fine to me â OK. The oacc-mem.c also looks okay, but I assume Thomas needs to rubber-stamp it. Tobias On 1/10/20 2:49 AM, Julian Brown wrote: > This patch fixes a bug with mapping Fortran components (i.e. with the > manual deep-copy support) which themselves have derived types. I've > also added a couple of new tests to make sure such mappings are lowered > correctly, and to check for the case that Tobias found in the message: > > https://gcc.gnu.org/ml/gcc-patches/2020-01/msg00215.html > > The previous incorrect mapping was causing the error condition mentioned > in that message to fail to trigger, and I think (my!) code in libgomp > (goacc_exit_data_internal) to handle GOMP_MAP_STRUCT specially was > papering over the bad mapping also. Oops! > > I haven't attempted to implement the (harder) sub-copy detection, if > that is indeed supposed to be forbidden by the spec. This patch should > get us to the same behaviour in Fortran as in C & C++ though. > > Tested with offloading to nvptx, also with the (uncommitted) > reference-count self-checking patch enabled. > > OK? > > Thanks, > > Julian > > ChangeLog > > gcc/fortran/ > * trans-openmp.c (gfc_trans_omp_clauses): Use inner not decl for > components with derived types. > > gcc/testsuite/ > * gfortran.dg/goacc/mapping-tests-3.f90: New test. > * gfortran.dg/goacc/mapping-tests-4.f90: New test. > > libgomp/ > * oacc-mem.c (goacc_exit_data_internal): Remove special (no-copyback) > behaviour for GOMP_MAP_STRUCT. > > component-mappings-derived-types-1.diff > > commit 5e9d8846bbaa33a9bdb08adcf1ee9f224a8e8fc0 > Author: Julian Brown<julian@codesourcery.com> > Date: Wed Jan 8 15:57:46 2020 -0800 > > Fix component mappings with derived types for OpenACC > > gcc/fortran/ > * trans-openmp.c (gfc_trans_omp_clauses): Use inner not decl for > components with derived types. > > gcc/testsuite/ > * gfortran.dg/goacc/mapping-tests-3.f90: New test. > * gfortran.dg/goacc/mapping-tests-4.f90: New test. > > libgomp/ > * oacc-mem.c (goacc_exit_data_internal): Remove special (no-copyback) > behaviour for GOMP_MAP_STRUCT. > > diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c > index 9835d2aae3c..efc392d7fa6 100644 > --- a/gcc/fortran/trans-openmp.c > +++ b/gcc/fortran/trans-openmp.c > @@ -2783,9 +2783,9 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, > } > else > { > - OMP_CLAUSE_DECL (node) = decl; > + OMP_CLAUSE_DECL (node) = inner; > OMP_CLAUSE_SIZE (node) > - = TYPE_SIZE_UNIT (TREE_TYPE (decl)); > + = TYPE_SIZE_UNIT (TREE_TYPE (inner)); > } > } > else if (lastcomp->next > diff --git a/gcc/testsuite/gfortran.dg/goacc/mapping-tests-3.f90 b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-3.f90 > new file mode 100644 > index 00000000000..312f596461e > --- /dev/null > +++ b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-3.f90 > @@ -0,0 +1,15 @@ > +! { dg-options "-fopenacc -fdump-tree-omplower" } > + > +subroutine foo > + type one > + integer i, j > + end type > + type two > + type(one) A, B > + end type > + > + type(two) x > + > + !$acc enter data copyin(x%A) > +! { dg-final { scan-tree-dump-times "omp target oacc_enter_exit_data map\\(struct:x \\\[len: 1\\\]\\) map\\(to:x.a \\\[len: \[0-9\]+\\\]\\)" 1 "omplower" } } > +end > diff --git a/gcc/testsuite/gfortran.dg/goacc/mapping-tests-4.f90 b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-4.f90 > new file mode 100644 > index 00000000000..6257af942df > --- /dev/null > +++ b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-4.f90 > @@ -0,0 +1,17 @@ > +subroutine foo > + type one > + integer i, j > + end type > + type two > + type(one) A, B > + end type > + > + type(two) x > + > +! This is accepted at present, although it represents a probably-unintentional > +! overlapping subcopy. > + !$acc enter data copyin(x%A, x%A%i) > +! But this raises an error. > + !$acc enter data copyin(x%A, x%A%i, x%A%i) > +! { dg-error ".x.a.i. appears more than once in map clauses" "" { target "*-*-*" } 15 } > +end > diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c > index 2d4bba78efd..232683a85f0 100644 > --- a/libgomp/oacc-mem.c > +++ b/libgomp/oacc-mem.c > @@ -1136,38 +1136,6 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum, > break; > > case GOMP_MAP_STRUCT: > - { > - int elems = sizes[i]; > - for (int j = 1; j <= elems; j++) > - { > - struct splay_tree_key_s k; > - k.host_start = (uintptr_t) hostaddrs[i + j]; > - k.host_end = k.host_start + sizes[i + j]; > - splay_tree_key str; > - str = splay_tree_lookup (&acc_dev->mem_map, &k); > - if (str) > - { > - if (finalize) > - { > - if (str->refcount != REFCOUNT_INFINITY) > - str->refcount -= str->virtual_refcount; > - str->virtual_refcount = 0; > - } > - if (str->virtual_refcount > 0) > - { > - if (str->refcount != REFCOUNT_INFINITY) > - str->refcount--; > - str->virtual_refcount--; > - } > - else if (str->refcount > 0 > - && str->refcount != REFCOUNT_INFINITY) > - str->refcount--; > - if (str->refcount == 0) > - gomp_remove_var_async (acc_dev, str, aq); > - } > - } > - i += elems; > - } > break; > > default: ^ permalink raw reply [flat|nested] 11+ messages in thread
* Re: [PATCH] Fix component mappings with derived types for OpenACC 2020-01-24 10:53 ` Tobias Burnus @ 2020-01-28 14:37 ` Julian Brown 2020-05-19 12:23 ` Thomas Schwinge 0 siblings, 1 reply; 11+ messages in thread From: Julian Brown @ 2020-01-28 14:37 UTC (permalink / raw) To: Tobias Burnus; +Cc: gcc-patches, fortran, Thomas Schwinge, Jakub Jelinek On Fri, 24 Jan 2020 10:58:49 +0100 Tobias Burnus <tobias@codesourcery.com> wrote: > Hi Julian, > > the gfortran part is rather obvious and it and the test case look > fine to me → OK. > The oacc-mem.c also looks okay, but I assume Thomas needs to > rubber-stamp it. I understand that Thomas is unavailable for the time being, so won't be able to use his rubber-stamp powers. I added the offending libgomp code to start with though, so I think I can go ahead and commit the patch. I'll hold off for 24 hours though in case there are any objections (Jakub?). Thanks, Julian > On 1/10/20 2:49 AM, Julian Brown wrote: > > This patch fixes a bug with mapping Fortran components (i.e. with > > the manual deep-copy support) which themselves have derived types. > > I've also added a couple of new tests to make sure such mappings > > are lowered correctly, and to check for the case that Tobias found > > in the message: > > > > https://gcc.gnu.org/ml/gcc-patches/2020-01/msg00215.html > > > > The previous incorrect mapping was causing the error condition > > mentioned in that message to fail to trigger, and I think (my!) > > code in libgomp (goacc_exit_data_internal) to handle > > GOMP_MAP_STRUCT specially was papering over the bad mapping also. > > Oops! > > > > I haven't attempted to implement the (harder) sub-copy detection, if > > that is indeed supposed to be forbidden by the spec. This patch > > should get us to the same behaviour in Fortran as in C & C++ though. > > > > Tested with offloading to nvptx, also with the (uncommitted) > > reference-count self-checking patch enabled. > > > > OK? > > > > Thanks, > > > > Julian > > > > ChangeLog > > > > gcc/fortran/ > > * trans-openmp.c (gfc_trans_omp_clauses): Use inner not decl > > for components with derived types. > > > > gcc/testsuite/ > > * gfortran.dg/goacc/mapping-tests-3.f90: New test. > > * gfortran.dg/goacc/mapping-tests-4.f90: New test. > > > > libgomp/ > > * oacc-mem.c (goacc_exit_data_internal): Remove special > > (no-copyback) behaviour for GOMP_MAP_STRUCT. > > > > component-mappings-derived-types-1.diff > > > > commit 5e9d8846bbaa33a9bdb08adcf1ee9f224a8e8fc0 > > Author: Julian Brown<julian@codesourcery.com> > > Date: Wed Jan 8 15:57:46 2020 -0800 > > > > Fix component mappings with derived types for OpenACC > > > > gcc/fortran/ > > * trans-openmp.c (gfc_trans_omp_clauses): Use inner > > not decl for components with derived types. > > > > gcc/testsuite/ > > * gfortran.dg/goacc/mapping-tests-3.f90: New test. > > * gfortran.dg/goacc/mapping-tests-4.f90: New test. > > > > libgomp/ > > * oacc-mem.c (goacc_exit_data_internal): Remove > > special (no-copyback) behaviour for GOMP_MAP_STRUCT. > > > > diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c > > index 9835d2aae3c..efc392d7fa6 100644 > > --- a/gcc/fortran/trans-openmp.c > > +++ b/gcc/fortran/trans-openmp.c > > @@ -2783,9 +2783,9 @@ gfc_trans_omp_clauses (stmtblock_t *block, > > gfc_omp_clauses *clauses, } > > else > > { > > - OMP_CLAUSE_DECL (node) = decl; > > + OMP_CLAUSE_DECL (node) = inner; > > OMP_CLAUSE_SIZE (node) > > - = TYPE_SIZE_UNIT (TREE_TYPE (decl)); > > + = TYPE_SIZE_UNIT (TREE_TYPE (inner)); > > } > > } > > else if (lastcomp->next > > diff --git a/gcc/testsuite/gfortran.dg/goacc/mapping-tests-3.f90 > > b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-3.f90 new file mode > > 100644 index 00000000000..312f596461e > > --- /dev/null > > +++ b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-3.f90 > > @@ -0,0 +1,15 @@ > > +! { dg-options "-fopenacc -fdump-tree-omplower" } > > + > > +subroutine foo > > + type one > > + integer i, j > > + end type > > + type two > > + type(one) A, B > > + end type > > + > > + type(two) x > > + > > + !$acc enter data copyin(x%A) > > +! { dg-final { scan-tree-dump-times "omp target > > oacc_enter_exit_data map\\(struct:x \\\[len: 1\\\]\\) map\\(to:x.a > > \\\[len: \[0-9\]+\\\]\\)" 1 "omplower" } } +end diff --git > > a/gcc/testsuite/gfortran.dg/goacc/mapping-tests-4.f90 > > b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-4.f90 new file mode > > 100644 index 00000000000..6257af942df --- /dev/null > > +++ b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-4.f90 > > @@ -0,0 +1,17 @@ > > +subroutine foo > > + type one > > + integer i, j > > + end type > > + type two > > + type(one) A, B > > + end type > > + > > + type(two) x > > + > > +! This is accepted at present, although it represents a > > probably-unintentional +! overlapping subcopy. > > + !$acc enter data copyin(x%A, x%A%i) > > +! But this raises an error. > > + !$acc enter data copyin(x%A, x%A%i, x%A%i) > > +! { dg-error ".x.a.i. appears more than once in map clauses" > > "" { target "*-*-*" } 15 } +end > > diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c > > index 2d4bba78efd..232683a85f0 100644 > > --- a/libgomp/oacc-mem.c > > +++ b/libgomp/oacc-mem.c > > @@ -1136,38 +1136,6 @@ goacc_exit_data_internal (struct > > gomp_device_descr *acc_dev, size_t mapnum, break; > > > > case GOMP_MAP_STRUCT: > > - { > > - int elems = sizes[i]; > > - for (int j = 1; j <= elems; j++) > > - { > > - struct splay_tree_key_s k; > > - k.host_start = (uintptr_t) hostaddrs[i + j]; > > - k.host_end = k.host_start + sizes[i + j]; > > - splay_tree_key str; > > - str = splay_tree_lookup (&acc_dev->mem_map, &k); > > - if (str) > > - { > > - if (finalize) > > - { > > - if (str->refcount != REFCOUNT_INFINITY) > > - str->refcount -= str->virtual_refcount; > > - str->virtual_refcount = 0; > > - } > > - if (str->virtual_refcount > 0) > > - { > > - if (str->refcount != REFCOUNT_INFINITY) > > - str->refcount--; > > - str->virtual_refcount--; > > - } > > - else if (str->refcount > 0 > > - && str->refcount != REFCOUNT_INFINITY) > > - str->refcount--; > > - if (str->refcount == 0) > > - gomp_remove_var_async (acc_dev, str, aq); > > - } > > - } > > - i += elems; > > - } > > break; > > > > default: ^ permalink raw reply [flat|nested] 11+ messages in thread
* Re: [PATCH] Fix component mappings with derived types for OpenACC 2020-01-28 14:37 ` Julian Brown @ 2020-05-19 12:23 ` Thomas Schwinge 2020-06-04 13:40 ` [PATCH 0/3] OpenACC "exit data" copyout, and Fortran derived-type members Julian Brown 2020-06-04 13:58 ` [PATCH] Fix component mappings with derived types for OpenACC Julian Brown 0 siblings, 2 replies; 11+ messages in thread From: Thomas Schwinge @ 2020-05-19 12:23 UTC (permalink / raw) To: Julian Brown; +Cc: gcc-patches, fortran, Jakub Jelinek, Tobias Burnus Hi! On 2020-01-28T13:41:00+0000, Julian Brown <julian@codesourcery.com> wrote: > On Fri, 24 Jan 2020 10:58:49 +0100 > Tobias Burnus <tobias@codesourcery.com> wrote: >> the gfortran part is rather obvious and it and the test case look >> fine to me → OK. >> The oacc-mem.c also looks okay, but I assume Thomas needs to >> rubber-stamp it. > > I understand that Thomas is unavailable for the time being, so won't be > able to use his rubber-stamp powers. I added the offending libgomp code > to start with though, so I think I can go ahead and commit the patch. > I'll hold off for 24 hours though in case there are any objections > (Jakub?). So, in the end you didn't commit this, and now we've got the "un-fixed" OpenACC/Fortran code generation in the GCC 10.1 release, so have to deal with it in one way or another going forward, regarding libgomp ABI compatibility. (Ideally), we need to make sure that "un-fixed" GCC 10.1-built executables continue to work "as good as before" when dynamically linked/running against "fixed" GCC 10.1+ shared libraries. Do we get such desired behavior by the patch quoted below? In particular: removing the 'GOMP_MAP_STRUCT' handling code, but leaving in the empty 'case GOMP_MAP_STRUCT:' as a no-op, so that we don't run into 'default:' case 'goacc_exit_data_internal UNHANDLED kind'? Is that sufficient? Is my understanding correct that "fixed" GCC won't generate such 'GOMP_MAP_STRUCT' anymore (I have't studied in detail), and this empty 'case GOMP_MAP_STRUCT:' only remains in here for backwards compatibility? In this case, please add a comment to the code, stating this. Otherwise, please add a comment why "do nothing" is appropriate for 'GOMP_MAP_STRUCT'. In particular, for both scenarios, why we don't need to skip the following 'sizes[i]' mappings? Grüße Thomas >> On 1/10/20 2:49 AM, Julian Brown wrote: >> > This patch fixes a bug with mapping Fortran components (i.e. with >> > the manual deep-copy support) which themselves have derived types. >> > I've also added a couple of new tests to make sure such mappings >> > are lowered correctly, and to check for the case that Tobias found >> > in the message: >> > >> > https://gcc.gnu.org/ml/gcc-patches/2020-01/msg00215.html >> > >> > The previous incorrect mapping was causing the error condition >> > mentioned in that message to fail to trigger, and I think (my!) >> > code in libgomp (goacc_exit_data_internal) to handle >> > GOMP_MAP_STRUCT specially was papering over the bad mapping also. >> > Oops! >> > >> > I haven't attempted to implement the (harder) sub-copy detection, if >> > that is indeed supposed to be forbidden by the spec. This patch >> > should get us to the same behaviour in Fortran as in C & C++ though. >> > >> > Tested with offloading to nvptx, also with the (uncommitted) >> > reference-count self-checking patch enabled. >> > >> > OK? >> > >> > Thanks, >> > >> > Julian >> > >> > ChangeLog >> > >> > gcc/fortran/ >> > * trans-openmp.c (gfc_trans_omp_clauses): Use inner not decl >> > for components with derived types. >> > >> > gcc/testsuite/ >> > * gfortran.dg/goacc/mapping-tests-3.f90: New test. >> > * gfortran.dg/goacc/mapping-tests-4.f90: New test. >> > >> > libgomp/ >> > * oacc-mem.c (goacc_exit_data_internal): Remove special >> > (no-copyback) behaviour for GOMP_MAP_STRUCT. >> > >> > component-mappings-derived-types-1.diff >> > >> > commit 5e9d8846bbaa33a9bdb08adcf1ee9f224a8e8fc0 >> > Author: Julian Brown<julian@codesourcery.com> >> > Date: Wed Jan 8 15:57:46 2020 -0800 >> > >> > Fix component mappings with derived types for OpenACC >> > >> > gcc/fortran/ >> > * trans-openmp.c (gfc_trans_omp_clauses): Use inner >> > not decl for components with derived types. >> > >> > gcc/testsuite/ >> > * gfortran.dg/goacc/mapping-tests-3.f90: New test. >> > * gfortran.dg/goacc/mapping-tests-4.f90: New test. >> > >> > libgomp/ >> > * oacc-mem.c (goacc_exit_data_internal): Remove >> > special (no-copyback) behaviour for GOMP_MAP_STRUCT. >> > >> > diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c >> > index 9835d2aae3c..efc392d7fa6 100644 >> > --- a/gcc/fortran/trans-openmp.c >> > +++ b/gcc/fortran/trans-openmp.c >> > @@ -2783,9 +2783,9 @@ gfc_trans_omp_clauses (stmtblock_t *block, >> > gfc_omp_clauses *clauses, } >> > else >> > { >> > - OMP_CLAUSE_DECL (node) = decl; >> > + OMP_CLAUSE_DECL (node) = inner; >> > OMP_CLAUSE_SIZE (node) >> > - = TYPE_SIZE_UNIT (TREE_TYPE (decl)); >> > + = TYPE_SIZE_UNIT (TREE_TYPE (inner)); >> > } >> > } >> > else if (lastcomp->next >> > diff --git a/gcc/testsuite/gfortran.dg/goacc/mapping-tests-3.f90 >> > b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-3.f90 new file mode >> > 100644 index 00000000000..312f596461e >> > --- /dev/null >> > +++ b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-3.f90 >> > @@ -0,0 +1,15 @@ >> > +! { dg-options "-fopenacc -fdump-tree-omplower" } >> > + >> > +subroutine foo >> > + type one >> > + integer i, j >> > + end type >> > + type two >> > + type(one) A, B >> > + end type >> > + >> > + type(two) x >> > + >> > + !$acc enter data copyin(x%A) >> > +! { dg-final { scan-tree-dump-times "omp target >> > oacc_enter_exit_data map\\(struct:x \\\[len: 1\\\]\\) map\\(to:x.a >> > \\\[len: \[0-9\]+\\\]\\)" 1 "omplower" } } +end diff --git >> > a/gcc/testsuite/gfortran.dg/goacc/mapping-tests-4.f90 >> > b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-4.f90 new file mode >> > 100644 index 00000000000..6257af942df --- /dev/null >> > +++ b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-4.f90 >> > @@ -0,0 +1,17 @@ >> > +subroutine foo >> > + type one >> > + integer i, j >> > + end type >> > + type two >> > + type(one) A, B >> > + end type >> > + >> > + type(two) x >> > + >> > +! This is accepted at present, although it represents a >> > probably-unintentional +! overlapping subcopy. >> > + !$acc enter data copyin(x%A, x%A%i) >> > +! But this raises an error. >> > + !$acc enter data copyin(x%A, x%A%i, x%A%i) >> > +! { dg-error ".x.a.i. appears more than once in map clauses" >> > "" { target "*-*-*" } 15 } +end >> > diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c >> > index 2d4bba78efd..232683a85f0 100644 >> > --- a/libgomp/oacc-mem.c >> > +++ b/libgomp/oacc-mem.c >> > @@ -1136,38 +1136,6 @@ goacc_exit_data_internal (struct >> > gomp_device_descr *acc_dev, size_t mapnum, break; >> > >> > case GOMP_MAP_STRUCT: >> > - { >> > - int elems = sizes[i]; >> > - for (int j = 1; j <= elems; j++) >> > - { >> > - struct splay_tree_key_s k; >> > - k.host_start = (uintptr_t) hostaddrs[i + j]; >> > - k.host_end = k.host_start + sizes[i + j]; >> > - splay_tree_key str; >> > - str = splay_tree_lookup (&acc_dev->mem_map, &k); >> > - if (str) >> > - { >> > - if (finalize) >> > - { >> > - if (str->refcount != REFCOUNT_INFINITY) >> > - str->refcount -= str->virtual_refcount; >> > - str->virtual_refcount = 0; >> > - } >> > - if (str->virtual_refcount > 0) >> > - { >> > - if (str->refcount != REFCOUNT_INFINITY) >> > - str->refcount--; >> > - str->virtual_refcount--; >> > - } >> > - else if (str->refcount > 0 >> > - && str->refcount != REFCOUNT_INFINITY) >> > - str->refcount--; >> > - if (str->refcount == 0) >> > - gomp_remove_var_async (acc_dev, str, aq); >> > - } >> > - } >> > - i += elems; >> > - } >> > break; >> > >> > default: ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter ^ permalink raw reply [flat|nested] 11+ messages in thread
* [PATCH 0/3] OpenACC "exit data" copyout, and Fortran derived-type members 2020-05-19 12:23 ` Thomas Schwinge @ 2020-06-04 13:40 ` Julian Brown 2020-06-04 13:40 ` [PATCH 1/3] OpenACC "exit data" copyout for struct members Julian Brown ` (2 more replies) 2020-06-04 13:58 ` [PATCH] Fix component mappings with derived types for OpenACC Julian Brown 1 sibling, 3 replies; 11+ messages in thread From: Julian Brown @ 2020-06-04 13:40 UTC (permalink / raw) To: gcc-patches; +Cc: Thomas Schwinge, Tobias Burnus, Catherine_Moore This is a split version of the patch previously posted here: https://gcc.gnu.org/legacy-ml/gcc-patches/2020-01/msg00512.html There were actually a couple of different issues addressed by that patch, and Thomas raised another (or so). This separated-out version should be more-obviously safe (in terms of maintaining backwards compatibility, where that makes sense). Further commentary on individual patches. Tested (as a series) with offloading to NVPTX. OK? Julian Julian Brown (3): OpenACC "exit data" copyout for struct members Strip GOMP_MAP_STRUCT from OpenACC exit data mappings Fortran derived-type mapping fix gcc/fortran/trans-openmp.c | 4 +- gcc/gimplify.c | 3 +- .../goacc/struct-enter-exit-data-1.c | 35 +++++++++++++++ .../gfortran.dg/goacc/mapping-tests-3.f90 | 15 +++++++ .../gfortran.dg/goacc/mapping-tests-4.f90 | 17 +++++++ libgomp/oacc-mem.c | 32 -------------- .../struct-copyout-1.c | 38 ++++++++++++++++ .../struct-copyout-2.c | 44 +++++++++++++++++++ 8 files changed, 153 insertions(+), 35 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c create mode 100644 gcc/testsuite/gfortran.dg/goacc/mapping-tests-3.f90 create mode 100644 gcc/testsuite/gfortran.dg/goacc/mapping-tests-4.f90 create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-1.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-2.c -- 2.23.0 ^ permalink raw reply [flat|nested] 11+ messages in thread
* [PATCH 1/3] OpenACC "exit data" copyout for struct members 2020-06-04 13:40 ` [PATCH 0/3] OpenACC "exit data" copyout, and Fortran derived-type members Julian Brown @ 2020-06-04 13:40 ` Julian Brown 2020-06-05 16:14 ` Add 'libgomp.oacc-c-c++-common/struct-copyout-{1, 2}.c' (was: [PATCH 1/3] OpenACC "exit data" copyout for struct members) Thomas Schwinge 2020-06-04 13:40 ` [PATCH 2/3] Strip GOMP_MAP_STRUCT from OpenACC exit data mappings Julian Brown 2020-06-04 13:40 ` [PATCH 3/3] Fortran derived-type mapping fix Julian Brown 2 siblings, 1 reply; 11+ messages in thread From: Julian Brown @ 2020-06-04 13:40 UTC (permalink / raw) To: gcc-patches; +Cc: Thomas Schwinge, Tobias Burnus, Catherine_Moore This patch removes unnecessary special-case code in oacc-mem.c which prevented copy-back of structure members on OpenACC "exit data" directives. I've added a couple of new testcases to verify the behaviour (which fail without the patch). (On editing this mail, I notice I omitted to add a comment about GOMP_MAP_STRUCT being a no-op, and/or no longer emitted for OpenACC exit data with the patch later in this series. I'll add such a comment before commit.) OK? Thanks, Julian libgomp/ * oacc-mem.c (goacc_exit_data_internal): Remove special-case (no-copyback) handling for GOMP_MAP_STRUCT. * testsuite/libgomp.oacc-c-c++-common/struct-copyout-1.c: New test. * testsuite/libgomp.oacc-c-c++-common/struct-copyout-2.c: New test. --- libgomp/oacc-mem.c | 32 -------------- .../struct-copyout-1.c | 38 ++++++++++++++++ .../struct-copyout-2.c | 44 +++++++++++++++++++ 3 files changed, 82 insertions(+), 32 deletions(-) create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-1.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-2.c diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index 2d4bba78efd..232683a85f0 100644 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.c @@ -1136,38 +1136,6 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum, break; case GOMP_MAP_STRUCT: - { - int elems = sizes[i]; - for (int j = 1; j <= elems; j++) - { - struct splay_tree_key_s k; - k.host_start = (uintptr_t) hostaddrs[i + j]; - k.host_end = k.host_start + sizes[i + j]; - splay_tree_key str; - str = splay_tree_lookup (&acc_dev->mem_map, &k); - if (str) - { - if (finalize) - { - if (str->refcount != REFCOUNT_INFINITY) - str->refcount -= str->virtual_refcount; - str->virtual_refcount = 0; - } - if (str->virtual_refcount > 0) - { - if (str->refcount != REFCOUNT_INFINITY) - str->refcount--; - str->virtual_refcount--; - } - else if (str->refcount > 0 - && str->refcount != REFCOUNT_INFINITY) - str->refcount--; - if (str->refcount == 0) - gomp_remove_var_async (acc_dev, str, aq); - } - } - i += elems; - } break; default: diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-1.c new file mode 100644 index 00000000000..b86f1c921a9 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-1.c @@ -0,0 +1,38 @@ +#include <assert.h> + +struct str1 { + int a; + int b; +}; + +struct str2 { + int c; + int d; + struct str1 s; +}; + +int +main (int argc, char *argv[]) +{ + struct str2 t; + + t.c = 1; + t.d = 2; + t.s.a = 3; + t.s.b = 4; + + #pragma acc enter data copyin(t.s) + + #pragma acc serial present(t.s) /* { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } } */ + { + t.s.a = 5; + t.s.b = 6; + } + + #pragma acc exit data copyout(t.s) + + assert (t.s.a == 5); + assert (t.s.b == 6); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-2.c new file mode 100644 index 00000000000..4dd8a3a7e17 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-2.c @@ -0,0 +1,44 @@ +#include <assert.h> +#include <stdlib.h> + +struct str1 { + int a; + int b; + int *c; +}; + +#define N 1024 + +int +main (int argc, char *argv[]) +{ + struct str1 s; + + s.a = 1; + s.b = 2; + s.c = (int *) malloc (sizeof (int) * N); + + for (int i = 0; i < N; i++) + s.c[i] = i + 10; + + #pragma acc enter data copyin(s.a, s.b, s.c[0:N]) + + #pragma acc serial present(s.a, s.b, s.c[0:N]) /* { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } } */ + { + s.a = 3; + s.b = 4; + for (int i = 0; i < N; i++) + s.c[i] = i + 20; + } + + #pragma acc exit data copyout(s.a, s.b, s.c[0:N]) + + assert (s.a == 3); + assert (s.b == 4); + for (int i = 0; i < N; i++) + assert (s.c[i] == i + 20); + + free (s.c); + + return 0; +} -- 2.23.0 ^ permalink raw reply [flat|nested] 11+ messages in thread
* Add 'libgomp.oacc-c-c++-common/struct-copyout-{1, 2}.c' (was: [PATCH 1/3] OpenACC "exit data" copyout for struct members) 2020-06-04 13:40 ` [PATCH 1/3] OpenACC "exit data" copyout for struct members Julian Brown @ 2020-06-05 16:14 ` Thomas Schwinge 0 siblings, 0 replies; 11+ messages in thread From: Thomas Schwinge @ 2020-06-05 16:14 UTC (permalink / raw) To: Julian Brown, gcc-patches [-- Attachment #1: Type: text/plain, Size: 879 bytes --] Hi! On 2020-06-04T06:40:53-0700, Julian Brown <julian@codesourcery.com> wrote: > [...] I've added a couple of new testcases to verify the behaviour > (which fail without the patch). Thanks. Given the preparational patches pushed yesterday, these now PASS already. > * testsuite/libgomp.oacc-c-c++-common/struct-copyout-1.c: New test. > * testsuite/libgomp.oacc-c-c++-common/struct-copyout-2.c: New test. Pushed "Add 'libgomp.oacc-c-c++-common/struct-copyout-{1,2}.c'" to master branch in commit 9643f5bbe237764cbefc975e934d1281f47ee3c2, and releases/gcc-10 branch in commit 52d737058897eb438099b57234d41330147d0b6f, see attached. Grüße Thomas ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter [-- Warning: decoded text below may be mangled, UTF-8 assumed --] [-- Attachment #2: 0001-Add-libgomp.oacc-c-c-common-struct-copyout-1-2-.c.patch --] [-- Type: text/x-diff, Size: 2753 bytes --] From 9643f5bbe237764cbefc975e934d1281f47ee3c2 Mon Sep 17 00:00:00 2001 From: Julian Brown <julian@codesourcery.com> Date: Thu, 4 Jun 2020 06:40:53 -0700 Subject: [PATCH] Add 'libgomp.oacc-c-c++-common/struct-copyout-{1,2}.c' libgomp/ * testsuite/libgomp.oacc-c-c++-common/struct-copyout-1.c: New test. * testsuite/libgomp.oacc-c-c++-common/struct-copyout-2.c: New test. Reviewed-by: Thomas Schwinge <thomas@codesourcery.com> --- .../struct-copyout-1.c | 38 ++++++++++++++++ .../struct-copyout-2.c | 44 +++++++++++++++++++ 2 files changed, 82 insertions(+) create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-1.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-2.c diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-1.c new file mode 100644 index 000000000000..b86f1c921a98 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-1.c @@ -0,0 +1,38 @@ +#include <assert.h> + +struct str1 { + int a; + int b; +}; + +struct str2 { + int c; + int d; + struct str1 s; +}; + +int +main (int argc, char *argv[]) +{ + struct str2 t; + + t.c = 1; + t.d = 2; + t.s.a = 3; + t.s.b = 4; + + #pragma acc enter data copyin(t.s) + + #pragma acc serial present(t.s) /* { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } } */ + { + t.s.a = 5; + t.s.b = 6; + } + + #pragma acc exit data copyout(t.s) + + assert (t.s.a == 5); + assert (t.s.b == 6); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-2.c new file mode 100644 index 000000000000..4dd8a3a7e175 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-2.c @@ -0,0 +1,44 @@ +#include <assert.h> +#include <stdlib.h> + +struct str1 { + int a; + int b; + int *c; +}; + +#define N 1024 + +int +main (int argc, char *argv[]) +{ + struct str1 s; + + s.a = 1; + s.b = 2; + s.c = (int *) malloc (sizeof (int) * N); + + for (int i = 0; i < N; i++) + s.c[i] = i + 10; + + #pragma acc enter data copyin(s.a, s.b, s.c[0:N]) + + #pragma acc serial present(s.a, s.b, s.c[0:N]) /* { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } } */ + { + s.a = 3; + s.b = 4; + for (int i = 0; i < N; i++) + s.c[i] = i + 20; + } + + #pragma acc exit data copyout(s.a, s.b, s.c[0:N]) + + assert (s.a == 3); + assert (s.b == 4); + for (int i = 0; i < N; i++) + assert (s.c[i] == i + 20); + + free (s.c); + + return 0; +} -- 2.26.2 [-- Warning: decoded text below may be mangled, UTF-8 assumed --] [-- Attachment #3: 0001-Add-libgomp.oacc-c-c-common-struct-copyout-1-2-..g10.patch --] [-- Type: text/x-diff, Size: 2822 bytes --] From 52d737058897eb438099b57234d41330147d0b6f Mon Sep 17 00:00:00 2001 From: Julian Brown <julian@codesourcery.com> Date: Thu, 4 Jun 2020 06:40:53 -0700 Subject: [PATCH] Add 'libgomp.oacc-c-c++-common/struct-copyout-{1,2}.c' libgomp/ * testsuite/libgomp.oacc-c-c++-common/struct-copyout-1.c: New test. * testsuite/libgomp.oacc-c-c++-common/struct-copyout-2.c: New test. Reviewed-by: Thomas Schwinge <thomas@codesourcery.com> (cherry picked from commit 9643f5bbe237764cbefc975e934d1281f47ee3c2) --- .../struct-copyout-1.c | 38 ++++++++++++++++ .../struct-copyout-2.c | 44 +++++++++++++++++++ 2 files changed, 82 insertions(+) create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-1.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-2.c diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-1.c new file mode 100644 index 000000000000..b86f1c921a98 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-1.c @@ -0,0 +1,38 @@ +#include <assert.h> + +struct str1 { + int a; + int b; +}; + +struct str2 { + int c; + int d; + struct str1 s; +}; + +int +main (int argc, char *argv[]) +{ + struct str2 t; + + t.c = 1; + t.d = 2; + t.s.a = 3; + t.s.b = 4; + + #pragma acc enter data copyin(t.s) + + #pragma acc serial present(t.s) /* { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } } */ + { + t.s.a = 5; + t.s.b = 6; + } + + #pragma acc exit data copyout(t.s) + + assert (t.s.a == 5); + assert (t.s.b == 6); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-2.c new file mode 100644 index 000000000000..4dd8a3a7e175 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-copyout-2.c @@ -0,0 +1,44 @@ +#include <assert.h> +#include <stdlib.h> + +struct str1 { + int a; + int b; + int *c; +}; + +#define N 1024 + +int +main (int argc, char *argv[]) +{ + struct str1 s; + + s.a = 1; + s.b = 2; + s.c = (int *) malloc (sizeof (int) * N); + + for (int i = 0; i < N; i++) + s.c[i] = i + 10; + + #pragma acc enter data copyin(s.a, s.b, s.c[0:N]) + + #pragma acc serial present(s.a, s.b, s.c[0:N]) /* { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } } */ + { + s.a = 3; + s.b = 4; + for (int i = 0; i < N; i++) + s.c[i] = i + 20; + } + + #pragma acc exit data copyout(s.a, s.b, s.c[0:N]) + + assert (s.a == 3); + assert (s.b == 4); + for (int i = 0; i < N; i++) + assert (s.c[i] == i + 20); + + free (s.c); + + return 0; +} -- 2.26.2 ^ permalink raw reply [flat|nested] 11+ messages in thread
* [PATCH 2/3] Strip GOMP_MAP_STRUCT from OpenACC exit data mappings 2020-06-04 13:40 ` [PATCH 0/3] OpenACC "exit data" copyout, and Fortran derived-type members Julian Brown 2020-06-04 13:40 ` [PATCH 1/3] OpenACC "exit data" copyout for struct members Julian Brown @ 2020-06-04 13:40 ` Julian Brown 2020-06-04 13:40 ` [PATCH 3/3] Fortran derived-type mapping fix Julian Brown 2 siblings, 0 replies; 11+ messages in thread From: Julian Brown @ 2020-06-04 13:40 UTC (permalink / raw) To: gcc-patches; +Cc: Thomas Schwinge, Tobias Burnus, Catherine_Moore As flagged by Thomas, the GOMP_MAP_STRUCT mapping is not itself necessary for OpenACC "exit data" directives, and is skipped over (now) in libgomp. We might as well not emit it to start with, in line with the equivalent OpenMP directive. We can keep the "no-op" handling in libgomp for the reason of backward compatibility. I've added a new scan test for the new behaviour. OK? Julian gcc/ * gimplify.c (gimplify_adjust_omp_clauses): Remove GOMP_MAP_STRUCT mapping from OpenACC exit data directives. --- gcc/gimplify.c | 3 +- .../goacc/struct-enter-exit-data-1.c | 35 +++++++++++++++++++ 2 files changed, 37 insertions(+), 1 deletion(-) create mode 100644 gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c diff --git a/gcc/gimplify.c b/gcc/gimplify.c index cb08b26dc65..e14932fafaf 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -10396,7 +10396,8 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, } } else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT - && code == OMP_TARGET_EXIT_DATA) + && (code == OMP_TARGET_EXIT_DATA + || code == OACC_EXIT_DATA)) remove = true; else if (DECL_SIZE (decl) && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST diff --git a/gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c b/gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c new file mode 100644 index 00000000000..4be5674b23e --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c @@ -0,0 +1,35 @@ +/* Check that extraneous GOMP_MAP_STRUCTs are removed from OpenACC exit data + directives. */ + +/* { dg-additional-options "-fdump-tree-omplower" } */ + +#include <stdlib.h> + +struct str { + int a; + int *b; + int *c; + int d; +}; + +#define N 1024 + +int +main (int argc, char *argv[]) +{ + struct str s; + + s.b = (int *) malloc (sizeof (int) * N); + s.c = (int *) malloc (sizeof (int) * N); + + #pragma acc enter data copyin(s.a, s.b[0:N], s.c[0:N], s.d) + /* { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_enter_exit_data map\(struct:s \[len: 4\]\) map\(to:s.a \[len: [0-9]+\]\) map\(alloc:s.b \[len: [0-9]+\]\) map\(alloc:s.c \[len: [0-9]+\]\) map\(to:s.d \[len: [0-9]+\]\) map\(to:\*[_0-9]+ \[len: [0-9]+\]\) map\(attach:s.b \[bias: 0\]\) map\(to:\*[_0-9]+ \[len: [0-9]+\]\) map\(attach:s.c \[bias: 0\]\)} omplower } } */ + + #pragma acc exit data copyout(s.a, s.b[0:N], s.c[0:N], s.d) + /* { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_enter_exit_data map\(from:s.a \[len: [0-9]+\]\) map\(release:s.b \[len: [0-9]+\]\) map\(release:s.c \[len: [0-9]+\]\) map\(from:s.d \[len: [0-9]+\]\) map\(from:\*[_0-9]+ \[len: [0-9]+\]\) map\(detach:s.b \[bias: 0\]\) map\(from:\*[_0-9]+ \[len: [0-9]+\]\) map\(detach:s.c \[bias: 0\]\)} omplower } } */ + + free (s.b); + free (s.c); + + return 0; +} -- 2.23.0 ^ permalink raw reply [flat|nested] 11+ messages in thread
* [PATCH 3/3] Fortran derived-type mapping fix 2020-06-04 13:40 ` [PATCH 0/3] OpenACC "exit data" copyout, and Fortran derived-type members Julian Brown 2020-06-04 13:40 ` [PATCH 1/3] OpenACC "exit data" copyout for struct members Julian Brown 2020-06-04 13:40 ` [PATCH 2/3] Strip GOMP_MAP_STRUCT from OpenACC exit data mappings Julian Brown @ 2020-06-04 13:40 ` Julian Brown 2020-06-05 16:59 ` Thomas Schwinge 2 siblings, 1 reply; 11+ messages in thread From: Julian Brown @ 2020-06-04 13:40 UTC (permalink / raw) To: gcc-patches; +Cc: Thomas Schwinge, Tobias Burnus, Catherine_Moore This patch provides the same fix for Fortran derived type component mappings with derived types as the patch posted previously: https://gcc.gnu.org/legacy-ml/gcc-patches/2020-01/msg00512.html IIUC this part was previously approved by Tobias. OK? Julian gcc/fortran/ * trans-openmp.c (gfc_trans_omp_clauses): Use 'inner' not 'decl' for derived type members which themselves have derived types. gcc/testsuite/ * gfortran.dg/goacc/mapping-tests-3.f90: New test. * gfortran.dg/goacc/mapping-tests-4.f90: New test. --- gcc/fortran/trans-openmp.c | 4 ++-- .../gfortran.dg/goacc/mapping-tests-3.f90 | 15 +++++++++++++++ .../gfortran.dg/goacc/mapping-tests-4.f90 | 17 +++++++++++++++++ 3 files changed, 34 insertions(+), 2 deletions(-) create mode 100644 gcc/testsuite/gfortran.dg/goacc/mapping-tests-3.f90 create mode 100644 gcc/testsuite/gfortran.dg/goacc/mapping-tests-4.f90 diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c index 7e2f6256c43..02c40fdc660 100644 --- a/gcc/fortran/trans-openmp.c +++ b/gcc/fortran/trans-openmp.c @@ -2774,9 +2774,9 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, } else { - OMP_CLAUSE_DECL (node) = decl; + OMP_CLAUSE_DECL (node) = inner; OMP_CLAUSE_SIZE (node) - = TYPE_SIZE_UNIT (TREE_TYPE (decl)); + = TYPE_SIZE_UNIT (TREE_TYPE (inner)); } } else if (lastcomp->next diff --git a/gcc/testsuite/gfortran.dg/goacc/mapping-tests-3.f90 b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-3.f90 new file mode 100644 index 00000000000..312f596461e --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-3.f90 @@ -0,0 +1,15 @@ +! { dg-options "-fopenacc -fdump-tree-omplower" } + +subroutine foo + type one + integer i, j + end type + type two + type(one) A, B + end type + + type(two) x + + !$acc enter data copyin(x%A) +! { dg-final { scan-tree-dump-times "omp target oacc_enter_exit_data map\\(struct:x \\\[len: 1\\\]\\) map\\(to:x.a \\\[len: \[0-9\]+\\\]\\)" 1 "omplower" } } +end diff --git a/gcc/testsuite/gfortran.dg/goacc/mapping-tests-4.f90 b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-4.f90 new file mode 100644 index 00000000000..6257af942df --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-4.f90 @@ -0,0 +1,17 @@ +subroutine foo + type one + integer i, j + end type + type two + type(one) A, B + end type + + type(two) x + +! This is accepted at present, although it represents a probably-unintentional +! overlapping subcopy. + !$acc enter data copyin(x%A, x%A%i) +! But this raises an error. + !$acc enter data copyin(x%A, x%A%i, x%A%i) +! { dg-error ".x.a.i. appears more than once in map clauses" "" { target "*-*-*" } 15 } +end -- 2.23.0 ^ permalink raw reply [flat|nested] 11+ messages in thread
* Re: [PATCH 3/3] Fortran derived-type mapping fix 2020-06-04 13:40 ` [PATCH 3/3] Fortran derived-type mapping fix Julian Brown @ 2020-06-05 16:59 ` Thomas Schwinge 0 siblings, 0 replies; 11+ messages in thread From: Thomas Schwinge @ 2020-06-05 16:59 UTC (permalink / raw) To: Julian Brown; +Cc: Jakub Jelinek, gcc-patches, Tobias Burnus Hi Julian! On 2020-06-04T06:40:55-0700, Julian Brown <julian@codesourcery.com> wrote: > This patch provides the same fix for Fortran derived type component > mappings with derived types as the patch posted previously: > > https://gcc.gnu.org/legacy-ml/gcc-patches/2020-01/msg00512.html > > IIUC this part was previously approved by Tobias. OK? I don't have in-depth knowledge of these Fortran front end OMP parts. Tobias had reviewed and approved your patch in January, and I'll assume his approval still holds. Or are you asking for further review of the patch? I suppose we want this backported to releases/gcc-10 branch, too? > --- /dev/null > +++ b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-3.f90 > @@ -0,0 +1,15 @@ > +! { dg-options "-fopenacc -fdump-tree-omplower" } As '-fopenacc' is the default for 'gfortran.dg/goacc/', you can avoid to repeat it by using '{ dg-additional-options "-fdump-tree-omplower" }'. Just curious: specific reason to scan '-fdump-tree-omplower' -- wouldn't the problem be visible in '-fdump-tree-original' or '-fdump-tree-gimplify' already? > + > +subroutine foo > + type one > + integer i, j > + end type > + type two > + type(one) A, B > + end type > + > + type(two) x > + > + !$acc enter data copyin(x%A) > +! { dg-final { scan-tree-dump-times "omp target oacc_enter_exit_data map\\(struct:x \\\[len: 1\\\]\\) map\\(to:x.a \\\[len: \[0-9\]+\\\]\\)" 1 "omplower" } } > +end > --- /dev/null > +++ b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-4.f90 > @@ -0,0 +1,17 @@ > +subroutine foo > + type one > + integer i, j > + end type > + type two > + type(one) A, B > + end type > + > + type(two) x > + > +! This is accepted at present, although it represents a probably-unintentional > +! overlapping subcopy. > + !$acc enter data copyin(x%A, x%A%i) Should that missing diagnostic get a PR filed, and that one be referenced here? > +! But this raises an error. > + !$acc enter data copyin(x%A, x%A%i, x%A%i) > +! { dg-error ".x.a.i. appears more than once in map clauses" "" { target "*-*-*" } 15 } > +end Not a problem, but unusual to see '"*-*-*"' quoted. Preferable to use relative '.-1' instead of absolute '15'. Grüße Thomas ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter ^ permalink raw reply [flat|nested] 11+ messages in thread
* Re: [PATCH] Fix component mappings with derived types for OpenACC 2020-05-19 12:23 ` Thomas Schwinge 2020-06-04 13:40 ` [PATCH 0/3] OpenACC "exit data" copyout, and Fortran derived-type members Julian Brown @ 2020-06-04 13:58 ` Julian Brown 1 sibling, 0 replies; 11+ messages in thread From: Julian Brown @ 2020-06-04 13:58 UTC (permalink / raw) To: Thomas Schwinge; +Cc: gcc-patches, fortran, Jakub Jelinek, Tobias Burnus On Tue, 19 May 2020 14:23:34 +0200 Thomas Schwinge <thomas@codesourcery.com> wrote: > Hi! > > On 2020-01-28T13:41:00+0000, Julian Brown <julian@codesourcery.com> > wrote: > > On Fri, 24 Jan 2020 10:58:49 +0100 > > Tobias Burnus <tobias@codesourcery.com> wrote: > >> the gfortran part is rather obvious and it and the test case look > >> fine to me → OK. > >> The oacc-mem.c also looks okay, but I assume Thomas needs to > >> rubber-stamp it. > > > > I understand that Thomas is unavailable for the time being, so > > won't be able to use his rubber-stamp powers. I added the offending > > libgomp code to start with though, so I think I can go ahead and > > commit the patch. I'll hold off for 24 hours though in case there > > are any objections (Jakub?). > > So, in the end you didn't commit this, and now we've got the > "un-fixed" OpenACC/Fortran code generation in the GCC 10.1 release, > so have to deal with it in one way or another going forward, > regarding libgomp ABI compatibility. (Ideally), we need to make sure > that "un-fixed" GCC 10.1-built executables continue to work "as good > as before" when dynamically linked/running against "fixed" GCC 10.1+ > shared libraries. > > Do we get such desired behavior by the patch quoted below? In > particular: removing the 'GOMP_MAP_STRUCT' handling code, but leaving > in the empty 'case GOMP_MAP_STRUCT:' as a no-op, so that we don't run > into 'default:' case 'goacc_exit_data_internal UNHANDLED kind'? Is > that sufficient? > > Is my understanding correct that "fixed" GCC won't generate such > 'GOMP_MAP_STRUCT' anymore (I have't studied in detail), and this empty > 'case GOMP_MAP_STRUCT:' only remains in here for backwards > compatibility? In this case, please add a comment to the code, > stating this. Otherwise, please add a comment why "do nothing" is > appropriate for 'GOMP_MAP_STRUCT'. In particular, for both > scenarios, why we don't need to skip the following 'sizes[i]' > mappings? I'm not sure if I got the threading right, but I've now followed up on this discussion here: https://gcc.gnu.org/pipermail/gcc-patches/2020-June/547300.html Thanks, Julian ^ permalink raw reply [flat|nested] 11+ messages in thread
end of thread, other threads:[~2020-06-05 16:59 UTC | newest] Thread overview: 11+ messages (download: mbox.gz / follow: Atom feed) -- links below jump to the message on this page -- 2020-01-10 1:51 [PATCH] Fix component mappings with derived types for OpenACC Julian Brown 2020-01-24 10:53 ` Tobias Burnus 2020-01-28 14:37 ` Julian Brown 2020-05-19 12:23 ` Thomas Schwinge 2020-06-04 13:40 ` [PATCH 0/3] OpenACC "exit data" copyout, and Fortran derived-type members Julian Brown 2020-06-04 13:40 ` [PATCH 1/3] OpenACC "exit data" copyout for struct members Julian Brown 2020-06-05 16:14 ` Add 'libgomp.oacc-c-c++-common/struct-copyout-{1, 2}.c' (was: [PATCH 1/3] OpenACC "exit data" copyout for struct members) Thomas Schwinge 2020-06-04 13:40 ` [PATCH 2/3] Strip GOMP_MAP_STRUCT from OpenACC exit data mappings Julian Brown 2020-06-04 13:40 ` [PATCH 3/3] Fortran derived-type mapping fix Julian Brown 2020-06-05 16:59 ` Thomas Schwinge 2020-06-04 13:58 ` [PATCH] Fix component mappings with derived types for OpenACC Julian Brown
This is a public inbox, see mirroring instructions for how to clone and mirror all data and code used for this inbox; as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).