+2018-06-20 Chung-Lin Tang <cltang@codesourcery.com>
+ Thomas Schwinge <thomas@codesourcery.com>
+ Cesar Philippidis <cesar@codesourcery.com>
+
+ * 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.
+
2018-06-20 Jakub Jelinek <jakub@redhat.com>
PR debug/86194
+2018-06-20 Chung-Lin Tang <cltang@codesourcery.com>
+ Thomas Schwinge <thomas@codesourcery.com>
+ Cesar Philippidis <cesar@codesourcery.com>
+
+ * c-pragma.h (enum pragma_omp_clause): Add
+ PRAGMA_OACC_CLAUSE_{FINALIZE,IF_PRESENT}. Remove
+ PRAGMA_OACC_CLAUSE_PRESENT_OR_{COPY,COPYIN,COPYOUT,CREATE}.
+
2018-06-20 Jakub Jelinek <jakub@redhat.com>
PR c++/86210
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,
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,
+2018-06-20 Chung-Lin Tang <cltang@codesourcery.com>
+ Thomas Schwinge <thomas@codesourcery.com>
+ Cesar Philippidis <cesar@codesourcery.com>
+
+ * 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.
+
2018-06-16 Kugan Vivekanandarajah <kuganv@linaro.org>
* c-typeck.c (build_unary_op): Handle ABSU_EXPR;
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))
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;
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))
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))
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))
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,
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;
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:
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 ();
}
return list;
}
-/* OpenACC:
+/* OpenACC 2.5:
auto
+ finalize
independent
nohost
seq */
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";
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);
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";
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);
| (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)
| (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)
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;
| (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 \
| (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
| (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) )
| (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) )
| (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
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;
+2018-06-20 Chung-Lin Tang <cltang@codesourcery.com>
+ Thomas Schwinge <thomas@codesourcery.com>
+ Cesar Philippidis <cesar@codesourcery.com>
+
+ * 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.
+
2018-06-20 Marek Polacek <polacek@redhat.com>
PR c++/86240
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))
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;
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))
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))
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,
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;
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:
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 ();
}
return list;
}
-/* OpenACC 2.0:
+/* OpenACC 2.5:
auto
+ finalize
independent
nohost
seq */
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);
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,
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);
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);
| (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)
| (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)
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;
| (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 \
| (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
| (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) )
| (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) \
| (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
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 ();
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:
+2018-06-20 Chung-Lin Tang <cltang@codesourcery.com>
+ Thomas Schwinge <thomas@codesourcery.com>
+ Cesar Philippidis <cesar@codesourcery.com>
+
+ * 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.
+
2018-06-18 Eric Botcazou <ebotcazou@adacore.com>
* trans-decl.c (gfc_get_fake_result_decl): Revert latest change.
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;
}
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,
OMP_CLAUSE_DELETE,
OMP_CLAUSE_AUTO,
OMP_CLAUSE_TILE,
+ OMP_CLAUSE_IF_PRESENT,
+ OMP_CLAUSE_FINALIZE,
/* This must come last. */
OMP_MASK2_LAST
};
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)
{
{
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 (",
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 (",
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':
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)
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 (",
&& 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],
}
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
}
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))
&& 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))
(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 \
#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 \
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);
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",
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;
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);
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:
case OMP_CLAUSE_AUTO:
case OMP_CLAUSE_SEQ:
case OMP_CLAUSE_TILE:
+ case OMP_CLAUSE_IF_PRESENT:
+ case OMP_CLAUSE_FINALIZE:
break;
default:
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;
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);
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;
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);
}
/* 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;
&& 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);
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:
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_:
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
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)
+2018-06-20 Chung-Lin Tang <cltang@codesourcery.com>
+ Thomas Schwinge <thomas@codesourcery.com>
+ Cesar Philippidis <cesar@codesourcery.com>
+
+ * 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.
+
2018-06-20 Jakub Jelinek <jakub@redhat.com>
PR c++/86210
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)
{
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)
#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)
#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" } */
}
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" } } */
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" } } */
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" } } */
--- /dev/null
+/* 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" } } */
+}
+
}
}
+/* 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" } } */
}
}
+/* 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" } } */
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 } } */
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" } */
/* "#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. */
#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. */
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. */
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)
{
return vb3;
}
-#pragma acc routine
+#pragma acc routine seq
int
func2 (int a)
{
static int vb7; /* { dg-error "clause used in" } */
#pragma acc declare link (vb7)
-#pragma acc routine
+#pragma acc routine seq
int
func3 (int a)
{
static int vb10;
#pragma acc declare create (vb10)
-#pragma acc routine
+#pragma acc routine seq
int
func4 (int a)
{
extern int vb13;
#pragma acc declare device_resident (vb13)
-#pragma acc routine
+#pragma acc routine seq
int
func5 (int a)
{
return vb13;
}
-#pragma acc routine
+#pragma acc routine seq
int
func6 (int a)
{
--- /dev/null
+/* 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" } } */
--- /dev/null
+/* 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++)
+ ;
+}
+
-#pragma acc routine
+#pragma acc routine seq
template <typename T> T
accDouble(int val)
{
#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;
#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;
}
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;
}
! { 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" } }
! { 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" } }
!$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
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;
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;
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;
!$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
--- /dev/null
+! 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
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" } }
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" } }
! { 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" } }
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)
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)
enddo
enddo
!$acc end kernels loop
+
+ !$acc exit data copyout(nonlocal_a) delete(nonlocal_i) finalize
end subroutine nonlocal
end program main
-! { dg-do compile }
-! { dg-additional-options "-fdump-tree-original" }
+! { dg-additional-options "-fdump-tree-original" }
! test for tree-dump-original and spaces-commas
!$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" } }
! { 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" } }
!$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" } }
--- /dev/null
+! 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" } }
--- /dev/null
+! 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
/* 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
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
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
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. */
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[] =
"num_workers",
"vector_length",
"tile",
- "_griddim_"
+ "_griddim_",
+ "if_present",
+ "finalize",
};
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:
+2018-06-20 Chung-Lin Tang <cltang@codesourcery.com>
+ Thomas Schwinge <thomas@codesourcery.com>
+ Cesar Philippidis <cesar@codesourcery.com>
+
+ * 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.
+ * 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.
+
2018-05-21 Janus Weil <janus@gcc.gnu.org>
PR fortran/85841
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;
};
};
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 **,
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 */
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;
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);
(void *) n->host_start, (int) host_size, (void *) h);
}
+ /* Mark for removal. */
+ n->refcount = 1;
+
t = n->tgt;
if (t->refcount == 2)
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))
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);
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 *
}
#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)
(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
}
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)
{
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;
}
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;
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);
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__);
}
#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);
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)
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);
}
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". */
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;
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",
}
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;
}
}
}
{
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",
}
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;
}
}
acc_dev->openacc.async_set_async_func (async);
+ bool update_device = false;
for (i = 0; i < mapnum; ++i)
{
unsigned char kind = kinds[i] & 0xff;
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;
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);
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;
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
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
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
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
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
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
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
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;
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
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
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;
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. */
+ 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)
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);
}
}
}
/* { 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 "" } */
}
/* { 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 "" } */
}
/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
-/* { dg-output "already mapped to" } */
-/* { dg-shouldfail "" } */
}
/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
-/* { dg-output "already mapped to" } */
-/* { dg-shouldfail "" } */
}
/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
-/* { dg-output "already mapped to" } */
-/* { dg-shouldfail "" } */
}
/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
-/* { dg-output "already mapped to" } */
-/* { dg-shouldfail "" } */
}
/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
-/* { dg-output "already mapped to" } */
-/* { dg-shouldfail "" } */
}
/* { 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 "" } */
-/* 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>
}
(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 "" } */
-/* 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>
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 "" } */
h[i] = i + 10;
}
- acc_copyout (h, S);
+ acc_copyout_finalize (h, S);
d = NULL;
if (!shared_mem)
if (acc_is_present (h, S))
abort ();
}
- acc_delete (h, S);
+ acc_delete_finalize (h, S);
d = NULL;
if (!shared_mem)
if (acc_is_present (h, S))
#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 ();
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);
--- /dev/null
+! { 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
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 "" }
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 "" }
END
! { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
-! { dg-output "already mapped to" }
-! { dg-shouldfail "" }
END
! { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
-! { dg-output "already mapped to" }
-! { dg-shouldfail "" }
END
! { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
-! { dg-output "already mapped to" }
-! { dg-shouldfail "" }
END
! { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
-! { dg-output "already mapped to" }
-! { dg-shouldfail "" }
END
! { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
-! { dg-output "already mapped to" }
-! { dg-shouldfail "" }
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 "" }
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
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
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
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
projects / platform / upstream / gcc.git / commitdiff