OpenE2K
/
gcc
2
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

Add OpenACC 2.6's no_create 

The clause makes any device code use the local memory address for each
of the variables specified unless the given variable is already present
on the current device.

2019-12-19  Julian Brown  <julian@codesourcery.com>
            Maciej W. Rozycki  <macro@codesourcery.com>
            Tobias Burnus  <tobias@codesourcery.com>
            Thomas Schwinge  <thomas@codesourcery.com>

        gcc/
        * omp-low.c (lower_omp_target): Support GOMP_MAP_NO_ALLOC.
        * tree-pretty-print.c (dump_omp_clause): Likewise.

        gcc/c-family/
        * c-pragma.h (pragma_omp_clause): Add
        PRAGMA_OACC_CLAUSE_NO_CREATE.

        gcc/c/
        * c-parser.c (c_parser_omp_clause_name): Support no_create.
        (c_parser_oacc_data_clause): Likewise.
        (c_parser_oacc_all_clauses): Likewise.
        (OACC_DATA_CLAUSE_MASK, OACC_KERNELS_CLAUSE_MASK)
        (OACC_PARALLEL_CLAUSE_MASK, OACC_SERIAL_CLAUSE_MASK): Add
        PRAGMA_OACC_CLAUSE_NO_CREATE.
        * c-typeck.c (handle_omp_array_sections): Support
        GOMP_MAP_NO_ALLOC.

        gcc/cp/
        * parser.c (cp_parser_omp_clause_name): Support no_create.
        (cp_parser_oacc_data_clause): Likewise.
        (cp_parser_oacc_all_clauses): Likewise.
        (OACC_DATA_CLAUSE_MASK, OACC_KERNELS_CLAUSE_MASK)
        (OACC_PARALLEL_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_NO_CREATE.
        * semantics.c (handle_omp_array_sections): Support no_create.

        gcc/fortran/
        * gfortran.h (gfc_omp_map_op): Add OMP_MAP_NO_ALLOC.
        * openmp.c (omp_mask2): Add OMP_CLAUSE_NO_CREATE.
        (gfc_match_omp_clauses): Support no_create.
        (OACC_PARALLEL_CLAUSES, OACC_KERNELS_CLAUSES)
        (OACC_DATA_CLAUSES): Add OMP_CLAUSE_NO_CREATE.
        * trans-openmp.c (gfc_trans_omp_clauses_1): Support
        OMP_MAP_NO_ALLOC.

        gcc/testsuite/
        * gfortran.dg/goacc/common-block-1.f90: Add no_create-clause tests.
        * gfortran.dg/goacc/common-block-1.f90: Likewise.
        * gfortran.dg/goacc/data-clauses.f95: Likewise.
        * gfortran.dg/goacc/data-tree.f95: Likewise.
        * gfortran.dg/goacc/kernels-tree.f95: Likewise.
        * gfortran.dg/goacc/parallel-tree.f95: Likewise.

        include/
        * gomp-constants.h (gomp_map_kind): Support GOMP_MAP_NO_ALLOC.

        libgomp/
        * target.c (gomp_map_vars_async): Support GOMP_MAP_NO_ALLOC.
        * testsuite/libgomp.oacc-c-c++-common/no_create-1.c: New test.
        * testsuite/libgomp.oacc-c-c++-common/no_create-2.c: New test.
        * testsuite/libgomp.oacc-c-c++-common/no_create-3.c: New test.
        * testsuite/libgomp.oacc-c-c++-common/no_create-4.c: New test.
        * testsuite/libgomp.oacc-c-c++-common/no_create-5.c: New test.
        * testsuite/libgomp.oacc-fortran/no_create-1.f90: New test.
        * testsuite/libgomp.oacc-fortran/no_create-2.f90: New test.
        * testsuite/libgomp.oacc-fortran/no_create-3.F90: New test.

Reviewed-by: Thomas Schwinge <thomas@codesourcery.com>

Co-Authored-By: Maciej W. Rozycki <macro@codesourcery.com>
Co-Authored-By: Thomas Schwinge <thomas@codesourcery.com>
Co-Authored-By: Tobias Burnus <tobias@codesourcery.com>

From-SVN: r279551
This commit is contained in:
Julian Brown 2019-12-18 23:10:08 +00:00 committed by Tobias Burnus
parent 11b8091fb3
commit a6163563f2
34 changed files with 616 additions and 19 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

8
gcc/ChangeLog
View File

@ -1,3 +1,11 @@
2019-12-19 Julian Brown <julian@codesourcery.com>
Maciej W. Rozycki <macro@codesourcery.com>
Tobias Burnus <tobias@codesourcery.com>
Thomas Schwinge <thomas@codesourcery.com>
* omp-low.c (lower_omp_target): Support GOMP_MAP_NO_ALLOC.
* tree-pretty-print.c (dump_omp_clause): Likewise.
2019-12-18 Eric Botcazou <ebotcazou@adacore.com>
* ira.c (ira): Use simple LRA algorithm when not optimizing.

8
gcc/c-family/ChangeLog
View File

@ -1,3 +1,11 @@
2019-12-19 Julian Brown <julian@codesourcery.com>
Maciej W. Rozycki <macro@codesourcery.com>
Tobias Burnus <tobias@codesourcery.com>
Thomas Schwinge <thomas@codesourcery.com>
* c-pragma.h (pragma_omp_clause): Add
PRAGMA_OACC_CLAUSE_NO_CREATE.
2019-12-17 Martin Sebor <msebor@redhat.com>
PR c++/61339

1
gcc/c-family/c-pragma.h
View File

@ -154,6 +154,7 @@ enum pragma_omp_clause {
PRAGMA_OACC_CLAUSE_GANG,
PRAGMA_OACC_CLAUSE_HOST,
PRAGMA_OACC_CLAUSE_INDEPENDENT,
PRAGMA_OACC_CLAUSE_NO_CREATE,
PRAGMA_OACC_CLAUSE_NUM_GANGS,
PRAGMA_OACC_CLAUSE_NUM_WORKERS,
PRAGMA_OACC_CLAUSE_PRESENT,

14
gcc/c/ChangeLog
View File

@ -1,3 +1,17 @@
2019-12-19 Julian Brown <julian@codesourcery.com>
Maciej W. Rozycki <macro@codesourcery.com>
Tobias Burnus <tobias@codesourcery.com>
Thomas Schwinge <thomas@codesourcery.com>
* c-parser.c (c_parser_omp_clause_name): Support no_create.
(c_parser_oacc_data_clause): Likewise.
(c_parser_oacc_all_clauses): Likewise.
(OACC_DATA_CLAUSE_MASK, OACC_KERNELS_CLAUSE_MASK)
(OACC_PARALLEL_CLAUSE_MASK, OACC_SERIAL_CLAUSE_MASK): Add
PRAGMA_OACC_CLAUSE_NO_CREATE.
* c-typeck.c (handle_omp_array_sections): Support
GOMP_MAP_NO_ALLOC.
2019-12-09 David Malcolm <dmalcolm@redhat.com>
* c-objc-common.c (range_label_for_type_mismatch::get_text):

20
gcc/c/c-parser.c
View File

@ -12650,7 +12650,9 @@ c_parser_omp_clause_name (c_parser *parser)
result = PRAGMA_OMP_CLAUSE_MERGEABLE;
break;
case 'n':
if (!strcmp ("nogroup", p))
if (!strcmp ("no_create", p))
result = PRAGMA_OACC_CLAUSE_NO_CREATE;
else if (!strcmp ("nogroup", p))
result = PRAGMA_OMP_CLAUSE_NOGROUP;
else if (!strcmp ("nontemporal", p))
result = PRAGMA_OMP_CLAUSE_NONTEMPORAL;
@ -13113,7 +13115,10 @@ c_parser_omp_var_list_parens (c_parser *parser, enum omp_clause_code kind,
copyout ( variable-list )
create ( variable-list )
delete ( variable-list )
present ( variable-list ) */
present ( variable-list )
OpenACC 2.6:
no_create ( variable-list ) */
static tree
c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind,
@ -13149,6 +13154,9 @@ c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind,
case PRAGMA_OACC_CLAUSE_LINK:
kind = GOMP_MAP_LINK;
break;
case PRAGMA_OACC_CLAUSE_NO_CREATE:
kind = GOMP_MAP_IF_PRESENT;
break;
case PRAGMA_OACC_CLAUSE_PRESENT:
kind = GOMP_MAP_FORCE_PRESENT;
break;
@ -15947,6 +15955,10 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
c_name = "link";
break;
case PRAGMA_OACC_CLAUSE_NO_CREATE:
clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
c_name = "no_create";
break;
case PRAGMA_OACC_CLAUSE_NUM_GANGS:
clauses = c_parser_oacc_single_int_clause (parser,
OMP_CLAUSE_NUM_GANGS,
@ -16415,6 +16427,7 @@ c_parser_oacc_cache (location_t loc, c_parser *parser)
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NO_CREATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT))
static tree
@ -16747,6 +16760,7 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NO_CREATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_GANGS) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_WORKERS) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \
@ -16762,6 +16776,7 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NO_CREATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRIVATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_FIRSTPRIVATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_GANGS) \
@ -16780,6 +16795,7 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NO_CREATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRIVATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_FIRSTPRIVATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \

