diff options
Diffstat (limited to 'gcc/c')
| -rw-r--r-- | gcc/c/ChangeLog | 14 | ||||
| -rw-r--r-- | gcc/c/c-parser.c | 41 | ||||
| -rw-r--r-- | gcc/c/c-tree.h | 1 | ||||
| -rw-r--r-- | gcc/c/c-typeck.c | 20 |
4 files changed, 76 insertions, 0 deletions
diff --git a/gcc/c/ChangeLog b/gcc/c/ChangeLog index 0655aef..43d1579 100644 --- a/gcc/c/ChangeLog +++ b/gcc/c/ChangeLog @@ -1,3 +1,17 @@ +2015-12-01 Julian Brown <julian@codesourcery.com> + Cesar Philippidis <cesar@codesourcery.com> + James Norris <James_Norris@mentor.com> + + * c-parser.c (c_parser_omp_clause_name): Add use_device support. + (c_parser_oacc_clause_use_device): New function. + (c_parser_oacc_all_clauses): Add use_device support. + (OACC_HOST_DATA_CLAUSE_MASK): New macro. + (c_parser_oacc_host_data): New function. + (c_parser_omp_construct): Add host_data support. + * c-tree.h (c_finish_oacc_host_data): Add prototype. + * c-typeck.c (c_finish_oacc_host_data): New function. + (c_finish_omp_clauses): Add use_device support. + 2015-11-29 Jan Hubicka <hubicka@ucw.cz> PR c/67106 diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index 0259f66..d4c512f 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -10279,6 +10279,8 @@ c_parser_omp_clause_name (c_parser *parser) result = PRAGMA_OMP_CLAUSE_UNTIED; else if (!strcmp ("use_device_ptr", p)) result = PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR; + else if (!strcmp ("use_device", p)) + result = PRAGMA_OACC_CLAUSE_USE_DEVICE; break; case 'v': if (!strcmp ("vector", p)) @@ -11631,6 +11633,15 @@ 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 ) */ @@ -12940,6 +12951,10 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask, clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); c_name = "self"; break; + case PRAGMA_OACC_CLAUSE_USE_DEVICE: + clauses = c_parser_oacc_clause_use_device (parser, clauses); + c_name = "use_device"; + break; case PRAGMA_OACC_CLAUSE_SEQ: clauses = c_parser_oacc_simple_clause (parser, OMP_CLAUSE_SEQ, clauses); @@ -13590,6 +13605,29 @@ c_parser_oacc_enter_exit_data (c_parser *parser, bool enter) /* OpenACC 2.0: + # pragma acc host_data oacc-data-clause[optseq] new-line + structured-block +*/ + +#define OACC_HOST_DATA_CLAUSE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_USE_DEVICE) ) + +static tree +c_parser_oacc_host_data (location_t loc, c_parser *parser) +{ + tree stmt, clauses, block; + + clauses = c_parser_oacc_all_clauses (parser, OACC_HOST_DATA_CLAUSE_MASK, + "#pragma acc host_data"); + + block = c_begin_omp_parallel (); + add_stmt (c_parser_omp_structured_block (parser)); + stmt = c_finish_oacc_host_data (loc, clauses, block); + return stmt; +} + + +/* OpenACC 2.0: # pragma acc loop oacc-loop-clause[optseq] new-line structured-block @@ -16897,6 +16935,9 @@ c_parser_omp_construct (c_parser *parser) case PRAGMA_OACC_DATA: stmt = c_parser_oacc_data (loc, parser); break; + case PRAGMA_OACC_HOST_DATA: + stmt = c_parser_oacc_host_data (loc, parser); + break; case PRAGMA_OACC_KERNELS: case PRAGMA_OACC_PARALLEL: strcpy (p_name, "#pragma acc"); diff --git a/gcc/c/c-tree.h b/gcc/c/c-tree.h index 6bc216a..848131e 100644 --- a/gcc/c/c-tree.h +++ b/gcc/c/c-tree.h @@ -653,6 +653,7 @@ extern tree c_finish_goto_ptr (location_t, tree); extern tree c_expr_to_decl (tree, bool *, bool *); extern tree c_finish_omp_construct (location_t, enum tree_code, tree, tree); extern tree c_finish_oacc_data (location_t, tree, tree); +extern tree c_finish_oacc_host_data (location_t, tree, tree); extern tree c_begin_omp_parallel (void); extern tree c_finish_omp_parallel (location_t, tree, tree); extern tree c_begin_omp_task (void); diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c index 741c75c..cc2e38e 100644 --- a/gcc/c/c-typeck.c +++ b/gcc/c/c-typeck.c @@ -11631,6 +11631,25 @@ c_finish_oacc_data (location_t loc, tree clauses, tree block) return add_stmt (stmt); } +/* Generate OACC_HOST_DATA, with CLAUSES and BLOCK as its compound + statement. LOC is the location of the OACC_HOST_DATA. */ + +tree +c_finish_oacc_host_data (location_t loc, tree clauses, tree block) +{ + tree stmt; + + block = c_end_compound_stmt (loc, block, true); + + stmt = make_node (OACC_HOST_DATA); + TREE_TYPE (stmt) = void_type_node; + OACC_HOST_DATA_CLAUSES (stmt) = clauses; + OACC_HOST_DATA_BODY (stmt) = block; + SET_EXPR_LOCATION (stmt, loc); + + return add_stmt (stmt); +} + /* Like c_begin_compound_stmt, except force the retention of the BLOCK. */ tree @@ -13074,6 +13093,7 @@ 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); |