2015-11-03 Cesar Philippidis Thomas Schwinge James Norris gcc/ * gimplify.c (gimplify_scan_omp_clauses): Add support for OMP_CLAUSE_TILE. Update handling of OMP_CLAUSE_INDEPENDENT. (gimplify_adjust_omp_clauses): Likewise. * omp-low.c (scan_sharing_clauses): Add support for OMP_CLAUSE_TILE. * tree-core.h (enum omp_clause_code): Add OMP_CLAUSE_TILE. * tree-pretty-print.c (dump_omp_clause): Handle OMP_CLAUSE_TILE. * tree.c (omp_clause_num_ops): Add an entry for OMP_CLAUSE_TILE. (omp_clause_code_name): Likewise. (walk_tree_1): Handle OMP_CLAUSE_TILE. * tree.h (OMP_TILE_LIST): New macro. gcc/c-family/ * c-omp.c (c_oacc_split_loop_clauses): Make TILE, GANG, WORKER, VECTOR, AUTO, SEQ and independent as loop clauses. Associate REDUCTION clauses with parallel and kernels. * c-pragma.h (enum pragma_omp_clause): Add entries for PRAGMA_OACC_CLAUSE_{INDEPENDENT,TILE,DEFAULT}. gcc/c/ * c-parser.c (c_parser_omp_clause_name): Add support for PRAGMA_OACC_CLAUSE_INDEPENDENT and PRAGMA_OACC_CLAUSE_TILE. (c_parser_omp_clause_default): Add is_oacc argument. Handle default(none) in OpenACC. (c_parser_oacc_clause_tile): New function. (c_parser_oacc_all_clauses): Add support for OMP_CLAUSE_DEFAULT, OMP_CLAUSE_INDEPENDENT and OMP_CLAUSE_TILE. (OACC_LOOP_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_{PRIVATE,INDEPENDENT, TILE}. (OACC_KERNELS_MASK): Add PRAGMA_OACC_CLAUSE_DEFAULT. (OACC_PARALLEL_MASK): Add PRAGMA_OACC_CLAUSE_{DEFAULT,PRIVATE, FIRSTPRIVATE}. (c_parser_omp_all_clauses): Update call to c_parser_omp_clause_default. * c-typeck.c (c_finish_omp_clauses): Add support for OMP_CLAUSE_TILE and OMP_CLAUSE_INDEPENDENT. gcc/cp/ * parser.c (cp_parser_omp_clause_name): Add support for PRAGMA_OACC_CLAUSE_INDEPENDENT and PRAGMA_OACC_CLAUSE_TILE. (cp_parser_oacc_clause_tile): New function. (cp_parser_omp_clause_default): Add is_oacc argument. Handle default(none) in OpenACC. (cp_parser_oacc_all_clauses): Add support for (cp_parser_omp_all_clauses): Update call to cp_parser_omp_clause_default. PRAGMA_OACC_CLAUSE_{DEFAULT,INDEPENDENT,TILE,PRIVATE,FIRSTPRIVATE}. (OACC_LOOP_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_{PRIVATE,INDEPENDENT, TILE}. (OACC_KERNELS_MASK): Add PRAGMA_OACC_CLAUSE_DEFAULT. (OACC_PARALLEL_MASK): Add PRAGMA_OACC_CLAUSE_{DEFAULT,PRIVATE, FIRSTPRIVATE}. * semantics.c (finish_omp_clauses): Add support for OMP_CLAUSE_INDEPENDENT and OMP_CLAUSE_TILE. gcc/testsuite/ * c-c++-common/goacc/combined-directives.c: New test. * c-c++-common/goacc/loop-clauses.c: New test. * c-c++-common/goacc/tile.c: New test. diff --git a/gcc/c-family/c-omp.c b/gcc/c-family/c-omp.c index 133d079..a04caf3 100644 --- a/gcc/c-family/c-omp.c +++ b/gcc/c-family/c-omp.c @@ -709,12 +709,21 @@ c_oacc_split_loop_clauses (tree clauses, tree *not_loop_clauses) switch (OMP_CLAUSE_CODE (clauses)) { + /* Loop clauses. */ case OMP_CLAUSE_COLLAPSE: - case OMP_CLAUSE_REDUCTION: + case OMP_CLAUSE_TILE: + case OMP_CLAUSE_GANG: + case OMP_CLAUSE_WORKER: + case OMP_CLAUSE_VECTOR: + case OMP_CLAUSE_AUTO: + case OMP_CLAUSE_SEQ: + case OMP_CLAUSE_INDEPENDENT: OMP_CLAUSE_CHAIN (clauses) = loop_clauses; loop_clauses = clauses; break; + /* Parallel/kernels clauses. */ + default: OMP_CLAUSE_CHAIN (clauses) = *not_loop_clauses; *not_loop_clauses = clauses; diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h index 69e7392..953c4e3 100644 --- a/gcc/c-family/c-pragma.h +++ b/gcc/c-family/c-pragma.h @@ -153,6 +153,7 @@ enum pragma_omp_clause { PRAGMA_OACC_CLAUSE_DEVICEPTR, PRAGMA_OACC_CLAUSE_GANG, PRAGMA_OACC_CLAUSE_HOST, + PRAGMA_OACC_CLAUSE_INDEPENDENT, PRAGMA_OACC_CLAUSE_NUM_GANGS, PRAGMA_OACC_CLAUSE_NUM_WORKERS, PRAGMA_OACC_CLAUSE_PRESENT, @@ -162,6 +163,7 @@ enum pragma_omp_clause { PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE, PRAGMA_OACC_CLAUSE_SELF, PRAGMA_OACC_CLAUSE_SEQ, + PRAGMA_OACC_CLAUSE_TILE, PRAGMA_OACC_CLAUSE_VECTOR, PRAGMA_OACC_CLAUSE_VECTOR_LENGTH, PRAGMA_OACC_CLAUSE_WAIT, @@ -169,6 +171,7 @@ enum pragma_omp_clause { PRAGMA_OACC_CLAUSE_COLLAPSE = PRAGMA_OMP_CLAUSE_COLLAPSE, PRAGMA_OACC_CLAUSE_COPYIN = PRAGMA_OMP_CLAUSE_COPYIN, PRAGMA_OACC_CLAUSE_DEVICE = PRAGMA_OMP_CLAUSE_DEVICE, + PRAGMA_OACC_CLAUSE_DEFAULT = PRAGMA_OMP_CLAUSE_DEFAULT, PRAGMA_OACC_CLAUSE_FIRSTPRIVATE = PRAGMA_OMP_CLAUSE_FIRSTPRIVATE, PRAGMA_OACC_CLAUSE_IF = PRAGMA_OMP_CLAUSE_IF, PRAGMA_OACC_CLAUSE_PRIVATE = PRAGMA_OMP_CLAUSE_PRIVATE, diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index ec88c65..ab2bd15 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -10006,6 +10006,8 @@ c_parser_omp_clause_name (c_parser *parser) case 'i': if (!strcmp ("inbranch", p)) result = PRAGMA_OMP_CLAUSE_INBRANCH; + else if (!strcmp ("independent", p)) + result = PRAGMA_OACC_CLAUSE_INDEPENDENT; else if (!strcmp ("is_device_ptr", p)) result = PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR; break; @@ -10102,6 +10104,8 @@ c_parser_omp_clause_name (c_parser *parser) result = PRAGMA_OMP_CLAUSE_THREAD_LIMIT; else if (!strcmp ("threads", p)) result = PRAGMA_OMP_CLAUSE_THREADS; + else if (!strcmp ("tile", p)) + result = PRAGMA_OACC_CLAUSE_TILE; else if (!strcmp ("to", p)) result = PRAGMA_OMP_CLAUSE_TO; break; @@ -10539,10 +10543,13 @@ c_parser_omp_clause_copyprivate (c_parser *parser, tree list) } /* OpenMP 2.5: - default ( shared | none ) */ + default ( shared | none ) + + OpenACC 2.0: + default (none) */ static tree -c_parser_omp_clause_default (c_parser *parser, tree list) +c_parser_omp_clause_default (c_parser *parser, tree list, bool is_oacc) { enum omp_clause_default_kind kind = OMP_CLAUSE_DEFAULT_UNSPECIFIED; location_t loc = c_parser_peek_token (parser)->location; @@ -10563,7 +10570,7 @@ c_parser_omp_clause_default (c_parser *parser, tree list) break; case 's': - if (strcmp ("shared", p) != 0) + if (strcmp ("shared", p) != 0 || is_oacc) goto invalid_kind; kind = OMP_CLAUSE_DEFAULT_SHARED; break; @@ -10577,7 +10584,10 @@ c_parser_omp_clause_default (c_parser *parser, tree list) else { invalid_kind: - c_parser_error (parser, "expected % or %"); + if (is_oacc) + c_parser_error (parser, "expected %"); + else + c_parser_error (parser, "expected % or %"); } c_parser_skip_until_found (parser, CPP_CLOSE_PAREN, "expected %<)%>"); @@ -11376,6 +11386,83 @@ c_parser_oacc_clause_async (c_parser *parser, tree list) return list; } +/* OpenACC 2.0: + tile ( size-expr-list ) */ + +static tree +c_parser_oacc_clause_tile (c_parser *parser, tree list) +{ + tree c, expr = error_mark_node; + location_t loc, expr_loc; + tree tile = NULL_TREE; + vec *tvec = make_tree_vector (); + + check_no_duplicate_clause (list, OMP_CLAUSE_TILE, "tile"); + + loc = c_parser_peek_token (parser)->location; + if (!c_parser_require (parser, CPP_OPEN_PAREN, "expected %<(%>")) + { + release_tree_vector (tvec); + return list; + } + + do + { + if (c_parser_next_token_is (parser, CPP_MULT)) + { + c_parser_consume_token (parser); + expr = integer_minus_one_node; + } + else + { + expr_loc = c_parser_peek_token (parser)->location; + expr = c_parser_expr_no_commas (parser, NULL).value; + + if (expr == error_mark_node) + goto cleanup_error; + + if (!INTEGRAL_TYPE_P (TREE_TYPE (expr))) + { + c_parser_error (parser, "% value must be integral"); + return list; + } + + mark_exp_read (expr); + expr = c_fully_fold (expr, false, NULL); + + /* Attempt to statically determine when expr isn't positive. */ + c = fold_build2_loc (expr_loc, LE_EXPR, boolean_type_node, expr, + build_int_cst (TREE_TYPE (expr), 0)); + protected_set_expr_location (c, expr_loc); + if (c == boolean_true_node) + { + warning_at (expr_loc, 0,"% value must be positive"); + expr = integer_one_node; + } + } + + vec_safe_push (tvec, expr); + if (c_parser_next_token_is (parser, CPP_COMMA)) + c_parser_consume_token (parser); + } + while (c_parser_next_token_is_not (parser, CPP_CLOSE_PAREN)); + + /* Consume the trailing ')'. */ + c_parser_consume_token (parser); + + c = build_omp_clause (loc, OMP_CLAUSE_TILE); + tile = build_tree_list_vec (tvec); + OMP_CLAUSE_TILE_LIST (c) = tile; + OMP_CLAUSE_CHAIN (c) = list; + release_tree_vector (tvec); + return c; + + cleanup_error: + c_parser_skip_until_found (parser, CPP_CLOSE_PAREN, "expected %<)%>"); + release_tree_vector (tvec); + return list; +} + /* OpenACC: wait ( int-expr-list ) */ @@ -12576,6 +12663,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 = "delete"; break; + case PRAGMA_OMP_CLAUSE_DEFAULT: + clauses = c_parser_omp_clause_default (parser, clauses, true); + c_name = "default"; + break; case PRAGMA_OACC_CLAUSE_DEVICE: clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); c_name = "device"; @@ -12601,6 +12692,11 @@ 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_INDEPENDENT: + clauses = c_parser_oacc_simple_clause (parser, OMP_CLAUSE_INDEPENDENT, + clauses); + c_name = "independent"; + break; case PRAGMA_OACC_CLAUSE_NUM_GANGS: clauses = c_parser_omp_clause_num_gangs (parser, clauses); c_name = "num_gangs"; @@ -12646,6 +12742,10 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask, clauses); c_name = "seq"; break; + case PRAGMA_OACC_CLAUSE_TILE: + clauses = c_parser_oacc_clause_tile (parser, clauses); + c_name = "tile"; + break; case PRAGMA_OACC_CLAUSE_VECTOR: c_name = "vector"; clauses = c_parser_oacc_shape_clause (parser, OMP_CLAUSE_VECTOR, @@ -12727,7 +12827,7 @@ c_parser_omp_all_clauses (c_parser *parser, omp_clause_mask mask, c_name = "copyprivate"; break; case PRAGMA_OMP_CLAUSE_DEFAULT: - clauses = c_parser_omp_clause_default (parser, clauses); + clauses = c_parser_omp_clause_default (parser, clauses, false); c_name = "default"; break; case PRAGMA_OMP_CLAUSE_FIRSTPRIVATE: @@ -13140,13 +13240,15 @@ c_parser_oacc_enter_exit_data (c_parser *parser, bool enter) #define OACC_LOOP_CLAUSE_MASK \ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COLLAPSE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRIVATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_GANG) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WORKER) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_AUTO) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_INDEPENDENT) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SEQ) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION) ) - + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_TILE) ) static tree c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name, omp_clause_mask mask, tree *cclauses) @@ -13191,6 +13293,7 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name, | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \ @@ -13206,8 +13309,11 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name, | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRIVATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_FIRSTPRIVATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_GANGS) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_WORKERS) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \ diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c index 2363b9b..24cedf3 100644 --- a/gcc/c/c-typeck.c +++ b/gcc/c/c-typeck.c @@ -12946,10 +12946,12 @@ c_finish_omp_clauses (tree clauses, bool is_omp, bool declare_simd) case OMP_CLAUSE_ASYNC: case OMP_CLAUSE_WAIT: case OMP_CLAUSE_AUTO: + case OMP_CLAUSE_INDEPENDENT: case OMP_CLAUSE_SEQ: case OMP_CLAUSE_GANG: case OMP_CLAUSE_WORKER: case OMP_CLAUSE_VECTOR: + case OMP_CLAUSE_TILE: pc = &OMP_CLAUSE_CHAIN (c); continue; diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index a90bf3b..200bf0c 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -29120,6 +29120,8 @@ cp_parser_omp_clause_name (cp_parser *parser) case 'i': if (!strcmp ("inbranch", p)) result = PRAGMA_OMP_CLAUSE_INBRANCH; + else if (!strcmp ("independent", p)) + result = PRAGMA_OACC_CLAUSE_INDEPENDENT; else if (!strcmp ("is_device_ptr", p)) result = PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR; break; @@ -29214,6 +29216,8 @@ cp_parser_omp_clause_name (cp_parser *parser) result = PRAGMA_OMP_CLAUSE_THREAD_LIMIT; else if (!strcmp ("threads", p)) result = PRAGMA_OMP_CLAUSE_THREADS; + else if (!strcmp ("tile", p)) + result = PRAGMA_OACC_CLAUSE_TILE; else if (!strcmp ("to", p)) result = PRAGMA_OMP_CLAUSE_TO; break; @@ -29714,6 +29718,58 @@ cp_parser_oacc_shape_clause (cp_parser *parser, omp_clause_code kind, return list; } +/* OpenACC 2.0: + tile ( size-expr-list ) */ + +static tree +cp_parser_oacc_clause_tile (cp_parser *parser, tree list, location_t here) +{ + tree c, expr = error_mark_node; + location_t loc; + tree tile = NULL_TREE; + vec *tvec = make_tree_vector (); + + check_no_duplicate_clause (list, OMP_CLAUSE_TILE, "tile", here); + + loc = cp_lexer_peek_token (parser->lexer)->location; + if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN)) + goto cleanup_error; + + do + { + if (cp_lexer_next_token_is (parser->lexer, CPP_MULT)) + { + cp_lexer_consume_token (parser->lexer); + expr = integer_minus_one_node; + } + else + expr = cp_parser_assignment_expression (parser, NULL, false, false); + + if (expr == error_mark_node) + goto cleanup_error; + + vec_safe_push (tvec, expr); + + if (cp_lexer_next_token_is (parser->lexer, CPP_COMMA)) + cp_lexer_consume_token (parser->lexer); + } + while (cp_lexer_next_token_is_not (parser->lexer, CPP_CLOSE_PAREN)); + + /* Consume the trailing ')'. */ + cp_lexer_consume_token (parser->lexer); + + c = build_omp_clause (loc, OMP_CLAUSE_TILE); + tile = build_tree_list_vec (tvec); + OMP_CLAUSE_TILE_LIST (c) = tile; + OMP_CLAUSE_CHAIN (c) = list; + release_tree_vector (tvec); + return c; + + cleanup_error: + release_tree_vector (tvec); + return list; +} + /* OpenACC: vector_length ( expression ) */ @@ -29860,10 +29916,14 @@ cp_parser_omp_clause_collapse (cp_parser *parser, tree list, location_t location } /* OpenMP 2.5: - default ( shared | none ) */ + default ( shared | none ) + + OpenACC 2.0 + default (none) */ static tree -cp_parser_omp_clause_default (cp_parser *parser, tree list, location_t location) +cp_parser_omp_clause_default (cp_parser *parser, tree list, + location_t location, bool is_oacc) { enum omp_clause_default_kind kind = OMP_CLAUSE_DEFAULT_UNSPECIFIED; tree c; @@ -29884,7 +29944,7 @@ cp_parser_omp_clause_default (cp_parser *parser, tree list, location_t location) break; case 's': - if (strcmp ("shared", p) != 0) + if (strcmp ("shared", p) != 0 || is_oacc) goto invalid_kind; kind = OMP_CLAUSE_DEFAULT_SHARED; break; @@ -29898,7 +29958,10 @@ cp_parser_omp_clause_default (cp_parser *parser, tree list, location_t location) else { invalid_kind: - cp_parser_error (parser, "expected % or %"); + if (is_oacc) + cp_parser_error (parser, "expected %"); + else + cp_parser_error (parser, "expected % or %"); } if (!cp_parser_require (parser, CPP_CLOSE_PAREN, RT_CLOSE_PAREN)) @@ -31467,6 +31530,10 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask, clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses); c_name = "delete"; break; + case PRAGMA_OMP_CLAUSE_DEFAULT: + clauses = cp_parser_omp_clause_default (parser, clauses, here, true); + c_name = "default"; + break; case PRAGMA_OACC_CLAUSE_DEVICE: clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses); c_name = "device"; @@ -31475,6 +31542,11 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask, clauses = cp_parser_oacc_data_clause_deviceptr (parser, clauses); c_name = "deviceptr"; break; + case PRAGMA_OACC_CLAUSE_FIRSTPRIVATE: + clauses = cp_parser_omp_var_list + (parser, OMP_CLAUSE_FIRSTPRIVATE, clauses); + c_name = "firstprivate"; + break; case PRAGMA_OACC_CLAUSE_GANG: c_name = "gang"; clauses = cp_parser_oacc_shape_clause (parser, OMP_CLAUSE_GANG, @@ -31488,6 +31560,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_INDEPENDENT: + clauses = cp_parser_oacc_simple_clause (parser, + OMP_CLAUSE_INDEPENDENT, + clauses, here); + c_name = "independent"; + break; case PRAGMA_OACC_CLAUSE_NUM_GANGS: clauses = cp_parser_omp_clause_num_gangs (parser, clauses); c_name = "num_gangs"; @@ -31516,6 +31594,11 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask, clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses); c_name = "present_or_create"; break; + case PRAGMA_OACC_CLAUSE_PRIVATE: + clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_PRIVATE, + clauses); + c_name = "private"; + break; case PRAGMA_OACC_CLAUSE_REDUCTION: clauses = cp_parser_omp_clause_reduction (parser, clauses); c_name = "reduction"; @@ -31529,6 +31612,10 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask, clauses, here); c_name = "seq"; break; + case PRAGMA_OACC_CLAUSE_TILE: + clauses = cp_parser_oacc_clause_tile (parser, clauses, here); + c_name = "tile"; + break; case PRAGMA_OACC_CLAUSE_VECTOR: c_name = "vector"; clauses = cp_parser_oacc_shape_clause (parser, OMP_CLAUSE_VECTOR, @@ -31616,7 +31703,7 @@ cp_parser_omp_all_clauses (cp_parser *parser, omp_clause_mask mask, break; case PRAGMA_OMP_CLAUSE_DEFAULT: clauses = cp_parser_omp_clause_default (parser, clauses, - token->location); + token->location, false); c_name = "default"; break; case PRAGMA_OMP_CLAUSE_FINAL: @@ -34460,12 +34547,15 @@ cp_parser_oacc_enter_exit_data (cp_parser *parser, cp_token *pragma_tok, #define OACC_LOOP_CLAUSE_MASK \ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COLLAPSE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRIVATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_GANG) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WORKER) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_AUTO) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_INDEPENDENT) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SEQ) \ - | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION) ) + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_TILE)) static tree cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name, @@ -34510,6 +34600,7 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name, | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \ @@ -34525,7 +34616,9 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name, | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_FIRSTPRIVATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_GANGS) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_WORKERS) \ @@ -34534,6 +34627,7 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name, | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYIN) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_COPYOUT) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT_OR_CREATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRIVATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) ) diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c index c73dcd0..14d006b 100644 --- a/gcc/cp/semantics.c +++ b/gcc/cp/semantics.c @@ -6704,9 +6704,51 @@ finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd) case OMP_CLAUSE_DEFAULTMAP: case OMP_CLAUSE__CILK_FOR_COUNT_: case OMP_CLAUSE_AUTO: + case OMP_CLAUSE_INDEPENDENT: case OMP_CLAUSE_SEQ: break; + case OMP_CLAUSE_TILE: + { + tree list = OMP_CLAUSE_TILE_LIST (c); + + while (list) + { + t = TREE_VALUE (list); + + if (t == error_mark_node) + remove = true; + else if (!type_dependent_expression_p (t) + && !INTEGRAL_TYPE_P (TREE_TYPE (t))) + { + error ("% value must be integral"); + remove = true; + } + else + { + t = mark_rvalue_use (t); + if (!processing_template_decl) + { + t = maybe_constant_value (t); + if (TREE_CODE (t) == INTEGER_CST + && tree_int_cst_sgn (t) != 1 + && t != integer_minus_one_node) + { + warning_at (OMP_CLAUSE_LOCATION (c), 0, + "% value must be positive"); + t = integer_one_node; + } + } + t = fold_build_cleanup_point_expr (TREE_TYPE (t), t); + } + + /* Update list item. */ + TREE_VALUE (list) = t; + list = TREE_CHAIN (list); + } + } + break; + case OMP_CLAUSE_INBRANCH: case OMP_CLAUSE_NOTINBRANCH: if (branch_seen) diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 03203c0..08b192d 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -6997,7 +6997,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, case OMP_CLAUSE_DEVICE_RESIDENT: case OMP_CLAUSE_USE_DEVICE: - case OMP_CLAUSE_INDEPENDENT: remove = true; break; @@ -7007,6 +7006,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, case OMP_CLAUSE_COLLAPSE: case OMP_CLAUSE_AUTO: case OMP_CLAUSE_SEQ: + case OMP_CLAUSE_INDEPENDENT: case OMP_CLAUSE_MERGEABLE: case OMP_CLAUSE_PROC_BIND: case OMP_CLAUSE_SAFELEN: @@ -7014,6 +7014,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_TILE: break; case OMP_CLAUSE_DEFAULTMAP: @@ -7482,6 +7483,7 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, tree *list_p, case OMP_CLAUSE_VECTOR: case OMP_CLAUSE_AUTO: case OMP_CLAUSE_SEQ: + case OMP_CLAUSE_TILE: break; default: diff --git a/gcc/omp-low.c b/gcc/omp-low.c index ccf0b63..9c87ac9 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -2230,6 +2230,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) case OMP_CLAUSE_GANG: case OMP_CLAUSE_WORKER: case OMP_CLAUSE_VECTOR: + case OMP_CLAUSE_TILE: break; case OMP_CLAUSE_ALIGNED: @@ -2408,6 +2409,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) case OMP_CLAUSE_GANG: case OMP_CLAUSE_WORKER: case OMP_CLAUSE_VECTOR: + case OMP_CLAUSE_TILE: break; case OMP_CLAUSE_DEVICE_RESIDENT: diff --git a/gcc/testsuite/c-c++-common/goacc/combined-directives.c b/gcc/testsuite/c-c++-common/goacc/combined-directives.c new file mode 100644 index 0000000..cefe712 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/combined-directives.c @@ -0,0 +1,117 @@ +// { dg-do compile } +// { dg-options "-fopenacc -fdump-tree-gimple" } + +// { dg-prune-output "sorry, unimplemented" } + +void +test () +{ + int a[100], i, j, z; + + // acc parallel + + #pragma acc parallel loop collapse (2) + for (i = 0; i < 100; i++) + for (j = 0; j < 10; j++) + ; + + #pragma acc parallel loop gang + for (i = 0; i < 100; i++) + ; + + #pragma acc parallel loop worker + for (i = 0; i < 100; i++) + for (j = 0; j < 10; j++) + ; + + #pragma acc parallel loop vector + for (i = 0; i < 100; i++) + for (j = 0; j < 10; j++) + ; + + #pragma acc parallel loop seq + for (i = 0; i < 100; i++) + for (j = 0; j < 10; j++) + ; + +#pragma acc parallel loop auto + for (i = 0; i < 100; i++) + for (j = 0; j < 10; j++) + ; + + #pragma acc parallel loop tile (2, 3) + for (i = 0; i < 100; i++) + for (j = 0; j < 10; j++) + ; + + #pragma acc parallel loop independent + for (i = 0; i < 100; i++) + ; + + #pragma acc parallel loop private (z) + for (i = 0; i < 100; i++) + z = 0; + + #pragma acc parallel loop reduction (+:z) copy (z) + for (i = 0; i < 100; i++) + ; + + // acc kernels + + #pragma acc kernels loop collapse (2) + for (i = 0; i < 100; i++) + for (j = 0; j < 10; j++) + ; + + #pragma acc kernels loop gang + for (i = 0; i < 100; i++) + ; + + #pragma acc kernels loop worker + for (i = 0; i < 100; i++) + for (j = 0; j < 10; j++) + ; + + #pragma acc kernels loop vector + for (i = 0; i < 100; i++) + for (j = 0; j < 10; j++) + ; + + #pragma acc kernels loop seq + for (i = 0; i < 100; i++) + for (j = 0; j < 10; j++) + ; + + #pragma acc kernels loop auto + for (i = 0; i < 100; i++) + for (j = 0; j < 10; j++) + ; + + #pragma acc kernels loop tile (2, 3) + for (i = 0; i < 100; i++) + for (j = 0; j < 10; j++) + ; + + #pragma acc kernels loop independent + for (i = 0; i < 100; i++) + ; + + #pragma acc kernels loop private (z) + for (i = 0; i < 100; i++) + z = 0; + + #pragma acc kernels loop reduction (+:z) copy (z) + for (i = 0; i < 100; i++) + ; +} + +// { dg-final { scan-tree-dump-times "acc loop collapse.2. private.j. private.i" 2 "gimple" } } +// { dg-final { scan-tree-dump-times "acc loop gang" 2 "gimple" } } +// { dg-final { scan-tree-dump-times "acc loop worker" 2 "gimple" } } +// { dg-final { scan-tree-dump-times "acc loop vector" 2 "gimple" } } +// { dg-final { scan-tree-dump-times "acc loop seq" 2 "gimple" } } +// { dg-final { scan-tree-dump-times "acc loop auto" 2 "gimple" } } +// { dg-final { scan-tree-dump-times "acc loop tile.2, 3" 2 "gimple" } } +// { dg-final { scan-tree-dump-times "acc loop independent private.i" 2 "gimple" } } +// { dg-final { scan-tree-dump-times "private.z" 2 "gimple" } } +// { dg-final { scan-tree-dump-times "reduction..:z. map.force_tofrom:z .len: 4.." 2 "gimple" } } diff --git a/gcc/testsuite/c-c++-common/goacc/loop-clauses.c b/gcc/testsuite/c-c++-common/goacc/loop-clauses.c new file mode 100644 index 0000000..97b8786 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/loop-clauses.c @@ -0,0 +1,86 @@ +/* { dg-do compile } */ + +/* { dg-prune-output "sorry, unimplemented" } */ + +int +main () +{ + int i, j; + +#pragma acc parallel firstprivate (j) private (i) + { +#pragma acc loop seq + for (i = 0; i < 10; i++) + { } + } + +#pragma acc parallel default (none) + { +#pragma acc loop auto private (j) + for (i = 0; i < 10; i++) + { } +#pragma acc loop gang + for (i = 0; i < 10; i++) + { } +#pragma acc loop gang(static:5) + for (i = 0; i < 10; i++) + { } +#pragma acc loop gang(static:*) + for (i = 0; i < 10; i++) + { } +#pragma acc loop vector + for (i = 0; i < 10; i++) + { } +#pragma acc loop worker + for (i = 0; i < 10; i++) + { } +#pragma acc loop auto + for (i = 0; i < 10; i++) + { } +#pragma acc loop independent + for (i = 0; i < 10; i++) + { } +#pragma acc loop seq + for (i = 0; i < 10; i++) + { } +#pragma acc loop gang worker vector + for (i = 0; i < 10; i++) + { } + } + +#pragma acc kernels default (none) + { +#pragma acc loop auto + for (i = 0; i < 10; i++) + { } +#pragma acc loop gang (num:5) + for (i = 0; i < 10; i++) + { } +#pragma acc loop gang(static:5) + for (i = 0; i < 10; i++) + { } +#pragma acc loop gang(static:*) + for (i = 0; i < 10; i++) + { } +#pragma acc loop vector(length:10) + for (i = 0; i < 10; i++) + { } +#pragma acc loop worker(num:5) + for (i = 0; i < 10; i++) + { } +#pragma acc loop auto + for (i = 0; i < 10; i++) + { } +#pragma acc loop independent + for (i = 0; i < 10; i++) + { } +#pragma acc loop seq + for (i = 0; i < 10; i++) + { } +#pragma acc loop gang worker vector + for (i = 0; i < 10; i++) + { } + } + + return 0; +} diff --git a/gcc/testsuite/c-c++-common/goacc/tile.c b/gcc/testsuite/c-c++-common/goacc/tile.c new file mode 100644 index 0000000..ce1391a --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/tile.c @@ -0,0 +1,61 @@ +/* { dg-do compile } */ + +int +main () +{ + int i; + +#pragma acc parallel loop tile (10) + for (i = 0; i < 100; i++) + ; + +#pragma acc parallel loop tile (*) + for (i = 0; i < 100; i++) + ; + +#pragma acc parallel loop tile (10, *) + for (i = 0; i < 100; i++) + ; + +#pragma acc parallel loop tile (10, *, i) + for (i = 0; i < 100; i++) + ; + +#pragma acc parallel loop tile // { dg-error "expected '\\\('" } + for (i = 0; i < 100; i++) + ; + +#pragma acc parallel loop tile () // { dg-error "" } + for (i = 0; i < 100; i++) + ; + +#pragma acc parallel loop tile (,1) // { dg-error "" } + for (i = 0; i < 100; i++) + ; + +#pragma acc parallel loop tile (,,) // { dg-error "" } + for (i = 0; i < 100; i++) + ; + +#pragma acc parallel loop tile (1.1) // { dg-error "'tile' value must be integral" } + for (i = 0; i < 100; i++) + ; + +#pragma acc parallel loop tile (-3) // { dg-warning "'tile' value must be positive" } + for (i = 0; i < 100; i++) + ; + +#pragma acc parallel loop tile (10,-3) // { dg-warning "'tile' value must be positive" } + for (i = 0; i < 100; i++) + ; + +#pragma acc parallel loop tile (-100,10,5) // { dg-warning "'tile' value must be positive" } + for (i = 0; i < 100; i++) + ; + +#pragma acc parallel loop tile (1,2.0,true) // { dg-error "" } + for (i = 0; i < 100; i++) + ; + + return 0; +} diff --git a/gcc/tree-core.h b/gcc/tree-core.h index 6b17da7..562f838 100644 --- a/gcc/tree-core.h +++ b/gcc/tree-core.h @@ -432,7 +432,10 @@ enum omp_clause_code { OMP_CLAUSE_NUM_WORKERS, /* OpenACC clause: vector_length (integer-expression). */ - OMP_CLAUSE_VECTOR_LENGTH + OMP_CLAUSE_VECTOR_LENGTH, + + /* OpenACC clause: tile ( size-expr-list ). */ + OMP_CLAUSE_TILE }; #undef DEFTREESTRUCT diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c index 8a2398d..6eeb12f 100644 --- a/gcc/tree-pretty-print.c +++ b/gcc/tree-pretty-print.c @@ -905,6 +905,12 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, int flags) case OMP_CLAUSE_INDEPENDENT: pp_string (pp, "independent"); break; + case OMP_CLAUSE_TILE: + pp_string (pp, "tile("); + dump_generic_node (pp, OMP_CLAUSE_TILE_LIST (clause), + spc, flags, false); + pp_right_paren (pp); + break; default: /* Should never happen. */ diff --git a/gcc/tree.c b/gcc/tree.c index 18d6544..5b9a7bd 100644 --- a/gcc/tree.c +++ b/gcc/tree.c @@ -328,6 +328,7 @@ unsigned const char omp_clause_num_ops[] = 1, /* OMP_CLAUSE_NUM_GANGS */ 1, /* OMP_CLAUSE_NUM_WORKERS */ 1, /* OMP_CLAUSE_VECTOR_LENGTH */ + 1, /* OMP_CLAUSE_TILE */ }; const char * const omp_clause_code_name[] = @@ -398,7 +399,8 @@ const char * const omp_clause_code_name[] = "vector", "num_gangs", "num_workers", - "vector_length" + "vector_length", + "tile" }; @@ -11595,6 +11597,7 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data, case OMP_CLAUSE_DEFAULTMAP: case OMP_CLAUSE_AUTO: case OMP_CLAUSE_SEQ: + case OMP_CLAUSE_TILE: WALK_SUBTREE_TAIL (OMP_CLAUSE_CHAIN (*tp)); case OMP_CLAUSE_LASTPRIVATE: diff --git a/gcc/tree.h b/gcc/tree.h index 65c3117..be18b64a 100644 --- a/gcc/tree.h +++ b/gcc/tree.h @@ -1561,6 +1561,9 @@ extern void protected_set_expr_location (tree, location_t); #define OMP_CLAUSE_DEFAULT_KIND(NODE) \ (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_DEFAULT)->omp_clause.subcode.default_kind) +#define OMP_CLAUSE_TILE_LIST(NODE) \ + OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_TILE), 0) + /* SSA_NAME accessors. */ /* Returns the IDENTIFIER_NODE giving the SSA name a name or NULL_TREE