Hi! On Mon, 26 Oct 2015 19:34:22 +0100, Jakub Jelinek wrote: > Your use_device sounds very similar to use_device_ptr clause in OpenMP, > which is allowed on #pragma omp target data construct and is implemented > quite a bit differently from this; it is unclear if the OpenACC standard > requires this kind of implementation, or you just chose to implement it this > way. In particular, the GOMP_target_data call puts the variables mentioned > in the use_device_ptr clauses into the mapping structures (similarly how > map clause appears) and the corresponding vars are privatized within the > target data region (which is a host region, basically a fancy { } braces), > where the private variables contain the offloading device's pointers. ACK. As the OpenACC use_device clause implementation now completely matches the OpenMP use_device_ptr clause implementation, there is no use anymore for a separate OMP_CLAUSE_USE_DEVICE, so in r231926 I cleaned that up, as obvious: commit 9d5fd7c608fef6e7a9efbfc940545d49452c4e01 Author: tschwinge Date: Wed Dec 23 11:01:18 2015 +0000 Merge OMP_CLAUSE_USE_DEVICE into OMP_CLAUSE_USE_DEVICE_PTR gcc/c/ * c-parser.c (c_parser_oacc_clause_use_device): Merge function into... (c_parser_omp_clause_use_device_ptr): ... this function. Adjust all users. gcc/ * tree-core.h (enum omp_clause_code): Merge OMP_CLAUSE_USE_DEVICE into OMP_CLAUSE_USE_DEVICE_PTR. Adjust all users. git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@231926 138bc75d-0d04-0410-961f-82ee72b054a4 --- gcc/ChangeLog | 5 +++++ gcc/c/ChangeLog | 7 +++++++ gcc/c/c-parser.c | 16 +++++----------- gcc/c/c-typeck.c | 1 - gcc/cp/parser.c | 2 +- gcc/cp/pt.c | 2 -- gcc/cp/semantics.c | 1 - gcc/fortran/trans-openmp.c | 2 +- gcc/gimplify.c | 2 -- gcc/omp-low.c | 11 ++--------- gcc/testsuite/gfortran.dg/goacc/host_data-tree.f95 | 2 +- gcc/tree-core.h | 6 ++---- gcc/tree-nested.c | 2 -- gcc/tree-pretty-print.c | 3 --- gcc/tree.c | 3 --- 15 files changed, 24 insertions(+), 41 deletions(-) diff --git gcc/ChangeLog gcc/ChangeLog index aa28d10..d67b9c6 100644 --- gcc/ChangeLog +++ gcc/ChangeLog @@ -1,3 +1,8 @@ +2015-12-23 Thomas Schwinge + + * tree-core.h (enum omp_clause_code): Merge OMP_CLAUSE_USE_DEVICE + into OMP_CLAUSE_USE_DEVICE_PTR. Adjust all users. + 2015-12-23 David Sherwood * config/arm/iterators.md (VMAXMINFNM): New int iterator. diff --git gcc/c/ChangeLog gcc/c/ChangeLog index f99f426..7b275d8 100644 --- gcc/c/ChangeLog +++ gcc/c/ChangeLog @@ -1,3 +1,10 @@ +2015-12-23 Thomas Schwinge + + * c-parser.c (c_parser_oacc_clause_use_device): Merge function + into... + (c_parser_omp_clause_use_device_ptr): ... this function. Adjust + all users. + 2015-12-22 Marek Polacek PR c/69002 diff --git gcc/c/c-parser.c gcc/c/c-parser.c index 353e3da..8e754d0 100644 --- gcc/c/c-parser.c +++ gcc/c/c-parser.c @@ -11395,7 +11395,10 @@ c_parser_omp_clause_defaultmap (c_parser *parser, tree list) return list; } -/* OpenMP 4.5: +/* OpenACC 2.0: + use_device ( variable-list ) + + OpenMP 4.5: use_device_ptr ( variable-list ) */ static tree @@ -11730,15 +11733,6 @@ c_parser_oacc_clause_tile (c_parser *parser, tree list) return c; } -/* OpenACC 2.0: - use_device ( variable-list ) */ - -static tree -c_parser_oacc_clause_use_device (c_parser *parser, tree list) -{ - return c_parser_omp_var_list_parens (parser, OMP_CLAUSE_USE_DEVICE, list); -} - /* OpenACC: wait ( int-expr-list ) */ @@ -13058,7 +13052,7 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask, c_name = "tile"; break; case PRAGMA_OACC_CLAUSE_USE_DEVICE: - clauses = c_parser_oacc_clause_use_device (parser, clauses); + clauses = c_parser_omp_clause_use_device_ptr (parser, clauses); c_name = "use_device"; break; case PRAGMA_OACC_CLAUSE_VECTOR: diff --git gcc/c/c-typeck.c gcc/c/c-typeck.c index 928fcd5..7406bd4 100644 --- gcc/c/c-typeck.c +++ gcc/c/c-typeck.c @@ -13125,7 +13125,6 @@ c_finish_omp_clauses (tree clauses, bool is_omp, bool declare_simd) bitmap_set_bit (&map_head, DECL_UID (t)); goto check_dup_generic; - case OMP_CLAUSE_USE_DEVICE: case OMP_CLAUSE_IS_DEVICE_PTR: case OMP_CLAUSE_USE_DEVICE_PTR: t = OMP_CLAUSE_DECL (c); diff --git gcc/cp/parser.c gcc/cp/parser.c index 842dded..4829a77 100644 --- gcc/cp/parser.c +++ gcc/cp/parser.c @@ -32097,7 +32097,7 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask, c_name = "tile"; break; case PRAGMA_OACC_CLAUSE_USE_DEVICE: - clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_USE_DEVICE, + clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_USE_DEVICE_PTR, clauses); c_name = "use_device"; break; diff --git gcc/cp/pt.c gcc/cp/pt.c index dab15bd..4555b32 100644 --- gcc/cp/pt.c +++ gcc/cp/pt.c @@ -14425,7 +14425,6 @@ tsubst_omp_clauses (tree clauses, bool declare_simd, bool allow_fields, case OMP_CLAUSE_FROM: case OMP_CLAUSE_TO: case OMP_CLAUSE_MAP: - case OMP_CLAUSE_USE_DEVICE: case OMP_CLAUSE_USE_DEVICE_PTR: case OMP_CLAUSE_IS_DEVICE_PTR: OMP_CLAUSE_DECL (nc) @@ -14552,7 +14551,6 @@ tsubst_omp_clauses (tree clauses, bool declare_simd, bool allow_fields, case OMP_CLAUSE_COPYPRIVATE: case OMP_CLAUSE_LINEAR: case OMP_CLAUSE_REDUCTION: - case OMP_CLAUSE_USE_DEVICE: case OMP_CLAUSE_USE_DEVICE_PTR: case OMP_CLAUSE_IS_DEVICE_PTR: /* tsubst_expr on SCOPE_REF results in returning diff --git gcc/cp/semantics.c gcc/cp/semantics.c index ab9989a..37bf050 100644 --- gcc/cp/semantics.c +++ gcc/cp/semantics.c @@ -6886,7 +6886,6 @@ finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd) } break; - case OMP_CLAUSE_USE_DEVICE: case OMP_CLAUSE_IS_DEVICE_PTR: case OMP_CLAUSE_USE_DEVICE_PTR: field_ok = allow_fields; diff --git gcc/fortran/trans-openmp.c gcc/fortran/trans-openmp.c index 227964c..70a7722 100644 --- gcc/fortran/trans-openmp.c +++ gcc/fortran/trans-openmp.c @@ -1771,7 +1771,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, clause_code = OMP_CLAUSE_UNIFORM; goto add_clause; case OMP_LIST_USE_DEVICE: - clause_code = OMP_CLAUSE_USE_DEVICE; + clause_code = OMP_CLAUSE_USE_DEVICE_PTR; goto add_clause; case OMP_LIST_DEVICE_RESIDENT: clause_code = OMP_CLAUSE_DEVICE_RESIDENT; diff --git gcc/gimplify.c gcc/gimplify.c index 62b0e64..bc90401 100644 --- gcc/gimplify.c +++ gcc/gimplify.c @@ -7139,7 +7139,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, } goto do_notice; - case OMP_CLAUSE_USE_DEVICE: case OMP_CLAUSE_USE_DEVICE_PTR: flags = GOVD_FIRSTPRIVATE | GOVD_EXPLICIT; goto do_add; @@ -8051,7 +8050,6 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, case OMP_CLAUSE_ASYNC: case OMP_CLAUSE_WAIT: case OMP_CLAUSE_DEVICE_RESIDENT: - case OMP_CLAUSE_USE_DEVICE: case OMP_CLAUSE_INDEPENDENT: case OMP_CLAUSE_NUM_GANGS: case OMP_CLAUSE_NUM_WORKERS: diff --git gcc/omp-low.c gcc/omp-low.c index 676b1df..a0c3e1c 100644 --- gcc/omp-low.c +++ gcc/omp-low.c @@ -1957,7 +1957,6 @@ scan_sharing_clauses (tree clauses, omp_context *ctx, install_var_local (decl, ctx); break; - case OMP_CLAUSE_USE_DEVICE: case OMP_CLAUSE_USE_DEVICE_PTR: decl = OMP_CLAUSE_DECL (c); if (TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE) @@ -2314,7 +2313,6 @@ scan_sharing_clauses (tree clauses, omp_context *ctx, case OMP_CLAUSE_SIMD: case OMP_CLAUSE_NOGROUP: case OMP_CLAUSE_DEFAULTMAP: - case OMP_CLAUSE_USE_DEVICE: case OMP_CLAUSE_USE_DEVICE_PTR: case OMP_CLAUSE__CILK_FOR_COUNT_: case OMP_CLAUSE_ASYNC: @@ -15288,7 +15286,6 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) } break; - case OMP_CLAUSE_USE_DEVICE: case OMP_CLAUSE_USE_DEVICE_PTR: case OMP_CLAUSE_IS_DEVICE_PTR: var = OMP_CLAUSE_DECL (c); @@ -15674,14 +15671,12 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) build_int_cstu (tkind_type, tkind)); break; - case OMP_CLAUSE_USE_DEVICE: case OMP_CLAUSE_USE_DEVICE_PTR: case OMP_CLAUSE_IS_DEVICE_PTR: ovar = OMP_CLAUSE_DECL (c); var = lookup_decl_in_outer_ctx (ovar, ctx); x = build_sender_ref (ovar, ctx); - if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR - || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE) + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR) tkind = GOMP_MAP_USE_DEVICE_PTR; else tkind = GOMP_MAP_FIRSTPRIVATE_INT; @@ -15884,12 +15879,10 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) gimple_build_assign (new_var, x)); } break; - case OMP_CLAUSE_USE_DEVICE: case OMP_CLAUSE_USE_DEVICE_PTR: case OMP_CLAUSE_IS_DEVICE_PTR: var = OMP_CLAUSE_DECL (c); - if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR - || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE) + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR) x = build_sender_ref (var, ctx); else x = build_receiver_ref (var, false, ctx); diff --git gcc/testsuite/gfortran.dg/goacc/host_data-tree.f95 gcc/testsuite/gfortran.dg/goacc/host_data-tree.f95 index 7a5eea6..23aba8c 100644 --- gcc/testsuite/gfortran.dg/goacc/host_data-tree.f95 +++ gcc/testsuite/gfortran.dg/goacc/host_data-tree.f95 @@ -8,4 +8,4 @@ program test !$acc host_data use_device(i) !$acc end host_data end program test -! { dg-final { scan-tree-dump-times "pragma acc host_data use_device\\(i\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "pragma acc host_data use_device_ptr\\(i\\)" 1 "original" } } diff --git gcc/tree-core.h gcc/tree-core.h index 9cc64d9..5371378 100644 --- gcc/tree-core.h +++ gcc/tree-core.h @@ -302,7 +302,8 @@ enum omp_clause_code { OpenMP clause: map ({alloc:,to:,from:,tofrom:,}variable-list). */ OMP_CLAUSE_MAP, - /* OpenMP clause: use_device_ptr (variable-list). */ + /* OpenACC clause: use_device (variable_list). + OpenMP clause: use_device_ptr (variable-list). */ OMP_CLAUSE_USE_DEVICE_PTR, /* OpenMP clause: is_device_ptr (variable-list). */ @@ -315,9 +316,6 @@ enum omp_clause_code { /* OpenACC clause: device_resident (variable_list). */ OMP_CLAUSE_DEVICE_RESIDENT, - /* OpenACC clause: use_device (variable_list). */ - OMP_CLAUSE_USE_DEVICE, - /* OpenACC clause: gang [(gang-argument-list)]. Where gang-argument-list: [gang-argument-list, ] gang-argument diff --git gcc/tree-nested.c gcc/tree-nested.c index 3a9479a..8211a12 100644 --- gcc/tree-nested.c +++ gcc/tree-nested.c @@ -1072,7 +1072,6 @@ convert_nonlocal_omp_clauses (tree *pclauses, struct walk_stmt_info *wi) case OMP_CLAUSE_SHARED: case OMP_CLAUSE_TO_DECLARE: case OMP_CLAUSE_LINK: - case OMP_CLAUSE_USE_DEVICE: case OMP_CLAUSE_USE_DEVICE_PTR: case OMP_CLAUSE_IS_DEVICE_PTR: do_decl_clause: @@ -1744,7 +1743,6 @@ convert_local_omp_clauses (tree *pclauses, struct walk_stmt_info *wi) case OMP_CLAUSE_SHARED: case OMP_CLAUSE_TO_DECLARE: case OMP_CLAUSE_LINK: - case OMP_CLAUSE_USE_DEVICE: case OMP_CLAUSE_USE_DEVICE_PTR: case OMP_CLAUSE_IS_DEVICE_PTR: do_decl_clause: diff --git gcc/tree-pretty-print.c gcc/tree-pretty-print.c index caec760..c4a180e 100644 --- gcc/tree-pretty-print.c +++ gcc/tree-pretty-print.c @@ -327,9 +327,6 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, int flags) case OMP_CLAUSE_DEVICE_RESIDENT: name = "device_resident"; goto print_remap; - case OMP_CLAUSE_USE_DEVICE: - name = "use_device"; - goto print_remap; case OMP_CLAUSE_TO_DECLARE: name = "to"; goto print_remap; diff --git gcc/tree.c gcc/tree.c index 2190cae..d837609 100644 --- gcc/tree.c +++ gcc/tree.c @@ -282,7 +282,6 @@ unsigned const char omp_clause_num_ops[] = 1, /* OMP_CLAUSE_IS_DEVICE_PTR */ 2, /* OMP_CLAUSE__CACHE_ */ 1, /* OMP_CLAUSE_DEVICE_RESIDENT */ - 1, /* OMP_CLAUSE_USE_DEVICE */ 2, /* OMP_CLAUSE_GANG */ 1, /* OMP_CLAUSE_ASYNC */ 1, /* OMP_CLAUSE_WAIT */ @@ -354,7 +353,6 @@ const char * const omp_clause_code_name[] = "is_device_ptr", "_cache_", "device_resident", - "use_device", "gang", "async", "wait", @@ -11612,7 +11610,6 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data, /* FALLTHRU */ case OMP_CLAUSE_DEVICE_RESIDENT: - case OMP_CLAUSE_USE_DEVICE: case OMP_CLAUSE_ASYNC: case OMP_CLAUSE_WAIT: case OMP_CLAUSE_WORKER: Grüße Thomas