From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from esa3.mentor.iphmx.com (esa3.mentor.iphmx.com [68.232.137.180]) by sourceware.org (Postfix) with ESMTPS id BFB393856B7D; Sat, 10 Dec 2022 12:10:37 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org BFB393856B7D Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com X-IronPort-AV: E=Sophos;i="5.96,232,1665475200"; d="diff'?scan'208";a="89061868" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa3.mentor.iphmx.com with ESMTP; 10 Dec 2022 04:10:36 -0800 IronPort-SDR: qPirwXY/hjNEedfDHqUFH3IhotymgTtOIWdkpsmKoRYbDDfieGERmjmNDQPvC0Ctdrw1DlhpLi 3p6O8pMRVRdVaEUHM6DeWc71bYymnuWIDA4iycB2Er1tYcjd1X+6pxPgASE4LWbwV9h9fCA67i qdBn/U3MulPIIuyMDp7WALfuQePfsGfOyKDCAOhYiLu4RE5hV5kStOjOTpHTRrD/cGI+ZXNesM ptLiCihL5mE+p6VJE/s+diRGrYzsD+OlDq2yFANr37Ajul6m3PR4YAz03zATwhw8+hPUZWHS1n HzA= Date: Sat, 10 Dec 2022 12:10:22 +0000 From: Julian Brown To: Tobias Burnus CC: , Jakub Jelinek , Subject: Re: [PATCH 2/2] OpenMP: Duplicate checking for map clauses in Fortran (PR107214) Message-ID: <20221210121022.73c2bb32@squid.athome> In-Reply-To: References: <20221020161414.7430-1-julian@codesourcery.com> <20221207190903.78a6b37f@squid.athome> <20221207191355.2e43ea14@squid.athome> Organization: Siemens Embedded X-Mailer: Claws Mail 4.1.1git7 (GTK 3.24.34; x86_64-pc-linux-gnu) MIME-Version: 1.0 Content-Type: multipart/mixed; boundary="MP_/Rh.WXeRsFOHIRUqmfygQ/jE" X-ClientProxiedBy: svr-orw-mbx-11.mgc.mentorg.com (147.34.90.211) To svr-ies-mbx-11.mgc.mentorg.com (139.181.222.11) X-Spam-Status: No, score=-10.5 required=5.0 tests=BAYES_00,GIT_PATCH_0,HEADER_FROM_DIFFERENT_DOMAINS,KAM_DMARC_STATUS,SPF_HELO_PASS,SPF_PASS,TXREP autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org List-Id: --MP_/Rh.WXeRsFOHIRUqmfygQ/jE Content-Type: text/plain; charset="UTF-8" Content-Transfer-Encoding: quoted-printable Content-Disposition: inline On Thu, 8 Dec 2022 13:04:20 +0100 Tobias Burnus wrote: > All in all, I am fine with the patch - but I spotted some issues. >=20 > First, I think you need to set for some error cases mark =3D 0 to avoid > duplicated errors. Namely: >=20 > ! Outputs the error twice ('Symbol =E2=80=98y=E2=80=99 present on mult= iple > clauses') !$omp target has_device_addr(y) firstprivate(y) > block; end block >=20 > * * * >=20 > Additionally, I think it would be good to have besides 'target' + > map/firstprivate (=E2=86=92 error) also a testcase for 'target simd' + > map/firstprivate =E2=86=92 error >=20 > And I think also gives-no-error checks all combined 'target ...' that > take firstprivate should be added, cf. your own patch - possibly with > looking at the original dump (scan-tree-dump) to see that the clause > is properly attached correctly. Example for 'target teams': >=20 > !$omp target teams map(x) firstprivate(x) > block; end block >=20 > (Works but no testcase.) >=20 > * * * >=20 > The following is not diagnosed and gives an ICE: >=20 > !$omp target in_reduction(+: x) private(x) > block; end block > end >=20 > The C testcase properly has: > error: =E2=80=98x=E2=80=99 appears more than once in data-sharing clau= ses >=20 > Note: Using 'firstprivate' instead of 'private' shows the proper > error also in Fortran. >=20 >=20 > The following does not ICE but does not make sense (and is rejected > in C): >=20 > 4 | #pragma omp target private(x) map(x) >=20 > vs. >=20 > !$omp target map(x) private(x) > block; end block >=20 > (The latter produces "#pragma omp target private(x.0) > map(tofrom:*x.0)", ups!) >=20 > * * * >=20 > I also note that 'simd' accepts private such that >=20 > #pragma omp target simd private(x) map(x) > for (int i=3D0; i < 0; i++) > ; >=20 > !$omp target simd map(x) private(x) > do i =3D 1, 0; end do >=20 > is valid. (It is accepted by gcc and gfortran, i.e. it just needs to > be added as testcase.) >=20 > * * * >=20 > I note that C rejects {map(x),firstprivate(x)} + > {has_device_addr(x),is_device_ptr(x)}', but gfortran + your patch > accepts: >=20 > !$omp target map(x) has_device_addr(x) > !$omp target map(x) is_device_ptr(x) >=20 > while >=20 > !$omp target firstprivate(x) has_device_addr(x) > !$omp target firstprivate(x) is_device_ptr(x) >=20 > is rejected =E2=80=93 showing the error message twice. >=20 > Expected: I think it should show an error in all four cases - but > only once. I believe this patch covers all the above cases (hopefully appropriately generalised), at least for Fortran. I haven't attempted to fix any missing cases for C, for now. Re-tested with offloading to NVPTX (with a few supporting patches, as before). Does this look OK now? Thanks, Julian --MP_/Rh.WXeRsFOHIRUqmfygQ/jE Content-Type: text/x-patch Content-Transfer-Encoding: 7bit Content-Disposition: attachment; filename="pr107214-3.diff" commit c5e0cb20a07fcce4822dba2731b7af7b372b3376 Author: Julian Brown Date: Tue Dec 6 23:10:58 2022 +0000 OpenMP: Duplicate checking for map clauses in Fortran (PR107214) This patch adds duplicate checking for OpenMP "map" clauses, taking some cues from the implementation for C in c-typeck.cc:c_finish_omp_clauses (and similar for C++). In addition to the existing use of the "mark" and "comp_mark" bitfields in the gfc_symbol structure, the patch adds several new bits handling duplicate checking within various categories of clause types. If "mark" is being used for map clauses, we need to use different bits for other clauses for cases where "map" and some other clause can refer to the same symbol (e.g. "map(n) shared(n)"). 2022-12-06 Julian Brown gcc/fortran/ PR fortran/107214 * gfortran.h (gfc_symbol): Add data_mark, dev_mark, gen_mark and reduc_mark bitfields. * openmp.cc (resolve_omp_clauses): Use above bitfields to improve duplicate clause detection. gcc/testsuite/ PR fortran/107214 * gfortran.dg/gomp/pr107214.f90: New test. * gfortran.dg/gomp/pr107214-2.f90: New test. * gfortran.dg/gomp/pr107214-3.f90: New test. * gfortran.dg/gomp/pr107214-4.f90: New test. * gfortran.dg/gomp/pr107214-5.f90: New test. * gfortran.dg/gomp/pr107214-6.f90: New test. * gfortran.dg/gomp/pr107214-7.f90: New test. * gfortran.dg/gomp/pr107214-8.f90: New test. diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h index 612a54f4cc49..990ad377e5b2 100644 --- a/gcc/fortran/gfortran.h +++ b/gcc/fortran/gfortran.h @@ -1871,16 +1871,6 @@ typedef struct gfc_symbol gfc_namelist *namelist, *namelist_tail; - /* Change management fields. Symbols that might be modified by the - current statement have the mark member nonzero. Of these symbols, - symbols with old_symbol equal to NULL are symbols created within - the current statement. Otherwise, old_symbol points to a copy of - the old symbol. gfc_new is used in symbol.cc to flag new symbols. - comp_mark is used to indicate variables which have component accesses - in OpenMP/OpenACC directive clauses. */ - struct gfc_symbol *old_symbol; - unsigned mark:1, comp_mark:1, gfc_new:1; - /* The tlink field is used in the front end to carry the module declaration of separate module procedures so that the characteristics can be compared with the corresponding declaration in a submodule. In @@ -1888,6 +1878,28 @@ typedef struct gfc_symbol deferred initialization. */ struct gfc_symbol *tlink; + /* Change management fields. Symbols that might be modified by the + current statement have the mark member nonzero. Of these symbols, + symbols with old_symbol equal to NULL are symbols created within + the current statement. Otherwise, old_symbol points to a copy of + the old symbol. gfc_new is used in symbol.cc to flag new symbols. + comp_mark is used to indicate variables which have component accesses + in OpenMP/OpenACC directive clauses (cf. c-typeck.cc:c_finish_omp_clauses, + map_field_head). + data_mark is used to check duplicate mappings for OpenMP data-sharing + clauses (see firstprivate_head/lastprivate_head in the above function). + dev_mark is used to check duplicate mappings for OpenMP + is_device_ptr/has_device_addr clauses (see is_on_device_head in above + function). + gen_mark is used to check duplicate mappings for OpenMP + use_device_ptr/use_device_addr/private/shared clauses (see generic_head in + above functon). + reduc_mark is used to check duplicate mappings for OpenMP reduction + clauses. */ + struct gfc_symbol *old_symbol; + unsigned mark:1, comp_mark:1, data_mark:1, dev_mark:1, gen_mark:1; + unsigned reduc_mark:1, gfc_new:1; + /* Nonzero if all equivalences associated with this symbol have been processed. */ unsigned equiv_built:1; diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc index 862c649b0b62..70edbe7595d1 100644 --- a/gcc/fortran/openmp.cc +++ b/gcc/fortran/openmp.cc @@ -7131,6 +7131,10 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, continue; n->sym->mark = 0; n->sym->comp_mark = 0; + n->sym->data_mark = 0; + n->sym->dev_mark = 0; + n->sym->gen_mark = 0; + n->sym->reduc_mark = 0; if (n->sym->attr.flavor == FL_VARIABLE || n->sym->attr.proc_pointer || (!code && (!n->sym->attr.dummy || n->sym->ns != ns))) @@ -7199,14 +7203,9 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, && list != OMP_LIST_LASTPRIVATE && list != OMP_LIST_ALIGNED && list != OMP_LIST_DEPEND - && (list != OMP_LIST_MAP || openacc) && list != OMP_LIST_FROM && list != OMP_LIST_TO && (list != OMP_LIST_REDUCTION || !openacc) - && list != OMP_LIST_REDUCTION_INSCAN - && list != OMP_LIST_REDUCTION_TASK - && list != OMP_LIST_IN_REDUCTION - && list != OMP_LIST_TASK_REDUCTION && list != OMP_LIST_ALLOCATE) for (n = omp_clauses->lists[list]; n; n = n->next) { @@ -7218,10 +7217,58 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, for (gfc_ref *ref = n->expr->ref; ref; ref = ref->next) if (ref->type == REF_COMPONENT) component_ref_p = true; - if ((!component_ref_p && n->sym->comp_mark) - || (component_ref_p && n->sym->mark)) - gfc_error ("Symbol %qs has mixed component and non-component " - "accesses at %L", n->sym->name, &n->where); + if ((list == OMP_LIST_IS_DEVICE_PTR + || list == OMP_LIST_HAS_DEVICE_ADDR) + && !component_ref_p) + { + if (n->sym->gen_mark + || n->sym->dev_mark + || n->sym->reduc_mark + || n->sym->mark) + gfc_error ("Symbol %qs present on multiple clauses at %L", + n->sym->name, &n->where); + else + n->sym->dev_mark = 1; + } + else if ((list == OMP_LIST_USE_DEVICE_PTR + || list == OMP_LIST_USE_DEVICE_ADDR + || list == OMP_LIST_PRIVATE + || list == OMP_LIST_SHARED) + && !component_ref_p) + { + if (n->sym->gen_mark || n->sym->dev_mark || n->sym->reduc_mark) + gfc_error ("Symbol %qs present on multiple clauses at %L", + n->sym->name, &n->where); + else + { + n->sym->gen_mark = 1; + /* Set both generic and device bits if we have + use_device_*(x) or shared(x). This allows us to diagnose + "map(x) private(x)" below. */ + if (list != OMP_LIST_PRIVATE) + n->sym->dev_mark = 1; + } + } + else if ((list == OMP_LIST_REDUCTION + || list == OMP_LIST_REDUCTION_TASK + || list == OMP_LIST_REDUCTION_INSCAN + || list == OMP_LIST_IN_REDUCTION + || list == OMP_LIST_TASK_REDUCTION) + && !component_ref_p) + { + /* Attempts to mix reduction types are diagnosed below. */ + if (n->sym->gen_mark || n->sym->dev_mark) + gfc_error ("Symbol %qs present on multiple clauses at %L", + n->sym->name, &n->where); + n->sym->reduc_mark = 1; + } + else if ((!component_ref_p && n->sym->comp_mark) + || (component_ref_p && n->sym->mark)) + { + if (openacc) + gfc_error ("Symbol %qs has mixed component and non-component " + "accesses at %L", n->sym->name, &n->where); + } else if (n->sym->mark) gfc_error ("Symbol %qs present on multiple clauses at %L", n->sym->name, &n->where); @@ -7234,34 +7281,62 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, } } + /* Detect specifically the case where we have "map(x) private(x)" and raise + an error. If we have "...simd" combined directives though, the "private" + applies to the simd part, so this is permitted though. */ + for (n = omp_clauses->lists[OMP_LIST_PRIVATE]; n; n = n->next) + if (n->sym->mark + && n->sym->gen_mark + && !n->sym->dev_mark + && !n->sym->reduc_mark + && code->op != EXEC_OMP_TARGET_SIMD + && code->op != EXEC_OMP_TARGET_PARALLEL_DO_SIMD + && code->op != EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD + && code->op != EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD) + gfc_error ("Symbol %qs present on multiple clauses at %L", + n->sym->name, &n->where); + gcc_assert (OMP_LIST_LASTPRIVATE == OMP_LIST_FIRSTPRIVATE + 1); for (list = OMP_LIST_FIRSTPRIVATE; list <= OMP_LIST_LASTPRIVATE; list++) for (n = omp_clauses->lists[list]; n; n = n->next) - if (n->sym->mark) + if (n->sym->data_mark || n->sym->gen_mark || n->sym->dev_mark) { gfc_error ("Symbol %qs present on multiple clauses at %L", n->sym->name, &n->where); - n->sym->mark = 0; + n->sym->data_mark = n->sym->gen_mark = n->sym->dev_mark = 0; } + else if (n->sym->mark + && code->op != EXEC_OMP_TARGET_TEAMS + && code->op != EXEC_OMP_TARGET_TEAMS_DISTRIBUTE + && code->op != EXEC_OMP_TARGET_TEAMS_LOOP + && code->op != EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD + && code->op != EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO + && code->op != EXEC_OMP_TARGET_PARALLEL + && code->op != EXEC_OMP_TARGET_PARALLEL_DO + && code->op != EXEC_OMP_TARGET_PARALLEL_LOOP + && code->op != EXEC_OMP_TARGET_PARALLEL_DO_SIMD + && code->op != EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD) + gfc_error ("Symbol %qs present on both data and map clauses " + "at %L", n->sym->name, &n->where); for (n = omp_clauses->lists[OMP_LIST_FIRSTPRIVATE]; n; n = n->next) { - if (n->sym->mark) + if (n->sym->data_mark || n->sym->gen_mark || n->sym->dev_mark) gfc_error ("Symbol %qs present on multiple clauses at %L", n->sym->name, &n->where); else - n->sym->mark = 1; + n->sym->data_mark = 1; } for (n = omp_clauses->lists[OMP_LIST_LASTPRIVATE]; n; n = n->next) - n->sym->mark = 0; + n->sym->data_mark = 0; for (n = omp_clauses->lists[OMP_LIST_LASTPRIVATE]; n; n = n->next) { - if (n->sym->mark) + if (n->sym->data_mark || n->sym->gen_mark || n->sym->dev_mark) gfc_error ("Symbol %qs present on multiple clauses at %L", n->sym->name, &n->where); else - n->sym->mark = 1; + n->sym->data_mark = 1; } for (n = omp_clauses->lists[OMP_LIST_ALIGNED]; n; n = n->next) diff --git a/gcc/testsuite/gfortran.dg/gomp/pr107214-2.f90 b/gcc/testsuite/gfortran.dg/gomp/pr107214-2.f90 new file mode 100644 index 000000000000..da47e40f359b --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/pr107214-2.f90 @@ -0,0 +1,6 @@ +integer :: y + +!$omp target has_device_addr(y) firstprivate(y) ! { dg-error "Symbol 'y' present on multiple clauses" } +!$omp end target + +end diff --git a/gcc/testsuite/gfortran.dg/gomp/pr107214-3.f90 b/gcc/testsuite/gfortran.dg/gomp/pr107214-3.f90 new file mode 100644 index 000000000000..526152e11018 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/pr107214-3.f90 @@ -0,0 +1,14 @@ +program p +integer :: y + +!$omp target map(y) firstprivate(y) ! { dg-error "Symbol 'y' present on both data and map clauses" } +y = y + 1 +!$omp end target + +!$omp target simd map(y) firstprivate(y) ! { dg-error "Symbol 'y' present on both data and map clauses" } +do i=1,1 + y = y + 1 +end do +!$omp end target simd + +end program p diff --git a/gcc/testsuite/gfortran.dg/gomp/pr107214-4.f90 b/gcc/testsuite/gfortran.dg/gomp/pr107214-4.f90 new file mode 100644 index 000000000000..b4f343a17acc --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/pr107214-4.f90 @@ -0,0 +1,147 @@ +! { dg-do compile } +! { dg-additional-options "-fdump-tree-original" } + +integer :: x, y + +! EXEC_OMP_TARGET_TEAMS + +!$omp target teams map(x) firstprivate(x) +x = x + 1 +!$omp end target teams + +!$omp target teams map(x) firstprivate(y) +x = y + 1 +!$omp end target teams + +! EXEC_OMP_TARGET_TEAMS_DISTRIBUTE + +!$omp target teams distribute map(x) firstprivate(x) +do i=1,1 + x = x + 1 +end do +!$omp end target teams distribute + +!$omp target teams distribute map(x) firstprivate(y) +do i=1,1 + x = y + 1 +end do +!$omp end target teams distribute + +! EXEC_OMP_TARGET_TEAMS_LOOP + +!$omp target teams loop map(x) firstprivate(x) +do i=1,1 + x = x + 1 +end do +!$omp end target teams loop + +!$omp target teams loop map(x) firstprivate(y) +do i=1,1 + x = y + 1 +end do +!$omp end target teams loop + +! EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD + +!$omp target teams distribute simd map(x) firstprivate(x) +do i=1,1 + x = x + 1 +end do +!$omp end target teams distribute simd + +!$omp target teams distribute simd map(x) firstprivate(y) +do i=1,1 + x = y + 1 +end do +!$omp end target teams distribute simd + +! EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO + +!$omp target teams distribute parallel do map(x) firstprivate(x) +do i=1,1 + x = x + 1 +end do +!$omp end target teams distribute parallel do + +!$omp target teams distribute parallel do map(x) firstprivate(y) +do i=1,1 + x = y + 1 +end do +!$omp end target teams distribute parallel do + +! EXEC_OMP_TARGET_PARALLEL + +!$omp target parallel map(x) firstprivate(x) +x = x + 1 +!$omp end target parallel + +!$omp target parallel map(x) firstprivate(y) +x = y + 1 +!$omp end target parallel + +! EXEC_OMP_TARGET_PARALLEL_DO + +!$omp target parallel do map(x) firstprivate(x) +do i=1,1 + x = x + 1 +end do +!$omp end target parallel do + +!$omp target parallel do map(x) firstprivate(y) +do i=1,1 + x = y + 1 +end do +!$omp end target parallel do + +! EXEC_OMP_TARGET_PARALLEL_LOOP + +!$omp target parallel loop map(x) firstprivate(x) +do i=1,1 + x = x + 1 +end do +!$omp end target parallel loop + +!$omp target parallel loop map(x) firstprivate(y) +do i=1,1 + x = y + 1 +end do +!$omp end target parallel loop + +! EXEC_OMP_TARGET_PARALLEL_DO_SIMD + +!$omp target parallel do simd map(x) firstprivate(x) +do i=1,1 + x = x + 1 +end do +!$omp end target parallel do simd + +!$omp target parallel do simd map(x) firstprivate(y) +do i=1,1 + x = y + 1 +end do +!$omp end target parallel do simd + +! EXEC_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_DO_SIMD + +!$omp target teams distribute parallel do simd map(x) firstprivate(x) +do i=1,1 + x = x + 1 +end do +!$omp end target teams distribute parallel do simd + +!$omp target teams distribute parallel do simd map(x) firstprivate(y) +do i=1,1 + x = y + 1 +end do +!$omp end target teams distribute parallel do simd + +! { dg-final { scan-tree-dump-times {omp target map\(tofrom:x\)} 10 "original" } } +! { dg-final { scan-tree-dump-times {omp target firstprivate\(y\) map\(tofrom:x\)} 10 "original" } } + +! { dg-final { scan-tree-dump-times {omp teams firstprivate\(x\)} 6 "original" } } +! { dg-final { scan-tree-dump-times {omp teams firstprivate\(y\)} 6 "original" } } + +! { dg-final { scan-tree-dump-times {omp parallel firstprivate\(x\)} 6 "original" } } +! { dg-final { scan-tree-dump-times {omp parallel firstprivate\(y\)} 6 "original" } } + +end diff --git a/gcc/testsuite/gfortran.dg/gomp/pr107214-5.f90 b/gcc/testsuite/gfortran.dg/gomp/pr107214-5.f90 new file mode 100644 index 000000000000..08a9f62b0881 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/pr107214-5.f90 @@ -0,0 +1,11 @@ +integer :: x, y + +!$omp target in_reduction(+: x) private(x) ! { dg-error "Symbol 'x' present on multiple clauses" } +x = x + 1 +!$omp end target + +!$omp target in_reduction(+: y) firstprivate(y) ! { dg-error "Symbol 'y' present on both data and map clauses" } +y = y + 1 +!$omp end target + +end diff --git a/gcc/testsuite/gfortran.dg/gomp/pr107214-6.f90 b/gcc/testsuite/gfortran.dg/gomp/pr107214-6.f90 new file mode 100644 index 000000000000..0a1270645518 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/pr107214-6.f90 @@ -0,0 +1,9 @@ +! { dg-do compile } + +integer :: x + +!$omp target map(x) private(x) ! { dg-error "Symbol 'x' present on multiple clauses" } +x = x + 1 +!$omp end target + +end diff --git a/gcc/testsuite/gfortran.dg/gomp/pr107214-7.f90 b/gcc/testsuite/gfortran.dg/gomp/pr107214-7.f90 new file mode 100644 index 000000000000..125d1bc4fed4 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/pr107214-7.f90 @@ -0,0 +1,33 @@ +! { dg-do compile } +! { dg-additional-options "-fdump-tree-original" } + +integer :: x + +!$omp target simd map(x) private(x) +do i=1,1 + x = x + 1 +end do +!$omp end target simd + +!$omp target teams distribute simd map(x) private(x) +do i=1,1 + x = x + 1 +end do +!$omp end target teams distribute simd + +!$omp target parallel do simd map(x) private(x) +do i=1,1 + x = x + 1 +end do +!$omp end target parallel do simd + +!$omp target teams distribute parallel do simd map(x) private(x) +do i=1,1 + x = x + 1 +end do +!$omp end target teams distribute parallel do simd + +! { dg-final { scan-tree-dump-times {omp target map\(tofrom:x\)} 4 "original" } } +! { dg-final { scan-tree-dump-times {(?n)omp simd.* private\(x\)} 4 "original" } } + +end diff --git a/gcc/testsuite/gfortran.dg/gomp/pr107214-8.f90 b/gcc/testsuite/gfortran.dg/gomp/pr107214-8.f90 new file mode 100644 index 000000000000..192c97a33e5c --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/pr107214-8.f90 @@ -0,0 +1,18 @@ +! { dg-do compile } + +integer, allocatable :: x +integer, pointer :: y + +!$omp target map(x) has_device_addr(x) ! { dg-error "Symbol 'x' present on multiple clauses" } +!$omp end target + +!$omp target map(y) is_device_ptr(y) ! { dg-error "Symbol 'y' present on multiple clauses" } +!$omp end target + +!$omp target firstprivate(x) has_device_addr(x) ! { dg-error "Symbol 'x' present on multiple clauses" } +!$omp end target + +!$omp target firstprivate(y) is_device_ptr(y) ! { dg-error "Symbol 'y' present on multiple clauses" } +!$omp end target + +end diff --git a/gcc/testsuite/gfortran.dg/gomp/pr107214.f90 b/gcc/testsuite/gfortran.dg/gomp/pr107214.f90 new file mode 100644 index 000000000000..25949934e840 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/pr107214.f90 @@ -0,0 +1,7 @@ +! { dg-do compile } + +program p + integer, allocatable :: a + !$omp target map(tofrom: a, a) ! { dg-error "Symbol 'a' present on multiple clauses" } + !$omp end target +end --MP_/Rh.WXeRsFOHIRUqmfygQ/jE--