2017-05-04 Cesar Philippidis <cesar@codesourcery.com>
gcc/c-family/
* c-pragma.h (enum pragma_omp_clause): Add
PRAGMA_OACC_CLAUSE_IF_PRESENT.
gcc/c/
* c-parser.c (c_parser_omp_clause_name): Add support for if_present.
(c_parser_oacc_all_clauses): Likewise.
(c_finish_oacc_routine): Likewise.
(OACC_UPDATE_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_IF_PRESENT.
* c-typeck.c (c_finish_omp_clauses): Add support for if_present.
gcc/cp/
* parser.c (cp_parser_omp_clause_name): Add support for if_present.
(cp_parser_oacc_all_clauses): Likewise.
(cp_parser_oacc_kernels_parallel): Likewise.
(OACC_UPDATE_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_IF_PRESENT.
* pt.c (tsubst_omp_clauses): Add support for if_present.
* semantics.c (finish_omp_clauses): Likewise.
gcc/fortran/
* gfortran.h (gfc_omp_clauses): Add if_present member.
* openmp.c (enum omp_mask2): Add OMP_CLAUSE_IF_PRESENT.
(gfc_match_omp_clauses): Handle it.
(OACC_UPDATE_CLAUSES): Add OMP_CLAUSE_IF_PRESENT.
* trans-openmp.c (gfc_trans_omp_clauses_1): Generate an omp clause for
if_present.
gcc/
* gimplify.c (gimplify_scan_omp_clauses): Handle OMP_CLAUSE_IF_PRESENT.
(gimplify_adjust_omp_clauses): Likewise.
* omp-low.c (scan_sharing_clauses): Likewise, but just ignore it for
now.
* tree-pretty-print.c (dump_omp_clause): Likewise.
* tree.c (omp_clause_num_ops): Add an entry for OMP_CLAUSE_IF_PRESENT.
(omp_clause_code_name): Likewise.
* tree-core.h (enum omp_clause_code): Likewise.
gcc/testsuite/
* 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/update-1.C: New test.
* gfortran.dg/goacc/update-if_present-1.f90: New test.
* gfortran.dg/goacc/update-if_present-2.f90: New test.
@@ -171,6 +171,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,
@@ -10393,7 +10393,9 @@ c_parser_omp_clause_name (c_parser *parser, bool consume_token = true)
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;
@@ -13268,6 +13270,12 @@ 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, here,
+ OMP_CLAUSE_IF_PRESENT,
+ clauses);
+ c_name = "if_present";
+ break;
case PRAGMA_OACC_CLAUSE_INDEPENDENT:
clauses = c_parser_oacc_simple_clause (parser, here,
OMP_CLAUSE_INDEPENDENT,
@@ -14344,6 +14352,7 @@ c_finish_oacc_routine (struct oacc_routine_data *data, tree fndecl,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_TYPE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_HOST) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF_PRESENT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
#define OACC_UPDATE_CLAUSE_DEVICE_TYPE_MASK \
@@ -13396,6 +13396,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
case OMP_CLAUSE_BIND:
case OMP_CLAUSE_NOHOST:
case OMP_CLAUSE_TILE:
+ case OMP_CLAUSE_IF_PRESENT:
pc = &OMP_CLAUSE_CHAIN (c);
continue;
@@ -29833,7 +29833,9 @@ cp_parser_omp_clause_name (cp_parser *parser, bool consume_token = true)
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;
@@ -32406,6 +32408,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,
@@ -35821,6 +35829,7 @@ cp_parser_oacc_kernels_parallel (cp_parser *parser, cp_token *pragma_tok,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICE_TYPE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_HOST) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF_PRESENT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT))
#define OACC_UPDATE_CLAUSE_DEVICE_TYPE_MASK \
@@ -14750,6 +14750,7 @@ 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_DEVICE_TYPE:
break;
case OMP_CLAUSE_BIND:
@@ -7106,6 +7106,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
case OMP_CLAUSE_SEQ:
case OMP_CLAUSE_BIND:
case OMP_CLAUSE_NOHOST:
+ case OMP_CLAUSE_IF_PRESENT:
break;
case OMP_CLAUSE_TILE:
@@ -1318,6 +1318,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, nohost:1, acc_collapse:1, bind:1;
+ unsigned if_present:1;
locus loc;
char bind_name[GFC_MAX_SYMBOL_LEN+1];
}
@@ -833,6 +833,7 @@ enum omp_mask2
OMP_CLAUSE_TILE,
OMP_CLAUSE_BIND,
OMP_CLAUSE_NOHOST,
+ OMP_CLAUSE_IF_PRESENT,
OMP_CLAUSE_DEVICE_TYPE,
/* This must come last. */
OMP_MASK2_LAST
@@ -1377,6 +1378,14 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, 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
@@ -2064,7 +2073,8 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, omp_mask mask,
| 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_TYPE)
+ | OMP_CLAUSE_DEVICE | OMP_CLAUSE_WAIT | OMP_CLAUSE_IF_PRESENT \
+ | OMP_CLAUSE_DEVICE_TYPE)
#define OACC_ENTER_DATA_CLAUSES \
(omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_WAIT \
| OMP_CLAUSE_COPYIN | OMP_CLAUSE_CREATE)
@@ -2931,6 +2931,11 @@ gfc_trans_omp_clauses_1 (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->independent)
{
c = build_omp_clause (where.lb->location, OMP_CLAUSE_INDEPENDENT);
@@ -7667,6 +7667,7 @@ 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_DEVICE_TYPE:
break;
@@ -8525,6 +8526,7 @@ 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_DEVICE_TYPE:
break;
@@ -2429,6 +2429,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx,
case OMP_CLAUSE_AUTO:
case OMP_CLAUSE_SEQ:
case OMP_CLAUSE_TILE:
+ case OMP_CLAUSE_IF_PRESENT:
case OMP_CLAUSE_DEVICE_TYPE:
break;
@@ -2603,6 +2604,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx,
case OMP_CLAUSE_SEQ:
case OMP_CLAUSE_TILE:
case OMP_CLAUSE__GRIDDIM_:
+ case OMP_CLAUSE_IF_PRESENT:
case OMP_CLAUSE_DEVICE_TYPE:
break;
new file mode 100644
@@ -0,0 +1,17 @@
+/* Test valid usages of the if_present clause. */
+
+/* { dg-compile } */
+/* { 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
+}
+
+/* { dg-final { scan-tree-dump-times "pragma omp target oacc_update if_present" 4 "omplower" } } */
new file mode 100644
@@ -0,0 +1,44 @@
+/* Test invalid usages of the if_present clause. */
+
+/* { dg-compile } */
+
+#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) device_type(nvidia) if_present /* { dg-error "'if_present' is not valid" } */
+#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++)
+ ;
+}
+
new file mode 100644
@@ -0,0 +1,18 @@
+/* Test valid usages of the if_present clause. */
+
+/* { dg-compile } */
+/* { dg-additional-options "-fdump-tree-omplower" } */
+
+template<typename T>
+void
+t ()
+{
+ T 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
+}
+
+/* { dg-final { scan-tree-dump-times "pragma omp target oacc_update if_present" 4 "omplower" } } */
new file mode 100644
@@ -0,0 +1,28 @@
+! Test valid usages of the if_present clause.
+
+! { dg-compile }
+! { 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" } }
new file mode 100644
@@ -0,0 +1,54 @@
+! Test invalid usages of the if_present clause.
+
+! { dg-compile }
+
+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
@@ -469,6 +469,9 @@ enum omp_clause_code {
kernel. */
OMP_CLAUSE__GRIDDIM_,
+ /* OpenACC clause: if_present. */
+ OMP_CLAUSE_IF_PRESENT,
+
/* OpenACC clause: device_type ( device-type-list). */
OMP_CLAUSE_DEVICE_TYPE
};
@@ -1079,7 +1079,6 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, int flags)
spc, flags, false);
pp_string (pp, ")");
break;
-
case OMP_CLAUSE__GRIDDIM_:
pp_string (pp, "_griddim_(");
pp_unsigned_wide_integer (pp, OMP_CLAUSE__GRIDDIM__DIMENSION (clause));
@@ -1091,6 +1090,9 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, int flags)
false);
pp_right_paren (pp);
break;
+ case OMP_CLAUSE_IF_PRESENT:
+ pp_string (pp, "if_present");
+ break;
default:
pp_string (pp, "unknown");
@@ -330,6 +330,7 @@ unsigned const char omp_clause_num_ops[] =
0, /* OMP_CLAUSE_NOHOST */
3, /* OMP_CLAUSE_TILE */
2, /* OMP_CLAUSE__GRIDDIM_ */
+ 0, /* OMP_CLAUSE_IF_PRESENT */
2 /* OMP_CLAUSE_DEVICE_TYPE */
};
@@ -404,6 +405,7 @@ const char * const omp_clause_code_name[] =
"nohost",
"tile",
"_griddim_",
+ "if_present",
"device_type"
};
@@ -11720,6 +11722,7 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data,
case OMP_CLAUSE_SEQ:
case OMP_CLAUSE_NOHOST:
case OMP_CLAUSE_TILE:
+ case OMP_CLAUSE_IF_PRESENT:
WALK_SUBTREE_TAIL (OMP_CLAUSE_CHAIN (*tp));
case OMP_CLAUSE_DEVICE_TYPE: