mirror of
git://gcc.gnu.org/git/gcc.git
synced 2025-04-14 20:51:17 +08:00
Update OpenACC data clause semantics to the 2.5 behavior
gcc/c-family/
* c-pragma.h (enum pragma_omp_clause): Add
PRAGMA_OACC_CLAUSE_{FINALIZE,IF_PRESENT}. Remove
PRAGMA_OACC_CLAUSE_PRESENT_OR_{COPY,COPYIN,COPYOUT,CREATE}.
gcc/c/
* c-parser.c (c_parser_omp_clause_name): Add support for finalize
and if_present. Make present_or_{copy,copyin,copyout,create} aliases
to their non-present_or_* counterparts. Make 'self' an alias to
PRAGMA_OACC_CLAUSE_HOST.
(c_parser_oacc_data_clause): Update GOMP mappings for
PRAGMA_OACC_CLAUSE_{COPY,COPYIN,COPYOUT,CREATE,DELETE}. Remove
PRAGMA_OACC_CLAUSE_{SELF,PRESENT_OR_*}.
(c_parser_oacc_all_clauses): Handle finalize and if_present clauses.
Remove support for present_or_* clauses.
(OACC_KERNELS_CLAUSE_MASK): Remove PRESENT_OR_* clauses.
(OACC_PARALLEL_CLAUSE_MASK): Likewise.
(OACC_DECLARE_CLAUSE_MASK): Likewise.
(OACC_DATA_CLAUSE_MASK): Likewise.
(OACC_ENTER_DATA_CLAUSE_MASK): Remove PRESENT_OR_* clauses.
(OACC_EXIT_DATA_CLAUSE_MASK): Add FINALIZE clause.
(OACC_UPDATE_CLAUSE_MASK): Remove SELF, add IF_PRESENT.
(c_parser_oacc_declare): Remove PRESENT_OR_* clauses.
* c-typeck.c (c_finish_omp_clauses): Handle IF_PRESENT and FINALIZE.
gcc/cp/
* parser.c (cp_parser_omp_clause_name): Add support for finalize
and if_present. Make present_or_{copy,copyin,copyout,create} aliases
to their non-present_or_* counterparts. Make 'self' an alias to
PRAGMA_OACC_CLAUSE_HOST.
(cp_parser_oacc_data_clause): Update GOMP mappings for
PRAGMA_OACC_CLAUSE_{COPY,COPYIN,COPYOUT,CREATE,DELETE}. Remove
PRAGMA_OACC_CLAUSE_{SELF,PRESENT_OR_*}.
(cp_parser_oacc_all_clauses): Handle finalize and if_present clauses.
Remove support for present_or_* clauses.
(OACC_KERNELS_CLAUSE_MASK): Remove PRESENT_OR_* clauses.
(OACC_PARALLEL_CLAUSE_MASK): Likewise.
(OACC_DECLARE_CLAUSE_MASK): Likewise.
(OACC_DATA_CLAUSE_MASK): Likewise.
(OACC_ENTER_DATA_CLAUSE_MASK): Remove PRESENT_OR_* clauses.
(OACC_EXIT_DATA_CLAUSE_MASK): Add FINALIZE clause.
(OACC_UPDATE_CLAUSE_MASK): Remove SELF, add IF_PRESENT.
(cp_parser_oacc_declare): Remove PRESENT_OR_* clauses.
* pt.c (tsubst_omp_clauses): Handle IF_PRESENT and FINALIZE.
* semantics.c (finish_omp_clauses): Handle IF_PRESENT and FINALIZE.
gcc/fortran/
* gfortran.h (gfc_omp_clauses): Add unsigned if_present, finalize
bitfields.
* openmp.c (enum omp_mask2): Remove OMP_CLAUSE_PRESENT_OR_*. Add
OMP_CLAUSE_{IF_PRESENT,FINALIZE}.
(gfc_match_omp_clauses): Update handling of copy, copyin, copyout,
create, deviceptr, present_of_*. Add support for finalize and
if_present.
(OACC_PARALLEL_CLAUSES): Remove PRESENT_OR_* clauses.
(OACC_KERNELS_CLAUSES): Likewise.
(OACC_DATA_CLAUSES): Likewise.
(OACC_DECLARE_CLAUSES): Likewise.
(OACC_UPDATE_CLAUSES): Add IF_PRESENT clause.
(OACC_ENTER_DATA_CLAUSES): Remove PRESENT_OR_* clauses.
(OACC_EXIT_DATA_CLAUSES): Add FINALIZE clause.
(gfc_match_oacc_declare): Update to OpenACC 2.5 semantics.
* trans-openmp.c (gfc_trans_omp_clauses): Add support for IF_PRESENT
and FINALIZE.
gcc/
* gimplify.c (gimplify_scan_omp_clauses): Add support for
OMP_CLAUSE_{IF_PRESENT,FINALIZE}.
(gimplify_adjust_omp_clauses): Likewise.
(gimplify_oacc_declare_1): Add support for GOMP_MAP_RELEASE, remove
support for GOMP_MAP_FORCE_{ALLOC,TO,FROM,TOFROM}.
(gimplify_omp_target_update): Update handling of acc update and
enter/exit data.
* omp-low.c (install_var_field): Remove unused parameter
base_pointers_restrict.
(scan_sharing_clauses): Remove base_pointers_restrict parameter.
Update call to install_var_field. Handle OMP_CLAUSE_{IF_PRESENT,
FINALIZE}
(omp_target_base_pointers_restrict_p): Delete.
(scan_omp_target): Update call to scan_sharing_clauses.
* tree-core.h (enum omp_clause_code): Add OMP_CLAUSE_{IF_PRESENT,
FINALIZE}.
* tree-nested.c (convert_nonlocal_omp_clauses): Handle
OMP_CLAUSE_{IF_PRESENT,FINALIZE}.
(convert_local_omp_clauses): Likewise.
* tree-pretty-print.c (dump_omp_clause): Likewise.
* tree.c (omp_clause_num_ops): Add entries for OMP_CLAUSE_{IF_PRESENT,
FINALIZE}.
(omp_clause_code_name): Likewise.
gcc/testsuite/
* c-c++-common/goacc/declare-1.c: Update test case to utilize OpenACC
2.5 data clause semantics.
* c-c++-common/goacc/declare-2.c: Likewise.
* c-c++-common/goacc/default-4.c: Likewise.
* c-c++-common/goacc/finalize-1.c: New test.
* c-c++-common/goacc/kernels-alias-2.c: Update test case to utilize
OpenACC 2.5 data clause semantics.
* c-c++-common/goacc/kernels-alias.c: Likewise.
* c-c++-common/goacc/routine-5.c: Likewise.
* c-c++-common/goacc/update-if_present-1.c: New test.
* c-c++-common/goacc/update-if_present-2.c: New test.
* g++.dg/goacc/template.C: Update test case to utilize OpenACC
2.5 data clause semantics.
* gfortran.dg/goacc/combined-directives.f90: Likewise.
* gfortran.dg/goacc/data-tree.f95: Likewise.
* gfortran.dg/goacc/declare-2.f95: Likewise.
* gfortran.dg/goacc/default-4.f: Likewise.
* gfortran.dg/goacc/enter-exit-data.f95: Likewise.
* gfortran.dg/goacc/finalize-1.f: New test.
* gfortran.dg/goacc/kernels-alias-2.f95: Update test case to utilize
OpenACC 2.5 data clause semantics.
* gfortran.dg/goacc/kernels-alias.f95: Likewise.
* gfortran.dg/goacc/kernels-tree.f95: Likewise.
* gfortran.dg/goacc/nested-function-1.f90: Likewise.
* gfortran.dg/goacc/parallel-tree.f95: Likewise.
* gfortran.dg/goacc/reduction-promotions.f90: Likewise.
* gfortran.dg/goacc/update-if_present-1.f90: New test.
* gfortran.dg/goacc/update-if_present-2.f90: New test.
libgomp/
* libgomp.h (struct splay_tree_key_s): Add dynamic_refcount member.
(gomp_acc_remove_pointer): Update declaration.
(gomp_acc_declare_allocate): Declare.
(gomp_remove_var): Declare.
* libgomp.map (OACC_2.5): Define.
* oacc-mem.c (acc_map_data): Update refcount.
(acc_unmap_data): Likewise.
(present_create_copy): Likewise.
(acc_create): Add FLAG_PRESENT when calling present_create_copy.
(acc_copyin): Likewise.
(FLAG_FINALIZE): Define.
(delete_copyout): Update dynamic refcounts, add support for FINALIZE.
(acc_delete_finalize): New function.
(acc_delete_finalize_async): New function.
(acc_copyout_finalize): New function.
(acc_copyout_finalize_async): New function.
(gomp_acc_insert_pointer): Update refcounts.
(gomp_acc_remove_pointer): Return if data is not present on the
accelerator.
* oacc-parallel.c (find_pset): Rename to find_pointer.
(find_pointer): Add support for GOMP_MAP_POINTER.
(handle_ftn_pointers): New function.
(GOACC_parallel_keyed): Update refcounts of variables.
(GOACC_enter_exit_data): Add support for finalized data mappings.
Add support for GOMP_MAP_{TO,ALLOC,RELESE,FROM}. Update handling
of fortran arrays.
(GOACC_update): Add support for GOMP_MAP_{ALWAYS_POINTER,TO,FROM}.
(GOACC_declare): Add support for GOMP_MAP_RELEASE, remove support
for GOMP_MAP_FORCE_FROM.
* openacc.f90 (module openacc_internal): Add
acc_copyout_finalize_{32_h,64_h,array_h,_l}, and
acc_delete_finalize_{32_h,64_h,array_h,_l}. Add interfaces for
acc_copyout_finalize and acc_delete_finalize.
(acc_copyout_finalize_32_h): New subroutine.
(acc_copyout_finalize_64_h): New subroutine.
(acc_copyout_finalize_array_h): New subroutine.
(acc_delete_finalize_32_h): New subroutine.
(acc_delete_finalize_64_h): New subroutine.
(acc_delete_finalize_array_h): New subroutine.
* openacc.h (acc_copyout_finalize): Declare.
(acc_copyout_finalize_async): Declare.
(acc_delete_finalize): Declare.
(acc_delete_finalize_async): Declare.
* openacc_lib.h (acc_copyout_finalize): New interface.
(acc_delete_finalize): New interface.
* target.c (gomp_map_vars): Update dynamic_refcount.
(gomp_remove_var): New function.
(gomp_unmap_vars): Use it.
(gomp_unload_image_from_device): Likewise.
* 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.
Co-Authored-By: Cesar Philippidis <cesar@codesourcery.com>
Co-Authored-By: Thomas Schwinge <thomas@codesourcery.com>
From-SVN: r261813
This commit is contained in:
Chung-Lin Tang
Cesar Philippidis
parent
f41b7612a9
commit
829c6349e9
@ -1,3 +1,31 @@
|
||||
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
|
||||
|
||||
@ -1,3 +1,11 @@
|
||||
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
|
||||
|
||||
@ -138,16 +138,13 @@ enum pragma_omp_clause {
|
||||
PRAGMA_OACC_CLAUSE_DELETE,
|
||||
PRAGMA_OACC_CLAUSE_DEVICEPTR,
|
||||
PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT,
|
||||
PRAGMA_OACC_CLAUSE_FINALIZE,
|
||||
PRAGMA_OACC_CLAUSE_GANG,
|
||||
PRAGMA_OACC_CLAUSE_HOST,
|
||||
PRAGMA_OACC_CLAUSE_INDEPENDENT,
|
||||
PRAGMA_OACC_CLAUSE_NUM_GANGS,
|
||||
PRAGMA_OACC_CLAUSE_NUM_WORKERS,
|
||||
PRAGMA_OACC_CLAUSE_PRESENT,
|
||||
PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY,
|
||||
PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN,
|
||||
PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT,
|
||||
PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE,
|
||||
PRAGMA_OACC_CLAUSE_SELF,
|
||||
PRAGMA_OACC_CLAUSE_SEQ,
|
||||
PRAGMA_OACC_CLAUSE_TILE,
|
||||
@ -156,6 +153,7 @@ enum pragma_omp_clause {
|
||||
PRAGMA_OACC_CLAUSE_VECTOR_LENGTH,
|
||||
PRAGMA_OACC_CLAUSE_WAIT,
|
||||
PRAGMA_OACC_CLAUSE_WORKER,
|
||||
PRAGMA_OACC_CLAUSE_IF_PRESENT,
|
||||
PRAGMA_OACC_CLAUSE_COLLAPSE = PRAGMA_OMP_CLAUSE_COLLAPSE,
|
||||
PRAGMA_OACC_CLAUSE_COPYIN = PRAGMA_OMP_CLAUSE_COPYIN,
|
||||
PRAGMA_OACC_CLAUSE_DEVICE = PRAGMA_OMP_CLAUSE_DEVICE,
|
||||
|
||||
@ -1,3 +1,26 @@
|
||||
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;
|
||||
|
||||
115
gcc/c/c-parser.c
115
gcc/c/c-parser.c
@ -11259,6 +11259,8 @@ c_parser_omp_clause_name (c_parser *parser)
|
||||
case 'f':
|
||||
if (!strcmp ("final", p))
|
||||
result = PRAGMA_OMP_CLAUSE_FINAL;
|
||||
else if (!strcmp ("finalize", p))
|
||||
result = PRAGMA_OACC_CLAUSE_FINALIZE;
|
||||
else if (!strcmp ("firstprivate", p))
|
||||
result = PRAGMA_OMP_CLAUSE_FIRSTPRIVATE;
|
||||
else if (!strcmp ("from", p))
|
||||
@ -11277,7 +11279,9 @@ c_parser_omp_clause_name (c_parser *parser)
|
||||
result = PRAGMA_OACC_CLAUSE_HOST;
|
||||
break;
|
||||
case 'i':
|
||||
if (!strcmp ("inbranch", p))
|
||||
if (!strcmp ("if_present", p))
|
||||
result = PRAGMA_OACC_CLAUSE_IF_PRESENT;
|
||||
else if (!strcmp ("inbranch", p))
|
||||
result = PRAGMA_OMP_CLAUSE_INBRANCH;
|
||||
else if (!strcmp ("independent", p))
|
||||
result = PRAGMA_OACC_CLAUSE_INDEPENDENT;
|
||||
@ -11325,18 +11329,20 @@ c_parser_omp_clause_name (c_parser *parser)
|
||||
result = PRAGMA_OMP_CLAUSE_PARALLEL;
|
||||
else if (!strcmp ("present", p))
|
||||
result = PRAGMA_OACC_CLAUSE_PRESENT;
|
||||
/* As of OpenACC 2.5, these are now aliases of the non-present_or
|
||||
clauses. */
|
||||
else if (!strcmp ("present_or_copy", p)
|
||||
|| !strcmp ("pcopy", p))
|
||||
result = PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY;
|
||||
result = PRAGMA_OACC_CLAUSE_COPY;
|
||||
else if (!strcmp ("present_or_copyin", p)
|
||||
|| !strcmp ("pcopyin", p))
|
||||
result = PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN;
|
||||
result = PRAGMA_OACC_CLAUSE_COPYIN;
|
||||
else if (!strcmp ("present_or_copyout", p)
|
||||
|| !strcmp ("pcopyout", p))
|
||||
result = PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT;
|
||||
result = PRAGMA_OACC_CLAUSE_COPYOUT;
|
||||
else if (!strcmp ("present_or_create", p)
|
||||
|| !strcmp ("pcreate", p))
|
||||
result = PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE;
|
||||
result = PRAGMA_OACC_CLAUSE_CREATE;
|
||||
else if (!strcmp ("priority", p))
|
||||
result = PRAGMA_OMP_CLAUSE_PRIORITY;
|
||||
else if (!strcmp ("private", p))
|
||||
@ -11355,6 +11361,8 @@ c_parser_omp_clause_name (c_parser *parser)
|
||||
result = PRAGMA_OMP_CLAUSE_SCHEDULE;
|
||||
else if (!strcmp ("sections", p))
|
||||
result = PRAGMA_OMP_CLAUSE_SECTIONS;
|
||||
else if (!strcmp ("self", p)) /* "self" is a synonym for "host". */
|
||||
result = PRAGMA_OACC_CLAUSE_HOST;
|
||||
else if (!strcmp ("seq", p))
|
||||
result = PRAGMA_OACC_CLAUSE_SEQ;
|
||||
else if (!strcmp ("shared", p))
|
||||
@ -11363,8 +11371,6 @@ c_parser_omp_clause_name (c_parser *parser)
|
||||
result = PRAGMA_OMP_CLAUSE_SIMD;
|
||||
else if (!strcmp ("simdlen", p))
|
||||
result = PRAGMA_OMP_CLAUSE_SIMDLEN;
|
||||
else if (!strcmp ("self", p))
|
||||
result = PRAGMA_OACC_CLAUSE_SELF;
|
||||
break;
|
||||
case 't':
|
||||
if (!strcmp ("taskgroup", p))
|
||||
@ -11646,15 +11652,7 @@ c_parser_omp_var_list_parens (c_parser *parser, enum omp_clause_code kind,
|
||||
copyout ( variable-list )
|
||||
create ( variable-list )
|
||||
delete ( variable-list )
|
||||
present ( variable-list )
|
||||
present_or_copy ( variable-list )
|
||||
pcopy ( variable-list )
|
||||
present_or_copyin ( variable-list )
|
||||
pcopyin ( variable-list )
|
||||
present_or_copyout ( variable-list )
|
||||
pcopyout ( variable-list )
|
||||
present_or_create ( variable-list )
|
||||
pcreate ( variable-list ) */
|
||||
present ( variable-list ) */
|
||||
|
||||
static tree
|
||||
c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind,
|
||||
@ -11664,19 +11662,19 @@ c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind,
|
||||
switch (c_kind)
|
||||
{
|
||||
case PRAGMA_OACC_CLAUSE_COPY:
|
||||
kind = GOMP_MAP_FORCE_TOFROM;
|
||||
kind = GOMP_MAP_TOFROM;
|
||||
break;
|
||||
case PRAGMA_OACC_CLAUSE_COPYIN:
|
||||
kind = GOMP_MAP_FORCE_TO;
|
||||
kind = GOMP_MAP_TO;
|
||||
break;
|
||||
case PRAGMA_OACC_CLAUSE_COPYOUT:
|
||||
kind = GOMP_MAP_FORCE_FROM;
|
||||
kind = GOMP_MAP_FROM;
|
||||
break;
|
||||
case PRAGMA_OACC_CLAUSE_CREATE:
|
||||
kind = GOMP_MAP_FORCE_ALLOC;
|
||||
kind = GOMP_MAP_ALLOC;
|
||||
break;
|
||||
case PRAGMA_OACC_CLAUSE_DELETE:
|
||||
kind = GOMP_MAP_DELETE;
|
||||
kind = GOMP_MAP_RELEASE;
|
||||
break;
|
||||
case PRAGMA_OACC_CLAUSE_DEVICE:
|
||||
kind = GOMP_MAP_FORCE_TO;
|
||||
@ -11685,7 +11683,6 @@ c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind,
|
||||
kind = GOMP_MAP_DEVICE_RESIDENT;
|
||||
break;
|
||||
case PRAGMA_OACC_CLAUSE_HOST:
|
||||
case PRAGMA_OACC_CLAUSE_SELF:
|
||||
kind = GOMP_MAP_FORCE_FROM;
|
||||
break;
|
||||
case PRAGMA_OACC_CLAUSE_LINK:
|
||||
@ -11694,18 +11691,6 @@ c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind,
|
||||
case PRAGMA_OACC_CLAUSE_PRESENT:
|
||||
kind = GOMP_MAP_FORCE_PRESENT;
|
||||
break;
|
||||
case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY:
|
||||
kind = GOMP_MAP_TOFROM;
|
||||
break;
|
||||
case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN:
|
||||
kind = GOMP_MAP_TO;
|
||||
break;
|
||||
case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT:
|
||||
kind = GOMP_MAP_FROM;
|
||||
break;
|
||||
case PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE:
|
||||
kind = GOMP_MAP_ALLOC;
|
||||
break;
|
||||
default:
|
||||
gcc_unreachable ();
|
||||
}
|
||||
@ -12597,8 +12582,9 @@ c_parser_oacc_shape_clause (c_parser *parser, omp_clause_code kind,
|
||||
return list;
|
||||
}
|
||||
|
||||
/* OpenACC:
|
||||
/* OpenACC 2.5:
|
||||
auto
|
||||
finalize
|
||||
independent
|
||||
nohost
|
||||
seq */
|
||||
@ -13955,6 +13941,11 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
|
||||
clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
|
||||
c_name = "device_resident";
|
||||
break;
|
||||
case PRAGMA_OACC_CLAUSE_FINALIZE:
|
||||
clauses = c_parser_oacc_simple_clause (parser, OMP_CLAUSE_FINALIZE,
|
||||
clauses);
|
||||
c_name = "finalize";
|
||||
break;
|
||||
case PRAGMA_OACC_CLAUSE_FIRSTPRIVATE:
|
||||
clauses = c_parser_omp_clause_firstprivate (parser, clauses);
|
||||
c_name = "firstprivate";
|
||||
@ -13972,6 +13963,11 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
|
||||
clauses = c_parser_omp_clause_if (parser, clauses, false);
|
||||
c_name = "if";
|
||||
break;
|
||||
case PRAGMA_OACC_CLAUSE_IF_PRESENT:
|
||||
clauses = c_parser_oacc_simple_clause (parser, OMP_CLAUSE_IF_PRESENT,
|
||||
clauses);
|
||||
c_name = "if_present";
|
||||
break;
|
||||
case PRAGMA_OACC_CLAUSE_INDEPENDENT:
|
||||
clauses = c_parser_oacc_simple_clause (parser, OMP_CLAUSE_INDEPENDENT,
|
||||
clauses);
|
||||
@ -13997,22 +13993,6 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
|
||||
clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
|
||||
c_name = "present";
|
||||
break;
|
||||
case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY:
|
||||
clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
|
||||
c_name = "present_or_copy";
|
||||
break;
|
||||
case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN:
|
||||
clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
|
||||
c_name = "present_or_copyin";
|
||||
break;
|
||||
case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT:
|
||||
clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
|
||||
c_name = "present_or_copyout";
|
||||
break;
|
||||
case PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE:
|
||||
clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
|
||||
c_name = "present_or_create";
|
||||
break;
|
||||
case PRAGMA_OACC_CLAUSE_PRIVATE:
|
||||
clauses = c_parser_omp_clause_private (parser, clauses);
|
||||
c_name = "private";
|
||||
@ -14021,10 +14001,6 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
|
||||
clauses = c_parser_omp_clause_reduction (parser, clauses);
|
||||
c_name = "reduction";
|
||||
break;
|
||||
case PRAGMA_OACC_CLAUSE_SELF:
|
||||
clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
|
||||
c_name = "self";
|
||||
break;
|
||||
case PRAGMA_OACC_CLAUSE_SEQ:
|
||||
clauses = c_parser_oacc_simple_clause (parser, OMP_CLAUSE_SEQ,
|
||||
clauses);
|
||||
@ -14417,11 +14393,7 @@ c_parser_oacc_cache (location_t loc, c_parser *parser)
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) )
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT))
|
||||
|
||||
static tree
|
||||
c_parser_oacc_data (location_t loc, c_parser *parser, bool *if_p)
|
||||
@ -14451,11 +14423,7 @@ c_parser_oacc_data (location_t loc, c_parser *parser, bool *if_p)
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_LINK) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) )
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT))
|
||||
|
||||
static void
|
||||
c_parser_oacc_declare (c_parser *parser)
|
||||
@ -14490,8 +14458,8 @@ c_parser_oacc_declare (c_parser *parser)
|
||||
switch (OMP_CLAUSE_MAP_KIND (t))
|
||||
{
|
||||
case GOMP_MAP_FIRSTPRIVATE_POINTER:
|
||||
case GOMP_MAP_FORCE_ALLOC:
|
||||
case GOMP_MAP_FORCE_TO:
|
||||
case GOMP_MAP_ALLOC:
|
||||
case GOMP_MAP_TO:
|
||||
case GOMP_MAP_FORCE_DEVICEPTR:
|
||||
case GOMP_MAP_DEVICE_RESIDENT:
|
||||
break;
|
||||
@ -14604,8 +14572,6 @@ c_parser_oacc_declare (c_parser *parser)
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
|
||||
|
||||
#define OACC_EXIT_DATA_CLAUSE_MASK \
|
||||
@ -14613,6 +14579,7 @@ c_parser_oacc_declare (c_parser *parser)
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DELETE) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_FINALIZE) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
|
||||
|
||||
static void
|
||||
@ -14756,10 +14723,6 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name,
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_GANGS) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_WORKERS) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
|
||||
|
||||
@ -14777,10 +14740,6 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name,
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_GANGS) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_WORKERS) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
|
||||
@ -15008,7 +14967,7 @@ c_finish_oacc_routine (struct oacc_routine_data *data, tree fndecl,
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_HOST) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SELF) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF_PRESENT) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
|
||||
|
||||
static void
|
||||
|
||||
@ -13897,6 +13897,8 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
|
||||
case OMP_CLAUSE_WORKER:
|
||||
case OMP_CLAUSE_VECTOR:
|
||||
case OMP_CLAUSE_TILE:
|
||||
case OMP_CLAUSE_IF_PRESENT:
|
||||
case OMP_CLAUSE_FINALIZE:
|
||||
pc = &OMP_CLAUSE_CHAIN (c);
|
||||
continue;
|
||||
|
||||
|
||||
@ -1,3 +1,27 @@
|
||||
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
|
||||
|
||||
114
gcc/cp/parser.c
114
gcc/cp/parser.c
@ -31375,6 +31375,8 @@ cp_parser_omp_clause_name (cp_parser *parser)
|
||||
case 'f':
|
||||
if (!strcmp ("final", p))
|
||||
result = PRAGMA_OMP_CLAUSE_FINAL;
|
||||
else if (!strcmp ("finalize", p))
|
||||
result = PRAGMA_OACC_CLAUSE_FINALIZE;
|
||||
else if (!strcmp ("firstprivate", p))
|
||||
result = PRAGMA_OMP_CLAUSE_FIRSTPRIVATE;
|
||||
else if (!strcmp ("from", p))
|
||||
@ -31393,7 +31395,9 @@ cp_parser_omp_clause_name (cp_parser *parser)
|
||||
result = PRAGMA_OACC_CLAUSE_HOST;
|
||||
break;
|
||||
case 'i':
|
||||
if (!strcmp ("inbranch", p))
|
||||
if (!strcmp ("if_present", p))
|
||||
result = PRAGMA_OACC_CLAUSE_IF_PRESENT;
|
||||
else if (!strcmp ("inbranch", p))
|
||||
result = PRAGMA_OMP_CLAUSE_INBRANCH;
|
||||
else if (!strcmp ("independent", p))
|
||||
result = PRAGMA_OACC_CLAUSE_INDEPENDENT;
|
||||
@ -31443,16 +31447,16 @@ cp_parser_omp_clause_name (cp_parser *parser)
|
||||
result = PRAGMA_OACC_CLAUSE_PRESENT;
|
||||
else if (!strcmp ("present_or_copy", p)
|
||||
|| !strcmp ("pcopy", p))
|
||||
result = PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY;
|
||||
result = PRAGMA_OACC_CLAUSE_COPY;
|
||||
else if (!strcmp ("present_or_copyin", p)
|
||||
|| !strcmp ("pcopyin", p))
|
||||
result = PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN;
|
||||
result = PRAGMA_OACC_CLAUSE_COPYIN;
|
||||
else if (!strcmp ("present_or_copyout", p)
|
||||
|| !strcmp ("pcopyout", p))
|
||||
result = PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT;
|
||||
result = PRAGMA_OACC_CLAUSE_COPYOUT;
|
||||
else if (!strcmp ("present_or_create", p)
|
||||
|| !strcmp ("pcreate", p))
|
||||
result = PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE;
|
||||
result = PRAGMA_OACC_CLAUSE_CREATE;
|
||||
else if (!strcmp ("priority", p))
|
||||
result = PRAGMA_OMP_CLAUSE_PRIORITY;
|
||||
else if (!strcmp ("proc_bind", p))
|
||||
@ -31469,8 +31473,8 @@ cp_parser_omp_clause_name (cp_parser *parser)
|
||||
result = PRAGMA_OMP_CLAUSE_SCHEDULE;
|
||||
else if (!strcmp ("sections", p))
|
||||
result = PRAGMA_OMP_CLAUSE_SECTIONS;
|
||||
else if (!strcmp ("self", p))
|
||||
result = PRAGMA_OACC_CLAUSE_SELF;
|
||||
else if (!strcmp ("self", p)) /* "self" is a synonym for "host". */
|
||||
result = PRAGMA_OACC_CLAUSE_HOST;
|
||||
else if (!strcmp ("seq", p))
|
||||
result = PRAGMA_OACC_CLAUSE_SEQ;
|
||||
else if (!strcmp ("shared", p))
|
||||
@ -31730,15 +31734,7 @@ cp_parser_omp_var_list (cp_parser *parser, enum omp_clause_code kind, tree list)
|
||||
copyout ( variable-list )
|
||||
create ( variable-list )
|
||||
delete ( variable-list )
|
||||
present ( variable-list )
|
||||
present_or_copy ( variable-list )
|
||||
pcopy ( variable-list )
|
||||
present_or_copyin ( variable-list )
|
||||
pcopyin ( variable-list )
|
||||
present_or_copyout ( variable-list )
|
||||
pcopyout ( variable-list )
|
||||
present_or_create ( variable-list )
|
||||
pcreate ( variable-list ) */
|
||||
present ( variable-list ) */
|
||||
|
||||
static tree
|
||||
cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind,
|
||||
@ -31748,19 +31744,19 @@ cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind,
|
||||
switch (c_kind)
|
||||
{
|
||||
case PRAGMA_OACC_CLAUSE_COPY:
|
||||
kind = GOMP_MAP_FORCE_TOFROM;
|
||||
kind = GOMP_MAP_TOFROM;
|
||||
break;
|
||||
case PRAGMA_OACC_CLAUSE_COPYIN:
|
||||
kind = GOMP_MAP_FORCE_TO;
|
||||
kind = GOMP_MAP_TO;
|
||||
break;
|
||||
case PRAGMA_OACC_CLAUSE_COPYOUT:
|
||||
kind = GOMP_MAP_FORCE_FROM;
|
||||
kind = GOMP_MAP_FROM;
|
||||
break;
|
||||
case PRAGMA_OACC_CLAUSE_CREATE:
|
||||
kind = GOMP_MAP_FORCE_ALLOC;
|
||||
kind = GOMP_MAP_ALLOC;
|
||||
break;
|
||||
case PRAGMA_OACC_CLAUSE_DELETE:
|
||||
kind = GOMP_MAP_DELETE;
|
||||
kind = GOMP_MAP_RELEASE;
|
||||
break;
|
||||
case PRAGMA_OACC_CLAUSE_DEVICE:
|
||||
kind = GOMP_MAP_FORCE_TO;
|
||||
@ -31769,7 +31765,6 @@ cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind,
|
||||
kind = GOMP_MAP_DEVICE_RESIDENT;
|
||||
break;
|
||||
case PRAGMA_OACC_CLAUSE_HOST:
|
||||
case PRAGMA_OACC_CLAUSE_SELF:
|
||||
kind = GOMP_MAP_FORCE_FROM;
|
||||
break;
|
||||
case PRAGMA_OACC_CLAUSE_LINK:
|
||||
@ -31778,18 +31773,6 @@ cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind,
|
||||
case PRAGMA_OACC_CLAUSE_PRESENT:
|
||||
kind = GOMP_MAP_FORCE_PRESENT;
|
||||
break;
|
||||
case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY:
|
||||
kind = GOMP_MAP_TOFROM;
|
||||
break;
|
||||
case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN:
|
||||
kind = GOMP_MAP_TO;
|
||||
break;
|
||||
case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT:
|
||||
kind = GOMP_MAP_FROM;
|
||||
break;
|
||||
case PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE:
|
||||
kind = GOMP_MAP_ALLOC;
|
||||
break;
|
||||
default:
|
||||
gcc_unreachable ();
|
||||
}
|
||||
@ -31828,8 +31811,9 @@ cp_parser_oacc_data_clause_deviceptr (cp_parser *parser, tree list)
|
||||
return list;
|
||||
}
|
||||
|
||||
/* OpenACC 2.0:
|
||||
/* OpenACC 2.5:
|
||||
auto
|
||||
finalize
|
||||
independent
|
||||
nohost
|
||||
seq */
|
||||
@ -33794,6 +33778,11 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
|
||||
clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
|
||||
c_name = "device_resident";
|
||||
break;
|
||||
case PRAGMA_OACC_CLAUSE_FINALIZE:
|
||||
clauses = cp_parser_oacc_simple_clause (parser, OMP_CLAUSE_FINALIZE,
|
||||
clauses, here);
|
||||
c_name = "finalize";
|
||||
break;
|
||||
case PRAGMA_OACC_CLAUSE_FIRSTPRIVATE:
|
||||
clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_FIRSTPRIVATE,
|
||||
clauses);
|
||||
@ -33812,6 +33801,12 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
|
||||
clauses = cp_parser_omp_clause_if (parser, clauses, here, false);
|
||||
c_name = "if";
|
||||
break;
|
||||
case PRAGMA_OACC_CLAUSE_IF_PRESENT:
|
||||
clauses = cp_parser_oacc_simple_clause (parser,
|
||||
OMP_CLAUSE_IF_PRESENT,
|
||||
clauses, here);
|
||||
c_name = "if_present";
|
||||
break;
|
||||
case PRAGMA_OACC_CLAUSE_INDEPENDENT:
|
||||
clauses = cp_parser_oacc_simple_clause (parser,
|
||||
OMP_CLAUSE_INDEPENDENT,
|
||||
@ -33838,22 +33833,6 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
|
||||
clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
|
||||
c_name = "present";
|
||||
break;
|
||||
case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY:
|
||||
clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
|
||||
c_name = "present_or_copy";
|
||||
break;
|
||||
case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN:
|
||||
clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
|
||||
c_name = "present_or_copyin";
|
||||
break;
|
||||
case PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT:
|
||||
clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
|
||||
c_name = "present_or_copyout";
|
||||
break;
|
||||
case PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE:
|
||||
clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
|
||||
c_name = "present_or_create";
|
||||
break;
|
||||
case PRAGMA_OACC_CLAUSE_PRIVATE:
|
||||
clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_PRIVATE,
|
||||
clauses);
|
||||
@ -33863,10 +33842,6 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
|
||||
clauses = cp_parser_omp_clause_reduction (parser, clauses);
|
||||
c_name = "reduction";
|
||||
break;
|
||||
case PRAGMA_OACC_CLAUSE_SELF:
|
||||
clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
|
||||
c_name = "self";
|
||||
break;
|
||||
case PRAGMA_OACC_CLAUSE_SEQ:
|
||||
clauses = cp_parser_oacc_simple_clause (parser, OMP_CLAUSE_SEQ,
|
||||
clauses, here);
|
||||
@ -36802,11 +36777,7 @@ cp_parser_oacc_cache (cp_parser *parser, cp_token *pragma_tok)
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE))
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) )
|
||||
|
||||
static tree
|
||||
cp_parser_oacc_data (cp_parser *parser, cp_token *pragma_tok, bool *if_p)
|
||||
@ -36861,11 +36832,7 @@ cp_parser_oacc_host_data (cp_parser *parser, cp_token *pragma_tok, bool *if_p)
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_LINK) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE))
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) )
|
||||
|
||||
static tree
|
||||
cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok)
|
||||
@ -36898,8 +36865,8 @@ cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok)
|
||||
switch (OMP_CLAUSE_MAP_KIND (t))
|
||||
{
|
||||
case GOMP_MAP_FIRSTPRIVATE_POINTER:
|
||||
case GOMP_MAP_FORCE_ALLOC:
|
||||
case GOMP_MAP_FORCE_TO:
|
||||
case GOMP_MAP_ALLOC:
|
||||
case GOMP_MAP_TO:
|
||||
case GOMP_MAP_FORCE_DEVICEPTR:
|
||||
case GOMP_MAP_DEVICE_RESIDENT:
|
||||
break;
|
||||
@ -37010,8 +36977,6 @@ cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok)
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
|
||||
|
||||
#define OACC_EXIT_DATA_CLAUSE_MASK \
|
||||
@ -37019,6 +36984,7 @@ cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok)
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DELETE) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_FINALIZE) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
|
||||
|
||||
static tree
|
||||
@ -37131,10 +37097,6 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name,
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_GANGS) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_WORKERS) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
|
||||
|
||||
@ -37151,10 +37113,6 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name,
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_GANGS) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_WORKERS) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPY) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRIVATE) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH) \
|
||||
@ -37215,7 +37173,7 @@ cp_parser_oacc_kernels_parallel (cp_parser *parser, cp_token *pragma_tok,
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_HOST) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SELF) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF_PRESENT) \
|
||||
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT))
|
||||
|
||||
static tree
|
||||
|
||||
@ -16107,6 +16107,8 @@ tsubst_omp_clauses (tree clauses, enum c_omp_region_type ort,
|
||||
case OMP_CLAUSE_INDEPENDENT:
|
||||
case OMP_CLAUSE_AUTO:
|
||||
case OMP_CLAUSE_SEQ:
|
||||
case OMP_CLAUSE_IF_PRESENT:
|
||||
case OMP_CLAUSE_FINALIZE:
|
||||
break;
|
||||
default:
|
||||
gcc_unreachable ();
|
||||
|
||||
@ -7091,6 +7091,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
|
||||
case OMP_CLAUSE_AUTO:
|
||||
case OMP_CLAUSE_INDEPENDENT:
|
||||
case OMP_CLAUSE_SEQ:
|
||||
case OMP_CLAUSE_IF_PRESENT:
|
||||
case OMP_CLAUSE_FINALIZE:
|
||||
break;
|
||||
|
||||
case OMP_CLAUSE_TILE:
|
||||
|
||||
@ -1,3 +1,25 @@
|
||||
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.
|
||||
|
||||
@ -1344,6 +1344,7 @@ typedef struct gfc_omp_clauses
|
||||
gfc_expr_list *tile_list;
|
||||
unsigned async:1, gang:1, worker:1, vector:1, seq:1, independent:1;
|
||||
unsigned wait:1, par_auto:1, gang_static:1;
|
||||
unsigned if_present:1, finalize:1;
|
||||
locus loc;
|
||||
|
||||
}
|
||||
|
||||
@ -796,10 +796,6 @@ enum omp_mask2
|
||||
OMP_CLAUSE_COPYOUT,
|
||||
OMP_CLAUSE_CREATE,
|
||||
OMP_CLAUSE_PRESENT,
|
||||
OMP_CLAUSE_PRESENT_OR_COPY,
|
||||
OMP_CLAUSE_PRESENT_OR_COPYIN,
|
||||
OMP_CLAUSE_PRESENT_OR_COPYOUT,
|
||||
OMP_CLAUSE_PRESENT_OR_CREATE,
|
||||
OMP_CLAUSE_DEVICEPTR,
|
||||
OMP_CLAUSE_GANG,
|
||||
OMP_CLAUSE_WORKER,
|
||||
@ -813,6 +809,8 @@ enum omp_mask2
|
||||
OMP_CLAUSE_DELETE,
|
||||
OMP_CLAUSE_AUTO,
|
||||
OMP_CLAUSE_TILE,
|
||||
OMP_CLAUSE_IF_PRESENT,
|
||||
OMP_CLAUSE_FINALIZE,
|
||||
/* This must come last. */
|
||||
OMP_MASK2_LAST
|
||||
};
|
||||
@ -1041,7 +1039,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
|
||||
if ((mask & OMP_CLAUSE_COPY)
|
||||
&& gfc_match ("copy ( ") == MATCH_YES
|
||||
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
|
||||
OMP_MAP_FORCE_TOFROM))
|
||||
OMP_MAP_TOFROM))
|
||||
continue;
|
||||
if (mask & OMP_CLAUSE_COPYIN)
|
||||
{
|
||||
@ -1049,7 +1047,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
|
||||
{
|
||||
if (gfc_match ("copyin ( ") == MATCH_YES
|
||||
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
|
||||
OMP_MAP_FORCE_TO))
|
||||
OMP_MAP_TO))
|
||||
continue;
|
||||
}
|
||||
else if (gfc_match_omp_variable_list ("copyin (",
|
||||
@ -1060,7 +1058,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
|
||||
if ((mask & OMP_CLAUSE_COPYOUT)
|
||||
&& gfc_match ("copyout ( ") == MATCH_YES
|
||||
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
|
||||
OMP_MAP_FORCE_FROM))
|
||||
OMP_MAP_FROM))
|
||||
continue;
|
||||
if ((mask & OMP_CLAUSE_COPYPRIVATE)
|
||||
&& gfc_match_omp_variable_list ("copyprivate (",
|
||||
@ -1070,7 +1068,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
|
||||
if ((mask & OMP_CLAUSE_CREATE)
|
||||
&& gfc_match ("create ( ") == MATCH_YES
|
||||
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
|
||||
OMP_MAP_FORCE_ALLOC))
|
||||
OMP_MAP_ALLOC))
|
||||
continue;
|
||||
break;
|
||||
case 'd':
|
||||
@ -1106,7 +1104,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
|
||||
if ((mask & OMP_CLAUSE_DELETE)
|
||||
&& gfc_match ("delete ( ") == MATCH_YES
|
||||
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
|
||||
OMP_MAP_DELETE))
|
||||
OMP_MAP_RELEASE))
|
||||
continue;
|
||||
if ((mask & OMP_CLAUSE_DEPEND)
|
||||
&& gfc_match ("depend ( ") == MATCH_YES)
|
||||
@ -1161,19 +1159,10 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
|
||||
OMP_MAP_FORCE_TO))
|
||||
continue;
|
||||
if ((mask & OMP_CLAUSE_DEVICEPTR)
|
||||
&& gfc_match ("deviceptr ( ") == MATCH_YES)
|
||||
{
|
||||
gfc_omp_namelist **list = &c->lists[OMP_LIST_MAP];
|
||||
gfc_omp_namelist **head = NULL;
|
||||
if (gfc_match_omp_variable_list ("", list, true, NULL,
|
||||
&head, false) == MATCH_YES)
|
||||
{
|
||||
gfc_omp_namelist *n;
|
||||
for (n = *head; n; n = n->next)
|
||||
n->u.map_op = OMP_MAP_FORCE_DEVICEPTR;
|
||||
continue;
|
||||
}
|
||||
}
|
||||
&& gfc_match ("deviceptr ( ") == MATCH_YES
|
||||
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
|
||||
OMP_MAP_FORCE_DEVICEPTR))
|
||||
continue;
|
||||
if ((mask & OMP_CLAUSE_DEVICE_RESIDENT)
|
||||
&& gfc_match_omp_variable_list
|
||||
("device_resident (",
|
||||
@ -1202,6 +1191,14 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
|
||||
&& c->final_expr == NULL
|
||||
&& gfc_match ("final ( %e )", &c->final_expr) == MATCH_YES)
|
||||
continue;
|
||||
if ((mask & OMP_CLAUSE_FINALIZE)
|
||||
&& !c->finalize
|
||||
&& gfc_match ("finalize") == MATCH_YES)
|
||||
{
|
||||
c->finalize = true;
|
||||
needs_space = true;
|
||||
continue;
|
||||
}
|
||||
if ((mask & OMP_CLAUSE_FIRSTPRIVATE)
|
||||
&& gfc_match_omp_variable_list ("firstprivate (",
|
||||
&c->lists[OMP_LIST_FIRSTPRIVATE],
|
||||
@ -1274,6 +1271,14 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
|
||||
}
|
||||
gfc_current_locus = old_loc;
|
||||
}
|
||||
if ((mask & OMP_CLAUSE_IF_PRESENT)
|
||||
&& !c->if_present
|
||||
&& gfc_match ("if_present") == MATCH_YES)
|
||||
{
|
||||
c->if_present = true;
|
||||
needs_space = true;
|
||||
continue;
|
||||
}
|
||||
if ((mask & OMP_CLAUSE_INBRANCH)
|
||||
&& !c->inbranch
|
||||
&& !c->notinbranch
|
||||
@ -1503,22 +1508,22 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
|
||||
}
|
||||
break;
|
||||
case 'p':
|
||||
if ((mask & OMP_CLAUSE_PRESENT_OR_COPY)
|
||||
if ((mask & OMP_CLAUSE_COPY)
|
||||
&& gfc_match ("pcopy ( ") == MATCH_YES
|
||||
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
|
||||
OMP_MAP_TOFROM))
|
||||
continue;
|
||||
if ((mask & OMP_CLAUSE_PRESENT_OR_COPYIN)
|
||||
if ((mask & OMP_CLAUSE_COPYIN)
|
||||
&& gfc_match ("pcopyin ( ") == MATCH_YES
|
||||
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
|
||||
OMP_MAP_TO))
|
||||
continue;
|
||||
if ((mask & OMP_CLAUSE_PRESENT_OR_COPYOUT)
|
||||
if ((mask & OMP_CLAUSE_COPYOUT)
|
||||
&& gfc_match ("pcopyout ( ") == MATCH_YES
|
||||
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
|
||||
OMP_MAP_FROM))
|
||||
continue;
|
||||
if ((mask & OMP_CLAUSE_PRESENT_OR_CREATE)
|
||||
if ((mask & OMP_CLAUSE_CREATE)
|
||||
&& gfc_match ("pcreate ( ") == MATCH_YES
|
||||
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
|
||||
OMP_MAP_ALLOC))
|
||||
@ -1528,22 +1533,22 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
|
||||
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
|
||||
OMP_MAP_FORCE_PRESENT))
|
||||
continue;
|
||||
if ((mask & OMP_CLAUSE_PRESENT_OR_COPY)
|
||||
if ((mask & OMP_CLAUSE_COPY)
|
||||
&& gfc_match ("present_or_copy ( ") == MATCH_YES
|
||||
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
|
||||
OMP_MAP_TOFROM))
|
||||
continue;
|
||||
if ((mask & OMP_CLAUSE_PRESENT_OR_COPYIN)
|
||||
if ((mask & OMP_CLAUSE_COPYIN)
|
||||
&& gfc_match ("present_or_copyin ( ") == MATCH_YES
|
||||
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
|
||||
OMP_MAP_TO))
|
||||
continue;
|
||||
if ((mask & OMP_CLAUSE_PRESENT_OR_COPYOUT)
|
||||
if ((mask & OMP_CLAUSE_COPYOUT)
|
||||
&& gfc_match ("present_or_copyout ( ") == MATCH_YES
|
||||
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
|
||||
OMP_MAP_FROM))
|
||||
continue;
|
||||
if ((mask & OMP_CLAUSE_PRESENT_OR_CREATE)
|
||||
if ((mask & OMP_CLAUSE_CREATE)
|
||||
&& gfc_match ("present_or_create ( ") == MATCH_YES
|
||||
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
|
||||
OMP_MAP_ALLOC))
|
||||
@ -1925,23 +1930,19 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
|
||||
(omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_NUM_GANGS \
|
||||
| OMP_CLAUSE_NUM_WORKERS | OMP_CLAUSE_VECTOR_LENGTH | OMP_CLAUSE_REDUCTION \
|
||||
| OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \
|
||||
| OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_PRESENT_OR_COPY \
|
||||
| OMP_CLAUSE_PRESENT_OR_COPYIN | OMP_CLAUSE_PRESENT_OR_COPYOUT \
|
||||
| OMP_CLAUSE_PRESENT_OR_CREATE | OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_PRIVATE \
|
||||
| OMP_CLAUSE_FIRSTPRIVATE | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_WAIT)
|
||||
| OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_DEVICEPTR \
|
||||
| OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE | OMP_CLAUSE_DEFAULT \
|
||||
| OMP_CLAUSE_WAIT)
|
||||
#define OACC_KERNELS_CLAUSES \
|
||||
(omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_NUM_GANGS \
|
||||
| OMP_CLAUSE_NUM_WORKERS | OMP_CLAUSE_VECTOR_LENGTH | OMP_CLAUSE_DEVICEPTR \
|
||||
| OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \
|
||||
| OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_PRESENT_OR_COPY \
|
||||
| OMP_CLAUSE_PRESENT_OR_COPYIN | OMP_CLAUSE_PRESENT_OR_COPYOUT \
|
||||
| OMP_CLAUSE_PRESENT_OR_CREATE | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_WAIT)
|
||||
| OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_DEFAULT \
|
||||
| OMP_CLAUSE_WAIT)
|
||||
#define OACC_DATA_CLAUSES \
|
||||
(omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_COPY \
|
||||
| OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT | OMP_CLAUSE_CREATE \
|
||||
| OMP_CLAUSE_PRESENT | OMP_CLAUSE_PRESENT_OR_COPY \
|
||||
| OMP_CLAUSE_PRESENT_OR_COPYIN | OMP_CLAUSE_PRESENT_OR_COPYOUT \
|
||||
| OMP_CLAUSE_PRESENT_OR_CREATE)
|
||||
| OMP_CLAUSE_PRESENT)
|
||||
#define OACC_LOOP_CLAUSES \
|
||||
(omp_mask (OMP_CLAUSE_COLLAPSE) | OMP_CLAUSE_GANG | OMP_CLAUSE_WORKER \
|
||||
| OMP_CLAUSE_VECTOR | OMP_CLAUSE_SEQ | OMP_CLAUSE_INDEPENDENT \
|
||||
@ -1955,19 +1956,17 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
|
||||
#define OACC_DECLARE_CLAUSES \
|
||||
(omp_mask (OMP_CLAUSE_COPY) | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \
|
||||
| OMP_CLAUSE_CREATE | OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_DEVICE_RESIDENT \
|
||||
| OMP_CLAUSE_PRESENT | OMP_CLAUSE_PRESENT_OR_COPY \
|
||||
| OMP_CLAUSE_PRESENT_OR_COPYIN | OMP_CLAUSE_PRESENT_OR_COPYOUT \
|
||||
| OMP_CLAUSE_PRESENT_OR_CREATE | OMP_CLAUSE_LINK)
|
||||
| OMP_CLAUSE_PRESENT \
|
||||
| OMP_CLAUSE_LINK)
|
||||
#define OACC_UPDATE_CLAUSES \
|
||||
(omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_HOST_SELF \
|
||||
| OMP_CLAUSE_DEVICE | OMP_CLAUSE_WAIT)
|
||||
| OMP_CLAUSE_DEVICE | OMP_CLAUSE_WAIT | OMP_CLAUSE_IF_PRESENT)
|
||||
#define OACC_ENTER_DATA_CLAUSES \
|
||||
(omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_WAIT \
|
||||
| OMP_CLAUSE_COPYIN | OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT_OR_COPYIN \
|
||||
| OMP_CLAUSE_PRESENT_OR_CREATE)
|
||||
| OMP_CLAUSE_COPYIN | OMP_CLAUSE_CREATE)
|
||||
#define OACC_EXIT_DATA_CLAUSES \
|
||||
(omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_WAIT \
|
||||
| OMP_CLAUSE_COPYOUT | OMP_CLAUSE_DELETE)
|
||||
| OMP_CLAUSE_COPYOUT | OMP_CLAUSE_DELETE | OMP_CLAUSE_FINALIZE)
|
||||
#define OACC_WAIT_CLAUSES \
|
||||
omp_mask (OMP_CLAUSE_ASYNC)
|
||||
#define OACC_ROUTINE_CLAUSES \
|
||||
@ -2061,8 +2060,7 @@ gfc_match_oacc_declare (void)
|
||||
|
||||
if (s->ns->proc_name && s->ns->proc_name->attr.proc == PROC_MODULE)
|
||||
{
|
||||
if (n->u.map_op != OMP_MAP_FORCE_ALLOC
|
||||
&& n->u.map_op != OMP_MAP_FORCE_TO)
|
||||
if (n->u.map_op != OMP_MAP_ALLOC && n->u.map_op != OMP_MAP_TO)
|
||||
{
|
||||
gfc_error ("Invalid clause in module with !$ACC DECLARE at %L",
|
||||
&where);
|
||||
@ -2072,6 +2070,13 @@ gfc_match_oacc_declare (void)
|
||||
module_var = true;
|
||||
}
|
||||
|
||||
if (ns->proc_name->attr.oacc_function)
|
||||
{
|
||||
gfc_error ("Invalid declare in routine with $!ACC DECLARE at %L",
|
||||
&where);
|
||||
return MATCH_ERROR;
|
||||
}
|
||||
|
||||
if (s->attr.use_assoc)
|
||||
{
|
||||
gfc_error ("Variable is USE-associated with !$ACC DECLARE at %L",
|
||||
@ -2090,10 +2095,12 @@ gfc_match_oacc_declare (void)
|
||||
switch (n->u.map_op)
|
||||
{
|
||||
case OMP_MAP_FORCE_ALLOC:
|
||||
case OMP_MAP_ALLOC:
|
||||
s->attr.oacc_declare_create = 1;
|
||||
break;
|
||||
|
||||
case OMP_MAP_FORCE_TO:
|
||||
case OMP_MAP_TO:
|
||||
s->attr.oacc_declare_copyin = 1;
|
||||
break;
|
||||
|
||||
|
||||
@ -2895,6 +2895,16 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
|
||||
c = build_omp_clause (where.lb->location, OMP_CLAUSE_AUTO);
|
||||
omp_clauses = gfc_trans_add_clause (c, omp_clauses);
|
||||
}
|
||||
if (clauses->if_present)
|
||||
{
|
||||
c = build_omp_clause (where.lb->location, OMP_CLAUSE_IF_PRESENT);
|
||||
omp_clauses = gfc_trans_add_clause (c, omp_clauses);
|
||||
}
|
||||
if (clauses->finalize)
|
||||
{
|
||||
c = build_omp_clause (where.lb->location, OMP_CLAUSE_FINALIZE);
|
||||
omp_clauses = gfc_trans_add_clause (c, omp_clauses);
|
||||
}
|
||||
if (clauses->independent)
|
||||
{
|
||||
c = build_omp_clause (where.lb->location, OMP_CLAUSE_INDEPENDENT);
|
||||
|
||||
@ -8524,6 +8524,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
|
||||
case OMP_CLAUSE_NOGROUP:
|
||||
case OMP_CLAUSE_THREADS:
|
||||
case OMP_CLAUSE_SIMD:
|
||||
case OMP_CLAUSE_IF_PRESENT:
|
||||
case OMP_CLAUSE_FINALIZE:
|
||||
break;
|
||||
|
||||
case OMP_CLAUSE_DEFAULTMAP:
|
||||
@ -9305,6 +9307,8 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
|
||||
case OMP_CLAUSE_AUTO:
|
||||
case OMP_CLAUSE_SEQ:
|
||||
case OMP_CLAUSE_TILE:
|
||||
case OMP_CLAUSE_IF_PRESENT:
|
||||
case OMP_CLAUSE_FINALIZE:
|
||||
break;
|
||||
|
||||
default:
|
||||
@ -9361,21 +9365,7 @@ gimplify_oacc_declare_1 (tree clause)
|
||||
switch (kind)
|
||||
{
|
||||
case GOMP_MAP_ALLOC:
|
||||
case GOMP_MAP_FORCE_ALLOC:
|
||||
case GOMP_MAP_FORCE_TO:
|
||||
new_op = GOMP_MAP_DELETE;
|
||||
ret = true;
|
||||
break;
|
||||
|
||||
case GOMP_MAP_FORCE_FROM:
|
||||
OMP_CLAUSE_SET_MAP_KIND (clause, GOMP_MAP_FORCE_ALLOC);
|
||||
new_op = GOMP_MAP_FORCE_FROM;
|
||||
ret = true;
|
||||
break;
|
||||
|
||||
case GOMP_MAP_FORCE_TOFROM:
|
||||
OMP_CLAUSE_SET_MAP_KIND (clause, GOMP_MAP_FORCE_TO);
|
||||
new_op = GOMP_MAP_FORCE_FROM;
|
||||
new_op = GOMP_MAP_RELEASE;
|
||||
ret = true;
|
||||
break;
|
||||
|
||||
@ -10817,6 +10807,53 @@ gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p)
|
||||
ort, TREE_CODE (expr));
|
||||
gimplify_adjust_omp_clauses (pre_p, NULL, &OMP_STANDALONE_CLAUSES (expr),
|
||||
TREE_CODE (expr));
|
||||
if (TREE_CODE (expr) == OACC_UPDATE
|
||||
&& omp_find_clause (OMP_STANDALONE_CLAUSES (expr),
|
||||
OMP_CLAUSE_IF_PRESENT))
|
||||
{
|
||||
/* The runtime uses GOMP_MAP_{TO,FROM} to denote the if_present
|
||||
clause. */
|
||||
for (tree c = OMP_STANDALONE_CLAUSES (expr); c; c = OMP_CLAUSE_CHAIN (c))
|
||||
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)
|
||||
switch (OMP_CLAUSE_MAP_KIND (c))
|
||||
{
|
||||
case GOMP_MAP_FORCE_TO:
|
||||
OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TO);
|
||||
break;
|
||||
case GOMP_MAP_FORCE_FROM:
|
||||
OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FROM);
|
||||
break;
|
||||
default:
|
||||
break;
|
||||
}
|
||||
}
|
||||
else if (TREE_CODE (expr) == OACC_EXIT_DATA
|
||||
&& omp_find_clause (OMP_STANDALONE_CLAUSES (expr),
|
||||
OMP_CLAUSE_FINALIZE))
|
||||
{
|
||||
/* Use GOMP_MAP_DELETE/GOMP_MAP_FORCE_FROM to denote that "finalize"
|
||||
semantics apply to all mappings of this OpenACC directive. */
|
||||
bool finalize_marked = false;
|
||||
for (tree c = OMP_STANDALONE_CLAUSES (expr); c; c = OMP_CLAUSE_CHAIN (c))
|
||||
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)
|
||||
switch (OMP_CLAUSE_MAP_KIND (c))
|
||||
{
|
||||
case GOMP_MAP_FROM:
|
||||
OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_FROM);
|
||||
finalize_marked = true;
|
||||
break;
|
||||
case GOMP_MAP_RELEASE:
|
||||
OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_DELETE);
|
||||
finalize_marked = true;
|
||||
break;
|
||||
default:
|
||||
/* Check consistency: libgomp relies on the very first data
|
||||
mapping clause being marked, so make sure we did that before
|
||||
any other mapping clauses. */
|
||||
gcc_assert (finalize_marked);
|
||||
break;
|
||||
}
|
||||
}
|
||||
stmt = gimple_build_omp_target (NULL, kind, OMP_STANDALONE_CLAUSES (expr));
|
||||
|
||||
gimplify_seq_add_stmt (pre_p, stmt);
|
||||
|
||||
@ -642,8 +642,7 @@ build_sender_ref (tree var, omp_context *ctx)
|
||||
BASE_POINTERS_RESTRICT, declare the field with restrict. */
|
||||
|
||||
static void
|
||||
install_var_field (tree var, bool by_ref, int mask, omp_context *ctx,
|
||||
bool base_pointers_restrict = false)
|
||||
install_var_field (tree var, bool by_ref, int mask, omp_context *ctx)
|
||||
{
|
||||
tree field, type, sfield = NULL_TREE;
|
||||
splay_tree_key key = (splay_tree_key) var;
|
||||
@ -674,11 +673,7 @@ install_var_field (tree var, bool by_ref, int mask, omp_context *ctx,
|
||||
type = build_pointer_type (build_pointer_type (type));
|
||||
}
|
||||
else if (by_ref)
|
||||
{
|
||||
type = build_pointer_type (type);
|
||||
if (base_pointers_restrict)
|
||||
type = build_qualified_type (type, TYPE_QUAL_RESTRICT);
|
||||
}
|
||||
type = build_pointer_type (type);
|
||||
else if ((mask & 3) == 1 && omp_is_reference (var))
|
||||
type = TREE_TYPE (type);
|
||||
|
||||
@ -992,12 +987,10 @@ fixup_child_record_type (omp_context *ctx)
|
||||
}
|
||||
|
||||
/* Instantiate decls as necessary in CTX to satisfy the data sharing
|
||||
specified by CLAUSES. If BASE_POINTERS_RESTRICT, install var field with
|
||||
restrict. */
|
||||
specified by CLAUSES. */
|
||||
|
||||
static void
|
||||
scan_sharing_clauses (tree clauses, omp_context *ctx,
|
||||
bool base_pointers_restrict = false)
|
||||
scan_sharing_clauses (tree clauses, omp_context *ctx)
|
||||
{
|
||||
tree c, decl;
|
||||
bool scan_array_reductions = false;
|
||||
@ -1256,8 +1249,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx,
|
||||
&& TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
|
||||
install_var_field (decl, true, 7, ctx);
|
||||
else
|
||||
install_var_field (decl, true, 3, ctx,
|
||||
base_pointers_restrict);
|
||||
install_var_field (decl, true, 3, ctx);
|
||||
if (is_gimple_omp_offloaded (ctx->stmt)
|
||||
&& !OMP_CLAUSE_MAP_IN_REDUCTION (c))
|
||||
install_var_local (decl, ctx);
|
||||
@ -1328,6 +1320,8 @@ scan_sharing_clauses (tree clauses, omp_context *ctx,
|
||||
case OMP_CLAUSE_TILE:
|
||||
case OMP_CLAUSE__SIMT_:
|
||||
case OMP_CLAUSE_DEFAULT:
|
||||
case OMP_CLAUSE_IF_PRESENT:
|
||||
case OMP_CLAUSE_FINALIZE:
|
||||
break;
|
||||
|
||||
case OMP_CLAUSE_ALIGNED:
|
||||
@ -1499,6 +1493,8 @@ scan_sharing_clauses (tree clauses, omp_context *ctx,
|
||||
case OMP_CLAUSE_TILE:
|
||||
case OMP_CLAUSE__GRIDDIM_:
|
||||
case OMP_CLAUSE__SIMT_:
|
||||
case OMP_CLAUSE_IF_PRESENT:
|
||||
case OMP_CLAUSE_FINALIZE:
|
||||
break;
|
||||
|
||||
case OMP_CLAUSE__CACHE_:
|
||||
@ -2266,68 +2262,6 @@ scan_omp_single (gomp_single *stmt, omp_context *outer_ctx)
|
||||
layout_type (ctx->record_type);
|
||||
}
|
||||
|
||||
/* Return true if the CLAUSES of an omp target guarantee that the base pointers
|
||||
used in the corresponding offloaded function are restrict. */
|
||||
|
||||
static bool
|
||||
omp_target_base_pointers_restrict_p (tree clauses)
|
||||
{
|
||||
/* The analysis relies on the GOMP_MAP_FORCE_* mapping kinds, which are only
|
||||
used by OpenACC. */
|
||||
if (flag_openacc == 0)
|
||||
return false;
|
||||
|
||||
/* I. Basic example:
|
||||
|
||||
void foo (void)
|
||||
{
|
||||
unsigned int a[2], b[2];
|
||||
|
||||
#pragma acc kernels \
|
||||
copyout (a) \
|
||||
copyout (b)
|
||||
{
|
||||
a[0] = 0;
|
||||
b[0] = 1;
|
||||
}
|
||||
}
|
||||
|
||||
After gimplification, we have:
|
||||
|
||||
#pragma omp target oacc_kernels \
|
||||
map(force_from:a [len: 8]) \
|
||||
map(force_from:b [len: 8])
|
||||
{
|
||||
a[0] = 0;
|
||||
b[0] = 1;
|
||||
}
|
||||
|
||||
Because both mappings have the force prefix, we know that they will be
|
||||
allocated when calling the corresponding offloaded function, which means we
|
||||
can mark the base pointers for a and b in the offloaded function as
|
||||
restrict. */
|
||||
|
||||
tree c;
|
||||
for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
|
||||
{
|
||||
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
|
||||
return false;
|
||||
|
||||
switch (OMP_CLAUSE_MAP_KIND (c))
|
||||
{
|
||||
case GOMP_MAP_FORCE_ALLOC:
|
||||
case GOMP_MAP_FORCE_TO:
|
||||
case GOMP_MAP_FORCE_FROM:
|
||||
case GOMP_MAP_FORCE_TOFROM:
|
||||
break;
|
||||
default:
|
||||
return false;
|
||||
}
|
||||
}
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
/* Scan a GIMPLE_OMP_TARGET. */
|
||||
|
||||
static void
|
||||
@ -2349,20 +2283,13 @@ scan_omp_target (gomp_target *stmt, omp_context *outer_ctx)
|
||||
TYPE_NAME (ctx->record_type) = name;
|
||||
TYPE_ARTIFICIAL (ctx->record_type) = 1;
|
||||
|
||||
bool base_pointers_restrict = false;
|
||||
if (offloaded)
|
||||
{
|
||||
create_omp_child_function (ctx, false);
|
||||
gimple_omp_target_set_child_fn (stmt, ctx->cb.dst_fn);
|
||||
|
||||
base_pointers_restrict = omp_target_base_pointers_restrict_p (clauses);
|
||||
if (base_pointers_restrict
|
||||
&& dump_file && (dump_flags & TDF_DETAILS))
|
||||
fprintf (dump_file,
|
||||
"Base pointers in offloaded function are restrict\n");
|
||||
}
|
||||
|
||||
scan_sharing_clauses (clauses, ctx, base_pointers_restrict);
|
||||
scan_sharing_clauses (clauses, ctx);
|
||||
scan_omp (gimple_omp_body_ptr (stmt), ctx);
|
||||
|
||||
if (TYPE_FIELDS (ctx->record_type) == NULL)
|
||||
|
||||
@ -1,3 +1,36 @@
|
||||
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
|
||||
|
||||
@ -19,6 +19,12 @@ int v4;
|
||||
int v5, v6, v7, v8;
|
||||
#pragma acc declare create(v5, v6) copyin(v7, v8)
|
||||
|
||||
int v9;
|
||||
#pragma acc declare present_or_copyin(v9)
|
||||
|
||||
int v10;
|
||||
#pragma acc declare present_or_create(v10)
|
||||
|
||||
void
|
||||
f (void)
|
||||
{
|
||||
@ -49,6 +55,12 @@ f (void)
|
||||
extern int ve4;
|
||||
#pragma acc declare link(ve4)
|
||||
|
||||
extern int ve5;
|
||||
#pragma acc declare present_or_copyin(ve5)
|
||||
|
||||
extern int ve6;
|
||||
#pragma acc declare present_or_create(ve6)
|
||||
|
||||
int va5;
|
||||
#pragma acc declare copy(va5)
|
||||
|
||||
|
||||
@ -29,13 +29,7 @@ int v6;
|
||||
#pragma acc declare present_or_copy(v6) /* { dg-error "at file scope" } */
|
||||
|
||||
int v7;
|
||||
#pragma acc declare present_or_copyin(v7) /* { dg-error "at file scope" } */
|
||||
|
||||
int v8;
|
||||
#pragma acc declare present_or_copyout(v8) /* { dg-error "at file scope" } */
|
||||
|
||||
int v9;
|
||||
#pragma acc declare present_or_create(v9) /* { dg-error "at file scope" } */
|
||||
#pragma acc declare present_or_copyout(v7) /* { dg-error "at file scope" } */
|
||||
|
||||
int va10;
|
||||
#pragma acc declare create (va10)
|
||||
@ -67,13 +61,7 @@ f (void)
|
||||
#pragma acc declare present_or_copy(ve3) /* { dg-error "invalid use of" } */
|
||||
|
||||
extern int ve4;
|
||||
#pragma acc declare present_or_copyin(ve4) /* { dg-error "invalid use of" } */
|
||||
#pragma acc declare present_or_copyout(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 (v9) /* { dg-error "invalid use of" } */
|
||||
#pragma acc declare present (v2) /* { dg-error "invalid use of" } */
|
||||
}
|
||||
|
||||
@ -8,7 +8,7 @@ void f1 ()
|
||||
float f1_b[2];
|
||||
|
||||
#pragma acc data copyin (f1_a) copyout (f1_b)
|
||||
/* { dg-final { scan-tree-dump-times "omp target oacc_data map\\(force_from:f1_b \[^\\)\]+\\) map\\(force_to:f1_a" 1 "gimple" } } */
|
||||
/* { dg-final { scan-tree-dump-times "omp target oacc_data map\\(from:f1_b \[^\\)\]+\\) map\\(to:f1_a" 1 "gimple" } } */
|
||||
{
|
||||
#pragma acc kernels
|
||||
/* { dg-final { scan-tree-dump-times "omp target oacc_kernels map\\(tofrom:f1_b \[^\\)\]+\\) map\\(tofrom:f1_a" 1 "gimple" } } */
|
||||
@ -29,7 +29,7 @@ void f2 ()
|
||||
float f2_b[2];
|
||||
|
||||
#pragma acc data copyin (f2_a) copyout (f2_b)
|
||||
/* { dg-final { scan-tree-dump-times "omp target oacc_data map\\(force_from:f2_b \[^\\)\]+\\) map\\(force_to:f2_a" 1 "gimple" } } */
|
||||
/* { dg-final { scan-tree-dump-times "omp target oacc_data map\\(from:f2_b \[^\\)\]+\\) map\\(to:f2_a" 1 "gimple" } } */
|
||||
{
|
||||
#pragma acc kernels default (none)
|
||||
/* { dg-final { scan-tree-dump-times "omp target oacc_kernels default\\(none\\) map\\(tofrom:f2_b \[^\\)\]+\\) map\\(tofrom:f2_a" 1 "gimple" } } */
|
||||
@ -50,7 +50,7 @@ void f3 ()
|
||||
float f3_b[2];
|
||||
|
||||
#pragma acc data copyin (f3_a) copyout (f3_b)
|
||||
/* { dg-final { scan-tree-dump-times "omp target oacc_data map\\(force_from:f3_b \[^\\)\]+\\) map\\(force_to:f3_a" 1 "gimple" } } */
|
||||
/* { dg-final { scan-tree-dump-times "omp target oacc_data map\\(from:f3_b \[^\\)\]+\\) map\\(to:f3_a" 1 "gimple" } } */
|
||||
{
|
||||
#pragma acc kernels default (present)
|
||||
/* { dg-final { scan-tree-dump-times "omp target oacc_kernels default\\(present\\) map\\(tofrom:f3_b \[^\\)\]+\\) map\\(tofrom:f3_a" 1 "gimple" } } */
|
||||
|
||||
28
gcc/testsuite/c-c++-common/goacc/finalize-1.c
Normal file
28
gcc/testsuite/c-c++-common/goacc/finalize-1.c
Normal file
@ -0,0 +1,28 @@
|
||||
/* Test valid usage and processing of the finalize clause. */
|
||||
|
||||
/* { dg-additional-options "-fdump-tree-original -fdump-tree-gimple" } */
|
||||
|
||||
extern int del_r;
|
||||
extern float del_f[3];
|
||||
extern double cpo_r[8];
|
||||
extern long cpo_f;
|
||||
|
||||
void f ()
|
||||
{
|
||||
#pragma acc exit data delete (del_r)
|
||||
/* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:del_r\\);$" 1 "original" } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(release:del_r \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } } */
|
||||
|
||||
#pragma acc exit data finalize delete (del_f)
|
||||
/* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:del_f\\) finalize;$" 1 "original" } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:del_f \\\[len: \[0-9\]+\\\]\\) finalize$" 1 "gimple" } } */
|
||||
|
||||
#pragma acc exit data copyout (cpo_r)
|
||||
/* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:cpo_r\\);$" 1 "original" } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(from:cpo_r \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } } */
|
||||
|
||||
#pragma acc exit data copyout (cpo_f) finalize
|
||||
/* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data finalize map\\(from:cpo_f\\);$" 1 "original" } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data finalize map\\(force_from:cpo_f \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } } */
|
||||
}
|
||||
|
||||
@ -18,10 +18,12 @@ foo (void)
|
||||
}
|
||||
}
|
||||
|
||||
/* The xfails occur due to the OpenACC 2.5 data semantics. */
|
||||
|
||||
/* { dg-final { scan-tree-dump-times "clique 1 base 1" 4 "ealias" } } */
|
||||
/* { dg-final { scan-tree-dump-times "clique 1 base 2" 1 "ealias" } } */
|
||||
/* { dg-final { scan-tree-dump-times "clique 1 base 3" 1 "ealias" } } */
|
||||
/* { dg-final { scan-tree-dump-times "clique 1 base 4" 1 "ealias" } } */
|
||||
/* { dg-final { scan-tree-dump-times "clique 1 base 5" 1 "ealias" } } */
|
||||
/* { dg-final { scan-tree-dump-times "clique 1 base 2" 1 "ealias" { xfail *-*-* } } } */
|
||||
/* { dg-final { scan-tree-dump-times "clique 1 base 3" 1 "ealias" { xfail *-*-* } } } */
|
||||
/* { dg-final { scan-tree-dump-times "clique 1 base 4" 1 "ealias" { xfail *-*-* } } } */
|
||||
/* { dg-final { scan-tree-dump-times "clique 1 base 5" 1 "ealias" { xfail *-*-* } } } */
|
||||
/* { dg-final { scan-tree-dump-times "(?n)clique .* base .*" 8 "ealias" } } */
|
||||
|
||||
|
||||
@ -20,10 +20,12 @@ foo (void)
|
||||
}
|
||||
}
|
||||
|
||||
/* The xfails occur due to the OpenACC 2.5 data semantics. */
|
||||
|
||||
/* { dg-final { scan-tree-dump-times "clique 1 base 1" 4 "ealias" } } */
|
||||
/* { dg-final { scan-tree-dump-times "clique 1 base 2" 1 "ealias" } } */
|
||||
/* { dg-final { scan-tree-dump-times "clique 1 base 3" 1 "ealias" } } */
|
||||
/* { dg-final { scan-tree-dump-times "clique 1 base 4" 1 "ealias" } } */
|
||||
/* { dg-final { scan-tree-dump-times "clique 1 base 5" 1 "ealias" } } */
|
||||
/* { dg-final { scan-tree-dump-times "clique 1 base 2" 1 "ealias" { xfail *-*-* } } } */
|
||||
/* { dg-final { scan-tree-dump-times "clique 1 base 3" 1 "ealias" { xfail *-*-* } } } */
|
||||
/* { dg-final { scan-tree-dump-times "clique 1 base 4" 1 "ealias" { xfail *-*-* } } } */
|
||||
/* { dg-final { scan-tree-dump-times "clique 1 base 5" 1 "ealias" { xfail *-*-* } } } */
|
||||
/* { dg-final { scan-tree-dump-times "(?n)clique .* base .*" 8 "ealias" } } */
|
||||
|
||||
|
||||
@ -4,11 +4,11 @@
|
||||
|
||||
struct PC
|
||||
{
|
||||
#pragma acc routine /* { dg-error ".#pragma acc routine. must be at file scope" } */
|
||||
#pragma acc routine seq /* { dg-error ".#pragma acc routine. must be at file scope" } */
|
||||
};
|
||||
|
||||
void PC1( /* { dg-bogus "variable or field .PC1. declared void" "TODO" { xfail c++ } } */
|
||||
#pragma acc routine
|
||||
#pragma acc routine seq
|
||||
/* { dg-error ".#pragma acc routine. must be at file scope" "" { target c } .-1 }
|
||||
{ dg-error ".#pragma. is not allowed here" "" { target c++ } .-2 } */
|
||||
) /* { dg-bogus "expected declaration specifiers or .\\.\\.\\.. before .\\). token" "TODO" { xfail c } } */
|
||||
@ -18,26 +18,26 @@ void PC1( /* { dg-bogus "variable or field .PC1. declared void" "TODO" { xfail c
|
||||
void PC2()
|
||||
{
|
||||
if (0)
|
||||
#pragma acc routine /* { dg-error ".#pragma acc routine. must be at file scope" } */
|
||||
#pragma acc routine seq /* { dg-error ".#pragma acc routine. must be at file scope" } */
|
||||
;
|
||||
}
|
||||
|
||||
void PC3()
|
||||
{
|
||||
#pragma acc routine /* { dg-error ".#pragma acc routine. must be at file scope" } */
|
||||
#pragma acc routine seq /* { dg-error ".#pragma acc routine. must be at file scope" } */
|
||||
}
|
||||
|
||||
|
||||
/* "( name )" syntax. */
|
||||
|
||||
#pragma acc routine ( /* { dg-error "expected (function name|unqualified-id) before end of line" } */
|
||||
#pragma acc routine () /* { dg-error "expected (function name|unqualified-id) before .\\). token" } */
|
||||
#pragma acc routine (+) /* { dg-error "expected (function name|unqualified-id) before .\\+. token" } */
|
||||
#pragma acc routine (?) /* { dg-error "expected (function name|unqualified-id) before .\\?. token" } */
|
||||
#pragma acc routine (:) /* { dg-error "expected (function name|unqualified-id) before .:. token" } */
|
||||
#pragma acc routine (4) /* { dg-error "expected (function name|unqualified-id) before numeric constant" } */
|
||||
#pragma acc routine () seq /* { dg-error "expected (function name|unqualified-id) before .\\). token" } */
|
||||
#pragma acc routine (+) seq /* { dg-error "expected (function name|unqualified-id) before .\\+. token" } */
|
||||
#pragma acc routine (?) seq /* { dg-error "expected (function name|unqualified-id) before .\\?. token" } */
|
||||
#pragma acc routine (:) seq /* { dg-error "expected (function name|unqualified-id) before .:. token" } */
|
||||
#pragma acc routine (4) seq /* { dg-error "expected (function name|unqualified-id) before numeric constant" } */
|
||||
#pragma acc routine ('4') /* { dg-error "expected (function name|unqualified-id) before .4." } */
|
||||
#pragma acc routine ("4") /* { dg-error "expected (function name|unqualified-id) before string constant" } */
|
||||
#pragma acc routine ("4") seq /* { dg-error "expected (function name|unqualified-id) before string constant" } */
|
||||
extern void R1(void);
|
||||
extern void R2(void);
|
||||
#pragma acc routine (R1, R2, R3) worker /* { dg-error "expected .\\). before .,. token" } */
|
||||
@ -49,84 +49,84 @@ extern void R2(void);
|
||||
/* "#pragma acc routine" not immediately followed by (a single) function
|
||||
declaration or definition. */
|
||||
|
||||
#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
|
||||
#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
|
||||
int a;
|
||||
|
||||
#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by a single function declaration or definition" } */
|
||||
#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by a single function declaration or definition" } */
|
||||
void fn1 (void), fn1b (void);
|
||||
|
||||
#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
|
||||
#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
|
||||
int b, fn2 (void);
|
||||
|
||||
#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
|
||||
#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
|
||||
int b_, fn2_ (void), B_;
|
||||
|
||||
#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by a single function declaration or definition" } */
|
||||
#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by a single function declaration or definition" } */
|
||||
int fn3 (void), b2;
|
||||
|
||||
#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
|
||||
#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
|
||||
typedef struct c c;
|
||||
|
||||
#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
|
||||
#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
|
||||
struct d {} d;
|
||||
|
||||
#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
|
||||
#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by a single function declaration or definition" } */
|
||||
#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
|
||||
#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by a single function declaration or definition" } */
|
||||
void fn1_2 (void), fn1b_2 (void);
|
||||
|
||||
#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
|
||||
#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
|
||||
#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
|
||||
int b_2, fn2_2 (void);
|
||||
|
||||
#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
|
||||
#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
|
||||
#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
|
||||
#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
|
||||
int b_2_, fn2_2_ (void), B_2_;
|
||||
|
||||
#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
|
||||
#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by a single function declaration or definition" } */
|
||||
#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
|
||||
#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by a single function declaration or definition" } */
|
||||
int fn3_2 (void), b2_2;
|
||||
|
||||
#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
|
||||
#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
|
||||
#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
|
||||
#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
|
||||
typedef struct c_2 c_2;
|
||||
|
||||
#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
|
||||
#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
|
||||
#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
|
||||
#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
|
||||
struct d_2 {} d_2;
|
||||
|
||||
#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
|
||||
#pragma acc routine
|
||||
#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
|
||||
#pragma acc routine seq
|
||||
int fn4 (void);
|
||||
|
||||
int fn5a (void);
|
||||
int fn5b (void);
|
||||
#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
|
||||
#pragma acc routine (fn5a)
|
||||
#pragma acc routine (fn5b)
|
||||
#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
|
||||
#pragma acc routine (fn5a) seq
|
||||
#pragma acc routine (fn5b) seq
|
||||
int fn5 (void);
|
||||
|
||||
#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
|
||||
#pragma acc routine (fn6a) /* { dg-error ".fn6a. has not been declared" } */
|
||||
#pragma acc routine (fn6b) /* { dg-error ".fn6b. has not been declared" } */
|
||||
#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
|
||||
#pragma acc routine (fn6a) seq /* { dg-error ".fn6a. has not been declared" } */
|
||||
#pragma acc routine (fn6b) seq /* { dg-error ".fn6b. has not been declared" } */
|
||||
int fn6 (void);
|
||||
|
||||
#ifdef __cplusplus
|
||||
|
||||
#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" "" { target c++ } } */
|
||||
#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" "" { target c++ } } */
|
||||
namespace f {}
|
||||
|
||||
namespace g {}
|
||||
|
||||
#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" "" { target c++ } } */
|
||||
#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" "" { target c++ } } */
|
||||
using namespace g;
|
||||
|
||||
#pragma acc routine (g) /* { dg-error ".g. does not refer to a function" "" { target c++ } } */
|
||||
#pragma acc routine (g) seq /* { dg-error ".g. does not refer to a function" "" { target c++ } } */
|
||||
|
||||
#endif /* __cplusplus */
|
||||
|
||||
#pragma acc routine (a) /* { dg-error ".a. does not refer to a function" } */
|
||||
#pragma acc routine (a) seq /* { dg-error ".a. does not refer to a function" } */
|
||||
|
||||
#pragma acc routine (c) /* { dg-error ".c. does not refer to a function" } */
|
||||
#pragma acc routine (c) seq /* { dg-error ".c. does not refer to a function" } */
|
||||
|
||||
|
||||
/* Static assert. */
|
||||
@ -143,66 +143,24 @@ static_assert(0, ""); /* { dg-error "static assertion failed" "" { target c++11
|
||||
#endif
|
||||
void f_static_assert();
|
||||
/* Check that we already recognized "f_static_assert" as an OpenACC routine. */
|
||||
#pragma acc routine (f_static_assert) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*f_static_assert" "TODO" { xfail *-*-* } } */
|
||||
#pragma acc routine (f_static_assert) seq /* { dg-error ".#pragma acc routine. already applied to .\[void \]*f_static_assert" "TODO" { xfail *-*-* } } */
|
||||
|
||||
|
||||
/* __extension__ usage. */
|
||||
|
||||
#pragma acc routine
|
||||
#pragma acc routine seq
|
||||
__extension__ extern void ex1();
|
||||
#pragma acc routine (ex1) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*ex1" } */
|
||||
|
||||
#pragma acc routine
|
||||
#pragma acc routine seq
|
||||
__extension__ __extension__ __extension__ __extension__ __extension__ void ex2()
|
||||
{
|
||||
}
|
||||
#pragma acc routine (ex2) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*ex2" } */
|
||||
|
||||
#pragma acc routine /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
|
||||
#pragma acc routine seq /* { dg-error ".#pragma acc routine. not immediately followed by function declaration or definition" } */
|
||||
__extension__ int ex3;
|
||||
#pragma acc routine (ex3) /* { dg-error ".ex3. does not refer to a function" } */
|
||||
|
||||
|
||||
/* "#pragma acc routine" already applied. */
|
||||
|
||||
extern void fungsi_1();
|
||||
#pragma acc routine(fungsi_1) gang
|
||||
#pragma acc routine(fungsi_1) gang /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_1" } */
|
||||
#pragma acc routine(fungsi_1) worker /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_1" } */
|
||||
#pragma acc routine(fungsi_1) vector /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_1" } */
|
||||
|
||||
#pragma acc routine seq
|
||||
extern void fungsi_2();
|
||||
#pragma acc routine(fungsi_2) seq /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_2." } */
|
||||
#pragma acc routine(fungsi_2) worker /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_2." } */
|
||||
#pragma acc routine(fungsi_2) /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_2." } */
|
||||
|
||||
#pragma acc routine vector
|
||||
extern void fungsi_3();
|
||||
#pragma acc routine vector /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_3." } */
|
||||
void fungsi_3()
|
||||
{
|
||||
}
|
||||
|
||||
extern void fungsi_4();
|
||||
#pragma acc routine (fungsi_4) worker
|
||||
#pragma acc routine gang /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_4." } */
|
||||
void fungsi_4()
|
||||
{
|
||||
}
|
||||
|
||||
#pragma acc routine gang
|
||||
void fungsi_5()
|
||||
{
|
||||
}
|
||||
#pragma acc routine (fungsi_5) worker /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_5." } */
|
||||
|
||||
#pragma acc routine seq
|
||||
void fungsi_6()
|
||||
{
|
||||
}
|
||||
#pragma acc routine seq /* { dg-error ".#pragma acc routine. already applied to .\[void \]*fungsi_6." } */
|
||||
extern void fungsi_6();
|
||||
#pragma acc routine (ex3) seq /* { dg-error ".ex3. does not refer to a function" } */
|
||||
|
||||
|
||||
/* "#pragma acc routine" must be applied before. */
|
||||
@ -214,11 +172,11 @@ void Foo ()
|
||||
Bar ();
|
||||
}
|
||||
|
||||
#pragma acc routine (Bar) // { dg-error ".#pragma acc routine. must be applied before use" }
|
||||
#pragma acc routine (Bar) seq // { dg-error ".#pragma acc routine. must be applied before use" }
|
||||
|
||||
#pragma acc routine (Foo) gang // { dg-error ".#pragma acc routine. must be applied before definition" }
|
||||
|
||||
#pragma acc routine (Baz) // { dg-error "not been declared" }
|
||||
#pragma acc routine (Baz) seq // { dg-error "not been declared" }
|
||||
|
||||
|
||||
/* OpenACC declare. */
|
||||
@ -227,7 +185,7 @@ int vb1; /* { dg-error "directive for use" } */
|
||||
extern int vb2; /* { dg-error "directive for use" } */
|
||||
static int vb3; /* { dg-error "directive for use" } */
|
||||
|
||||
#pragma acc routine
|
||||
#pragma acc routine seq
|
||||
int
|
||||
func1 (int a)
|
||||
{
|
||||
@ -238,7 +196,7 @@ func1 (int a)
|
||||
return vb3;
|
||||
}
|
||||
|
||||
#pragma acc routine
|
||||
#pragma acc routine seq
|
||||
int
|
||||
func2 (int a)
|
||||
{
|
||||
@ -256,7 +214,7 @@ extern int vb6; /* { dg-error "clause used in" } */
|
||||
static int vb7; /* { dg-error "clause used in" } */
|
||||
#pragma acc declare link (vb7)
|
||||
|
||||
#pragma acc routine
|
||||
#pragma acc routine seq
|
||||
int
|
||||
func3 (int a)
|
||||
{
|
||||
@ -273,7 +231,7 @@ extern int vb9;
|
||||
static int vb10;
|
||||
#pragma acc declare create (vb10)
|
||||
|
||||
#pragma acc routine
|
||||
#pragma acc routine seq
|
||||
int
|
||||
func4 (int a)
|
||||
{
|
||||
@ -291,7 +249,7 @@ extern int vb12;
|
||||
extern int vb13;
|
||||
#pragma acc declare device_resident (vb13)
|
||||
|
||||
#pragma acc routine
|
||||
#pragma acc routine seq
|
||||
int
|
||||
func5 (int a)
|
||||
{
|
||||
@ -302,7 +260,7 @@ func5 (int a)
|
||||
return vb13;
|
||||
}
|
||||
|
||||
#pragma acc routine
|
||||
#pragma acc routine seq
|
||||
int
|
||||
func6 (int a)
|
||||
{
|
||||
|
||||
28
gcc/testsuite/c-c++-common/goacc/update-if_present-1.c
Normal file
28
gcc/testsuite/c-c++-common/goacc/update-if_present-1.c
Normal file
@ -0,0 +1,28 @@
|
||||
/* Test valid usages of the if_present clause. */
|
||||
|
||||
/* { dg-additional-options "-fdump-tree-omplower" } */
|
||||
|
||||
void
|
||||
t ()
|
||||
{
|
||||
int a, b, c[10];
|
||||
|
||||
#pragma acc update self(a) if_present
|
||||
#pragma acc update device(b) async if_present
|
||||
#pragma acc update host(c[1:3]) wait(4) if_present
|
||||
#pragma acc update self(c) device(b) host (a) async(10) if (a == 5) if_present
|
||||
|
||||
#pragma acc update self(a)
|
||||
#pragma acc update device(b) async
|
||||
#pragma acc update host(c[1:3]) wait(4)
|
||||
#pragma acc update self(c) device(b) host (a) async(10) if (a == 5)
|
||||
}
|
||||
|
||||
/* { dg-final { scan-tree-dump-times "omp target oacc_update if_present map.from:a .len: 4.." 1 "omplower" } } */
|
||||
/* { dg-final { scan-tree-dump-times "omp target oacc_update if_present async.-1. map.to:b .len: 4.." 1 "omplower" } } */
|
||||
/* { dg-final { scan-tree-dump-times "omp target oacc_update if_present wait.4. map.from:c.1. .len: 12.." 1 "omplower" } } */
|
||||
/* { dg-final { scan-tree-dump-times "omp target oacc_update if_present if.... async.10. map.from:a .len: 4.. map.to:b .len: 4.. map.from:c .len: 40.." 1 "omplower" } } */
|
||||
/* { dg-final { scan-tree-dump-times "omp target oacc_update map.force_from:a .len: 4.." 1 "omplower" } } */
|
||||
/* { dg-final { scan-tree-dump-times "omp target oacc_update async.-1. map.force_to:b .len: 4.." 1 "omplower" } } */
|
||||
/* { dg-final { scan-tree-dump-times "omp target oacc_update wait.4. map.force_from:c.1. .len: 12.." 1 "omplower" } } */
|
||||
/* { dg-final { scan-tree-dump-times "omp target oacc_update if.... async.10. map.force_from:a .len: 4.. map.force_to:b .len: 4.. map.force_from:c .len: 40.." 1 "omplower" } } */
|
||||
42
gcc/testsuite/c-c++-common/goacc/update-if_present-2.c
Normal file
42
gcc/testsuite/c-c++-common/goacc/update-if_present-2.c
Normal file
@ -0,0 +1,42 @@
|
||||
/* Test invalid usages of the if_present clause. */
|
||||
|
||||
#pragma acc routine gang if_present /* { dg-error "'if_present' is not valid" } */
|
||||
void
|
||||
t1 ()
|
||||
{
|
||||
int a, b, c[10];
|
||||
|
||||
#pragma acc enter data copyin(a) if_present /* { dg-error "'if_present' is not valid" } */
|
||||
#pragma acc exit data copyout(a) if_present /* { dg-error "'if_present' is not valid" } */
|
||||
|
||||
#pragma acc data copy(a) if_present /* { dg-error "'if_present' is not valid" } */
|
||||
{
|
||||
}
|
||||
|
||||
#pragma acc declare create(c) if_present /* { dg-error "'if_present' is not valid" } */
|
||||
|
||||
#pragma acc init if_present
|
||||
#pragma acc shutdown if_present
|
||||
}
|
||||
|
||||
void
|
||||
t2 ()
|
||||
{
|
||||
int a, b, c[10];
|
||||
|
||||
#pragma acc update self(a)
|
||||
#pragma acc parallel
|
||||
#pragma acc loop if_present /* { dg-error "'if_present' is not valid" } */
|
||||
for (b = 1; b < 10; b++)
|
||||
;
|
||||
#pragma acc end parallel
|
||||
|
||||
#pragma acc kernels loop if_present /* { dg-error "'if_present' is not valid" } */
|
||||
for (b = 1; b < 10; b++)
|
||||
;
|
||||
|
||||
#pragma acc parallel loop if_present /* { dg-error "'if_present' is not valid" } */
|
||||
for (b = 1; b < 10; b++)
|
||||
;
|
||||
}
|
||||
|
||||
@ -1,4 +1,4 @@
|
||||
#pragma acc routine
|
||||
#pragma acc routine seq
|
||||
template <typename T> T
|
||||
accDouble(int val)
|
||||
{
|
||||
@ -31,7 +31,7 @@ oacc_parallel_copy (T a)
|
||||
|
||||
#pragma acc parallel num_gangs (a) if (1)
|
||||
{
|
||||
#pragma acc loop independent collapse (2) gang
|
||||
#pragma acc loop independent collapse (2)
|
||||
for (int i = 0; i < a; i++)
|
||||
for (int j = 0; j < 5; j++)
|
||||
b = a;
|
||||
@ -86,6 +86,8 @@ oacc_parallel_copy (T a)
|
||||
#pragma acc update self (b)
|
||||
#pragma acc update device (b)
|
||||
#pragma acc exit data delete (b)
|
||||
#pragma acc exit data finalize copyout (b)
|
||||
#pragma acc exit data delete (b) finalize
|
||||
|
||||
return b;
|
||||
}
|
||||
@ -133,6 +135,13 @@ oacc_kernels_copy (T a)
|
||||
b = a;
|
||||
}
|
||||
|
||||
#pragma acc update host (b)
|
||||
#pragma acc update self (b)
|
||||
#pragma acc update device (b)
|
||||
#pragma acc exit data delete (b)
|
||||
#pragma acc exit data finalize copyout (b)
|
||||
#pragma acc exit data delete (b) finalize
|
||||
|
||||
return b;
|
||||
}
|
||||
|
||||
|
||||
@ -146,5 +146,5 @@ end subroutine test
|
||||
! { dg-final { scan-tree-dump-times "acc loop private.i. private.j. tile.2, 3" 2 "gimple" } }
|
||||
! { dg-final { scan-tree-dump-times "acc loop private.i. independent" 2 "gimple" } }
|
||||
! { dg-final { scan-tree-dump-times "private.z" 2 "gimple" } }
|
||||
! { dg-final { scan-tree-dump-times "omp target oacc_\[^ \]+ map.force_tofrom:y" 2 "gimple" } }
|
||||
! { dg-final { scan-tree-dump-times "omp target oacc_\[^ \]+ map.tofrom:y" 2 "gimple" } }
|
||||
! { dg-final { scan-tree-dump-times "acc loop private.i. reduction..:y." 2 "gimple" } }
|
||||
|
||||
@ -15,10 +15,10 @@ end program test
|
||||
! { dg-final { scan-tree-dump-times "pragma acc data" 1 "original" } }
|
||||
|
||||
! { dg-final { scan-tree-dump-times "if" 1 "original" } }
|
||||
! { dg-final { scan-tree-dump-times "map\\(force_tofrom:i\\)" 1 "original" } }
|
||||
! { dg-final { scan-tree-dump-times "map\\(force_to:j\\)" 1 "original" } }
|
||||
! { dg-final { scan-tree-dump-times "map\\(force_from:k\\)" 1 "original" } }
|
||||
! { dg-final { scan-tree-dump-times "map\\(force_alloc:m\\)" 1 "original" } }
|
||||
! { dg-final { scan-tree-dump-times "map\\(tofrom:i\\)" 1 "original" } }
|
||||
! { dg-final { scan-tree-dump-times "map\\(to:j\\)" 1 "original" } }
|
||||
! { dg-final { scan-tree-dump-times "map\\(from:k\\)" 1 "original" } }
|
||||
! { dg-final { scan-tree-dump-times "map\\(alloc:m\\)" 1 "original" } }
|
||||
|
||||
! { dg-final { scan-tree-dump-times "map\\(force_present:o\\)" 1 "original" } }
|
||||
! { dg-final { scan-tree-dump-times "map\\(tofrom:p\\)" 1 "original" } }
|
||||
|
||||
@ -11,11 +11,11 @@ subroutine asubr (b)
|
||||
!$acc declare copyout (b) ! { dg-error "Invalid clause in module" }
|
||||
!$acc declare present (b) ! { dg-error "Invalid clause in module" }
|
||||
!$acc declare present_or_copy (b) ! { dg-error "Invalid clause in module" }
|
||||
!$acc declare present_or_copyin (b) ! { dg-error "Invalid clause in module" }
|
||||
!$acc declare present_or_copyin (b) ! { dg-error "present on multiple" }
|
||||
!$acc declare present_or_copyout (b) ! { dg-error "Invalid clause in module" }
|
||||
!$acc declare present_or_create (b) ! { dg-error "Invalid clause in module" }
|
||||
!$acc declare present_or_create (b) ! { dg-error "present on multiple" }
|
||||
!$acc declare deviceptr (b) ! { dg-error "Invalid clause in module" }
|
||||
!$acc declare create (b) copyin (b) ! { dg-error "present on multiple clauses" }
|
||||
!$acc declare create (b) copyin (b) ! { dg-error "present on multiple" }
|
||||
|
||||
end subroutine
|
||||
|
||||
|
||||
@ -8,7 +8,7 @@
|
||||
REAL, DIMENSION (2) :: F1_B
|
||||
|
||||
!$ACC DATA COPYIN (F1_A) COPYOUT (F1_B)
|
||||
! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(force_to:f1_a \[^\\)\]+\\) map\\(force_from:f1_b" 1 "gimple" } }
|
||||
! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(to:f1_a \[^\\)\]+\\) map\\(from:f1_b" 1 "gimple" } }
|
||||
!$ACC KERNELS
|
||||
! { dg-final { scan-tree-dump-times "omp target oacc_kernels map\\(tofrom:f1_b \[^\\)\]+\\) map\\(tofrom:f1_a" 1 "gimple" } }
|
||||
F1_B(1) = F1_A;
|
||||
@ -26,7 +26,7 @@
|
||||
REAL, DIMENSION (2) :: F2_B
|
||||
|
||||
!$ACC DATA COPYIN (F2_A) COPYOUT (F2_B)
|
||||
! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(force_to:f2_a \[^\\)\]+\\) map\\(force_from:f2_b" 1 "gimple" } }
|
||||
! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(to:f2_a \[^\\)\]+\\) map\\(from:f2_b" 1 "gimple" } }
|
||||
!$ACC KERNELS DEFAULT (NONE)
|
||||
! { dg-final { scan-tree-dump-times "omp target oacc_kernels default\\(none\\) map\\(tofrom:f2_b \[^\\)\]+\\) map\\(tofrom:f2_a" 1 "gimple" } }
|
||||
F2_B(1) = F2_A;
|
||||
@ -44,7 +44,7 @@
|
||||
REAL, DIMENSION (2) :: F3_B
|
||||
|
||||
!$ACC DATA COPYIN (F3_A) COPYOUT (F3_B)
|
||||
! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(force_to:f3_a \[^\\)\]+\\) map\\(force_from:f3_b" 1 "gimple" } }
|
||||
! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(to:f3_a \[^\\)\]+\\) map\\(from:f3_b" 1 "gimple" } }
|
||||
!$ACC KERNELS DEFAULT (PRESENT)
|
||||
! { dg-final { scan-tree-dump-times "omp target oacc_kernels default\\(present\\) map\\(tofrom:f3_b \[^\\)\]+\\) map\\(tofrom:f3_a" 1 "gimple" } }
|
||||
F3_B(1) = F3_A;
|
||||
|
||||
@ -84,5 +84,8 @@ contains
|
||||
!$acc exit data delete (tip) ! { dg-error "POINTER" }
|
||||
!$acc exit data delete (tia) ! { dg-error "ALLOCATABLE" }
|
||||
!$acc exit data copyout (i) delete (i) ! { dg-error "multiple clauses" }
|
||||
!$acc exit data finalize
|
||||
!$acc exit data finalize copyout (i)
|
||||
!$acc exit data finalize delete (i)
|
||||
end subroutine foo
|
||||
end module test
|
||||
|
||||
27
gcc/testsuite/gfortran.dg/goacc/finalize-1.f
Normal file
27
gcc/testsuite/gfortran.dg/goacc/finalize-1.f
Normal file
@ -0,0 +1,27 @@
|
||||
! Test valid usage and processing of the finalize clause.
|
||||
|
||||
! { dg-additional-options "-fdump-tree-original -fdump-tree-gimple" }
|
||||
|
||||
SUBROUTINE f
|
||||
IMPLICIT NONE
|
||||
INTEGER :: del_r
|
||||
REAL, DIMENSION (3) :: del_f
|
||||
DOUBLE PRECISION, DIMENSION (8) :: cpo_r
|
||||
LOGICAL :: cpo_f
|
||||
|
||||
!$ACC EXIT DATA DELETE (del_r)
|
||||
! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:del_r\\);$" 1 "original" } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(release:del_r \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } }
|
||||
|
||||
!$ACC EXIT DATA FINALIZE DELETE (del_f)
|
||||
! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:del_f\\) finalize;$" 1 "original" } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:del_f \\\[len: \[0-9\]+\\\]\\) finalize$" 1 "gimple" } }
|
||||
|
||||
!$ACC EXIT DATA COPYOUT (cpo_r)
|
||||
! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:cpo_r\\);$" 1 "original" } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(from:cpo_r \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } }
|
||||
|
||||
!$ACC EXIT DATA COPYOUT (cpo_f) FINALIZE
|
||||
! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:cpo_f\\) finalize;$" 1 "original" } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_from:cpo_f \\\[len: \[0-9\]+\\\]\\) finalize$" 1 "gimple" } }
|
||||
END SUBROUTINE f
|
||||
@ -15,9 +15,11 @@ program main
|
||||
|
||||
end program main
|
||||
|
||||
! The xfails occur in light of the new OpenACC data semantics.
|
||||
|
||||
! { dg-final { scan-tree-dump-times "clique 1 base 1" 4 "ealias" } }
|
||||
! { dg-final { scan-tree-dump-times "clique 1 base 2" 1 "ealias" } }
|
||||
! { dg-final { scan-tree-dump-times "clique 1 base 3" 1 "ealias" } }
|
||||
! { dg-final { scan-tree-dump-times "clique 1 base 4" 1 "ealias" } }
|
||||
! { dg-final { scan-tree-dump-times "clique 1 base 5" 1 "ealias" } }
|
||||
! { dg-final { scan-tree-dump-times "clique 1 base 2" 1 "ealias" { xfail *-*-* } } }
|
||||
! { dg-final { scan-tree-dump-times "clique 1 base 3" 1 "ealias" { xfail *-*-* } } }
|
||||
! { dg-final { scan-tree-dump-times "clique 1 base 4" 1 "ealias" { xfail *-*-* } } }
|
||||
! { dg-final { scan-tree-dump-times "clique 1 base 5" 1 "ealias" { xfail *-*-* } } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)clique .* base .*" 8 "ealias" } }
|
||||
|
||||
@ -15,9 +15,11 @@ program main
|
||||
|
||||
end program main
|
||||
|
||||
! The xfails occur in light of the new OpenACC data semantics.
|
||||
|
||||
! { dg-final { scan-tree-dump-times "clique 1 base 1" 4 "ealias" } }
|
||||
! { dg-final { scan-tree-dump-times "clique 1 base 2" 1 "ealias" } }
|
||||
! { dg-final { scan-tree-dump-times "clique 1 base 3" 1 "ealias" } }
|
||||
! { dg-final { scan-tree-dump-times "clique 1 base 4" 1 "ealias" } }
|
||||
! { dg-final { scan-tree-dump-times "clique 1 base 5" 1 "ealias" } }
|
||||
! { dg-final { scan-tree-dump-times "clique 1 base 2" 1 "ealias" { xfail *-*-* } } }
|
||||
! { dg-final { scan-tree-dump-times "clique 1 base 3" 1 "ealias" { xfail *-*-* } } }
|
||||
! { dg-final { scan-tree-dump-times "clique 1 base 4" 1 "ealias" { xfail *-*-* } } }
|
||||
! { dg-final { scan-tree-dump-times "clique 1 base 5" 1 "ealias" { xfail *-*-* } } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)clique .* base .*" 8 "ealias" } }
|
||||
|
||||
@ -21,10 +21,10 @@ end program test
|
||||
! { dg-final { scan-tree-dump-times "num_workers" 1 "original" } }
|
||||
! { dg-final { scan-tree-dump-times "vector_length" 1 "original" } }
|
||||
|
||||
! { dg-final { scan-tree-dump-times "map\\(force_tofrom:i\\)" 1 "original" } }
|
||||
! { dg-final { scan-tree-dump-times "map\\(force_to:j\\)" 1 "original" } }
|
||||
! { dg-final { scan-tree-dump-times "map\\(force_from:k\\)" 1 "original" } }
|
||||
! { dg-final { scan-tree-dump-times "map\\(force_alloc:m\\)" 1 "original" } }
|
||||
! { dg-final { scan-tree-dump-times "map\\(tofrom:i\\)" 1 "original" } }
|
||||
! { dg-final { scan-tree-dump-times "map\\(to:j\\)" 1 "original" } }
|
||||
! { dg-final { scan-tree-dump-times "map\\(from:k\\)" 1 "original" } }
|
||||
! { dg-final { scan-tree-dump-times "map\\(alloc:m\\)" 1 "original" } }
|
||||
|
||||
! { dg-final { scan-tree-dump-times "map\\(force_present:o\\)" 1 "original" } }
|
||||
! { dg-final { scan-tree-dump-times "map\\(tofrom:p\\)" 1 "original" } }
|
||||
|
||||
@ -25,6 +25,8 @@ contains
|
||||
local_a (:) = 5
|
||||
local_arg = 5
|
||||
|
||||
!$acc update device(local_a) if_present
|
||||
|
||||
!$acc kernels loop &
|
||||
!$acc gang(num:local_arg) worker(local_arg) vector(local_arg) &
|
||||
!$acc wait async(local_arg)
|
||||
@ -54,12 +56,16 @@ contains
|
||||
enddo
|
||||
enddo
|
||||
!$acc end kernels loop
|
||||
|
||||
!$acc exit data copyout(local_a) delete(local_i) finalize
|
||||
end subroutine local
|
||||
|
||||
subroutine nonlocal ()
|
||||
nonlocal_a (:) = 5
|
||||
nonlocal_arg = 5
|
||||
|
||||
!$acc update device(nonlocal_a) if_present
|
||||
|
||||
!$acc kernels loop &
|
||||
!$acc gang(num:nonlocal_arg) worker(nonlocal_arg) vector(nonlocal_arg) &
|
||||
!$acc wait async(nonlocal_arg)
|
||||
@ -89,5 +95,7 @@ contains
|
||||
enddo
|
||||
enddo
|
||||
!$acc end kernels loop
|
||||
|
||||
!$acc exit data copyout(nonlocal_a) delete(nonlocal_i) finalize
|
||||
end subroutine nonlocal
|
||||
end program main
|
||||
|
||||
@ -1,5 +1,4 @@
|
||||
! { dg-do compile }
|
||||
! { dg-additional-options "-fdump-tree-original" }
|
||||
! { dg-additional-options "-fdump-tree-original" }
|
||||
|
||||
! test for tree-dump-original and spaces-commas
|
||||
|
||||
@ -15,6 +14,7 @@ program test
|
||||
!$acc end parallel
|
||||
|
||||
end program test
|
||||
|
||||
! { dg-final { scan-tree-dump-times "pragma acc parallel" 1 "original" } }
|
||||
|
||||
! { dg-final { scan-tree-dump-times "if" 1 "original" } }
|
||||
@ -24,10 +24,10 @@ end program test
|
||||
! { dg-final { scan-tree-dump-times "vector_length" 1 "original" } }
|
||||
|
||||
! { dg-final { scan-tree-dump-times "reduction\\(max:q\\)" 1 "original" } }
|
||||
! { dg-final { scan-tree-dump-times "map\\(force_tofrom:i\\)" 1 "original" } }
|
||||
! { dg-final { scan-tree-dump-times "map\\(force_to:j\\)" 1 "original" } }
|
||||
! { dg-final { scan-tree-dump-times "map\\(force_from:k\\)" 1 "original" } }
|
||||
! { dg-final { scan-tree-dump-times "map\\(force_alloc:m\\)" 1 "original" } }
|
||||
! { dg-final { scan-tree-dump-times "map\\(tofrom:i\\)" 1 "original" } }
|
||||
! { dg-final { scan-tree-dump-times "map\\(to:j\\)" 1 "original" } }
|
||||
! { dg-final { scan-tree-dump-times "map\\(from:k\\)" 1 "original" } }
|
||||
! { dg-final { scan-tree-dump-times "map\\(alloc:m\\)" 1 "original" } }
|
||||
|
||||
! { dg-final { scan-tree-dump-times "map\\(force_present:o\\)" 1 "original" } }
|
||||
! { dg-final { scan-tree-dump-times "map\\(tofrom:p\\)" 1 "original" } }
|
||||
|
||||
@ -38,9 +38,7 @@ program test
|
||||
!$acc end parallel
|
||||
end program test
|
||||
|
||||
! { dg-final { scan-tree-dump-times "map.tofrom:v1" 8 "gimple" } }
|
||||
! { dg-final { scan-tree-dump-times "map.tofrom:v2" 8 "gimple" } }
|
||||
! { dg-final { scan-tree-dump-times "map.force_tofrom:v1" 1 "gimple" } }
|
||||
! { dg-final { scan-tree-dump-times "map.force_tofrom:v2" 1 "gimple" } }
|
||||
! { dg-final { scan-tree-dump-times "map.tofrom:v1" 9 "gimple" } }
|
||||
! { dg-final { scan-tree-dump-times "map.tofrom:v2" 9 "gimple" } }
|
||||
! { dg-final { scan-tree-dump-times "map.force_present:v1" 1 "gimple" } }
|
||||
! { dg-final { scan-tree-dump-times "map.force_present:v2" 1 "gimple" } }
|
||||
|
||||
27
gcc/testsuite/gfortran.dg/goacc/update-if_present-1.f90
Normal file
27
gcc/testsuite/gfortran.dg/goacc/update-if_present-1.f90
Normal file
@ -0,0 +1,27 @@
|
||||
! Test valid usages of the if_present clause.
|
||||
|
||||
! { dg-additional-options "-fdump-tree-omplower" }
|
||||
|
||||
subroutine t
|
||||
implicit none
|
||||
integer a, b, c(10)
|
||||
real, allocatable :: x, y, z(:)
|
||||
|
||||
a = 5
|
||||
b = 10
|
||||
c(:) = -1
|
||||
|
||||
allocate (x, y, z(100))
|
||||
|
||||
!$acc update self(a) if_present
|
||||
!$acc update device(b) if_present async
|
||||
!$acc update host(c(1:3)) wait(4) if_present
|
||||
!$acc update self(c) device(a) host(b) if_present async(10) if(a == 10)
|
||||
|
||||
!$acc update self(x) if_present
|
||||
!$acc update device(y) if_present async
|
||||
!$acc update host(z(1:3)) wait(3) if_present
|
||||
!$acc update self(z) device(y) host(x) if_present async(4) if(a == 1)
|
||||
end subroutine t
|
||||
|
||||
! { dg-final { scan-tree-dump-times " if_present" 8 "omplower" } }
|
||||
52
gcc/testsuite/gfortran.dg/goacc/update-if_present-2.f90
Normal file
52
gcc/testsuite/gfortran.dg/goacc/update-if_present-2.f90
Normal file
@ -0,0 +1,52 @@
|
||||
! Test invalid usages of the if_present clause.
|
||||
|
||||
subroutine t1
|
||||
implicit none
|
||||
!$acc routine gang if_present ! { dg-error "Unclassifiable OpenACC directive" }
|
||||
integer a, b, c(10)
|
||||
real, allocatable :: x, y, z(:)
|
||||
|
||||
a = 5
|
||||
b = 10
|
||||
c(:) = -1
|
||||
|
||||
allocate (x, y, z(100))
|
||||
|
||||
!$acc enter data copyin(a) if_present ! { dg-error "Unclassifiable OpenACC directive" }
|
||||
!$acc exit data copyout(a) if_present ! { dg-error "Unclassifiable OpenACC directive" }
|
||||
|
||||
!$acc data copy(a) if_present ! { dg-error "Unclassifiable OpenACC directive" }
|
||||
!$acc end data ! { dg-error "Unexpected ..ACC END DATA statement" }
|
||||
|
||||
!$acc declare link(a) if_present ! { dg-error "Unexpected junk after" }
|
||||
|
||||
!$acc init if_present ! { dg-error "Unclassifiable OpenACC directive" }
|
||||
!$acc shutdown if_present ! { dg-error "Unclassifiable OpenACC directive" }
|
||||
|
||||
!$acc update self(a) device_type(nvidia) device(b) if_present ! { dg-error "Unclassifiable OpenACC directive" }
|
||||
end subroutine t1
|
||||
|
||||
subroutine t2
|
||||
implicit none
|
||||
integer a, b, c(10)
|
||||
|
||||
a = 5
|
||||
b = 10
|
||||
c(:) = -1
|
||||
|
||||
!$acc parallel
|
||||
!$acc loop if_present ! { dg-error "Unclassifiable OpenACC directive" }
|
||||
do b = 1, 10
|
||||
end do
|
||||
!$acc end parallel
|
||||
|
||||
!$acc kernels loop if_present ! { dg-error "Unclassifiable OpenACC directive" }
|
||||
do b = 1, 10
|
||||
end do
|
||||
!$acc end kernels loop ! { dg-error "Unexpected ..ACC END KERNELS LOOP statement" }
|
||||
|
||||
!$acc parallel loop if_present ! { dg-error "Unclassifiable OpenACC directive" }
|
||||
do b = 1, 10
|
||||
end do
|
||||
!$acc end parallel loop ! { dg-error "Unexpected ..ACC END PARALLEL LOOP statement" }
|
||||
end subroutine t2
|
||||
@ -454,7 +454,13 @@ enum omp_clause_code {
|
||||
|
||||
/* OpenMP internal-only clause to specify grid dimensions of a gridified
|
||||
kernel. */
|
||||
OMP_CLAUSE__GRIDDIM_
|
||||
OMP_CLAUSE__GRIDDIM_,
|
||||
|
||||
/* OpenACC clause: if_present. */
|
||||
OMP_CLAUSE_IF_PRESENT,
|
||||
|
||||
/* OpenACC clause: finalize. */
|
||||
OMP_CLAUSE_FINALIZE
|
||||
};
|
||||
|
||||
#undef DEFTREESTRUCT
|
||||
|
||||
@ -1333,6 +1333,8 @@ convert_nonlocal_omp_clauses (tree *pclauses, struct walk_stmt_info *wi)
|
||||
case OMP_CLAUSE_SEQ:
|
||||
case OMP_CLAUSE_INDEPENDENT:
|
||||
case OMP_CLAUSE_AUTO:
|
||||
case OMP_CLAUSE_IF_PRESENT:
|
||||
case OMP_CLAUSE_FINALIZE:
|
||||
break;
|
||||
|
||||
/* The following clause belongs to the OpenACC cache directive, which
|
||||
@ -2022,6 +2024,8 @@ convert_local_omp_clauses (tree *pclauses, struct walk_stmt_info *wi)
|
||||
case OMP_CLAUSE_SEQ:
|
||||
case OMP_CLAUSE_INDEPENDENT:
|
||||
case OMP_CLAUSE_AUTO:
|
||||
case OMP_CLAUSE_IF_PRESENT:
|
||||
case OMP_CLAUSE_FINALIZE:
|
||||
break;
|
||||
|
||||
/* The following clause belongs to the OpenACC cache directive, which
|
||||
|
||||
@ -1045,6 +1045,12 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
|
||||
false);
|
||||
pp_right_paren (pp);
|
||||
break;
|
||||
case OMP_CLAUSE_IF_PRESENT:
|
||||
pp_string (pp, "if_present");
|
||||
break;
|
||||
case OMP_CLAUSE_FINALIZE:
|
||||
pp_string (pp, "finalize");
|
||||
break;
|
||||
|
||||
default:
|
||||
/* Should never happen. */
|
||||
|
||||
@ -343,6 +343,8 @@ unsigned const char omp_clause_num_ops[] =
|
||||
1, /* OMP_CLAUSE_VECTOR_LENGTH */
|
||||
3, /* OMP_CLAUSE_TILE */
|
||||
2, /* OMP_CLAUSE__GRIDDIM_ */
|
||||
0, /* OMP_CLAUSE_IF_PRESENT */
|
||||
0, /* OMP_CLAUSE_FINALIZE */
|
||||
};
|
||||
|
||||
const char * const omp_clause_code_name[] =
|
||||
@ -413,7 +415,9 @@ const char * const omp_clause_code_name[] =
|
||||
"num_workers",
|
||||
"vector_length",
|
||||
"tile",
|
||||
"_griddim_"
|
||||
"_griddim_",
|
||||
"if_present",
|
||||
"finalize",
|
||||
};
|
||||
|
||||
|
||||
@ -11594,6 +11598,8 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data,
|
||||
case OMP_CLAUSE_SEQ:
|
||||
case OMP_CLAUSE_TILE:
|
||||
case OMP_CLAUSE__SIMT_:
|
||||
case OMP_CLAUSE_IF_PRESENT:
|
||||
case OMP_CLAUSE_FINALIZE:
|
||||
WALK_SUBTREE_TAIL (OMP_CLAUSE_CHAIN (*tp));
|
||||
|
||||
case OMP_CLAUSE_LASTPRIVATE:
|
||||
|
||||
@ -1,3 +1,82 @@
|
||||
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
|
||||
|
||||
@ -853,6 +853,8 @@ struct splay_tree_key_s {
|
||||
uintptr_t tgt_offset;
|
||||
/* Reference count. */
|
||||
uintptr_t refcount;
|
||||
/* Dynamic reference count. */
|
||||
uintptr_t dynamic_refcount;
|
||||
/* Pointer to the original mapping of "omp declare target link" object. */
|
||||
splay_tree_key link_key;
|
||||
};
|
||||
@ -991,7 +993,9 @@ enum gomp_map_vars_kind
|
||||
};
|
||||
|
||||
extern void gomp_acc_insert_pointer (size_t, void **, size_t *, void *);
|
||||
extern void gomp_acc_remove_pointer (void *, bool, int, int);
|
||||
extern void gomp_acc_remove_pointer (void *, size_t, bool, int, int, int);
|
||||
extern void gomp_acc_declare_allocate (bool, size_t, void **, size_t *,
|
||||
unsigned short *);
|
||||
|
||||
extern struct target_mem_desc *gomp_map_vars (struct gomp_device_descr *,
|
||||
size_t, void **, void **,
|
||||
@ -1001,6 +1005,7 @@ extern void gomp_unmap_vars (struct target_mem_desc *, bool);
|
||||
extern void gomp_init_device (struct gomp_device_descr *);
|
||||
extern void gomp_free_memmap (struct splay_tree_s *);
|
||||
extern void gomp_unload_device (struct gomp_device_descr *);
|
||||
extern bool gomp_remove_var (struct gomp_device_descr *, splay_tree_key);
|
||||
|
||||
/* work.c */
|
||||
|
||||
|
||||
@ -386,6 +386,18 @@ OACC_2.0.1 {
|
||||
acc_pcreate;
|
||||
} OACC_2.0;
|
||||
|
||||
OACC_2.5 {
|
||||
global:
|
||||
acc_copyout_finalize;
|
||||
acc_copyout_finalize_32_h_;
|
||||
acc_copyout_finalize_64_h_;
|
||||
acc_copyout_finalize_array_h_;
|
||||
acc_delete_finalize;
|
||||
acc_delete_finalize_32_h_;
|
||||
acc_delete_finalize_64_h_;
|
||||
acc_delete_finalize_array_h_;
|
||||
} OACC_2.0.1;
|
||||
|
||||
GOACC_2.0 {
|
||||
global:
|
||||
GOACC_data_end;
|
||||
|
||||
@ -347,6 +347,7 @@ acc_map_data (void *h, void *d, size_t s)
|
||||
|
||||
tgt = gomp_map_vars (acc_dev, mapnum, &hostaddrs, &devaddrs, &sizes,
|
||||
&kinds, true, GOMP_MAP_VARS_OPENACC);
|
||||
tgt->list[0].key->refcount = REFCOUNT_INFINITY;
|
||||
}
|
||||
|
||||
gomp_mutex_lock (&acc_dev->lock);
|
||||
@ -389,6 +390,9 @@ acc_unmap_data (void *h)
|
||||
(void *) n->host_start, (int) host_size, (void *) h);
|
||||
}
|
||||
|
||||
/* Mark for removal. */
|
||||
n->refcount = 1;
|
||||
|
||||
t = n->tgt;
|
||||
|
||||
if (t->refcount == 2)
|
||||
@ -460,6 +464,11 @@ present_create_copy (unsigned f, void *h, size_t s)
|
||||
gomp_fatal ("[%p,+%d] not mapped", (void *)h, (int)s);
|
||||
}
|
||||
|
||||
if (n->refcount != REFCOUNT_INFINITY)
|
||||
{
|
||||
n->refcount++;
|
||||
n->dynamic_refcount++;
|
||||
}
|
||||
gomp_mutex_unlock (&acc_dev->lock);
|
||||
}
|
||||
else if (!(f & FLAG_CREATE))
|
||||
@ -483,6 +492,8 @@ present_create_copy (unsigned f, void *h, size_t s)
|
||||
|
||||
tgt = gomp_map_vars (acc_dev, mapnum, &hostaddrs, NULL, &s, &kinds, true,
|
||||
GOMP_MAP_VARS_OPENACC);
|
||||
/* Initialize dynamic refcount. */
|
||||
tgt->list[0].key->dynamic_refcount = 1;
|
||||
|
||||
gomp_mutex_lock (&acc_dev->lock);
|
||||
|
||||
@ -499,13 +510,13 @@ present_create_copy (unsigned f, void *h, size_t s)
|
||||
void *
|
||||
acc_create (void *h, size_t s)
|
||||
{
|
||||
return present_create_copy (FLAG_CREATE, h, s);
|
||||
return present_create_copy (FLAG_PRESENT | FLAG_CREATE, h, s);
|
||||
}
|
||||
|
||||
void *
|
||||
acc_copyin (void *h, size_t s)
|
||||
{
|
||||
return present_create_copy (FLAG_CREATE | FLAG_COPY, h, s);
|
||||
return present_create_copy (FLAG_PRESENT | FLAG_CREATE | FLAG_COPY, h, s);
|
||||
}
|
||||
|
||||
void *
|
||||
@ -542,7 +553,8 @@ acc_pcopyin (void *h, size_t s)
|
||||
}
|
||||
#endif
|
||||
|
||||
#define FLAG_COPYOUT (1 << 0)
|
||||
#define FLAG_COPYOUT (1 << 0)
|
||||
#define FLAG_FINALIZE (1 << 1)
|
||||
|
||||
static void
|
||||
delete_copyout (unsigned f, void *h, size_t s, const char *libfnname)
|
||||
@ -581,15 +593,52 @@ delete_copyout (unsigned f, void *h, size_t s, const char *libfnname)
|
||||
(void *) n->host_start, (int) host_size, (void *) h, (int) s);
|
||||
}
|
||||
|
||||
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_FINALIZE)
|
||||
{
|
||||
n->refcount -= n->dynamic_refcount;
|
||||
n->dynamic_refcount = 0;
|
||||
}
|
||||
else if (n->dynamic_refcount)
|
||||
{
|
||||
n->dynamic_refcount--;
|
||||
n->refcount--;
|
||||
}
|
||||
|
||||
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 (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);
|
||||
|
||||
if (f & FLAG_COPYOUT)
|
||||
acc_dev->dev2host_func (acc_dev->target_id, h, d, s);
|
||||
|
||||
acc_unmap_data (h);
|
||||
|
||||
if (!acc_dev->free_func (acc_dev->target_id, d))
|
||||
gomp_fatal ("error in freeing device memory in %s", libfnname);
|
||||
}
|
||||
|
||||
void
|
||||
@ -598,12 +647,36 @@ acc_delete (void *h , size_t s)
|
||||
delete_copyout (0, h, s, __FUNCTION__);
|
||||
}
|
||||
|
||||
void
|
||||
acc_delete_finalize (void *h , size_t s)
|
||||
{
|
||||
delete_copyout (FLAG_FINALIZE, h, s, __FUNCTION__);
|
||||
}
|
||||
|
||||
void
|
||||
acc_delete_finalize_async (void *h , size_t s, int async)
|
||||
{
|
||||
delete_copyout (FLAG_FINALIZE, h, s, __FUNCTION__);
|
||||
}
|
||||
|
||||
void
|
||||
acc_copyout (void *h, size_t s)
|
||||
{
|
||||
delete_copyout (FLAG_COPYOUT, h, s, __FUNCTION__);
|
||||
}
|
||||
|
||||
void
|
||||
acc_copyout_finalize (void *h, size_t s)
|
||||
{
|
||||
delete_copyout (FLAG_COPYOUT | FLAG_FINALIZE, h, s, __FUNCTION__);
|
||||
}
|
||||
|
||||
void
|
||||
acc_copyout_finalize_async (void *h, size_t s, int async)
|
||||
{
|
||||
delete_copyout (FLAG_COPYOUT | FLAG_FINALIZE, h, s, __FUNCTION__);
|
||||
}
|
||||
|
||||
static void
|
||||
update_dev_host (int is_dev, void *h, size_t s)
|
||||
{
|
||||
@ -659,11 +732,37 @@ gomp_acc_insert_pointer (size_t mapnum, void **hostaddrs, size_t *sizes,
|
||||
struct goacc_thread *thr = goacc_thread ();
|
||||
struct gomp_device_descr *acc_dev = thr->dev;
|
||||
|
||||
if (acc_is_present (*hostaddrs, *sizes))
|
||||
{
|
||||
splay_tree_key n;
|
||||
gomp_mutex_lock (&acc_dev->lock);
|
||||
n = lookup_host (acc_dev, *hostaddrs, *sizes);
|
||||
gomp_mutex_unlock (&acc_dev->lock);
|
||||
|
||||
tgt = n->tgt;
|
||||
for (size_t i = 0; i < tgt->list_count; i++)
|
||||
if (tgt->list[i].key == n)
|
||||
{
|
||||
for (size_t j = 0; j < mapnum; j++)
|
||||
if (i + j < tgt->list_count && tgt->list[i + j].key)
|
||||
{
|
||||
tgt->list[i + j].key->refcount++;
|
||||
tgt->list[i + j].key->dynamic_refcount++;
|
||||
}
|
||||
return;
|
||||
}
|
||||
/* Should not reach here. */
|
||||
gomp_fatal ("Dynamic refcount incrementing failed for pointer/pset");
|
||||
}
|
||||
|
||||
gomp_debug (0, " %s: prepare mappings\n", __FUNCTION__);
|
||||
tgt = gomp_map_vars (acc_dev, mapnum, hostaddrs,
|
||||
NULL, sizes, kinds, true, GOMP_MAP_VARS_OPENACC);
|
||||
gomp_debug (0, " %s: mappings prepared\n", __FUNCTION__);
|
||||
|
||||
/* Initialize dynamic refcount. */
|
||||
tgt->list[0].key->dynamic_refcount = 1;
|
||||
|
||||
gomp_mutex_lock (&acc_dev->lock);
|
||||
tgt->prev = acc_dev->openacc.data_environ;
|
||||
acc_dev->openacc.data_environ = tgt;
|
||||
@ -671,7 +770,8 @@ gomp_acc_insert_pointer (size_t mapnum, void **hostaddrs, size_t *sizes,
|
||||
}
|
||||
|
||||
void
|
||||
gomp_acc_remove_pointer (void *h, bool force_copyfrom, int async, int mapnum)
|
||||
gomp_acc_remove_pointer (void *h, size_t s, bool force_copyfrom, int async,
|
||||
int finalize, int mapnum)
|
||||
{
|
||||
struct goacc_thread *thr = goacc_thread ();
|
||||
struct gomp_device_descr *acc_dev = thr->dev;
|
||||
@ -679,6 +779,9 @@ gomp_acc_remove_pointer (void *h, bool force_copyfrom, int async, int mapnum)
|
||||
struct target_mem_desc *t;
|
||||
int minrefs = (mapnum == 1) ? 2 : 3;
|
||||
|
||||
if (!acc_is_present (h, s))
|
||||
return;
|
||||
|
||||
gomp_mutex_lock (&acc_dev->lock);
|
||||
|
||||
n = lookup_host (acc_dev, h, 1);
|
||||
@ -693,40 +796,65 @@ gomp_acc_remove_pointer (void *h, bool force_copyfrom, int async, int mapnum)
|
||||
|
||||
t = n->tgt;
|
||||
|
||||
struct target_mem_desc *tp;
|
||||
|
||||
if (t->refcount == minrefs)
|
||||
if (n->refcount < n->dynamic_refcount)
|
||||
{
|
||||
/* 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;
|
||||
|
||||
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;
|
||||
}
|
||||
}
|
||||
gomp_mutex_unlock (&acc_dev->lock);
|
||||
gomp_fatal ("Dynamic reference counting assert fail\n");
|
||||
}
|
||||
|
||||
if (force_copyfrom)
|
||||
t->list[0].copy_from = 1;
|
||||
if (finalize)
|
||||
{
|
||||
n->refcount -= n->dynamic_refcount;
|
||||
n->dynamic_refcount = 0;
|
||||
}
|
||||
else if (n->dynamic_refcount)
|
||||
{
|
||||
n->dynamic_refcount--;
|
||||
n->refcount--;
|
||||
}
|
||||
|
||||
gomp_mutex_unlock (&acc_dev->lock);
|
||||
|
||||
/* 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);
|
||||
if (n->refcount == 0)
|
||||
{
|
||||
if (t->refcount == minrefs)
|
||||
{
|
||||
/* 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 (n->tgt == t)
|
||||
{
|
||||
if (tp)
|
||||
tp->prev = t->prev;
|
||||
else
|
||||
acc_dev->openacc.data_environ = t->prev;
|
||||
break;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
/* 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;
|
||||
}
|
||||
|
||||
/* 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);
|
||||
}
|
||||
|
||||
gomp_mutex_unlock (&acc_dev->lock);
|
||||
|
||||
gomp_debug (0, " %s: mappings restored\n", __FUNCTION__);
|
||||
}
|
||||
|
||||
@ -38,15 +38,68 @@
|
||||
#include <stdarg.h>
|
||||
#include <assert.h>
|
||||
|
||||
/* Returns the number of mappings associated with the pointer or pset. PSET
|
||||
have three mappings, whereas pointer have two. */
|
||||
|
||||
static int
|
||||
find_pset (int pos, size_t mapnum, unsigned short *kinds)
|
||||
find_pointer (int pos, size_t mapnum, unsigned short *kinds)
|
||||
{
|
||||
if (pos + 1 >= mapnum)
|
||||
return 0;
|
||||
|
||||
unsigned char kind = kinds[pos+1] & 0xff;
|
||||
|
||||
return kind == GOMP_MAP_TO_PSET;
|
||||
if (kind == GOMP_MAP_TO_PSET)
|
||||
return 3;
|
||||
else if (kind == GOMP_MAP_POINTER)
|
||||
return 2;
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
/* Handle the mapping pair that are presented when a
|
||||
deviceptr clause is used with Fortran. */
|
||||
|
||||
static void
|
||||
handle_ftn_pointers (size_t mapnum, void **hostaddrs, size_t *sizes,
|
||||
unsigned short *kinds)
|
||||
{
|
||||
int i;
|
||||
|
||||
for (i = 0; i < mapnum; i++)
|
||||
{
|
||||
unsigned short kind1 = kinds[i] & 0xff;
|
||||
|
||||
/* Handle Fortran deviceptr clause. */
|
||||
if (kind1 == GOMP_MAP_FORCE_DEVICEPTR)
|
||||
{
|
||||
unsigned short kind2;
|
||||
|
||||
if (i < (signed)mapnum - 1)
|
||||
kind2 = kinds[i + 1] & 0xff;
|
||||
else
|
||||
kind2 = 0xffff;
|
||||
|
||||
if (sizes[i] == sizeof (void *))
|
||||
continue;
|
||||
|
||||
/* At this point, we're dealing with a Fortran deviceptr.
|
||||
If the next element is not what we're expecting, then
|
||||
this is an instance of where the deviceptr variable was
|
||||
not used within the region and the pointer was removed
|
||||
by the gimplifier. */
|
||||
if (kind2 == GOMP_MAP_POINTER
|
||||
&& sizes[i + 1] == 0
|
||||
&& hostaddrs[i] == *(void **)hostaddrs[i + 1])
|
||||
{
|
||||
kinds[i+1] = kinds[i];
|
||||
sizes[i+1] = sizeof (void *);
|
||||
}
|
||||
|
||||
/* Invalidate the entry. */
|
||||
hostaddrs[i] = NULL;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
static void goacc_wait (int async, int num_waits, va_list *ap);
|
||||
@ -88,6 +141,8 @@ GOACC_parallel_keyed (int device, void (*fn) (void *),
|
||||
thr = goacc_thread ();
|
||||
acc_dev = thr->dev;
|
||||
|
||||
handle_ftn_pointers (mapnum, hostaddrs, sizes, kinds);
|
||||
|
||||
/* Host fallback if "if" clause is false or if the current device is set to
|
||||
the host. */
|
||||
if (host_fallback)
|
||||
@ -183,10 +238,29 @@ GOACC_parallel_keyed (int device, void (*fn) (void *),
|
||||
async, dims, tgt);
|
||||
|
||||
/* If running synchronously, unmap immediately. */
|
||||
bool copyfrom = true;
|
||||
if (async_synchronous_p (async))
|
||||
gomp_unmap_vars (tgt, true);
|
||||
else
|
||||
tgt->device_descr->openacc.register_async_cleanup_func (tgt, async);
|
||||
{
|
||||
bool async_unmap = false;
|
||||
for (size_t i = 0; i < tgt->list_count; i++)
|
||||
{
|
||||
splay_tree_key k = tgt->list[i].key;
|
||||
if (k && k->refcount == 1)
|
||||
{
|
||||
async_unmap = true;
|
||||
break;
|
||||
}
|
||||
}
|
||||
if (async_unmap)
|
||||
tgt->device_descr->openacc.register_async_cleanup_func (tgt, async);
|
||||
else
|
||||
{
|
||||
copyfrom = false;
|
||||
gomp_unmap_vars (tgt, copyfrom);
|
||||
}
|
||||
}
|
||||
|
||||
acc_dev->openacc.async_set_async_func (acc_async_sync);
|
||||
}
|
||||
@ -286,6 +360,17 @@ GOACC_enter_exit_data (int device, size_t mapnum,
|
||||
va_end (ap);
|
||||
}
|
||||
|
||||
/* Determine whether "finalize" semantics apply to all mappings of this
|
||||
OpenACC directive. */
|
||||
bool finalize = false;
|
||||
if (mapnum > 0)
|
||||
{
|
||||
unsigned char kind = kinds[0] & 0xff;
|
||||
if (kind == GOMP_MAP_DELETE
|
||||
|| kind == GOMP_MAP_FORCE_FROM)
|
||||
finalize = true;
|
||||
}
|
||||
|
||||
acc_dev->openacc.async_set_async_func (async);
|
||||
|
||||
/* Determine if this is an "acc enter data". */
|
||||
@ -298,13 +383,17 @@ GOACC_enter_exit_data (int device, size_t mapnum,
|
||||
|
||||
if (kind == GOMP_MAP_FORCE_ALLOC
|
||||
|| kind == GOMP_MAP_FORCE_PRESENT
|
||||
|| kind == GOMP_MAP_FORCE_TO)
|
||||
|| kind == GOMP_MAP_FORCE_TO
|
||||
|| kind == GOMP_MAP_TO
|
||||
|| kind == GOMP_MAP_ALLOC)
|
||||
{
|
||||
data_enter = true;
|
||||
break;
|
||||
}
|
||||
|
||||
if (kind == GOMP_MAP_DELETE
|
||||
if (kind == GOMP_MAP_RELEASE
|
||||
|| kind == GOMP_MAP_DELETE
|
||||
|| kind == GOMP_MAP_FROM
|
||||
|| kind == GOMP_MAP_FORCE_FROM)
|
||||
break;
|
||||
|
||||
@ -312,31 +401,39 @@ GOACC_enter_exit_data (int device, size_t mapnum,
|
||||
kind);
|
||||
}
|
||||
|
||||
/* In c, non-pointers and arrays are represented by a single data clause.
|
||||
Dynamically allocated arrays and subarrays are represented by a data
|
||||
clause followed by an internal GOMP_MAP_POINTER.
|
||||
|
||||
In fortran, scalars and not allocated arrays are represented by a
|
||||
single data clause. Allocated arrays and subarrays have three mappings:
|
||||
1) the original data clause, 2) a PSET 3) a pointer to the array data.
|
||||
*/
|
||||
|
||||
if (data_enter)
|
||||
{
|
||||
for (i = 0; i < mapnum; i++)
|
||||
{
|
||||
unsigned char kind = kinds[i] & 0xff;
|
||||
|
||||
/* Scan for PSETs. */
|
||||
int psets = find_pset (i, mapnum, kinds);
|
||||
/* Scan for pointers and PSETs. */
|
||||
int pointer = find_pointer (i, mapnum, kinds);
|
||||
|
||||
if (!psets)
|
||||
if (!pointer)
|
||||
{
|
||||
switch (kind)
|
||||
{
|
||||
case GOMP_MAP_POINTER:
|
||||
gomp_acc_insert_pointer (1, &hostaddrs[i], &sizes[i],
|
||||
&kinds[i]);
|
||||
case GOMP_MAP_ALLOC:
|
||||
acc_present_or_create (hostaddrs[i], sizes[i]);
|
||||
break;
|
||||
case GOMP_MAP_FORCE_ALLOC:
|
||||
acc_create (hostaddrs[i], sizes[i]);
|
||||
break;
|
||||
case GOMP_MAP_FORCE_PRESENT:
|
||||
case GOMP_MAP_TO:
|
||||
acc_present_or_copyin (hostaddrs[i], sizes[i]);
|
||||
break;
|
||||
case GOMP_MAP_FORCE_TO:
|
||||
acc_present_or_copyin (hostaddrs[i], sizes[i]);
|
||||
acc_copyin (hostaddrs[i], sizes[i]);
|
||||
break;
|
||||
default:
|
||||
gomp_fatal (">>>> GOACC_enter_exit_data UNHANDLED kind 0x%.2x",
|
||||
@ -346,12 +443,13 @@ GOACC_enter_exit_data (int device, size_t mapnum,
|
||||
}
|
||||
else
|
||||
{
|
||||
gomp_acc_insert_pointer (3, &hostaddrs[i], &sizes[i], &kinds[i]);
|
||||
gomp_acc_insert_pointer (pointer, &hostaddrs[i],
|
||||
&sizes[i], &kinds[i]);
|
||||
/* Increment 'i' by two because OpenACC requires fortran
|
||||
arrays to be contiguous, so each PSET is associated with
|
||||
one of MAP_FORCE_ALLOC/MAP_FORCE_PRESET/MAP_FORCE_TO, and
|
||||
one MAP_POINTER. */
|
||||
i += 2;
|
||||
i += pointer - 1;
|
||||
}
|
||||
}
|
||||
}
|
||||
@ -360,22 +458,28 @@ GOACC_enter_exit_data (int device, size_t mapnum,
|
||||
{
|
||||
unsigned char kind = kinds[i] & 0xff;
|
||||
|
||||
int psets = find_pset (i, mapnum, kinds);
|
||||
int pointer = find_pointer (i, mapnum, kinds);
|
||||
|
||||
if (!psets)
|
||||
if (!pointer)
|
||||
{
|
||||
switch (kind)
|
||||
{
|
||||
case GOMP_MAP_POINTER:
|
||||
gomp_acc_remove_pointer (hostaddrs[i], (kinds[i] & 0xff)
|
||||
== GOMP_MAP_FORCE_FROM,
|
||||
async, 1);
|
||||
break;
|
||||
case GOMP_MAP_RELEASE:
|
||||
case GOMP_MAP_DELETE:
|
||||
acc_delete (hostaddrs[i], sizes[i]);
|
||||
if (acc_is_present (hostaddrs[i], sizes[i]))
|
||||
{
|
||||
if (finalize)
|
||||
acc_delete_finalize (hostaddrs[i], sizes[i]);
|
||||
else
|
||||
acc_delete (hostaddrs[i], sizes[i]);
|
||||
}
|
||||
break;
|
||||
case GOMP_MAP_FROM:
|
||||
case GOMP_MAP_FORCE_FROM:
|
||||
acc_copyout (hostaddrs[i], sizes[i]);
|
||||
if (finalize)
|
||||
acc_copyout_finalize (hostaddrs[i], sizes[i]);
|
||||
else
|
||||
acc_copyout (hostaddrs[i], sizes[i]);
|
||||
break;
|
||||
default:
|
||||
gomp_fatal (">>>> GOACC_enter_exit_data UNHANDLED kind 0x%.2x",
|
||||
@ -385,10 +489,12 @@ GOACC_enter_exit_data (int device, size_t mapnum,
|
||||
}
|
||||
else
|
||||
{
|
||||
gomp_acc_remove_pointer (hostaddrs[i], (kinds[i] & 0xff)
|
||||
== GOMP_MAP_FORCE_FROM, async, 3);
|
||||
bool copyfrom = (kind == GOMP_MAP_FORCE_FROM
|
||||
|| kind == GOMP_MAP_FROM);
|
||||
gomp_acc_remove_pointer (hostaddrs[i], sizes[i], copyfrom, async,
|
||||
finalize, pointer);
|
||||
/* See the above comment. */
|
||||
i += 2;
|
||||
i += pointer - 1;
|
||||
}
|
||||
}
|
||||
|
||||
@ -447,6 +553,7 @@ GOACC_update (int device, size_t mapnum,
|
||||
|
||||
acc_dev->openacc.async_set_async_func (async);
|
||||
|
||||
bool update_device = false;
|
||||
for (i = 0; i < mapnum; ++i)
|
||||
{
|
||||
unsigned char kind = kinds[i] & 0xff;
|
||||
@ -457,11 +564,46 @@ GOACC_update (int device, size_t mapnum,
|
||||
case GOMP_MAP_TO_PSET:
|
||||
break;
|
||||
|
||||
case GOMP_MAP_ALWAYS_POINTER:
|
||||
if (update_device)
|
||||
{
|
||||
/* Save the contents of the host pointer. */
|
||||
void *dptr = acc_deviceptr (hostaddrs[i-1]);
|
||||
uintptr_t t = *(uintptr_t *) hostaddrs[i];
|
||||
|
||||
/* Update the contents of the host pointer to reflect
|
||||
the value of the allocated device memory in the
|
||||
previous pointer. */
|
||||
*(uintptr_t *) hostaddrs[i] = (uintptr_t)dptr;
|
||||
acc_update_device (hostaddrs[i], sizeof (uintptr_t));
|
||||
|
||||
/* Restore the host pointer. */
|
||||
*(uintptr_t *) hostaddrs[i] = t;
|
||||
update_device = false;
|
||||
}
|
||||
break;
|
||||
|
||||
case GOMP_MAP_TO:
|
||||
if (!acc_is_present (hostaddrs[i], sizes[i]))
|
||||
{
|
||||
update_device = false;
|
||||
break;
|
||||
}
|
||||
/* Fallthru */
|
||||
case GOMP_MAP_FORCE_TO:
|
||||
update_device = true;
|
||||
acc_update_device (hostaddrs[i], sizes[i]);
|
||||
break;
|
||||
|
||||
case GOMP_MAP_FROM:
|
||||
if (!acc_is_present (hostaddrs[i], sizes[i]))
|
||||
{
|
||||
update_device = false;
|
||||
break;
|
||||
}
|
||||
/* Fallthru */
|
||||
case GOMP_MAP_FORCE_FROM:
|
||||
update_device = false;
|
||||
acc_update_self (hostaddrs[i], sizes[i]);
|
||||
break;
|
||||
|
||||
@ -522,6 +664,7 @@ GOACC_declare (int device, size_t mapnum,
|
||||
case GOMP_MAP_FORCE_FROM:
|
||||
case GOMP_MAP_FORCE_TO:
|
||||
case GOMP_MAP_POINTER:
|
||||
case GOMP_MAP_RELEASE:
|
||||
case GOMP_MAP_DELETE:
|
||||
GOACC_enter_exit_data (device, 1, &hostaddrs[i], &sizes[i],
|
||||
&kinds[i], GOMP_ASYNC_SYNC, 0);
|
||||
@ -543,7 +686,6 @@ GOACC_declare (int device, size_t mapnum,
|
||||
break;
|
||||
|
||||
case GOMP_MAP_FROM:
|
||||
kinds[i] = GOMP_MAP_FORCE_FROM;
|
||||
GOACC_enter_exit_data (device, 1, &hostaddrs[i], &sizes[i],
|
||||
&kinds[i], GOMP_ASYNC_SYNC, 0);
|
||||
break;
|
||||
|
||||
@ -222,6 +222,24 @@ module openacc_internal
|
||||
type (*), dimension (..), contiguous :: a
|
||||
end subroutine
|
||||
|
||||
subroutine acc_copyout_finalize_32_h (a, len)
|
||||
use iso_c_binding, only: c_int32_t
|
||||
!GCC$ ATTRIBUTES NO_ARG_CHECK :: a
|
||||
type (*), dimension (*) :: a
|
||||
integer (c_int32_t) len
|
||||
end subroutine
|
||||
|
||||
subroutine acc_copyout_finalize_64_h (a, len)
|
||||
use iso_c_binding, only: c_int64_t
|
||||
!GCC$ ATTRIBUTES NO_ARG_CHECK :: a
|
||||
type (*), dimension (*) :: a
|
||||
integer (c_int64_t) len
|
||||
end subroutine
|
||||
|
||||
subroutine acc_copyout_finalize_array_h (a)
|
||||
type (*), dimension (..), contiguous :: a
|
||||
end subroutine
|
||||
|
||||
subroutine acc_delete_32_h (a, len)
|
||||
use iso_c_binding, only: c_int32_t
|
||||
!GCC$ ATTRIBUTES NO_ARG_CHECK :: a
|
||||
@ -240,6 +258,24 @@ module openacc_internal
|
||||
type (*), dimension (..), contiguous :: a
|
||||
end subroutine
|
||||
|
||||
subroutine acc_delete_finalize_32_h (a, len)
|
||||
use iso_c_binding, only: c_int32_t
|
||||
!GCC$ ATTRIBUTES NO_ARG_CHECK :: a
|
||||
type (*), dimension (*) :: a
|
||||
integer (c_int32_t) len
|
||||
end subroutine
|
||||
|
||||
subroutine acc_delete_finalize_64_h (a, len)
|
||||
use iso_c_binding, only: c_int64_t
|
||||
!GCC$ ATTRIBUTES NO_ARG_CHECK :: a
|
||||
type (*), dimension (*) :: a
|
||||
integer (c_int64_t) len
|
||||
end subroutine
|
||||
|
||||
subroutine acc_delete_finalize_array_h (a)
|
||||
type (*), dimension (..), contiguous :: a
|
||||
end subroutine
|
||||
|
||||
subroutine acc_update_device_32_h (a, len)
|
||||
use iso_c_binding, only: c_int32_t
|
||||
!GCC$ ATTRIBUTES NO_ARG_CHECK :: a
|
||||
@ -426,6 +462,14 @@ module openacc_internal
|
||||
integer (c_size_t), value :: len
|
||||
end subroutine
|
||||
|
||||
subroutine acc_copyout_finalize_l (a, len) &
|
||||
bind (C, name = "acc_copyout_finalize")
|
||||
use iso_c_binding, only: c_size_t
|
||||
!GCC$ ATTRIBUTES NO_ARG_CHECK :: a
|
||||
type (*), dimension (*) :: a
|
||||
integer (c_size_t), value :: len
|
||||
end subroutine
|
||||
|
||||
subroutine acc_delete_l (a, len) &
|
||||
bind (C, name = "acc_delete")
|
||||
use iso_c_binding, only: c_size_t
|
||||
@ -434,6 +478,14 @@ module openacc_internal
|
||||
integer (c_size_t), value :: len
|
||||
end subroutine
|
||||
|
||||
subroutine acc_delete_finalize_l (a, len) &
|
||||
bind (C, name = "acc_delete_finalize")
|
||||
use iso_c_binding, only: c_size_t
|
||||
!GCC$ ATTRIBUTES NO_ARG_CHECK :: a
|
||||
type (*), dimension (*) :: a
|
||||
integer (c_size_t), value :: len
|
||||
end subroutine
|
||||
|
||||
subroutine acc_update_device_l (a, len) &
|
||||
bind (C, name = "acc_update_device")
|
||||
use iso_c_binding, only: c_size_t
|
||||
@ -598,12 +650,24 @@ module openacc
|
||||
procedure :: acc_copyout_array_h
|
||||
end interface
|
||||
|
||||
interface acc_copyout_finalize
|
||||
procedure :: acc_copyout_finalize_32_h
|
||||
procedure :: acc_copyout_finalize_64_h
|
||||
procedure :: acc_copyout_finalize_array_h
|
||||
end interface
|
||||
|
||||
interface acc_delete
|
||||
procedure :: acc_delete_32_h
|
||||
procedure :: acc_delete_64_h
|
||||
procedure :: acc_delete_array_h
|
||||
end interface
|
||||
|
||||
interface acc_delete_finalize
|
||||
procedure :: acc_delete_finalize_32_h
|
||||
procedure :: acc_delete_finalize_64_h
|
||||
procedure :: acc_delete_finalize_array_h
|
||||
end interface
|
||||
|
||||
interface acc_update_device
|
||||
procedure :: acc_update_device_32_h
|
||||
procedure :: acc_update_device_64_h
|
||||
@ -860,6 +924,30 @@ subroutine acc_copyout_array_h (a)
|
||||
call acc_copyout_l (a, sizeof (a))
|
||||
end subroutine
|
||||
|
||||
subroutine acc_copyout_finalize_32_h (a, len)
|
||||
use iso_c_binding, only: c_int32_t, c_size_t
|
||||
use openacc_internal, only: acc_copyout_finalize_l
|
||||
!GCC$ ATTRIBUTES NO_ARG_CHECK :: a
|
||||
type (*), dimension (*) :: a
|
||||
integer (c_int32_t) len
|
||||
call acc_copyout_finalize_l (a, int (len, kind = c_size_t))
|
||||
end subroutine
|
||||
|
||||
subroutine acc_copyout_finalize_64_h (a, len)
|
||||
use iso_c_binding, only: c_int64_t, c_size_t
|
||||
use openacc_internal, only: acc_copyout_finalize_l
|
||||
!GCC$ ATTRIBUTES NO_ARG_CHECK :: a
|
||||
type (*), dimension (*) :: a
|
||||
integer (c_int64_t) len
|
||||
call acc_copyout_finalize_l (a, int (len, kind = c_size_t))
|
||||
end subroutine
|
||||
|
||||
subroutine acc_copyout_finalize_array_h (a)
|
||||
use openacc_internal, only: acc_copyout_finalize_l
|
||||
type (*), dimension (..), contiguous :: a
|
||||
call acc_copyout_finalize_l (a, sizeof (a))
|
||||
end subroutine
|
||||
|
||||
subroutine acc_delete_32_h (a, len)
|
||||
use iso_c_binding, only: c_int32_t, c_size_t
|
||||
use openacc_internal, only: acc_delete_l
|
||||
@ -884,6 +972,30 @@ subroutine acc_delete_array_h (a)
|
||||
call acc_delete_l (a, sizeof (a))
|
||||
end subroutine
|
||||
|
||||
subroutine acc_delete_finalize_32_h (a, len)
|
||||
use iso_c_binding, only: c_int32_t, c_size_t
|
||||
use openacc_internal, only: acc_delete_finalize_l
|
||||
!GCC$ ATTRIBUTES NO_ARG_CHECK :: a
|
||||
type (*), dimension (*) :: a
|
||||
integer (c_int32_t) len
|
||||
call acc_delete_finalize_l (a, int (len, kind = c_size_t))
|
||||
end subroutine
|
||||
|
||||
subroutine acc_delete_finalize_64_h (a, len)
|
||||
use iso_c_binding, only: c_int64_t, c_size_t
|
||||
use openacc_internal, only: acc_delete_finalize_l
|
||||
!GCC$ ATTRIBUTES NO_ARG_CHECK :: a
|
||||
type (*), dimension (*) :: a
|
||||
integer (c_int64_t) len
|
||||
call acc_delete_finalize_l (a, int (len, kind = c_size_t))
|
||||
end subroutine
|
||||
|
||||
subroutine acc_delete_finalize_array_h (a)
|
||||
use openacc_internal, only: acc_delete_finalize_l
|
||||
type (*), dimension (..), contiguous :: a
|
||||
call acc_delete_finalize_l (a, sizeof (a))
|
||||
end subroutine
|
||||
|
||||
subroutine acc_update_device_32_h (a, len)
|
||||
use iso_c_binding, only: c_int32_t, c_size_t
|
||||
use openacc_internal, only: acc_update_device_l
|
||||
|
||||
@ -109,6 +109,12 @@ int acc_is_present (void *, size_t) __GOACC_NOTHROW;
|
||||
void acc_memcpy_to_device (void *, void *, size_t) __GOACC_NOTHROW;
|
||||
void acc_memcpy_from_device (void *, void *, size_t) __GOACC_NOTHROW;
|
||||
|
||||
/* Finalize versions of copyout/delete functions, specified in OpenACC 2.5. */
|
||||
void acc_copyout_finalize (void *, size_t) __GOACC_NOTHROW;
|
||||
void acc_copyout_finalize_async (void *, size_t, int) __GOACC_NOTHROW;
|
||||
void acc_delete_finalize (void *, size_t) __GOACC_NOTHROW;
|
||||
void acc_delete_finalize_async (void *, size_t, int) __GOACC_NOTHROW;
|
||||
|
||||
/* CUDA-specific routines. */
|
||||
void *acc_get_current_cuda_device (void) __GOACC_NOTHROW;
|
||||
void *acc_get_current_cuda_context (void) __GOACC_NOTHROW;
|
||||
|
||||
@ -273,6 +273,26 @@
|
||||
end subroutine
|
||||
end interface
|
||||
|
||||
interface acc_copyout_finalize
|
||||
subroutine acc_copyout_finalize_32_h (a, len)
|
||||
use iso_c_binding, only: c_int32_t
|
||||
!GCC$ ATTRIBUTES NO_ARG_CHECK :: a
|
||||
type (*), dimension (*) :: a
|
||||
integer (c_int32_t) len
|
||||
end subroutine
|
||||
|
||||
subroutine acc_copyout_finalize_64_h (a, len)
|
||||
use iso_c_binding, only: c_int64_t
|
||||
!GCC$ ATTRIBUTES NO_ARG_CHECK :: a
|
||||
type (*), dimension (*) :: a
|
||||
integer (c_int64_t) len
|
||||
end subroutine
|
||||
|
||||
subroutine acc_copyout_finalize_array_h (a)
|
||||
type (*), dimension (..), contiguous :: a
|
||||
end subroutine
|
||||
end interface
|
||||
|
||||
interface acc_delete
|
||||
subroutine acc_delete_32_h (a, len)
|
||||
use iso_c_binding, only: c_int32_t
|
||||
@ -293,6 +313,26 @@
|
||||
end subroutine
|
||||
end interface
|
||||
|
||||
interface acc_delete_finalize
|
||||
subroutine acc_delete_finalize_32_h (a, len)
|
||||
use iso_c_binding, only: c_int32_t
|
||||
!GCC$ ATTRIBUTES NO_ARG_CHECK :: a
|
||||
type (*), dimension (*) :: a
|
||||
integer (c_int32_t) len
|
||||
end subroutine
|
||||
|
||||
subroutine acc_delete_finalize_64_h (a, len)
|
||||
use iso_c_binding, only: c_int64_t
|
||||
!GCC$ ATTRIBUTES NO_ARG_CHECK :: a
|
||||
type (*), dimension (*) :: a
|
||||
integer (c_int64_t) len
|
||||
end subroutine
|
||||
|
||||
subroutine acc_delete_finalize_array_h (a)
|
||||
type (*), dimension (..), contiguous :: a
|
||||
end subroutine
|
||||
end interface
|
||||
|
||||
interface acc_update_device
|
||||
subroutine acc_update_device_32_h (a, len)
|
||||
use iso_c_binding, only: c_int32_t
|
||||
|
||||
@ -859,6 +859,7 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
|
||||
tgt->list[i].offset = 0;
|
||||
tgt->list[i].length = k->host_end - k->host_start;
|
||||
k->refcount = 1;
|
||||
k->dynamic_refcount = 0;
|
||||
tgt->refcount++;
|
||||
array->left = NULL;
|
||||
array->right = NULL;
|
||||
@ -1011,6 +1012,23 @@ gomp_unmap_tgt (struct target_mem_desc *tgt)
|
||||
free (tgt);
|
||||
}
|
||||
|
||||
attribute_hidden bool
|
||||
gomp_remove_var (struct gomp_device_descr *devicep, splay_tree_key k)
|
||||
{
|
||||
bool is_tgt_unmapped = false;
|
||||
splay_tree_remove (&devicep->mem_map, k);
|
||||
if (k->link_key)
|
||||
splay_tree_insert (&devicep->mem_map, (splay_tree_node) k->link_key);
|
||||
if (k->tgt->refcount > 1)
|
||||
k->tgt->refcount--;
|
||||
else
|
||||
{
|
||||
is_tgt_unmapped = true;
|
||||
gomp_unmap_tgt (k->tgt);
|
||||
}
|
||||
return is_tgt_unmapped;
|
||||
}
|
||||
|
||||
/* Unmap variables described by TGT. If DO_COPYFROM is true, copy relevant
|
||||
variables back from device to host: if it is false, it is assumed that this
|
||||
has been done already. */
|
||||
@ -1059,16 +1077,7 @@ gomp_unmap_vars (struct target_mem_desc *tgt, bool do_copyfrom)
|
||||
+ tgt->list[i].offset),
|
||||
tgt->list[i].length);
|
||||
if (do_unmap)
|
||||
{
|
||||
splay_tree_remove (&devicep->mem_map, k);
|
||||
if (k->link_key)
|
||||
splay_tree_insert (&devicep->mem_map,
|
||||
(splay_tree_node) k->link_key);
|
||||
if (k->tgt->refcount > 1)
|
||||
k->tgt->refcount--;
|
||||
else
|
||||
gomp_unmap_tgt (k->tgt);
|
||||
}
|
||||
gomp_remove_var (devicep, k);
|
||||
}
|
||||
|
||||
if (tgt->refcount > 1)
|
||||
@ -1298,17 +1307,7 @@ gomp_unload_image_from_device (struct gomp_device_descr *devicep,
|
||||
else
|
||||
{
|
||||
splay_tree_key n = splay_tree_lookup (&devicep->mem_map, &k);
|
||||
splay_tree_remove (&devicep->mem_map, n);
|
||||
if (n->link_key)
|
||||
{
|
||||
if (n->tgt->refcount > 1)
|
||||
n->tgt->refcount--;
|
||||
else
|
||||
{
|
||||
is_tgt_unmapped = true;
|
||||
gomp_unmap_tgt (n->tgt);
|
||||
}
|
||||
}
|
||||
is_tgt_unmapped = gomp_remove_var (devicep, n);
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
@ -18,5 +18,3 @@ main (int argc, char *argv[])
|
||||
}
|
||||
|
||||
/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
|
||||
/* { dg-output "Trying to map into device \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) object when \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) is already mapped" } */
|
||||
/* { dg-shouldfail "" } */
|
||||
|
||||
@ -18,5 +18,3 @@ main (int argc, char *argv[])
|
||||
}
|
||||
|
||||
/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
|
||||
/* { dg-output "Trying to map into device \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) object when \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) is already mapped" } */
|
||||
/* { dg-shouldfail "" } */
|
||||
|
||||
@ -18,5 +18,3 @@ main (int argc, char *argv[])
|
||||
}
|
||||
|
||||
/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
|
||||
/* { dg-output "already mapped to" } */
|
||||
/* { dg-shouldfail "" } */
|
||||
|
||||
@ -16,5 +16,3 @@ main (int argc, char *argv[])
|
||||
}
|
||||
|
||||
/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
|
||||
/* { dg-output "already mapped to" } */
|
||||
/* { dg-shouldfail "" } */
|
||||
|
||||
@ -16,5 +16,3 @@ main (int argc, char *argv[])
|
||||
}
|
||||
|
||||
/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
|
||||
/* { dg-output "already mapped to" } */
|
||||
/* { dg-shouldfail "" } */
|
||||
|
||||
@ -16,5 +16,3 @@ main (int argc, char *argv[])
|
||||
}
|
||||
|
||||
/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
|
||||
/* { dg-output "already mapped to" } */
|
||||
/* { dg-shouldfail "" } */
|
||||
|
||||
@ -16,5 +16,3 @@ main (int argc, char *argv[])
|
||||
}
|
||||
|
||||
/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
|
||||
/* { dg-output "already mapped to" } */
|
||||
/* { dg-shouldfail "" } */
|
||||
|
||||
@ -18,5 +18,3 @@ main (int argc, char *argv[])
|
||||
}
|
||||
|
||||
/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
|
||||
/* { dg-output "Trying to map into device \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) object when \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) is already mapped" } */
|
||||
/* { dg-shouldfail "" } */
|
||||
|
||||
@ -1,8 +1,5 @@
|
||||
/* Test if duplicate data mappings with acc_copy_in. */
|
||||
/* Test if acc_copyin has present_or_ and reference counting behavior. */
|
||||
|
||||
/* { dg-do run { target openacc_nvidia_accel_selected } } */
|
||||
|
||||
#include <stdio.h>
|
||||
#include <stdlib.h>
|
||||
#include <openacc.h>
|
||||
|
||||
@ -21,15 +18,21 @@ main (int argc, char **argv)
|
||||
}
|
||||
|
||||
(void) acc_copyin (h, N);
|
||||
|
||||
fprintf (stderr, "CheCKpOInT\n");
|
||||
(void) acc_copyin (h, N);
|
||||
|
||||
acc_copyout (h, N);
|
||||
|
||||
if (!acc_is_present (h, N))
|
||||
abort ();
|
||||
|
||||
acc_copyout (h, N);
|
||||
|
||||
#if !ACC_MEM_SHARED
|
||||
if (acc_is_present (h, N))
|
||||
abort ();
|
||||
#endif
|
||||
|
||||
free (h);
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
|
||||
/* { dg-output "\\\[\[0-9a-fA-FxX\]+,\\\+256\\\] already mapped to \\\[\[0-9a-fA-FxX\]+,\\\+256\\\]" } */
|
||||
/* { dg-shouldfail "" } */
|
||||
|
||||
@ -1,8 +1,5 @@
|
||||
/* Exercise acc_create and acc_delete on nvidia targets. */
|
||||
/* Exercise acc_create and acc_delete. */
|
||||
|
||||
/* { dg-do run { target openacc_nvidia_accel_selected } } */
|
||||
|
||||
#include <stdio.h>
|
||||
#include <stdlib.h>
|
||||
#include <openacc.h>
|
||||
|
||||
@ -19,18 +16,23 @@ main (int argc, char **argv)
|
||||
if (!d)
|
||||
abort ();
|
||||
|
||||
fprintf (stderr, "CheCKpOInT\n");
|
||||
d = acc_create (h, N);
|
||||
if (!d)
|
||||
abort ();
|
||||
|
||||
acc_delete (h, N);
|
||||
|
||||
if (!acc_is_present (h, N))
|
||||
abort ();
|
||||
|
||||
acc_delete (h, N);
|
||||
|
||||
#if !ACC_MEM_SHARED
|
||||
if (acc_is_present (h, N))
|
||||
abort ();
|
||||
#endif
|
||||
|
||||
free (h);
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */
|
||||
/* { dg-output "\\\[\[0-9a-fA-FxX\]+,\\\+256\\\] already mapped to \\\[\[0-9a-fA-FxX\]+,\\\+256\\\]" } */
|
||||
/* { dg-shouldfail "" } */
|
||||
|
||||
@ -127,7 +127,7 @@ main (int argc, char **argv)
|
||||
h[i] = i + 10;
|
||||
}
|
||||
|
||||
acc_copyout (h, S);
|
||||
acc_copyout_finalize (h, S);
|
||||
d = NULL;
|
||||
if (!shared_mem)
|
||||
if (acc_is_present (h, S))
|
||||
@ -236,7 +236,7 @@ main (int argc, char **argv)
|
||||
abort ();
|
||||
}
|
||||
|
||||
acc_delete (h, S);
|
||||
acc_delete_finalize (h, S);
|
||||
d = NULL;
|
||||
if (!shared_mem)
|
||||
if (acc_is_present (h, S))
|
||||
|
||||
@ -5,21 +5,19 @@
|
||||
#include <stdlib.h>
|
||||
#include <unistd.h>
|
||||
#include <openacc.h>
|
||||
#include "timer.h"
|
||||
#include <cuda.h>
|
||||
#include <sys/time.h>
|
||||
|
||||
int
|
||||
main (int argc, char **argv)
|
||||
{
|
||||
float atime;
|
||||
CUstream stream;
|
||||
CUresult r;
|
||||
struct timeval tv1, tv2;
|
||||
time_t t1;
|
||||
|
||||
acc_init (acc_device_nvidia);
|
||||
|
||||
(void) acc_get_device_num (acc_device_nvidia);
|
||||
|
||||
init_timers (1);
|
||||
|
||||
stream = (CUstream) acc_get_cuda_stream (0);
|
||||
if (stream != NULL)
|
||||
abort ();
|
||||
@ -34,22 +32,22 @@ main (int argc, char **argv)
|
||||
if (!acc_set_cuda_stream (0, stream))
|
||||
abort ();
|
||||
|
||||
start_timer (0);
|
||||
gettimeofday (&tv1, NULL);
|
||||
|
||||
acc_wait_all_async (0);
|
||||
|
||||
acc_wait (0);
|
||||
|
||||
atime = stop_timer (0);
|
||||
gettimeofday (&tv2, NULL);
|
||||
|
||||
if (0.010 < atime)
|
||||
t1 = ((tv2.tv_sec - tv1.tv_sec) * 1000000) + (tv2.tv_usec - tv1.tv_usec);
|
||||
|
||||
if (t1 > 1000)
|
||||
{
|
||||
fprintf (stderr, "actual time too long\n");
|
||||
fprintf (stderr, "too long\n");
|
||||
abort ();
|
||||
}
|
||||
|
||||
fini_timers ();
|
||||
|
||||
acc_shutdown (acc_device_nvidia);
|
||||
|
||||
exit (0);
|
||||
|
||||
56
libgomp/testsuite/libgomp.oacc-fortran/data-5.f90
Normal file
56
libgomp/testsuite/libgomp.oacc-fortran/data-5.f90
Normal file
@ -0,0 +1,56 @@
|
||||
! { dg-do run }
|
||||
! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
|
||||
|
||||
program refcount_test
|
||||
use openacc
|
||||
integer, allocatable :: h(:)
|
||||
integer i, N
|
||||
|
||||
N = 256
|
||||
allocate (h(N))
|
||||
|
||||
do i = 1, N
|
||||
h(i) = i
|
||||
end do
|
||||
|
||||
!$acc enter data create (h(1:N))
|
||||
!$acc enter data copyin (h(1:N))
|
||||
!$acc enter data copyin (h(1:N))
|
||||
!$acc enter data copyin (h(1:N))
|
||||
|
||||
call acc_update_self (h)
|
||||
do i = 1, N
|
||||
if (h(i) .eq. i) c = c + 1
|
||||
end do
|
||||
! h[] should be filled with uninitialized device values,
|
||||
! abort if it's not.
|
||||
if (c .eq. N) call abort
|
||||
|
||||
h(:) = 0
|
||||
|
||||
!$acc parallel present (h(1:N))
|
||||
do i = 1, N
|
||||
h(i) = 111
|
||||
end do
|
||||
!$acc end parallel
|
||||
|
||||
! No actual copyout should happen.
|
||||
call acc_copyout (h)
|
||||
do i = 1, N
|
||||
if (h(i) .ne. 0) call abort
|
||||
end do
|
||||
|
||||
!$acc exit data delete (h(1:N))
|
||||
|
||||
! This should not actually be deleted yet.
|
||||
if (acc_is_present (h) .eqv. .FALSE.) call abort
|
||||
|
||||
!$acc exit data copyout (h(1:N)) finalize
|
||||
|
||||
do i = 1, N
|
||||
if (h(i) .ne. 111) call abort
|
||||
end do
|
||||
|
||||
if (acc_is_present (h) .eqv. .TRUE.) call abort
|
||||
|
||||
end program refcount_test
|
||||
@ -14,5 +14,3 @@
|
||||
END
|
||||
|
||||
! { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
|
||||
! { dg-output "Trying to map into device \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) object when \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) is already mapped" }
|
||||
! { dg-shouldfail "" }
|
||||
|
||||
@ -14,5 +14,3 @@
|
||||
END
|
||||
|
||||
! { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
|
||||
! { dg-output "Trying to map into device \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) object when \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) is already mapped" }
|
||||
! { dg-shouldfail "" }
|
||||
|
||||
@ -13,5 +13,3 @@
|
||||
END
|
||||
|
||||
! { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
|
||||
! { dg-output "already mapped to" }
|
||||
! { dg-shouldfail "" }
|
||||
|
||||
@ -12,5 +12,3 @@
|
||||
END
|
||||
|
||||
! { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
|
||||
! { dg-output "already mapped to" }
|
||||
! { dg-shouldfail "" }
|
||||
|
||||
@ -12,5 +12,3 @@
|
||||
END
|
||||
|
||||
! { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
|
||||
! { dg-output "already mapped to" }
|
||||
! { dg-shouldfail "" }
|
||||
|
||||
@ -12,5 +12,3 @@
|
||||
END
|
||||
|
||||
! { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
|
||||
! { dg-output "already mapped to" }
|
||||
! { dg-shouldfail "" }
|
||||
|
||||
@ -12,5 +12,3 @@
|
||||
END
|
||||
|
||||
! { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
|
||||
! { dg-output "already mapped to" }
|
||||
! { dg-shouldfail "" }
|
||||
|
||||
@ -14,5 +14,3 @@
|
||||
END
|
||||
|
||||
! { dg-output "CheCKpOInT(\n|\r\n|\r).*" }
|
||||
! { dg-output "Trying to map into device \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) object when \\\[\[0-9a-fA-FxX\]+..\[0-9a-fA-FxX\]+\\\) is already mapped" }
|
||||
! { dg-shouldfail "" }
|
||||
|
||||
@ -90,7 +90,7 @@
|
||||
H(I) = I + 10
|
||||
END DO
|
||||
|
||||
CALL ACC_COPYOUT (H, INT (SIZEOF (H), 4))
|
||||
CALL ACC_COPYOUT_FINALIZE (H, INT (SIZEOF (H), 4))
|
||||
IF (.NOT. SHARED_MEM) THEN
|
||||
IF (ACC_IS_PRESENT (H, INT (SIZEOF (H), 8))) STOP 11
|
||||
ENDIF
|
||||
@ -163,7 +163,7 @@
|
||||
IF (H(I) .NE. I + MERGE (18, 17, SHARED_MEM)) STOP 23
|
||||
END DO
|
||||
|
||||
CALL ACC_DELETE (H)
|
||||
CALL ACC_DELETE_FINALIZE (H)
|
||||
IF (.NOT. SHARED_MEM) THEN
|
||||
IF (ACC_IS_PRESENT (H, INT (SIZEOF (H), 4))) STOP 24
|
||||
ENDIF
|
||||
|
||||
@ -90,7 +90,7 @@
|
||||
H(I) = I + 10
|
||||
END DO
|
||||
|
||||
CALL ACC_COPYOUT (H, INT (SIZEOF (H), 4))
|
||||
CALL ACC_COPYOUT_FINALIZE (H, INT (SIZEOF (H), 4))
|
||||
IF (.NOT. SHARED_MEM) THEN
|
||||
IF (ACC_IS_PRESENT (H, INT (SIZEOF (H), 8))) STOP 11
|
||||
ENDIF
|
||||
@ -163,7 +163,7 @@
|
||||
IF (H(I) .NE. I + MERGE (18, 17, SHARED_MEM)) STOP 23
|
||||
END DO
|
||||
|
||||
CALL ACC_DELETE (H)
|
||||
CALL ACC_DELETE_FINALIZE (H)
|
||||
IF (.NOT. SHARED_MEM) THEN
|
||||
IF (ACC_IS_PRESENT (H, INT (SIZEOF (H), 4))) STOP 24
|
||||
ENDIF
|
||||
|
||||
Loading…
x
Reference in New Issue
Block a user