* Re: [PATCH, OpenACC 2.7] Implement self clause for compute constructs [not found] <e1a5d579-f28c-9417-44ed-ba6248850398@siemens.com> @ 2023-10-25 8:57 ` Thomas Schwinge 2023-10-25 9:10 ` Disentangle handling of OpenACC 'host', 'self' pragma tokens (was: [PATCH, OpenACC 2.7] Implement self clause for compute constructs) Thomas Schwinge ` (2 more replies) 0 siblings, 3 replies; 5+ messages in thread From: Thomas Schwinge @ 2023-10-25 8:57 UTC (permalink / raw) To: Chung-Lin Tang, gcc-patches, fortran; +Cc: Catherine Moore, Tobias Burnus [-- Attachment #1: Type: text/plain, Size: 1641 bytes --] Hi! On 2023-06-13T23:52:25+0800, Chung-Lin Tang via Gcc-patches <gcc-patches@gcc.gnu.org> wrote: > This patch implements the compiler side for the 'self' clause for compute constructs: > parallel, kernels, and serial. > > As you know, the actual "local device" device type for libgomp is not yet implemented, > so the libgomp side is basically just a simple duplicate of what host-fallback is doing, Thanks, and ACK. > though everything else should be completed by this patch. What also is missing is allowing nested OpenACC compute constructs, which GCC currently rejects. (Just removing the nesting restriction isn't sufficient, I think: will also have to think about explicit/implicit data (and other?) clauses in nested compute constructs, for example, so this doesn't seem entirely trivial to implement.) I'm fine to defer that item until actual multicore CPU "device" support emerges (for avoidance of doubt: we're not currently working on that). > Tested on powerpc64le-linux/nvptx, x64_64-linux/amdgcn tests pending. > Is this okay for trunk? With minor textual conflicts resolved, I've pushed this to master branch in commit 3a3596389c2e539cb8fd5dc5784a4e2afe193a2a "OpenACC 2.7: Implement self clause for compute constructs", see attached. I'll then apply/submit a number of follow-on commits. Grüße Thomas ----------------- Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955 [-- Warning: decoded text below may be mangled, UTF-8 assumed --] [-- Attachment #2: 0001-OpenACC-2.7-Implement-self-clause-for-compute-constr.patch --] [-- Type: text/x-diff, Size: 49315 bytes --] From 3a3596389c2e539cb8fd5dc5784a4e2afe193a2a Mon Sep 17 00:00:00 2001 From: Chung-Lin Tang <cltang@codesourcery.com> Date: Tue, 13 Jun 2023 08:44:31 -0700 Subject: [PATCH] OpenACC 2.7: Implement self clause for compute constructs This patch implements the 'self' clause for compute constructs: parallel, kernels, and serial. This clause conditionally uses the local device (the host mult-core CPU) as the executing device of the compute region. The actual implementation of the "local device" device type inside libgomp (presumably using pthreads) is still not yet completed, so the libgomp side is still implemented the exact same as host-fallback mode. (so as of now, it essentially behaves like the 'if' clause with the condition inverted) gcc/c/ChangeLog: * c-parser.cc (c_parser_oacc_compute_clause_self): New function. (c_parser_oacc_all_clauses): Add new 'bool compute_p = false' parameter, add parsing of self clause when compute_p is true. (OACC_KERNELS_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_SELF. (OACC_PARALLEL_CLAUSE_MASK): Likewise, (OACC_SERIAL_CLAUSE_MASK): Likewise. (c_parser_oacc_compute): Adjust call to c_parser_oacc_all_clauses to set compute_p argument to true. * c-typeck.cc (c_finish_omp_clauses): Add OMP_CLAUSE_SELF case. gcc/cp/ChangeLog: * parser.cc (cp_parser_oacc_compute_clause_self): New function. (cp_parser_oacc_all_clauses): Add new 'bool compute_p = false' parameter, add parsing of self clause when compute_p is true. (OACC_KERNELS_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_SELF. (OACC_PARALLEL_CLAUSE_MASK): Likewise, (OACC_SERIAL_CLAUSE_MASK): Likewise. (cp_parser_oacc_compute): Adjust call to c_parser_oacc_all_clauses to set compute_p argument to true. * pt.cc (tsubst_omp_clauses): Add OMP_CLAUSE_SELF case. * semantics.cc (c_finish_omp_clauses): Add OMP_CLAUSE_SELF case, merged with OMP_CLAUSE_IF case. gcc/fortran/ChangeLog: * gfortran.h (typedef struct gfc_omp_clauses): Add self_expr field. * openmp.cc (enum omp_mask2): Add OMP_CLAUSE_SELF. (gfc_match_omp_clauses): Add handling for OMP_CLAUSE_SELF. (OACC_PARALLEL_CLAUSES): Add OMP_CLAUSE_SELF. (OACC_KERNELS_CLAUSES): Likewise. (OACC_SERIAL_CLAUSES): Likewise. (resolve_omp_clauses): Add handling for omp_clauses->self_expr. * trans-openmp.cc (gfc_trans_omp_clauses): Add handling of clauses->self_expr and building of OMP_CLAUSE_SELF tree clause. (gfc_split_omp_clauses): Add handling of self_expr field copy. gcc/ChangeLog: * gimplify.cc (gimplify_scan_omp_clauses): Add OMP_CLAUSE_SELF case. (gimplify_adjust_omp_clauses): Likewise. * omp-expand.cc (expand_omp_target): Add OMP_CLAUSE_SELF expansion code, * omp-low.cc (scan_sharing_clauses): Add OMP_CLAUSE_SELF case. * tree-core.h (enum omp_clause_code): Add OMP_CLAUSE_SELF enum. * tree-nested.cc (convert_nonlocal_omp_clauses): Add OMP_CLAUSE_SELF case. (convert_local_omp_clauses): Likewise. * tree-pretty-print.cc (dump_omp_clause): Add OMP_CLAUSE_SELF case. * tree.cc (omp_clause_num_ops): Add OMP_CLAUSE_SELF entry. (omp_clause_code_name): Likewise. * tree.h (OMP_CLAUSE_SELF_EXPR): New macro. gcc/testsuite/ChangeLog: * c-c++-common/goacc/self-clause-1.c: New test. * c-c++-common/goacc/self-clause-2.c: New test. * gfortran.dg/goacc/self.f95: New test. include/ChangeLog: * gomp-constants.h (GOACC_FLAG_LOCAL_DEVICE): New flag bit value. libgomp/ChangeLog: * oacc-parallel.c (GOACC_parallel_keyed): Add code to handle GOACC_FLAG_LOCAL_DEVICE case. * testsuite/libgomp.oacc-c-c++-common/self-1.c: New test. --- gcc/c/c-parser.cc | 60 +- gcc/c/c-typeck.cc | 1 + gcc/cp/parser.cc | 64 +- gcc/cp/pt.cc | 1 + gcc/cp/semantics.cc | 5 +- gcc/fortran/gfortran.h | 1 + gcc/fortran/openmp.cc | 40 +- gcc/fortran/trans-openmp.cc | 18 + gcc/gimplify.cc | 2 + gcc/omp-expand.cc | 41 + gcc/omp-low.cc | 2 + .../c-c++-common/goacc/self-clause-1.c | 22 + .../c-c++-common/goacc/self-clause-2.c | 17 + gcc/testsuite/gfortran.dg/goacc/self.f95 | 53 + gcc/tree-core.h | 3 + gcc/tree-nested.cc | 2 + gcc/tree-pretty-print.cc | 7 +- gcc/tree.cc | 2 + gcc/tree.h | 2 + include/gomp-constants.h | 2 + libgomp/oacc-parallel.c | 11 + .../libgomp.oacc-c-c++-common/self-1.c | 962 ++++++++++++++++++ 22 files changed, 1306 insertions(+), 12 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/goacc/self-clause-1.c create mode 100644 gcc/testsuite/c-c++-common/goacc/self-clause-2.c create mode 100644 gcc/testsuite/gfortran.dg/goacc/self.f95 create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/self-1.c diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc index 0d468b86bd8..a82f5afeff7 100644 --- a/gcc/c/c-parser.cc +++ b/gcc/c/c-parser.cc @@ -15923,6 +15923,41 @@ c_parser_oacc_clause_wait (c_parser *parser, tree list) return list; } +/* OpenACC 2.7: + self [( expression )] */ + +static tree +c_parser_oacc_compute_clause_self (c_parser *parser, tree list) +{ + tree t; + location_t location = c_parser_peek_token (parser)->location; + if (c_parser_peek_token (parser)->type == CPP_OPEN_PAREN) + { + matching_parens parens; + parens.consume_open (parser); + + location_t loc = c_parser_peek_token (parser)->location; + c_expr expr = c_parser_expr_no_commas (parser, NULL); + expr = convert_lvalue_to_rvalue (loc, expr, true, true); + t = c_objc_common_truthvalue_conversion (loc, expr.value); + t = c_fully_fold (t, false, NULL); + parens.skip_until_found_close (parser); + } + else + t = truthvalue_true_node; + + for (tree c = list; c; c = OMP_CLAUSE_CHAIN (c)) + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_SELF) + { + error_at (location, "too many %<self%> clauses"); + return list; + } + + tree c = build_omp_clause (location, OMP_CLAUSE_SELF); + OMP_CLAUSE_SELF_EXPR (c) = t; + OMP_CLAUSE_CHAIN (c) = list; + return c; +} /* OpenMP 5.0: order ( concurrent ) @@ -18048,7 +18083,8 @@ c_parser_omp_clause_detach (c_parser *parser, tree list) static tree c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask, - const char *where, bool finish_p = true) + const char *where, bool finish_p = true, + bool compute_p = false) { tree clauses = NULL; bool first = true; @@ -18064,7 +18100,18 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask, c_parser_consume_token (parser); here = c_parser_peek_token (parser)->location; - c_kind = c_parser_omp_clause_name (parser); + + /* For OpenACC compute directives */ + if (compute_p + && c_parser_next_token_is (parser, CPP_NAME) + && !strcmp (IDENTIFIER_POINTER (c_parser_peek_token (parser)->value), + "self")) + { + c_kind = PRAGMA_OACC_CLAUSE_SELF; + c_parser_consume_token (parser); + } + else + c_kind = c_parser_omp_clause_name (parser); switch (c_kind) { @@ -18196,6 +18243,10 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask, false, clauses); c_name = "reduction"; break; + case PRAGMA_OACC_CLAUSE_SELF: + clauses = c_parser_oacc_compute_clause_self (parser, clauses); + c_name = "self"; + break; case PRAGMA_OACC_CLAUSE_SEQ: clauses = c_parser_oacc_simple_clause (here, OMP_CLAUSE_SEQ, clauses); @@ -19032,6 +19083,7 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name, | (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) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SELF) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) ) @@ -19052,6 +19104,7 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name, | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_WORKERS) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SELF) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) ) @@ -19070,6 +19123,7 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name, | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_FIRSTPRIVATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SELF) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) ) static tree @@ -19112,7 +19166,7 @@ c_parser_oacc_compute (location_t loc, c_parser *parser, } } - tree clauses = c_parser_oacc_all_clauses (parser, mask, p_name); + tree clauses = c_parser_oacc_all_clauses (parser, mask, p_name, true, true); tree block = c_begin_omp_parallel (); add_stmt (c_parser_omp_structured_block (parser, if_p)); diff --git a/gcc/c/c-typeck.cc b/gcc/c/c-typeck.cc index 112d28fd656..1eee653b3b8 100644 --- a/gcc/c/c-typeck.cc +++ b/gcc/c/c-typeck.cc @@ -15845,6 +15845,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) continue; case OMP_CLAUSE_IF: + case OMP_CLAUSE_SELF: case OMP_CLAUSE_NUM_THREADS: case OMP_CLAUSE_NUM_TEAMS: case OMP_CLAUSE_THREAD_LIMIT: diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc index 5483121f51c..c5a9928ad27 100644 --- a/gcc/cp/parser.cc +++ b/gcc/cp/parser.cc @@ -41192,13 +41192,51 @@ cp_parser_oacc_clause_async (cp_parser *parser, tree list) return list; } +/* OpenACC 2.7: + self [( expression )] */ + +static tree +cp_parser_oacc_compute_clause_self (cp_parser *parser, tree list) +{ + tree t; + location_t location = cp_lexer_peek_token (parser->lexer)->location; + if (cp_lexer_peek_token (parser->lexer)->type == CPP_OPEN_PAREN) + { + matching_parens parens; + parens.consume_open (parser); + t = cp_parser_assignment_expression (parser); + if (t == error_mark_node + || !parens.require_close (parser)) + { + cp_parser_skip_to_closing_parenthesis (parser, /*recovering=*/true, + /*or_comma=*/false, + /*consume_paren=*/true); + return list; + } + } + else + t = truthvalue_true_node; + + for (tree c = list; c; c = OMP_CLAUSE_CHAIN (c)) + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_SELF) + { + error_at (location, "too many %<self%> clauses"); + return list; + } + + tree c = build_omp_clause (location, OMP_CLAUSE_SELF); + OMP_CLAUSE_SELF_EXPR (c) = t; + OMP_CLAUSE_CHAIN (c) = list; + return c; +} + /* Parse all OpenACC clauses. The set clauses allowed by the directive is a bitmask in MASK. Return the list of clauses found. */ static tree cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask, const char *where, cp_token *pragma_tok, - bool finish_p = true) + bool finish_p = true, bool compute_p = false) { tree clauses = NULL; bool first = true; @@ -41218,7 +41256,19 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask, cp_lexer_consume_token (parser->lexer); here = cp_lexer_peek_token (parser->lexer)->location; - c_kind = cp_parser_omp_clause_name (parser); + + /* For OpenACC compute directives */ + if (compute_p + && cp_lexer_next_token_is (parser->lexer, CPP_NAME) + && !strcmp (IDENTIFIER_POINTER + (cp_lexer_peek_token (parser->lexer)->u.value), + "self")) + { + c_kind = PRAGMA_OACC_CLAUSE_SELF; + cp_lexer_consume_token (parser->lexer); + } + else + c_kind = cp_parser_omp_clause_name (parser); switch (c_kind) { @@ -41352,6 +41402,10 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask, false, clauses); c_name = "reduction"; break; + case PRAGMA_OACC_CLAUSE_SELF: + clauses = cp_parser_oacc_compute_clause_self (parser, clauses); + c_name = "self"; + break; case PRAGMA_OACC_CLAUSE_SEQ: clauses = cp_parser_oacc_simple_clause (here, OMP_CLAUSE_SEQ, clauses); @@ -46866,6 +46920,7 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name, | (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) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SELF) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) ) @@ -46886,6 +46941,7 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name, | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRIVATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SELF) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) ) @@ -46904,6 +46960,7 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name, | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_FIRSTPRIVATE) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SELF) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) ) static tree @@ -46949,7 +47006,8 @@ cp_parser_oacc_compute (cp_parser *parser, cp_token *pragma_tok, } } - tree clauses = cp_parser_oacc_all_clauses (parser, mask, p_name, pragma_tok); + tree clauses = cp_parser_oacc_all_clauses (parser, mask, p_name, pragma_tok, + true, true); tree block = begin_omp_parallel (); unsigned int save = cp_parser_begin_omp_structured_block (parser); diff --git a/gcc/cp/pt.cc b/gcc/cp/pt.cc index 1c1c9313118..86c95b278ba 100644 --- a/gcc/cp/pt.cc +++ b/gcc/cp/pt.cc @@ -17418,6 +17418,7 @@ tsubst_omp_clauses (tree clauses, enum c_omp_region_type ort, /* FALLTHRU */ case OMP_CLAUSE_TILE: case OMP_CLAUSE_IF: + case OMP_CLAUSE_SELF: case OMP_CLAUSE_NUM_THREADS: case OMP_CLAUSE_SCHEDULE: case OMP_CLAUSE_COLLAPSE: diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc index dc3c11461fb..72ec72de690 100644 --- a/gcc/cp/semantics.cc +++ b/gcc/cp/semantics.cc @@ -7377,13 +7377,14 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) goto handle_field_decl; case OMP_CLAUSE_IF: - t = OMP_CLAUSE_IF_EXPR (c); + case OMP_CLAUSE_SELF: + t = OMP_CLAUSE_OPERAND (c, 0); t = maybe_convert_cond (t); if (t == error_mark_node) remove = true; else if (!processing_template_decl) t = fold_build_cleanup_point_expr (TREE_TYPE (t), t); - OMP_CLAUSE_IF_EXPR (c) = t; + OMP_CLAUSE_OPERAND (c, 0) = t; break; case OMP_CLAUSE_FINAL: diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h index 9c1a39a19de..ceb2873c03f 100644 --- a/gcc/fortran/gfortran.h +++ b/gcc/fortran/gfortran.h @@ -1546,6 +1546,7 @@ typedef struct gfc_omp_clauses gfc_omp_namelist *lists[OMP_LIST_NUM]; struct gfc_expr *if_expr; struct gfc_expr *if_exprs[OMP_IF_LAST]; + struct gfc_expr *self_expr; struct gfc_expr *final_expr; struct gfc_expr *num_threads; struct gfc_expr *chunk_size; diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc index caa5c9e15a3..083c15e5599 100644 --- a/gcc/fortran/openmp.cc +++ b/gcc/fortran/openmp.cc @@ -1094,6 +1094,7 @@ enum omp_mask2 OMP_CLAUSE_DOACROSS, /* OpenMP 5.2 */ OMP_CLAUSE_ASSUMPTIONS, /* OpenMP 5.1. */ OMP_CLAUSE_USES_ALLOCATORS, /* OpenMP 5.0 */ + OMP_CLAUSE_SELF, /* OpenACC 2.7 */ /* This must come last. */ OMP_MASK2_LAST }; @@ -3519,6 +3520,27 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, else gfc_current_locus = old_loc; } + if ((mask & OMP_CLAUSE_SELF) + && (m = gfc_match_dupl_check (!c->self_expr, "self")) + != MATCH_NO) + { + gcc_assert (!(mask & OMP_CLAUSE_HOST_SELF)); + if (m == MATCH_ERROR) + goto error; + m = gfc_match (" ( %e )", &c->self_expr); + if (m == MATCH_ERROR) + { + gfc_current_locus = old_loc; + break; + } + else if (m == MATCH_NO) + { + c->self_expr = gfc_get_logical_expr (gfc_default_logical_kind, + NULL, true); + needs_space = true; + } + continue; + } if ((mask & OMP_CLAUSE_HOST_SELF) && gfc_match ("self ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], @@ -3791,19 +3813,22 @@ error: | OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \ | OMP_CLAUSE_CREATE | OMP_CLAUSE_NO_CREATE | OMP_CLAUSE_PRESENT \ | OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE \ - | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_WAIT | OMP_CLAUSE_ATTACH) + | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_WAIT | OMP_CLAUSE_ATTACH \ + | OMP_CLAUSE_SELF) #define OACC_KERNELS_CLAUSES \ (omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_NUM_GANGS \ | OMP_CLAUSE_NUM_WORKERS | OMP_CLAUSE_VECTOR_LENGTH | OMP_CLAUSE_DEVICEPTR \ | OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \ | OMP_CLAUSE_CREATE | OMP_CLAUSE_NO_CREATE | OMP_CLAUSE_PRESENT \ - | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_WAIT | OMP_CLAUSE_ATTACH) + | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_WAIT | OMP_CLAUSE_ATTACH \ + | OMP_CLAUSE_SELF) #define OACC_SERIAL_CLAUSES \ (omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_REDUCTION \ | OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \ | OMP_CLAUSE_CREATE | OMP_CLAUSE_NO_CREATE | OMP_CLAUSE_PRESENT \ | OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE \ - | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_WAIT | OMP_CLAUSE_ATTACH) + | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_WAIT | OMP_CLAUSE_ATTACH \ + | OMP_CLAUSE_SELF) #define OACC_DATA_CLAUSES \ (omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_COPY \ | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT | OMP_CLAUSE_CREATE \ @@ -7540,6 +7565,15 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, } } + if (omp_clauses->self_expr) + { + gfc_expr *expr = omp_clauses->self_expr; + if (!gfc_resolve_expr (expr) + || expr->ts.type != BT_LOGICAL || expr->rank != 0) + gfc_error ("SELF clause at %L requires a scalar LOGICAL expression", + &expr->where); + } + if (omp_clauses->final_expr) { gfc_expr *expr = omp_clauses->final_expr; diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc index 7930f2fd5d1..82bbc41b388 100644 --- a/gcc/fortran/trans-openmp.cc +++ b/gcc/fortran/trans-openmp.cc @@ -3966,6 +3966,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, OMP_CLAUSE_IF_EXPR (c) = if_var; omp_clauses = gfc_trans_add_clause (c, omp_clauses); } + for (ifc = 0; ifc < OMP_IF_LAST; ifc++) if (clauses->if_exprs[ifc]) { @@ -4017,6 +4018,21 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, omp_clauses = gfc_trans_add_clause (c, omp_clauses); } + if (clauses->self_expr) + { + tree self_var; + + gfc_init_se (&se, NULL); + gfc_conv_expr (&se, clauses->self_expr); + gfc_add_block_to_block (block, &se.pre); + self_var = gfc_evaluate_now (se.expr, block); + gfc_add_block_to_block (block, &se.post); + + c = build_omp_clause (gfc_get_location (&where), OMP_CLAUSE_SELF); + OMP_CLAUSE_SELF_EXPR (c) = self_var; + omp_clauses = gfc_trans_add_clause (c, omp_clauses); + } + if (clauses->final_expr) { tree final_var; @@ -6615,6 +6631,8 @@ gfc_split_omp_clauses (gfc_code *code, /* And this is copied to all. */ clausesa[GFC_OMP_SPLIT_TARGET].if_expr = code->ext.omp_clauses->if_expr; + clausesa[GFC_OMP_SPLIT_TARGET].self_expr + = code->ext.omp_clauses->self_expr; clausesa[GFC_OMP_SPLIT_TARGET].nowait = code->ext.omp_clauses->nowait; } diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index 22ff1075abb..77f07af6ec5 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -12121,6 +12121,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, } /* Fall through. */ + case OMP_CLAUSE_SELF: case OMP_CLAUSE_FINAL: OMP_CLAUSE_OPERAND (c, 0) = gimple_boolify (OMP_CLAUSE_OPERAND (c, 0)); @@ -13342,6 +13343,7 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, case OMP_CLAUSE_COPYIN: case OMP_CLAUSE_COPYPRIVATE: case OMP_CLAUSE_IF: + case OMP_CLAUSE_SELF: case OMP_CLAUSE_NUM_THREADS: case OMP_CLAUSE_NUM_TEAMS: case OMP_CLAUSE_THREAD_LIMIT: diff --git a/gcc/omp-expand.cc b/gcc/omp-expand.cc index 6172839c719..8576b938102 100644 --- a/gcc/omp-expand.cc +++ b/gcc/omp-expand.cc @@ -10332,6 +10332,47 @@ expand_omp_target (struct omp_region *region) } } + if ((c = omp_find_clause (clauses, OMP_CLAUSE_SELF)) != NULL_TREE) + { + gcc_assert (is_gimple_omp_oacc (entry_stmt) && offloaded); + + edge e = split_block_after_labels (new_bb); + basic_block cond_bb = e->src; + new_bb = e->dest; + remove_edge (e); + + basic_block then_bb = create_empty_bb (cond_bb); + basic_block else_bb = create_empty_bb (then_bb); + set_immediate_dominator (CDI_DOMINATORS, then_bb, cond_bb); + set_immediate_dominator (CDI_DOMINATORS, else_bb, cond_bb); + + tree self_cond = gimple_boolify (OMP_CLAUSE_SELF_EXPR (c)); + stmt = gimple_build_cond_empty (self_cond); + gsi = gsi_last_bb (cond_bb); + gsi_insert_after (&gsi, stmt, GSI_CONTINUE_LINKING); + + tree tmp_var = create_tmp_var (TREE_TYPE (goacc_flags)); + stmt = gimple_build_assign (tmp_var, BIT_IOR_EXPR, goacc_flags, + build_int_cst (integer_type_node, + GOACC_FLAG_LOCAL_DEVICE)); + gsi = gsi_start_bb (then_bb); + gsi_insert_after (&gsi, stmt, GSI_CONTINUE_LINKING); + + gsi = gsi_start_bb (else_bb); + stmt = gimple_build_assign (tmp_var, goacc_flags); + gsi_insert_after (&gsi, stmt, GSI_CONTINUE_LINKING); + + make_edge (cond_bb, then_bb, EDGE_TRUE_VALUE); + make_edge (cond_bb, else_bb, EDGE_FALSE_VALUE); + add_bb_to_loop (then_bb, cond_bb->loop_father); + add_bb_to_loop (else_bb, cond_bb->loop_father); + make_edge (then_bb, new_bb, EDGE_FALLTHRU); + make_edge (else_bb, new_bb, EDGE_FALLTHRU); + + goacc_flags = tmp_var; + gsi = gsi_last_nondebug_bb (new_bb); + } + if (need_device_adjustment) { tree uns = fold_convert (unsigned_type_node, device); diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc index 91ef74f1f6a..161bcfeec05 100644 --- a/gcc/omp-low.cc +++ b/gcc/omp-low.cc @@ -1493,6 +1493,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) case OMP_CLAUSE_FINAL: case OMP_CLAUSE_IF: + case OMP_CLAUSE_SELF: case OMP_CLAUSE_NUM_THREADS: case OMP_CLAUSE_NUM_TEAMS: case OMP_CLAUSE_THREAD_LIMIT: @@ -1920,6 +1921,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) case OMP_CLAUSE_COPYIN: case OMP_CLAUSE_DEFAULT: case OMP_CLAUSE_IF: + case OMP_CLAUSE_SELF: case OMP_CLAUSE_NUM_THREADS: case OMP_CLAUSE_NUM_TEAMS: case OMP_CLAUSE_THREAD_LIMIT: diff --git a/gcc/testsuite/c-c++-common/goacc/self-clause-1.c b/gcc/testsuite/c-c++-common/goacc/self-clause-1.c new file mode 100644 index 00000000000..ed5d072e81f --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/self-clause-1.c @@ -0,0 +1,22 @@ +/* { dg-skip-if "not yet" { c++ } } */ + +void +f (int b) +{ + struct { int i; } *p; + +#pragma acc parallel self self(b) /* { dg-error "too many 'self' clauses" } */ + ; +#pragma acc parallel self(*p) /* { dg-error "used struct type value where scalar is required" } */ + ; + +#pragma acc kernels self self(b) /* { dg-error "too many 'self' clauses" } */ + ; +#pragma acc kernels self(*p) /* { dg-error "used struct type value where scalar is required" } */ + ; + +#pragma acc serial self self(b) /* { dg-error "too many 'self' clauses" } */ + ; +#pragma acc serial self(*p) /* { dg-error "used struct type value where scalar is required" } */ + ; +} diff --git a/gcc/testsuite/c-c++-common/goacc/self-clause-2.c b/gcc/testsuite/c-c++-common/goacc/self-clause-2.c new file mode 100644 index 00000000000..d932ac9a4a6 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/self-clause-2.c @@ -0,0 +1,17 @@ +/* { dg-additional-options "-fdump-tree-gimple" } */ + +void +f (short c) +{ +#pragma acc parallel self(c) copy(c) + /* { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_parallel map\(tofrom:c \[len: [0-9]+\]\) self\(_[0-9]+\)$} 1 "gimple" } } */ + ++c; + +#pragma acc kernels self(c) copy(c) + /* { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_kernels map\(tofrom:c \[len: [0-9]+\]\) self\(_[0-9]+\)$} 1 "gimple" } } */ + ++c; + +#pragma acc serial self(c) copy(c) + /* { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_serial map\(tofrom:c \[len: [0-9]+\]\) self\(_[0-9]+\)$} 1 "gimple" } } */ + ++c; +} diff --git a/gcc/testsuite/gfortran.dg/goacc/self.f95 b/gcc/testsuite/gfortran.dg/goacc/self.f95 new file mode 100644 index 00000000000..4817f16be56 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/self.f95 @@ -0,0 +1,53 @@ +! { dg-do compile } + +program test + implicit none + + logical :: x + integer :: i + + !$acc parallel self () ! { dg-error "Invalid character" } + !$acc parallel self (i) ! { dg-error "scalar LOGICAL expression" } + !$acc end parallel + !$acc parallel self (1) ! { dg-error "scalar LOGICAL expression" } + !$acc end parallel + + !$acc kernels self () ! { dg-error "Invalid character" } + !$acc kernels self (i) ! { dg-error "scalar LOGICAL expression" } + !$acc end kernels + !$acc kernels self (1) ! { dg-error "scalar LOGICAL expression" } + !$acc end kernels + + !$acc serial self () ! { dg-error "Invalid character" } + !$acc serial self (i) ! { dg-error "scalar LOGICAL expression" } + !$acc end serial + !$acc serial self (1) ! { dg-error "scalar LOGICAL expression" } + !$acc end serial + + ! at most one self clause may appear + !$acc parallel self (.false.) self (.false.) { dg-error "Duplicated 'self' clause" } + !$acc kernels self (.false.) self (.false.) { dg-error "Duplicated 'self' clause" } + !$acc serial self (.false.) self (.false.) { dg-error "Duplicated 'self' clause" } + + !$acc parallel self (x) + !$acc end parallel + !$acc parallel self (.true.) + !$acc end parallel + !$acc parallel self (i.gt.1) + !$acc end parallel + + !$acc kernels self (x) + !$acc end kernels + !$acc kernels self (.true.) + !$acc end kernels + !$acc kernels self (i.gt.1) + !$acc end kernels + + !$acc serial self (x) + !$acc end serial + !$acc serial self (.true.) + !$acc end serial + !$acc serial self (i.gt.1) + !$acc end serial + +end program test diff --git a/gcc/tree-core.h b/gcc/tree-core.h index 2c89b655691..cfe37c1d627 100644 --- a/gcc/tree-core.h +++ b/gcc/tree-core.h @@ -527,6 +527,9 @@ enum omp_clause_code { /* OpenACC clause: nohost. */ OMP_CLAUSE_NOHOST, + + /* OpenACC clause: self. */ + OMP_CLAUSE_SELF, }; #undef DEFTREESTRUCT diff --git a/gcc/tree-nested.cc b/gcc/tree-nested.cc index d2fe3fca8af..51a7a032bc3 100644 --- a/gcc/tree-nested.cc +++ b/gcc/tree-nested.cc @@ -1374,6 +1374,7 @@ convert_nonlocal_omp_clauses (tree *pclauses, struct walk_stmt_info *wi) /* FALLTHRU */ case OMP_CLAUSE_FINAL: case OMP_CLAUSE_IF: + case OMP_CLAUSE_SELF: case OMP_CLAUSE_NUM_THREADS: case OMP_CLAUSE_DEPEND: case OMP_CLAUSE_DOACROSS: @@ -2165,6 +2166,7 @@ convert_local_omp_clauses (tree *pclauses, struct walk_stmt_info *wi) /* FALLTHRU */ case OMP_CLAUSE_FINAL: case OMP_CLAUSE_IF: + case OMP_CLAUSE_SELF: case OMP_CLAUSE_NUM_THREADS: case OMP_CLAUSE_DEPEND: case OMP_CLAUSE_DOACROSS: diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc index 58705c57ad5..39ec1df9394 100644 --- a/gcc/tree-pretty-print.cc +++ b/gcc/tree-pretty-print.cc @@ -1453,7 +1453,12 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) false); pp_right_paren (pp); break; - + case OMP_CLAUSE_SELF: + pp_string (pp, "self("); + dump_generic_node (pp, OMP_CLAUSE_SELF_EXPR (clause), + spc, flags, false); + pp_right_paren (pp); + break; default: gcc_unreachable (); } diff --git a/gcc/tree.cc b/gcc/tree.cc index f9fa7b78fff..c38b09c431b 100644 --- a/gcc/tree.cc +++ b/gcc/tree.cc @@ -326,6 +326,7 @@ unsigned const char omp_clause_num_ops[] = 0, /* OMP_CLAUSE_IF_PRESENT */ 0, /* OMP_CLAUSE_FINALIZE */ 0, /* OMP_CLAUSE_NOHOST */ + 1, /* OMP_CLAUSE_SELF */ }; const char * const omp_clause_code_name[] = @@ -417,6 +418,7 @@ const char * const omp_clause_code_name[] = "if_present", "finalize", "nohost", + "self", }; /* Unless specific to OpenACC, we tend to internally maintain OpenMP-centric diff --git a/gcc/tree.h b/gcc/tree.h index 781297c7e77..aaf744c060e 100644 --- a/gcc/tree.h +++ b/gcc/tree.h @@ -1734,6 +1734,8 @@ class auto_suppress_location_wrappers OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_HINT), 0) #define OMP_CLAUSE_FILTER_EXPR(NODE) \ OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_FILTER), 0) +#define OMP_CLAUSE_SELF_EXPR(NODE) \ + OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_SELF), 0) #define OMP_CLAUSE_GRAINSIZE_EXPR(NODE) \ OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_GRAINSIZE),0) diff --git a/include/gomp-constants.h b/include/gomp-constants.h index 8d4e8e81303..89b966e63c6 100644 --- a/include/gomp-constants.h +++ b/include/gomp-constants.h @@ -304,6 +304,8 @@ enum gomp_map_kind /* Force host fallback execution. */ #define GOACC_FLAG_HOST_FALLBACK (1 << 0) +/* Execute on local device (i.e. host multicore CPU). */ +#define GOACC_FLAG_LOCAL_DEVICE (1 << 1) /* For legacy reasons, in the ABI, the GOACC_FLAGs are encoded as an inverted bitmask. */ diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c index 363e6656982..cf37a1bdd7d 100644 --- a/libgomp/oacc-parallel.c +++ b/libgomp/oacc-parallel.c @@ -193,6 +193,17 @@ GOACC_parallel_keyed (int flags_m, void (*fn) (void *), goacc_restore_bind (); goto out_prof; } + else if (flags & GOACC_FLAG_LOCAL_DEVICE) + { + /* TODO: a proper pthreads based "multi-core CPU" local device + implementation. Currently, this is still the same as host-fallback. */ + prof_info.device_type = acc_device_host; + api_info.device_type = prof_info.device_type; + goacc_save_and_set_bind (acc_device_host); + fn (hostaddrs); + goacc_restore_bind (); + goto out_prof; + } else if (acc_device_type (acc_dev->type) == acc_device_host) { fn (hostaddrs); diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/self-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/self-1.c new file mode 100644 index 00000000000..752e16e8545 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/self-1.c @@ -0,0 +1,962 @@ +#include <openacc.h> +#include <stdlib.h> +#include <stdbool.h> + +#define N 32 + +int +main(int argc, char **argv) +{ + float *a, *b, *d_a, *d_b, exp, exp2; + int i; + const int one = 1; + const int zero = 0; + int n; + + a = (float *) malloc (N * sizeof (float)); + b = (float *) malloc (N * sizeof (float)); + d_a = (float *) acc_malloc (N * sizeof (float)); + d_b = (float *) acc_malloc (N * sizeof (float)); + + for (i = 0; i < N; i++) + a[i] = 4.0; + +#pragma acc parallel copyin(a[0:N]) copyout(b[0:N]) self(0) + { + int ii; + + for (ii = 0; ii < N; ii++) + { + if (acc_on_device (acc_device_host)) + b[ii] = a[ii] + 1; + else + b[ii] = a[ii]; + } + } + +#if ACC_MEM_SHARED + exp = 5.0; +#else + exp = 4.0; +#endif + + for (i = 0; i < N; i++) + { + if (b[i] != exp) + abort(); + } + + for (i = 0; i < N; i++) + a[i] = 16.0; + +#pragma acc parallel self(1) + { + int ii; + + for (ii = 0; ii < N; ii++) + { + if (acc_on_device (acc_device_host)) + b[ii] = a[ii] + 1; + else + b[ii] = a[ii]; + } + } + + for (i = 0; i < N; i++) + { + if (b[i] != 17.0) + abort(); + } + + for (i = 0; i < N; i++) + a[i] = 8.0; + +#pragma acc parallel copyin(a[0:N]) copyout(b[0:N]) self(!one) + { + int ii; + + for (ii = 0; ii < N; ii++) + { + if (acc_on_device (acc_device_host)) + b[ii] = a[ii] + 1; + else + b[ii] = a[ii]; + } + } + +#if ACC_MEM_SHARED + exp = 9.0; +#else + exp = 8.0; +#endif + + for (i = 0; i < N; i++) + { + if (b[i] != exp) + abort(); + } + + for (i = 0; i < N; i++) + a[i] = 22.0; + +#pragma acc parallel self(!zero) + { + int ii; + + for (ii = 0; ii < N; ii++) + { + if (acc_on_device (acc_device_host)) + b[ii] = a[ii] + 1; + else + b[ii] = a[ii]; + } + } + + for (i = 0; i < N; i++) + { + if (b[i] != 23.0) + abort(); + } + + for (i = 0; i < N; i++) + a[i] = 16.0; + +#pragma acc parallel copyin(a[0:N]) copyout(b[0:N]) self(false) + { + int ii; + + for (ii = 0; ii < N; ii++) + { + if (acc_on_device (acc_device_host)) + b[ii] = a[ii] + 1; + else + b[ii] = a[ii]; + } + } + +#if ACC_MEM_SHARED + exp = 17.0; +#else + exp = 16.0; +#endif + + for (i = 0; i < N; i++) + { + if (b[i] != exp) + abort(); + } + + for (i = 0; i < N; i++) + a[i] = 76.0; + +#pragma acc parallel self(true) + { + int ii; + + for (ii = 0; ii < N; ii++) + { + if (acc_on_device (acc_device_host)) + b[ii] = a[ii] + 1; + else + b[ii] = a[ii]; + } + } + + for (i = 0; i < N; i++) + { + if (b[i] != 77.0) + abort(); + } + + for (i = 0; i < N; i++) + a[i] = 22.0; + + n = 1; + +#pragma acc parallel copyin(a[0:N]) copyout(b[0:N]) self(!n) + { + int ii; + + for (ii = 0; ii < N; ii++) + { + if (acc_on_device (acc_device_host)) + b[ii] = a[ii] + 1; + else + b[ii] = a[ii]; + } + } + +#if ACC_MEM_SHARED + exp = 23.0; +#else + exp = 22.0; +#endif + + for (i = 0; i < N; i++) + { + if (b[i] != exp) + abort(); + } + + for (i = 0; i < N; i++) + a[i] = 18.0; + + n = 0; + +#pragma acc parallel self(!n) + { + int ii; + + for (ii = 0; ii < N; ii++) + { + if (acc_on_device (acc_device_host)) + b[ii] = a[ii] + 1; + else + b[ii] = a[ii]; + } + } + + for (i = 0; i < N; i++) + { + if (b[i] != 19.0) + abort(); + } + + for (i = 0; i < N; i++) + a[i] = 49.0; + + n = 1; + +#pragma acc parallel copyin(a[0:N]) copyout(b[0:N]) self(!(n + n)) + { + int ii; + + for (ii = 0; ii < N; ii++) + { + if (acc_on_device (acc_device_host)) + b[ii] = a[ii] + 1; + else + b[ii] = a[ii]; + } + } + +#if ACC_MEM_SHARED + exp = 50.0; +#else + exp = 49.0; +#endif + + for (i = 0; i < N; i++) + { + if (b[i] != exp) + abort(); + } + + for (i = 0; i < N; i++) + a[i] = 38.0; + + n = 0; + +#pragma acc parallel self(!(n + n)) + { + int ii; + + for (ii = 0; ii < N; ii++) + { + if (acc_on_device (acc_device_host)) + b[ii] = a[ii] + 1; + else + b[ii] = a[ii]; + } + } + + for (i = 0; i < N; i++) + { + if (b[i] != 39.0) + abort(); + } + + for (i = 0; i < N; i++) + a[i] = 91.0; + +#pragma acc parallel copyin(a[0:N]) copyout(b[0:N]) self(!(-2)) + { + int ii; + + for (ii = 0; ii < N; ii++) + { + if (acc_on_device (acc_device_host)) + b[ii] = a[ii] + 1; + else + b[ii] = a[ii]; + } + } + +#if ACC_MEM_SHARED + exp = 92.0; +#else + exp = 91.0; +#endif + + for (i = 0; i < N; i++) + { + if (b[i] != exp) + abort(); + } + + for (i = 0; i < N; i++) + a[i] = 43.0; + +#pragma acc parallel copyin(a[0:N]) copyout(b[0:N]) self(one != 1) + { + int ii; + + for (ii = 0; ii < N; ii++) + { + if (acc_on_device (acc_device_host)) + b[ii] = a[ii] + 1; + else + b[ii] = a[ii]; + } + } + +#if ACC_MEM_SHARED + exp = 44.0; +#else + exp = 43.0; +#endif + + for (i = 0; i < N; i++) + { + if (b[i] != exp) + abort(); + } + + for (i = 0; i < N; i++) + a[i] = 87.0; + +#pragma acc parallel self(one != 0) + { + int ii; + + for (ii = 0; ii < N; ii++) + { + if (acc_on_device (acc_device_host)) + b[ii] = a[ii] + 1; + else + b[ii] = a[ii]; + } + } + + for (i = 0; i < N; i++) + { + if (b[i] != 88.0) + abort(); + } + + for (i = 0; i < N; i++) + { + a[i] = 3.0; + b[i] = 9.0; + } + +#if ACC_MEM_SHARED + exp = 0.0; + exp2 = 0.0; +#else + acc_map_data (a, d_a, N * sizeof (float)); + acc_map_data (b, d_b, N * sizeof (float)); + exp = 3.0; + exp2 = 9.0; +#endif + +#pragma acc update device(a[0:N], b[0:N]) if(1) + + for (i = 0; i < N; i++) + { + a[i] = 0.0; + b[i] = 0.0; + } + +#pragma acc update host(a[0:N], b[0:N]) if(1) + + for (i = 0; i < N; i++) + { + if (a[i] != exp) + abort(); + + if (b[i] != exp2) + abort(); + } + + for (i = 0; i < N; i++) + { + a[i] = 6.0; + b[i] = 12.0; + } + +#pragma acc update device(a[0:N], b[0:N]) if(0) + + for (i = 0; i < N; i++) + { + a[i] = 0.0; + b[i] = 0.0; + } + +#pragma acc update host(a[0:N], b[0:N]) if(1) + + for (i = 0; i < N; i++) + { + if (a[i] != exp) + abort(); + + if (b[i] != exp2) + abort(); + } + + for (i = 0; i < N; i++) + { + a[i] = 26.0; + b[i] = 21.0; + } + +#pragma acc update device(a[0:N], b[0:N]) if(1) + + for (i = 0; i < N; i++) + { + a[i] = 0.0; + b[i] = 0.0; + } + +#pragma acc update host(a[0:N], b[0:N]) if(0) + + for (i = 0; i < N; i++) + { + if (a[i] != 0.0) + abort(); + + if (b[i] != 0.0) + abort(); + } + +#if !ACC_MEM_SHARED + acc_unmap_data (a); + acc_unmap_data (b); +#endif + + acc_free (d_a); + acc_free (d_b); + + for (i = 0; i < N; i++) + { + a[i] = 4.0; + b[i] = 0.0; + } + +#pragma acc data copyin(a[0:N]) copyout(b[0:N]) if(1) +{ +#pragma acc parallel present(a[0:N]) + { + int ii; + + for (ii = 0; ii < N; ii++) + { + b[ii] = a[ii]; + } + } +} + + for (i = 0; i < N; i++) + { + if (b[i] != 4.0) + abort(); + } + + for (i = 0; i < N; i++) + { + a[i] = 8.0; + b[i] = 1.0; + } + +#pragma acc data copyin(a[0:N]) copyout(b[0:N]) if(0) +{ +#if !ACC_MEM_SHARED + if (acc_is_present (a, N * sizeof (float))) + abort (); +#endif + +#if !ACC_MEM_SHARED + if (acc_is_present (b, N * sizeof (float))) + abort (); +#endif +} + + for (i = 0; i < N; i++) + { + a[i] = 18.0; + b[i] = 21.0; + } + +#pragma acc data copyin(a[0:N]) if(1) +{ +#if !ACC_MEM_SHARED + if (!acc_is_present (a, N * sizeof (float))) + abort (); +#endif + +#pragma acc data copyout(b[0:N]) if(0) + { +#if !ACC_MEM_SHARED + if (acc_is_present (b, N * sizeof (float))) + abort (); +#endif + +#pragma acc data copyout(b[0:N]) if(1) + { +#pragma acc parallel present(a[0:N]) present(b[0:N]) + { + int ii; + + for (ii = 0; ii < N; ii++) + { + b[ii] = a[ii]; + } + } + } + +#if !ACC_MEM_SHARED + if (acc_is_present (b, N * sizeof (float))) + abort (); +#endif + } +} + + for (i = 0; i < N; i++) + { + if (b[i] != 18.0) + abort (); + } + +#pragma acc enter data copyin (b[0:N]) if (0) + +#if !ACC_MEM_SHARED + if (acc_is_present (b, N * sizeof (float))) + abort (); +#endif + +#pragma acc exit data delete (b[0:N]) if (0) + +#pragma acc enter data copyin (b[0:N]) if (1) + +#if !ACC_MEM_SHARED + if (!acc_is_present (b, N * sizeof (float))) + abort (); +#endif + +#pragma acc exit data delete (b[0:N]) if (1) + +#if !ACC_MEM_SHARED + if (acc_is_present (b, N * sizeof (float))) + abort (); +#endif + +#pragma acc enter data copyin (b[0:N]) if (zero) + +#if !ACC_MEM_SHARED + if (acc_is_present (b, N * sizeof (float))) + abort (); +#endif + +#pragma acc exit data delete (b[0:N]) if (zero) + +#pragma acc enter data copyin (b[0:N]) if (one) + +#if !ACC_MEM_SHARED + if (!acc_is_present (b, N * sizeof (float))) + abort (); +#endif + +#pragma acc exit data delete (b[0:N]) if (one) + +#if !ACC_MEM_SHARED + if (acc_is_present (b, N * sizeof (float))) + abort (); +#endif + +#pragma acc enter data copyin (b[0:N]) if (one == 0) + +#if !ACC_MEM_SHARED + if (acc_is_present (b, N * sizeof (float))) + abort (); +#endif + +#pragma acc exit data delete (b[0:N]) if (one == 0) + +#pragma acc enter data copyin (b[0:N]) if (one == 1) + +#if !ACC_MEM_SHARED + if (!acc_is_present (b, N * sizeof (float))) + abort (); +#endif + +#pragma acc exit data delete (b[0:N]) if (one == 1) + +#if !ACC_MEM_SHARED + if (acc_is_present (b, N * sizeof (float))) + abort (); +#endif + + for (i = 0; i < N; i++) + a[i] = 4.0; + +#pragma acc kernels copyin(a[0:N]) copyout(b[0:N]) self(0) + { + int ii; + + for (ii = 0; ii < N; ii++) + { + if (acc_on_device (acc_device_host)) + b[ii] = a[ii] + 1; + else + b[ii] = a[ii]; + } + } + +#if ACC_MEM_SHARED + exp = 5.0; +#else + exp = 4.0; +#endif + + for (i = 0; i < N; i++) + { + if (b[i] != exp) + abort(); + } + + for (i = 0; i < N; i++) + a[i] = 16.0; + +#pragma acc kernels self(1) + { + int ii; + + for (ii = 0; ii < N; ii++) + { + if (acc_on_device (acc_device_host)) + b[ii] = a[ii] + 1; + else + b[ii] = a[ii]; + } + } + + for (i = 0; i < N; i++) + { + if (b[i] != 17.0) + abort(); + } + + for (i = 0; i < N; i++) + a[i] = 8.0; + +#pragma acc kernels copyin(a[0:N]) copyout(b[0:N]) self(!one) + { + int ii; + + for (ii = 0; ii < N; ii++) + { + if (acc_on_device (acc_device_host)) + b[ii] = a[ii] + 1; + else + b[ii] = a[ii]; + } + } + +#if ACC_MEM_SHARED + exp = 9.0; +#else + exp = 8.0; +#endif + + for (i = 0; i < N; i++) + { + if (b[i] != exp) + abort(); + } + + for (i = 0; i < N; i++) + a[i] = 22.0; + +#pragma acc kernels self(!zero) + { + int ii; + + for (ii = 0; ii < N; ii++) + { + if (acc_on_device (acc_device_host)) + b[ii] = a[ii] + 1; + else + b[ii] = a[ii]; + } + } + + for (i = 0; i < N; i++) + { + if (b[i] != 23.0) + abort(); + } + + for (i = 0; i < N; i++) + a[i] = 16.0; + +#pragma acc kernels copyin(a[0:N]) copyout(b[0:N]) self(false) + { + int ii; + + for (ii = 0; ii < N; ii++) + { + if (acc_on_device (acc_device_host)) + b[ii] = a[ii] + 1; + else + b[ii] = a[ii]; + } + } + +#if ACC_MEM_SHARED + exp = 17.0; +#else + exp = 16.0; +#endif + + for (i = 0; i < N; i++) + { + if (b[i] != exp) + abort(); + } + + for (i = 0; i < N; i++) + a[i] = 76.0; + +#pragma acc kernels self(true) + { + int ii; + + for (ii = 0; ii < N; ii++) + { + if (acc_on_device (acc_device_host)) + b[ii] = a[ii] + 1; + else + b[ii] = a[ii]; + } + } + + for (i = 0; i < N; i++) + { + if (b[i] != 77.0) + abort(); + } + + for (i = 0; i < N; i++) + a[i] = 22.0; + + n = 1; + +#pragma acc kernels copyin(a[0:N]) copyout(b[0:N]) self(!n) + { + int ii; + + for (ii = 0; ii < N; ii++) + { + if (acc_on_device (acc_device_host)) + b[ii] = a[ii] + 1; + else + b[ii] = a[ii]; + } + } + +#if ACC_MEM_SHARED + exp = 23.0; +#else + exp = 22.0; +#endif + + for (i = 0; i < N; i++) + { + if (b[i] != exp) + abort(); + } + + for (i = 0; i < N; i++) + a[i] = 18.0; + + n = 0; + +#pragma acc kernels self(!n) + { + int ii; + + for (ii = 0; ii < N; ii++) + { + if (acc_on_device (acc_device_host)) + b[ii] = a[ii] + 1; + else + b[ii] = a[ii]; + } + } + + for (i = 0; i < N; i++) + { + if (b[i] != 19.0) + abort(); + } + + for (i = 0; i < N; i++) + a[i] = 49.0; + + n = 1; + +#pragma acc kernels copyin(a[0:N]) copyout(b[0:N]) self((n + n) == 0) + { + int ii; + + for (ii = 0; ii < N; ii++) + { + if (acc_on_device (acc_device_host)) + b[ii] = a[ii] + 1; + else + b[ii] = a[ii]; + } + } + +#if ACC_MEM_SHARED + exp = 50.0; +#else + exp = 49.0; +#endif + + for (i = 0; i < N; i++) + { + if (b[i] != exp) + abort(); + } + + for (i = 0; i < N; i++) + a[i] = 38.0; + + n = 0; + +#pragma acc kernels self(!(n + n)) + { + int ii; + + for (ii = 0; ii < N; ii++) + { + if (acc_on_device (acc_device_host)) + b[ii] = a[ii] + 1; + else + b[ii] = a[ii]; + } + } + + for (i = 0; i < N; i++) + { + if (b[i] != 39.0) + abort(); + } + + for (i = 0; i < N; i++) + a[i] = 91.0; + +#pragma acc kernels copyin(a[0:N]) copyout(b[0:N]) self(!(-2)) + { + int ii; + + for (ii = 0; ii < N; ii++) + { + if (acc_on_device (acc_device_host)) + b[ii] = a[ii] + 1; + else + b[ii] = a[ii]; + } + } + +#if ACC_MEM_SHARED + exp = 92.0; +#else + exp = 91.0; +#endif + + for (i = 0; i < N; i++) + { + if (b[i] != exp) + abort(); + } + + for (i = 0; i < N; i++) + a[i] = 43.0; + +#pragma acc kernels copyin(a[0:N]) copyout(b[0:N]) self(one != 1) + { + int ii; + + for (ii = 0; ii < N; ii++) + { + if (acc_on_device (acc_device_host)) + b[ii] = a[ii] + 1; + else + b[ii] = a[ii]; + } + } + +#if ACC_MEM_SHARED + exp = 44.0; +#else + exp = 43.0; +#endif + + for (i = 0; i < N; i++) + { + if (b[i] != exp) + abort(); + } + + for (i = 0; i < N; i++) + a[i] = 87.0; + +#pragma acc kernels self(one != 0) + { + int ii; + + for (ii = 0; ii < N; ii++) + { + if (acc_on_device (acc_device_host)) + b[ii] = a[ii] + 1; + else + b[ii] = a[ii]; + } + } + + for (i = 0; i < N; i++) + { + if (b[i] != 88.0) + abort(); + } + + for (i = 0; i < N; i++) + { + a[i] = 3.0; + b[i] = 9.0; + } + +#if ACC_MEM_SHARED + exp = 0.0; + exp2 = 0.0; +#else + acc_map_data (a, d_a, N * sizeof (float)); + acc_map_data (b, d_b, N * sizeof (float)); + exp = 3.0; + exp2 = 9.0; +#endif + + return 0; +} -- 2.34.1 ^ permalink raw reply [flat|nested] 5+ messages in thread
* Disentangle handling of OpenACC 'host', 'self' pragma tokens (was: [PATCH, OpenACC 2.7] Implement self clause for compute constructs) 2023-10-25 8:57 ` [PATCH, OpenACC 2.7] Implement self clause for compute constructs Thomas Schwinge @ 2023-10-25 9:10 ` Thomas Schwinge 2023-10-25 9:29 ` Extend test suite coverage for OpenACC 'self' clause for compute constructs " Thomas Schwinge 2023-10-25 9:44 ` Minor fixes for OpenACC/Fortran 'self' clause for compute constructs (was: [PATCH, OpenACC 2.7] Implement self clause for compute constructs) Thomas Schwinge 2 siblings, 0 replies; 5+ messages in thread From: Thomas Schwinge @ 2023-10-25 9:10 UTC (permalink / raw) To: Chung-Lin Tang, gcc-patches, fortran; +Cc: Catherine Moore, Tobias Burnus [-- Attachment #1: Type: text/plain, Size: 1951 bytes --] Hi! On 2023-10-25T10:57:06+0200, I wrote: > With minor textual conflicts resolved, I've pushed this to master branch > in commit 3a3596389c2e539cb8fd5dc5784a4e2afe193a2a > "OpenACC 2.7: Implement self clause for compute constructs", see > attached. > > > I'll then apply/submit a number of follow-on commits. I found this: > --- a/gcc/c/c-parser.cc > +++ b/gcc/c/c-parser.cc > static tree > c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask, > - const char *where, bool finish_p = true) > + const char *where, bool finish_p = true, > + bool compute_p = false) > { > tree clauses = NULL; > bool first = true; > @@ -18064,7 +18100,18 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask, > c_parser_consume_token (parser); > > here = c_parser_peek_token (parser)->location; > - c_kind = c_parser_omp_clause_name (parser); > + > + /* For OpenACC compute directives */ > + if (compute_p > + && c_parser_next_token_is (parser, CPP_NAME) > + && !strcmp (IDENTIFIER_POINTER (c_parser_peek_token (parser)->value), > + "self")) > + { > + c_kind = PRAGMA_OACC_CLAUSE_SELF; > + c_parser_consume_token (parser); > + } > + else > + c_kind = c_parser_omp_clause_name (parser); ..., and similarly in the C++ and (to a lesser extent) Fortran front ends a bit twisted, and pushed to master branch commit c92509d9fd98e02d17ab1610f696c88f606dcdf4 "Disentangle handling of OpenACC 'host', 'self' pragma tokens", see attached. Grüße Thomas ----------------- Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955 [-- Warning: decoded text below may be mangled, UTF-8 assumed --] [-- Attachment #2: 0001-Disentangle-handling-of-OpenACC-host-self-pragma-tok.patch --] [-- Type: text/x-diff, Size: 13693 bytes --] From c92509d9fd98e02d17ab1610f696c88f606dcdf4 Mon Sep 17 00:00:00 2001 From: Thomas Schwinge <thomas@codesourcery.com> Date: Fri, 20 Oct 2023 14:47:58 +0200 Subject: [PATCH] Disentangle handling of OpenACC 'host', 'self' pragma tokens 'gcc/c-family/c-pragma.h:pragma_omp_clause' already defines 'PRAGMA_OACC_CLAUSE_SELF', but it has no longer been used for the 'update' directive's 'self' clause as of 2018 commit 829c6349e96c5bfa8603aaef8858b38e237a2f33 (Subversion r261813) "Update OpenACC data clause semantics to the 2.5 behavior". That one instead mapped the 'self' pragma token to the 'host' one (same semantics). That means that we're later not able to tell whether originally we had seen 'self' or 'host', which was OK as long as only the 'update' directive had a 'self' clause. However, as of recent commit 3a3596389c2e539cb8fd5dc5784a4e2afe193a2a "OpenACC 2.7: Implement self clause for compute constructs", also OpenACC compute constructs may have a 'self' clause -- with different semantics. That means, we need to know which OpenACC directive we're parsing clauses for, which can be done in a simpler way than in that commit, similar to how the OpenMP 'to' clause is handled. While at that, clarify that (already in OpenACC 2.0a) "The 'host' clause is a synonym for the 'self' clause." -- not the other way round. gcc/c/ * c-parser.cc (c_parser_omp_clause_name): Return 'PRAGMA_OACC_CLAUSE_SELF' for "self". (c_parser_oacc_data_clause, OACC_UPDATE_CLAUSE_MASK): Adjust. (c_parser_oacc_all_clauses): Remove 'bool compute_p' formal parameter, and instead locally determine whether we're called for an OpenACC compute construct or OpenACC 'update' directive. (c_parser_oacc_compute): Adjust. gcc/cp/ * parser.cc (cp_parser_omp_clause_name): Return 'PRAGMA_OACC_CLAUSE_SELF' for "self". (cp_parser_oacc_data_clause, OACC_UPDATE_CLAUSE_MASK): Adjust. (cp_parser_oacc_all_clauses): Remove 'bool compute_p' formal parameter, and instead locally determine whether we're called for an OpenACC compute construct or OpenACC 'update' directive. (cp_parser_oacc_compute): Adjust. gcc/fortran/ * openmp.cc (omp_mask2): Split 'OMP_CLAUSE_HOST_SELF' into 'OMP_CLAUSE_SELF', 'OMP_CLAUSE_HOST'. (gfc_match_omp_clauses, OACC_UPDATE_CLAUSES): Adjust. --- gcc/c/c-parser.cc | 38 +++++++++++++++++--------------------- gcc/cp/parser.cc | 39 +++++++++++++++++---------------------- gcc/fortran/openmp.cc | 27 ++++++++++++++------------- 3 files changed, 48 insertions(+), 56 deletions(-) diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc index a82f5afeff7..5213a57a1ec 100644 --- a/gcc/c/c-parser.cc +++ b/gcc/c/c-parser.cc @@ -14061,8 +14061,8 @@ c_parser_omp_clause_name (c_parser *parser) result = PRAGMA_OMP_CLAUSE_SCHEDULE; else if (!strcmp ("sections", p)) result = PRAGMA_OMP_CLAUSE_SECTIONS; - else if (!strcmp ("self", p)) /* "self" is a synonym for "host". */ - result = PRAGMA_OACC_CLAUSE_HOST; + else if (!strcmp ("self", p)) + result = PRAGMA_OACC_CLAUSE_SELF; else if (!strcmp ("seq", p)) result = PRAGMA_OACC_CLAUSE_SEQ; else if (!strcmp ("shared", p)) @@ -14583,9 +14583,6 @@ c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind, case PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT: kind = GOMP_MAP_DEVICE_RESIDENT; break; - case PRAGMA_OACC_CLAUSE_HOST: - kind = GOMP_MAP_FORCE_FROM; - break; case PRAGMA_OACC_CLAUSE_LINK: kind = GOMP_MAP_LINK; break; @@ -14595,6 +14592,11 @@ c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind, case PRAGMA_OACC_CLAUSE_PRESENT: kind = GOMP_MAP_FORCE_PRESENT; break; + case PRAGMA_OACC_CLAUSE_SELF: + /* "The 'host' clause is a synonym for the 'self' clause." */ + case PRAGMA_OACC_CLAUSE_HOST: + kind = GOMP_MAP_FORCE_FROM; + break; default: gcc_unreachable (); } @@ -18083,8 +18085,7 @@ c_parser_omp_clause_detach (c_parser *parser, tree list) static tree c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask, - const char *where, bool finish_p = true, - bool compute_p = false) + const char *where, bool finish_p = true) { tree clauses = NULL; bool first = true; @@ -18100,18 +18101,7 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask, c_parser_consume_token (parser); here = c_parser_peek_token (parser)->location; - - /* For OpenACC compute directives */ - if (compute_p - && c_parser_next_token_is (parser, CPP_NAME) - && !strcmp (IDENTIFIER_POINTER (c_parser_peek_token (parser)->value), - "self")) - { - c_kind = PRAGMA_OACC_CLAUSE_SELF; - c_parser_consume_token (parser); - } - else - c_kind = c_parser_omp_clause_name (parser); + c_kind = c_parser_omp_clause_name (parser); switch (c_kind) { @@ -18244,7 +18234,12 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask, c_name = "reduction"; break; case PRAGMA_OACC_CLAUSE_SELF: - clauses = c_parser_oacc_compute_clause_self (parser, clauses); + if ((mask & (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_HOST)) == 0) + /* OpenACC compute construct */ + clauses = c_parser_oacc_compute_clause_self (parser, clauses); + else + /* OpenACC 'update' directive */ + clauses = c_parser_oacc_data_clause (parser, c_kind, clauses); c_name = "self"; break; case PRAGMA_OACC_CLAUSE_SEQ: @@ -19166,7 +19161,7 @@ c_parser_oacc_compute (location_t loc, c_parser *parser, } } - tree clauses = c_parser_oacc_all_clauses (parser, mask, p_name, true, true); + tree clauses = c_parser_oacc_all_clauses (parser, mask, p_name); tree block = c_begin_omp_parallel (); add_stmt (c_parser_omp_structured_block (parser, if_p)); @@ -19366,6 +19361,7 @@ c_finish_oacc_routine (struct oacc_routine_data *data, tree fndecl, | (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_SELF) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) ) static void diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc index c5a9928ad27..5a6c416d932 100644 --- a/gcc/cp/parser.cc +++ b/gcc/cp/parser.cc @@ -37566,8 +37566,8 @@ cp_parser_omp_clause_name (cp_parser *parser) result = PRAGMA_OMP_CLAUSE_SCHEDULE; else if (!strcmp ("sections", p)) result = PRAGMA_OMP_CLAUSE_SECTIONS; - else if (!strcmp ("self", p)) /* "self" is a synonym for "host". */ - result = PRAGMA_OACC_CLAUSE_HOST; + else if (!strcmp ("self", p)) + result = PRAGMA_OACC_CLAUSE_SELF; else if (!strcmp ("seq", p)) result = PRAGMA_OACC_CLAUSE_SEQ; else if (!strcmp ("shared", p)) @@ -38004,9 +38004,6 @@ cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind, case PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT: kind = GOMP_MAP_DEVICE_RESIDENT; break; - case PRAGMA_OACC_CLAUSE_HOST: - kind = GOMP_MAP_FORCE_FROM; - break; case PRAGMA_OACC_CLAUSE_LINK: kind = GOMP_MAP_LINK; break; @@ -38016,6 +38013,11 @@ cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind, case PRAGMA_OACC_CLAUSE_PRESENT: kind = GOMP_MAP_FORCE_PRESENT; break; + case PRAGMA_OACC_CLAUSE_SELF: + /* "The 'host' clause is a synonym for the 'self' clause." */ + case PRAGMA_OACC_CLAUSE_HOST: + kind = GOMP_MAP_FORCE_FROM; + break; default: gcc_unreachable (); } @@ -41236,7 +41238,7 @@ cp_parser_oacc_compute_clause_self (cp_parser *parser, tree list) static tree cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask, const char *where, cp_token *pragma_tok, - bool finish_p = true, bool compute_p = false) + bool finish_p = true) { tree clauses = NULL; bool first = true; @@ -41256,19 +41258,7 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask, cp_lexer_consume_token (parser->lexer); here = cp_lexer_peek_token (parser->lexer)->location; - - /* For OpenACC compute directives */ - if (compute_p - && cp_lexer_next_token_is (parser->lexer, CPP_NAME) - && !strcmp (IDENTIFIER_POINTER - (cp_lexer_peek_token (parser->lexer)->u.value), - "self")) - { - c_kind = PRAGMA_OACC_CLAUSE_SELF; - cp_lexer_consume_token (parser->lexer); - } - else - c_kind = cp_parser_omp_clause_name (parser); + c_kind = cp_parser_omp_clause_name (parser); switch (c_kind) { @@ -41403,7 +41393,12 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask, c_name = "reduction"; break; case PRAGMA_OACC_CLAUSE_SELF: - clauses = cp_parser_oacc_compute_clause_self (parser, clauses); + if ((mask & (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_HOST)) == 0) + /* OpenACC compute construct */ + clauses = cp_parser_oacc_compute_clause_self (parser, clauses); + else + /* OpenACC 'update' directive */ + clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses); c_name = "self"; break; case PRAGMA_OACC_CLAUSE_SEQ: @@ -47006,8 +47001,7 @@ cp_parser_oacc_compute (cp_parser *parser, cp_token *pragma_tok, } } - tree clauses = cp_parser_oacc_all_clauses (parser, mask, p_name, pragma_tok, - true, true); + tree clauses = cp_parser_oacc_all_clauses (parser, mask, p_name, pragma_tok); tree block = begin_omp_parallel (); unsigned int save = cp_parser_begin_omp_structured_block (parser); @@ -47026,6 +47020,7 @@ cp_parser_oacc_compute (cp_parser *parser, cp_token *pragma_tok, | (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_SELF) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT)) static tree diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc index 083c15e5599..2e2e23d567b 100644 --- a/gcc/fortran/openmp.cc +++ b/gcc/fortran/openmp.cc @@ -1080,7 +1080,8 @@ enum omp_mask2 OMP_CLAUSE_INDEPENDENT, OMP_CLAUSE_USE_DEVICE, OMP_CLAUSE_DEVICE_RESIDENT, - OMP_CLAUSE_HOST_SELF, + OMP_CLAUSE_SELF, + OMP_CLAUSE_HOST, OMP_CLAUSE_WAIT, OMP_CLAUSE_DELETE, OMP_CLAUSE_AUTO, @@ -1094,7 +1095,6 @@ enum omp_mask2 OMP_CLAUSE_DOACROSS, /* OpenMP 5.2 */ OMP_CLAUSE_ASSUMPTIONS, /* OpenMP 5.1. */ OMP_CLAUSE_USES_ALLOCATORS, /* OpenMP 5.0 */ - OMP_CLAUSE_SELF, /* OpenACC 2.7 */ /* This must come last. */ OMP_MASK2_LAST }; @@ -1877,8 +1877,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, that should work. */ bool allow_derived = (openacc && ((mask & OMP_CLAUSE_ATTACH) - || (mask & OMP_CLAUSE_DETACH) - || (mask & OMP_CLAUSE_HOST_SELF))); + || (mask & OMP_CLAUSE_DETACH))); gcc_checking_assert (OMP_MASK1_LAST <= 64 && OMP_MASK2_LAST <= 64); *cp = NULL; @@ -2550,7 +2549,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, && gfc_match ("device ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], OMP_MAP_FORCE_TO, true, - allow_derived)) + /* allow_derived = */ true)) continue; if ((mask & OMP_CLAUSE_DEVICEPTR) && gfc_match ("deviceptr ( ") == MATCH_YES @@ -2725,11 +2724,11 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, c->assume->holds = el; continue; } - if ((mask & OMP_CLAUSE_HOST_SELF) + if ((mask & OMP_CLAUSE_HOST) && gfc_match ("host ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], OMP_MAP_FORCE_FROM, true, - allow_derived)) + /* allow_derived = */ true)) continue; break; case 'i': @@ -3521,10 +3520,10 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, gfc_current_locus = old_loc; } if ((mask & OMP_CLAUSE_SELF) + && !(mask & OMP_CLAUSE_HOST) /* OpenACC compute construct */ && (m = gfc_match_dupl_check (!c->self_expr, "self")) != MATCH_NO) { - gcc_assert (!(mask & OMP_CLAUSE_HOST_SELF)); if (m == MATCH_ERROR) goto error; m = gfc_match (" ( %e )", &c->self_expr); @@ -3541,11 +3540,12 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, } continue; } - if ((mask & OMP_CLAUSE_HOST_SELF) + if ((mask & OMP_CLAUSE_SELF) + && (mask & OMP_CLAUSE_HOST) /* OpenACC 'update' directive */ && gfc_match ("self ( ") == MATCH_YES && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP], OMP_MAP_FORCE_FROM, true, - allow_derived)) + /* allow_derived = */ true)) continue; if ((mask & OMP_CLAUSE_SEQ) && (m = gfc_match_dupl_check (!c->seq, "seq")) != MATCH_NO) @@ -3854,9 +3854,10 @@ error: | OMP_CLAUSE_CREATE | OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_DEVICE_RESIDENT \ | 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_IF_PRESENT) +#define OACC_UPDATE_CLAUSES \ + (omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_HOST \ + | OMP_CLAUSE_DEVICE | OMP_CLAUSE_WAIT | OMP_CLAUSE_IF_PRESENT \ + | OMP_CLAUSE_SELF) #define OACC_ENTER_DATA_CLAUSES \ (omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_WAIT \ | OMP_CLAUSE_COPYIN | OMP_CLAUSE_CREATE | OMP_CLAUSE_ATTACH) -- 2.34.1 ^ permalink raw reply [flat|nested] 5+ messages in thread
* Extend test suite coverage for OpenACC 'self' clause for compute constructs (was: [PATCH, OpenACC 2.7] Implement self clause for compute constructs) 2023-10-25 8:57 ` [PATCH, OpenACC 2.7] Implement self clause for compute constructs Thomas Schwinge 2023-10-25 9:10 ` Disentangle handling of OpenACC 'host', 'self' pragma tokens (was: [PATCH, OpenACC 2.7] Implement self clause for compute constructs) Thomas Schwinge @ 2023-10-25 9:29 ` Thomas Schwinge 2023-10-25 9:35 ` Handle OpenACC 'self' clause for compute constructs in OpenACC 'kernels' decomposition (was: Extend test suite coverage for OpenACC 'self' clause for compute constructs (was: [PATCH, OpenACC 2.7] Implement self clause for compute constructs)) Thomas Schwinge 2023-10-25 9:44 ` Minor fixes for OpenACC/Fortran 'self' clause for compute constructs (was: [PATCH, OpenACC 2.7] Implement self clause for compute constructs) Thomas Schwinge 2 siblings, 1 reply; 5+ messages in thread From: Thomas Schwinge @ 2023-10-25 9:29 UTC (permalink / raw) To: Chung-Lin Tang, gcc-patches, fortran; +Cc: Catherine Moore, Tobias Burnus [-- Attachment #1: Type: text/plain, Size: 1511 bytes --] Hi! On 2023-10-25T10:57:06+0200, I wrote: > With minor textual conflicts resolved, I've pushed this to master branch > in commit 3a3596389c2e539cb8fd5dc5784a4e2afe193a2a > "OpenACC 2.7: Implement self clause for compute constructs", see > attached. > > > I'll then apply/submit a number of follow-on commits. > From 3a3596389c2e539cb8fd5dc5784a4e2afe193a2a Mon Sep 17 00:00:00 2001 > From: Chung-Lin Tang <cltang@codesourcery.com> > Date: Tue, 13 Jun 2023 08:44:31 -0700 > Subject: [PATCH] OpenACC 2.7: Implement self clause for compute constructs > .../c-c++-common/goacc/self-clause-1.c | 22 + > .../c-c++-common/goacc/self-clause-2.c | 17 + > gcc/testsuite/gfortran.dg/goacc/self.f95 | 53 + > .../libgomp.oacc-c-c++-common/self-1.c | 962 ++++++++++++++++++ I found that insufficient, and added some more. Pushed to master branch commit 047841a68ebf5f991e842961f9e54f3c10b94f2c "Extend test suite coverage for OpenACC 'self' clause for compute constructs", see attached. This is mostly just adapting and cross-linking some existing 'if' clause test cases. (..., which turned up a problem when the 'self' clause is used with OpenACC 'kernels'.) Grüße Thomas ----------------- Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955 [-- Warning: decoded text below may be mangled, UTF-8 assumed --] [-- Attachment #2: 0001-Extend-test-suite-coverage-for-OpenACC-self-clause-f.patch --] [-- Type: text/x-diff, Size: 43846 bytes --] From 047841a68ebf5f991e842961f9e54f3c10b94f2c Mon Sep 17 00:00:00 2001 From: Thomas Schwinge <thomas@codesourcery.com> Date: Mon, 23 Oct 2023 14:53:29 +0200 Subject: [PATCH] Extend test suite coverage for OpenACC 'self' clause for compute constructs ... on top of what was provided in recent commit 3a3596389c2e539cb8fd5dc5784a4e2afe193a2a "OpenACC 2.7: Implement self clause for compute constructs". gcc/testsuite/ * c-c++-common/goacc/if-clause-2.c: Enhance. * c-c++-common/goacc/self-clause-1.c: Likewise. * c-c++-common/goacc/self-clause-2.c: Likewise. * gfortran.dg/goacc/if.f95: Likewise. * gfortran.dg/goacc/kernels-tree.f95: Likewise. * gfortran.dg/goacc/parallel-tree.f95: Likewise. * gfortran.dg/goacc/self.f95: Likewise. libgomp/ * testsuite/libgomp.oacc-c-c++-common/if-1.c: Enhance. * testsuite/libgomp.oacc-c-c++-common/self-1.c: Likewise. * testsuite/libgomp.oacc-fortran/if-1.f90: Likewise. * testsuite/libgomp.oacc-c-c++-common/if-self-1.c: New. * testsuite/libgomp.oacc-fortran/self-1.f90: Likewise. --- .../c-c++-common/goacc/if-clause-2.c | 2 + .../c-c++-common/goacc/self-clause-1.c | 6 + .../c-c++-common/goacc/self-clause-2.c | 20 + gcc/testsuite/gfortran.dg/goacc/if.f95 | 10 +- .../gfortran.dg/goacc/kernels-tree.f95 | 5 +- .../gfortran.dg/goacc/parallel-tree.f95 | 3 +- gcc/testsuite/gfortran.dg/goacc/self.f95 | 8 + .../libgomp.oacc-c-c++-common/if-1.c | 4 + .../libgomp.oacc-c-c++-common/if-self-1.c | 36 + .../libgomp.oacc-c-c++-common/self-1.c | 5 + .../testsuite/libgomp.oacc-fortran/if-1.f90 | 4 + .../testsuite/libgomp.oacc-fortran/self-1.f90 | 996 ++++++++++++++++++ 12 files changed, 1094 insertions(+), 5 deletions(-) create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/if-self-1.c create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/self-1.f90 diff --git a/gcc/testsuite/c-c++-common/goacc/if-clause-2.c b/gcc/testsuite/c-c++-common/goacc/if-clause-2.c index a48072509e1..71475521758 100644 --- a/gcc/testsuite/c-c++-common/goacc/if-clause-2.c +++ b/gcc/testsuite/c-c++-common/goacc/if-clause-2.c @@ -1,3 +1,5 @@ +/* See also 'self-clause-2.c'. */ + /* { dg-additional-options "-fdump-tree-gimple" } */ /* { dg-additional-options "--param=openacc-kernels=decompose" } { dg-additional-options "-fdump-tree-omp_oacc_kernels_decompose" } */ diff --git a/gcc/testsuite/c-c++-common/goacc/self-clause-1.c b/gcc/testsuite/c-c++-common/goacc/self-clause-1.c index fe892bea210..28de3dc0584 100644 --- a/gcc/testsuite/c-c++-common/goacc/self-clause-1.c +++ b/gcc/testsuite/c-c++-common/goacc/self-clause-1.c @@ -5,6 +5,8 @@ f (int b) { struct { int i; } *p; +#pragma acc parallel self(0) self(b) /* { dg-error "too many 'self' clauses" } */ + ; #pragma acc parallel self self(b) /* { dg-error "too many 'self' clauses" } */ ; #pragma acc parallel self(*p) @@ -12,6 +14,8 @@ f (int b) { dg-error {could not convert '\* p' from 'f\(int\)::<unnamed struct>' to 'bool'} {} { target c++ } .-2 } */ ; +#pragma acc kernels self(0) self(b) /* { dg-error "too many 'self' clauses" } */ + ; #pragma acc kernels self self(b) /* { dg-error "too many 'self' clauses" } */ ; #pragma acc kernels self(*p) @@ -19,6 +23,8 @@ f (int b) { dg-error {could not convert '\* p' from 'f\(int\)::<unnamed struct>' to 'bool'} {} { target c++ } .-2 } */ ; +#pragma acc serial self(0) self(b) /* { dg-error "too many 'self' clauses" } */ + ; #pragma acc serial self self(b) /* { dg-error "too many 'self' clauses" } */ ; #pragma acc serial self(*p) diff --git a/gcc/testsuite/c-c++-common/goacc/self-clause-2.c b/gcc/testsuite/c-c++-common/goacc/self-clause-2.c index d932ac9a4a6..769694baec9 100644 --- a/gcc/testsuite/c-c++-common/goacc/self-clause-2.c +++ b/gcc/testsuite/c-c++-common/goacc/self-clause-2.c @@ -1,3 +1,5 @@ +/* See also 'if-clause-2.c'. */ + /* { dg-additional-options "-fdump-tree-gimple" } */ void @@ -15,3 +17,21 @@ f (short c) /* { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_serial map\(tofrom:c \[len: [0-9]+\]\) self\(_[0-9]+\)$} 1 "gimple" } } */ ++c; } + +/* The same, but with implicit 'true' condition-argument. */ + +void +g (short d) +{ +#pragma acc parallel self copy(d) + /* { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_parallel map\(tofrom:d \[len: [0-9]+\]\) self\(1\)$} 1 "gimple" } } */ + ++d; + +#pragma acc kernels self copy(d) + /* { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_kernels map\(tofrom:d \[len: [0-9]+\]\) self\(1\)$} 1 "gimple" } } */ + ++d; + +#pragma acc serial self copy(d) + /* { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_serial map\(tofrom:d \[len: [0-9]+\]\) self\(1\)$} 1 "gimple" } } */ + ++d; +} diff --git a/gcc/testsuite/gfortran.dg/goacc/if.f95 b/gcc/testsuite/gfortran.dg/goacc/if.f95 index 56f3711f320..753ef8251c2 100644 --- a/gcc/testsuite/gfortran.dg/goacc/if.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/if.f95 @@ -1,3 +1,5 @@ +! See also 'self.f95'. + ! { dg-do compile } program test @@ -12,12 +14,14 @@ program test !$acc end parallel !$acc parallel if (1) ! { dg-error "scalar LOGICAL expression" } !$acc end parallel - !$acc kernels if (i) ! { dg-error "scalar LOGICAL expression" } - !$acc end kernels + !$acc kernels if ! { dg-error "Expected '\\(' after 'if'" } !$acc kernels if () ! { dg-error "Invalid character" } + !$acc kernels if (i) ! { dg-error "scalar LOGICAL expression" } + !$acc end kernels !$acc kernels if (1) ! { dg-error "scalar LOGICAL expression" } !$acc end kernels + !$acc data if ! { dg-error "Expected '\\(' after 'if'" } !$acc data if () ! { dg-error "Invalid character" } !$acc data if (i) ! { dg-error "scalar LOGICAL expression" } @@ -36,12 +40,14 @@ program test !$acc end parallel !$acc parallel if (i.gt.1) !$acc end parallel + !$acc kernels if (x) !$acc end kernels !$acc kernels if (.true.) !$acc end kernels !$acc kernels if (i.gt.1) !$acc end kernels + !$acc data if (x) !$acc end data !$acc data if (.true.) diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 index ceb07fbb9e9..1ba04a84e12 100644 --- a/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 @@ -12,6 +12,7 @@ program test logical :: l = .true. !$acc kernels if(l) async num_gangs(i) num_workers(i) vector_length(i) & + !$acc self & !$acc copy(i), copyin(j), copyout(k), create(m) & !$acc no_create(n) & !$acc present(o), pcopy(p), pcopyin(r), pcopyout(s), pcreate(t) & @@ -27,7 +28,7 @@ end program test ! { dg-final { scan-tree-dump-times "num_gangs" 1 "original" } } ! { dg-final { scan-tree-dump-times "num_workers" 1 "original" } } ! { dg-final { scan-tree-dump-times "vector_length" 1 "original" } } - +! { dg-final { scan-tree-dump-times "self\\(1\\)" 1 "original" } } ! { dg-final { scan-tree-dump-times "map\\(tofrom:i\\)" 1 "original" } } ! { dg-final { scan-tree-dump-times "map\\(to:j\\)" 1 "original" } } ! { dg-final { scan-tree-dump-times "map\\(from:k\\)" 1 "original" } } @@ -42,4 +43,4 @@ end program test ! { dg-final { scan-tree-dump-times "map\\(force_deviceptr:u\\)" 1 "original" } } ! { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_data_kernels if\((?:D\.|_)[0-9]+\)$} 1 "omp_oacc_kernels_decompose" } } -! { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_parallel_kernels_gang_single num_gangs\(1\) if\((?:D\.|_)[0-9]+\) async\(-1\)$} 1 "omp_oacc_kernels_decompose" } } +! { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_parallel_kernels_gang_single num_gangs\(1\) if\((?:D\.|_)[0-9]+\) self\(1\) async\(-1\)$} 1 "omp_oacc_kernels_decompose" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 b/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 index 6110d93b91e..0d4ec1133af 100644 --- a/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 @@ -14,6 +14,7 @@ program test logical :: l = .true. !$acc parallel if(l) async num_gangs(i) num_workers(i) vector_length(i) & + !$acc self & !$acc reduction(max:q), copy(i), copyin(j), copyout(k), create(m) & !$acc no_create(n) & !$acc present(o), pcopy(p), pcopyin(r), pcopyout(s), pcreate(t) & @@ -33,7 +34,7 @@ end program test ! { dg-final { scan-tree-dump-times "num_gangs" 1 "original" } } ! { dg-final { scan-tree-dump-times "num_workers" 1 "original" } } ! { dg-final { scan-tree-dump-times "vector_length" 1 "original" } } - +! { dg-final { scan-tree-dump-times "self\\(1\\)" 1 "original" } } ! { dg-final { scan-tree-dump-times "reduction\\(max:q\\)" 1 "original" } } ! { dg-final { scan-tree-dump-times "map\\(tofrom:i\\)" 1 "original" } } ! { dg-final { scan-tree-dump-times "map\\(to:j\\)" 1 "original" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/self.f95 b/gcc/testsuite/gfortran.dg/goacc/self.f95 index 4817f16be56..aa0f6fe88c5 100644 --- a/gcc/testsuite/gfortran.dg/goacc/self.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/self.f95 @@ -1,3 +1,5 @@ +! See also 'if.f95'. + ! { dg-do compile } program test @@ -29,6 +31,8 @@ program test !$acc kernels self (.false.) self (.false.) { dg-error "Duplicated 'self' clause" } !$acc serial self (.false.) self (.false.) { dg-error "Duplicated 'self' clause" } + !$acc parallel self + !$acc end parallel !$acc parallel self (x) !$acc end parallel !$acc parallel self (.true.) @@ -36,6 +40,8 @@ program test !$acc parallel self (i.gt.1) !$acc end parallel + !$acc kernels self + !$acc end kernels !$acc kernels self (x) !$acc end kernels !$acc kernels self (.true.) @@ -43,6 +49,8 @@ program test !$acc kernels self (i.gt.1) !$acc end kernels + !$acc serial self + !$acc end serial !$acc serial self (x) !$acc end serial !$acc serial self (.true.) diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c index 5398905f411..f04c02fdb89 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c @@ -1,3 +1,7 @@ +/* OpenACC 'if' clause. */ + +/* See also 'self-1.c'. */ + #include <openacc.h> #include <stdlib.h> #include <stdbool.h> diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/if-self-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/if-self-1.c new file mode 100644 index 00000000000..f77edda83e3 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/if-self-1.c @@ -0,0 +1,36 @@ +/* Test 'if' and 'self' clause appearing together. */ + +#include <openacc.h> + +static int test(float i, long double s) +{ + int ret; +#pragma acc serial copyout(ret) if(i) self(s) + /* { dg-bogus {using 'vector_length \(32\)', ignoring 1} {} { xfail openacc_nvidia_accel_selected } .-1 } */ + { + ret = acc_on_device(acc_device_host); + } + return ret; +} + +int main() +{ + if (!test(0, 0)) + __builtin_abort(); + + if (!test(0, 1)) + __builtin_abort(); + +#if ACC_MEM_SHARED + if (!test(1, 0)) + __builtin_abort(); +#else + if (test(1, 0)) + __builtin_abort(); +#endif + + if (!test(1, 1)) + __builtin_abort(); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/self-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/self-1.c index 752e16e8545..d17e27c8af3 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/self-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/self-1.c @@ -1,3 +1,8 @@ +/* OpenACC 'self' clause. */ + +/* This is 'if-1.c' with 'self(!cond)' instead of 'if(cond)' on compute + constructs. */ + #include <openacc.h> #include <stdlib.h> #include <stdbool.h> diff --git a/libgomp/testsuite/libgomp.oacc-fortran/if-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/if-1.f90 index e0cfd912d0f..5ba6509749e 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/if-1.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/if-1.f90 @@ -1,3 +1,7 @@ +! OpenACC 'if' clause. + +! See also 'self-1.f90'. + ! { dg-do run } ! { dg-additional-options "-cpp" } diff --git a/libgomp/testsuite/libgomp.oacc-fortran/self-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/self-1.f90 new file mode 100644 index 00000000000..b9ec9de08d9 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/self-1.f90 @@ -0,0 +1,996 @@ +! OpenACC 'self' clause. + +! This is 'if-1.f90' with 'self(!cond)' instead of 'if(cond)' on compute +! constructs. +! ..., which the exception of certain 'kernels' constructs. + +! { dg-do run } +! { dg-additional-options "-cpp" } + +! { dg-additional-options "--param=openacc-kernels=decompose" } + +! { dg-additional-options "-fopt-info-note-omp" } +! { dg-additional-options "-foffload=-fopt-info-note-omp" } + +! { dg-additional-options "--param=openacc-privatization=noisy" } +! { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } +! Prune a few: uninteresting, and potentially varying depending on GCC configuration (data types): +! { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } + +! It's only with Tcl 8.5 (released in 2007) that "the variable 'varName' +! passed to 'incr' may be unset, and in that case, it will be set to [...]", +! so to maintain compatibility with earlier Tcl releases, we manually +! initialize counter variables: +! { dg-line l_dummy[variable c_compute 0] } +! { dg-message "dummy" "" { target iN-VAl-Id } l_dummy } to avoid +! "WARNING: dg-line var l_dummy defined, but not used". */ + +program main + use openacc + implicit none + + integer, parameter :: N = 8 + integer, parameter :: one = 1 + integer, parameter :: zero = 0 + integer i, nn + real, allocatable :: a(:), b(:) + real exp, exp2 + + i = 0 + + allocate (a(N)) + allocate (b(N)) + + a(:) = 4.0 + + !$acc parallel copyin (a(1:N)) copyout (b(1:N)) self (1 /= 1) ! { dg-line l_compute[incr c_compute] } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + !TODO Unhandled 'CONST_DECL' instances for constant argument in 'acc_on_device' call. + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end parallel + +#if ACC_MEM_SHARED + exp = 5.0 +#else + exp = 4.0 +#endif + + do i = 1, N + if (b(i) .ne. exp) STOP 1 + end do + + a(:) = 16.0 + + !$acc parallel self (0 /= 1) ! { dg-line l_compute[incr c_compute] } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end parallel + + do i = 1, N + if (b(i) .ne. 17.0) STOP 2 + end do + + a(:) = 8.0 + + !$acc parallel copyin (a(1:N)) copyout (b(1:N)) self (one /= 1) ! { dg-line l_compute[incr c_compute] } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end parallel + +#if ACC_MEM_SHARED + exp = 9.0 +#else + exp = 8.0 +#endif + + do i = 1, N + if (b(i) .ne. exp) STOP 3 + end do + + a(:) = 22.0 + + !$acc parallel self (zero /= 1) ! { dg-line l_compute[incr c_compute] } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end parallel + + do i = 1, N + if (b(i) .ne. 23.0) STOP 4 + end do + + a(:) = 16.0 + + !$acc parallel copyin (a(1:N)) copyout (b(1:N)) self (.FALSE.) ! { dg-line l_compute[incr c_compute] } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end parallel + +#if ACC_MEM_SHARED + exp = 17.0; +#else + exp = 16.0; +#endif + + do i = 1, N + if (b(i) .ne. exp) STOP 5 + end do + + a(:) = 76.0 + + !$acc parallel self (.TRUE.) ! { dg-line l_compute[incr c_compute] } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end parallel + + do i = 1, N + if (b(i) .ne. 77.0) STOP 6 + end do + + a(:) = 22.0 + + nn = 1 + + !$acc parallel copyin (a(1:N)) copyout (b(1:N)) self (nn /= 1) ! { dg-line l_compute[incr c_compute] } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end parallel + +#if ACC_MEM_SHARED + exp = 23.0; +#else + exp = 22.0; +#endif + + do i = 1, N + if (b(i) .ne. exp) STOP 7 + end do + + a(:) = 18.0 + + nn = 0 + + !$acc parallel self (nn /= 1) ! { dg-line l_compute[incr c_compute] } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end parallel + + do i = 1, N + if (b(i) .ne. 19.0) STOP 8 + end do + + a(:) = 49.0 + + nn = 1 + + !$acc parallel copyin (a(1:N)) copyout (b(1:N)) self (.NOT. ((nn + nn) > 0)) ! { dg-line l_compute[incr c_compute] } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end parallel + +#if ACC_MEM_SHARED + exp = 50.0 +#else + exp = 49.0 +#endif + + do i = 1, N + if (b(i) .ne. exp) STOP 9 + end do + + a(:) = 38.0 + + nn = 0; + + !$acc parallel copyin (a(1:N)) copyout (b(1:N)) self (.NOT. ((nn + nn) > 0)) ! { dg-line l_compute[incr c_compute] } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end parallel + + do i = 1, N + if (b(i) .ne. 39.0) STOP 10 + end do + + a(:) = 91.0 + + !$acc parallel copyin (a(1:N)) copyout (b(1:N)) self (.NOT. (-2 > 0)) ! { dg-line l_compute[incr c_compute] } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end parallel + + do i = 1, N + if (b(i) .ne. 92.0) STOP 11 + end do + + a(:) = 43.0 + + !$acc parallel copyin (a(1:N)) copyout (b(1:N)) self (one /= 1) ! { dg-line l_compute[incr c_compute] } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end parallel + +#if ACC_MEM_SHARED + exp = 44.0 +#else + exp = 43.0 +#endif + + do i = 1, N + if (b(i) .ne. exp) STOP 12 + end do + + a(:) = 87.0 + + !$acc parallel self (one /= 0) ! { dg-line l_compute[incr c_compute] } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end parallel + + do i = 1, N + if (b(i) .ne. 88.0) STOP 13 + end do + + a(:) = 3.0 + b(:) = 9.0 + +#if ACC_MEM_SHARED + exp = 0.0 + exp2 = 0.0 +#else + call acc_copyin (a, sizeof (a)) + call acc_copyin (b, sizeof (b)) + exp = 3.0; + exp2 = 9.0; +#endif + + !$acc update device (a(1:N), b(1:N)) if (1 == 1) + + a(:) = 0.0 + b(:) = 0.0 + + !$acc update host (a(1:N), b(1:N)) if (1 == 1) + + do i = 1, N + if (a(i) .ne. exp) STOP 14 + if (b(i) .ne. exp2) STOP 15 + end do + + a(:) = 6.0 + b(:) = 12.0 + + !$acc update device (a(1:N), b(1:N)) if (0 == 1) + + a(:) = 0.0 + b(:) = 0.0 + + !$acc update host (a(1:N), b(1:N)) if (1 == 1) + + do i = 1, N + if (a(i) .ne. exp) STOP 16 + if (b(i) .ne. exp2) STOP 17 + end do + + a(:) = 26.0 + b(:) = 21.0 + + !$acc update device (a(1:N), b(1:N)) if (1 == 1) + + a(:) = 0.0 + b(:) = 0.0 + + !$acc update host (a(1:N), b(1:N)) if (0 == 1) + + do i = 1, N + if (a(i) .ne. 0.0) STOP 18 + if (b(i) .ne. 0.0) STOP 19 + end do + +#if !ACC_MEM_SHARED + call acc_copyout (a, sizeof (a)) + call acc_copyout (b, sizeof (b)) +#endif + + a(:) = 4.0 + b(:) = 0.0 + + !$acc data copyin (a(1:N)) copyout (b(1:N)) if (1 == 1) + ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 } + + !$acc parallel present (a(1:N)) + do i = 1, N + b(i) = a(i) + end do + !$acc end parallel + !$acc end data + + do i = 1, N + if (b(i) .ne. 4.0) STOP 20 + end do + + a(:) = 8.0 + b(:) = 1.0 + + !$acc data copyin (a(1:N)) copyout (b(1:N)) if (0 == 1) + ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-1 } + +#if !ACC_MEM_SHARED + if (acc_is_present (a) .eqv. .TRUE.) STOP 21 + if (acc_is_present (b) .eqv. .TRUE.) STOP 22 +#endif + + !$acc end data + + a(:) = 18.0 + b(:) = 21.0 + + !$acc data copyin (a(1:N)) if (1 == 1) + ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 } + ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-2 } + +#if !ACC_MEM_SHARED + if (acc_is_present (a) .eqv. .FALSE.) STOP 23 +#endif + + !$acc data copyout (b(1:N)) if (0 == 1) + ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 } + ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-2 } +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .TRUE.) STOP 24 +#endif + !$acc data copyout (b(1:N)) if (1 == 1) + ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 } + + !$acc parallel present (a(1:N)) present (b(1:N)) + do i = 1, N + b(i) = a(i) + end do + !$acc end parallel + + !$acc end data + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .TRUE.) STOP 25 +#endif + !$acc end data + !$acc end data + + do i = 1, N + if (b(1) .ne. 18.0) STOP 26 + end do + + !$acc enter data copyin (b(1:N)) if (0 == 1) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .TRUE.) STOP 27 +#endif + + !$acc exit data delete (b(1:N)) if (0 == 1) + + !$acc enter data copyin (b(1:N)) if (1 == 1) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .FALSE.) STOP 28 +#endif + + !$acc exit data delete (b(1:N)) if (1 == 1) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .TRUE.) STOP 29 +#endif + + !$acc enter data copyin (b(1:N)) if (zero == 1) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .TRUE.) STOP 30 +#endif + + !$acc exit data delete (b(1:N)) if (zero == 1) + + !$acc enter data copyin (b(1:N)) if (one == 1) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .FALSE.) STOP 31 +#endif + + !$acc exit data delete (b(1:N)) if (one == 1) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .TRUE.) STOP 32 +#endif + + !$acc enter data copyin (b(1:N)) if (one == 0) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .TRUE.) STOP 33 +#endif + + !$acc exit data delete (b(1:N)) if (one == 0) + + !$acc enter data copyin (b(1:N)) if (one == 1) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .FALSE.) STOP 34 +#endif + + !$acc exit data delete (b(1:N)) if (one == 1) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .TRUE.) STOP 35 +#endif + + a(:) = 4.0 + + !$acc kernels copyin (a(1:N)) copyout (b(1:N)) self (1 /= 1) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end kernels + +#if ACC_MEM_SHARED + exp = 5.0 +#else + exp = 4.0 +#endif + + do i = 1, N + if (b(i) .ne. exp) STOP 36 + end do + + a(:) = 16.0 + + !$acc kernels if (0 == 1) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end kernels + + do i = 1, N + if (b(i) .ne. 17.0) STOP 37 + end do + + a(:) = 8.0 + + !$acc kernels copyin (a(1:N)) copyout (b(1:N)) self (one /= 1) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end kernels + +#if ACC_MEM_SHARED + exp = 9.0 +#else + exp = 8.0 +#endif + + do i = 1, N + if (b(i) .ne. exp) STOP 38 + end do + + a(:) = 22.0 + + !$acc kernels if (zero == 1) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end kernels + + do i = 1, N + if (b(i) .ne. 23.0) STOP 39 + end do + + a(:) = 16.0 + + !$acc kernels copyin (a(1:N)) copyout (b(1:N)) self (.FALSE.) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end kernels + +#if ACC_MEM_SHARED + exp = 17.0; +#else + exp = 16.0; +#endif + + do i = 1, N + if (b(i) .ne. exp) STOP 40 + end do + + a(:) = 76.0 + + !$acc kernels if (.FALSE.) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end kernels + + do i = 1, N + if (b(i) .ne. 77.0) STOP 41 + end do + + a(:) = 22.0 + + nn = 1 + + !$acc kernels copyin (a(1:N)) copyout (b(1:N)) self (nn /= 1) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end kernels + +#if ACC_MEM_SHARED + exp = 23.0; +#else + exp = 22.0; +#endif + + do i = 1, N + if (b(i) .ne. exp) STOP 42 + end do + + a(:) = 18.0 + + nn = 0 + + !$acc kernels if (nn == 1) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end kernels + + do i = 1, N + if (b(i) .ne. 19.0) STOP 43 + end do + + a(:) = 49.0 + + nn = 1 + + !$acc kernels copyin (a(1:N)) copyout (b(1:N)) self (.NOT. ((nn + nn) > 0)) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end kernels + +#if ACC_MEM_SHARED + exp = 50.0 +#else + exp = 49.0 +#endif + + do i = 1, N + if (b(i) .ne. exp) STOP 44 + end do + + a(:) = 38.0 + + nn = 0; + + !$acc kernels copyin (a(1:N)) copyout (b(1:N)) if ((nn + nn) > 0) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end kernels + + do i = 1, N + if (b(i) .ne. 39.0) STOP 45 + end do + + a(:) = 91.0 + + !$acc kernels copyin (a(1:N)) copyout (b(1:N)) if (-2 > 0) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end kernels + + do i = 1, N + if (b(i) .ne. 92.0) STOP 46 + end do + + a(:) = 43.0 + + !$acc kernels copyin (a(1:N)) copyout (b(1:N)) self (one /= 1) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end kernels + +#if ACC_MEM_SHARED + exp = 44.0 +#else + exp = 43.0 +#endif + + do i = 1, N + if (b(i) .ne. exp) STOP 47 + end do + + a(:) = 87.0 + + !$acc kernels if (one == 0) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end kernels + + do i = 1, N + if (b(i) .ne. 88.0) STOP 48 + end do + + a(:) = 3.0 + b(:) = 9.0 + +#if ACC_MEM_SHARED + exp = 0.0 + exp2 = 0.0 +#else + call acc_copyin (a, sizeof (a)) + call acc_copyin (b, sizeof (b)) + exp = 3.0; + exp2 = 9.0; +#endif + + !$acc update device (a(1:N), b(1:N)) if (1 == 1) + + a(:) = 0.0 + b(:) = 0.0 + + !$acc update host (a(1:N), b(1:N)) if (1 == 1) + + do i = 1, N + if (a(i) .ne. exp) STOP 49 + if (b(i) .ne. exp2) STOP 50 + end do + + a(:) = 6.0 + b(:) = 12.0 + + !$acc update device (a(1:N), b(1:N)) if (0 == 1) + + a(:) = 0.0 + b(:) = 0.0 + + !$acc update host (a(1:N), b(1:N)) if (1 == 1) + + do i = 1, N + if (a(i) .ne. exp) STOP 51 + if (b(i) .ne. exp2) STOP 52 + end do + + a(:) = 26.0 + b(:) = 21.0 + + !$acc update device (a(1:N), b(1:N)) if (1 == 1) + + a(:) = 0.0 + b(:) = 0.0 + + !$acc update host (a(1:N), b(1:N)) if (0 == 1) + + do i = 1, N + if (a(i) .ne. 0.0) STOP 53 + if (b(i) .ne. 0.0) STOP 54 + end do + +#if !ACC_MEM_SHARED + call acc_copyout (a, sizeof (a)) + call acc_copyout (b, sizeof (b)) +#endif + + a(:) = 4.0 + b(:) = 0.0 + + !$acc data copyin (a(1:N)) copyout (b(1:N)) if (1 == 1) + ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 } + + !$acc kernels present (a(1:N)) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } + do i = 1, N + b(i) = a(i) + end do + !$acc end kernels + !$acc end data + + do i = 1, N + if (b(i) .ne. 4.0) STOP 55 + end do + + a(:) = 8.0 + b(:) = 1.0 + + !$acc data copyin (a(1:N)) copyout (b(1:N)) if (0 == 1) + ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-1 } + +#if !ACC_MEM_SHARED + if (acc_is_present (a) .eqv. .TRUE.) STOP 56 + if (acc_is_present (b) .eqv. .TRUE.) STOP 57 +#endif + + !$acc end data + + a(:) = 18.0 + b(:) = 21.0 + + !$acc data copyin (a(1:N)) if (1 == 1) + ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 } + ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-2 } + +#if !ACC_MEM_SHARED + if (acc_is_present (a) .eqv. .FALSE.) STOP 58 +#endif + + !$acc data copyout (b(1:N)) if (0 == 1) + ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 } + ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-2 } +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .TRUE.) STOP 59 +#endif + !$acc data copyout (b(1:N)) if (1 == 1) + ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 } + + !$acc kernels present (a(1:N)) present (b(1:N)) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } + do i = 1, N + b(i) = a(i) + end do + !$acc end kernels + + !$acc end data + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .TRUE.) STOP 60 +#endif + !$acc end data + !$acc end data + + do i = 1, N + if (b(1) .ne. 18.0) STOP 61 + end do + + !$acc enter data copyin (b(1:N)) if (0 == 1) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .TRUE.) STOP 62 +#endif + + !$acc exit data delete (b(1:N)) if (0 == 1) + + !$acc enter data copyin (b(1:N)) if (1 == 1) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .FALSE.) STOP 63 +#endif + + !$acc exit data delete (b(1:N)) if (1 == 1) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .TRUE.) STOP 64 +#endif + + !$acc enter data copyin (b(1:N)) if (zero == 1) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .TRUE.) STOP 65 +#endif + + !$acc exit data delete (b(1:N)) if (zero == 1) + + !$acc enter data copyin (b(1:N)) if (one == 1) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .FALSE.) STOP 66 +#endif + + !$acc exit data delete (b(1:N)) if (one == 1) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .TRUE.) STOP 67 +#endif + + !$acc enter data copyin (b(1:N)) if (one == 0) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .TRUE.) STOP 68 +#endif + + !$acc exit data delete (b(1:N)) if (one == 0) + + !$acc enter data copyin (b(1:N)) if (one == 1) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .FALSE.) STOP 69 +#endif + + !$acc exit data delete (b(1:N)) if (one == 1) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .TRUE.) STOP 70 +#endif + +end program main -- 2.34.1 ^ permalink raw reply [flat|nested] 5+ messages in thread
* Handle OpenACC 'self' clause for compute constructs in OpenACC 'kernels' decomposition (was: Extend test suite coverage for OpenACC 'self' clause for compute constructs (was: [PATCH, OpenACC 2.7] Implement self clause for compute constructs)) 2023-10-25 9:29 ` Extend test suite coverage for OpenACC 'self' clause for compute constructs " Thomas Schwinge @ 2023-10-25 9:35 ` Thomas Schwinge 0 siblings, 0 replies; 5+ messages in thread From: Thomas Schwinge @ 2023-10-25 9:35 UTC (permalink / raw) To: Chung-Lin Tang, gcc-patches, fortran; +Cc: Catherine Moore, Tobias Burnus [-- Attachment #1: Type: text/plain, Size: 2081 bytes --] Hi! On 2023-10-25T11:29:52+0200, I wrote: > On 2023-10-25T10:57:06+0200, I wrote: >> With minor textual conflicts resolved, I've pushed this to master branch >> in commit 3a3596389c2e539cb8fd5dc5784a4e2afe193a2a >> "OpenACC 2.7: Implement self clause for compute constructs", see >> attached. >> >> >> I'll then apply/submit a number of follow-on commits. > >> From 3a3596389c2e539cb8fd5dc5784a4e2afe193a2a Mon Sep 17 00:00:00 2001 >> From: Chung-Lin Tang <cltang@codesourcery.com> >> Date: Tue, 13 Jun 2023 08:44:31 -0700 >> Subject: [PATCH] OpenACC 2.7: Implement self clause for compute constructs > >> .../c-c++-common/goacc/self-clause-1.c | 22 + >> .../c-c++-common/goacc/self-clause-2.c | 17 + >> gcc/testsuite/gfortran.dg/goacc/self.f95 | 53 + > >> .../libgomp.oacc-c-c++-common/self-1.c | 962 ++++++++++++++++++ > > I found that insufficient, and added some more. Pushed to > master branch commit 047841a68ebf5f991e842961f9e54f3c10b94f2c > "Extend test suite coverage for OpenACC 'self' clause for compute constructs", > see attached. This is mostly just adapting and cross-linking some > existing 'if' clause test cases. (..., which turned up a problem when > the 'self' clause is used with OpenACC 'kernels'.) > --- /dev/null > +++ b/libgomp/testsuite/libgomp.oacc-fortran/self-1.f90 > @@ -0,0 +1,996 @@ > +! OpenACC 'self' clause. > + > +! This is 'if-1.f90' with 'self(!cond)' instead of 'if(cond)' on compute > +! constructs. > +! ..., which the exception of certain 'kernels' constructs. ..., which I've then fixed up per master branch commit 7b2ae64b68132c1c643cb34d58cd5eab6f9de652 "Handle OpenACC 'self' clause for compute constructs in OpenACC 'kernels' decomposition", see attached. Grüße Thomas ----------------- Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955 [-- Warning: decoded text below may be mangled, UTF-8 assumed --] [-- Attachment #2: 0001-Handle-OpenACC-self-clause-for-compute-constructs-in.patch --] [-- Type: text/x-diff, Size: 12143 bytes --] From 7b2ae64b68132c1c643cb34d58cd5eab6f9de652 Mon Sep 17 00:00:00 2001 From: Thomas Schwinge <thomas@codesourcery.com> Date: Mon, 23 Oct 2023 15:28:30 +0200 Subject: [PATCH] Handle OpenACC 'self' clause for compute constructs in OpenACC 'kernels' decomposition ... to fix up recent commit 3a3596389c2e539cb8fd5dc5784a4e2afe193a2a "OpenACC 2.7: Implement self clause for compute constructs" for that case. gcc/ * omp-oacc-kernels-decompose.cc (omp_oacc_kernels_decompose_1): Handle 'OMP_CLAUSE_SELF' like 'OMP_CLAUSE_IF'. * omp-expand.cc (expand_omp_target): Handle 'OMP_CLAUSE_SELF' for 'GF_OMP_TARGET_KIND_OACC_DATA_KERNELS'. gcc/testsuite/ * c-c++-common/goacc/self-clause-2.c: Verify '--param=openacc-kernels=decompose'. * gfortran.dg/goacc/kernels-tree.f95: Adjust. libgomp/ * oacc-parallel.c (GOACC_data_start): Handle 'GOACC_FLAG_LOCAL_DEVICE'. (GOACC_parallel_keyed): Simplify accordingly. * testsuite/libgomp.oacc-fortran/self-1.f90: Adjust. --- gcc/omp-expand.cc | 14 ++++++++++++-- gcc/omp-oacc-kernels-decompose.cc | 15 ++++++++------- .../c-c++-common/goacc/self-clause-2.c | 6 ++++++ .../gfortran.dg/goacc/kernels-tree.f95 | 2 +- libgomp/oacc-parallel.c | 17 +++++------------ .../testsuite/libgomp.oacc-fortran/self-1.f90 | 15 +++++++-------- 6 files changed, 39 insertions(+), 30 deletions(-) diff --git a/gcc/omp-expand.cc b/gcc/omp-expand.cc index 8576b938102..5c6a7f2e381 100644 --- a/gcc/omp-expand.cc +++ b/gcc/omp-expand.cc @@ -10334,9 +10334,19 @@ expand_omp_target (struct omp_region *region) if ((c = omp_find_clause (clauses, OMP_CLAUSE_SELF)) != NULL_TREE) { - gcc_assert (is_gimple_omp_oacc (entry_stmt) && offloaded); + gcc_assert ((is_gimple_omp_oacc (entry_stmt) && offloaded) + || (gimple_omp_target_kind (entry_stmt) + == GF_OMP_TARGET_KIND_OACC_DATA_KERNELS)); - edge e = split_block_after_labels (new_bb); + edge e; + if (offloaded) + e = split_block_after_labels (new_bb); + else + { + gsi = gsi_last_nondebug_bb (new_bb); + gsi_prev (&gsi); + e = split_block (new_bb, gsi_stmt (gsi)); + } basic_block cond_bb = e->src; new_bb = e->dest; remove_edge (e); diff --git a/gcc/omp-oacc-kernels-decompose.cc b/gcc/omp-oacc-kernels-decompose.cc index ffc0a8f813e..dfbb34935d0 100644 --- a/gcc/omp-oacc-kernels-decompose.cc +++ b/gcc/omp-oacc-kernels-decompose.cc @@ -1519,17 +1519,18 @@ omp_oacc_kernels_decompose_1 (gimple *kernels_stmt) break; } } - else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IF) + else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_IF + || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_SELF) { - /* If there is an 'if' clause, it must be duplicated to the - enclosing data region. Temporarily remove the if clause's - chain to avoid copying it. */ + /* If there is an 'if' or 'self' clause, it must be duplicated to the + enclosing data region. Temporarily remove its chain to avoid + copying it. */ tree saved_chain = OMP_CLAUSE_CHAIN (c); OMP_CLAUSE_CHAIN (c) = NULL; - tree new_if_clause = unshare_expr (c); + tree new_clause = unshare_expr (c); OMP_CLAUSE_CHAIN (c) = saved_chain; - OMP_CLAUSE_CHAIN (new_if_clause) = data_clauses; - data_clauses = new_if_clause; + OMP_CLAUSE_CHAIN (new_clause) = data_clauses; + data_clauses = new_clause; } } /* Restore the original order of the clauses. */ diff --git a/gcc/testsuite/c-c++-common/goacc/self-clause-2.c b/gcc/testsuite/c-c++-common/goacc/self-clause-2.c index 769694baec9..3ac29a03bc4 100644 --- a/gcc/testsuite/c-c++-common/goacc/self-clause-2.c +++ b/gcc/testsuite/c-c++-common/goacc/self-clause-2.c @@ -1,6 +1,8 @@ /* See also 'if-clause-2.c'. */ /* { dg-additional-options "-fdump-tree-gimple" } */ +/* { dg-additional-options "--param=openacc-kernels=decompose" } + { dg-additional-options "-fdump-tree-omp_oacc_kernels_decompose" } */ void f (short c) @@ -11,6 +13,8 @@ f (short c) #pragma acc kernels self(c) copy(c) /* { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_kernels map\(tofrom:c \[len: [0-9]+\]\) self\(_[0-9]+\)$} 1 "gimple" } } */ + /* { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_data_kernels map\(tofrom:c \[len: [0-9]+\]\) self\(_[0-9]+\)$} 1 "omp_oacc_kernels_decompose" } } + { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_parallel_kernels_gang_single async\(-1\) num_gangs\(1\) map\(force_present:c \[len: [0-9]+\]\) self\(_[0-9]+\)$} 1 "omp_oacc_kernels_decompose" } } */ ++c; #pragma acc serial self(c) copy(c) @@ -29,6 +33,8 @@ g (short d) #pragma acc kernels self copy(d) /* { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_kernels map\(tofrom:d \[len: [0-9]+\]\) self\(1\)$} 1 "gimple" } } */ + /* { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_data_kernels map\(tofrom:d \[len: [0-9]+\]\) self\(1\)$} 1 "omp_oacc_kernels_decompose" } } + { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_parallel_kernels_gang_single async\(-1\) num_gangs\(1\) map\(force_present:d \[len: [0-9]+\]\) self\(1+\)$} 1 "omp_oacc_kernels_decompose" } } */ ++d; #pragma acc serial self copy(d) diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 index 1ba04a84e12..2ee578f7f32 100644 --- a/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 @@ -42,5 +42,5 @@ end program test ! { dg-final { scan-tree-dump-times "map\\(force_deviceptr:u\\)" 1 "original" } } -! { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_data_kernels if\((?:D\.|_)[0-9]+\)$} 1 "omp_oacc_kernels_decompose" } } +! { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_data_kernels if\((?:D\.|_)[0-9]+\) self\(1\)$} 1 "omp_oacc_kernels_decompose" } } ! { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_parallel_kernels_gang_single num_gangs\(1\) if\((?:D\.|_)[0-9]+\) self\(1\) async\(-1\)$} 1 "omp_oacc_kernels_decompose" } } diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c index cf37a1bdd7d..16cf3948e2d 100644 --- a/libgomp/oacc-parallel.c +++ b/libgomp/oacc-parallel.c @@ -184,19 +184,11 @@ GOACC_parallel_keyed (int flags_m, void (*fn) (void *), /* Host fallback if "if" clause is false or if the current device is set to the host. */ - if (flags & GOACC_FLAG_HOST_FALLBACK) - { - prof_info.device_type = acc_device_host; - api_info.device_type = prof_info.device_type; - goacc_save_and_set_bind (acc_device_host); - fn (hostaddrs); - goacc_restore_bind (); - goto out_prof; - } - else if (flags & GOACC_FLAG_LOCAL_DEVICE) - { + if ((flags & GOACC_FLAG_HOST_FALLBACK) /* TODO: a proper pthreads based "multi-core CPU" local device implementation. Currently, this is still the same as host-fallback. */ + || (flags & GOACC_FLAG_LOCAL_DEVICE)) + { prof_info.device_type = acc_device_host; api_info.device_type = prof_info.device_type; goacc_save_and_set_bind (acc_device_host); @@ -457,7 +449,8 @@ GOACC_data_start (int flags_m, size_t mapnum, /* Host fallback or 'do nothing'. */ if ((acc_dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) - || (flags & GOACC_FLAG_HOST_FALLBACK)) + || (flags & GOACC_FLAG_HOST_FALLBACK) + || (flags & GOACC_FLAG_LOCAL_DEVICE)) { prof_info.device_type = acc_device_host; api_info.device_type = prof_info.device_type; diff --git a/libgomp/testsuite/libgomp.oacc-fortran/self-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/self-1.f90 index b9ec9de08d9..6c1233d6cf5 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/self-1.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/self-1.f90 @@ -2,7 +2,6 @@ ! This is 'if-1.f90' with 'self(!cond)' instead of 'if(cond)' on compute ! constructs. -! ..., which the exception of certain 'kernels' constructs. ! { dg-do run } ! { dg-additional-options "-cpp" } @@ -523,7 +522,7 @@ program main a(:) = 16.0 - !$acc kernels if (0 == 1) ! { dg-line l_compute[incr c_compute] } + !$acc kernels self (0 /= 1) ! { dg-line l_compute[incr c_compute] } ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } @@ -569,7 +568,7 @@ program main a(:) = 22.0 - !$acc kernels if (zero == 1) ! { dg-line l_compute[incr c_compute] } + !$acc kernels self (zero /= 1) ! { dg-line l_compute[incr c_compute] } ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } @@ -615,7 +614,7 @@ program main a(:) = 76.0 - !$acc kernels if (.FALSE.) ! { dg-line l_compute[incr c_compute] } + !$acc kernels self (.TRUE.) ! { dg-line l_compute[incr c_compute] } ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } @@ -665,7 +664,7 @@ program main nn = 0 - !$acc kernels if (nn == 1) ! { dg-line l_compute[incr c_compute] } + !$acc kernels self (nn /= 1) ! { dg-line l_compute[incr c_compute] } ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } @@ -715,7 +714,7 @@ program main nn = 0; - !$acc kernels copyin (a(1:N)) copyout (b(1:N)) if ((nn + nn) > 0) ! { dg-line l_compute[incr c_compute] } + !$acc kernels copyin (a(1:N)) copyout (b(1:N)) self (.NOT. ((nn + nn) > 0)) ! { dg-line l_compute[incr c_compute] } ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } @@ -735,7 +734,7 @@ program main a(:) = 91.0 - !$acc kernels copyin (a(1:N)) copyout (b(1:N)) if (-2 > 0) ! { dg-line l_compute[incr c_compute] } + !$acc kernels copyin (a(1:N)) copyout (b(1:N)) self (.NOT. (-2 > 0)) ! { dg-line l_compute[incr c_compute] } ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } @@ -781,7 +780,7 @@ program main a(:) = 87.0 - !$acc kernels if (one == 0) ! { dg-line l_compute[incr c_compute] } + !$acc kernels self (one /= 0) ! { dg-line l_compute[incr c_compute] } ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } -- 2.34.1 ^ permalink raw reply [flat|nested] 5+ messages in thread
* Minor fixes for OpenACC/Fortran 'self' clause for compute constructs (was: [PATCH, OpenACC 2.7] Implement self clause for compute constructs) 2023-10-25 8:57 ` [PATCH, OpenACC 2.7] Implement self clause for compute constructs Thomas Schwinge 2023-10-25 9:10 ` Disentangle handling of OpenACC 'host', 'self' pragma tokens (was: [PATCH, OpenACC 2.7] Implement self clause for compute constructs) Thomas Schwinge 2023-10-25 9:29 ` Extend test suite coverage for OpenACC 'self' clause for compute constructs " Thomas Schwinge @ 2023-10-25 9:44 ` Thomas Schwinge 2 siblings, 0 replies; 5+ messages in thread From: Thomas Schwinge @ 2023-10-25 9:44 UTC (permalink / raw) To: Chung-Lin Tang, gcc-patches, fortran, Tobias Burnus; +Cc: Catherine Moore [-- Attachment #1: Type: text/plain, Size: 2330 bytes --] Hi! On 2023-10-25T10:57:06+0200, I wrote: > With minor textual conflicts resolved, I've pushed this to master branch > in commit 3a3596389c2e539cb8fd5dc5784a4e2afe193a2a > "OpenACC 2.7: Implement self clause for compute constructs", see > attached. > > > I'll then apply/submit a number of follow-on commits. Regarding the Fortran front end changes: > From 3a3596389c2e539cb8fd5dc5784a4e2afe193a2a Mon Sep 17 00:00:00 2001 > From: Chung-Lin Tang <cltang@codesourcery.com> > Date: Tue, 13 Jun 2023 08:44:31 -0700 > Subject: [PATCH] OpenACC 2.7: Implement self clause for compute constructs > --- a/gcc/fortran/gfortran.h > +++ b/gcc/fortran/gfortran.h > @@ -1546,6 +1546,7 @@ typedef struct gfc_omp_clauses > gfc_omp_namelist *lists[OMP_LIST_NUM]; > struct gfc_expr *if_expr; > struct gfc_expr *if_exprs[OMP_IF_LAST]; > + struct gfc_expr *self_expr; > struct gfc_expr *final_expr; > struct gfc_expr *num_threads; > struct gfc_expr *chunk_size; ..., this needs to be handled in a few more places, I think... > --- a/gcc/fortran/trans-openmp.cc > +++ b/gcc/fortran/trans-openmp.cc > @@ -6615,6 +6631,8 @@ gfc_split_omp_clauses (gfc_code *code, > /* And this is copied to all. */ > clausesa[GFC_OMP_SPLIT_TARGET].if_expr > = code->ext.omp_clauses->if_expr; > + clausesa[GFC_OMP_SPLIT_TARGET].self_expr > + = code->ext.omp_clauses->self_expr; > clausesa[GFC_OMP_SPLIT_TARGET].nowait > = code->ext.omp_clauses->nowait; > } ..., but this change isn't necessary: that function is for OpenMP only, and generally doesn't (have to) care about OpenACC-only clauses. OK to push the attached "Minor fixes for OpenACC/Fortran 'self' clause for compute constructs", or is anything more needed? Also, I've filed <https://gcc.gnu.org/PR111938> "Missing OpenACC/Fortran handling in 'gcc/fortran/frontend-passes.c'", which applies generally, not just to the OpenACC 'self' clause on compute constructs. Grüße Thomas ----------------- Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955 [-- Warning: decoded text below may be mangled, UTF-8 assumed --] [-- Attachment #2: 0001-Minor-fixes-for-OpenACC-Fortran-self-clause-for-comp.patch --] [-- Type: text/x-diff, Size: 2313 bytes --] From 943de6d5f1498aabfc343bf5e9dd6c2b63fc55ed Mon Sep 17 00:00:00 2001 From: Thomas Schwinge <thomas@codesourcery.com> Date: Fri, 20 Oct 2023 15:49:35 +0200 Subject: [PATCH] Minor fixes for OpenACC/Fortran 'self' clause for compute constructs ... to fix up recent commit 3a3596389c2e539cb8fd5dc5784a4e2afe193a2a "OpenACC 2.7: Implement self clause for compute constructs". gcc/fortran/ * dump-parse-tree.cc (show_omp_clauses): Handle 'self_expr'. * openmp.cc (gfc_free_omp_clauses): Likewise. * trans-openmp.cc (gfc_split_omp_clauses): Don't handle 'self_expr'. --- gcc/fortran/dump-parse-tree.cc | 6 ++++++ gcc/fortran/openmp.cc | 1 + gcc/fortran/trans-openmp.cc | 2 -- 3 files changed, 7 insertions(+), 2 deletions(-) diff --git a/gcc/fortran/dump-parse-tree.cc b/gcc/fortran/dump-parse-tree.cc index cc4846e5d74..26391df46e6 100644 --- a/gcc/fortran/dump-parse-tree.cc +++ b/gcc/fortran/dump-parse-tree.cc @@ -1614,6 +1614,12 @@ show_omp_clauses (gfc_omp_clauses *omp_clauses) show_expr (omp_clauses->if_exprs[i]); fputc (')', dumpfile); } + if (omp_clauses->self_expr) + { + fputs (" SELF(", dumpfile); + show_expr (omp_clauses->self_expr); + fputc (')', dumpfile); + } if (omp_clauses->final_expr) { fputs (" FINAL(", dumpfile); diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc index 2e2e23d567b..5e3cd0570bb 100644 --- a/gcc/fortran/openmp.cc +++ b/gcc/fortran/openmp.cc @@ -163,6 +163,7 @@ gfc_free_omp_clauses (gfc_omp_clauses *c) gfc_free_expr (c->if_expr); for (i = 0; i < OMP_IF_LAST; i++) gfc_free_expr (c->if_exprs[i]); + gfc_free_expr (c->self_expr); gfc_free_expr (c->final_expr); gfc_free_expr (c->num_threads); gfc_free_expr (c->chunk_size); diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc index 82bbc41b388..00782ee1716 100644 --- a/gcc/fortran/trans-openmp.cc +++ b/gcc/fortran/trans-openmp.cc @@ -6631,8 +6631,6 @@ gfc_split_omp_clauses (gfc_code *code, /* And this is copied to all. */ clausesa[GFC_OMP_SPLIT_TARGET].if_expr = code->ext.omp_clauses->if_expr; - clausesa[GFC_OMP_SPLIT_TARGET].self_expr - = code->ext.omp_clauses->self_expr; clausesa[GFC_OMP_SPLIT_TARGET].nowait = code->ext.omp_clauses->nowait; } -- 2.34.1 ^ permalink raw reply [flat|nested] 5+ messages in thread
end of thread, other threads:[~2023-10-25 9:44 UTC | newest]
Thread overview: 5+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
[not found] <e1a5d579-f28c-9417-44ed-ba6248850398@siemens.com>
2023-10-25 8:57 ` [PATCH, OpenACC 2.7] Implement self clause for compute constructs Thomas Schwinge
2023-10-25 9:10 ` Disentangle handling of OpenACC 'host', 'self' pragma tokens (was: [PATCH, OpenACC 2.7] Implement self clause for compute constructs) Thomas Schwinge
2023-10-25 9:29 ` Extend test suite coverage for OpenACC 'self' clause for compute constructs " Thomas Schwinge
2023-10-25 9:35 ` Handle OpenACC 'self' clause for compute constructs in OpenACC 'kernels' decomposition (was: Extend test suite coverage for OpenACC 'self' clause for compute constructs (was: [PATCH, OpenACC 2.7] Implement self clause for compute constructs)) Thomas Schwinge
2023-10-25 9:44 ` Minor fixes for OpenACC/Fortran 'self' clause for compute constructs (was: [PATCH, OpenACC 2.7] Implement self clause for compute constructs) Thomas Schwinge
This is a public inbox, see mirroring instructions for how to clone and mirror all data and code used for this inbox; as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).