This patch make the c, c++ and fortran FEs aware of the new OpenACC 2.5 if_present clause for the update directive. The ME and runtime support will come in a separate followup patch.
Thomas, for some reason I'm seeing a couple of new UNRESOLVED tests for update-1.C. The c++ tests running with goacc.exp are built with -fopenacc, but for some reason the tests in g++.dg/goacc/ are still ran without -fopenacc for g++.dg/dg.exp. Maybe there's something wrong with g++.dg/goacc/goacc.exp handling of .C files? This patch has been committed to gomp-4_0-branch. Cesar
2017-05-04 Cesar Philippidis <ce...@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. diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h index 7b77dca..f1716ad 100644 --- a/gcc/c-family/c-pragma.h +++ b/gcc/c-family/c-pragma.h @@ -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, diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index b1af31f..957007e 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -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 \ diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c index b04db44..70c15be 100644 --- a/gcc/c/c-typeck.c +++ b/gcc/c/c-typeck.c @@ -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; diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index b082feb..b9c9747 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -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 \ diff --git a/gcc/cp/pt.c b/gcc/cp/pt.c index 56758d6..84f64d8 100644 --- a/gcc/cp/pt.c +++ b/gcc/cp/pt.c @@ -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: diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c index 6e8fb17..c5cf002 100644 --- a/gcc/cp/semantics.c +++ b/gcc/cp/semantics.c @@ -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: diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h index 3c762a8..a9a6856 100644 --- a/gcc/fortran/gfortran.h +++ b/gcc/fortran/gfortran.h @@ -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]; } diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c index 77a9c8d..12b2430 100644 --- a/gcc/fortran/openmp.c +++ b/gcc/fortran/openmp.c @@ -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) diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c index 3718da2..1ed2d68 100644 --- a/gcc/fortran/trans-openmp.c +++ b/gcc/fortran/trans-openmp.c @@ -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); diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 0250e4a..af908f4 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -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; diff --git a/gcc/omp-low.c b/gcc/omp-low.c index 9e9a363..a681800 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -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; diff --git a/gcc/testsuite/c-c++-common/goacc/update-if_present-1.c b/gcc/testsuite/c-c++-common/goacc/update-if_present-1.c new file mode 100644 index 0000000..519393cf --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/update-if_present-1.c @@ -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" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/update-if_present-2.c b/gcc/testsuite/c-c++-common/goacc/update-if_present-2.c new file mode 100644 index 0000000..be44b3e --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/update-if_present-2.c @@ -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++) + ; +} + diff --git a/gcc/testsuite/g++.dg/goacc/update-1.C b/gcc/testsuite/g++.dg/goacc/update-1.C new file mode 100644 index 0000000..10c8020 --- /dev/null +++ b/gcc/testsuite/g++.dg/goacc/update-1.C @@ -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" } } */ diff --git a/gcc/testsuite/gfortran.dg/goacc/update-if_present-1.f90 b/gcc/testsuite/gfortran.dg/goacc/update-if_present-1.f90 new file mode 100644 index 0000000..8302e06 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/update-if_present-1.f90 @@ -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" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/update-if_present-2.f90 b/gcc/testsuite/gfortran.dg/goacc/update-if_present-2.f90 new file mode 100644 index 0000000..bdf3112 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/update-if_present-2.f90 @@ -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 diff --git a/gcc/tree-core.h b/gcc/tree-core.h index d0f490c..494a725 100644 --- a/gcc/tree-core.h +++ b/gcc/tree-core.h @@ -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 }; diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c index a36c938..a208e8f 100644 --- a/gcc/tree-pretty-print.c +++ b/gcc/tree-pretty-print.c @@ -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"); diff --git a/gcc/tree.c b/gcc/tree.c index 6b74880..58c3375 100644 --- a/gcc/tree.c +++ b/gcc/tree.c @@ -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: