From: Gergö Barany Date: Fri, 21 Dec 2018 09:12:44 +0000 (-0800) Subject: Add OpenACC 2.6 if and if_present clauses on host_data construct: GOACC_FLAG_HOST_DAT... X-Git-Url: http://git.ipfire.org/cgi-bin/gitweb.cgi?a=commitdiff_plain;h=9c306c69a2f6f60f9a06c208106adf259c1351cf;p=thirdparty%2Fgcc.git Add OpenACC 2.6 if and if_present clauses on host_data construct: GOACC_FLAG_HOST_DATA_IF_PRESENT gcc/c/ * c-parser.c (OACC_HOST_DATA_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_IF and PRAGMA_OACC_CLAUSE_IF_PRESENT. gcc/cp/ * parser.c (OACC_HOST_DATA_CLAUSE_MASK): Likewise. gcc/fortran/ * openmp.c (OACC_HOST_DATA_CLAUSES): Add OMP_CLAUSE_IF and OMP_CLAUSE_IF_PRESENT. gcc/ * omp-expand.c (expand_omp_target): Handle if_present flag on OpenACC host_data construct. gcc/testsuite/ * c-c++-common/goacc/host_data-1.c: Add tests of if and if_present clauses on host_data. * gfortran.dg/goacc/host_data-tree.f95: Likewise. include/ * gomp-constants.h (GOACC_FLAG_HOST_DATA_IF_PRESENT): New constant. libgomp/ * libgomp.h (enum gomp_map_vars_kind): Add GOMP_MAP_VARS_OPENACC_IF_PRESENT. * oacc-parallel.c (GOACC_data_start): Handle GOACC_FLAG_HOST_DATA_IF_PRESENT flag. * target.c (gomp_map_vars_async): Handle GOMP_MAP_VARS_OPENACC_IF_PRESENT mapping kind. * testsuite/libgomp.oacc-c-c++-common/host_data-6.c: New test. Reviewed-by: Thomas Schwinge (cherry picked from openacc-gcc-9-branch commit b52c8d006581b2a657ac3d81ab2527bb398b6615) --- diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp index 797a62e1af29..78f4042c5e24 100644 --- a/gcc/ChangeLog.omp +++ b/gcc/ChangeLog.omp @@ -1,3 +1,8 @@ +2018-12-21 Gergö Barany + + * omp-expand.c (expand_omp_target): Handle if_present flag on + OpenACC host_data construct. + 2018-12-20 Gergö Barany * omp-low.c (struct omp_context): New fields diff --git a/gcc/c/ChangeLog.omp b/gcc/c/ChangeLog.omp index db92b2422816..e085e879f673 100644 --- a/gcc/c/ChangeLog.omp +++ b/gcc/c/ChangeLog.omp @@ -1,3 +1,8 @@ +2018-12-21 Gergö Barany + + * c-parser.c (OACC_HOST_DATA_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_IF + and PRAGMA_OACC_CLAUSE_IF_PRESENT. + 2018-12-13 Cesar Philippidis Nathan Sidwell Julian Brown diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index 0e49167381d3..d11216472d7b 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -15616,7 +15616,9 @@ c_parser_oacc_enter_exit_data (c_parser *parser, bool enter) */ #define OACC_HOST_DATA_CLAUSE_MASK \ - ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_USE_DEVICE) ) + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_USE_DEVICE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF_PRESENT) ) static tree c_parser_oacc_host_data (location_t loc, c_parser *parser, bool *if_p) diff --git a/gcc/cp/ChangeLog.omp b/gcc/cp/ChangeLog.omp index 370d895d2c5f..b2c7cfc20dbd 100644 --- a/gcc/cp/ChangeLog.omp +++ b/gcc/cp/ChangeLog.omp @@ -1,3 +1,7 @@ +2018-12-21 Gergö Barany + + * parser.c (OACC_HOST_DATA_CLAUSE_MASK): Likewise. + 2018-12-13 Cesar Philippidis Nathan Sidwell Julian Brown diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index b8de1427650a..3c42f43742ed 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -38767,7 +38767,9 @@ cp_parser_oacc_data (cp_parser *parser, cp_token *pragma_tok, bool *if_p) structured-block */ #define OACC_HOST_DATA_CLAUSE_MASK \ - ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_USE_DEVICE) ) + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_USE_DEVICE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF_PRESENT) ) static tree cp_parser_oacc_host_data (cp_parser *parser, cp_token *pragma_tok, bool *if_p) diff --git a/gcc/fortran/ChangeLog.omp b/gcc/fortran/ChangeLog.omp index b01f1ba83b53..00ccd26462c4 100644 --- a/gcc/fortran/ChangeLog.omp +++ b/gcc/fortran/ChangeLog.omp @@ -1,3 +1,8 @@ +2018-12-21 Gergö Barany + + * openmp.c (OACC_HOST_DATA_CLAUSES): Add OMP_CLAUSE_IF and + OMP_CLAUSE_IF_PRESENT. + 2019-01-30 Kwok Cheung Yeung * trans-openmp.c (gfc_build_conditional_assign): New. diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c index 81b69b9c7af7..2a6dce7d65f8 100644 --- a/gcc/fortran/openmp.c +++ b/gcc/fortran/openmp.c @@ -2001,7 +2001,10 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, (OACC_LOOP_CLAUSES | OACC_PARALLEL_CLAUSES) #define OACC_KERNELS_LOOP_CLAUSES \ (OACC_LOOP_CLAUSES | OACC_KERNELS_CLAUSES) -#define OACC_HOST_DATA_CLAUSES omp_mask (OMP_CLAUSE_USE_DEVICE) +#define OACC_HOST_DATA_CLAUSES \ + (omp_mask (OMP_CLAUSE_USE_DEVICE) \ + | OMP_CLAUSE_IF \ + | OMP_CLAUSE_IF_PRESENT) #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 \ diff --git a/gcc/omp-expand.c b/gcc/omp-expand.c index 27de3be190d6..6dd87689e5ab 100644 --- a/gcc/omp-expand.c +++ b/gcc/omp-expand.c @@ -7591,7 +7591,17 @@ expand_omp_target (struct omp_region *region) if (is_gimple_omp_oacc (entry_stmt)) { /* By default, no GOACC_FLAGs are set. */ - goacc_flags = integer_zero_node; + int goacc_flags_i = 0; + + if (start_ix != BUILT_IN_GOACC_UPDATE + && omp_find_clause (clauses, OMP_CLAUSE_IF_PRESENT)) + { + gcc_checking_assert (gimple_omp_target_kind (entry_stmt) + == GF_OMP_TARGET_KIND_OACC_HOST_DATA); + goacc_flags_i |= GOACC_FLAG_HOST_DATA_IF_PRESENT; + } + + goacc_flags = build_int_cst (integer_type_node, goacc_flags_i); } else { diff --git a/gcc/testsuite/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp index 10b52bee15d0..c9917439f699 100644 --- a/gcc/testsuite/ChangeLog.omp +++ b/gcc/testsuite/ChangeLog.omp @@ -1,3 +1,9 @@ +2018-12-21 Gergö Barany + + * c-c++-common/goacc/host_data-1.c: Add tests of if and if_present + clauses on host_data. + * gfortran.dg/goacc/host_data-tree.f95: Likewise. + 2018-12-20 Gergö Barany * c-c++-common/goacc/nested-reductions-fail.c: New test. diff --git a/gcc/testsuite/c-c++-common/goacc/host_data-1.c b/gcc/testsuite/c-c++-common/goacc/host_data-1.c index 0c7a857d0bf3..658b7a677bc7 100644 --- a/gcc/testsuite/c-c++-common/goacc/host_data-1.c +++ b/gcc/testsuite/c-c++-common/goacc/host_data-1.c @@ -7,6 +7,9 @@ f (void) { #pragma acc host_data use_device(v1) ; + +#pragma acc host_data use_device(v1) if_present + ; } @@ -16,9 +19,32 @@ void foo (float *x, float *y) { int n = 1 << 10; -#pragma acc data create(x[0:n]) copyout(y[0:n]) +#pragma acc data create(x[0:n]) { + bar (x, y); + + /* This should fail at run time because y is not mapped. */ #pragma acc host_data use_device(x,y) bar (x, y); + + /* y is still not mapped, but this should not fail at run time but + continue execution with y remaining as the host address. */ +#pragma acc host_data use_device(x,y) if_present + bar (x, y); + +#pragma acc data copyout(y[0:n]) + { +#pragma acc host_data use_device(x,y) + bar (x, y); + +#pragma acc host_data use_device(x,y) if_present + bar (x, y); + +#pragma acc host_data use_device(x,y) if(x != y) + bar (x, y); + +#pragma acc host_data use_device(x,y) if_present if(x != y) + bar (x, y); + } } } diff --git a/gcc/testsuite/gfortran.dg/goacc/host_data-tree.f95 b/gcc/testsuite/gfortran.dg/goacc/host_data-tree.f95 index d44ca5870516..2ac1c0d66d62 100644 --- a/gcc/testsuite/gfortran.dg/goacc/host_data-tree.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/host_data-tree.f95 @@ -7,5 +7,15 @@ program test !$acc host_data use_device(p) !$acc end host_data + + !$acc host_data use_device(p) if (p == 42) + !$acc end host_data + + !$acc host_data use_device(p) if_present if (p == 43) + !$acc end host_data end program test -! { dg-final { scan-tree-dump-times "pragma acc host_data use_device_ptr\\(p\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "pragma acc host_data use_device_ptr\\(p\\)" 3 "original" } } +! { dg-final { scan-tree-dump-times "D.\[0-9\]+ = \\*p == 42;" 1 "original" } } +! { dg-final { scan-tree-dump-times "pragma acc host_data use_device_ptr\\(p\\) if\\(D.\[0-9\]+\\)" 2 "original" } } +! { dg-final { scan-tree-dump-times "D.\[0-9\]+ = \\*p == 43;" 1 "original" } } +! { dg-final { scan-tree-dump-times "pragma acc host_data use_device_ptr\\(p\\) if\\(D.\[0-9\]+\\) if_present" 1 "original" } } diff --git a/include/ChangeLog.omp b/include/ChangeLog.omp index 3cc2e9ace8de..e632fdb4ebf9 100644 --- a/include/ChangeLog.omp +++ b/include/ChangeLog.omp @@ -1,3 +1,7 @@ +2018-12-21 Gergö Barany + + * gomp-constants.h (GOACC_FLAG_HOST_DATA_IF_PRESENT): New constant. + 2018-10-04 Cesar Philippidis Julian Brown diff --git a/include/gomp-constants.h b/include/gomp-constants.h index 2bbba6117123..90b18091a135 100644 --- a/include/gomp-constants.h +++ b/include/gomp-constants.h @@ -240,6 +240,8 @@ enum gomp_map_kind /* Force host fallback execution. */ #define GOACC_FLAG_HOST_FALLBACK (1 << 0) +/* "if_present" semantics for OpenACC "host_data" constructs. */ +#define GOACC_FLAG_HOST_DATA_IF_PRESENT (1 << 1) /* For legacy reasons, in the ABI, the GOACC_FLAGs are encoded as an inverted bitmask. */ diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp index 7f1642187843..5dc06978255f 100644 --- a/libgomp/ChangeLog.omp +++ b/libgomp/ChangeLog.omp @@ -1,3 +1,13 @@ +2018-12-21 Gergö Barany + + * libgomp.h (enum gomp_map_vars_kind): Add + GOMP_MAP_VARS_OPENACC_IF_PRESENT. + * oacc-parallel.c (GOACC_data_start): Handle + GOACC_FLAG_HOST_DATA_IF_PRESENT flag. + * target.c (gomp_map_vars_async): Handle + GOMP_MAP_VARS_OPENACC_IF_PRESENT mapping kind. + * testsuite/libgomp.oacc-c-c++-common/host_data-6.c: New test. + 2018-12-20 Gergö Barany * testsuite/libgomp.oacc-c-c++-common/par-loop-comb-reduction-1.c: diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h index cf3682a09eee..6bc39ce9aae4 100644 --- a/libgomp/libgomp.h +++ b/libgomp/libgomp.h @@ -1082,6 +1082,9 @@ struct gomp_device_descr enum gomp_map_vars_kind { GOMP_MAP_VARS_OPENACC, + /* Like "GOMP_MAP_VARS_OPENACC", but with "GOACC_FLAG_HOST_DATA_IF_PRESENT" + semantics. */ + GOMP_MAP_VARS_OPENACC_IF_PRESENT, GOMP_MAP_VARS_OPENACC_ENTER_DATA, GOMP_MAP_VARS_TARGET, GOMP_MAP_VARS_DATA, diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c index 3369a3656b7a..98302f1f7adb 100644 --- a/libgomp/oacc-parallel.c +++ b/libgomp/oacc-parallel.c @@ -331,12 +331,17 @@ GOACC_data_start (int flags_m, size_t mapnum, handle_ftn_pointers (mapnum, hostaddrs, sizes, kinds); + enum gomp_map_vars_kind pragma_kind; + if (flags & GOACC_FLAG_HOST_DATA_IF_PRESENT) + pragma_kind = GOMP_MAP_VARS_OPENACC_IF_PRESENT; + else + pragma_kind = GOMP_MAP_VARS_OPENACC; + /* Host fallback or 'do nothing'. */ if ((acc_dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) || (flags & GOACC_FLAG_HOST_FALLBACK)) { - tgt = gomp_map_vars (NULL, 0, NULL, NULL, NULL, NULL, true, - GOMP_MAP_VARS_OPENACC); + tgt = gomp_map_vars (NULL, 0, NULL, NULL, NULL, NULL, true, pragma_kind); tgt->prev = thr->mapped_data; thr->mapped_data = tgt; @@ -345,7 +350,7 @@ GOACC_data_start (int flags_m, size_t mapnum, gomp_debug (0, " %s: prepare mappings\n", __FUNCTION__); tgt = gomp_map_vars (acc_dev, mapnum, hostaddrs, NULL, sizes, kinds, true, - GOMP_MAP_VARS_OPENACC); + pragma_kind); gomp_debug (0, " %s: mappings prepared\n", __FUNCTION__); tgt->prev = thr->mapped_data; thr->mapped_data = tgt; diff --git a/libgomp/target.c b/libgomp/target.c index 320942742b35..b34043f31423 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -1102,6 +1102,9 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, splay_tree_key n = gomp_map_lookup (mem_map, &cur_node); if (n == NULL) { + if (pragma_kind == GOMP_MAP_VARS_OPENACC_IF_PRESENT) + /* No error, continue using the host address. */ + continue; gomp_mutex_unlock (&devicep->lock); gomp_fatal ("use_device_ptr pointer wasn't mapped"); } diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-6.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-6.c new file mode 100644 index 000000000000..c5744feca58f --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-6.c @@ -0,0 +1,53 @@ +/* Test if, if_present clauses on host_data construct. */ + +#include +#include + +void +foo (float *p, intptr_t host_p, int shared_mem_p, int cond) +{ + assert (p == (float *) host_p); + +#pragma acc data copyin(host_p) + { +#pragma acc host_data use_device(p) if_present + /* p not mapped yet, so it will be equal to the host pointer. */ + assert (p == (float *) host_p); + +#pragma acc data copy(p[0:100]) + { + /* Not inside a host_data construct, so p is still the host pointer. */ + assert (p == (float *) host_p); + + if (!shared_mem_p) + { +#pragma acc host_data use_device(p) + /* The device address is different from the host address. */ + assert (p != (float *) host_p); + +#pragma acc host_data use_device(p) if_present + /* p is present now, so this is the same as above. */ + assert (p != (float *) host_p); + } + +#pragma acc host_data use_device(p) if(cond) + /* p is the device pointer iff cond is true and device memory is + separate from host memory. */ + assert ((p != (float *) host_p) == (cond && !shared_mem_p)); + } + } +} + +int +main (void) +{ + float arr[100]; + int shared_mem_p = 0; +#if ACC_MEM_SHARED + shared_mem_p = 1; +#endif + foo (arr, (intptr_t) arr, shared_mem_p, 0); + foo (arr, (intptr_t) arr, shared_mem_p, 1); + + return 0; +}