1
gcc/c/c-typeck.c
View File

@ -13422,6 +13422,7 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
switch (OMP_CLAUSE_MAP_KIND (c))
{
case GOMP_MAP_ALLOC:
case GOMP_MAP_IF_PRESENT:
case GOMP_MAP_TO:
case GOMP_MAP_FROM:
case GOMP_MAP_TOFROM:

12
gcc/cp/ChangeLog
View File

@ -1,3 +1,15 @@
2019-12-19 Julian Brown <julian@codesourcery.com>
Maciej W. Rozycki <macro@codesourcery.com>
Tobias Burnus <tobias@codesourcery.com>
Thomas Schwinge <thomas@codesourcery.com>
* parser.c (cp_parser_omp_clause_name): Support no_create.
(cp_parser_oacc_data_clause): Likewise.
(cp_parser_oacc_all_clauses): Likewise.
(OACC_DATA_CLAUSE_MASK, OACC_KERNELS_CLAUSE_MASK)
(OACC_PARALLEL_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_NO_CREATE.
* semantics.c (handle_omp_array_sections): Support no_create.
2019-12-18 Paolo Carlini <paolo.carlini@oracle.com>
* typeck.c (cxx_sizeof_or_alignof_type): Add location_t parameter

22
gcc/cp/parser.c
View File

@ -33622,7 +33622,9 @@ cp_parser_omp_clause_name (cp_parser *parser)
result = PRAGMA_OMP_CLAUSE_MERGEABLE;
break;
case 'n':
if (!strcmp ("nogroup", p))
if (!strcmp ("no_create", p))
result = PRAGMA_OACC_CLAUSE_NO_CREATE;
else if (!strcmp ("nogroup", p))
result = PRAGMA_OMP_CLAUSE_NOGROUP;
else if (!strcmp ("nontemporal", p))
result = PRAGMA_OMP_CLAUSE_NONTEMPORAL;
@ -33988,7 +33990,10 @@ cp_parser_omp_var_list (cp_parser *parser, enum omp_clause_code kind, tree list)
copyout ( variable-list )
create ( variable-list )
delete ( variable-list )
present ( variable-list ) */
present ( variable-list )
OpenACC 2.6:
no_create ( variable-list ) */
static tree
cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind,
@ -34024,6 +34029,9 @@ cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind,
case PRAGMA_OACC_CLAUSE_LINK:
kind = GOMP_MAP_LINK;
break;
case PRAGMA_OACC_CLAUSE_NO_CREATE:
kind = GOMP_MAP_IF_PRESENT;
break;
case PRAGMA_OACC_CLAUSE_PRESENT:
kind = GOMP_MAP_FORCE_PRESENT;
break;
@ -36586,6 +36594,10 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
c_name = "link";
break;
case PRAGMA_OACC_CLAUSE_NO_CREATE:
clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
c_name = "no_create";
break;
case PRAGMA_OACC_CLAUSE_NUM_GANGS:
code = OMP_CLAUSE_NUM_GANGS;
c_name = "num_gangs";
@ -40391,6 +40403,7 @@ cp_parser_oacc_cache (cp_parser *parser, cp_token *pragma_tok)
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NO_CREATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) )
static tree
@ -40712,6 +40725,7 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NO_CREATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_GANGS) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_WORKERS) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \
@ -40726,8 +40740,9 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_FIRSTPRIVATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_FIRSTPRIVATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NO_CREATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_GANGS) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_WORKERS) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \
@ -40745,6 +40760,7 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NO_CREATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRIVATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_FIRSTPRIVATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \

