* Re: [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior - C++
2018-06-19 16:56 ` Cesar Philippidis
@ 2018-06-19 16:58 ` Cesar Philippidis
2018-06-20 16:16 ` Jakub Jelinek
2018-06-19 16:59 ` [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior - C Cesar Philippidis
` (5 subsequent siblings)
6 siblings, 1 reply; 31+ messages in thread
From: Cesar Philippidis @ 2018-06-19 16:58 UTC (permalink / raw)
To: gcc-patches, Jakub Jelinek
[-- Attachment #1: Type: text/plain, Size: 103 bytes --]
This patch implements the OpenACC 2.5 data clause semantics in the C++ FE.
Is it OK for trunk?
Cesar
[-- Attachment #2: 0003-c++-front-end.patch --]
[-- Type: text/x-patch, Size: 14105 bytes --]
2018-06-19 Chung-Lin Tang <cltang@codesourcery.com>
Thomas Schwinge <thomas@codesourcery.com>
Cesar Philippidis <cesar@codesourcery.com>
gcc/cp/
* parser.c (cp_parser_omp_clause_name): Add support for finalize
and if_present. Make present_or_{copy,copyin,copyout,create} aliases
to their non-present_or_* counterparts. Make 'self' an alias to
PRAGMA_OACC_CLAUSE_HOST.
(cp_parser_oacc_data_clause): Update GOMP mappings for
PRAGMA_OACC_CLAUSE_{COPY,COPYIN,COPYOUT,CREATE,DELETE}. Remove
PRAGMA_OACC_CLAUSE_{SELF,PRESENT_OR_*}.
(cp_parser_oacc_all_clauses): Handle finalize and if_present clauses.
Remove support for present_or_* clauses.
(OACC_KERNELS_CLAUSE_MASK): Remove PRESENT_OR_* clauses.
(OACC_PARALLEL_CLAUSE_MASK): Likewise.
(OACC_DECLARE_CLAUSE_MASK): Likewise.
(OACC_DATA_CLAUSE_MASK): Likewise.
(OACC_ENTER_DATA_CLAUSE_MASK): Remove PRESENT_OR_* clauses.
(OACC_EXIT_DATA_CLAUSE_MASK): Add FINALIZE clause.
(OACC_UPDATE_CLAUSE_MASK): Remove SELF, add IF_PRESENT.
(cp_parser_oacc_declare): Remove PRESENT_OR_* clauses.
* pt.c (tsubst_omp_clauses): Handle IF_PRESENT and FINALIZE.
* semantics.c (finish_omp_clauses): Handle IF_PRESENT and FINALIZE.
From 8d4e76446e128d74d67c5eb367445e6e1f62b46c Mon Sep 17 00:00:00 2001
From: Cesar Philippidis <cesar@codesourcery.com>
Date: Tue, 19 Jun 2018 09:30:20 -0700
Subject: [PATCH 3/7] c++ front end
---
gcc/cp/parser.c | 114 ++++++++++++++-------------------------------
gcc/cp/pt.c | 2 +
gcc/cp/semantics.c | 2 +
3 files changed, 40 insertions(+), 78 deletions(-)
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index c6206fc353c..902ec9c205c 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -31372,6 +31372,8 @@ cp_parser_omp_clause_name (cp_parser *parser)
case 'f':
if (!strcmp ("final", p))
result = PRAGMA_OMP_CLAUSE_FINAL;
+ else if (!strcmp ("finalize", p))
+ result = PRAGMA_OACC_CLAUSE_FINALIZE;
else if (!strcmp ("firstprivate", p))
result = PRAGMA_OMP_CLAUSE_FIRSTPRIVATE;
else if (!strcmp ("from", p))
@@ -31390,7 +31392,9 @@ cp_parser_omp_clause_name (cp_parser *parser)
result = PRAGMA_OACC_CLAUSE_HOST;
break;
case 'i':
- if (!strcmp ("inbranch", p))
+ if (!strcmp ("if_present", p))
+ result = PRAGMA_OACC_CLAUSE_IF_PRESENT;
+ else if (!strcmp ("inbranch", p))
result = PRAGMA_OMP_CLAUSE_INBRANCH;
else if (!strcmp ("independent", p))
result = PRAGMA_OACC_CLAUSE_INDEPENDENT;
@@ -31440,16 +31444,16 @@ cp_parser_omp_clause_name (cp_parser *parser)
result = PRAGMA_OACC_CLAUSE_PRESENT;
else if (!strcmp ("present_or_copy", p)
|| !strcmp ("pcopy", p))
- result = PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY;
+ result = PRAGMA_OACC_CLAUSE_COPY;
else if (!strcmp ("present_or_copyin", p)
|| !strcmp ("pcopyin", p))
- result = PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN;
+ result = PRAGMA_OACC_CLAUSE_COPYIN;
else if (!strcmp ("present_or_copyout", p)
|| !strcmp ("pcopyout", p))
- result = PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT;
+ result = PRAGMA_OACC_CLAUSE_COPYOUT;
else if (!strcmp ("present_or_create", p)
|| !strcmp ("pcreate", p))
- result = PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE;
+ result = PRAGMA_OACC_CLAUSE_CREATE;
else if (!strcmp ("priority", p))
result = PRAGMA_OMP_CLAUSE_PRIORITY;
else if (!strcmp ("proc_bind", p))
@@ -31466,8 +31470,8 @@ cp_parser_omp_clause_name (cp_parser *parser)
result = PRAGMA_OMP_CLAUSE_SCHEDULE;
else if (!strcmp ("sections", p))
result = PRAGMA_OMP_CLAUSE_SECTIONS;
- else if (!strcmp ("self", p))
- result = PRAGMA_OACC_CLAUSE_SELF;
+ else if (!strcmp ("self", p)) /* "self" is a synonym for "host". */
+ result = PRAGMA_OACC_CLAUSE_HOST;
else if (!strcmp ("seq", p))
result = PRAGMA_OACC_CLAUSE_SEQ;
else if (!strcmp ("shared", p))
@@ -31727,15 +31731,7 @@ 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_or_copy ( variable-list )
- pcopy ( variable-list )
- present_or_copyin ( variable-list )
- pcopyin ( variable-list )
- present_or_copyout ( variable-list )
- pcopyout ( variable-list )
- present_or_create ( variable-list )
- pcreate ( variable-list ) */
+ present ( variable-list ) */
static tree
cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind,
@@ -31745,19 +31741,19 @@ cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind,
switch (c_kind)
{
case PRAGMA_OACC_CLAUSE_COPY:
- kind = GOMP_MAP_FORCE_TOFROM;
+ kind = GOMP_MAP_TOFROM;
break;
case PRAGMA_OACC_CLAUSE_COPYIN:
- kind = GOMP_MAP_FORCE_TO;
+ kind = GOMP_MAP_TO;
break;
case PRAGMA_OACC_CLAUSE_COPYOUT:
- kind = GOMP_MAP_FORCE_FROM;
+ kind = GOMP_MAP_FROM;
break;
case PRAGMA_OACC_CLAUSE_CREATE:
- kind = GOMP_MAP_FORCE_ALLOC;
+ kind = GOMP_MAP_ALLOC;
break;
case PRAGMA_OACC_CLAUSE_DELETE:
- kind = GOMP_MAP_DELETE;
+ kind = GOMP_MAP_RELEASE;
break;
case PRAGMA_OACC_CLAUSE_DEVICE:
kind = GOMP_MAP_FORCE_TO;
@@ -31766,7 +31762,6 @@ cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind,
kind = GOMP_MAP_DEVICE_RESIDENT;
break;
case PRAGMA_OACC_CLAUSE_HOST:
- case PRAGMA_OACC_CLAUSE_SELF:
kind = GOMP_MAP_FORCE_FROM;
break;
case PRAGMA_OACC_CLAUSE_LINK:
@@ -31775,18 +31770,6 @@ cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind,
case PRAGMA_OACC_CLAUSE_PRESENT:
kind = GOMP_MAP_FORCE_PRESENT;
break;
- case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY:
- kind = GOMP_MAP_TOFROM;
- break;
- case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN:
- kind = GOMP_MAP_TO;
- break;
- case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT:
- kind = GOMP_MAP_FROM;
- break;
- case PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE:
- kind = GOMP_MAP_ALLOC;
- break;
default:
gcc_unreachable ();
}
@@ -31825,8 +31808,9 @@ cp_parser_oacc_data_clause_deviceptr (cp_parser *parser, tree list)
return list;
}
-/* OpenACC 2.0:
+/* OpenACC 2.5:
auto
+ finalize
independent
nohost
seq */
@@ -33791,6 +33775,11 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
c_name = "device_resident";
break;
+ case PRAGMA_OACC_CLAUSE_FINALIZE:
+ clauses = cp_parser_oacc_simple_clause (parser, OMP_CLAUSE_FINALIZE,
+ clauses, here);
+ c_name = "finalize";
+ break;
case PRAGMA_OACC_CLAUSE_FIRSTPRIVATE:
clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_FIRSTPRIVATE,
clauses);
@@ -33809,6 +33798,12 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
clauses = cp_parser_omp_clause_if (parser, clauses, here, false);
c_name = "if";
break;
+ case PRAGMA_OACC_CLAUSE_IF_PRESENT:
+ clauses = cp_parser_oacc_simple_clause (parser,
+ OMP_CLAUSE_IF_PRESENT,
+ clauses, here);
+ c_name = "if_present";
+ break;
case PRAGMA_OACC_CLAUSE_INDEPENDENT:
clauses = cp_parser_oacc_simple_clause (parser,
OMP_CLAUSE_INDEPENDENT,
@@ -33835,22 +33830,6 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
c_name = "present";
break;
- case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY:
- clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
- c_name = "present_or_copy";
- break;
- case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN:
- clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
- c_name = "present_or_copyin";
- break;
- case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT:
- clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
- c_name = "present_or_copyout";
- break;
- case PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE:
- clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
- c_name = "present_or_create";
- break;
case PRAGMA_OACC_CLAUSE_PRIVATE:
clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_PRIVATE,
clauses);
@@ -33860,10 +33839,6 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
clauses = cp_parser_omp_clause_reduction (parser, clauses);
c_name = "reduction";
break;
- case PRAGMA_OACC_CLAUSE_SELF:
- clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
- c_name = "self";
- break;
case PRAGMA_OACC_CLAUSE_SEQ:
clauses = cp_parser_oacc_simple_clause (parser, OMP_CLAUSE_SEQ,
clauses, here);
@@ -36799,11 +36774,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_PRESENT) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE))
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) )
static tree
cp_parser_oacc_data (cp_parser *parser, cp_token *pragma_tok, bool *if_p)
@@ -36858,11 +36829,7 @@ cp_parser_oacc_host_data (cp_parser *parser, cp_token *pragma_tok, bool *if_p)
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_LINK) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE))
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) )
static tree
cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok)
@@ -36895,8 +36862,8 @@ cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok)
switch (OMP_CLAUSE_MAP_KIND (t))
{
case GOMP_MAP_FIRSTPRIVATE_POINTER:
- case GOMP_MAP_FORCE_ALLOC:
- case GOMP_MAP_FORCE_TO:
+ case GOMP_MAP_ALLOC:
+ case GOMP_MAP_TO:
case GOMP_MAP_FORCE_DEVICEPTR:
case GOMP_MAP_DEVICE_RESIDENT:
break;
@@ -37007,8 +36974,6 @@ cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok)
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
#define OACC_EXIT_DATA_CLAUSE_MASK \
@@ -37016,6 +36981,7 @@ cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok)
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DELETE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_FINALIZE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
static tree
@@ -37128,10 +37094,6 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name,
| (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) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
@@ -37148,10 +37110,6 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name,
| (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) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRIVATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH) \
@@ -37212,7 +37170,7 @@ cp_parser_oacc_kernels_parallel (cp_parser *parser, cp_token *pragma_tok,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_HOST) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SELF) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF_PRESENT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT))
static tree
diff --git a/gcc/cp/pt.c b/gcc/cp/pt.c
index 1ecc6fb373d..d2f98325a0a 100644
--- a/gcc/cp/pt.c
+++ b/gcc/cp/pt.c
@@ -16110,6 +16110,8 @@ tsubst_omp_clauses (tree clauses, enum c_omp_region_type ort,
case OMP_CLAUSE_INDEPENDENT:
case OMP_CLAUSE_AUTO:
case OMP_CLAUSE_SEQ:
+ case OMP_CLAUSE_IF_PRESENT:
+ case OMP_CLAUSE_FINALIZE:
break;
default:
gcc_unreachable ();
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index 2356940a6bb..14958089e45 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -7093,6 +7093,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
case OMP_CLAUSE_AUTO:
case OMP_CLAUSE_INDEPENDENT:
case OMP_CLAUSE_SEQ:
+ case OMP_CLAUSE_IF_PRESENT:
+ case OMP_CLAUSE_FINALIZE:
break;
case OMP_CLAUSE_TILE:
--
2.17.1
^ permalink raw reply [flat|nested] 31+ messages in thread* Re: [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior - C++
2018-06-19 16:58 ` [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior - C++ Cesar Philippidis
@ 2018-06-20 16:16 ` Jakub Jelinek
0 siblings, 0 replies; 31+ messages in thread
From: Jakub Jelinek @ 2018-06-20 16:16 UTC (permalink / raw)
To: Cesar Philippidis; +Cc: gcc-patches
On Tue, Jun 19, 2018 at 09:58:26AM -0700, Cesar Philippidis wrote:
> This patch implements the OpenACC 2.5 data clause semantics in the C++ FE.
>
> Is it OK for trunk?
>
> Cesar
> 2018-06-19 Chung-Lin Tang <cltang@codesourcery.com>
> Thomas Schwinge <thomas@codesourcery.com>
> Cesar Philippidis <cesar@codesourcery.com>
>
> gcc/cp/
> * parser.c (cp_parser_omp_clause_name): Add support for finalize
> and if_present. Make present_or_{copy,copyin,copyout,create} aliases
> to their non-present_or_* counterparts. Make 'self' an alias to
> PRAGMA_OACC_CLAUSE_HOST.
> (cp_parser_oacc_data_clause): Update GOMP mappings for
> PRAGMA_OACC_CLAUSE_{COPY,COPYIN,COPYOUT,CREATE,DELETE}. Remove
> PRAGMA_OACC_CLAUSE_{SELF,PRESENT_OR_*}.
> (cp_parser_oacc_all_clauses): Handle finalize and if_present clauses.
> Remove support for present_or_* clauses.
> (OACC_KERNELS_CLAUSE_MASK): Remove PRESENT_OR_* clauses.
> (OACC_PARALLEL_CLAUSE_MASK): Likewise.
> (OACC_DECLARE_CLAUSE_MASK): Likewise.
> (OACC_DATA_CLAUSE_MASK): Likewise.
> (OACC_ENTER_DATA_CLAUSE_MASK): Remove PRESENT_OR_* clauses.
> (OACC_EXIT_DATA_CLAUSE_MASK): Add FINALIZE clause.
> (OACC_UPDATE_CLAUSE_MASK): Remove SELF, add IF_PRESENT.
> (cp_parser_oacc_declare): Remove PRESENT_OR_* clauses.
> * pt.c (tsubst_omp_clauses): Handle IF_PRESENT and FINALIZE.
> * semantics.c (finish_omp_clauses): Handle IF_PRESENT and FINALIZE.
Ok, thanks.
Jakub
^ permalink raw reply [flat|nested] 31+ messages in thread
* Re: [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior - C
2018-06-19 16:56 ` Cesar Philippidis
2018-06-19 16:58 ` [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior - C++ Cesar Philippidis
@ 2018-06-19 16:59 ` Cesar Philippidis
2018-06-20 16:14 ` Jakub Jelinek
2018-06-19 17:00 ` [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior - Fortran Cesar Philippidis
` (4 subsequent siblings)
6 siblings, 1 reply; 31+ messages in thread
From: Cesar Philippidis @ 2018-06-19 16:59 UTC (permalink / raw)
To: gcc-patches, Jakub Jelinek
[-- Attachment #1: Type: text/plain, Size: 101 bytes --]
This patch implements the OpenACC 2.5 data clause semantics in the C FE.
Is it OK for trunk?
Cesar
[-- Attachment #2: 0004-c-front-end.patch --]
[-- Type: text/x-patch, Size: 13818 bytes --]
2018-06-19 Chung-Lin Tang <cltang@codesourcery.com>
Thomas Schwinge <thomas@codesourcery.com>
Cesar Philippidis <cesar@codesourcery.com>
gcc/c/
* c-parser.c (c_parser_omp_clause_name): Add support for finalize
and if_present. Make present_or_{copy,copyin,copyout,create} aliases
to their non-present_or_* counterparts. Make 'self' an alias to
PRAGMA_OACC_CLAUSE_HOST.
(c_parser_oacc_data_clause): Update GOMP mappings for
PRAGMA_OACC_CLAUSE_{COPY,COPYIN,COPYOUT,CREATE,DELETE}. Remove
PRAGMA_OACC_CLAUSE_{SELF,PRESENT_OR_*}.
(c_parser_oacc_all_clauses): Handle finalize and if_present clauses.
Remove support for present_or_* clauses.
(OACC_KERNELS_CLAUSE_MASK): Remove PRESENT_OR_* clauses.
(OACC_PARALLEL_CLAUSE_MASK): Likewise.
(OACC_DECLARE_CLAUSE_MASK): Likewise.
(OACC_DATA_CLAUSE_MASK): Likewise.
(OACC_ENTER_DATA_CLAUSE_MASK): Remove PRESENT_OR_* clauses.
(OACC_EXIT_DATA_CLAUSE_MASK): Add FINALIZE clause.
(OACC_UPDATE_CLAUSE_MASK): Remove SELF, add IF_PRESENT.
(c_parser_oacc_declare): Remove PRESENT_OR_* clauses.
* c-typeck.c (c_finish_omp_clauses): Handle IF_PRESENT and FINALIZE.
From 03ce56dc85f6b35e873688d482fd0b9570242c3f Mon Sep 17 00:00:00 2001
From: Cesar Philippidis <cesar@codesourcery.com>
Date: Tue, 19 Jun 2018 09:31:14 -0700
Subject: [PATCH 4/7] c front end
---
gcc/c/c-parser.c | 115 +++++++++++++++--------------------------------
gcc/c/c-typeck.c | 2 +
2 files changed, 39 insertions(+), 78 deletions(-)
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index 6b41a615dbd..7a926285f3a 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -11259,6 +11259,8 @@ c_parser_omp_clause_name (c_parser *parser)
case 'f':
if (!strcmp ("final", p))
result = PRAGMA_OMP_CLAUSE_FINAL;
+ else if (!strcmp ("finalize", p))
+ result = PRAGMA_OACC_CLAUSE_FINALIZE;
else if (!strcmp ("firstprivate", p))
result = PRAGMA_OMP_CLAUSE_FIRSTPRIVATE;
else if (!strcmp ("from", p))
@@ -11277,7 +11279,9 @@ c_parser_omp_clause_name (c_parser *parser)
result = PRAGMA_OACC_CLAUSE_HOST;
break;
case 'i':
- if (!strcmp ("inbranch", p))
+ if (!strcmp ("if_present", p))
+ result = PRAGMA_OACC_CLAUSE_IF_PRESENT;
+ else if (!strcmp ("inbranch", p))
result = PRAGMA_OMP_CLAUSE_INBRANCH;
else if (!strcmp ("independent", p))
result = PRAGMA_OACC_CLAUSE_INDEPENDENT;
@@ -11325,18 +11329,20 @@ c_parser_omp_clause_name (c_parser *parser)
result = PRAGMA_OMP_CLAUSE_PARALLEL;
else if (!strcmp ("present", p))
result = PRAGMA_OACC_CLAUSE_PRESENT;
+ /* As of OpenACC 2.5, these are now aliases of the non-present_or
+ clauses. */
else if (!strcmp ("present_or_copy", p)
|| !strcmp ("pcopy", p))
- result = PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY;
+ result = PRAGMA_OACC_CLAUSE_COPY;
else if (!strcmp ("present_or_copyin", p)
|| !strcmp ("pcopyin", p))
- result = PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN;
+ result = PRAGMA_OACC_CLAUSE_COPYIN;
else if (!strcmp ("present_or_copyout", p)
|| !strcmp ("pcopyout", p))
- result = PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT;
+ result = PRAGMA_OACC_CLAUSE_COPYOUT;
else if (!strcmp ("present_or_create", p)
|| !strcmp ("pcreate", p))
- result = PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE;
+ result = PRAGMA_OACC_CLAUSE_CREATE;
else if (!strcmp ("priority", p))
result = PRAGMA_OMP_CLAUSE_PRIORITY;
else if (!strcmp ("private", p))
@@ -11355,6 +11361,8 @@ c_parser_omp_clause_name (c_parser *parser)
result = PRAGMA_OMP_CLAUSE_SCHEDULE;
else if (!strcmp ("sections", p))
result = PRAGMA_OMP_CLAUSE_SECTIONS;
+ else if (!strcmp ("self", p)) /* "self" is a synonym for "host". */
+ result = PRAGMA_OACC_CLAUSE_HOST;
else if (!strcmp ("seq", p))
result = PRAGMA_OACC_CLAUSE_SEQ;
else if (!strcmp ("shared", p))
@@ -11363,8 +11371,6 @@ c_parser_omp_clause_name (c_parser *parser)
result = PRAGMA_OMP_CLAUSE_SIMD;
else if (!strcmp ("simdlen", p))
result = PRAGMA_OMP_CLAUSE_SIMDLEN;
- else if (!strcmp ("self", p))
- result = PRAGMA_OACC_CLAUSE_SELF;
break;
case 't':
if (!strcmp ("taskgroup", p))
@@ -11646,15 +11652,7 @@ 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_or_copy ( variable-list )
- pcopy ( variable-list )
- present_or_copyin ( variable-list )
- pcopyin ( variable-list )
- present_or_copyout ( variable-list )
- pcopyout ( variable-list )
- present_or_create ( variable-list )
- pcreate ( variable-list ) */
+ present ( variable-list ) */
static tree
c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind,
@@ -11664,19 +11662,19 @@ c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind,
switch (c_kind)
{
case PRAGMA_OACC_CLAUSE_COPY:
- kind = GOMP_MAP_FORCE_TOFROM;
+ kind = GOMP_MAP_TOFROM;
break;
case PRAGMA_OACC_CLAUSE_COPYIN:
- kind = GOMP_MAP_FORCE_TO;
+ kind = GOMP_MAP_TO;
break;
case PRAGMA_OACC_CLAUSE_COPYOUT:
- kind = GOMP_MAP_FORCE_FROM;
+ kind = GOMP_MAP_FROM;
break;
case PRAGMA_OACC_CLAUSE_CREATE:
- kind = GOMP_MAP_FORCE_ALLOC;
+ kind = GOMP_MAP_ALLOC;
break;
case PRAGMA_OACC_CLAUSE_DELETE:
- kind = GOMP_MAP_DELETE;
+ kind = GOMP_MAP_RELEASE;
break;
case PRAGMA_OACC_CLAUSE_DEVICE:
kind = GOMP_MAP_FORCE_TO;
@@ -11685,7 +11683,6 @@ c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind,
kind = GOMP_MAP_DEVICE_RESIDENT;
break;
case PRAGMA_OACC_CLAUSE_HOST:
- case PRAGMA_OACC_CLAUSE_SELF:
kind = GOMP_MAP_FORCE_FROM;
break;
case PRAGMA_OACC_CLAUSE_LINK:
@@ -11694,18 +11691,6 @@ c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind,
case PRAGMA_OACC_CLAUSE_PRESENT:
kind = GOMP_MAP_FORCE_PRESENT;
break;
- case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY:
- kind = GOMP_MAP_TOFROM;
- break;
- case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN:
- kind = GOMP_MAP_TO;
- break;
- case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT:
- kind = GOMP_MAP_FROM;
- break;
- case PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE:
- kind = GOMP_MAP_ALLOC;
- break;
default:
gcc_unreachable ();
}
@@ -12597,8 +12582,9 @@ c_parser_oacc_shape_clause (c_parser *parser, omp_clause_code kind,
return list;
}
-/* OpenACC:
+/* OpenACC 2.5:
auto
+ finalize
independent
nohost
seq */
@@ -13955,6 +13941,11 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
c_name = "device_resident";
break;
+ case PRAGMA_OACC_CLAUSE_FINALIZE:
+ clauses = c_parser_oacc_simple_clause (parser, OMP_CLAUSE_FINALIZE,
+ clauses);
+ c_name = "finalize";
+ break;
case PRAGMA_OACC_CLAUSE_FIRSTPRIVATE:
clauses = c_parser_omp_clause_firstprivate (parser, clauses);
c_name = "firstprivate";
@@ -13972,6 +13963,11 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
clauses = c_parser_omp_clause_if (parser, clauses, false);
c_name = "if";
break;
+ case PRAGMA_OACC_CLAUSE_IF_PRESENT:
+ clauses = c_parser_oacc_simple_clause (parser, OMP_CLAUSE_IF_PRESENT,
+ clauses);
+ c_name = "if_present";
+ break;
case PRAGMA_OACC_CLAUSE_INDEPENDENT:
clauses = c_parser_oacc_simple_clause (parser, OMP_CLAUSE_INDEPENDENT,
clauses);
@@ -13997,22 +13993,6 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
c_name = "present";
break;
- case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY:
- clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
- c_name = "present_or_copy";
- break;
- case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN:
- clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
- c_name = "present_or_copyin";
- break;
- case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT:
- clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
- c_name = "present_or_copyout";
- break;
- case PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE:
- clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
- c_name = "present_or_create";
- break;
case PRAGMA_OACC_CLAUSE_PRIVATE:
clauses = c_parser_omp_clause_private (parser, clauses);
c_name = "private";
@@ -14021,10 +14001,6 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
clauses = c_parser_omp_clause_reduction (parser, clauses);
c_name = "reduction";
break;
- case PRAGMA_OACC_CLAUSE_SELF:
- clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
- c_name = "self";
- break;
case PRAGMA_OACC_CLAUSE_SEQ:
clauses = c_parser_oacc_simple_clause (parser, OMP_CLAUSE_SEQ,
clauses);
@@ -14417,11 +14393,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_PRESENT) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) )
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT))
static tree
c_parser_oacc_data (location_t loc, c_parser *parser, bool *if_p)
@@ -14451,11 +14423,7 @@ c_parser_oacc_data (location_t loc, c_parser *parser, bool *if_p)
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_LINK) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) )
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT))
static void
c_parser_oacc_declare (c_parser *parser)
@@ -14490,8 +14458,8 @@ c_parser_oacc_declare (c_parser *parser)
switch (OMP_CLAUSE_MAP_KIND (t))
{
case GOMP_MAP_FIRSTPRIVATE_POINTER:
- case GOMP_MAP_FORCE_ALLOC:
- case GOMP_MAP_FORCE_TO:
+ case GOMP_MAP_ALLOC:
+ case GOMP_MAP_TO:
case GOMP_MAP_FORCE_DEVICEPTR:
case GOMP_MAP_DEVICE_RESIDENT:
break;
@@ -14604,8 +14572,6 @@ c_parser_oacc_declare (c_parser *parser)
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
#define OACC_EXIT_DATA_CLAUSE_MASK \
@@ -14613,6 +14579,7 @@ c_parser_oacc_declare (c_parser *parser)
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DELETE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_FINALIZE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
static void
@@ -14756,10 +14723,6 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name,
| (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) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
@@ -14777,10 +14740,6 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name,
| (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) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
@@ -15008,7 +14967,7 @@ c_finish_oacc_routine (struct oacc_routine_data *data, tree fndecl,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_HOST) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SELF) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF_PRESENT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
static void
diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index aa70b23ff10..90ae306c99a 100644
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -13897,6 +13897,8 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
case OMP_CLAUSE_WORKER:
case OMP_CLAUSE_VECTOR:
case OMP_CLAUSE_TILE:
+ case OMP_CLAUSE_IF_PRESENT:
+ case OMP_CLAUSE_FINALIZE:
pc = &OMP_CLAUSE_CHAIN (c);
continue;
--
2.17.1
^ permalink raw reply [flat|nested] 31+ messages in thread* Re: [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior - C
2018-06-19 16:59 ` [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior - C Cesar Philippidis
@ 2018-06-20 16:14 ` Jakub Jelinek
0 siblings, 0 replies; 31+ messages in thread
From: Jakub Jelinek @ 2018-06-20 16:14 UTC (permalink / raw)
To: Cesar Philippidis; +Cc: gcc-patches
On Tue, Jun 19, 2018 at 09:59:09AM -0700, Cesar Philippidis wrote:
> This patch implements the OpenACC 2.5 data clause semantics in the C FE.
>
> Is it OK for trunk?
>
> Cesar
> 2018-06-19 Chung-Lin Tang <cltang@codesourcery.com>
> Thomas Schwinge <thomas@codesourcery.com>
> Cesar Philippidis <cesar@codesourcery.com>
>
> gcc/c/
> * c-parser.c (c_parser_omp_clause_name): Add support for finalize
> and if_present. Make present_or_{copy,copyin,copyout,create} aliases
> to their non-present_or_* counterparts. Make 'self' an alias to
> PRAGMA_OACC_CLAUSE_HOST.
> (c_parser_oacc_data_clause): Update GOMP mappings for
> PRAGMA_OACC_CLAUSE_{COPY,COPYIN,COPYOUT,CREATE,DELETE}. Remove
> PRAGMA_OACC_CLAUSE_{SELF,PRESENT_OR_*}.
> (c_parser_oacc_all_clauses): Handle finalize and if_present clauses.
> Remove support for present_or_* clauses.
> (OACC_KERNELS_CLAUSE_MASK): Remove PRESENT_OR_* clauses.
> (OACC_PARALLEL_CLAUSE_MASK): Likewise.
> (OACC_DECLARE_CLAUSE_MASK): Likewise.
> (OACC_DATA_CLAUSE_MASK): Likewise.
> (OACC_ENTER_DATA_CLAUSE_MASK): Remove PRESENT_OR_* clauses.
> (OACC_EXIT_DATA_CLAUSE_MASK): Add FINALIZE clause.
> (OACC_UPDATE_CLAUSE_MASK): Remove SELF, add IF_PRESENT.
> (c_parser_oacc_declare): Remove PRESENT_OR_* clauses.
> * c-typeck.c (c_finish_omp_clauses): Handle IF_PRESENT and FINALIZE.
Ok.
Jakub
^ permalink raw reply [flat|nested] 31+ messages in thread
* Re: [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior - Fortran
2018-06-19 16:56 ` Cesar Philippidis
2018-06-19 16:58 ` [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior - C++ Cesar Philippidis
2018-06-19 16:59 ` [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior - C Cesar Philippidis
@ 2018-06-19 17:00 ` Cesar Philippidis
2018-06-20 16:17 ` Jakub Jelinek
2018-06-19 17:01 ` [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior - runtime Cesar Philippidis
` (3 subsequent siblings)
6 siblings, 1 reply; 31+ messages in thread
From: Cesar Philippidis @ 2018-06-19 17:00 UTC (permalink / raw)
To: gcc-patches, Jakub Jelinek, Fortran List
[-- Attachment #1: Type: text/plain, Size: 107 bytes --]
This patch implements the OpenACC 2.5 data clause semantics in the
Fortran FE.
Is it OK for trunk?
Cesar
[-- Attachment #2: 0005-fortran-front-end.patch --]
[-- Type: text/x-patch, Size: 13427 bytes --]
2018-06-19 Chung-Lin Tang <cltang@codesourcery.com>
Thomas Schwinge <thomas@codesourcery.com>
Cesar Philippidis <cesar@codesourcery.com>
gcc/fortran/
* gfortran.h (gfc_omp_clauses): Add unsigned if_present, finalize
bitfields.
* openmp.c (enum omp_mask2): Remove OMP_CLAUSE_PRESENT_OR_*. Add
OMP_CLAUSE_{IF_PRESENT,FINALIZE}.
(gfc_match_omp_clauses): Update handling of copy, copyin, copyout,
create, deviceptr, present_of_*. Add support for finalize and
if_present.
(OACC_PARALLEL_CLAUSES): Remove PRESENT_OR_* clauses.
(OACC_KERNELS_CLAUSES): Likewise.
(OACC_DATA_CLAUSES): Likewise.
(OACC_DECLARE_CLAUSES): Likewise.
(OACC_UPDATE_CLAUSES): Add IF_PRESENT clause.
(OACC_ENTER_DATA_CLAUSES): Remove PRESENT_OR_* clauses.
(OACC_EXIT_DATA_CLAUSES): Add FINALIZE clause.
(gfc_match_oacc_declare): Update to OpenACC 2.5 semantics.
* trans-openmp.c (gfc_trans_omp_clauses): Add support for IF_PRESENT
and FINALIZE.
From 5a74cac327d44a6674a608c3fd3b80f2e04d0b47 Mon Sep 17 00:00:00 2001
From: Cesar Philippidis <cesar@codesourcery.com>
Date: Tue, 19 Jun 2018 09:31:57 -0700
Subject: [PATCH 5/7] fortran front end
---
gcc/fortran/gfortran.h | 1 +
gcc/fortran/openmp.c | 105 ++++++++++++++++++++-----------------
gcc/fortran/trans-openmp.c | 10 ++++
3 files changed, 67 insertions(+), 49 deletions(-)
diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index 1d98d2554c7..0b89f8de950 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -1344,6 +1344,7 @@ typedef struct gfc_omp_clauses
gfc_expr_list *tile_list;
unsigned async:1, gang:1, worker:1, vector:1, seq:1, independent:1;
unsigned wait:1, par_auto:1, gang_static:1;
+ unsigned if_present:1, finalize:1;
locus loc;
}
diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index 97d6e782373..94a7f7eaa50 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -796,10 +796,6 @@ enum omp_mask2
OMP_CLAUSE_COPYOUT,
OMP_CLAUSE_CREATE,
OMP_CLAUSE_PRESENT,
- OMP_CLAUSE_PRESENT_OR_COPY,
- OMP_CLAUSE_PRESENT_OR_COPYIN,
- OMP_CLAUSE_PRESENT_OR_COPYOUT,
- OMP_CLAUSE_PRESENT_OR_CREATE,
OMP_CLAUSE_DEVICEPTR,
OMP_CLAUSE_GANG,
OMP_CLAUSE_WORKER,
@@ -813,6 +809,8 @@ enum omp_mask2
OMP_CLAUSE_DELETE,
OMP_CLAUSE_AUTO,
OMP_CLAUSE_TILE,
+ OMP_CLAUSE_IF_PRESENT,
+ OMP_CLAUSE_FINALIZE,
/* This must come last. */
OMP_MASK2_LAST
};
@@ -1041,7 +1039,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
if ((mask & OMP_CLAUSE_COPY)
&& gfc_match ("copy ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_FORCE_TOFROM))
+ OMP_MAP_TOFROM))
continue;
if (mask & OMP_CLAUSE_COPYIN)
{
@@ -1049,7 +1047,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
{
if (gfc_match ("copyin ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_FORCE_TO))
+ OMP_MAP_TO))
continue;
}
else if (gfc_match_omp_variable_list ("copyin (",
@@ -1060,7 +1058,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
if ((mask & OMP_CLAUSE_COPYOUT)
&& gfc_match ("copyout ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_FORCE_FROM))
+ OMP_MAP_FROM))
continue;
if ((mask & OMP_CLAUSE_COPYPRIVATE)
&& gfc_match_omp_variable_list ("copyprivate (",
@@ -1070,7 +1068,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
if ((mask & OMP_CLAUSE_CREATE)
&& gfc_match ("create ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_FORCE_ALLOC))
+ OMP_MAP_ALLOC))
continue;
break;
case 'd':
@@ -1106,7 +1104,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
if ((mask & OMP_CLAUSE_DELETE)
&& gfc_match ("delete ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_DELETE))
+ OMP_MAP_RELEASE))
continue;
if ((mask & OMP_CLAUSE_DEPEND)
&& gfc_match ("depend ( ") == MATCH_YES)
@@ -1161,19 +1159,10 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
OMP_MAP_FORCE_TO))
continue;
if ((mask & OMP_CLAUSE_DEVICEPTR)
- && gfc_match ("deviceptr ( ") == MATCH_YES)
- {
- gfc_omp_namelist **list = &c->lists[OMP_LIST_MAP];
- gfc_omp_namelist **head = NULL;
- if (gfc_match_omp_variable_list ("", list, true, NULL,
- &head, false) == MATCH_YES)
- {
- gfc_omp_namelist *n;
- for (n = *head; n; n = n->next)
- n->u.map_op = OMP_MAP_FORCE_DEVICEPTR;
- continue;
- }
- }
+ && gfc_match ("deviceptr ( ") == MATCH_YES
+ && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
+ OMP_MAP_FORCE_DEVICEPTR))
+ continue;
if ((mask & OMP_CLAUSE_DEVICE_RESIDENT)
&& gfc_match_omp_variable_list
("device_resident (",
@@ -1202,6 +1191,14 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
&& c->final_expr == NULL
&& gfc_match ("final ( %e )", &c->final_expr) == MATCH_YES)
continue;
+ if ((mask & OMP_CLAUSE_FINALIZE)
+ && !c->finalize
+ && gfc_match ("finalize") == MATCH_YES)
+ {
+ c->finalize = true;
+ needs_space = true;
+ continue;
+ }
if ((mask & OMP_CLAUSE_FIRSTPRIVATE)
&& gfc_match_omp_variable_list ("firstprivate (",
&c->lists[OMP_LIST_FIRSTPRIVATE],
@@ -1274,6 +1271,14 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
}
gfc_current_locus = old_loc;
}
+ if ((mask & OMP_CLAUSE_IF_PRESENT)
+ && !c->if_present
+ && gfc_match ("if_present") == MATCH_YES)
+ {
+ c->if_present = true;
+ needs_space = true;
+ continue;
+ }
if ((mask & OMP_CLAUSE_INBRANCH)
&& !c->inbranch
&& !c->notinbranch
@@ -1503,22 +1508,22 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
}
break;
case 'p':
- if ((mask & OMP_CLAUSE_PRESENT_OR_COPY)
+ if ((mask & OMP_CLAUSE_COPY)
&& gfc_match ("pcopy ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
OMP_MAP_TOFROM))
continue;
- if ((mask & OMP_CLAUSE_PRESENT_OR_COPYIN)
+ if ((mask & OMP_CLAUSE_COPYIN)
&& gfc_match ("pcopyin ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
OMP_MAP_TO))
continue;
- if ((mask & OMP_CLAUSE_PRESENT_OR_COPYOUT)
+ if ((mask & OMP_CLAUSE_COPYOUT)
&& gfc_match ("pcopyout ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
OMP_MAP_FROM))
continue;
- if ((mask & OMP_CLAUSE_PRESENT_OR_CREATE)
+ if ((mask & OMP_CLAUSE_CREATE)
&& gfc_match ("pcreate ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
OMP_MAP_ALLOC))
@@ -1528,22 +1533,22 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
OMP_MAP_FORCE_PRESENT))
continue;
- if ((mask & OMP_CLAUSE_PRESENT_OR_COPY)
+ if ((mask & OMP_CLAUSE_COPY)
&& gfc_match ("present_or_copy ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
OMP_MAP_TOFROM))
continue;
- if ((mask & OMP_CLAUSE_PRESENT_OR_COPYIN)
+ if ((mask & OMP_CLAUSE_COPYIN)
&& gfc_match ("present_or_copyin ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
OMP_MAP_TO))
continue;
- if ((mask & OMP_CLAUSE_PRESENT_OR_COPYOUT)
+ if ((mask & OMP_CLAUSE_COPYOUT)
&& gfc_match ("present_or_copyout ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
OMP_MAP_FROM))
continue;
- if ((mask & OMP_CLAUSE_PRESENT_OR_CREATE)
+ if ((mask & OMP_CLAUSE_CREATE)
&& gfc_match ("present_or_create ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
OMP_MAP_ALLOC))
@@ -1925,23 +1930,19 @@ 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_PRESENT_OR_COPY \
- | OMP_CLAUSE_PRESENT_OR_COPYIN | OMP_CLAUSE_PRESENT_OR_COPYOUT \
- | OMP_CLAUSE_PRESENT_OR_CREATE | OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_PRIVATE \
- | OMP_CLAUSE_FIRSTPRIVATE | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_WAIT)
+ | OMP_CLAUSE_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_PRESENT_OR_COPY \
- | OMP_CLAUSE_PRESENT_OR_COPYIN | OMP_CLAUSE_PRESENT_OR_COPYOUT \
- | OMP_CLAUSE_PRESENT_OR_CREATE | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_WAIT)
+ | OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT | 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_PRESENT_OR_COPY \
- | OMP_CLAUSE_PRESENT_OR_COPYIN | OMP_CLAUSE_PRESENT_OR_COPYOUT \
- | OMP_CLAUSE_PRESENT_OR_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 \
@@ -1955,19 +1956,17 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
#define OACC_DECLARE_CLAUSES \
(omp_mask (OMP_CLAUSE_COPY) | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \
| OMP_CLAUSE_CREATE | OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_DEVICE_RESIDENT \
- | OMP_CLAUSE_PRESENT | OMP_CLAUSE_PRESENT_OR_COPY \
- | OMP_CLAUSE_PRESENT_OR_COPYIN | OMP_CLAUSE_PRESENT_OR_COPYOUT \
- | OMP_CLAUSE_PRESENT_OR_CREATE | OMP_CLAUSE_LINK)
+ | OMP_CLAUSE_PRESENT \
+ | OMP_CLAUSE_LINK)
#define OACC_UPDATE_CLAUSES \
(omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_HOST_SELF \
- | OMP_CLAUSE_DEVICE | OMP_CLAUSE_WAIT)
+ | OMP_CLAUSE_DEVICE | OMP_CLAUSE_WAIT | OMP_CLAUSE_IF_PRESENT)
#define OACC_ENTER_DATA_CLAUSES \
(omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_WAIT \
- | OMP_CLAUSE_COPYIN | OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT_OR_COPYIN \
- | OMP_CLAUSE_PRESENT_OR_CREATE)
+ | OMP_CLAUSE_COPYIN | OMP_CLAUSE_CREATE)
#define OACC_EXIT_DATA_CLAUSES \
(omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_WAIT \
- | OMP_CLAUSE_COPYOUT | OMP_CLAUSE_DELETE)
+ | OMP_CLAUSE_COPYOUT | OMP_CLAUSE_DELETE | OMP_CLAUSE_FINALIZE)
#define OACC_WAIT_CLAUSES \
omp_mask (OMP_CLAUSE_ASYNC)
#define OACC_ROUTINE_CLAUSES \
@@ -2061,8 +2060,7 @@ gfc_match_oacc_declare (void)
if (s->ns->proc_name && s->ns->proc_name->attr.proc == PROC_MODULE)
{
- if (n->u.map_op != OMP_MAP_FORCE_ALLOC
- && n->u.map_op != OMP_MAP_FORCE_TO)
+ if (n->u.map_op != OMP_MAP_ALLOC && n->u.map_op != OMP_MAP_TO)
{
gfc_error ("Invalid clause in module with !$ACC DECLARE at %L",
&where);
@@ -2072,6 +2070,13 @@ gfc_match_oacc_declare (void)
module_var = true;
}
+ if (ns->proc_name->attr.oacc_function)
+ {
+ gfc_error ("Invalid declare in routine with $!ACC DECLARE at %L",
+ &where);
+ return MATCH_ERROR;
+ }
+
if (s->attr.use_assoc)
{
gfc_error ("Variable is USE-associated with !$ACC DECLARE at %L",
@@ -2090,10 +2095,12 @@ gfc_match_oacc_declare (void)
switch (n->u.map_op)
{
case OMP_MAP_FORCE_ALLOC:
+ case OMP_MAP_ALLOC:
s->attr.oacc_declare_create = 1;
break;
case OMP_MAP_FORCE_TO:
+ case OMP_MAP_TO:
s->attr.oacc_declare_copyin = 1;
break;
diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index 795175d701a..f038f4c5bf8 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -2895,6 +2895,16 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
c = build_omp_clause (where.lb->location, OMP_CLAUSE_AUTO);
omp_clauses = gfc_trans_add_clause (c, omp_clauses);
}
+ if (clauses->if_present)
+ {
+ c = build_omp_clause (where.lb->location, OMP_CLAUSE_IF_PRESENT);
+ omp_clauses = gfc_trans_add_clause (c, omp_clauses);
+ }
+ if (clauses->finalize)
+ {
+ c = build_omp_clause (where.lb->location, OMP_CLAUSE_FINALIZE);
+ omp_clauses = gfc_trans_add_clause (c, omp_clauses);
+ }
if (clauses->independent)
{
c = build_omp_clause (where.lb->location, OMP_CLAUSE_INDEPENDENT);
--
2.17.1
^ permalink raw reply [flat|nested] 31+ messages in thread* Re: [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior - Fortran
2018-06-19 17:00 ` [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior - Fortran Cesar Philippidis
@ 2018-06-20 16:17 ` Jakub Jelinek
0 siblings, 0 replies; 31+ messages in thread
From: Jakub Jelinek @ 2018-06-20 16:17 UTC (permalink / raw)
To: Cesar Philippidis; +Cc: gcc-patches, Fortran List
On Tue, Jun 19, 2018 at 09:59:57AM -0700, Cesar Philippidis wrote:
> This patch implements the OpenACC 2.5 data clause semantics in the
> Fortran FE.
>
> Is it OK for trunk?
>
> Cesar
> 2018-06-19 Chung-Lin Tang <cltang@codesourcery.com>
> Thomas Schwinge <thomas@codesourcery.com>
> Cesar Philippidis <cesar@codesourcery.com>
>
> gcc/fortran/
> * gfortran.h (gfc_omp_clauses): Add unsigned if_present, finalize
> bitfields.
> * openmp.c (enum omp_mask2): Remove OMP_CLAUSE_PRESENT_OR_*. Add
> OMP_CLAUSE_{IF_PRESENT,FINALIZE}.
> (gfc_match_omp_clauses): Update handling of copy, copyin, copyout,
> create, deviceptr, present_of_*. Add support for finalize and
> if_present.
> (OACC_PARALLEL_CLAUSES): Remove PRESENT_OR_* clauses.
> (OACC_KERNELS_CLAUSES): Likewise.
> (OACC_DATA_CLAUSES): Likewise.
> (OACC_DECLARE_CLAUSES): Likewise.
> (OACC_UPDATE_CLAUSES): Add IF_PRESENT clause.
> (OACC_ENTER_DATA_CLAUSES): Remove PRESENT_OR_* clauses.
> (OACC_EXIT_DATA_CLAUSES): Add FINALIZE clause.
> (gfc_match_oacc_declare): Update to OpenACC 2.5 semantics.
> * trans-openmp.c (gfc_trans_omp_clauses): Add support for IF_PRESENT
> and FINALIZE.
Ok.
Jakub
^ permalink raw reply [flat|nested] 31+ messages in thread
* Re: [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior - runtime
2018-06-19 16:56 ` Cesar Philippidis
` (2 preceding siblings ...)
2018-06-19 17:00 ` [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior - Fortran Cesar Philippidis
@ 2018-06-19 17:01 ` Cesar Philippidis
2018-06-20 16:45 ` Jakub Jelinek
` (4 more replies)
2018-06-19 17:01 ` [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior - middle end Cesar Philippidis
` (2 subsequent siblings)
6 siblings, 5 replies; 31+ messages in thread
From: Cesar Philippidis @ 2018-06-19 17:01 UTC (permalink / raw)
To: gcc-patches, Jakub Jelinek
[-- Attachment #1: Type: text/plain, Size: 100 bytes --]
This patch implements the OpenACC 2.5 data clause semantics in libgomp.
Is it OK for trunk?
Cesar
[-- Attachment #2: 0007-runtime-changes.patch --]
[-- Type: text/x-patch, Size: 33752 bytes --]
2018-06-19 Chung-Lin Tang <cltang@codesourcery.com>
Thomas Schwinge <thomas@codesourcery.com>
Cesar Philippidis <cesar@codesourcery.com>
libgomp/
* libgomp.h (struct splay_tree_key_s): Add dynamic_refcount member.
(gomp_acc_remove_pointer): Update declaration.
(gomp_acc_declare_allocate): Declare.
(gomp_remove_var): Declare.
* libgomp.map (OACC_2.5): Define.
* oacc-mem.c (acc_map_data): Update refcount.
(acc_unmap_data): Likewise.
(present_create_copy): Likewise.
(acc_create): Add FLAG_PRESENT when calling present_create_copy.
(acc_copyin): Likewise.
(FLAG_FINALIZE): Define.
(delete_copyout): Update dynamic refcounts, add support for FINALIZE.
(acc_delete_finalize): New function.
(acc_delete_finalize_async): New function.
(acc_copyout_finalize): New function.
(acc_copyout_finalize_async): New function.
(gomp_acc_insert_pointer): Update refcounts.
(gomp_acc_remove_pointer): Return if data is not present on the
accelerator.
* oacc-parallel.c (find_pset): Rename to find_pointer.
(find_pointer): Add support for GOMP_MAP_POINTER.
(handle_ftn_pointers): New function.
(GOACC_parallel_keyed): Update refcounts of variables.
(GOACC_enter_exit_data): Add support for finalized data mappings.
Add support for GOMP_MAP_{TO,ALLOC,RELESE,FROM}. Update handling
of fortran arrays.
(GOACC_update): Add support for GOMP_MAP_{ALWAYS_POINTER,TO,FROM}.
(GOACC_declare): Add support for GOMP_MAP_RELEASE, remove support
for GOMP_MAP_FORCE_FROM.
* openacc.f90 (module openacc_internal): Add
acc_copyout_finalize_{32_h,64_h,array_h,_l}, and
acc_delete_finalize_{32_h,64_h,array_h,_l}. Add interfaces for
acc_copyout_finalize and acc_delete_finalize.
(acc_copyout_finalize_32_h): New subroutine.
(acc_copyout_finalize_64_h): New subroutine.
(acc_copyout_finalize_array_h): New subroutine.
(acc_delete_finalize_32_h): New subroutine.
(acc_delete_finalize_64_h): New subroutine.
(acc_delete_finalize_array_h): New subroutine.
* openacc.h (acc_copyout_finalize): Declare.
(acc_copyout_finalize_async): Declare.
(acc_delete_finalize): Declare.
(acc_delete_finalize_async): Declare.
* openacc_lib.h (acc_copyout_finalize): New interface.
(acc_delete_finalize): New interface.
* target.c (gomp_map_vars): Update dynamic_refcount.
(gomp_remove_var): New function.
(gomp_unmap_vars): Use it.
(gomp_unload_image_from_device): Likewise.
From 53ee03231c5e6e4747b4ef01335079a2d4a98480 Mon Sep 17 00:00:00 2001
From: Cesar Philippidis <cesar@codesourcery.com>
Date: Tue, 19 Jun 2018 09:33:04 -0700
Subject: [PATCH 7/7] runtime changes
---
libgomp/libgomp.h | 7 +-
libgomp/libgomp.map | 12 +++
libgomp/oacc-mem.c | 196 ++++++++++++++++++++++++++++++++-------
libgomp/oacc-parallel.c | 198 ++++++++++++++++++++++++++++++++++------
libgomp/openacc.f90 | 112 +++++++++++++++++++++++
libgomp/openacc.h | 6 ++
libgomp/openacc_lib.h | 40 ++++++++
libgomp/target.c | 41 ++++-----
8 files changed, 528 insertions(+), 84 deletions(-)
diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h
index 10ea8940c96..3a8cc2bd7d6 100644
--- a/libgomp/libgomp.h
+++ b/libgomp/libgomp.h
@@ -853,6 +853,8 @@ struct splay_tree_key_s {
uintptr_t tgt_offset;
/* Reference count. */
uintptr_t refcount;
+ /* Dynamic reference count. */
+ uintptr_t dynamic_refcount;
/* Pointer to the original mapping of "omp declare target link" object. */
splay_tree_key link_key;
};
@@ -991,7 +993,9 @@ enum gomp_map_vars_kind
};
extern void gomp_acc_insert_pointer (size_t, void **, size_t *, void *);
-extern void gomp_acc_remove_pointer (void *, bool, int, int);
+extern void gomp_acc_remove_pointer (void *, size_t, bool, int, int, int);
+extern void gomp_acc_declare_allocate (bool, size_t, void **, size_t *,
+ unsigned short *);
extern struct target_mem_desc *gomp_map_vars (struct gomp_device_descr *,
size_t, void **, void **,
@@ -1001,6 +1005,7 @@ extern void gomp_unmap_vars (struct target_mem_desc *, bool);
extern void gomp_init_device (struct gomp_device_descr *);
extern void gomp_free_memmap (struct splay_tree_s *);
extern void gomp_unload_device (struct gomp_device_descr *);
+extern bool gomp_remove_var (struct gomp_device_descr *, splay_tree_key);
/* work.c */
diff --git a/libgomp/libgomp.map b/libgomp/libgomp.map
index 8752348fbf2..2cd3bf524bc 100644
--- a/libgomp/libgomp.map
+++ b/libgomp/libgomp.map
@@ -386,6 +386,18 @@ OACC_2.0.1 {
acc_pcreate;
} OACC_2.0;
+OACC_2.5 {
+ global:
+ acc_copyout_finalize;
+ acc_copyout_finalize_32_h_;
+ acc_copyout_finalize_64_h_;
+ acc_copyout_finalize_array_h_;
+ acc_delete_finalize;
+ acc_delete_finalize_32_h_;
+ acc_delete_finalize_64_h_;
+ acc_delete_finalize_array_h_;
+} OACC_2.0.1;
+
GOACC_2.0 {
global:
GOACC_data_end;
diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index 158f0862018..3787ce49e38 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -347,6 +347,7 @@ acc_map_data (void *h, void *d, size_t s)
tgt = gomp_map_vars (acc_dev, mapnum, &hostaddrs, &devaddrs, &sizes,
&kinds, true, GOMP_MAP_VARS_OPENACC);
+ tgt->list[0].key->refcount = REFCOUNT_INFINITY;
}
gomp_mutex_lock (&acc_dev->lock);
@@ -389,6 +390,9 @@ acc_unmap_data (void *h)
(void *) n->host_start, (int) host_size, (void *) h);
}
+ /* Mark for removal. */
+ n->refcount = 1;
+
t = n->tgt;
if (t->refcount == 2)
@@ -460,6 +464,11 @@ present_create_copy (unsigned f, void *h, size_t s)
gomp_fatal ("[%p,+%d] not mapped", (void *)h, (int)s);
}
+ if (n->refcount != REFCOUNT_INFINITY)
+ {
+ n->refcount++;
+ n->dynamic_refcount++;
+ }
gomp_mutex_unlock (&acc_dev->lock);
}
else if (!(f & FLAG_CREATE))
@@ -483,6 +492,8 @@ present_create_copy (unsigned f, void *h, size_t s)
tgt = gomp_map_vars (acc_dev, mapnum, &hostaddrs, NULL, &s, &kinds, true,
GOMP_MAP_VARS_OPENACC);
+ /* Initialize dynamic refcount. */
+ tgt->list[0].key->dynamic_refcount = 1;
gomp_mutex_lock (&acc_dev->lock);
@@ -499,13 +510,13 @@ present_create_copy (unsigned f, void *h, size_t s)
void *
acc_create (void *h, size_t s)
{
- return present_create_copy (FLAG_CREATE, h, s);
+ return present_create_copy (FLAG_PRESENT | FLAG_CREATE, h, s);
}
void *
acc_copyin (void *h, size_t s)
{
- return present_create_copy (FLAG_CREATE | FLAG_COPY, h, s);
+ return present_create_copy (FLAG_PRESENT | FLAG_CREATE | FLAG_COPY, h, s);
}
void *
@@ -542,7 +553,8 @@ acc_pcopyin (void *h, size_t s)
}
#endif
-#define FLAG_COPYOUT (1 << 0)
+#define FLAG_COPYOUT (1 << 0)
+#define FLAG_FINALIZE (1 << 1)
static void
delete_copyout (unsigned f, void *h, size_t s, const char *libfnname)
@@ -581,15 +593,52 @@ delete_copyout (unsigned f, void *h, size_t s, const char *libfnname)
(void *) n->host_start, (int) host_size, (void *) h, (int) s);
}
- gomp_mutex_unlock (&acc_dev->lock);
+ if (n->refcount == REFCOUNT_INFINITY)
+ {
+ n->refcount = 0;
+ n->dynamic_refcount = 0;
+ }
+ if (n->refcount < n->dynamic_refcount)
+ {
+ gomp_mutex_unlock (&acc_dev->lock);
+ gomp_fatal ("Dynamic reference counting assert fail\n");
+ }
- if (f & FLAG_COPYOUT)
- acc_dev->dev2host_func (acc_dev->target_id, h, d, s);
+ if (f & FLAG_FINALIZE)
+ {
+ n->refcount -= n->dynamic_refcount;
+ n->dynamic_refcount = 0;
+ }
+ else if (n->dynamic_refcount)
+ {
+ n->dynamic_refcount--;
+ n->refcount--;
+ }
- acc_unmap_data (h);
+ if (n->refcount == 0)
+ {
+ if (n->tgt->refcount == 2)
+ {
+ struct target_mem_desc *tp, *t;
+ for (tp = NULL, t = acc_dev->openacc.data_environ; t != NULL;
+ tp = t, t = t->prev)
+ if (n->tgt == t)
+ {
+ if (tp)
+ tp->prev = t->prev;
+ else
+ acc_dev->openacc.data_environ = t->prev;
+ break;
+ }
+ }
- if (!acc_dev->free_func (acc_dev->target_id, d))
- gomp_fatal ("error in freeing device memory in %s", libfnname);
+ if (f & FLAG_COPYOUT)
+ acc_dev->dev2host_func (acc_dev->target_id, h, d, s);
+
+ gomp_remove_var (acc_dev, n);
+ }
+
+ gomp_mutex_unlock (&acc_dev->lock);
}
void
@@ -598,12 +647,36 @@ acc_delete (void *h , size_t s)
delete_copyout (0, h, s, __FUNCTION__);
}
+void
+acc_delete_finalize (void *h , size_t s)
+{
+ delete_copyout (FLAG_FINALIZE, h, s, __FUNCTION__);
+}
+
+void
+acc_delete_finalize_async (void *h , size_t s, int async)
+{
+ delete_copyout (FLAG_FINALIZE, h, s, __FUNCTION__);
+}
+
void
acc_copyout (void *h, size_t s)
{
delete_copyout (FLAG_COPYOUT, h, s, __FUNCTION__);
}
+void
+acc_copyout_finalize (void *h, size_t s)
+{
+ delete_copyout (FLAG_COPYOUT | FLAG_FINALIZE, h, s, __FUNCTION__);
+}
+
+void
+acc_copyout_finalize_async (void *h, size_t s, int async)
+{
+ delete_copyout (FLAG_COPYOUT | FLAG_FINALIZE, h, s, __FUNCTION__);
+}
+
static void
update_dev_host (int is_dev, void *h, size_t s)
{
@@ -659,11 +732,37 @@ gomp_acc_insert_pointer (size_t mapnum, void **hostaddrs, size_t *sizes,
struct goacc_thread *thr = goacc_thread ();
struct gomp_device_descr *acc_dev = thr->dev;
+ if (acc_is_present (*hostaddrs, *sizes))
+ {
+ splay_tree_key n;
+ gomp_mutex_lock (&acc_dev->lock);
+ n = lookup_host (acc_dev, *hostaddrs, *sizes);
+ gomp_mutex_unlock (&acc_dev->lock);
+
+ tgt = n->tgt;
+ for (size_t i = 0; i < tgt->list_count; i++)
+ if (tgt->list[i].key == n)
+ {
+ for (size_t j = 0; j < mapnum; j++)
+ if (i + j < tgt->list_count && tgt->list[i + j].key)
+ {
+ tgt->list[i + j].key->refcount++;
+ tgt->list[i + j].key->dynamic_refcount++;
+ }
+ return;
+ }
+ /* Should not reach here. */
+ gomp_fatal ("Dynamic refcount incrementing failed for pointer/pset");
+ }
+
gomp_debug (0, " %s: prepare mappings\n", __FUNCTION__);
tgt = gomp_map_vars (acc_dev, mapnum, hostaddrs,
NULL, sizes, kinds, true, GOMP_MAP_VARS_OPENACC);
gomp_debug (0, " %s: mappings prepared\n", __FUNCTION__);
+ /* Initialize dynamic refcount. */
+ tgt->list[0].key->dynamic_refcount = 1;
+
gomp_mutex_lock (&acc_dev->lock);
tgt->prev = acc_dev->openacc.data_environ;
acc_dev->openacc.data_environ = tgt;
@@ -671,7 +770,8 @@ gomp_acc_insert_pointer (size_t mapnum, void **hostaddrs, size_t *sizes,
}
void
-gomp_acc_remove_pointer (void *h, bool force_copyfrom, int async, int mapnum)
+gomp_acc_remove_pointer (void *h, size_t s, bool force_copyfrom, int async,
+ int finalize, int mapnum)
{
struct goacc_thread *thr = goacc_thread ();
struct gomp_device_descr *acc_dev = thr->dev;
@@ -679,6 +779,9 @@ gomp_acc_remove_pointer (void *h, bool force_copyfrom, int async, int mapnum)
struct target_mem_desc *t;
int minrefs = (mapnum == 1) ? 2 : 3;
+ if (!acc_is_present (h, s))
+ return;
+
gomp_mutex_lock (&acc_dev->lock);
n = lookup_host (acc_dev, h, 1);
@@ -693,40 +796,65 @@ gomp_acc_remove_pointer (void *h, bool force_copyfrom, int async, int mapnum)
t = n->tgt;
- struct target_mem_desc *tp;
+ if (n->refcount < n->dynamic_refcount)
+ {
+ gomp_mutex_unlock (&acc_dev->lock);
+ gomp_fatal ("Dynamic reference counting assert fail\n");
+ }
- if (t->refcount == minrefs)
+ if (finalize)
{
- /* This is the last reference, so pull the descriptor off the
- chain. This avoids gomp_unmap_vars via gomp_unmap_tgt from
- freeing the device memory. */
- t->tgt_end = 0;
- t->to_free = 0;
+ n->refcount -= n->dynamic_refcount;
+ n->dynamic_refcount = 0;
+ }
+ else if (n->dynamic_refcount)
+ {
+ n->dynamic_refcount--;
+ n->refcount--;
+ }
- for (tp = NULL, t = acc_dev->openacc.data_environ; t != NULL;
- tp = t, t = t->prev)
+ gomp_mutex_unlock (&acc_dev->lock);
+
+ if (n->refcount == 0)
+ {
+ if (t->refcount == minrefs)
{
- if (n->tgt == t)
+ /* This is the last reference, so pull the descriptor off the
+ chain. This prevents gomp_unmap_vars via gomp_unmap_tgt from
+ freeing the device memory. */
+ struct target_mem_desc *tp;
+ for (tp = NULL, t = acc_dev->openacc.data_environ; t != NULL;
+ tp = t, t = t->prev)
{
- if (tp)
- tp->prev = t->prev;
- else
- acc_dev->openacc.data_environ = t->prev;
- break;
+ if (n->tgt == t)
+ {
+ if (tp)
+ tp->prev = t->prev;
+ else
+ acc_dev->openacc.data_environ = t->prev;
+ break;
+ }
}
}
- }
- if (force_copyfrom)
- t->list[0].copy_from = 1;
+ /* Set refcount to 1 to allow gomp_unmap_vars to unmap it. */
+ n->refcount = 1;
+ t->refcount = minrefs;
+ for (size_t i = 0; i < t->list_count; i++)
+ if (t->list[i].key == n)
+ {
+ t->list[i].copy_from = force_copyfrom ? 1 : 0;
+ break;
+ }
- gomp_mutex_unlock (&acc_dev->lock);
+ /* If running synchronously, unmap immediately. */
+ if (async < acc_async_noval)
+ gomp_unmap_vars (t, true);
+ else
+ t->device_descr->openacc.register_async_cleanup_func (t, async);
+ }
- /* If running synchronously, unmap immediately. */
- if (async_synchronous_p (async))
- gomp_unmap_vars (t, true);
- else
- t->device_descr->openacc.register_async_cleanup_func (t, async);
+ gomp_mutex_unlock (&acc_dev->lock);
gomp_debug (0, " %s: mappings restored\n", __FUNCTION__);
}
diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c
index 9eae43131f8..b80ace58590 100644
--- a/libgomp/oacc-parallel.c
+++ b/libgomp/oacc-parallel.c
@@ -38,15 +38,68 @@
#include <stdarg.h>
#include <assert.h>
+/* Returns the number of mappings associated with the pointer or pset. PSET
+ have three mappings, whereas pointer have two. */
+
static int
-find_pset (int pos, size_t mapnum, unsigned short *kinds)
+find_pointer (int pos, size_t mapnum, unsigned short *kinds)
{
if (pos + 1 >= mapnum)
return 0;
unsigned char kind = kinds[pos+1] & 0xff;
- return kind == GOMP_MAP_TO_PSET;
+ if (kind == GOMP_MAP_TO_PSET)
+ return 3;
+ else if (kind == GOMP_MAP_POINTER)
+ return 2;
+
+ return 0;
+}
+
+/* Handle the mapping pair that are presented when a
+ deviceptr clause is used with Fortran. */
+
+static void
+handle_ftn_pointers (size_t mapnum, void **hostaddrs, size_t *sizes,
+ unsigned short *kinds)
+{
+ int i;
+
+ for (i = 0; i < mapnum; i++)
+ {
+ unsigned short kind1 = kinds[i] & 0xff;
+
+ /* Handle Fortran deviceptr clause. */
+ if (kind1 == GOMP_MAP_FORCE_DEVICEPTR)
+ {
+ unsigned short kind2;
+
+ if (i < (signed)mapnum - 1)
+ kind2 = kinds[i + 1] & 0xff;
+ else
+ kind2 = 0xffff;
+
+ if (sizes[i] == sizeof (void *))
+ continue;
+
+ /* At this point, we're dealing with a Fortran deviceptr.
+ If the next element is not what we're expecting, then
+ this is an instance of where the deviceptr variable was
+ not used within the region and the pointer was removed
+ by the gimplifier. */
+ if (kind2 == GOMP_MAP_POINTER
+ && sizes[i + 1] == 0
+ && hostaddrs[i] == *(void **)hostaddrs[i + 1])
+ {
+ kinds[i+1] = kinds[i];
+ sizes[i+1] = sizeof (void *);
+ }
+
+ /* Invalidate the entry. */
+ hostaddrs[i] = NULL;
+ }
+ }
}
static void goacc_wait (int async, int num_waits, va_list *ap);
@@ -88,6 +141,8 @@ GOACC_parallel_keyed (int device, void (*fn) (void *),
thr = goacc_thread ();
acc_dev = thr->dev;
+ handle_ftn_pointers (mapnum, hostaddrs, sizes, kinds);
+
/* Host fallback if "if" clause is false or if the current device is set to
the host. */
if (host_fallback)
@@ -183,10 +238,29 @@ GOACC_parallel_keyed (int device, void (*fn) (void *),
async, dims, tgt);
/* If running synchronously, unmap immediately. */
+ bool copyfrom = true;
if (async_synchronous_p (async))
gomp_unmap_vars (tgt, true);
else
- tgt->device_descr->openacc.register_async_cleanup_func (tgt, async);
+ {
+ bool async_unmap = false;
+ for (size_t i = 0; i < tgt->list_count; i++)
+ {
+ splay_tree_key k = tgt->list[i].key;
+ if (k && k->refcount == 1)
+ {
+ async_unmap = true;
+ break;
+ }
+ }
+ if (async_unmap)
+ tgt->device_descr->openacc.register_async_cleanup_func (tgt, async);
+ else
+ {
+ copyfrom = false;
+ gomp_unmap_vars (tgt, copyfrom);
+ }
+ }
acc_dev->openacc.async_set_async_func (acc_async_sync);
}
@@ -286,6 +360,17 @@ GOACC_enter_exit_data (int device, size_t mapnum,
va_end (ap);
}
+ /* Determine whether "finalize" semantics apply to all mappings of this
+ OpenACC directive. */
+ bool finalize = false;
+ if (mapnum > 0)
+ {
+ unsigned char kind = kinds[0] & 0xff;
+ if (kind == GOMP_MAP_DELETE
+ || kind == GOMP_MAP_FORCE_FROM)
+ finalize = true;
+ }
+
acc_dev->openacc.async_set_async_func (async);
/* Determine if this is an "acc enter data". */
@@ -298,13 +383,17 @@ GOACC_enter_exit_data (int device, size_t mapnum,
if (kind == GOMP_MAP_FORCE_ALLOC
|| kind == GOMP_MAP_FORCE_PRESENT
- || kind == GOMP_MAP_FORCE_TO)
+ || kind == GOMP_MAP_FORCE_TO
+ || kind == GOMP_MAP_TO
+ || kind == GOMP_MAP_ALLOC)
{
data_enter = true;
break;
}
- if (kind == GOMP_MAP_DELETE
+ if (kind == GOMP_MAP_RELEASE
+ || kind == GOMP_MAP_DELETE
+ || kind == GOMP_MAP_FROM
|| kind == GOMP_MAP_FORCE_FROM)
break;
@@ -312,31 +401,39 @@ GOACC_enter_exit_data (int device, size_t mapnum,
kind);
}
+ /* In c, non-pointers and arrays are represented by a single data clause.
+ Dynamically allocated arrays and subarrays are represented by a data
+ clause followed by an internal GOMP_MAP_POINTER.
+
+ In fortran, scalars and not allocated arrays are represented by a
+ single data clause. Allocated arrays and subarrays have three mappings:
+ 1) the original data clause, 2) a PSET 3) a pointer to the array data.
+ */
+
if (data_enter)
{
for (i = 0; i < mapnum; i++)
{
unsigned char kind = kinds[i] & 0xff;
- /* Scan for PSETs. */
- int psets = find_pset (i, mapnum, kinds);
+ /* Scan for pointers and PSETs. */
+ int pointer = find_pointer (i, mapnum, kinds);
- if (!psets)
+ if (!pointer)
{
switch (kind)
{
- case GOMP_MAP_POINTER:
- gomp_acc_insert_pointer (1, &hostaddrs[i], &sizes[i],
- &kinds[i]);
+ case GOMP_MAP_ALLOC:
+ acc_present_or_create (hostaddrs[i], sizes[i]);
break;
case GOMP_MAP_FORCE_ALLOC:
acc_create (hostaddrs[i], sizes[i]);
break;
- case GOMP_MAP_FORCE_PRESENT:
+ case GOMP_MAP_TO:
acc_present_or_copyin (hostaddrs[i], sizes[i]);
break;
case GOMP_MAP_FORCE_TO:
- acc_present_or_copyin (hostaddrs[i], sizes[i]);
+ acc_copyin (hostaddrs[i], sizes[i]);
break;
default:
gomp_fatal (">>>> GOACC_enter_exit_data UNHANDLED kind 0x%.2x",
@@ -346,12 +443,13 @@ GOACC_enter_exit_data (int device, size_t mapnum,
}
else
{
- gomp_acc_insert_pointer (3, &hostaddrs[i], &sizes[i], &kinds[i]);
+ gomp_acc_insert_pointer (pointer, &hostaddrs[i],
+ &sizes[i], &kinds[i]);
/* Increment 'i' by two because OpenACC requires fortran
arrays to be contiguous, so each PSET is associated with
one of MAP_FORCE_ALLOC/MAP_FORCE_PRESET/MAP_FORCE_TO, and
one MAP_POINTER. */
- i += 2;
+ i += pointer - 1;
}
}
}
@@ -360,22 +458,28 @@ GOACC_enter_exit_data (int device, size_t mapnum,
{
unsigned char kind = kinds[i] & 0xff;
- int psets = find_pset (i, mapnum, kinds);
+ int pointer = find_pointer (i, mapnum, kinds);
- if (!psets)
+ if (!pointer)
{
switch (kind)
{
- case GOMP_MAP_POINTER:
- gomp_acc_remove_pointer (hostaddrs[i], (kinds[i] & 0xff)
- == GOMP_MAP_FORCE_FROM,
- async, 1);
- break;
+ case GOMP_MAP_RELEASE:
case GOMP_MAP_DELETE:
- acc_delete (hostaddrs[i], sizes[i]);
+ if (acc_is_present (hostaddrs[i], sizes[i]))
+ {
+ if (finalize)
+ acc_delete_finalize (hostaddrs[i], sizes[i]);
+ else
+ acc_delete (hostaddrs[i], sizes[i]);
+ }
break;
+ case GOMP_MAP_FROM:
case GOMP_MAP_FORCE_FROM:
- acc_copyout (hostaddrs[i], sizes[i]);
+ if (finalize)
+ acc_copyout_finalize (hostaddrs[i], sizes[i]);
+ else
+ acc_copyout (hostaddrs[i], sizes[i]);
break;
default:
gomp_fatal (">>>> GOACC_enter_exit_data UNHANDLED kind 0x%.2x",
@@ -385,10 +489,12 @@ GOACC_enter_exit_data (int device, size_t mapnum,
}
else
{
- gomp_acc_remove_pointer (hostaddrs[i], (kinds[i] & 0xff)
- == GOMP_MAP_FORCE_FROM, async, 3);
+ bool copyfrom = (kind == GOMP_MAP_FORCE_FROM
+ || kind == GOMP_MAP_FROM);
+ gomp_acc_remove_pointer (hostaddrs[i], sizes[i], copyfrom, async,
+ finalize, pointer);
/* See the above comment. */
- i += 2;
+ i += pointer - 1;
}
}
@@ -447,6 +553,7 @@ GOACC_update (int device, size_t mapnum,
acc_dev->openacc.async_set_async_func (async);
+ bool update_device = false;
for (i = 0; i < mapnum; ++i)
{
unsigned char kind = kinds[i] & 0xff;
@@ -457,11 +564,46 @@ GOACC_update (int device, size_t mapnum,
case GOMP_MAP_TO_PSET:
break;
+ case GOMP_MAP_ALWAYS_POINTER:
+ if (update_device)
+ {
+ /* Save the contents of the host pointer. */
+ void *dptr = acc_deviceptr (hostaddrs[i-1]);
+ uintptr_t t = *(uintptr_t *) hostaddrs[i];
+
+ /* Update the contents of the host pointer to reflect
+ the value of the allocated device memory in the
+ previous pointer. */
+ *(uintptr_t *) hostaddrs[i] = (uintptr_t)dptr;
+ acc_update_device (hostaddrs[i], sizeof (uintptr_t));
+
+ /* Restore the host pointer. */
+ *(uintptr_t *) hostaddrs[i] = t;
+ update_device = false;
+ }
+ break;
+
+ case GOMP_MAP_TO:
+ if (!acc_is_present (hostaddrs[i], sizes[i]))
+ {
+ update_device = false;
+ break;
+ }
+ /* Fallthru */
case GOMP_MAP_FORCE_TO:
+ update_device = true;
acc_update_device (hostaddrs[i], sizes[i]);
break;
+ case GOMP_MAP_FROM:
+ if (!acc_is_present (hostaddrs[i], sizes[i]))
+ {
+ update_device = false;
+ break;
+ }
+ /* Fallthru */
case GOMP_MAP_FORCE_FROM:
+ update_device = false;
acc_update_self (hostaddrs[i], sizes[i]);
break;
@@ -522,6 +664,7 @@ GOACC_declare (int device, size_t mapnum,
case GOMP_MAP_FORCE_FROM:
case GOMP_MAP_FORCE_TO:
case GOMP_MAP_POINTER:
+ case GOMP_MAP_RELEASE:
case GOMP_MAP_DELETE:
GOACC_enter_exit_data (device, 1, &hostaddrs[i], &sizes[i],
&kinds[i], GOMP_ASYNC_SYNC, 0);
@@ -543,7 +686,6 @@ GOACC_declare (int device, size_t mapnum,
break;
case GOMP_MAP_FROM:
- kinds[i] = GOMP_MAP_FORCE_FROM;
GOACC_enter_exit_data (device, 1, &hostaddrs[i], &sizes[i],
&kinds[i], GOMP_ASYNC_SYNC, 0);
break;
diff --git a/libgomp/openacc.f90 b/libgomp/openacc.f90
index d201d1dde6f..84a8700f072 100644
--- a/libgomp/openacc.f90
+++ b/libgomp/openacc.f90
@@ -222,6 +222,24 @@ module openacc_internal
type (*), dimension (..), contiguous :: a
end subroutine
+ subroutine acc_copyout_finalize_32_h (a, len)
+ use iso_c_binding, only: c_int32_t
+ !GCC$ ATTRIBUTES NO_ARG_CHECK :: a
+ type (*), dimension (*) :: a
+ integer (c_int32_t) len
+ end subroutine
+
+ subroutine acc_copyout_finalize_64_h (a, len)
+ use iso_c_binding, only: c_int64_t
+ !GCC$ ATTRIBUTES NO_ARG_CHECK :: a
+ type (*), dimension (*) :: a
+ integer (c_int64_t) len
+ end subroutine
+
+ subroutine acc_copyout_finalize_array_h (a)
+ type (*), dimension (..), contiguous :: a
+ end subroutine
+
subroutine acc_delete_32_h (a, len)
use iso_c_binding, only: c_int32_t
!GCC$ ATTRIBUTES NO_ARG_CHECK :: a
@@ -240,6 +258,24 @@ module openacc_internal
type (*), dimension (..), contiguous :: a
end subroutine
+ subroutine acc_delete_finalize_32_h (a, len)
+ use iso_c_binding, only: c_int32_t
+ !GCC$ ATTRIBUTES NO_ARG_CHECK :: a
+ type (*), dimension (*) :: a
+ integer (c_int32_t) len
+ end subroutine
+
+ subroutine acc_delete_finalize_64_h (a, len)
+ use iso_c_binding, only: c_int64_t
+ !GCC$ ATTRIBUTES NO_ARG_CHECK :: a
+ type (*), dimension (*) :: a
+ integer (c_int64_t) len
+ end subroutine
+
+ subroutine acc_delete_finalize_array_h (a)
+ type (*), dimension (..), contiguous :: a
+ end subroutine
+
subroutine acc_update_device_32_h (a, len)
use iso_c_binding, only: c_int32_t
!GCC$ ATTRIBUTES NO_ARG_CHECK :: a
@@ -426,6 +462,14 @@ module openacc_internal
integer (c_size_t), value :: len
end subroutine
+ subroutine acc_copyout_finalize_l (a, len) &
+ bind (C, name = "acc_copyout_finalize")
+ use iso_c_binding, only: c_size_t
+ !GCC$ ATTRIBUTES NO_ARG_CHECK :: a
+ type (*), dimension (*) :: a
+ integer (c_size_t), value :: len
+ end subroutine
+
subroutine acc_delete_l (a, len) &
bind (C, name = "acc_delete")
use iso_c_binding, only: c_size_t
@@ -434,6 +478,14 @@ module openacc_internal
integer (c_size_t), value :: len
end subroutine
+ subroutine acc_delete_finalize_l (a, len) &
+ bind (C, name = "acc_delete_finalize")
+ use iso_c_binding, only: c_size_t
+ !GCC$ ATTRIBUTES NO_ARG_CHECK :: a
+ type (*), dimension (*) :: a
+ integer (c_size_t), value :: len
+ end subroutine
+
subroutine acc_update_device_l (a, len) &
bind (C, name = "acc_update_device")
use iso_c_binding, only: c_size_t
@@ -598,12 +650,24 @@ module openacc
procedure :: acc_copyout_array_h
end interface
+ interface acc_copyout_finalize
+ procedure :: acc_copyout_finalize_32_h
+ procedure :: acc_copyout_finalize_64_h
+ procedure :: acc_copyout_finalize_array_h
+ end interface
+
interface acc_delete
procedure :: acc_delete_32_h
procedure :: acc_delete_64_h
procedure :: acc_delete_array_h
end interface
+ interface acc_delete_finalize
+ procedure :: acc_delete_finalize_32_h
+ procedure :: acc_delete_finalize_64_h
+ procedure :: acc_delete_finalize_array_h
+ end interface
+
interface acc_update_device
procedure :: acc_update_device_32_h
procedure :: acc_update_device_64_h
@@ -860,6 +924,30 @@ subroutine acc_copyout_array_h (a)
call acc_copyout_l (a, sizeof (a))
end subroutine
+subroutine acc_copyout_finalize_32_h (a, len)
+ use iso_c_binding, only: c_int32_t, c_size_t
+ use openacc_internal, only: acc_copyout_finalize_l
+ !GCC$ ATTRIBUTES NO_ARG_CHECK :: a
+ type (*), dimension (*) :: a
+ integer (c_int32_t) len
+ call acc_copyout_finalize_l (a, int (len, kind = c_size_t))
+end subroutine
+
+subroutine acc_copyout_finalize_64_h (a, len)
+ use iso_c_binding, only: c_int64_t, c_size_t
+ use openacc_internal, only: acc_copyout_finalize_l
+ !GCC$ ATTRIBUTES NO_ARG_CHECK :: a
+ type (*), dimension (*) :: a
+ integer (c_int64_t) len
+ call acc_copyout_finalize_l (a, int (len, kind = c_size_t))
+end subroutine
+
+subroutine acc_copyout_finalize_array_h (a)
+ use openacc_internal, only: acc_copyout_finalize_l
+ type (*), dimension (..), contiguous :: a
+ call acc_copyout_finalize_l (a, sizeof (a))
+end subroutine
+
subroutine acc_delete_32_h (a, len)
use iso_c_binding, only: c_int32_t, c_size_t
use openacc_internal, only: acc_delete_l
@@ -884,6 +972,30 @@ subroutine acc_delete_array_h (a)
call acc_delete_l (a, sizeof (a))
end subroutine
+subroutine acc_delete_finalize_32_h (a, len)
+ use iso_c_binding, only: c_int32_t, c_size_t
+ use openacc_internal, only: acc_delete_finalize_l
+ !GCC$ ATTRIBUTES NO_ARG_CHECK :: a
+ type (*), dimension (*) :: a
+ integer (c_int32_t) len
+ call acc_delete_finalize_l (a, int (len, kind = c_size_t))
+end subroutine
+
+subroutine acc_delete_finalize_64_h (a, len)
+ use iso_c_binding, only: c_int64_t, c_size_t
+ use openacc_internal, only: acc_delete_finalize_l
+ !GCC$ ATTRIBUTES NO_ARG_CHECK :: a
+ type (*), dimension (*) :: a
+ integer (c_int64_t) len
+ call acc_delete_finalize_l (a, int (len, kind = c_size_t))
+end subroutine
+
+subroutine acc_delete_finalize_array_h (a)
+ use openacc_internal, only: acc_delete_finalize_l
+ type (*), dimension (..), contiguous :: a
+ call acc_delete_finalize_l (a, sizeof (a))
+end subroutine
+
subroutine acc_update_device_32_h (a, len)
use iso_c_binding, only: c_int32_t, c_size_t
use openacc_internal, only: acc_update_device_l
diff --git a/libgomp/openacc.h b/libgomp/openacc.h
index b8572574f13..02a85a09ddb 100644
--- a/libgomp/openacc.h
+++ b/libgomp/openacc.h
@@ -109,6 +109,12 @@ int acc_is_present (void *, size_t) __GOACC_NOTHROW;
void acc_memcpy_to_device (void *, void *, size_t) __GOACC_NOTHROW;
void acc_memcpy_from_device (void *, void *, size_t) __GOACC_NOTHROW;
+/* Finalize versions of copyout/delete functions, specified in OpenACC 2.5. */
+void acc_copyout_finalize (void *, size_t) __GOACC_NOTHROW;
+void acc_copyout_finalize_async (void *, size_t, int) __GOACC_NOTHROW;
+void acc_delete_finalize (void *, size_t) __GOACC_NOTHROW;
+void acc_delete_finalize_async (void *, size_t, int) __GOACC_NOTHROW;
+
/* CUDA-specific routines. */
void *acc_get_current_cuda_device (void) __GOACC_NOTHROW;
void *acc_get_current_cuda_context (void) __GOACC_NOTHROW;
diff --git a/libgomp/openacc_lib.h b/libgomp/openacc_lib.h
index 5cf743c2491..737c582041d 100644
--- a/libgomp/openacc_lib.h
+++ b/libgomp/openacc_lib.h
@@ -273,6 +273,26 @@
end subroutine
end interface
+ interface acc_copyout_finalize
+ subroutine acc_copyout_finalize_32_h (a, len)
+ use iso_c_binding, only: c_int32_t
+ !GCC$ ATTRIBUTES NO_ARG_CHECK :: a
+ type (*), dimension (*) :: a
+ integer (c_int32_t) len
+ end subroutine
+
+ subroutine acc_copyout_finalize_64_h (a, len)
+ use iso_c_binding, only: c_int64_t
+ !GCC$ ATTRIBUTES NO_ARG_CHECK :: a
+ type (*), dimension (*) :: a
+ integer (c_int64_t) len
+ end subroutine
+
+ subroutine acc_copyout_finalize_array_h (a)
+ type (*), dimension (..), contiguous :: a
+ end subroutine
+ end interface
+
interface acc_delete
subroutine acc_delete_32_h (a, len)
use iso_c_binding, only: c_int32_t
@@ -293,6 +313,26 @@
end subroutine
end interface
+ interface acc_delete_finalize
+ subroutine acc_delete_finalize_32_h (a, len)
+ use iso_c_binding, only: c_int32_t
+ !GCC$ ATTRIBUTES NO_ARG_CHECK :: a
+ type (*), dimension (*) :: a
+ integer (c_int32_t) len
+ end subroutine
+
+ subroutine acc_delete_finalize_64_h (a, len)
+ use iso_c_binding, only: c_int64_t
+ !GCC$ ATTRIBUTES NO_ARG_CHECK :: a
+ type (*), dimension (*) :: a
+ integer (c_int64_t) len
+ end subroutine
+
+ subroutine acc_delete_finalize_array_h (a)
+ type (*), dimension (..), contiguous :: a
+ end subroutine
+ end interface
+
interface acc_update_device
subroutine acc_update_device_32_h (a, len)
use iso_c_binding, only: c_int32_t
diff --git a/libgomp/target.c b/libgomp/target.c
index 509776d17a8..dda041cdbef 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -859,6 +859,7 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
tgt->list[i].offset = 0;
tgt->list[i].length = k->host_end - k->host_start;
k->refcount = 1;
+ k->dynamic_refcount = 0;
tgt->refcount++;
array->left = NULL;
array->right = NULL;
@@ -1011,6 +1012,23 @@ gomp_unmap_tgt (struct target_mem_desc *tgt)
free (tgt);
}
+attribute_hidden bool
+gomp_remove_var (struct gomp_device_descr *devicep, splay_tree_key k)
+{
+ bool is_tgt_unmapped = false;
+ splay_tree_remove (&devicep->mem_map, k);
+ if (k->link_key)
+ splay_tree_insert (&devicep->mem_map, (splay_tree_node) k->link_key);
+ if (k->tgt->refcount > 1)
+ k->tgt->refcount--;
+ else
+ {
+ is_tgt_unmapped = true;
+ gomp_unmap_tgt (k->tgt);
+ }
+ return is_tgt_unmapped;
+}
+
/* Unmap variables described by TGT. If DO_COPYFROM is true, copy relevant
variables back from device to host: if it is false, it is assumed that this
has been done already. */
@@ -1059,16 +1077,7 @@ gomp_unmap_vars (struct target_mem_desc *tgt, bool do_copyfrom)
+ tgt->list[i].offset),
tgt->list[i].length);
if (do_unmap)
- {
- splay_tree_remove (&devicep->mem_map, k);
- if (k->link_key)
- splay_tree_insert (&devicep->mem_map,
- (splay_tree_node) k->link_key);
- if (k->tgt->refcount > 1)
- k->tgt->refcount--;
- else
- gomp_unmap_tgt (k->tgt);
- }
+ gomp_remove_var (devicep, k);
}
if (tgt->refcount > 1)
@@ -1298,17 +1307,7 @@ gomp_unload_image_from_device (struct gomp_device_descr *devicep,
else
{
splay_tree_key n = splay_tree_lookup (&devicep->mem_map, &k);
- splay_tree_remove (&devicep->mem_map, n);
- if (n->link_key)
- {
- if (n->tgt->refcount > 1)
- n->tgt->refcount--;
- else
- {
- is_tgt_unmapped = true;
- gomp_unmap_tgt (n->tgt);
- }
- }
+ is_tgt_unmapped = gomp_remove_var (devicep, n);
}
}
--
2.17.1
^ permalink raw reply [flat|nested] 31+ messages in thread* Re: [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior - runtime
2018-06-19 17:01 ` [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior - runtime Cesar Philippidis
@ 2018-06-20 16:45 ` Jakub Jelinek
2018-06-20 16:59 ` Cesar Philippidis
2019-05-02 14:03 ` Thomas Schwinge
` (3 subsequent siblings)
4 siblings, 1 reply; 31+ messages in thread
From: Jakub Jelinek @ 2018-06-20 16:45 UTC (permalink / raw)
To: Cesar Philippidis; +Cc: gcc-patches
On Tue, Jun 19, 2018 at 10:01:20AM -0700, Cesar Philippidis wrote:
> This patch implements the OpenACC 2.5 data clause semantics in libgomp.
>
> Is it OK for trunk?
> 2018-06-19 Chung-Lin Tang <cltang@codesourcery.com>
> Thomas Schwinge <thomas@codesourcery.com>
> Cesar Philippidis <cesar@codesourcery.com>
>
> libgomp/
> * libgomp.h (struct splay_tree_key_s): Add dynamic_refcount member.
> (gomp_acc_remove_pointer): Update declaration.
> (gomp_acc_declare_allocate): Declare.
> (gomp_remove_var): Declare.
> * libgomp.map (OACC_2.5): Define.
> * oacc-mem.c (acc_map_data): Update refcount.
> (acc_unmap_data): Likewise.
> (present_create_copy): Likewise.
> (acc_create): Add FLAG_PRESENT when calling present_create_copy.
> (acc_copyin): Likewise.
> (FLAG_FINALIZE): Define.
> (delete_copyout): Update dynamic refcounts, add support for FINALIZE.
> (acc_delete_finalize): New function.
> (acc_delete_finalize_async): New function.
> (acc_copyout_finalize): New function.
> (acc_copyout_finalize_async): New function.
> (gomp_acc_insert_pointer): Update refcounts.
> (gomp_acc_remove_pointer): Return if data is not present on the
> accelerator.
> * oacc-parallel.c (find_pset): Rename to find_pointer.
> (find_pointer): Add support for GOMP_MAP_POINTER.
> (handle_ftn_pointers): New function.
> (GOACC_parallel_keyed): Update refcounts of variables.
> (GOACC_enter_exit_data): Add support for finalized data mappings.
> Add support for GOMP_MAP_{TO,ALLOC,RELESE,FROM}. Update handling
> of fortran arrays.
> (GOACC_update): Add support for GOMP_MAP_{ALWAYS_POINTER,TO,FROM}.
> (GOACC_declare): Add support for GOMP_MAP_RELEASE, remove support
> for GOMP_MAP_FORCE_FROM.
> * openacc.f90 (module openacc_internal): Add
> acc_copyout_finalize_{32_h,64_h,array_h,_l}, and
> acc_delete_finalize_{32_h,64_h,array_h,_l}. Add interfaces for
> acc_copyout_finalize and acc_delete_finalize.
> (acc_copyout_finalize_32_h): New subroutine.
> (acc_copyout_finalize_64_h): New subroutine.
> (acc_copyout_finalize_array_h): New subroutine.
> (acc_delete_finalize_32_h): New subroutine.
> (acc_delete_finalize_64_h): New subroutine.
> (acc_delete_finalize_array_h): New subroutine.
> * openacc.h (acc_copyout_finalize): Declare.
> (acc_copyout_finalize_async): Declare.
> (acc_delete_finalize): Declare.
> (acc_delete_finalize_async): Declare.
> * openacc_lib.h (acc_copyout_finalize): New interface.
> (acc_delete_finalize): New interface.
> * target.c (gomp_map_vars): Update dynamic_refcount.
> (gomp_remove_var): New function.
> (gomp_unmap_vars): Use it.
> (gomp_unload_image_from_device): Likewise.
>
>
> >From 53ee03231c5e6e4747b4ef01335079a2d4a98480 Mon Sep 17 00:00:00 2001
> From: Cesar Philippidis <cesar@codesourcery.com>
> Date: Tue, 19 Jun 2018 09:33:04 -0700
> Subject: [PATCH 7/7] runtime changes
>
> ---
> libgomp/libgomp.h | 7 +-
> libgomp/libgomp.map | 12 +++
> libgomp/oacc-mem.c | 196 ++++++++++++++++++++++++++++++++-------
> libgomp/oacc-parallel.c | 198 ++++++++++++++++++++++++++++++++++------
> libgomp/openacc.f90 | 112 +++++++++++++++++++++++
> libgomp/openacc.h | 6 ++
> libgomp/openacc_lib.h | 40 ++++++++
> libgomp/target.c | 41 ++++-----
> 8 files changed, 528 insertions(+), 84 deletions(-)
>
> diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h
> index 10ea8940c96..3a8cc2bd7d6 100644
> --- a/libgomp/libgomp.h
> +++ b/libgomp/libgomp.h
> @@ -853,6 +853,8 @@ struct splay_tree_key_s {
> uintptr_t tgt_offset;
> /* Reference count. */
> uintptr_t refcount;
> + /* Dynamic reference count. */
> + uintptr_t dynamic_refcount;
> /* Pointer to the original mapping of "omp declare target link" object. */
> splay_tree_key link_key;
> };
I'm not entirely happy about this, it grows the structure for not just
OpenACC, but also OpenMP which will never use it. Are there any fields
not used by OpenACC? E.g. is link_key used?
Or could the dynamic refcounts be an array allocated (for OpenACC mappings
only) after the tgt->array array, accessed using
key->tgt->dynamic_refcounts[key - key->tgt->array] ?
Jakub
^ permalink raw reply [flat|nested] 31+ messages in thread* Re: [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior - runtime
2018-06-20 16:45 ` Jakub Jelinek
@ 2018-06-20 16:59 ` Cesar Philippidis
2018-06-20 17:03 ` Jakub Jelinek
0 siblings, 1 reply; 31+ messages in thread
From: Cesar Philippidis @ 2018-06-20 16:59 UTC (permalink / raw)
To: Jakub Jelinek, Chung-Lin Tang; +Cc: gcc-patches
On 06/20/2018 09:45 AM, Jakub Jelinek wrote:
> On Tue, Jun 19, 2018 at 10:01:20AM -0700, Cesar Philippidis wrote:
>> >From 53ee03231c5e6e4747b4ef01335079a2d4a98480 Mon Sep 17 00:00:00 2001
>> From: Cesar Philippidis <cesar@codesourcery.com>
>> Date: Tue, 19 Jun 2018 09:33:04 -0700
>> Subject: [PATCH 7/7] runtime changes
>>
>> ---
>> libgomp/libgomp.h | 7 +-
>> libgomp/libgomp.map | 12 +++
>> libgomp/oacc-mem.c | 196 ++++++++++++++++++++++++++++++++-------
>> libgomp/oacc-parallel.c | 198 ++++++++++++++++++++++++++++++++++------
>> libgomp/openacc.f90 | 112 +++++++++++++++++++++++
>> libgomp/openacc.h | 6 ++
>> libgomp/openacc_lib.h | 40 ++++++++
>> libgomp/target.c | 41 ++++-----
>> 8 files changed, 528 insertions(+), 84 deletions(-)
>>
>> diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h
>> index 10ea8940c96..3a8cc2bd7d6 100644
>> --- a/libgomp/libgomp.h
>> +++ b/libgomp/libgomp.h
>> @@ -853,6 +853,8 @@ struct splay_tree_key_s {
>> uintptr_t tgt_offset;
>> /* Reference count. */
>> uintptr_t refcount;
>> + /* Dynamic reference count. */
>> + uintptr_t dynamic_refcount;
>> /* Pointer to the original mapping of "omp declare target link" object. */
>> splay_tree_key link_key;
>> };
>
> I'm not entirely happy about this, it grows the structure for not just
> OpenACC, but also OpenMP which will never use it. Are there any fields
> not used by OpenACC? E.g. is link_key used?
> Or could the dynamic refcounts be an array allocated (for OpenACC mappings
> only) after the tgt->array array, accessed using
> key->tgt->dynamic_refcounts[key - key->tgt->array] ?
Sorry, I mistakenly committed this patch with the front end changes. Can
I address this issue in a follow up patch?
If it means anything, we have a significant async change that removes
the async_refcount field in that struct.
Cesar
^ permalink raw reply [flat|nested] 31+ messages in thread* Re: [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior - runtime
2018-06-20 16:59 ` Cesar Philippidis
@ 2018-06-20 17:03 ` Jakub Jelinek
2018-06-20 17:07 ` Cesar Philippidis
0 siblings, 1 reply; 31+ messages in thread
From: Jakub Jelinek @ 2018-06-20 17:03 UTC (permalink / raw)
To: Cesar Philippidis; +Cc: Chung-Lin Tang, gcc-patches
On Wed, Jun 20, 2018 at 09:59:29AM -0700, Cesar Philippidis wrote:
> > I'm not entirely happy about this, it grows the structure for not just
> > OpenACC, but also OpenMP which will never use it. Are there any fields
> > not used by OpenACC? E.g. is link_key used?
> > Or could the dynamic refcounts be an array allocated (for OpenACC mappings
> > only) after the tgt->array array, accessed using
> > key->tgt->dynamic_refcounts[key - key->tgt->array] ?
> Sorry, I mistakenly committed this patch with the front end changes. Can
> I address this issue in a follow up patch?
Yes. If it isn't possible to get rid of it, I can live with it, but would
appreciate if you tried to avoid it if possible.
> If it means anything, we have a significant async change that removes
> the async_refcount field in that struct.
Wasn't async_refcount removed 2 years ago?
Jakub
^ permalink raw reply [flat|nested] 31+ messages in thread
* Re: [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior - runtime
2018-06-20 17:03 ` Jakub Jelinek
@ 2018-06-20 17:07 ` Cesar Philippidis
0 siblings, 0 replies; 31+ messages in thread
From: Cesar Philippidis @ 2018-06-20 17:07 UTC (permalink / raw)
To: Jakub Jelinek; +Cc: Chung-Lin Tang, gcc-patches
On 06/20/2018 10:03 AM, Jakub Jelinek wrote:
> On Wed, Jun 20, 2018 at 09:59:29AM -0700, Cesar Philippidis wrote:
>> If it means anything, we have a significant async change that removes
>> the async_refcount field in that struct.
>
> Wasn't async_refcount removed 2 years ago?
You're right. I was looking at the og8 history. I'm juggling a lot of
patches right now.
Cesar
^ permalink raw reply [flat|nested] 31+ messages in thread
* Re: [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior - runtime
2018-06-19 17:01 ` [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior - runtime Cesar Philippidis
2018-06-20 16:45 ` Jakub Jelinek
@ 2019-05-02 14:03 ` Thomas Schwinge
2019-05-29 14:32 ` Thomas Schwinge
2019-12-09 11:44 ` In 'libgomp/target.c:gomp_exit_data', remove open-coded 'gomp_remove_var' (was: [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior - runtime) Thomas Schwinge
` (2 subsequent siblings)
4 siblings, 1 reply; 31+ messages in thread
From: Thomas Schwinge @ 2019-05-02 14:03 UTC (permalink / raw)
To: Jakub Jelinek; +Cc: gcc-patches
[-- Attachment #1: Type: text/plain, Size: 2897 bytes --]
Hi Jakub!
I'm currently working on other pending OpenACC 'deviceptr' clause patches
from our backlog, and I noticed the following, which I don't understand.
You reviewed and approved this patch, could you please help?
On Tue, 19 Jun 2018 10:01:20 -0700, Cesar Philippidis <cesar@codesourcery.com> wrote:
> --- a/libgomp/oacc-parallel.c
> +++ b/libgomp/oacc-parallel.c
> +/* Handle the mapping pair that are presented when a
> + deviceptr clause is used with Fortran. */
> +
> +static void
> +handle_ftn_pointers (size_t mapnum, void **hostaddrs, size_t *sizes,
> + unsigned short *kinds)
> +{
> + int i;
> +
> + for (i = 0; i < mapnum; i++)
> + {
> + unsigned short kind1 = kinds[i] & 0xff;
> +
> + /* Handle Fortran deviceptr clause. */
> + if (kind1 == GOMP_MAP_FORCE_DEVICEPTR)
> + {
> + unsigned short kind2;
> +
> + if (i < (signed)mapnum - 1)
> + kind2 = kinds[i + 1] & 0xff;
> + else
> + kind2 = 0xffff;
> +
> + if (sizes[i] == sizeof (void *))
> + continue;
> +
> + /* At this point, we're dealing with a Fortran deviceptr.
> + If the next element is not what we're expecting, then
> + this is an instance of where the deviceptr variable was
> + not used within the region and the pointer was removed
> + by the gimplifier. */
> + if (kind2 == GOMP_MAP_POINTER
> + && sizes[i + 1] == 0
> + && hostaddrs[i] == *(void **)hostaddrs[i + 1])
> + {
> + kinds[i+1] = kinds[i];
> + sizes[i+1] = sizeof (void *);
> + }
> +
> + /* Invalidate the entry. */
> + hostaddrs[i] = NULL;
> + }
> + }
> }
This is used for rewriting the mappings for OpenACC 'parallel'
etc. constructs:
> @@ -88,6 +141,8 @@ GOACC_parallel_keyed (int device, void (*fn) (void *),
> thr = goacc_thread ();
> acc_dev = thr->dev;
>
> + handle_ftn_pointers (mapnum, hostaddrs, sizes, kinds);
> +
> /* Host fallback if "if" clause is false or if the current device is set to
> the host. */
> if (host_fallback)
..., and on our OpenACC development branch likewise for OpenACC 'data'
constructs ('GOACC_data_start').
What this function seems to be doing, as I understand this, is that when
there is an 'GOMP_MAP_FORCE_DEVICEPTR' with a size not eqal to pointer
size (which should never happen, as per the information given
'include/gomp-constants.h'?), which is followed by a 'GOMP_MAP_POINTER',
then preserve the 'GOMP_MAP_FORCE_DEVICEPTR' (by storing it into the slot
of the 'GOMP_MAP_POINTER'), and unconditionally remove the
'GOMP_MAP_POINTER'. This seems like a strange choice of a GCC/libgomp
ABI to me -- or am I not understanding this correctly?
Instead of rewriting the mappings at run time, why isn't (presumably) the
gimplifier changed to just emit the correct mappings?
Grüße
Thomas
[-- Attachment #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 658 bytes --]
^ permalink raw reply [flat|nested] 31+ messages in thread* Re: [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior - runtime
2019-05-02 14:03 ` Thomas Schwinge
@ 2019-05-29 14:32 ` Thomas Schwinge
0 siblings, 0 replies; 31+ messages in thread
From: Thomas Schwinge @ 2019-05-29 14:32 UTC (permalink / raw)
To: Jakub Jelinek; +Cc: gcc-patches
[-- Attachment #1: Type: text/plain, Size: 3144 bytes --]
Hi Jakub!
Any comments on my questions, please?
On Thu, 02 May 2019 16:03:09 +0200, I wrote:
> I'm currently working on other pending OpenACC 'deviceptr' clause patches
> from our backlog, and I noticed the following, which I don't understand.
> You reviewed and approved this patch, could you please help?
>
> On Tue, 19 Jun 2018 10:01:20 -0700, Cesar Philippidis <cesar@codesourcery.com> wrote:
> > --- a/libgomp/oacc-parallel.c
> > +++ b/libgomp/oacc-parallel.c
>
> > +/* Handle the mapping pair that are presented when a
> > + deviceptr clause is used with Fortran. */
> > +
> > +static void
> > +handle_ftn_pointers (size_t mapnum, void **hostaddrs, size_t *sizes,
> > + unsigned short *kinds)
> > +{
> > + int i;
> > +
> > + for (i = 0; i < mapnum; i++)
> > + {
> > + unsigned short kind1 = kinds[i] & 0xff;
> > +
> > + /* Handle Fortran deviceptr clause. */
> > + if (kind1 == GOMP_MAP_FORCE_DEVICEPTR)
> > + {
> > + unsigned short kind2;
> > +
> > + if (i < (signed)mapnum - 1)
> > + kind2 = kinds[i + 1] & 0xff;
> > + else
> > + kind2 = 0xffff;
> > +
> > + if (sizes[i] == sizeof (void *))
> > + continue;
> > +
> > + /* At this point, we're dealing with a Fortran deviceptr.
> > + If the next element is not what we're expecting, then
> > + this is an instance of where the deviceptr variable was
> > + not used within the region and the pointer was removed
> > + by the gimplifier. */
> > + if (kind2 == GOMP_MAP_POINTER
> > + && sizes[i + 1] == 0
> > + && hostaddrs[i] == *(void **)hostaddrs[i + 1])
> > + {
> > + kinds[i+1] = kinds[i];
> > + sizes[i+1] = sizeof (void *);
> > + }
> > +
> > + /* Invalidate the entry. */
> > + hostaddrs[i] = NULL;
> > + }
> > + }
> > }
>
> This is used for rewriting the mappings for OpenACC 'parallel'
> etc. constructs:
>
> > @@ -88,6 +141,8 @@ GOACC_parallel_keyed (int device, void (*fn) (void *),
> > thr = goacc_thread ();
> > acc_dev = thr->dev;
> >
> > + handle_ftn_pointers (mapnum, hostaddrs, sizes, kinds);
> > +
> > /* Host fallback if "if" clause is false or if the current device is set to
> > the host. */
> > if (host_fallback)
>
> ..., and on our OpenACC development branch likewise for OpenACC 'data'
> constructs ('GOACC_data_start').
>
> What this function seems to be doing, as I understand this, is that when
> there is an 'GOMP_MAP_FORCE_DEVICEPTR' with a size not eqal to pointer
> size (which should never happen, as per the information given
> 'include/gomp-constants.h'?), which is followed by a 'GOMP_MAP_POINTER',
> then preserve the 'GOMP_MAP_FORCE_DEVICEPTR' (by storing it into the slot
> of the 'GOMP_MAP_POINTER'), and unconditionally remove the
> 'GOMP_MAP_POINTER'. This seems like a strange choice of a GCC/libgomp
> ABI to me -- or am I not understanding this correctly?
>
> Instead of rewriting the mappings at run time, why isn't (presumably) the
> gimplifier changed to just emit the correct mappings?
Grüße
Thomas
[-- Attachment #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 658 bytes --]
^ permalink raw reply [flat|nested] 31+ messages in thread
* In 'libgomp/target.c:gomp_exit_data', remove open-coded 'gomp_remove_var' (was: [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior - runtime)
2018-06-19 17:01 ` [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior - runtime Cesar Philippidis
2018-06-20 16:45 ` Jakub Jelinek
2019-05-02 14:03 ` Thomas Schwinge
@ 2019-12-09 11:44 ` Thomas Schwinge
2019-12-11 16:59 ` [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior - runtime Thomas Schwinge
2020-05-19 13:58 ` 'gomp_map_vars' locking protocol (was: [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior - runtime) Thomas Schwinge
4 siblings, 0 replies; 31+ messages in thread
From: Thomas Schwinge @ 2019-12-09 11:44 UTC (permalink / raw)
To: gcc-patches, Jakub Jelinek; +Cc: Julian Brown
[-- Attachment #1.1: Type: text/plain, Size: 1896 bytes --]
Hi!
On 2018-06-19T10:01:20-0700, Cesar Philippidis <cesar@codesourcery.com> wrote:
> --- a/libgomp/target.c
> +++ b/libgomp/target.c
> +attribute_hidden bool
> +gomp_remove_var (struct gomp_device_descr *devicep, splay_tree_key k)
> +{
> + bool is_tgt_unmapped = false;
> + splay_tree_remove (&devicep->mem_map, k);
> + if (k->link_key)
> + splay_tree_insert (&devicep->mem_map, (splay_tree_node) k->link_key);
> + if (k->tgt->refcount > 1)
> + k->tgt->refcount--;
> + else
> + {
> + is_tgt_unmapped = true;
> + gomp_unmap_tgt (k->tgt);
> + }
> + return is_tgt_unmapped;
This new function, can, like done here:
> @@ -1059,16 +1077,7 @@ gomp_unmap_vars (struct target_mem_desc *tgt, bool do_copyfrom)
> + tgt->list[i].offset),
> tgt->list[i].length);
> if (do_unmap)
> - {
> - splay_tree_remove (&devicep->mem_map, k);
> - if (k->link_key)
> - splay_tree_insert (&devicep->mem_map,
> - (splay_tree_node) k->link_key);
> - if (k->tgt->refcount > 1)
> - k->tgt->refcount--;
> - else
> - gomp_unmap_tgt (k->tgt);
> - }
> + gomp_remove_var (devicep, k);
> }
..., and here:
> @@ -1298,17 +1307,7 @@ gomp_unload_image_from_device (struct gomp_device_descr *devicep,
> else
> {
> splay_tree_key n = splay_tree_lookup (&devicep->mem_map, &k);
> - splay_tree_remove (&devicep->mem_map, n);
> - if (n->link_key)
> - {
> - if (n->tgt->refcount > 1)
> - n->tgt->refcount--;
> - else
> - {
> - is_tgt_unmapped = true;
> - gomp_unmap_tgt (n->tgt);
> - }
> - }
> + is_tgt_unmapped = gomp_remove_var (devicep, n);
> }
..., also be used in 'gomp_exit_data', see attached "In
'libgomp/target.c:gomp_exit_data', remove open-coded 'gomp_remove_var'",
committed to trunk in r279118.
Grüße
Thomas
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #1.2: 0001-In-libgomp-target.c-gomp_exit_data-remove-open.trunk.patch --]
[-- Type: text/x-diff, Size: 1652 bytes --]
From bbfdb255a0b5cb6e183e11026c2a482d4eeba981 Mon Sep 17 00:00:00 2001
From: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Mon, 9 Dec 2019 11:39:57 +0000
Subject: [PATCH] In 'libgomp/target.c:gomp_exit_data', remove open-coded
'gomp_remove_var'
libgomp/
* target.c (gomp_exit_data): Use 'gomp_remove_var'.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@279118 138bc75d-0d04-0410-961f-82ee72b054a4
---
libgomp/ChangeLog | 4 ++++
libgomp/target.c | 11 +----------
2 files changed, 5 insertions(+), 10 deletions(-)
diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index a0bd25177d1..c5541bcec81 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,7 @@
+2019-12-09 Thomas Schwinge <thomas@codesourcery.com>
+
+ * target.c (gomp_exit_data): Use 'gomp_remove_var'.
+
2019-12-09 Tobias Burnus <tobias@codesourcery.com>
* testsuite/libgomp.fortran/use_device_addr-3.f90: Make 'stop' codes
diff --git a/libgomp/target.c b/libgomp/target.c
index 84d6daa76ca..13f7921651f 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -2095,16 +2095,7 @@ gomp_exit_data (struct gomp_device_descr *devicep, size_t mapnum,
- k->host_start),
cur_node.host_end - cur_node.host_start);
if (k->refcount == 0)
- {
- splay_tree_remove (&devicep->mem_map, k);
- if (k->link_key)
- splay_tree_insert (&devicep->mem_map,
- (splay_tree_node) k->link_key);
- if (k->tgt->refcount > 1)
- k->tgt->refcount--;
- else
- gomp_unmap_tgt (k->tgt);
- }
+ gomp_remove_var (devicep, k);
break;
default:
--
2.17.1
[-- Attachment #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 658 bytes --]
^ permalink raw reply [flat|nested] 31+ messages in thread* Re: [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior - runtime
2018-06-19 17:01 ` [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior - runtime Cesar Philippidis
` (2 preceding siblings ...)
2019-12-09 11:44 ` In 'libgomp/target.c:gomp_exit_data', remove open-coded 'gomp_remove_var' (was: [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior - runtime) Thomas Schwinge
@ 2019-12-11 16:59 ` Thomas Schwinge
2020-05-19 13:58 ` 'gomp_map_vars' locking protocol (was: [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior - runtime) Thomas Schwinge
4 siblings, 0 replies; 31+ messages in thread
From: Thomas Schwinge @ 2019-12-11 16:59 UTC (permalink / raw)
To: gcc-patches, Jakub Jelinek; +Cc: Julian Brown
[-- Attachment #1.1: Type: text/plain, Size: 828 bytes --]
Hi!
On 2018-06-19T10:01:20-0700, Cesar Philippidis <cesar@codesourcery.com> wrote:
> This patch implements the OpenACC 2.5 data clause semantics in libgomp.
> --- a/libgomp/libgomp.h
> +++ b/libgomp/libgomp.h
> @@ -853,6 +853,8 @@ struct splay_tree_key_s {
> uintptr_t tgt_offset;
> /* Reference count. */
> uintptr_t refcount;
> + /* Dynamic reference count. */
> + uintptr_t dynamic_refcount;
> /* Pointer to the original mapping of "omp declare target link" object. */
> splay_tree_key link_key;
> };
See attached "[OpenACC] Initialize 'dynamic_refcount' whenever we
initialize 'refcount'" for 'Cases missed in r261813 "Update OpenACC data
clause semantics to the 2.5 behavior"'; committed to trunk in r279230,
and backported to gcc-9-branch in r279238.
Grüße
Thomas
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #1.2: 0001-OpenACC-Initialize-dynamic_refcount-whenever-w.trunk.patch --]
[-- Type: text/x-diff, Size: 2439 bytes --]
From 20d0998b970ba693b23ee24bd0c94ecb57adf281 Mon Sep 17 00:00:00 2001
From: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Wed, 11 Dec 2019 16:48:44 +0000
Subject: [PATCH] [OpenACC] Initialize 'dynamic_refcount' whenever we
initialize 'refcount'
Cases missed in r261813 "Update OpenACC data clause semantics to the 2.5
behavior".
libgomp/
* target.c (gomp_load_image_to_device, omp_target_associate_ptr):
Initialize 'dynamic_refcount' whenever we initialize 'refcount'.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@279230 138bc75d-0d04-0410-961f-82ee72b054a4
---
libgomp/ChangeLog | 6 ++++++
libgomp/target.c | 3 +++
2 files changed, 9 insertions(+)
diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index 6cefeba5f5f..6635ed7b44b 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,9 @@
+2019-12-11 Thomas Schwinge <thomas@codesourcery.com>
+ Julian Brown <julian@codesourcery.com>
+
+ * target.c (gomp_load_image_to_device, omp_target_associate_ptr):
+ Initialize 'dynamic_refcount' whenever we initialize 'refcount'.
+
2019-12-11 Tobias Burnus <tobias@codesourcery.com>
* omp_lib.h.in: Fix spelling of function declaration
diff --git a/libgomp/target.c b/libgomp/target.c
index 39a24f56395..1151debf256 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -1334,6 +1334,7 @@ gomp_load_image_to_device (struct gomp_device_descr *devicep, unsigned version,
k->tgt = tgt;
k->tgt_offset = target_table[i].start;
k->refcount = REFCOUNT_INFINITY;
+ k->dynamic_refcount = 0;
k->link_key = NULL;
array->left = NULL;
array->right = NULL;
@@ -1366,6 +1367,7 @@ gomp_load_image_to_device (struct gomp_device_descr *devicep, unsigned version,
k->tgt = tgt;
k->tgt_offset = target_var->start;
k->refcount = target_size & link_bit ? REFCOUNT_LINK : REFCOUNT_INFINITY;
+ k->dynamic_refcount = 0;
k->link_key = NULL;
array->left = NULL;
array->right = NULL;
@@ -2627,6 +2629,7 @@ omp_target_associate_ptr (const void *host_ptr, const void *device_ptr,
k->tgt = tgt;
k->tgt_offset = (uintptr_t) device_ptr + device_offset;
k->refcount = REFCOUNT_INFINITY;
+ k->dynamic_refcount = 0;
array->left = NULL;
array->right = NULL;
splay_tree_insert (&devicep->mem_map, array);
--
2.17.1
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #1.3: 0001-OpenACC-libgomp-Initialize-dynamic_refc.gcc-9-branch.patch --]
[-- Type: text/x-diff, Size: 2463 bytes --]
From f301776d131dd584f1259a4e6bfa5662451407c4 Mon Sep 17 00:00:00 2001
From: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Wed, 11 Dec 2019 16:51:31 +0000
Subject: [PATCH] [OpenACC, libgomp] Initialize 'dynamic_refcount' whenever we
initialize 'refcount'
Cases missed in r261813 "Update OpenACC data clause semantics to the 2.5
behavior".
libgomp/
* target.c (gomp_load_image_to_device, omp_target_associate_ptr):
Initialize 'dynamic_refcount' whenever we initialize 'refcount'.
Backport trunk r279230.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gcc-9-branch@279238 138bc75d-0d04-0410-961f-82ee72b054a4
---
libgomp/ChangeLog | 6 ++++++
libgomp/target.c | 3 +++
2 files changed, 9 insertions(+)
diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index 70a7f50c22b..c1959a44b8c 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,9 @@
+2019-12-11 Thomas Schwinge <thomas@codesourcery.com>
+ Julian Brown <julian@codesourcery.com>
+
+ * target.c (gomp_load_image_to_device, omp_target_associate_ptr):
+ Initialize 'dynamic_refcount' whenever we initialize 'refcount'.
+
2019-12-11 Tobias Burnus <tobias@codesourcery.com>
Backported from mainline
diff --git a/libgomp/target.c b/libgomp/target.c
index 31148003d0a..97fc1ee2ddc 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -1214,6 +1214,7 @@ gomp_load_image_to_device (struct gomp_device_descr *devicep, unsigned version,
k->tgt = tgt;
k->tgt_offset = target_table[i].start;
k->refcount = REFCOUNT_INFINITY;
+ k->dynamic_refcount = 0;
k->link_key = NULL;
array->left = NULL;
array->right = NULL;
@@ -1246,6 +1247,7 @@ gomp_load_image_to_device (struct gomp_device_descr *devicep, unsigned version,
k->tgt = tgt;
k->tgt_offset = target_var->start;
k->refcount = target_size & link_bit ? REFCOUNT_LINK : REFCOUNT_INFINITY;
+ k->dynamic_refcount = 0;
k->link_key = NULL;
array->left = NULL;
array->right = NULL;
@@ -2501,6 +2503,7 @@ omp_target_associate_ptr (const void *host_ptr, const void *device_ptr,
k->tgt = tgt;
k->tgt_offset = (uintptr_t) device_ptr + device_offset;
k->refcount = REFCOUNT_INFINITY;
+ k->dynamic_refcount = 0;
array->left = NULL;
array->right = NULL;
splay_tree_insert (&devicep->mem_map, array);
--
2.17.1
[-- Attachment #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 658 bytes --]
^ permalink raw reply [flat|nested] 31+ messages in thread* 'gomp_map_vars' locking protocol (was: [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior - runtime)
2018-06-19 17:01 ` [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior - runtime Cesar Philippidis
` (3 preceding siblings ...)
2019-12-11 16:59 ` [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior - runtime Thomas Schwinge
@ 2020-05-19 13:58 ` Thomas Schwinge
4 siblings, 0 replies; 31+ messages in thread
From: Thomas Schwinge @ 2020-05-19 13:58 UTC (permalink / raw)
To: Jakub Jelinek, Julian Brown; +Cc: gcc-patches
Hi Jakub, Julian!
Can you please help me understand the following:
On 2018-06-19T10:01:20-0700, Cesar Philippidis <cesar@codesourcery.com> wrote:
> This patch implements the OpenACC 2.5 data clause semantics in libgomp.
(This got committed as r261813, 2018-06-20. The code has seen some
changes in the mean time, but the underlying issue remains.)
> --- a/libgomp/oacc-mem.c
> +++ b/libgomp/oacc-mem.c
> @@ -347,6 +347,7 @@ acc_map_data (void *h, void *d, size_t s)
>
> tgt = gomp_map_vars (acc_dev, mapnum, &hostaddrs, &devaddrs, &sizes,
> &kinds, true, GOMP_MAP_VARS_OPENACC);
> + tgt->list[0].key->refcount = REFCOUNT_INFINITY;
> }
>
> gomp_mutex_lock (&acc_dev->lock);
Without 'acc_dev->lock' locked, we here touch the 'refcount' (via
'tgt->list[0].key->refcount'), as returned from 'gomp_map_vars' in the
case when entering a new mapping.
> @@ -483,6 +492,8 @@ present_create_copy (unsigned f, void *h, size_t s)
>
> tgt = gomp_map_vars (acc_dev, mapnum, &hostaddrs, NULL, &s, &kinds, true,
> GOMP_MAP_VARS_OPENACC);
> + /* Initialize dynamic refcount. */
> + tgt->list[0].key->dynamic_refcount = 1;
>
> gomp_mutex_lock (&acc_dev->lock);
Likewise here, for the 'dynamic_refcount' (via
'tgt->list[0].key->dynamic_refcount'), as returned from 'gomp_map_vars'
in the case when entering a new mapping.
By construction, it is safe to assume that 'tgt->list[0].key' is the 'n'
we're looking to modify: this is the case where we're entering a new
mapping.
But: is it safe to access this unlocked? It may seem so, as the new
mapping has not yet been exposed to user code, so only exists internal in
the respective libgomp functions. Yet, could still a concurrent
'acc_unmap_data'/'acc_delete'/etc. already "see" it (that is, look it up,
as it already has been entered into the mapping table), and unmap while
we're accessing 'tgt->list[0].key' here?
(Hmm, and actually a similar issue, if we consider the case of two
'gomp_map_vars' running concurrently?)
In that case, I suppose we should change the 'gomp_map_vars' interface
and all callers so that 'gomp_map_vars' always takes the device locked.
That doesn't appear problematic: locking the device is one of the first
things 'gomp_map_vars' does anyway (just not in the 'mapnum == 0' case,
but I suppose it's OK to pessimize that one?), and it remains locked
until the end of the function.
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] 31+ messages in thread
* Re: [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior - middle end
2018-06-19 16:56 ` Cesar Philippidis
` (3 preceding siblings ...)
2018-06-19 17:01 ` [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior - runtime Cesar Philippidis
@ 2018-06-19 17:01 ` Cesar Philippidis
2018-06-20 16:20 ` Jakub Jelinek
2018-06-19 17:02 ` [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior - compiler tests Cesar Philippidis
2018-06-19 17:03 ` [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior - runtime tests Cesar Philippidis
6 siblings, 1 reply; 31+ messages in thread
From: Cesar Philippidis @ 2018-06-19 17:01 UTC (permalink / raw)
To: gcc-patches, Jakub Jelinek
[-- Attachment #1: Type: text/plain, Size: 107 bytes --]
This patch implements the OpenACC 2.5 data clause semantics in the
middle end.
Is it OK for trunk?
Cesar
[-- Attachment #2: 0006-gcc-middle-end.patch --]
[-- Type: text/x-patch, Size: 13696 bytes --]
2018-06-19 Chung-Lin Tang <cltang@codesourcery.com>
Thomas Schwinge <thomas@codesourcery.com>
Cesar Philippidis <cesar@codesourcery.com>
gcc/c-family/
* c-pragma.h (enum pragma_omp_clause): Add
PRAGMA_OACC_CLAUSE_{FINALIZE,IF_PRESENT}. Remove
PRAGMA_OACC_CLAUSE_PRESENT_OR_{COPY,COPYIN,COPYOUT,CREATE}.
gcc/
* gimplify.c (gimplify_scan_omp_clauses): Add support for
OMP_CLAUSE_{IF_PRESENT,FINALIZE}.
(gimplify_adjust_omp_clauses): Likewise.
(gimplify_oacc_declare_1): Add support for GOMP_MAP_RELEASE, remove
support for GOMP_MAP_FORCE_{ALLOC,TO,FROM,TOFROM}.
(gimplify_omp_target_update): Update handling of acc update and
enter/exit data.
* omp-low.c (install_var_field): Remove unused parameter
base_pointers_restrict.
(scan_sharing_clauses): Remove base_pointers_restrict parameter.
Update call to install_var_field. Handle OMP_CLAUSE_{IF_PRESENT,
FINALIZE}
(omp_target_base_pointers_restrict_p): Delete.
(scan_omp_target): Update call to scan_sharing_clauses.
* tree-core.h (enum omp_clause_code): Add OMP_CLAUSE_{IF_PRESENT,
FINALIZE}.
* tree-nested.c (convert_nonlocal_omp_clauses): Handle
OMP_CLAUSE_{IF_PRESENT,FINALIZE}.
(convert_local_omp_clauses): Likewise.
* tree-pretty-print.c (dump_omp_clause): Likewise.
* tree.c (omp_clause_num_ops): Add entries for OMP_CLAUSE_{IF_PRESENT,
FINALIZE}.
(omp_clause_code_name): Likewise.
From f79b0e6f0d796dc18ef1faf20b9fad0b7feeaa94 Mon Sep 17 00:00:00 2001
From: Cesar Philippidis <cesar@codesourcery.com>
Date: Tue, 19 Jun 2018 09:32:20 -0700
Subject: [PATCH 6/7] gcc middle end
---
gcc/c-family/c-pragma.h | 6 +--
gcc/gimplify.c | 67 ++++++++++++++++++++++-------
gcc/omp-low.c | 93 +++++------------------------------------
gcc/tree-core.h | 8 +++-
gcc/tree-nested.c | 4 ++
gcc/tree-pretty-print.c | 6 +++
gcc/tree.c | 8 +++-
7 files changed, 88 insertions(+), 104 deletions(-)
diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h
index c70380c211b..b322547b11a 100644
--- a/gcc/c-family/c-pragma.h
+++ b/gcc/c-family/c-pragma.h
@@ -138,16 +138,13 @@ enum pragma_omp_clause {
PRAGMA_OACC_CLAUSE_DELETE,
PRAGMA_OACC_CLAUSE_DEVICEPTR,
PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT,
+ PRAGMA_OACC_CLAUSE_FINALIZE,
PRAGMA_OACC_CLAUSE_GANG,
PRAGMA_OACC_CLAUSE_HOST,
PRAGMA_OACC_CLAUSE_INDEPENDENT,
PRAGMA_OACC_CLAUSE_NUM_GANGS,
PRAGMA_OACC_CLAUSE_NUM_WORKERS,
PRAGMA_OACC_CLAUSE_PRESENT,
- PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY,
- PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN,
- PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT,
- PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE,
PRAGMA_OACC_CLAUSE_SELF,
PRAGMA_OACC_CLAUSE_SEQ,
PRAGMA_OACC_CLAUSE_TILE,
@@ -156,6 +153,7 @@ enum pragma_omp_clause {
PRAGMA_OACC_CLAUSE_VECTOR_LENGTH,
PRAGMA_OACC_CLAUSE_WAIT,
PRAGMA_OACC_CLAUSE_WORKER,
+ PRAGMA_OACC_CLAUSE_IF_PRESENT,
PRAGMA_OACC_CLAUSE_COLLAPSE = PRAGMA_OMP_CLAUSE_COLLAPSE,
PRAGMA_OACC_CLAUSE_COPYIN = PRAGMA_OMP_CLAUSE_COPYIN,
PRAGMA_OACC_CLAUSE_DEVICE = PRAGMA_OMP_CLAUSE_DEVICE,
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 1523a27e828..97543ed5f70 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -8524,6 +8524,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
case OMP_CLAUSE_NOGROUP:
case OMP_CLAUSE_THREADS:
case OMP_CLAUSE_SIMD:
+ case OMP_CLAUSE_IF_PRESENT:
+ case OMP_CLAUSE_FINALIZE:
break;
case OMP_CLAUSE_DEFAULTMAP:
@@ -9305,6 +9307,8 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
case OMP_CLAUSE_AUTO:
case OMP_CLAUSE_SEQ:
case OMP_CLAUSE_TILE:
+ case OMP_CLAUSE_IF_PRESENT:
+ case OMP_CLAUSE_FINALIZE:
break;
default:
@@ -9361,21 +9365,7 @@ gimplify_oacc_declare_1 (tree clause)
switch (kind)
{
case GOMP_MAP_ALLOC:
- case GOMP_MAP_FORCE_ALLOC:
- case GOMP_MAP_FORCE_TO:
- new_op = GOMP_MAP_DELETE;
- ret = true;
- break;
-
- case GOMP_MAP_FORCE_FROM:
- OMP_CLAUSE_SET_MAP_KIND (clause, GOMP_MAP_FORCE_ALLOC);
- new_op = GOMP_MAP_FORCE_FROM;
- ret = true;
- break;
-
- case GOMP_MAP_FORCE_TOFROM:
- OMP_CLAUSE_SET_MAP_KIND (clause, GOMP_MAP_FORCE_TO);
- new_op = GOMP_MAP_FORCE_FROM;
+ new_op = GOMP_MAP_RELEASE;
ret = true;
break;
@@ -10817,6 +10807,53 @@ gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p)
ort, TREE_CODE (expr));
gimplify_adjust_omp_clauses (pre_p, NULL, &OMP_STANDALONE_CLAUSES (expr),
TREE_CODE (expr));
+ if (TREE_CODE (expr) == OACC_UPDATE
+ && omp_find_clause (OMP_STANDALONE_CLAUSES (expr),
+ OMP_CLAUSE_IF_PRESENT))
+ {
+ /* The runtime uses GOMP_MAP_{TO,FROM} to denote the if_present
+ clause. */
+ for (tree c = OMP_STANDALONE_CLAUSES (expr); c; c = OMP_CLAUSE_CHAIN (c))
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)
+ switch (OMP_CLAUSE_MAP_KIND (c))
+ {
+ case GOMP_MAP_FORCE_TO:
+ OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TO);
+ break;
+ case GOMP_MAP_FORCE_FROM:
+ OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FROM);
+ break;
+ default:
+ break;
+ }
+ }
+ else if (TREE_CODE (expr) == OACC_EXIT_DATA
+ && omp_find_clause (OMP_STANDALONE_CLAUSES (expr),
+ OMP_CLAUSE_FINALIZE))
+ {
+ /* Use GOMP_MAP_DELETE/GOMP_MAP_FORCE_FROM to denote that "finalize"
+ semantics apply to all mappings of this OpenACC directive. */
+ bool finalize_marked = false;
+ for (tree c = OMP_STANDALONE_CLAUSES (expr); c; c = OMP_CLAUSE_CHAIN (c))
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)
+ switch (OMP_CLAUSE_MAP_KIND (c))
+ {
+ case GOMP_MAP_FROM:
+ OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_FROM);
+ finalize_marked = true;
+ break;
+ case GOMP_MAP_RELEASE:
+ OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_DELETE);
+ finalize_marked = true;
+ break;
+ default:
+ /* Check consistency: libgomp relies on the very first data
+ mapping clause being marked, so make sure we did that before
+ any other mapping clauses. */
+ gcc_assert (finalize_marked);
+ break;
+ }
+ }
stmt = gimple_build_omp_target (NULL, kind, OMP_STANDALONE_CLAUSES (expr));
gimplify_seq_add_stmt (pre_p, stmt);
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index ba6c705cf8b..c591231d8f1 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -642,8 +642,7 @@ build_sender_ref (tree var, omp_context *ctx)
BASE_POINTERS_RESTRICT, declare the field with restrict. */
static void
-install_var_field (tree var, bool by_ref, int mask, omp_context *ctx,
- bool base_pointers_restrict = false)
+install_var_field (tree var, bool by_ref, int mask, omp_context *ctx)
{
tree field, type, sfield = NULL_TREE;
splay_tree_key key = (splay_tree_key) var;
@@ -674,11 +673,7 @@ install_var_field (tree var, bool by_ref, int mask, omp_context *ctx,
type = build_pointer_type (build_pointer_type (type));
}
else if (by_ref)
- {
- type = build_pointer_type (type);
- if (base_pointers_restrict)
- type = build_qualified_type (type, TYPE_QUAL_RESTRICT);
- }
+ type = build_pointer_type (type);
else if ((mask & 3) == 1 && omp_is_reference (var))
type = TREE_TYPE (type);
@@ -992,12 +987,10 @@ fixup_child_record_type (omp_context *ctx)
}
/* Instantiate decls as necessary in CTX to satisfy the data sharing
- specified by CLAUSES. If BASE_POINTERS_RESTRICT, install var field with
- restrict. */
+ specified by CLAUSES. */
static void
-scan_sharing_clauses (tree clauses, omp_context *ctx,
- bool base_pointers_restrict = false)
+scan_sharing_clauses (tree clauses, omp_context *ctx)
{
tree c, decl;
bool scan_array_reductions = false;
@@ -1256,8 +1249,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx,
&& TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
install_var_field (decl, true, 7, ctx);
else
- install_var_field (decl, true, 3, ctx,
- base_pointers_restrict);
+ install_var_field (decl, true, 3, ctx);
if (is_gimple_omp_offloaded (ctx->stmt)
&& !OMP_CLAUSE_MAP_IN_REDUCTION (c))
install_var_local (decl, ctx);
@@ -1328,6 +1320,8 @@ scan_sharing_clauses (tree clauses, omp_context *ctx,
case OMP_CLAUSE_TILE:
case OMP_CLAUSE__SIMT_:
case OMP_CLAUSE_DEFAULT:
+ case OMP_CLAUSE_IF_PRESENT:
+ case OMP_CLAUSE_FINALIZE:
break;
case OMP_CLAUSE_ALIGNED:
@@ -1499,6 +1493,8 @@ scan_sharing_clauses (tree clauses, omp_context *ctx,
case OMP_CLAUSE_TILE:
case OMP_CLAUSE__GRIDDIM_:
case OMP_CLAUSE__SIMT_:
+ case OMP_CLAUSE_IF_PRESENT:
+ case OMP_CLAUSE_FINALIZE:
break;
case OMP_CLAUSE__CACHE_:
@@ -2266,68 +2262,6 @@ scan_omp_single (gomp_single *stmt, omp_context *outer_ctx)
layout_type (ctx->record_type);
}
-/* Return true if the CLAUSES of an omp target guarantee that the base pointers
- used in the corresponding offloaded function are restrict. */
-
-static bool
-omp_target_base_pointers_restrict_p (tree clauses)
-{
- /* The analysis relies on the GOMP_MAP_FORCE_* mapping kinds, which are only
- used by OpenACC. */
- if (flag_openacc == 0)
- return false;
-
- /* I. Basic example:
-
- void foo (void)
- {
- unsigned int a[2], b[2];
-
- #pragma acc kernels \
- copyout (a) \
- copyout (b)
- {
- a[0] = 0;
- b[0] = 1;
- }
- }
-
- After gimplification, we have:
-
- #pragma omp target oacc_kernels \
- map(force_from:a [len: 8]) \
- map(force_from:b [len: 8])
- {
- a[0] = 0;
- b[0] = 1;
- }
-
- Because both mappings have the force prefix, we know that they will be
- allocated when calling the corresponding offloaded function, which means we
- can mark the base pointers for a and b in the offloaded function as
- restrict. */
-
- tree c;
- for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
- {
- if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
- return false;
-
- switch (OMP_CLAUSE_MAP_KIND (c))
- {
- case GOMP_MAP_FORCE_ALLOC:
- case GOMP_MAP_FORCE_TO:
- case GOMP_MAP_FORCE_FROM:
- case GOMP_MAP_FORCE_TOFROM:
- break;
- default:
- return false;
- }
- }
-
- return true;
-}
-
/* Scan a GIMPLE_OMP_TARGET. */
static void
@@ -2349,20 +2283,13 @@ scan_omp_target (gomp_target *stmt, omp_context *outer_ctx)
TYPE_NAME (ctx->record_type) = name;
TYPE_ARTIFICIAL (ctx->record_type) = 1;
- bool base_pointers_restrict = false;
if (offloaded)
{
create_omp_child_function (ctx, false);
gimple_omp_target_set_child_fn (stmt, ctx->cb.dst_fn);
-
- base_pointers_restrict = omp_target_base_pointers_restrict_p (clauses);
- if (base_pointers_restrict
- && dump_file && (dump_flags & TDF_DETAILS))
- fprintf (dump_file,
- "Base pointers in offloaded function are restrict\n");
}
- scan_sharing_clauses (clauses, ctx, base_pointers_restrict);
+ scan_sharing_clauses (clauses, ctx);
scan_omp (gimple_omp_body_ptr (stmt), ctx);
if (TYPE_FIELDS (ctx->record_type) == NULL)
diff --git a/gcc/tree-core.h b/gcc/tree-core.h
index 2bebb22a7e9..4a04e9e8b26 100644
--- a/gcc/tree-core.h
+++ b/gcc/tree-core.h
@@ -454,7 +454,13 @@ enum omp_clause_code {
/* OpenMP internal-only clause to specify grid dimensions of a gridified
kernel. */
- OMP_CLAUSE__GRIDDIM_
+ OMP_CLAUSE__GRIDDIM_,
+
+ /* OpenACC clause: if_present. */
+ OMP_CLAUSE_IF_PRESENT,
+
+ /* OpenACC clause: finalize. */
+ OMP_CLAUSE_FINALIZE
};
#undef DEFTREESTRUCT
diff --git a/gcc/tree-nested.c b/gcc/tree-nested.c
index b335d6b0afe..257ceae6f2d 100644
--- a/gcc/tree-nested.c
+++ b/gcc/tree-nested.c
@@ -1333,6 +1333,8 @@ convert_nonlocal_omp_clauses (tree *pclauses, struct walk_stmt_info *wi)
case OMP_CLAUSE_SEQ:
case OMP_CLAUSE_INDEPENDENT:
case OMP_CLAUSE_AUTO:
+ case OMP_CLAUSE_IF_PRESENT:
+ case OMP_CLAUSE_FINALIZE:
break;
/* The following clause belongs to the OpenACC cache directive, which
@@ -2022,6 +2024,8 @@ convert_local_omp_clauses (tree *pclauses, struct walk_stmt_info *wi)
case OMP_CLAUSE_SEQ:
case OMP_CLAUSE_INDEPENDENT:
case OMP_CLAUSE_AUTO:
+ case OMP_CLAUSE_IF_PRESENT:
+ case OMP_CLAUSE_FINALIZE:
break;
/* The following clause belongs to the OpenACC cache directive, which
diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c
index 63ec823c0ba..e65c40a41a3 100644
--- a/gcc/tree-pretty-print.c
+++ b/gcc/tree-pretty-print.c
@@ -1045,6 +1045,12 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
false);
pp_right_paren (pp);
break;
+ case OMP_CLAUSE_IF_PRESENT:
+ pp_string (pp, "if_present");
+ break;
+ case OMP_CLAUSE_FINALIZE:
+ pp_string (pp, "finalize");
+ break;
default:
/* Should never happen. */
diff --git a/gcc/tree.c b/gcc/tree.c
index 889d88c50b4..cf0b5f6dad3 100644
--- a/gcc/tree.c
+++ b/gcc/tree.c
@@ -343,6 +343,8 @@ unsigned const char omp_clause_num_ops[] =
1, /* OMP_CLAUSE_VECTOR_LENGTH */
3, /* OMP_CLAUSE_TILE */
2, /* OMP_CLAUSE__GRIDDIM_ */
+ 0, /* OMP_CLAUSE_IF_PRESENT */
+ 0, /* OMP_CLAUSE_FINALIZE */
};
const char * const omp_clause_code_name[] =
@@ -413,7 +415,9 @@ const char * const omp_clause_code_name[] =
"num_workers",
"vector_length",
"tile",
- "_griddim_"
+ "_griddim_",
+ "if_present",
+ "finalize",
};
@@ -11579,6 +11583,8 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data,
case OMP_CLAUSE_SEQ:
case OMP_CLAUSE_TILE:
case OMP_CLAUSE__SIMT_:
+ case OMP_CLAUSE_IF_PRESENT:
+ case OMP_CLAUSE_FINALIZE:
WALK_SUBTREE_TAIL (OMP_CLAUSE_CHAIN (*tp));
case OMP_CLAUSE_LASTPRIVATE:
--
2.17.1
^ permalink raw reply [flat|nested] 31+ messages in thread* Re: [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior - middle end
2018-06-19 17:01 ` [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior - middle end Cesar Philippidis
@ 2018-06-20 16:20 ` Jakub Jelinek
0 siblings, 0 replies; 31+ messages in thread
From: Jakub Jelinek @ 2018-06-20 16:20 UTC (permalink / raw)
To: Cesar Philippidis; +Cc: gcc-patches
On Tue, Jun 19, 2018 at 10:00:37AM -0700, Cesar Philippidis wrote:
> This patch implements the OpenACC 2.5 data clause semantics in the
> middle end.
>
> Is it OK for trunk?
>
> Cesar
> 2018-06-19 Chung-Lin Tang <cltang@codesourcery.com>
> Thomas Schwinge <thomas@codesourcery.com>
> Cesar Philippidis <cesar@codesourcery.com>
>
> gcc/c-family/
> * c-pragma.h (enum pragma_omp_clause): Add
> PRAGMA_OACC_CLAUSE_{FINALIZE,IF_PRESENT}. Remove
> PRAGMA_OACC_CLAUSE_PRESENT_OR_{COPY,COPYIN,COPYOUT,CREATE}.
>
> gcc/
> * gimplify.c (gimplify_scan_omp_clauses): Add support for
> OMP_CLAUSE_{IF_PRESENT,FINALIZE}.
> (gimplify_adjust_omp_clauses): Likewise.
> (gimplify_oacc_declare_1): Add support for GOMP_MAP_RELEASE, remove
> support for GOMP_MAP_FORCE_{ALLOC,TO,FROM,TOFROM}.
> (gimplify_omp_target_update): Update handling of acc update and
> enter/exit data.
> * omp-low.c (install_var_field): Remove unused parameter
> base_pointers_restrict.
> (scan_sharing_clauses): Remove base_pointers_restrict parameter.
> Update call to install_var_field. Handle OMP_CLAUSE_{IF_PRESENT,
> FINALIZE}
> (omp_target_base_pointers_restrict_p): Delete.
> (scan_omp_target): Update call to scan_sharing_clauses.
> * tree-core.h (enum omp_clause_code): Add OMP_CLAUSE_{IF_PRESENT,
> FINALIZE}.
> * tree-nested.c (convert_nonlocal_omp_clauses): Handle
> OMP_CLAUSE_{IF_PRESENT,FINALIZE}.
> (convert_local_omp_clauses): Likewise.
> * tree-pretty-print.c (dump_omp_clause): Likewise.
> * tree.c (omp_clause_num_ops): Add entries for OMP_CLAUSE_{IF_PRESENT,
> FINALIZE}.
> (omp_clause_code_name): Likewise.
Ok.
Jakub
^ permalink raw reply [flat|nested] 31+ messages in thread
* Re: [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior - compiler tests
2018-06-19 16:56 ` Cesar Philippidis
` (4 preceding siblings ...)
2018-06-19 17:01 ` [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior - middle end Cesar Philippidis
@ 2018-06-19 17:02 ` Cesar Philippidis
2018-06-19 17:03 ` [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior - runtime tests Cesar Philippidis
6 siblings, 0 replies; 31+ messages in thread
From: Cesar Philippidis @ 2018-06-19 17:02 UTC (permalink / raw)
To: gcc-patches, Jakub Jelinek
[-- Attachment #1: Type: text/plain, Size: 131 bytes --]
This patch updates the existing OpenACC compiler tests with the new
OpenACC 2.5 data clause semantics.
Is it OK for trunk?
Cesar
[-- Attachment #2: 0001-compiler-tests.patch --]
[-- Type: text/x-patch, Size: 49392 bytes --]
2018-06-19 Chung-Lin Tang <cltang@codesourcery.com>
Thomas Schwinge <thomas@codesourcery.com>
Cesar Philippidis <cesar@codesourcery.com>
gcc/testsuite/
* c-c++-common/goacc/declare-1.c: Update test case to utilize OpenACC
2.5 data clause semantics.
* c-c++-common/goacc/declare-2.c: Likewise.
* c-c++-common/goacc/default-4.c: Likewise.
* c-c++-common/goacc/finalize-1.c: New test.
* c-c++-common/goacc/kernels-alias-2.c: Update test case to utilize
OpenACC 2.5 data clause semantics.
* c-c++-common/goacc/kernels-alias.c: Likewise.
* c-c++-common/goacc/routine-5.c: Likewise.
* c-c++-common/goacc/update-if_present-1.c: New test.
* c-c++-common/goacc/update-if_present-2.c: New test.
* g++.dg/goacc/template.C: Update test case to utilize OpenACC
2.5 data clause semantics.
* gfortran.dg/goacc/combined-directives.f90: Likewise.
* gfortran.dg/goacc/data-tree.f95: Likewise.
* gfortran.dg/goacc/declare-2.f95: Likewise.
* gfortran.dg/goacc/default-4.f: Likewise.
* gfortran.dg/goacc/enter-exit-data.f95: Likewise.
* gfortran.dg/goacc/finalize-1.f: New test.
* gfortran.dg/goacc/kernels-alias-2.f95: Update test case to utilize
OpenACC 2.5 data clause semantics.
* gfortran.dg/goacc/kernels-alias.f95: Likewise.
* gfortran.dg/goacc/kernels-tree.f95: Likewise.
* gfortran.dg/goacc/nested-function-1.f90: Likewise.
* gfortran.dg/goacc/parallel-tree.f95: Likewise.
* gfortran.dg/goacc/reduction-promotions.f90: Likewise.
* gfortran.dg/goacc/update-if_present-1.f90: New test.
* gfortran.dg/goacc/update-if_present-2.f90: New test.
From 87819f5846a3d4aae3983740e09a0ac4e1eb866f Mon Sep 17 00:00:00 2001
From: Cesar Philippidis <cesar@codesourcery.com>
Date: Tue, 19 Jun 2018 09:28:54 -0700
Subject: [PATCH 1/7] compiler tests
---
gcc/testsuite/c-c++-common/goacc/declare-1.c | 12 ++
gcc/testsuite/c-c++-common/goacc/declare-2.c | 18 +--
gcc/testsuite/c-c++-common/goacc/default-4.c | 6 +-
gcc/testsuite/c-c++-common/goacc/finalize-1.c | 28 ++++
.../c-c++-common/goacc/kernels-alias-2.c | 10 +-
.../c-c++-common/goacc/kernels-alias.c | 10 +-
gcc/testsuite/c-c++-common/goacc/routine-5.c | 150 +++++++-----------
.../c-c++-common/goacc/update-if_present-1.c | 28 ++++
.../c-c++-common/goacc/update-if_present-2.c | 42 +++++
gcc/testsuite/g++.dg/goacc/template.C | 13 +-
.../gfortran.dg/goacc/combined-directives.f90 | 2 +-
gcc/testsuite/gfortran.dg/goacc/data-tree.f95 | 8 +-
gcc/testsuite/gfortran.dg/goacc/declare-2.f95 | 6 +-
gcc/testsuite/gfortran.dg/goacc/default-4.f | 6 +-
.../gfortran.dg/goacc/enter-exit-data.f95 | 3 +
gcc/testsuite/gfortran.dg/goacc/finalize-1.f | 27 ++++
.../gfortran.dg/goacc/kernels-alias-2.f95 | 10 +-
.../gfortran.dg/goacc/kernels-alias.f95 | 10 +-
.../gfortran.dg/goacc/kernels-tree.f95 | 8 +-
.../gfortran.dg/goacc/nested-function-1.f90 | 8 +
.../gfortran.dg/goacc/parallel-tree.f95 | 12 +-
.../goacc/reduction-promotions.f90 | 6 +-
.../gfortran.dg/goacc/update-if_present-1.f90 | 27 ++++
.../gfortran.dg/goacc/update-if_present-2.f90 | 52 ++++++
24 files changed, 345 insertions(+), 157 deletions(-)
create mode 100644 gcc/testsuite/c-c++-common/goacc/finalize-1.c
create mode 100644 gcc/testsuite/c-c++-common/goacc/update-if_present-1.c
create mode 100644 gcc/testsuite/c-c++-common/goacc/update-if_present-2.c
create mode 100644 gcc/testsuite/gfortran.dg/goacc/finalize-1.f
create mode 100644 gcc/testsuite/gfortran.dg/goacc/update-if_present-1.f90
create mode 100644 gcc/testsuite/gfortran.dg/goacc/update-if_present-2.f90
diff --git a/gcc/testsuite/c-c++-common/goacc/declare-1.c b/gcc/testsuite/c-c++-common/goacc/declare-1.c
index b036c636166..35b1ccd367b 100644
--- a/gcc/testsuite/c-c++-common/goacc/declare-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/declare-1.c
@@ -19,6 +19,12 @@ int v4;
int v5, v6, v7, v8;
#pragma acc declare create(v5, v6) copyin(v7, v8)
+int v9;
+#pragma acc declare present_or_copyin(v9)
+
+int v10;
+#pragma acc declare present_or_create(v10)
+
void
f (void)
{
@@ -49,6 +55,12 @@ f (void)
extern int ve4;
#pragma acc declare link(ve4)
+ extern int ve5;
+#pragma acc declare present_or_copyin(ve5)
+
+ extern int ve6;
+#pragma acc declare present_or_create(ve6)
+
int va5;
#pragma acc declare copy(va5)
diff --git a/gcc/testsuite/c-c++-common/goacc/declare-2.c b/gcc/testsuite/c-c++-common/goacc/declare-2.c
index e41a0f59537..33b82459bfc 100644
--- a/gcc/testsuite/c-c++-common/goacc/declare-2.c
+++ b/gcc/testsuite/c-c++-common/goacc/declare-2.c
@@ -29,13 +29,7 @@ int v6;
#pragma acc declare present_or_copy(v6) /* { dg-error "at file scope" } */
int v7;
-#pragma acc declare present_or_copyin(v7) /* { dg-error "at file scope" } */
-
-int v8;
-#pragma acc declare present_or_copyout(v8) /* { dg-error "at file scope" } */
-
-int v9;
-#pragma acc declare present_or_create(v9) /* { dg-error "at file scope" } */
+#pragma acc declare present_or_copyout(v7) /* { dg-error "at file scope" } */
int va10;
#pragma acc declare create (va10)
@@ -67,13 +61,7 @@ f (void)
#pragma acc declare present_or_copy(ve3) /* { dg-error "invalid use of" } */
extern int ve4;
-#pragma acc declare present_or_copyin(ve4) /* { dg-error "invalid use of" } */
-
- extern int ve5;
-#pragma acc declare present_or_copyout(ve5) /* { dg-error "invalid use of" } */
-
- extern int ve6;
-#pragma acc declare present_or_create(ve6) /* { dg-error "invalid use of" } */
+#pragma acc declare present_or_copyout(ve4) /* { dg-error "invalid use of" } */
-#pragma acc declare present (v9) /* { dg-error "invalid use of" } */
+#pragma acc declare present (v2) /* { dg-error "invalid use of" } */
}
diff --git a/gcc/testsuite/c-c++-common/goacc/default-4.c b/gcc/testsuite/c-c++-common/goacc/default-4.c
index dfa79bbbe6e..867175d4847 100644
--- a/gcc/testsuite/c-c++-common/goacc/default-4.c
+++ b/gcc/testsuite/c-c++-common/goacc/default-4.c
@@ -8,7 +8,7 @@ void f1 ()
float f1_b[2];
#pragma acc data copyin (f1_a) copyout (f1_b)
- /* { dg-final { scan-tree-dump-times "omp target oacc_data map\\(force_from:f1_b \[^\\)\]+\\) map\\(force_to:f1_a" 1 "gimple" } } */
+ /* { dg-final { scan-tree-dump-times "omp target oacc_data map\\(from:f1_b \[^\\)\]+\\) map\\(to:f1_a" 1 "gimple" } } */
{
#pragma acc kernels
/* { dg-final { scan-tree-dump-times "omp target oacc_kernels map\\(tofrom:f1_b \[^\\)\]+\\) map\\(tofrom:f1_a" 1 "gimple" } } */
@@ -29,7 +29,7 @@ void f2 ()
float f2_b[2];
#pragma acc data copyin (f2_a) copyout (f2_b)
- /* { dg-final { scan-tree-dump-times "omp target oacc_data map\\(force_from:f2_b \[^\\)\]+\\) map\\(force_to:f2_a" 1 "gimple" } } */
+ /* { dg-final { scan-tree-dump-times "omp target oacc_data map\\(from:f2_b \[^\\)\]+\\) map\\(to:f2_a" 1 "gimple" } } */
{
#pragma acc kernels default (none)
/* { dg-final { scan-tree-dump-times "omp target oacc_kernels default\\(none\\) map\\(tofrom:f2_b \[^\\)\]+\\) map\\(tofrom:f2_a" 1 "gimple" } } */
@@ -50,7 +50,7 @@ void f3 ()
float f3_b[2];
#pragma acc data copyin (f3_a) copyout (f3_b)
- /* { dg-final { scan-tree-dump-times "omp target oacc_data map\\(force_from:f3_b \[^\\)\]+\\) map\\(force_to:f3_a" 1 "gimple" } } */
+ /* { dg-final { scan-tree-dump-times "omp target oacc_data map\\(from:f3_b \[^\\)\]+\\) map\\(to:f3_a" 1 "gimple" } } */
{
#pragma acc kernels default (present)
/* { dg-final { scan-tree-dump-times "omp target oacc_kernels default\\(present\\) map\\(tofrom:f3_b \[^\\)\]+\\) map\\(tofrom:f3_a" 1 "gimple" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/finalize-1.c b/gcc/testsuite/c-c++-common/goacc/finalize-1.c
new file mode 100644
index 00000000000..94820290b94
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/finalize-1.c
@@ -0,0 +1,28 @@
+/* Test valid usage and processing of the finalize clause. */
+
+/* { dg-additional-options "-fdump-tree-original -fdump-tree-gimple" } */
+
+extern int del_r;
+extern float del_f[3];
+extern double cpo_r[8];
+extern long cpo_f;
+
+void f ()
+{
+#pragma acc exit data delete (del_r)
+/* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:del_r\\);$" 1 "original" } }
+ { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(release:del_r \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } } */
+
+#pragma acc exit data finalize delete (del_f)
+/* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:del_f\\) finalize;$" 1 "original" } }
+ { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:del_f \\\[len: \[0-9\]+\\\]\\) finalize$" 1 "gimple" } } */
+
+#pragma acc exit data copyout (cpo_r)
+/* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:cpo_r\\);$" 1 "original" } }
+ { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(from:cpo_r \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } } */
+
+#pragma acc exit data copyout (cpo_f) finalize
+/* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data finalize map\\(from:cpo_f\\);$" 1 "original" } }
+ { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data finalize map\\(force_from:cpo_f \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } } */
+}
+
diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-alias-2.c b/gcc/testsuite/c-c++-common/goacc/kernels-alias-2.c
index d437c47779d..7576a6484f1 100644
--- a/gcc/testsuite/c-c++-common/goacc/kernels-alias-2.c
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-alias-2.c
@@ -18,10 +18,12 @@ foo (void)
}
}
+/* The xfails occur due to the OpenACC 2.5 data semantics. */
+
/* { dg-final { scan-tree-dump-times "clique 1 base 1" 4 "ealias" } } */
-/* { dg-final { scan-tree-dump-times "clique 1 base 2" 1 "ealias" } } */
-/* { dg-final { scan-tree-dump-times "clique 1 base 3" 1 "ealias" } } */
-/* { dg-final { scan-tree-dump-times "clique 1 base 4" 1 "ealias" } } */
-/* { dg-final { scan-tree-dump-times "clique 1 base 5" 1 "ealias" } } */
+/* { dg-final { scan-tree-dump-times "clique 1 base 2" 1 "ealias" { xfail *-*-* } } } */
+/* { dg-final { scan-tree-dump-times "clique 1 base 3" 1 "ealias" { xfail *-*-* } } } */
+/* { dg-final { scan-tree-dump-times "clique 1 base 4" 1 "ealias" { xfail *-*-* } } } */
+/* { dg-final { scan-tree-dump-times "clique 1 base 5" 1 "ealias" { xfail *-*-* } } } */
/* { dg-final { scan-tree-dump-times "(?n)clique .* base .*" 8 "ealias" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-alias.c b/gcc/testsuite/c-c++-common/goacc/kernels-alias.c
index 25821ab2aea..e8ff018d512 100644
--- a/gcc/testsuite/c-c++-common/goacc/kernels-alias.c
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-alias.c
@@ -20,10 +20,12 @@ foo (void)
}
}
+/* The xfails occur due to the OpenACC 2.5 data semantics. */
+
/* { dg-final { scan-tree-dump-times "clique 1 base 1" 4 "ealias" } } */
-/* { dg-final { scan-tree-dump-times "clique 1 base 2" 1 "ealias" } } */
-/* { dg-final { scan-tree-dump-times "clique 1 base 3" 1 "ealias" } } */
-/* { dg-final { scan-tree-dump-times "clique 1 base 4" 1 "ealias" } } */
-/* { dg-final { scan-tree-dump-times "clique 1 base 5" 1 "ealias" } } */
+/* { dg-final { scan-tree-dump-times "clique 1 base 2" 1 "ealias" { xfail *-*-* } } } */
+/* { dg-final { scan-tree-dump-times "clique 1 base 3" 1 "ealias" { xfail *-*-* } } } */
+/* { dg-final { scan-tree-dump-times "clique 1 base 4" 1 "ealias" { xfail *-*-* } } } */
+/* { dg-final { scan-tree-dump-times "clique 1 base 5" 1 "ealias" { xfail *-*-* } } } */
/* { dg-final { scan-tree-dump-times "(?n)clique .* base .*" 8 "ealias" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/routine-5.c b/gcc/testsuite/c-c++-common/goacc/routine-5.c
index b967a7447bd..b759db3292d 100644
--- a/gcc/testsuite/c-c++-common/goacc/routine-5.c
+++ b/gcc/testsuite/c-c++-common/goacc/routine-5.c
@@ -4,11 +4,11 @@
struct PC
{
-#pragma acc routine /* { dg-error ".#pragma acc routine. must be at file scope" } */
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. must be at file scope" } */
};
void PC1( /* { dg-bogus "variable or field .PC1. declared void" "TODO" { xfail c++ } } */
-#pragma acc routine
+#pragma acc routine seq
/* { dg-error ".#pragma acc routine. must be at file scope" "" { target c } .-1 }
{ dg-error ".#pragma. is not allowed here" "" { target c++ } .-2 } */
) /* { dg-bogus "expected declaration specifiers or .\\.\\.\\.. before .\\). token" "TODO" { xfail c } } */
@@ -18,26 +18,26 @@ void PC1( /* { dg-bogus "variable or field .PC1. declared void" "TODO" { xfail c
void PC2()
{
if (0)
-#pragma acc routine /* { dg-error ".#pragma acc routine. must be at file scope" } */
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. must be at file scope" } */
;
}
void PC3()
{
-#pragma acc routine /* { dg-error ".#pragma acc routine. must be at file scope" } */
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. must be at file scope" } */
}
/* "( name )" syntax. */
#pragma acc routine ( /* { dg-error "expected (function name|unqualified-id) before end of line" } */
-#pragma acc routine () /* { dg-error "expected (function name|unqualified-id) before .\\). token" } */
-#pragma acc routine (+) /* { dg-error "expected (function name|unqualified-id) before .\\+. token" } */
-#pragma acc routine (?) /* { dg-error "expected (function name|unqualified-id) before .\\?. token" } */
-#pragma acc routine (:) /* { dg-error "expected (function name|unqualified-id) before .:. token" } */
-#pragma acc routine (4) /* { dg-error "expected (function name|unqualified-id) before numeric constant" } */
+#pragma acc routine () seq /* { dg-error "expected (function name|unqualified-id) before .\\). token" } */
+#pragma acc routine (+) seq /* { dg-error "expected (function name|unqualified-id) before .\\+. token" } */
+#pragma acc routine (?) seq /* { dg-error "expected (function name|unqualified-id) before .\\?. token" } */
+#pragma acc routine (:) seq /* { dg-error "expected (function name|unqualified-id) before .:. token" } */
+#pragma acc routine (4) seq /* { dg-error "expected (function name|unqualified-id) before numeric constant" } */
#pragma acc routine ('4') /* { dg-error "expected (function name|unqualified-id) before .4." } */
-#pragma acc routine ("4") /* { dg-error "expected (function name|unqualified-id) before string constant" } */
+#pragma acc routine ("4") seq /* { dg-error "expected (function name|unqualified-id) before string constant" } */
extern void R1(void);
extern void R2(void);
#pragma acc routine (R1, R2, R3) worker /* { dg-error "expected .\\). before .,. token" } */
@@ -49,84 +49,84 @@ extern void R2(void);
/* "#pragma acc routine" not immediately followed by (a single) function
declaration or definition. */
-#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
int a;
-#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by a single function declaration or definition" } */
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by a single function declaration or definition" } */
void fn1 (void), fn1b (void);
-#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
int b, fn2 (void);
-#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
int b_, fn2_ (void), B_;
-#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by a single function declaration or definition" } */
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by a single function declaration or definition" } */
int fn3 (void), b2;
-#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
typedef struct c c;
-#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
struct d {} d;
-#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
-#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by a single function declaration or definition" } */
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by a single function declaration or definition" } */
void fn1_2 (void), fn1b_2 (void);
-#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
int b_2, fn2_2 (void);
-#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
-#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
int b_2_, fn2_2_ (void), B_2_;
-#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
-#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by a single function declaration or definition" } */
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by a single function declaration or definition" } */
int fn3_2 (void), b2_2;
-#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
-#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
typedef struct c_2 c_2;
-#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
-#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
struct d_2 {} d_2;
-#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
-#pragma acc routine
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
+#pragma acc routine seq
int fn4 (void);
int fn5a (void);
int fn5b (void);
-#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
-#pragma acc routine (fn5a)
-#pragma acc routine (fn5b)
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
+#pragma acc routine (fn5a) seq
+#pragma acc routine (fn5b) seq
int fn5 (void);
-#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
-#pragma acc routine (fn6a) /* { dg-error ".fn6a. has not been declared" } */
-#pragma acc routine (fn6b) /* { dg-error ".fn6b. has not been declared" } */
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
+#pragma acc routine (fn6a) seq /* { dg-error ".fn6a. has not been declared" } */
+#pragma acc routine (fn6b) seq /* { dg-error ".fn6b. has not been declared" } */
int fn6 (void);
#ifdef __cplusplus
-#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" "" { target c++ } } */
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" "" { target c++ } } */
namespace f {}
namespace g {}
-#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" "" { target c++ } } */
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" "" { target c++ } } */
using namespace g;
-#pragma acc routine (g) /* { dg-error ".g. does not refer to a function" "" { target c++ } } */
+#pragma acc routine (g) seq /* { dg-error ".g. does not refer to a function" "" { target c++ } } */
#endif /* __cplusplus */
-#pragma acc routine (a) /* { dg-error ".a. does not refer to a function" } */
+#pragma acc routine (a) seq /* { dg-error ".a. does not refer to a function" } */
-#pragma acc routine (c) /* { dg-error ".c. does not refer to a function" } */
+#pragma acc routine (c) seq /* { dg-error ".c. does not refer to a function" } */
/* Static assert. */
@@ -143,66 +143,24 @@ static_assert(0, ""); /* { dg-error "static assertion failed" "" { target c++11
#endif
void f_static_assert();
/* Check that we already recognized "f_static_assert" as an OpenACC routine. */
-#pragma acc routine (f_static_assert) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*f_static_assert" "TODO" { xfail *-*-* } } */
+#pragma acc routine (f_static_assert) seq /* { dg-error ".#pragma acc routine. already applied to .\[void \]*f_static_assert" "TODO" { xfail *-*-* } } */
/* __extension__ usage. */
-#pragma acc routine
+#pragma acc routine seq
__extension__ extern void ex1();
#pragma acc routine (ex1) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*ex1" } */
-#pragma acc routine
+#pragma acc routine seq
__extension__ __extension__ __extension__ __extension__ __extension__ void ex2()
{
}
#pragma acc routine (ex2) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*ex2" } */
-#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
+#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
__extension__ int ex3;
-#pragma acc routine (ex3) /* { dg-error ".ex3. does not refer to a function" } */
-
-
-/* "#pragma acc routine" already applied. */
-
-extern void fungsi_1();
-#pragma acc routine(fungsi_1) gang
-#pragma acc routine(fungsi_1) gang /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_1" } */
-#pragma acc routine(fungsi_1) worker /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_1" } */
-#pragma acc routine(fungsi_1) vector /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_1" } */
-
-#pragma acc routine seq
-extern void fungsi_2();
-#pragma acc routine(fungsi_2) seq /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_2." } */
-#pragma acc routine(fungsi_2) worker /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_2." } */
-#pragma acc routine(fungsi_2) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_2." } */
-
-#pragma acc routine vector
-extern void fungsi_3();
-#pragma acc routine vector /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_3." } */
-void fungsi_3()
-{
-}
-
-extern void fungsi_4();
-#pragma acc routine (fungsi_4) worker
-#pragma acc routine gang /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_4." } */
-void fungsi_4()
-{
-}
-
-#pragma acc routine gang
-void fungsi_5()
-{
-}
-#pragma acc routine (fungsi_5) worker /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_5." } */
-
-#pragma acc routine seq
-void fungsi_6()
-{
-}
-#pragma acc routine seq /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_6." } */
-extern void fungsi_6();
+#pragma acc routine (ex3) seq /* { dg-error ".ex3. does not refer to a function" } */
/* "#pragma acc routine" must be applied before. */
@@ -214,11 +172,11 @@ void Foo ()
Bar ();
}
-#pragma acc routine (Bar) // { dg-error ".#pragma acc routine. must be applied before use" }
+#pragma acc routine (Bar) seq // { dg-error ".#pragma acc routine. must be applied before use" }
#pragma acc routine (Foo) gang // { dg-error ".#pragma acc routine. must be applied before definition" }
-#pragma acc routine (Baz) // { dg-error "not been declared" }
+#pragma acc routine (Baz) seq // { dg-error "not been declared" }
/* OpenACC declare. */
@@ -227,7 +185,7 @@ int vb1; /* { dg-error "directive for use" } */
extern int vb2; /* { dg-error "directive for use" } */
static int vb3; /* { dg-error "directive for use" } */
-#pragma acc routine
+#pragma acc routine seq
int
func1 (int a)
{
@@ -238,7 +196,7 @@ func1 (int a)
return vb3;
}
-#pragma acc routine
+#pragma acc routine seq
int
func2 (int a)
{
@@ -256,7 +214,7 @@ extern int vb6; /* { dg-error "clause used in" } */
static int vb7; /* { dg-error "clause used in" } */
#pragma acc declare link (vb7)
-#pragma acc routine
+#pragma acc routine seq
int
func3 (int a)
{
@@ -273,7 +231,7 @@ extern int vb9;
static int vb10;
#pragma acc declare create (vb10)
-#pragma acc routine
+#pragma acc routine seq
int
func4 (int a)
{
@@ -291,7 +249,7 @@ extern int vb12;
extern int vb13;
#pragma acc declare device_resident (vb13)
-#pragma acc routine
+#pragma acc routine seq
int
func5 (int a)
{
@@ -302,7 +260,7 @@ func5 (int a)
return vb13;
}
-#pragma acc routine
+#pragma acc routine seq
int
func6 (int a)
{
diff --git a/gcc/testsuite/c-c++-common/goacc/update-if_present-1.c b/gcc/testsuite/c-c++-common/goacc/update-if_present-1.c
new file mode 100644
index 00000000000..c34a0e48065
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/update-if_present-1.c
@@ -0,0 +1,28 @@
+/* Test valid usages of the if_present clause. */
+
+/* { dg-additional-options "-fdump-tree-omplower" } */
+
+void
+t ()
+{
+ int a, b, c[10];
+
+#pragma acc update self(a) if_present
+#pragma acc update device(b) async if_present
+#pragma acc update host(c[1:3]) wait(4) if_present
+#pragma acc update self(c) device(b) host (a) async(10) if (a == 5) if_present
+
+#pragma acc update self(a)
+#pragma acc update device(b) async
+#pragma acc update host(c[1:3]) wait(4)
+#pragma acc update self(c) device(b) host (a) async(10) if (a == 5)
+}
+
+/* { dg-final { scan-tree-dump-times "omp target oacc_update if_present map.from:a .len: 4.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "omp target oacc_update if_present async.-1. map.to:b .len: 4.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "omp target oacc_update if_present wait.4. map.from:c.1. .len: 12.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "omp target oacc_update if_present if.... async.10. map.from:a .len: 4.. map.to:b .len: 4.. map.from:c .len: 40.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "omp target oacc_update map.force_from:a .len: 4.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "omp target oacc_update async.-1. map.force_to:b .len: 4.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "omp target oacc_update wait.4. map.force_from:c.1. .len: 12.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "omp target oacc_update if.... async.10. map.force_from:a .len: 4.. map.force_to:b .len: 4.. map.force_from:c .len: 40.." 1 "omplower" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/update-if_present-2.c b/gcc/testsuite/c-c++-common/goacc/update-if_present-2.c
new file mode 100644
index 00000000000..974f1b8c427
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/update-if_present-2.c
@@ -0,0 +1,42 @@
+/* Test invalid usages of the if_present clause. */
+
+#pragma acc routine gang if_present /* { dg-error "'if_present' is not valid" } */
+void
+t1 ()
+{
+ int a, b, c[10];
+
+#pragma acc enter data copyin(a) if_present /* { dg-error "'if_present' is not valid" } */
+#pragma acc exit data copyout(a) if_present /* { dg-error "'if_present' is not valid" } */
+
+#pragma acc data copy(a) if_present /* { dg-error "'if_present' is not valid" } */
+ {
+ }
+
+#pragma acc declare create(c) if_present /* { dg-error "'if_present' is not valid" } */
+
+#pragma acc init if_present
+#pragma acc shutdown if_present
+}
+
+void
+t2 ()
+{
+ int a, b, c[10];
+
+#pragma acc update self(a)
+#pragma acc parallel
+#pragma acc loop if_present /* { dg-error "'if_present' is not valid" } */
+ for (b = 1; b < 10; b++)
+ ;
+#pragma acc end parallel
+
+#pragma acc kernels loop if_present /* { dg-error "'if_present' is not valid" } */
+ for (b = 1; b < 10; b++)
+ ;
+
+#pragma acc parallel loop if_present /* { dg-error "'if_present' is not valid" } */
+ for (b = 1; b < 10; b++)
+ ;
+}
+
diff --git a/gcc/testsuite/g++.dg/goacc/template.C b/gcc/testsuite/g++.dg/goacc/template.C
index 852f42f2b42..dae92b08987 100644
--- a/gcc/testsuite/g++.dg/goacc/template.C
+++ b/gcc/testsuite/g++.dg/goacc/template.C
@@ -1,4 +1,4 @@
-#pragma acc routine
+#pragma acc routine seq
template <typename T> T
accDouble(int val)
{
@@ -31,7 +31,7 @@ oacc_parallel_copy (T a)
#pragma acc parallel num_gangs (a) if (1)
{
-#pragma acc loop independent collapse (2) gang
+#pragma acc loop independent collapse (2)
for (int i = 0; i < a; i++)
for (int j = 0; j < 5; j++)
b = a;
@@ -86,6 +86,8 @@ oacc_parallel_copy (T a)
#pragma acc update self (b)
#pragma acc update device (b)
#pragma acc exit data delete (b)
+#pragma acc exit data finalize copyout (b)
+#pragma acc exit data delete (b) finalize
return b;
}
@@ -133,6 +135,13 @@ oacc_kernels_copy (T a)
b = a;
}
+#pragma acc update host (b)
+#pragma acc update self (b)
+#pragma acc update device (b)
+#pragma acc exit data delete (b)
+#pragma acc exit data finalize copyout (b)
+#pragma acc exit data delete (b) finalize
+
return b;
}
diff --git a/gcc/testsuite/gfortran.dg/goacc/combined-directives.f90 b/gcc/testsuite/gfortran.dg/goacc/combined-directives.f90
index 42a447ad06b..956349204f4 100644
--- a/gcc/testsuite/gfortran.dg/goacc/combined-directives.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/combined-directives.f90
@@ -146,5 +146,5 @@ end subroutine test
! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. tile.2, 3" 2 "gimple" } }
! { dg-final { scan-tree-dump-times "acc loop private.i. independent" 2 "gimple" } }
! { dg-final { scan-tree-dump-times "private.z" 2 "gimple" } }
-! { dg-final { scan-tree-dump-times "omp target oacc_\[^ \]+ map.force_tofrom:y" 2 "gimple" } }
+! { dg-final { scan-tree-dump-times "omp target oacc_\[^ \]+ map.tofrom:y" 2 "gimple" } }
! { dg-final { scan-tree-dump-times "acc loop private.i. reduction..:y." 2 "gimple" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/data-tree.f95 b/gcc/testsuite/gfortran.dg/goacc/data-tree.f95
index 44efc8a670b..f16d62cce69 100644
--- a/gcc/testsuite/gfortran.dg/goacc/data-tree.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/data-tree.f95
@@ -15,10 +15,10 @@ end program test
! { dg-final { scan-tree-dump-times "pragma acc data" 1 "original" } }
! { dg-final { scan-tree-dump-times "if" 1 "original" } }
-! { dg-final { scan-tree-dump-times "map\\(force_tofrom:i\\)" 1 "original" } }
-! { dg-final { scan-tree-dump-times "map\\(force_to:j\\)" 1 "original" } }
-! { dg-final { scan-tree-dump-times "map\\(force_from:k\\)" 1 "original" } }
-! { dg-final { scan-tree-dump-times "map\\(force_alloc:m\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "map\\(tofrom:i\\)" 1 "original" } }
+! { 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\\(force_present:o\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "map\\(tofrom:p\\)" 1 "original" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/declare-2.f95 b/gcc/testsuite/gfortran.dg/goacc/declare-2.f95
index aa1704f77d0..7aa3dab4707 100644
--- a/gcc/testsuite/gfortran.dg/goacc/declare-2.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/declare-2.f95
@@ -11,11 +11,11 @@ subroutine asubr (b)
!$acc declare copyout (b) ! { dg-error "Invalid clause in module" }
!$acc declare present (b) ! { dg-error "Invalid clause in module" }
!$acc declare present_or_copy (b) ! { dg-error "Invalid clause in module" }
- !$acc declare present_or_copyin (b) ! { dg-error "Invalid clause in module" }
+ !$acc declare present_or_copyin (b) ! { dg-error "present on multiple" }
!$acc declare present_or_copyout (b) ! { dg-error "Invalid clause in module" }
- !$acc declare present_or_create (b) ! { dg-error "Invalid clause in module" }
+ !$acc declare present_or_create (b) ! { dg-error "present on multiple" }
!$acc declare deviceptr (b) ! { dg-error "Invalid clause in module" }
- !$acc declare create (b) copyin (b) ! { dg-error "present on multiple clauses" }
+ !$acc declare create (b) copyin (b) ! { dg-error "present on multiple" }
end subroutine
diff --git a/gcc/testsuite/gfortran.dg/goacc/default-4.f b/gcc/testsuite/gfortran.dg/goacc/default-4.f
index 77291f43eff..30f411f70ab 100644
--- a/gcc/testsuite/gfortran.dg/goacc/default-4.f
+++ b/gcc/testsuite/gfortran.dg/goacc/default-4.f
@@ -8,7 +8,7 @@
REAL, DIMENSION (2) :: F1_B
!$ACC DATA COPYIN (F1_A) COPYOUT (F1_B)
-! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(force_to:f1_a \[^\\)\]+\\) map\\(force_from:f1_b" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(to:f1_a \[^\\)\]+\\) map\\(from:f1_b" 1 "gimple" } }
!$ACC KERNELS
! { dg-final { scan-tree-dump-times "omp target oacc_kernels map\\(tofrom:f1_b \[^\\)\]+\\) map\\(tofrom:f1_a" 1 "gimple" } }
F1_B(1) = F1_A;
@@ -26,7 +26,7 @@
REAL, DIMENSION (2) :: F2_B
!$ACC DATA COPYIN (F2_A) COPYOUT (F2_B)
-! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(force_to:f2_a \[^\\)\]+\\) map\\(force_from:f2_b" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(to:f2_a \[^\\)\]+\\) map\\(from:f2_b" 1 "gimple" } }
!$ACC KERNELS DEFAULT (NONE)
! { dg-final { scan-tree-dump-times "omp target oacc_kernels default\\(none\\) map\\(tofrom:f2_b \[^\\)\]+\\) map\\(tofrom:f2_a" 1 "gimple" } }
F2_B(1) = F2_A;
@@ -44,7 +44,7 @@
REAL, DIMENSION (2) :: F3_B
!$ACC DATA COPYIN (F3_A) COPYOUT (F3_B)
-! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(force_to:f3_a \[^\\)\]+\\) map\\(force_from:f3_b" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(to:f3_a \[^\\)\]+\\) map\\(from:f3_b" 1 "gimple" } }
!$ACC KERNELS DEFAULT (PRESENT)
! { dg-final { scan-tree-dump-times "omp target oacc_kernels default\\(present\\) map\\(tofrom:f3_b \[^\\)\]+\\) map\\(tofrom:f3_a" 1 "gimple" } }
F3_B(1) = F3_A;
diff --git a/gcc/testsuite/gfortran.dg/goacc/enter-exit-data.f95 b/gcc/testsuite/gfortran.dg/goacc/enter-exit-data.f95
index 8f1715e659d..805459c1bb0 100644
--- a/gcc/testsuite/gfortran.dg/goacc/enter-exit-data.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/enter-exit-data.f95
@@ -84,5 +84,8 @@ contains
!$acc exit data delete (tip) ! { dg-error "POINTER" }
!$acc exit data delete (tia) ! { dg-error "ALLOCATABLE" }
!$acc exit data copyout (i) delete (i) ! { dg-error "multiple clauses" }
+ !$acc exit data finalize
+ !$acc exit data finalize copyout (i)
+ !$acc exit data finalize delete (i)
end subroutine foo
end module test
diff --git a/gcc/testsuite/gfortran.dg/goacc/finalize-1.f b/gcc/testsuite/gfortran.dg/goacc/finalize-1.f
new file mode 100644
index 00000000000..5c7a921a2e3
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/finalize-1.f
@@ -0,0 +1,27 @@
+! Test valid usage and processing of the finalize clause.
+
+! { dg-additional-options "-fdump-tree-original -fdump-tree-gimple" }
+
+ SUBROUTINE f
+ IMPLICIT NONE
+ INTEGER :: del_r
+ REAL, DIMENSION (3) :: del_f
+ DOUBLE PRECISION, DIMENSION (8) :: cpo_r
+ LOGICAL :: cpo_f
+
+!$ACC EXIT DATA DELETE (del_r)
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:del_r\\);$" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(release:del_r \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } }
+
+!$ACC EXIT DATA FINALIZE DELETE (del_f)
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:del_f\\) finalize;$" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:del_f \\\[len: \[0-9\]+\\\]\\) finalize$" 1 "gimple" } }
+
+!$ACC EXIT DATA COPYOUT (cpo_r)
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:cpo_r\\);$" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(from:cpo_r \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } }
+
+!$ACC EXIT DATA COPYOUT (cpo_f) FINALIZE
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:cpo_f\\) finalize;$" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_from:cpo_f \\\[len: \[0-9\]+\\\]\\) finalize$" 1 "gimple" } }
+ END SUBROUTINE f
diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-alias-2.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-alias-2.f95
index 7e348dde2bd..6a9f241a596 100644
--- a/gcc/testsuite/gfortran.dg/goacc/kernels-alias-2.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-alias-2.f95
@@ -15,9 +15,11 @@ program main
end program main
+! The xfails occur in light of the new OpenACC data semantics.
+
! { dg-final { scan-tree-dump-times "clique 1 base 1" 4 "ealias" } }
-! { dg-final { scan-tree-dump-times "clique 1 base 2" 1 "ealias" } }
-! { dg-final { scan-tree-dump-times "clique 1 base 3" 1 "ealias" } }
-! { dg-final { scan-tree-dump-times "clique 1 base 4" 1 "ealias" } }
-! { dg-final { scan-tree-dump-times "clique 1 base 5" 1 "ealias" } }
+! { dg-final { scan-tree-dump-times "clique 1 base 2" 1 "ealias" { xfail *-*-* } } }
+! { dg-final { scan-tree-dump-times "clique 1 base 3" 1 "ealias" { xfail *-*-* } } }
+! { dg-final { scan-tree-dump-times "clique 1 base 4" 1 "ealias" { xfail *-*-* } } }
+! { dg-final { scan-tree-dump-times "clique 1 base 5" 1 "ealias" { xfail *-*-* } } }
! { dg-final { scan-tree-dump-times "(?n)clique .* base .*" 8 "ealias" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-alias.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-alias.f95
index 8d6ccb338b9..62f9a713991 100644
--- a/gcc/testsuite/gfortran.dg/goacc/kernels-alias.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-alias.f95
@@ -15,9 +15,11 @@ program main
end program main
+! The xfails occur in light of the new OpenACC data semantics.
+
! { dg-final { scan-tree-dump-times "clique 1 base 1" 4 "ealias" } }
-! { dg-final { scan-tree-dump-times "clique 1 base 2" 1 "ealias" } }
-! { dg-final { scan-tree-dump-times "clique 1 base 3" 1 "ealias" } }
-! { dg-final { scan-tree-dump-times "clique 1 base 4" 1 "ealias" } }
-! { dg-final { scan-tree-dump-times "clique 1 base 5" 1 "ealias" } }
+! { dg-final { scan-tree-dump-times "clique 1 base 2" 1 "ealias" { xfail *-*-* } } }
+! { dg-final { scan-tree-dump-times "clique 1 base 3" 1 "ealias" { xfail *-*-* } } }
+! { dg-final { scan-tree-dump-times "clique 1 base 4" 1 "ealias" { xfail *-*-* } } }
+! { dg-final { scan-tree-dump-times "clique 1 base 5" 1 "ealias" { xfail *-*-* } } }
! { dg-final { scan-tree-dump-times "(?n)clique .* base .*" 8 "ealias" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95
index 7daca59020e..a70f1e737bd 100644
--- a/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95
@@ -21,10 +21,10 @@ end program test
! { dg-final { scan-tree-dump-times "num_workers" 1 "original" } }
! { dg-final { scan-tree-dump-times "vector_length" 1 "original" } }
-! { dg-final { scan-tree-dump-times "map\\(force_tofrom:i\\)" 1 "original" } }
-! { dg-final { scan-tree-dump-times "map\\(force_to:j\\)" 1 "original" } }
-! { dg-final { scan-tree-dump-times "map\\(force_from:k\\)" 1 "original" } }
-! { dg-final { scan-tree-dump-times "map\\(force_alloc:m\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "map\\(tofrom:i\\)" 1 "original" } }
+! { 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\\(force_present:o\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "map\\(tofrom:p\\)" 1 "original" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/nested-function-1.f90 b/gcc/testsuite/gfortran.dg/goacc/nested-function-1.f90
index 2fcaa400ee3..005193f30a7 100644
--- a/gcc/testsuite/gfortran.dg/goacc/nested-function-1.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/nested-function-1.f90
@@ -25,6 +25,8 @@ contains
local_a (:) = 5
local_arg = 5
+ !$acc update device(local_a) if_present
+
!$acc kernels loop &
!$acc gang(num:local_arg) worker(local_arg) vector(local_arg) &
!$acc wait async(local_arg)
@@ -54,12 +56,16 @@ contains
enddo
enddo
!$acc end kernels loop
+
+ !$acc exit data copyout(local_a) delete(local_i) finalize
end subroutine local
subroutine nonlocal ()
nonlocal_a (:) = 5
nonlocal_arg = 5
+ !$acc update device(nonlocal_a) if_present
+
!$acc kernels loop &
!$acc gang(num:nonlocal_arg) worker(nonlocal_arg) vector(nonlocal_arg) &
!$acc wait async(nonlocal_arg)
@@ -89,5 +95,7 @@ contains
enddo
enddo
!$acc end kernels loop
+
+ !$acc exit data copyout(nonlocal_a) delete(nonlocal_i) finalize
end subroutine nonlocal
end program main
diff --git a/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 b/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95
index 5b2e01d4878..2697bb79e7f 100644
--- a/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95
@@ -1,5 +1,4 @@
-! { dg-do compile }
-! { dg-additional-options "-fdump-tree-original" }
+! { dg-additional-options "-fdump-tree-original" }
! test for tree-dump-original and spaces-commas
@@ -15,6 +14,7 @@ program test
!$acc end parallel
end program test
+
! { dg-final { scan-tree-dump-times "pragma acc parallel" 1 "original" } }
! { dg-final { scan-tree-dump-times "if" 1 "original" } }
@@ -24,10 +24,10 @@ end program test
! { dg-final { scan-tree-dump-times "vector_length" 1 "original" } }
! { dg-final { scan-tree-dump-times "reduction\\(max:q\\)" 1 "original" } }
-! { dg-final { scan-tree-dump-times "map\\(force_tofrom:i\\)" 1 "original" } }
-! { dg-final { scan-tree-dump-times "map\\(force_to:j\\)" 1 "original" } }
-! { dg-final { scan-tree-dump-times "map\\(force_from:k\\)" 1 "original" } }
-! { dg-final { scan-tree-dump-times "map\\(force_alloc:m\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "map\\(tofrom:i\\)" 1 "original" } }
+! { 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\\(force_present:o\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "map\\(tofrom:p\\)" 1 "original" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/reduction-promotions.f90 b/gcc/testsuite/gfortran.dg/goacc/reduction-promotions.f90
index 6ff913ade8d..1d247ca238e 100644
--- a/gcc/testsuite/gfortran.dg/goacc/reduction-promotions.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/reduction-promotions.f90
@@ -38,9 +38,7 @@ program test
!$acc end parallel
end program test
-! { dg-final { scan-tree-dump-times "map.tofrom:v1" 8 "gimple" } }
-! { dg-final { scan-tree-dump-times "map.tofrom:v2" 8 "gimple" } }
-! { dg-final { scan-tree-dump-times "map.force_tofrom:v1" 1 "gimple" } }
-! { dg-final { scan-tree-dump-times "map.force_tofrom:v2" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "map.tofrom:v1" 9 "gimple" } }
+! { dg-final { scan-tree-dump-times "map.tofrom:v2" 9 "gimple" } }
! { dg-final { scan-tree-dump-times "map.force_present:v1" 1 "gimple" } }
! { dg-final { scan-tree-dump-times "map.force_present:v2" 1 "gimple" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/update-if_present-1.f90 b/gcc/testsuite/gfortran.dg/goacc/update-if_present-1.f90
new file mode 100644
index 00000000000..a183aae44c5
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/update-if_present-1.f90
@@ -0,0 +1,27 @@
+! Test valid usages of the if_present clause.
+
+! { dg-additional-options "-fdump-tree-omplower" }
+
+subroutine t
+ implicit none
+ integer a, b, c(10)
+ real, allocatable :: x, y, z(:)
+
+ a = 5
+ b = 10
+ c(:) = -1
+
+ allocate (x, y, z(100))
+
+ !$acc update self(a) if_present
+ !$acc update device(b) if_present async
+ !$acc update host(c(1:3)) wait(4) if_present
+ !$acc update self(c) device(a) host(b) if_present async(10) if(a == 10)
+
+ !$acc update self(x) if_present
+ !$acc update device(y) if_present async
+ !$acc update host(z(1:3)) wait(3) if_present
+ !$acc update self(z) device(y) host(x) if_present async(4) if(a == 1)
+end subroutine t
+
+! { dg-final { scan-tree-dump-times " if_present" 8 "omplower" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/update-if_present-2.f90 b/gcc/testsuite/gfortran.dg/goacc/update-if_present-2.f90
new file mode 100644
index 00000000000..e73c2dc0875
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/update-if_present-2.f90
@@ -0,0 +1,52 @@
+! Test invalid usages of the if_present clause.
+
+subroutine t1
+ implicit none
+ !$acc routine gang if_present ! { dg-error "Unclassifiable OpenACC directive" }
+ integer a, b, c(10)
+ real, allocatable :: x, y, z(:)
+
+ a = 5
+ b = 10
+ c(:) = -1
+
+ allocate (x, y, z(100))
+
+ !$acc enter data copyin(a) if_present ! { dg-error "Unclassifiable OpenACC directive" }
+ !$acc exit data copyout(a) if_present ! { dg-error "Unclassifiable OpenACC directive" }
+
+ !$acc data copy(a) if_present ! { dg-error "Unclassifiable OpenACC directive" }
+ !$acc end data ! { dg-error "Unexpected ..ACC END DATA statement" }
+
+ !$acc declare link(a) if_present ! { dg-error "Unexpected junk after" }
+
+ !$acc init if_present ! { dg-error "Unclassifiable OpenACC directive" }
+ !$acc shutdown if_present ! { dg-error "Unclassifiable OpenACC directive" }
+
+ !$acc update self(a) device_type(nvidia) device(b) if_present ! { dg-error "Unclassifiable OpenACC directive" }
+end subroutine t1
+
+subroutine t2
+ implicit none
+ integer a, b, c(10)
+
+ a = 5
+ b = 10
+ c(:) = -1
+
+ !$acc parallel
+ !$acc loop if_present ! { dg-error "Unclassifiable OpenACC directive" }
+ do b = 1, 10
+ end do
+ !$acc end parallel
+
+ !$acc kernels loop if_present ! { dg-error "Unclassifiable OpenACC directive" }
+ do b = 1, 10
+ end do
+ !$acc end kernels loop ! { dg-error "Unexpected ..ACC END KERNELS LOOP statement" }
+
+ !$acc parallel loop if_present ! { dg-error "Unclassifiable OpenACC directive" }
+ do b = 1, 10
+ end do
+ !$acc end parallel loop ! { dg-error "Unexpected ..ACC END PARALLEL LOOP statement" }
+end subroutine t2
--
2.17.1
^ permalink raw reply [flat|nested] 31+ messages in thread* Re: [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior - runtime tests
2018-06-19 16:56 ` Cesar Philippidis
` (5 preceding siblings ...)
2018-06-19 17:02 ` [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior - compiler tests Cesar Philippidis
@ 2018-06-19 17:03 ` Cesar Philippidis
6 siblings, 0 replies; 31+ messages in thread
From: Cesar Philippidis @ 2018-06-19 17:03 UTC (permalink / raw)
To: gcc-patches, Jakub Jelinek
[-- Attachment #1: Type: text/plain, Size: 138 bytes --]
This patch updates the existing OpenACC libgomp runtime tests with the
new OpenACC 2.5 data clause semantics.
Is it OK for trunk?
Cesar
[-- Attachment #2: 0002-libgomp-tests.patch --]
[-- Type: text/x-patch, Size: 18523 bytes --]
2018-06-19 Chung-Lin Tang <cltang@codesourcery.com>
Thomas Schwinge <thomas@codesourcery.com>
Cesar Philippidis <cesar@codesourcery.com>
libgomp/
* testsuite/libgomp.oacc-c-c++-common/data-already-1.c: Update test
case to utilize OpenACC 2.5 data clause semantics.
* testsuite/libgomp.oacc-c-c++-common/data-already-2.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/data-already-3.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/data-already-4.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/data-already-5.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/data-already-6.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/data-already-7.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/data-already-8.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/lib-16.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/lib-25.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/lib-32.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/lib-83.c: Likewise.
* testsuite/libgomp.oacc-fortran/data-5.f90: New test.
* testsuite/libgomp.oacc-fortran/data-already-1.f: Update test case to
utilize OpenACC 2.5 data clause semantics.
* testsuite/libgomp.oacc-fortran/data-already-2.f: Likewise.
* testsuite/libgomp.oacc-fortran/data-already-3.f: Likewise.
* testsuite/libgomp.oacc-fortran/data-already-4.f: Likewise.
* testsuite/libgomp.oacc-fortran/data-already-5.f: Likewise.
* testsuite/libgomp.oacc-fortran/data-already-6.f: Likewise.
* testsuite/libgomp.oacc-fortran/data-already-7.f: Likewise.
* testsuite/libgomp.oacc-fortran/data-already-8.f: Likewise.
* testsuite/libgomp.oacc-fortran/lib-32-1.f: Likewise.
* testsuite/libgomp.oacc-fortran/lib-32-2.f: Likewise.
From 43ba3412e27dda1fa619f39d2720bf109b83508c Mon Sep 17 00:00:00 2001
From: Cesar Philippidis <cesar@codesourcery.com>
Date: Tue, 19 Jun 2018 09:29:21 -0700
Subject: [PATCH 2/7] libgomp tests
---
.../data-already-1.c | 2 -
.../data-already-2.c | 2 -
.../data-already-3.c | 2 -
.../data-already-4.c | 2 -
.../data-already-5.c | 2 -
.../data-already-6.c | 2 -
.../data-already-7.c | 2 -
.../data-already-8.c | 2 -
.../libgomp.oacc-c-c++-common/lib-16.c | 23 ++++----
.../libgomp.oacc-c-c++-common/lib-25.c | 20 ++++---
.../libgomp.oacc-c-c++-common/lib-32.c | 4 +-
.../libgomp.oacc-c-c++-common/lib-83.c | 22 ++++----
.../testsuite/libgomp.oacc-fortran/data-5.f90 | 56 +++++++++++++++++++
.../libgomp.oacc-fortran/data-already-1.f | 2 -
.../libgomp.oacc-fortran/data-already-2.f | 2 -
.../libgomp.oacc-fortran/data-already-3.f | 2 -
.../libgomp.oacc-fortran/data-already-4.f | 2 -
.../libgomp.oacc-fortran/data-already-5.f | 2 -
.../libgomp.oacc-fortran/data-already-6.f | 2 -
.../libgomp.oacc-fortran/data-already-7.f | 2 -
.../libgomp.oacc-fortran/data-already-8.f | 2 -
.../testsuite/libgomp.oacc-fortran/lib-32-1.f | 4 +-
.../testsuite/libgomp.oacc-fortran/lib-32-2.f | 4 +-
23 files changed, 96 insertions(+), 69 deletions(-)
create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/data-5.f90
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-1.c
index 0ed53a41a96..fd3b77dcff5 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-1.c
@@ -18,5 +18,3 @@ main (int argc, char *argv[])
}
/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
-/* { dg-output "Trying to map into device \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) object when \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) is already mapped" } */
-/* { dg-shouldfail "" } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-2.c
index 00adf2a2bf4..0118b2568e2 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-2.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-2.c
@@ -18,5 +18,3 @@ main (int argc, char *argv[])
}
/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
-/* { dg-output "Trying to map into device \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) object when \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) is already mapped" } */
-/* { dg-shouldfail "" } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-3.c
index 04073e323d9..b346c69826f 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-3.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-3.c
@@ -18,5 +18,3 @@ main (int argc, char *argv[])
}
/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
-/* { dg-output "already mapped to" } */
-/* { dg-shouldfail "" } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-4.c
index cfa5cb2c44a..e99ad33d9be 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-4.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-4.c
@@ -16,5 +16,3 @@ main (int argc, char *argv[])
}
/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
-/* { dg-output "already mapped to" } */
-/* { dg-shouldfail "" } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-5.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-5.c
index e15c3fb1aaa..f8370c006df 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-5.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-5.c
@@ -16,5 +16,3 @@ main (int argc, char *argv[])
}
/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
-/* { dg-output "already mapped to" } */
-/* { dg-shouldfail "" } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-6.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-6.c
index 4570c74965c..d7f4deb18e4 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-6.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-6.c
@@ -16,5 +16,3 @@ main (int argc, char *argv[])
}
/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
-/* { dg-output "already mapped to" } */
-/* { dg-shouldfail "" } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-7.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-7.c
index 467cf39aa5d..54be59507ca 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-7.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-7.c
@@ -16,5 +16,3 @@ main (int argc, char *argv[])
}
/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
-/* { dg-output "already mapped to" } */
-/* { dg-shouldfail "" } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-8.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-8.c
index f41431c1418..e5c0f9cfb32 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-8.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-8.c
@@ -18,5 +18,3 @@ main (int argc, char *argv[])
}
/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
-/* { dg-output "Trying to map into device \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) object when \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) is already mapped" } */
-/* { dg-shouldfail "" } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-16.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-16.c
index c81a78de26d..9a1c9d30bef 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-16.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-16.c
@@ -1,8 +1,5 @@
-/* Test if duplicate data mappings with acc_copy_in. */
+/* Test if acc_copyin has present_or_ and reference counting behavior. */
-/* { dg-do run { target openacc_nvidia_accel_selected } } */
-
-#include <stdio.h>
#include <stdlib.h>
#include <openacc.h>
@@ -21,15 +18,21 @@ main (int argc, char **argv)
}
(void) acc_copyin (h, N);
-
- fprintf (stderr, "CheCKpOInT\n");
(void) acc_copyin (h, N);
+ acc_copyout (h, N);
+
+ if (!acc_is_present (h, N))
+ abort ();
+
+ acc_copyout (h, N);
+
+#if !ACC_MEM_SHARED
+ if (acc_is_present (h, N))
+ abort ();
+#endif
+
free (h);
return 0;
}
-
-/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
-/* { dg-output "\\\[\[0-9a-fA-FxX\]+,\\\+256\\\] already mapped to \\\[\[0-9a-fA-FxX\]+,\\\+256\\\]" } */
-/* { dg-shouldfail "" } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-25.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-25.c
index 5f00ccb3885..9b42dee9d87 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-25.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-25.c
@@ -1,8 +1,5 @@
-/* Exercise acc_create and acc_delete on nvidia targets. */
+/* Exercise acc_create and acc_delete. */
-/* { dg-do run { target openacc_nvidia_accel_selected } } */
-
-#include <stdio.h>
#include <stdlib.h>
#include <openacc.h>
@@ -19,18 +16,23 @@ main (int argc, char **argv)
if (!d)
abort ();
- fprintf (stderr, "CheCKpOInT\n");
d = acc_create (h, N);
if (!d)
abort ();
acc_delete (h, N);
+ if (!acc_is_present (h, N))
+ abort ();
+
+ acc_delete (h, N);
+
+#if !ACC_MEM_SHARED
+ if (acc_is_present (h, N))
+ abort ();
+#endif
+
free (h);
return 0;
}
-
-/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
-/* { dg-output "\\\[\[0-9a-fA-FxX\]+,\\\+256\\\] already mapped to \\\[\[0-9a-fA-FxX\]+,\\\+256\\\]" } */
-/* { dg-shouldfail "" } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-32.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-32.c
index 1696fb6f9ef..9ec345361d8 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-32.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-32.c
@@ -127,7 +127,7 @@ main (int argc, char **argv)
h[i] = i + 10;
}
- acc_copyout (h, S);
+ acc_copyout_finalize (h, S);
d = NULL;
if (!shared_mem)
if (acc_is_present (h, S))
@@ -236,7 +236,7 @@ main (int argc, char **argv)
abort ();
}
- acc_delete (h, S);
+ acc_delete_finalize (h, S);
d = NULL;
if (!shared_mem)
if (acc_is_present (h, S))
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-83.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-83.c
index 1c2e52b4c5f..51b7ee73b9c 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-83.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-83.c
@@ -5,21 +5,19 @@
#include <stdlib.h>
#include <unistd.h>
#include <openacc.h>
-#include "timer.h"
+#include <cuda.h>
+#include <sys/time.h>
int
main (int argc, char **argv)
{
- float atime;
CUstream stream;
CUresult r;
+ struct timeval tv1, tv2;
+ time_t t1;
acc_init (acc_device_nvidia);
- (void) acc_get_device_num (acc_device_nvidia);
-
- init_timers (1);
-
stream = (CUstream) acc_get_cuda_stream (0);
if (stream != NULL)
abort ();
@@ -34,22 +32,22 @@ main (int argc, char **argv)
if (!acc_set_cuda_stream (0, stream))
abort ();
- start_timer (0);
+ gettimeofday (&tv1, NULL);
acc_wait_all_async (0);
acc_wait (0);
- atime = stop_timer (0);
+ gettimeofday (&tv2, NULL);
- if (0.010 < atime)
+ t1 = ((tv2.tv_sec - tv1.tv_sec) * 1000000) + (tv2.tv_usec - tv1.tv_usec);
+
+ if (t1 > 1000)
{
- fprintf (stderr, "actual time too long\n");
+ fprintf (stderr, "too long\n");
abort ();
}
- fini_timers ();
-
acc_shutdown (acc_device_nvidia);
exit (0);
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/data-5.f90 b/libgomp/testsuite/libgomp.oacc-fortran/data-5.f90
new file mode 100644
index 00000000000..a8843dedc22
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/data-5.f90
@@ -0,0 +1,56 @@
+! { dg-do run }
+! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
+
+program refcount_test
+ use openacc
+ integer, allocatable :: h(:)
+ integer i, N
+
+ N = 256
+ allocate (h(N))
+
+ do i = 1, N
+ h(i) = i
+ end do
+
+ !$acc enter data create (h(1:N))
+ !$acc enter data copyin (h(1:N))
+ !$acc enter data copyin (h(1:N))
+ !$acc enter data copyin (h(1:N))
+
+ call acc_update_self (h)
+ do i = 1, N
+ if (h(i) .eq. i) c = c + 1
+ end do
+ ! h[] should be filled with uninitialized device values,
+ ! abort if it's not.
+ if (c .eq. N) call abort
+
+ h(:) = 0
+
+ !$acc parallel present (h(1:N))
+ do i = 1, N
+ h(i) = 111
+ end do
+ !$acc end parallel
+
+ ! No actual copyout should happen.
+ call acc_copyout (h)
+ do i = 1, N
+ if (h(i) .ne. 0) call abort
+ end do
+
+ !$acc exit data delete (h(1:N))
+
+ ! This should not actually be deleted yet.
+ if (acc_is_present (h) .eqv. .FALSE.) call abort
+
+ !$acc exit data copyout (h(1:N)) finalize
+
+ do i = 1, N
+ if (h(i) .ne. 111) call abort
+ end do
+
+ if (acc_is_present (h) .eqv. .TRUE.) call abort
+
+end program refcount_test
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/data-already-1.f b/libgomp/testsuite/libgomp.oacc-fortran/data-already-1.f
index 9e99cc60be5..fab0ffc99cc 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/data-already-1.f
+++ b/libgomp/testsuite/libgomp.oacc-fortran/data-already-1.f
@@ -14,5 +14,3 @@
END
! { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
-! { dg-output "Trying to map into device \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) object when \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) is already mapped" }
-! { dg-shouldfail "" }
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/data-already-2.f b/libgomp/testsuite/libgomp.oacc-fortran/data-already-2.f
index b908a0c0702..bd03062670f 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/data-already-2.f
+++ b/libgomp/testsuite/libgomp.oacc-fortran/data-already-2.f
@@ -14,5 +14,3 @@
END
! { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
-! { dg-output "Trying to map into device \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) object when \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) is already mapped" }
-! { dg-shouldfail "" }
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/data-already-3.f b/libgomp/testsuite/libgomp.oacc-fortran/data-already-3.f
index d93e1c5cedd..60ea3864e4e 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/data-already-3.f
+++ b/libgomp/testsuite/libgomp.oacc-fortran/data-already-3.f
@@ -13,5 +13,3 @@
END
! { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
-! { dg-output "already mapped to" }
-! { dg-shouldfail "" }
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/data-already-4.f b/libgomp/testsuite/libgomp.oacc-fortran/data-already-4.f
index ea76e058d9c..2abdbf0f868 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/data-already-4.f
+++ b/libgomp/testsuite/libgomp.oacc-fortran/data-already-4.f
@@ -12,5 +12,3 @@
END
! { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
-! { dg-output "already mapped to" }
-! { dg-shouldfail "" }
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/data-already-5.f b/libgomp/testsuite/libgomp.oacc-fortran/data-already-5.f
index 19df1f8bde2..f361d8c1772 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/data-already-5.f
+++ b/libgomp/testsuite/libgomp.oacc-fortran/data-already-5.f
@@ -12,5 +12,3 @@
END
! { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
-! { dg-output "already mapped to" }
-! { dg-shouldfail "" }
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/data-already-6.f b/libgomp/testsuite/libgomp.oacc-fortran/data-already-6.f
index 2bd1079087d..a864737c692 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/data-already-6.f
+++ b/libgomp/testsuite/libgomp.oacc-fortran/data-already-6.f
@@ -12,5 +12,3 @@
END
! { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
-! { dg-output "already mapped to" }
-! { dg-shouldfail "" }
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/data-already-7.f b/libgomp/testsuite/libgomp.oacc-fortran/data-already-7.f
index 1342360f53a..0d893280e40 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/data-already-7.f
+++ b/libgomp/testsuite/libgomp.oacc-fortran/data-already-7.f
@@ -12,5 +12,3 @@
END
! { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
-! { dg-output "already mapped to" }
-! { dg-shouldfail "" }
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/data-already-8.f b/libgomp/testsuite/libgomp.oacc-fortran/data-already-8.f
index b206547bca7..7a41f67225a 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/data-already-8.f
+++ b/libgomp/testsuite/libgomp.oacc-fortran/data-already-8.f
@@ -14,5 +14,3 @@
END
! { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
-! { dg-output "Trying to map into device \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) object when \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) is already mapped" }
-! { dg-shouldfail "" }
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/lib-32-1.f b/libgomp/testsuite/libgomp.oacc-fortran/lib-32-1.f
index 99e8f357764..3f979eba034 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/lib-32-1.f
+++ b/libgomp/testsuite/libgomp.oacc-fortran/lib-32-1.f
@@ -90,7 +90,7 @@
H(I) = I + 10
END DO
- CALL ACC_COPYOUT (H, INT (SIZEOF (H), 4))
+ CALL ACC_COPYOUT_FINALIZE (H, INT (SIZEOF (H), 4))
IF (.NOT. SHARED_MEM) THEN
IF (ACC_IS_PRESENT (H, INT (SIZEOF (H), 8))) STOP 11
ENDIF
@@ -163,7 +163,7 @@
IF (H(I) .NE. I + MERGE (18, 17, SHARED_MEM)) STOP 23
END DO
- CALL ACC_DELETE (H)
+ CALL ACC_DELETE_FINALIZE (H)
IF (.NOT. SHARED_MEM) THEN
IF (ACC_IS_PRESENT (H, INT (SIZEOF (H), 4))) STOP 24
ENDIF
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/lib-32-2.f b/libgomp/testsuite/libgomp.oacc-fortran/lib-32-2.f
index 514c04e7bca..642552cae60 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/lib-32-2.f
+++ b/libgomp/testsuite/libgomp.oacc-fortran/lib-32-2.f
@@ -90,7 +90,7 @@
H(I) = I + 10
END DO
- CALL ACC_COPYOUT (H, INT (SIZEOF (H), 4))
+ CALL ACC_COPYOUT_FINALIZE (H, INT (SIZEOF (H), 4))
IF (.NOT. SHARED_MEM) THEN
IF (ACC_IS_PRESENT (H, INT (SIZEOF (H), 8))) STOP 11
ENDIF
@@ -163,7 +163,7 @@
IF (H(I) .NE. I + MERGE (18, 17, SHARED_MEM)) STOP 23
END DO
- CALL ACC_DELETE (H)
+ CALL ACC_DELETE_FINALIZE (H)
IF (.NOT. SHARED_MEM) THEN
IF (ACC_IS_PRESENT (H, INT (SIZEOF (H), 4))) STOP 24
ENDIF
--
2.17.1
^ permalink raw reply [flat|nested] 31+ messages in thread