1
gcc/cp/semantics.c
View File

@ -5292,6 +5292,7 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
switch (OMP_CLAUSE_MAP_KIND (c))
{
case GOMP_MAP_ALLOC:
case GOMP_MAP_IF_PRESENT:
case GOMP_MAP_TO:
case GOMP_MAP_FROM:
case GOMP_MAP_TOFROM:

13
gcc/fortran/ChangeLog
View File

@ -1,3 +1,16 @@
2019-12-19 Julian Brown <julian@codesourcery.com>
Maciej W. Rozycki <macro@codesourcery.com>
Tobias Burnus <tobias@codesourcery.com>
Thomas Schwinge <thomas@codesourcery.com>
* gfortran.h (gfc_omp_map_op): Add OMP_MAP_NO_ALLOC.
* openmp.c (omp_mask2): Add OMP_CLAUSE_NO_CREATE.
(gfc_match_omp_clauses): Support no_create.
(OACC_PARALLEL_CLAUSES, OACC_KERNELS_CLAUSES)
(OACC_DATA_CLAUSES): Add OMP_CLAUSE_NO_CREATE.
* trans-openmp.c (gfc_trans_omp_clauses_1): Support
OMP_MAP_NO_ALLOC.
2019-12-18 Harald Anlauf <anlauf@gmx.de>
PR fortran/70853

1
gcc/fortran/gfortran.h
View File

@ -1192,6 +1192,7 @@ enum gfc_omp_depend_op
enum gfc_omp_map_op
{
OMP_MAP_ALLOC,
OMP_MAP_IF_PRESENT,
OMP_MAP_TO,
OMP_MAP_FROM,
OMP_MAP_TOFROM,

28
gcc/fortran/openmp.c
View File

@ -807,6 +807,7 @@ enum omp_mask2
OMP_CLAUSE_COPY,
OMP_CLAUSE_COPYOUT,
OMP_CLAUSE_CREATE,
OMP_CLAUSE_NO_CREATE,
OMP_CLAUSE_PRESENT,
OMP_CLAUSE_DEVICEPTR,
OMP_CLAUSE_GANG,
@ -1445,6 +1446,11 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
}
break;
case 'n':
if ((mask & OMP_CLAUSE_NO_CREATE)
&& gfc_match ("no_create ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
OMP_MAP_IF_PRESENT, true))
continue;
if ((mask & OMP_CLAUSE_NOGROUP)
&& !c->nogroup
&& gfc_match ("nogroup") == MATCH_YES)
@ -1955,25 +1961,25 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
(omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_NUM_GANGS \
| OMP_CLAUSE_NUM_WORKERS | OMP_CLAUSE_VECTOR_LENGTH | OMP_CLAUSE_REDUCTION \
| OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \
| OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_DEVICEPTR \
| OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE | OMP_CLAUSE_DEFAULT \
| OMP_CLAUSE_WAIT)
| OMP_CLAUSE_CREATE | OMP_CLAUSE_NO_CREATE | OMP_CLAUSE_PRESENT \
| OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE \
| OMP_CLAUSE_DEFAULT | OMP_CLAUSE_WAIT)
#define OACC_KERNELS_CLAUSES \
(omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_NUM_GANGS \
| OMP_CLAUSE_NUM_WORKERS | OMP_CLAUSE_VECTOR_LENGTH | OMP_CLAUSE_DEVICEPTR \
| OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \
| OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_DEFAULT \
| OMP_CLAUSE_WAIT)
| OMP_CLAUSE_CREATE | OMP_CLAUSE_NO_CREATE | OMP_CLAUSE_PRESENT \
| OMP_CLAUSE_DEFAULT | OMP_CLAUSE_WAIT)
#define OACC_SERIAL_CLAUSES \
(omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_REDUCTION \
| OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \
| OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_DEVICEPTR \
| OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE | OMP_CLAUSE_DEFAULT \
| OMP_CLAUSE_WAIT)
| OMP_CLAUSE_CREATE | OMP_CLAUSE_NO_CREATE | OMP_CLAUSE_PRESENT \
| OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE \
| OMP_CLAUSE_DEFAULT | OMP_CLAUSE_WAIT)
#define OACC_DATA_CLAUSES \
(omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_COPY \
| OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT | OMP_CLAUSE_CREATE \
| OMP_CLAUSE_PRESENT)
| OMP_CLAUSE_NO_CREATE | OMP_CLAUSE_PRESENT)
#define OACC_LOOP_CLAUSES \
(omp_mask (OMP_CLAUSE_COLLAPSE) | OMP_CLAUSE_GANG | OMP_CLAUSE_WORKER \
| OMP_CLAUSE_VECTOR | OMP_CLAUSE_SEQ | OMP_CLAUSE_INDEPENDENT \
@ -2509,7 +2515,7 @@ cleanup:
#define OMP_TASKLOOP_CLAUSES \
(omp_mask (OMP_CLAUSE_PRIVATE) | OMP_CLAUSE_FIRSTPRIVATE \
| OMP_CLAUSE_LASTPRIVATE | OMP_CLAUSE_SHARED | OMP_CLAUSE_IF \
| OMP_CLAUSE_DEFAULT | OMP_CLAUSE_UNTIED | OMP_CLAUSE_FINAL \
| OMP_CLAUSE_DEFAULT | OMP_CLAUSE_UNTIED | OMP_CLAUSE_FINAL \
| OMP_CLAUSE_MERGEABLE | OMP_CLAUSE_PRIORITY | OMP_CLAUSE_GRAINSIZE \
| OMP_CLAUSE_NUM_TASKS | OMP_CLAUSE_COLLAPSE | OMP_CLAUSE_NOGROUP)
#define OMP_TARGET_CLAUSES \
@ -2531,7 +2537,7 @@ cleanup:
| OMP_CLAUSE_FROM | OMP_CLAUSE_DEPEND | OMP_CLAUSE_NOWAIT)
#define OMP_TEAMS_CLAUSES \
(omp_mask (OMP_CLAUSE_NUM_TEAMS) | OMP_CLAUSE_THREAD_LIMIT \
| OMP_CLAUSE_DEFAULT | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE \
| OMP_CLAUSE_DEFAULT | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE \
| OMP_CLAUSE_SHARED | OMP_CLAUSE_REDUCTION)
#define OMP_DISTRIBUTE_CLAUSES \
(omp_mask (OMP_CLAUSE_PRIVATE) | OMP_CLAUSE_FIRSTPRIVATE \

3
gcc/fortran/trans-openmp.c
View File

@ -2624,6 +2624,9 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
case OMP_MAP_ALLOC:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALLOC);
break;
case OMP_MAP_IF_PRESENT:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_IF_PRESENT);
break;
case OMP_MAP_TO:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_TO);
break;

2
gcc/omp-low.c
View File

@ -11431,6 +11431,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
case GOMP_MAP_STRUCT:
case GOMP_MAP_ALWAYS_POINTER:
break;
case GOMP_MAP_IF_PRESENT:
case GOMP_MAP_FORCE_ALLOC:
case GOMP_MAP_FORCE_TO:
case GOMP_MAP_FORCE_FROM:
@ -11842,6 +11843,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
switch (tkind)
{
case GOMP_MAP_ALLOC:
case GOMP_MAP_IF_PRESENT:
case GOMP_MAP_TO:
case GOMP_MAP_FROM:
case GOMP_MAP_TOFROM:

12
gcc/testsuite/ChangeLog
View File

@ -1,3 +1,15 @@
2019-12-19 Julian Brown <julian@codesourcery.com>
Maciej W. Rozycki <macro@codesourcery.com>
Tobias Burnus <tobias@codesourcery.com>
Thomas Schwinge <thomas@codesourcery.com>
* gfortran.dg/goacc/common-block-1.f90: Add no_create-clause tests.
* gfortran.dg/goacc/common-block-1.f90: Likewise.
* gfortran.dg/goacc/data-clauses.f95: Likewise.
* gfortran.dg/goacc/data-tree.f95: Likewise.
* gfortran.dg/goacc/kernels-tree.f95: Likewise.
* gfortran.dg/goacc/parallel-tree.f95: Likewise.
2019-12-18 Paolo Carlini <paolo.carlini@oracle.com>
* g++.dg/diagnostic/alignof2.C: New.

3
gcc/testsuite/gfortran.dg/goacc/common-block-1.f90
View File

@ -51,6 +51,9 @@ program test
!$acc data pcopyout(/blockA/, /blockB/, e, v)
!$acc end data
!$acc data no_create(/blockA/, /blockB/, e, v)
!$acc end data
!$acc parallel private(/blockA/, /blockB/, e, v)
!$acc end parallel

3
gcc/testsuite/gfortran.dg/goacc/common-block-2.f90
View File

@ -39,6 +39,9 @@ program test
!$acc data pcopyout(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
!$acc end data
!$acc data no_create(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
!$acc end data
!$acc parallel private(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
!$acc end parallel

21
gcc/testsuite/gfortran.dg/goacc/data-clauses.f95
View File

@ -111,6 +111,27 @@ contains
!$acc end data
!$acc parallel no_create (tip) ! { dg-error "POINTER" }
!$acc end parallel
!$acc parallel no_create (tia) ! { dg-error "ALLOCATABLE" }
!$acc end parallel
!$acc parallel deviceptr (i) no_create (i) ! { dg-error "multiple clauses" }
!$acc end parallel
!$acc parallel copy (i) no_create (i) ! { dg-error "multiple clauses" }
!$acc end parallel
!$acc parallel copyin (i) no_create (i) ! { dg-error "multiple clauses" }
!$acc end parallel
!$acc parallel copyout (i) no_create (i) ! { dg-error "multiple clauses" }
!$acc end parallel
!$acc parallel no_create (i, c, r, ia, ca, ra, asa, rp, ti, vi, aa)
!$acc end parallel
!$acc kernels no_create (i, c, r, ia, ca, ra, asa, rp, ti, vi, aa)
!$acc end kernels
!$acc data no_create (i, c, r, ia, ca, ra, asa, rp, ti, vi, aa)
!$acc end data
!$acc parallel present (tip) ! { dg-error "POINTER" }
!$acc end parallel
!$acc parallel present (tia) ! { dg-error "ALLOCATABLE" }

3
gcc/testsuite/gfortran.dg/goacc/data-tree.f95
View File

@ -7,6 +7,7 @@ program test
logical :: l = .true.
!$acc data if(l) copy(i), copyin(j), copyout(k), create(m) &
!$acc no_create(n) &
!$acc present(o), pcopy(p), pcopyin(r), pcopyout(s), pcreate(t) &
!$acc deviceptr(u)
!$acc end data
@ -19,7 +20,7 @@ end program test
! { dg-final { scan-tree-dump-times "map\\(to:j\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "map\\(from:k\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "map\\(alloc:m\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "map\\(no_alloc:n\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "map\\(force_present:o\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "map\\(tofrom:p\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "map\\(to:r\\)" 1 "original" } }

3
gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95
View File

@ -8,6 +8,7 @@ program test
!$acc kernels if(l) async num_gangs(i) num_workers(i) vector_length(i) &
!$acc copy(i), copyin(j), copyout(k), create(m) &
!$acc no_create(n) &
!$acc present(o), pcopy(p), pcopyin(r), pcopyout(s), pcreate(t) &
!$acc deviceptr(u)
!$acc end kernels
@ -25,7 +26,7 @@ end program test
! { dg-final { scan-tree-dump-times "map\\(to:j\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "map\\(from:k\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "map\\(alloc:m\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "map\\(no_alloc:n\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "map\\(force_present:o\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "map\\(tofrom:p\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "map\\(to:r\\)" 1 "original" } }

3
gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95
View File

@ -9,6 +9,7 @@ program test
!$acc parallel if(l) async num_gangs(i) num_workers(i) vector_length(i) &
!$acc reduction(max:q), copy(i), copyin(j), copyout(k), create(m) &
!$acc no_create(n) &
!$acc present(o), pcopy(p), pcopyin(r), pcopyout(s), pcreate(t) &
!$acc deviceptr(u), private(v), firstprivate(w)
!$acc end parallel
@ -28,7 +29,7 @@ end program test
! { dg-final { scan-tree-dump-times "map\\(to:j\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "map\\(from:k\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "map\\(alloc:m\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "map\\(no_alloc:n\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "map\\(force_present:o\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "map\\(tofrom:p\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "map\\(to:r\\)" 1 "original" } }

3
gcc/tree-pretty-print.c
View File

@ -788,6 +788,9 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
case GOMP_MAP_POINTER:
pp_string (pp, "alloc");
break;
case GOMP_MAP_IF_PRESENT:
pp_string (pp, "no_alloc");
break;
case GOMP_MAP_TO:
case GOMP_MAP_TO_PSET:
pp_string (pp, "to");

7
include/ChangeLog
View File

@ -1,3 +1,10 @@
2019-12-19 Julian Brown <julian@codesourcery.com>
Maciej W. Rozycki <macro@codesourcery.com>
Tobias Burnus <tobias@codesourcery.com>
Thomas Schwinge <thomas@codesourcery.com>
* gomp-constants.h (gomp_map_kind): Support GOMP_MAP_NO_ALLOC.
2019-11-16 Tim Ruehsen <tim.ruehsen@gmx.de>
* demangle.h (struct demangle_component): Add member

2
include/gomp-constants.h
View File

@ -75,6 +75,8 @@ enum gomp_map_kind
GOMP_MAP_DEVICE_RESIDENT = (GOMP_MAP_FLAG_SPECIAL_1 | 1),
/* OpenACC link. */
GOMP_MAP_LINK = (GOMP_MAP_FLAG_SPECIAL_1 | 2),
/* Use device data if present, fall back to host address otherwise. */
GOMP_MAP_IF_PRESENT = (GOMP_MAP_FLAG_SPECIAL_1 | 3),
/* Do not map, copy bits for firstprivate instead. */
GOMP_MAP_FIRSTPRIVATE = (GOMP_MAP_FLAG_SPECIAL | 0),
/* Similarly, but store the value in the pointer rather than

15
libgomp/ChangeLog
View File

@ -1,3 +1,18 @@
2019-12-19 Julian Brown <julian@codesourcery.com>
Maciej W. Rozycki <macro@codesourcery.com>
Tobias Burnus <tobias@codesourcery.com>
Thomas Schwinge <thomas@codesourcery.com>
* target.c (gomp_map_vars_async): Support GOMP_MAP_NO_ALLOC.
* testsuite/libgomp.oacc-c-c++-common/no_create-1.c: New test.
* testsuite/libgomp.oacc-c-c++-common/no_create-2.c: New test.
* testsuite/libgomp.oacc-c-c++-common/no_create-3.c: New test.
* testsuite/libgomp.oacc-c-c++-common/no_create-4.c: New test.
* testsuite/libgomp.oacc-c-c++-common/no_create-5.c: New test.
* testsuite/libgomp.oacc-fortran/no_create-1.f90: New test.
* testsuite/libgomp.oacc-fortran/no_create-2.f90: New test.
* testsuite/libgomp.oacc-fortran/no_create-3.F90: New test.
2019-12-18 Thomas Schwinge <thomas@codesourcery.com>
* oacc-mem.c (goacc_enter_data): Refactor, so that it can be

23
libgomp/target.c
View File

@ -707,6 +707,21 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
{
tgt->list[i].key = NULL;
if ((kind & typemask) == GOMP_MAP_IF_PRESENT)
{
/* Not present, hence, skip entry - including its MAP_POINTER,
when existing. */
tgt->list[i].offset = 0;
if (i + 1 < mapnum
&& ((typemask & get_kind (short_mapkind, kinds, i + 1))
== GOMP_MAP_POINTER))
{
++i;
tgt->list[i].key = NULL;
tgt->list[i].offset = 0;
}
continue;
}
size_t align = (size_t) 1 << (kind >> rshift);
not_found_cnt++;
if (tgt_align < align)
@ -893,6 +908,14 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
cur_node.tgt_offset = n->tgt->tgt_start + n->tgt_offset
+ cur_node.host_start - n->host_start;
continue;
case GOMP_MAP_IF_PRESENT:
/* Not present - otherwise handled above. Skip over its
MAP_POINTER as well. */
if (i + 1 < mapnum
&& ((typemask & get_kind (short_mapkind, kinds, i + 1))
== GOMP_MAP_POINTER))
++i;
continue;
default:
break;
}

49
libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-1.c Normal file
View File

@ -0,0 +1,49 @@
/* Test 'no_create' clause on compute construct, with data present on the
device. */
#include <stdlib.h>
#include <stdio.h>
#include <openacc.h>
#define N 128
int
main (int argc, char *argv[])
{
int var;
int *arr = (int *) malloc (N * sizeof (*arr));
int *devptr[2];
acc_copyin (&var, sizeof (var));
acc_copyin (arr, N * sizeof (*arr));
#pragma acc parallel no_create(var, arr[0:N]) copyout(devptr)
{
devptr[0] = &var;
devptr[1] = &arr[2];
}
if (acc_hostptr (devptr[0]) != (void *) &var)
__builtin_abort ();
if (acc_hostptr (devptr[1]) != (void *) &arr[2])
__builtin_abort ();
acc_delete (&var, sizeof (var));
acc_delete (arr, N * sizeof (*arr));
#if ACC_MEM_SHARED
if (devptr[0] != &var)
__builtin_abort ();
if (devptr[1] != &arr[2])
__builtin_abort ();
#else
if (devptr[0] == &var)
__builtin_abort ();
if (devptr[1] == &arr[2])
__builtin_abort ();
#endif
free (arr);
return 0;
}

30
libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-2.c Normal file
View File

@ -0,0 +1,30 @@
/* Test 'no_create' clause on compute construct, with data not present on the
device. */
#include <stdlib.h>
#include <stdio.h>
#define N 128
int
main (int argc, char *argv[])
{
int var;
int *arr = (int *) malloc (N * sizeof (*arr));
int *devptr[2];
#pragma acc parallel no_create(var, arr[0:N]) copyout(devptr)
{
devptr[0] = &var;
devptr[1] = &arr[2];
}
if (devptr[0] != &var)
__builtin_abort ();
if (devptr[1] != &arr[2])
__builtin_abort ();
free (arr);
return 0;
}

25
libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-3.c Normal file
View File

@ -0,0 +1,25 @@
#include <float.h> /* For FLT_EPSILON. */
#include <math.h> /* For fabs. */
#include <stdlib.h> /* For abort. */
int main()
{
#define N 100
float b[N];
float c[N];
#pragma acc enter data create(b)
#pragma acc parallel loop no_create(b) no_create(c)
for (int i = 0; i < N; ++i)
b[i] = i;
#pragma acc exit data copyout(b)
for (int i = 0; i < N; ++i)
if (fabs (b[i] - i) > 10.0*FLT_EPSILON)
abort ();
return 0;
}

82
libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-4.c Normal file
View File

@ -0,0 +1,82 @@
/* Test 'no_create' clause on 'data' construct and nested compute construct,
with data present on the device. */
#include <stdlib.h>
#include <stdio.h>
#include <openacc.h>
#define N 128
int
main (int argc, char *argv[])
{
int var;
int *arr = (int *) malloc (N * sizeof (*arr));
int *devptr[2];
acc_copyin (&var, sizeof (var));
acc_copyin (arr, N * sizeof (*arr));
#pragma acc data no_create(var, arr[0:N])
{
devptr[0] = (int *) acc_deviceptr (&var);
devptr[1] = (int *) acc_deviceptr (&arr[2]);
if (devptr[0] == NULL)
__builtin_abort ();
if (devptr[1] == NULL)
__builtin_abort ();
if (acc_hostptr (devptr[0]) != (void *) &var)
__builtin_abort ();
if (acc_hostptr (devptr[1]) != (void *) &arr[2])
__builtin_abort ();
#if ACC_MEM_SHARED
if (devptr[0] != &var)
__builtin_abort ();
if (devptr[1] != &arr[2])
__builtin_abort ();
#else
if (devptr[0] == &var)
__builtin_abort ();
if (devptr[1] == &arr[2])
__builtin_abort ();
#endif
#pragma acc parallel copyout(devptr)
{
devptr[0] = &var;
devptr[1] = &arr[2];
}
if (devptr[0] == NULL)
__builtin_abort ();
if (devptr[1] == NULL)
__builtin_abort ();
if (acc_hostptr (devptr[0]) != (void *) &var)
__builtin_abort ();
if (acc_hostptr (devptr[1]) != (void *) &arr[2])
__builtin_abort ();
#if ACC_MEM_SHARED
if (devptr[0] != &var)
__builtin_abort ();
if (devptr[1] != &arr[2])
__builtin_abort ();
#else
if (devptr[0] == &var)
__builtin_abort ();
if (devptr[1] == &arr[2])
__builtin_abort ();
#endif
}
acc_delete (&var, sizeof (var));
acc_delete (arr, N * sizeof (*arr));
free (arr);
return 0;
}

49
libgomp/testsuite/libgomp.oacc-c-c++-common/no_create-5.c Normal file
View File

@ -0,0 +1,49 @@
/* Test 'no_create' clause on 'data' construct and nested compute construct,
with data not present on the device. */
#include <stdlib.h>
#include <stdio.h>
#include <openacc.h>
#define N 128
int
main (int argc, char *argv[])
{
int var;
int *arr = (int *) malloc (N * sizeof (*arr));
int *devptr[2];
#pragma acc data no_create(var, arr[0:N])
{
devptr[0] = (int *) acc_deviceptr (&var);
devptr[1] = (int *) acc_deviceptr (&arr[2]);
#if ACC_MEM_SHARED
if (devptr[0] == NULL)
__builtin_abort ();
if (devptr[1] == NULL)
__builtin_abort ();
#else
if (devptr[0] != NULL)
__builtin_abort ();
if (devptr[1] != NULL)
__builtin_abort ();
#endif
#pragma acc parallel copyout(devptr) // TODO implicit 'copy(var)' -- huh?!
{
devptr[0] = &var;
devptr[1] = &arr[2];
}
if (devptr[0] != &var)
__builtin_abort (); // { dg-xfail-run-if "TODO" { *-*-* } { "-DACC_MEM_SHARED=0" } }
if (devptr[1] != &arr[2])
__builtin_abort ();
}
free (arr);
return 0;
}

39
libgomp/testsuite/libgomp.oacc-fortran/no_create-1.f90 Normal file
View File

@ -0,0 +1,39 @@
! { dg-do run }
! Test no_create clause with data construct when data is present/not present.
program no_create
use openacc
implicit none
logical :: shared_memory
integer, parameter :: n = 512
integer :: myvar, myarr(n)
integer i
shared_memory = .false.
!$acc kernels copyin (shared_memory)
shared_memory = .true.
!$acc end kernels
myvar = 77
do i = 1, n
myarr(i) = 0
end do
!$acc data no_create (myvar, myarr)
if (acc_is_present (myvar) .neqv. shared_memory) stop 10
if (acc_is_present (myarr) .neqv. shared_memory) stop 11
!$acc end data
!$acc enter data copyin (myvar, myarr)
!$acc data no_create (myvar, myarr)
if (acc_is_present (myvar) .eqv. .false.) stop 20
if (acc_is_present (myarr) .eqv. .false.) stop 21
!$acc end data
!$acc exit data copyout (myvar, myarr)
if (myvar .ne. 77) stop 30
do i = 1, n
if (myarr(i) .ne. 0) stop 31
end do
end program no_create

90
libgomp/testsuite/libgomp.oacc-fortran/no_create-2.f90 Normal file
View File

@ -0,0 +1,90 @@
! { dg-do run }
! Test no_create clause with data/parallel constructs.
program no_create
use openacc
implicit none
logical :: shared_memory
integer, parameter :: n = 512
integer :: myvar, myarr(n)
integer i
shared_memory = .false.
!$acc kernels copyin (shared_memory)
shared_memory = .true.
!$acc end kernels
myvar = 55
do i = 1, n
myarr(i) = 0
end do
call do_on_target(myvar, n, myarr)
if (shared_memory) then
if (myvar .ne. 44) stop 10
else
if (myvar .ne. 33) stop 11
end if
do i = 1, n
if (shared_memory) then
if (myarr(i) .ne. i * 2) stop 20
else
if (myarr(i) .ne. i) stop 21
end if
end do
myvar = 55
do i = 1, n
myarr(i) = 0
end do
!$acc enter data copyin(myvar, myarr)
call do_on_target(myvar, n, myarr)
!$acc exit data copyout(myvar, myarr)
if (myvar .ne. 44) stop 30
do i = 1, n
if (myarr(i) .ne. i * 2) stop 31
end do
end program no_create
subroutine do_on_target (var, n, arr)
use openacc
implicit none
integer :: var, n, arr(n)
integer :: i
!$acc data no_create (var, arr)
if (acc_is_present(var)) then
! The no_create clause is meant for partially shared-memory machines. This
! test is written to work on non-shared-memory machines, though this is not
! necessarily a useful way to use the no_create clause in practice.
!$acc parallel !no_create (var)
var = 44
!$acc end parallel
else
var = 33
end if
if (acc_is_present(arr)) then
! The no_create clause is meant for partially shared-memory machines. This
! test is written to work on non-shared-memory machines, though this is not
! necessarily a useful way to use the no_create clause in practice.
!$acc parallel loop !no_create (arr)
do i = 1, n
arr(i) = i * 2
end do
!$acc end parallel loop
else
do i = 1, n
arr(i) = i
end do
end if
!$acc end data
end subroutine do_on_target

39
libgomp/testsuite/libgomp.oacc-fortran/no_create-3.F90 Normal file
View File

@ -0,0 +1,39 @@
! { dg-do run }
program main
use iso_c_binding, only: c_sizeof
use openacc, only: acc_is_present
implicit none
integer i
integer, parameter :: n = 100
real*4 b(n), c(n)
real :: d(n), e(n)
common /BLOCK/ d, e
!$acc enter data create(b) create(d)
if (.not. acc_is_present(b, c_sizeof(b))) stop 1
if (.not. acc_is_present(d, c_sizeof(d))) stop 2
#if !ACC_MEM_SHARED
if (acc_is_present(c, 1) .or. acc_is_present(c, c_sizeof(c))) stop 3
if (acc_is_present(e, 1) .or. acc_is_present(e, c_sizeof(d))) stop 4
#endif
!$acc parallel loop no_create(b) no_create(c) no_create(/BLOCK/)
do i = 1, n
b(i) = i
d(i) = -i
end do
!$acc end parallel loop
if (.not. acc_is_present(b, c_sizeof(b))) stop 5
if (.not. acc_is_present(d, c_sizeof(d))) stop 6
#if !ACC_MEM_SHARED
if (acc_is_present(c, 1) .or. acc_is_present(c, c_sizeof(c))) stop 7
if (acc_is_present(e, 1) .or. acc_is_present(e, c_sizeof(e))) stop 8
#endif
!$acc exit data copyout(b) copyout(d)
if (any(abs(b - [(real(i), i = 1, n)]) > 10*epsilon(b))) stop 9
if (any(abs(d - [(real(-i), i = 1, n)]) > 10*epsilon(d))) stop 10
end program main