From patchwork Thu Nov 7 09:52:13 2019 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Frederik Harwath X-Patchwork-Id: 1191027 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=209.132.180.131; helo=sourceware.org; envelope-from=gcc-patches-return-512683-incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b="RT5cvJpT"; dkim-atps=neutral Received: from sourceware.org (server1.sourceware.org [209.132.180.131]) (using TLSv1.2 with cipher ECDHE-RSA-AES256-GCM-SHA384 (256/256 bits)) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 477zHQ4zztz9sPk for ; Thu, 7 Nov 2019 20:53:20 +1100 (AEDT) DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:subject:date:message-id:mime-version:content-type; q=dns; s= default; b=ERyroLcPV1KPpOTclSmyD7wu1SIn+Yemd6oIkWwCJSgPR+QFptqlY aqflMlLF2jHn3fJvtnJtF4WR7TpLoS5pFU+PWr2xJNflPIqlTEoLxNa0c5dsTYny PrkaQXz6fId4nhyWDv91TVHLQJpAYDC0VgiREyna/64oKPxfCTG5xQ= DKIM-Signature: v=1; a=rsa-sha1; c=relaxed; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:subject:date:message-id:mime-version:content-type; s= default; bh=kc9qNYHdrHPSnLTiZUKLCC/x0jk=; b=RT5cvJpTB84PRoL0owNV ubXTGKZFCBwEFKHiXFoYjxiz0krhw5Em1tFHyKkc1Vhip1NFMFc6zCseq8zls/hL sugykxsTjq89qtOhG6jYHsHzk9txaKsRFooa13C8Al70mhORNyprXFyOfKKgH8Kx ByKL5rYf5JUAr2L+CdKBbCY= Received: (qmail 129046 invoked by alias); 7 Nov 2019 09:53:10 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Unsubscribe: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Delivered-To: mailing list gcc-patches@gcc.gnu.org Received: (qmail 129024 invoked by uid 89); 7 Nov 2019 09:53:09 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-18.5 required=5.0 tests=AWL, BAYES_00, GIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3, SPF_PASS autolearn=ham version=3.3.1 spammy=constexprc, UD:constexpr.c, constexpr.c, 50859 X-HELO: esa3.mentor.iphmx.com Received: from esa3.mentor.iphmx.com (HELO esa3.mentor.iphmx.com) (68.232.137.180) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Thu, 07 Nov 2019 09:53:03 +0000 IronPort-SDR: 1uxU/rzQ3yNOgzSEp2TE08ZFwEpa3VlUO/dgBEHf1hJkR1nEN0XWl5JmTRvZSQS1TG/TOodNz+ 2aluGnO/AJF6ufqhj7z8jmX9Qjj7aLTSlbd/eCP5luOqiX83TFGDRwIJDbDsAA89jcaOyj4ARb +E0mgpjx0nrGfr4Tk06kS22duZ10HKeM0c2SL/Q+h2w/7sEvC6E9fXZFBMApUaD/qAHXi+uh4+ 38NT4dkKa+sv8fyFOsAeWOPikzaKBjuLZgWg+VkDdDQGJCVWYgi6/1EC8Rp7lzn2Nri4TA8oRo B8c= Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa3.mentor.iphmx.com with ESMTP; 07 Nov 2019 01:53:01 -0800 IronPort-SDR: bAWuoZWBRff1Dspd8v/YDBrHtd/wtDNVbcDiMC9gk3axj7dLY/tA/QGzNpkg1HhtuPDJjRqRdA j7eaGXgYJ8Mx+6PncZ6EXYojBuaLm9Q1tGppDkaqc2R9DMNKX7gJI7VNVpwS5v+Oj1+Ns//Wtm msNrDZhpNa+EnUTCwWE8Kp16c0wUWrjXSYcF5kDAZI2rhEREffKT87MjnFT5B2P6tOuBXH60qs X6PtSqZXQUSYO/2d9NIjWhVj1XTyf96M8ivL13HVIuvoSGErSGepd7eyX3Y/0x0FAbWcIFfviD LRc= From: Frederik Harwath To: , , , , Subject: [PATCH] Add OpenACC 2.6 `serial' construct support Date: Thu, 7 Nov 2019 10:52:13 +0100 Message-ID: <20191107095213.11618-1-frederik@codesourcery.com> MIME-Version: 1.0 X-IsSubscribed: yes Hi, this patch implements the OpenACC 2.6 "serial" construct. It has been tested by running the testsuite with nvptx-none offloading on x86_64-pc-linux-gnu. Best regards, Frederik -------------------- 8< ------------------- The `serial' construct (cf. section 2.5.3 of the OpenACC 2.6 standard) is equivalent to a `parallel' construct with clauses `num_gangs(1) num_workers(1) vector_length(1)' implied. These clauses are therefore not supported with the `serial' construct. All the remaining clauses accepted with `parallel' are also accepted with `serial'. The `serial' construct is implemented like `parallel', except for hardcoding dimensions rather than taking them from the relevant clauses, in `expand_omp_target'. Separate codes are used to denote the `serial' construct throughout the middle end, even though the mapping of `serial' to an equivalent `parallel' construct could have been done in the individual language frontends. In particular, this allows to distinguish between `parallel' and `serial' in warnings, error messages, dumps etc. 2019-11-07 Maciej W. Rozycki Tobias Burnus Frederik Harwath gcc/ * gimple.h (gf_mask): Add GF_OMP_TARGET_KIND_OACC_SERIAL enumeration constant. (is_gimple_omp_oacc): Handle GF_OMP_TARGET_KIND_OACC_SERIAL. (is_gimple_omp_offloaded): Likewise. * gimplify.c (omp_region_type): Add ORT_ACC_SERIAL enumeration constant. Adjust the value of ORT_NONE accordingly. (is_gimple_stmt): Handle OACC_SERIAL. (oacc_default_clause): Handle ORT_ACC_SERIAL. (gomp_needs_data_present): Likewise. (gimplify_adjust_omp_clauses): Likewise. (gimplify_omp_workshare): Handle OACC_SERIAL. (gimplify_expr): Likewise. * omp-builtins.def (BUILT_IN_GOACC_PARALLEL): Add parameter. * omp-expand.c (expand_omp_target): Handle GF_OMP_TARGET_KIND_OACC_SERIAL. (build_omp_regions_1, omp_make_gimple_edges): Likewise. * omp-low.c (is_oacc_parallel): Rename function to... (is_oacc_parallel_or_serial): ... this. Handle GF_OMP_TARGET_KIND_OACC_SERIAL. (scan_sharing_clauses): Adjust accordingly. (scan_omp_for): Likewise. (lower_oacc_head_mark): Likewise. (convert_from_firstprivate_int): Likewise. (lower_omp_target): Likewise. (check_omp_nesting_restrictions): Handle GF_OMP_TARGET_KIND_OACC_SERIAL. (lower_oacc_reductions): Likewise. (lower_omp_target): Likewise. * tree.def (OACC_SERIAL): New tree code. * tree-pretty-print.c (dump_generic_node): Handle OACC_SERIAL. * doc/generic.texi (OpenACC): Document OACC_SERIAL. gcc/c-family/ * c-pragma.h (pragma_kind): Add PRAGMA_OACC_SERIAL enumeration constant. * c-pragma.c (oacc_pragmas): Add "serial" entry. gcc/c/ * c-parser.c (OACC_SERIAL_CLAUSE_MASK): New macro. (c_parser_oacc_kernels_parallel): Rename function to... (c_parser_oacc_compute): ... this. Handle PRAGMA_OACC_SERIAL. (c_parser_omp_construct): Update accordingly. gcc/cp/ * constexpr.c (potential_constant_expression_1): Handle OACC_SERIAL. * parser.c (OACC_SERIAL_CLAUSE_MASK): New macro. (cp_parser_oacc_kernels_parallel): Rename function to... (cp_parser_oacc_compute): ... this. Handle PRAGMA_OACC_SERIAL. (cp_parser_omp_construct): Update accordingly. (cp_parser_pragma): Handle PRAGMA_OACC_SERIAL. Fix alphabetic order. * pt.c (tsubst_expr): Handle OACC_SERIAL. gcc/fortran/ * gfortran.h (gfc_statement): Add ST_OACC_SERIAL_LOOP, ST_OACC_END_SERIAL_LOOP, ST_OACC_SERIAL and ST_OACC_END_SERIAL enumeration constants. (gfc_exec_op): Add EXEC_OACC_SERIAL_LOOP and EXEC_OACC_SERIAL enumeration constants. * match.h (gfc_match_oacc_serial): New prototype. (gfc_match_oacc_serial_loop): Likewise. * dump-parse-tree.c (show_omp_node, show_code_node): Handle EXEC_OACC_SERIAL_LOOP and EXEC_OACC_SERIAL. * match.c (match_exit_cycle): Handle EXEC_OACC_SERIAL_LOOP. * openmp.c (OACC_SERIAL_CLAUSES): New macro. (gfc_match_oacc_serial_loop): New function. (gfc_match_oacc_serial): Likewise. (oacc_is_loop): Handle EXEC_OACC_SERIAL_LOOP. (resolve_omp_clauses): Handle EXEC_OACC_SERIAL. (oacc_code_to_statement): Handle EXEC_OACC_SERIAL and EXEC_OACC_SERIAL_LOOP. (gfc_resolve_oacc_directive): Likewise. * parse.c (decode_oacc_directive) <'s'>: Add case for "serial" and "serial loop". (next_statement): Handle ST_OACC_SERIAL_LOOP and ST_OACC_SERIAL. (gfc_ascii_statement): Likewise. Handle ST_OACC_END_SERIAL_LOOP and ST_OACC_END_SERIAL. (parse_oacc_structured_block): Handle ST_OACC_SERIAL. (parse_oacc_loop): Handle ST_OACC_SERIAL_LOOP and ST_OACC_END_SERIAL_LOOP. (parse_executable): Handle ST_OACC_SERIAL_LOOP and ST_OACC_SERIAL. (is_oacc): Handle EXEC_OACC_SERIAL_LOOP and EXEC_OACC_SERIAL. * resolve.c (gfc_resolve_blocks, gfc_resolve_code): Likewise. * st.c (gfc_free_statement): Likewise. * trans-openmp.c (gfc_trans_oacc_construct): Handle EXEC_OACC_SERIAL. (gfc_trans_oacc_combined_directive): Handle EXEC_OACC_SERIAL_LOOP. (gfc_trans_oacc_directive): Handle EXEC_OACC_SERIAL_LOOP and EXEC_OACC_SERIAL. * trans.c (trans_code): Likewise. gcc/testsuite/ * c-c++-common/goacc/serial-dims.c: New test. * gfortran.dg/goacc/serial-dims.f90: New test. libgomp/ * testsuite/libgomp.oacc-c-c++-common/serial-dims.c: New test. * testsuite/libgomp.oacc-fortran/serial-dims-aux.c: New test. * testsuite/libgomp.oacc-fortran/serial-dims.f90: New test. --- gcc/c-family/c-pragma.c | 1 + gcc/c-family/c-pragma.h | 1 + gcc/c/c-parser.c | 34 ++++++- gcc/cp/constexpr.c | 1 + gcc/cp/parser.c | 35 ++++++- gcc/cp/pt.c | 1 + gcc/doc/generic.texi | 5 + gcc/fortran/dump-parse-tree.c | 6 ++ gcc/fortran/gfortran.h | 13 +-- gcc/fortran/match.c | 3 +- gcc/fortran/match.h | 2 + gcc/fortran/openmp.c | 35 ++++++- gcc/fortran/parse.c | 30 +++++- gcc/fortran/resolve.c | 6 ++ gcc/fortran/st.c | 2 + gcc/fortran/trans-openmp.c | 13 ++- gcc/fortran/trans.c | 2 + gcc/gimple-pretty-print.c | 3 + gcc/gimple.h | 3 + gcc/gimplify.c | 20 +++- gcc/omp-expand.c | 47 ++++++++-- gcc/omp-low.c | 33 ++++--- .../c-c++-common/goacc/serial-dims.c | 12 +++ .../gfortran.dg/goacc/serial-dims.f90 | 40 ++++++++ gcc/tree-pretty-print.c | 4 + gcc/tree.def | 6 ++ .../libgomp.oacc-c-c++-common/serial-dims.c | 92 +++++++++++++++++++ .../libgomp.oacc-fortran/serial-dims-aux.c | 41 +++++++++ .../libgomp.oacc-fortran/serial-dims.f90 | 89 ++++++++++++++++++ 29 files changed, 535 insertions(+), 45 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/goacc/serial-dims.c create mode 100644 gcc/testsuite/gfortran.dg/goacc/serial-dims.f90 create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/serial-dims.c create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/serial-dims-aux.c create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/serial-dims.f90 diff --git a/gcc/c-family/c-pragma.c b/gcc/c-family/c-pragma.c index 9fee84b22383..158154ec1294 100644 --- a/gcc/c-family/c-pragma.c +++ b/gcc/c-family/c-pragma.c @@ -1291,6 +1291,7 @@ static const struct omp_pragma_def oacc_pragmas[] = { { "loop", PRAGMA_OACC_LOOP }, { "parallel", PRAGMA_OACC_PARALLEL }, { "routine", PRAGMA_OACC_ROUTINE }, + { "serial", PRAGMA_OACC_SERIAL }, { "update", PRAGMA_OACC_UPDATE }, { "wait", PRAGMA_OACC_WAIT } }; diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h index e0aa774555a3..bfe681bb430a 100644 --- a/gcc/c-family/c-pragma.h +++ b/gcc/c-family/c-pragma.h @@ -38,6 +38,7 @@ enum pragma_kind { PRAGMA_OACC_LOOP, PRAGMA_OACC_PARALLEL, PRAGMA_OACC_ROUTINE, + PRAGMA_OACC_SERIAL, PRAGMA_OACC_UPDATE, PRAGMA_OACC_WAIT, diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c index 4f044127a7e2..f5d217d0b7a4 100644 --- a/gcc/c/c-parser.c +++ b/gcc/c/c-parser.c @@ -16060,6 +16060,11 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name, # pragma acc parallel oacc-parallel-clause[optseq] new-line structured-block + OpenACC 2.6: + + # pragma acc serial oacc-serial-clause[optseq] new-line + structured-block + LOC is the location of the #pragma token. */ @@ -16096,10 +16101,24 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name, | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) ) +#define OACC_SERIAL_CLAUSE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRIVATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_FIRSTPRIVATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) ) + static tree -c_parser_oacc_kernels_parallel (location_t loc, c_parser *parser, - enum pragma_kind p_kind, char *p_name, - bool *if_p) +c_parser_oacc_compute (location_t loc, c_parser *parser, + enum pragma_kind p_kind, char *p_name, bool *if_p) { omp_clause_mask mask; enum tree_code code; @@ -16115,6 +16134,11 @@ c_parser_oacc_kernels_parallel (location_t loc, c_parser *parser, mask = OACC_PARALLEL_CLAUSE_MASK; code = OACC_PARALLEL; break; + case PRAGMA_OACC_SERIAL: + strcat (p_name, " serial"); + mask = OACC_SERIAL_CLAUSE_MASK; + code = OACC_SERIAL; + break; default: gcc_unreachable (); } @@ -20578,9 +20602,9 @@ c_parser_omp_construct (c_parser *parser, bool *if_p) break; case PRAGMA_OACC_KERNELS: case PRAGMA_OACC_PARALLEL: + case PRAGMA_OACC_SERIAL: strcpy (p_name, "#pragma acc"); - stmt = c_parser_oacc_kernels_parallel (loc, parser, p_kind, p_name, - if_p); + stmt = c_parser_oacc_compute (loc, parser, p_kind, p_name, if_p); break; case PRAGMA_OACC_LOOP: strcpy (p_name, "#pragma acc"); diff --git a/gcc/cp/constexpr.c b/gcc/cp/constexpr.c index 20fddc57825a..8c79b0484fce 100644 --- a/gcc/cp/constexpr.c +++ b/gcc/cp/constexpr.c @@ -6986,6 +6986,7 @@ potential_constant_expression_1 (tree t, bool want_rval, bool strict, bool now, case OMP_DEPOBJ: case OACC_PARALLEL: case OACC_KERNELS: + case OACC_SERIAL: case OACC_DATA: case OACC_HOST_DATA: case OACC_LOOP: diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c index 7138aebebced..c45bfccf6e61 100644 --- a/gcc/cp/parser.c +++ b/gcc/cp/parser.c @@ -40150,6 +40150,10 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name, # pragma acc parallel oacc-parallel-clause[optseq] new-line structured-block + + OpenACC 2.6: + + # pragma acc serial oacc-serial-clause[optseq] new-line */ #define OACC_KERNELS_CLAUSE_MASK \ @@ -40185,9 +40189,24 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, char *p_name, | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) ) +#define OACC_SERIAL_CLAUSE_MASK \ + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRIVATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_FIRSTPRIVATE) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) ) + static tree -cp_parser_oacc_kernels_parallel (cp_parser *parser, cp_token *pragma_tok, - char *p_name, bool *if_p) +cp_parser_oacc_compute (cp_parser *parser, cp_token *pragma_tok, + char *p_name, bool *if_p) { omp_clause_mask mask; enum tree_code code; @@ -40203,6 +40222,11 @@ cp_parser_oacc_kernels_parallel (cp_parser *parser, cp_token *pragma_tok, mask = OACC_PARALLEL_CLAUSE_MASK; code = OACC_PARALLEL; break; + case PRAGMA_OACC_SERIAL: + strcat (p_name, " serial"); + mask = OACC_SERIAL_CLAUSE_MASK; + code = OACC_SERIAL; + break; default: gcc_unreachable (); } @@ -42022,9 +42046,9 @@ cp_parser_omp_construct (cp_parser *parser, cp_token *pragma_tok, bool *if_p) break; case PRAGMA_OACC_KERNELS: case PRAGMA_OACC_PARALLEL: + case PRAGMA_OACC_SERIAL: strcpy (p_name, "#pragma acc"); - stmt = cp_parser_oacc_kernels_parallel (parser, pragma_tok, p_name, - if_p); + stmt = cp_parser_oacc_compute (parser, pragma_tok, p_name, if_p); break; case PRAGMA_OACC_LOOP: strcpy (p_name, "#pragma acc"); @@ -42691,8 +42715,9 @@ cp_parser_pragma (cp_parser *parser, enum pragma_context context, bool *if_p) case PRAGMA_OACC_DATA: case PRAGMA_OACC_HOST_DATA: case PRAGMA_OACC_KERNELS: - case PRAGMA_OACC_PARALLEL: case PRAGMA_OACC_LOOP: + case PRAGMA_OACC_PARALLEL: + case PRAGMA_OACC_SERIAL: case PRAGMA_OMP_ATOMIC: case PRAGMA_OMP_CRITICAL: case PRAGMA_OMP_DISTRIBUTE: diff --git a/gcc/cp/pt.c b/gcc/cp/pt.c index 8bacb3952ff2..5a0efaa86c8b 100644 --- a/gcc/cp/pt.c +++ b/gcc/cp/pt.c @@ -17991,6 +17991,7 @@ tsubst_expr (tree t, tree args, tsubst_flags_t complain, tree in_decl, case OACC_KERNELS: case OACC_PARALLEL: + case OACC_SERIAL: tmp = tsubst_omp_clauses (OMP_CLAUSES (t), C_ORT_ACC, args, complain, in_decl); stmt = begin_omp_parallel (); diff --git a/gcc/doc/generic.texi b/gcc/doc/generic.texi index 94e339c15ee8..badaaec38979 100644 --- a/gcc/doc/generic.texi +++ b/gcc/doc/generic.texi @@ -2388,6 +2388,7 @@ compilation. @tindex OACC_KERNELS @tindex OACC_LOOP @tindex OACC_PARALLEL +@tindex OACC_SERIAL @tindex OACC_UPDATE All the statements starting with @code{OACC_} represent directives and @@ -2432,6 +2433,10 @@ See the description of the @code{OMP_FOR} code. Represents @code{#pragma acc parallel [clause1 @dots{} clauseN]}. +@item OACC_SERIAL + +Represents @code{#pragma acc serial [clause1 @dots{} clauseN]}. + @item OACC_UPDATE Represents @code{#pragma acc update [clause1 @dots{} clauseN]}. diff --git a/gcc/fortran/dump-parse-tree.c b/gcc/fortran/dump-parse-tree.c index 9d7aad19e2f5..253fe15b201d 100644 --- a/gcc/fortran/dump-parse-tree.c +++ b/gcc/fortran/dump-parse-tree.c @@ -1654,6 +1654,8 @@ show_omp_node (int level, gfc_code *c) case EXEC_OACC_PARALLEL: name = "PARALLEL"; is_oacc = true; break; case EXEC_OACC_KERNELS_LOOP: name = "KERNELS LOOP"; is_oacc = true; break; case EXEC_OACC_KERNELS: name = "KERNELS"; is_oacc = true; break; + case EXEC_OACC_SERIAL_LOOP: name = "SERIAL LOOP"; is_oacc = true; break; + case EXEC_OACC_SERIAL: name = "SERIAL"; is_oacc = true; break; case EXEC_OACC_DATA: name = "DATA"; is_oacc = true; break; case EXEC_OACC_HOST_DATA: name = "HOST_DATA"; is_oacc = true; break; case EXEC_OACC_LOOP: name = "LOOP"; is_oacc = true; break; @@ -1729,6 +1731,8 @@ show_omp_node (int level, gfc_code *c) case EXEC_OACC_PARALLEL: case EXEC_OACC_KERNELS_LOOP: case EXEC_OACC_KERNELS: + case EXEC_OACC_SERIAL_LOOP: + case EXEC_OACC_SERIAL: case EXEC_OACC_DATA: case EXEC_OACC_HOST_DATA: case EXEC_OACC_LOOP: @@ -2918,6 +2922,8 @@ show_code_node (int level, gfc_code *c) case EXEC_OACC_PARALLEL: case EXEC_OACC_KERNELS_LOOP: case EXEC_OACC_KERNELS: + case EXEC_OACC_SERIAL_LOOP: + case EXEC_OACC_SERIAL: case EXEC_OACC_DATA: case EXEC_OACC_HOST_DATA: case EXEC_OACC_LOOP: diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h index 920acdafc6b7..e962db59bc59 100644 --- a/gcc/fortran/gfortran.h +++ b/gcc/fortran/gfortran.h @@ -223,7 +223,8 @@ enum gfc_statement ST_OACC_END_DATA, ST_OACC_HOST_DATA, ST_OACC_END_HOST_DATA, ST_OACC_LOOP, ST_OACC_END_LOOP, ST_OACC_DECLARE, ST_OACC_UPDATE, ST_OACC_WAIT, ST_OACC_CACHE, ST_OACC_KERNELS_LOOP, ST_OACC_END_KERNELS_LOOP, - ST_OACC_ENTER_DATA, ST_OACC_EXIT_DATA, ST_OACC_ROUTINE, + ST_OACC_SERIAL_LOOP, ST_OACC_END_SERIAL_LOOP, ST_OACC_SERIAL, + ST_OACC_END_SERIAL, ST_OACC_ENTER_DATA, ST_OACC_EXIT_DATA, ST_OACC_ROUTINE, ST_OACC_ATOMIC, ST_OACC_END_ATOMIC, ST_OMP_ATOMIC, ST_OMP_BARRIER, ST_OMP_CRITICAL, ST_OMP_END_ATOMIC, ST_OMP_END_CRITICAL, ST_OMP_END_DO, ST_OMP_END_MASTER, ST_OMP_END_ORDERED, @@ -2572,11 +2573,11 @@ enum gfc_exec_op EXEC_BACKSPACE, EXEC_ENDFILE, EXEC_INQUIRE, EXEC_REWIND, EXEC_FLUSH, EXEC_FORM_TEAM, EXEC_CHANGE_TEAM, EXEC_END_TEAM, EXEC_SYNC_TEAM, EXEC_LOCK, EXEC_UNLOCK, EXEC_EVENT_POST, EXEC_EVENT_WAIT, EXEC_FAIL_IMAGE, - EXEC_OACC_KERNELS_LOOP, EXEC_OACC_PARALLEL_LOOP, EXEC_OACC_ROUTINE, - EXEC_OACC_PARALLEL, EXEC_OACC_KERNELS, EXEC_OACC_DATA, EXEC_OACC_HOST_DATA, - EXEC_OACC_LOOP, EXEC_OACC_UPDATE, EXEC_OACC_WAIT, EXEC_OACC_CACHE, - EXEC_OACC_ENTER_DATA, EXEC_OACC_EXIT_DATA, EXEC_OACC_ATOMIC, - EXEC_OACC_DECLARE, + EXEC_OACC_KERNELS_LOOP, EXEC_OACC_PARALLEL_LOOP, EXEC_OACC_SERIAL_LOOP, + EXEC_OACC_ROUTINE, EXEC_OACC_PARALLEL, EXEC_OACC_KERNELS, EXEC_OACC_SERIAL, + EXEC_OACC_DATA, EXEC_OACC_HOST_DATA, EXEC_OACC_LOOP, EXEC_OACC_UPDATE, + EXEC_OACC_WAIT, EXEC_OACC_CACHE, EXEC_OACC_ENTER_DATA, EXEC_OACC_EXIT_DATA, + EXEC_OACC_ATOMIC, EXEC_OACC_DECLARE, EXEC_OMP_CRITICAL, EXEC_OMP_DO, EXEC_OMP_FLUSH, EXEC_OMP_MASTER, EXEC_OMP_ORDERED, EXEC_OMP_PARALLEL, EXEC_OMP_PARALLEL_DO, EXEC_OMP_PARALLEL_SECTIONS, EXEC_OMP_PARALLEL_WORKSHARE, diff --git a/gcc/fortran/match.c b/gcc/fortran/match.c index 4a31080a2856..b5945049de55 100644 --- a/gcc/fortran/match.c +++ b/gcc/fortran/match.c @@ -2860,7 +2860,8 @@ match_exit_cycle (gfc_statement st, gfc_exec_op op) && o != NULL && o->state == COMP_OMP_STRUCTURED_BLOCK && (o->head->op == EXEC_OACC_LOOP - || o->head->op == EXEC_OACC_PARALLEL_LOOP)) + || o->head->op == EXEC_OACC_PARALLEL_LOOP + || o->head->op == EXEC_OACC_SERIAL_LOOP)) { int collapse = 1; gcc_assert (o->head->next != NULL diff --git a/gcc/fortran/match.h b/gcc/fortran/match.h index 611d79646458..954af72f0e07 100644 --- a/gcc/fortran/match.h +++ b/gcc/fortran/match.h @@ -147,6 +147,8 @@ match gfc_match_oacc_kernels_loop (void); match gfc_match_oacc_parallel (void); match gfc_match_oacc_parallel_loop (void); match gfc_match_oacc_enter_data (void); +match gfc_match_oacc_serial (void); +match gfc_match_oacc_serial_loop (void); match gfc_match_oacc_exit_data (void); match gfc_match_oacc_routine (void); diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c index ca3427885457..198facce636d 100644 --- a/gcc/fortran/openmp.c +++ b/gcc/fortran/openmp.c @@ -1964,6 +1964,15 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, | OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \ | OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_DEFAULT \ | OMP_CLAUSE_WAIT) +#define OACC_SERIAL_CLAUSES \ + (omp_mask (OMP_CLAUSE_ASYNC) | OMP_CLAUSE_WAIT \ + | OMP_CLAUSE_IF \ + | OMP_CLAUSE_REDUCTION \ + | OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \ + | OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT \ + | OMP_CLAUSE_DEVICEPTR \ + | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE \ + | OMP_CLAUSE_DEFAULT) #define OACC_DATA_CLAUSES \ (omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_COPY \ | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT | OMP_CLAUSE_CREATE \ @@ -2038,6 +2047,21 @@ gfc_match_oacc_kernels (void) } +match +gfc_match_oacc_serial_loop (void) +{ + return match_acc (EXEC_OACC_SERIAL_LOOP, + OACC_SERIAL_CLAUSES | OACC_LOOP_CLAUSES); +} + + +match +gfc_match_oacc_serial (void) +{ + return match_acc (EXEC_OACC_SERIAL, OACC_SERIAL_CLAUSES); +} + + match gfc_match_oacc_data (void) { @@ -3783,6 +3807,7 @@ oacc_is_loop (gfc_code *code) { return code->op == EXEC_OACC_PARALLEL_LOOP || code->op == EXEC_OACC_KERNELS_LOOP + || code->op == EXEC_OACC_SERIAL_LOOP || code->op == EXEC_OACC_LOOP; } @@ -4626,7 +4651,9 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, n->sym->name, name, &n->where); } if (code - && (oacc_is_loop (code) || code->op == EXEC_OACC_PARALLEL)) + && (oacc_is_loop (code) + || code->op == EXEC_OACC_PARALLEL + || code->op == EXEC_OACC_SERIAL)) check_array_not_assumed (n->sym, n->where, name); else if (n->sym->as && n->sym->as->type == AS_ASSUMED_SIZE) gfc_error ("Assumed size array %qs in %s clause at %L", @@ -5818,6 +5845,8 @@ oacc_code_to_statement (gfc_code *code) return ST_OACC_PARALLEL; case EXEC_OACC_KERNELS: return ST_OACC_KERNELS; + case EXEC_OACC_SERIAL: + return ST_OACC_SERIAL; case EXEC_OACC_DATA: return ST_OACC_DATA; case EXEC_OACC_HOST_DATA: @@ -5826,6 +5855,8 @@ oacc_code_to_statement (gfc_code *code) return ST_OACC_PARALLEL_LOOP; case EXEC_OACC_KERNELS_LOOP: return ST_OACC_KERNELS_LOOP; + case EXEC_OACC_SERIAL_LOOP: + return ST_OACC_SERIAL_LOOP; case EXEC_OACC_LOOP: return ST_OACC_LOOP; case EXEC_OACC_ATOMIC: @@ -6163,6 +6194,7 @@ gfc_resolve_oacc_directive (gfc_code *code, gfc_namespace *ns ATTRIBUTE_UNUSED) { case EXEC_OACC_PARALLEL: case EXEC_OACC_KERNELS: + case EXEC_OACC_SERIAL: case EXEC_OACC_DATA: case EXEC_OACC_HOST_DATA: case EXEC_OACC_UPDATE: @@ -6174,6 +6206,7 @@ gfc_resolve_oacc_directive (gfc_code *code, gfc_namespace *ns ATTRIBUTE_UNUSED) break; case EXEC_OACC_PARALLEL_LOOP: case EXEC_OACC_KERNELS_LOOP: + case EXEC_OACC_SERIAL_LOOP: case EXEC_OACC_LOOP: resolve_oacc_loop (code); break; diff --git a/gcc/fortran/parse.c b/gcc/fortran/parse.c index 15f6bf2937c4..1a38606682ca 100644 --- a/gcc/fortran/parse.c +++ b/gcc/fortran/parse.c @@ -683,6 +683,9 @@ decode_oacc_directive (void) matcha ("end parallel loop", gfc_match_omp_eos_error, ST_OACC_END_PARALLEL_LOOP); matcha ("end parallel", gfc_match_omp_eos_error, ST_OACC_END_PARALLEL); + matcha ("end serial loop", gfc_match_omp_eos_error, + ST_OACC_END_SERIAL_LOOP); + matcha ("end serial", gfc_match_omp_eos_error, ST_OACC_END_SERIAL); matcha ("enter data", gfc_match_oacc_enter_data, ST_OACC_ENTER_DATA); matcha ("exit data", gfc_match_oacc_exit_data, ST_OACC_EXIT_DATA); break; @@ -705,6 +708,10 @@ decode_oacc_directive (void) case 'r': match ("routine", gfc_match_oacc_routine, ST_OACC_ROUTINE); break; + case 's': + matcha ("serial loop", gfc_match_oacc_serial_loop, ST_OACC_SERIAL_LOOP); + matcha ("serial", gfc_match_oacc_serial, ST_OACC_SERIAL); + break; case 'u': matcha ("update", gfc_match_oacc_update, ST_OACC_UPDATE); break; @@ -1583,7 +1590,8 @@ next_statement (void) case ST_CRITICAL: \ case ST_OACC_PARALLEL_LOOP: case ST_OACC_PARALLEL: case ST_OACC_KERNELS: \ case ST_OACC_DATA: case ST_OACC_HOST_DATA: case ST_OACC_LOOP: \ - case ST_OACC_KERNELS_LOOP: case ST_OACC_ATOMIC + case ST_OACC_KERNELS_LOOP: case ST_OACC_SERIAL_LOOP: case ST_OACC_SERIAL: \ + case ST_OACC_ATOMIC /* Declaration statements */ @@ -2157,6 +2165,18 @@ gfc_ascii_statement (gfc_statement st) case ST_OACC_END_KERNELS_LOOP: p = "!$ACC END KERNELS LOOP"; break; + case ST_OACC_SERIAL_LOOP: + p = "!$ACC SERIAL LOOP"; + break; + case ST_OACC_END_SERIAL_LOOP: + p = "!$ACC END SERIAL LOOP"; + break; + case ST_OACC_SERIAL: + p = "!$ACC SERIAL"; + break; + case ST_OACC_END_SERIAL: + p = "!$ACC END SERIAL"; + break; case ST_OACC_DATA: p = "!$ACC DATA"; break; @@ -5065,6 +5085,9 @@ parse_oacc_structured_block (gfc_statement acc_st) case ST_OACC_KERNELS: acc_end_st = ST_OACC_END_KERNELS; break; + case ST_OACC_SERIAL: + acc_end_st = ST_OACC_END_SERIAL; + break; case ST_OACC_DATA: acc_end_st = ST_OACC_END_DATA; break; @@ -5149,6 +5172,7 @@ parse_oacc_loop (gfc_statement acc_st) gfc_warning (0, "Redundant !$ACC END LOOP at %C"); if ((acc_st == ST_OACC_PARALLEL_LOOP && st == ST_OACC_END_PARALLEL_LOOP) || (acc_st == ST_OACC_KERNELS_LOOP && st == ST_OACC_END_KERNELS_LOOP) || + (acc_st == ST_OACC_SERIAL_LOOP && st == ST_OACC_END_SERIAL_LOOP) || (acc_st == ST_OACC_LOOP && st == ST_OACC_END_LOOP)) { gcc_assert (new_st.op == EXEC_NOP); @@ -5488,6 +5512,7 @@ parse_executable (gfc_statement st) case ST_OACC_PARALLEL_LOOP: case ST_OACC_KERNELS_LOOP: + case ST_OACC_SERIAL_LOOP: case ST_OACC_LOOP: st = parse_oacc_loop (st); if (st == ST_IMPLIED_ENDDO) @@ -5496,6 +5521,7 @@ parse_executable (gfc_statement st) case ST_OACC_PARALLEL: case ST_OACC_KERNELS: + case ST_OACC_SERIAL: case ST_OACC_DATA: case ST_OACC_HOST_DATA: parse_oacc_structured_block (st); @@ -6544,6 +6570,8 @@ is_oacc (gfc_state_data *sd) case EXEC_OACC_PARALLEL: case EXEC_OACC_KERNELS_LOOP: case EXEC_OACC_KERNELS: + case EXEC_OACC_SERIAL_LOOP: + case EXEC_OACC_SERIAL: case EXEC_OACC_DATA: case EXEC_OACC_HOST_DATA: case EXEC_OACC_LOOP: diff --git a/gcc/fortran/resolve.c b/gcc/fortran/resolve.c index 218c2edba57c..9b1437d70327 100644 --- a/gcc/fortran/resolve.c +++ b/gcc/fortran/resolve.c @@ -10576,6 +10576,8 @@ gfc_resolve_blocks (gfc_code *b, gfc_namespace *ns) case EXEC_OACC_PARALLEL: case EXEC_OACC_KERNELS_LOOP: case EXEC_OACC_KERNELS: + case EXEC_OACC_SERIAL_LOOP: + case EXEC_OACC_SERIAL: case EXEC_OACC_DATA: case EXEC_OACC_HOST_DATA: case EXEC_OACC_LOOP: @@ -11527,6 +11529,8 @@ gfc_resolve_code (gfc_code *code, gfc_namespace *ns) case EXEC_OACC_PARALLEL: case EXEC_OACC_KERNELS_LOOP: case EXEC_OACC_KERNELS: + case EXEC_OACC_SERIAL_LOOP: + case EXEC_OACC_SERIAL: case EXEC_OACC_DATA: case EXEC_OACC_HOST_DATA: case EXEC_OACC_LOOP: @@ -11940,6 +11944,8 @@ start: case EXEC_OACC_PARALLEL: case EXEC_OACC_KERNELS_LOOP: case EXEC_OACC_KERNELS: + case EXEC_OACC_SERIAL_LOOP: + case EXEC_OACC_SERIAL: case EXEC_OACC_DATA: case EXEC_OACC_HOST_DATA: case EXEC_OACC_LOOP: diff --git a/gcc/fortran/st.c b/gcc/fortran/st.c index ee18d7aea8ad..12eed71e3a26 100644 --- a/gcc/fortran/st.c +++ b/gcc/fortran/st.c @@ -202,6 +202,8 @@ gfc_free_statement (gfc_code *p) case EXEC_OACC_PARALLEL: case EXEC_OACC_KERNELS_LOOP: case EXEC_OACC_KERNELS: + case EXEC_OACC_SERIAL_LOOP: + case EXEC_OACC_SERIAL: case EXEC_OACC_DATA: case EXEC_OACC_HOST_DATA: case EXEC_OACC_LOOP: diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c index 14a3c3e42843..0d5a5a9615d6 100644 --- a/gcc/fortran/trans-openmp.c +++ b/gcc/fortran/trans-openmp.c @@ -3141,7 +3141,7 @@ gfc_trans_omp_code (gfc_code *code, bool force_empty) } /* Trans OpenACC directives. */ -/* parallel, kernels, data and host_data. */ +/* parallel, serial, kernels, data and host_data. */ static tree gfc_trans_oacc_construct (gfc_code *code) { @@ -3157,6 +3157,9 @@ gfc_trans_oacc_construct (gfc_code *code) case EXEC_OACC_KERNELS: construct_code = OACC_KERNELS; break; + case EXEC_OACC_SERIAL: + construct_code = OACC_SERIAL; + break; case EXEC_OACC_DATA: construct_code = OACC_DATA; break; @@ -3964,7 +3967,8 @@ gfc_trans_omp_do (gfc_code *code, gfc_exec_op op, stmtblock_t *pblock, return gfc_finish_block (&block); } -/* parallel loop and kernels loop. */ +/* Combined OpenACC parallel loop, kernels loop and serial loop. */ + static tree gfc_trans_oacc_combined_directive (gfc_code *code) { @@ -3982,6 +3986,9 @@ gfc_trans_oacc_combined_directive (gfc_code *code) case EXEC_OACC_KERNELS_LOOP: construct_code = OACC_KERNELS; break; + case EXEC_OACC_SERIAL_LOOP: + construct_code = OACC_SERIAL; + break; default: gcc_unreachable (); } @@ -5214,9 +5221,11 @@ gfc_trans_oacc_directive (gfc_code *code) { case EXEC_OACC_PARALLEL_LOOP: case EXEC_OACC_KERNELS_LOOP: + case EXEC_OACC_SERIAL_LOOP: return gfc_trans_oacc_combined_directive (code); case EXEC_OACC_PARALLEL: case EXEC_OACC_KERNELS: + case EXEC_OACC_SERIAL: case EXEC_OACC_DATA: case EXEC_OACC_HOST_DATA: return gfc_trans_oacc_construct (code); diff --git a/gcc/fortran/trans.c b/gcc/fortran/trans.c index 2f878f6b1185..d9b278199b75 100644 --- a/gcc/fortran/trans.c +++ b/gcc/fortran/trans.c @@ -2137,6 +2137,8 @@ trans_code (gfc_code * code, tree cond) case EXEC_OACC_KERNELS_LOOP: case EXEC_OACC_PARALLEL: case EXEC_OACC_PARALLEL_LOOP: + case EXEC_OACC_SERIAL: + case EXEC_OACC_SERIAL_LOOP: case EXEC_OACC_ENTER_DATA: case EXEC_OACC_EXIT_DATA: case EXEC_OACC_ATOMIC: diff --git a/gcc/gimple-pretty-print.c b/gcc/gimple-pretty-print.c index 2d5ece068053..f59cc2aa3188 100644 --- a/gcc/gimple-pretty-print.c +++ b/gcc/gimple-pretty-print.c @@ -1676,6 +1676,9 @@ dump_gimple_omp_target (pretty_printer *buffer, gomp_target *gs, case GF_OMP_TARGET_KIND_OACC_PARALLEL: kind = " oacc_parallel"; break; + case GF_OMP_TARGET_KIND_OACC_SERIAL: + kind = " oacc_serial"; + break; case GF_OMP_TARGET_KIND_OACC_DATA: kind = " oacc_data"; break; diff --git a/gcc/gimple.h b/gcc/gimple.h index cf1f8da5ae24..83a449be3643 100644 --- a/gcc/gimple.h +++ b/gcc/gimple.h @@ -182,6 +182,7 @@ enum gf_mask { GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA = 9, GF_OMP_TARGET_KIND_OACC_DECLARE = 10, GF_OMP_TARGET_KIND_OACC_HOST_DATA = 11, + GF_OMP_TARGET_KIND_OACC_SERIAL = 12, GF_OMP_TEAMS_GRID_PHONY = 1 << 0, GF_OMP_TEAMS_HOST = 1 << 1, @@ -6476,6 +6477,7 @@ is_gimple_omp_oacc (const gimple *stmt) { case GF_OMP_TARGET_KIND_OACC_PARALLEL: case GF_OMP_TARGET_KIND_OACC_KERNELS: + case GF_OMP_TARGET_KIND_OACC_SERIAL: case GF_OMP_TARGET_KIND_OACC_DATA: case GF_OMP_TARGET_KIND_OACC_UPDATE: case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA: @@ -6505,6 +6507,7 @@ is_gimple_omp_offloaded (const gimple *stmt) case GF_OMP_TARGET_KIND_REGION: case GF_OMP_TARGET_KIND_OACC_PARALLEL: case GF_OMP_TARGET_KIND_OACC_KERNELS: + case GF_OMP_TARGET_KIND_OACC_SERIAL: return true; default: return false; diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 5fa0ba6dda60..94a69643aaab 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -161,6 +161,7 @@ enum omp_region_type ORT_ACC_DATA = ORT_ACC | ORT_TARGET_DATA, /* Data construct. */ ORT_ACC_PARALLEL = ORT_ACC | ORT_TARGET, /* Parallel construct */ ORT_ACC_KERNELS = ORT_ACC | ORT_TARGET | 2, /* Kernels construct. */ + ORT_ACC_SERIAL = ORT_ACC | ORT_TARGET | 4, /* Serial construct. */ ORT_ACC_HOST_DATA = ORT_ACC | ORT_TARGET_DATA | 2, /* Host data. */ /* Dummy OpenMP region, used to disable expansion of @@ -5551,6 +5552,7 @@ is_gimple_stmt (tree t) case STATEMENT_LIST: case OACC_PARALLEL: case OACC_KERNELS: + case OACC_SERIAL: case OACC_DATA: case OACC_HOST_DATA: case OACC_DECLARE: @@ -7289,7 +7291,8 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags) break; case ORT_ACC_PARALLEL: - rkind = "parallel"; + case ORT_ACC_SERIAL: + rkind = ctx->region_type == ORT_ACC_PARALLEL ? "parallel" : "serial"; if (is_private) flags |= GOVD_FIRSTPRIVATE; @@ -10101,7 +10104,8 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, /* Data clauses associated with acc parallel reductions must be compatible with present_or_copy. Warn and adjust the clause if that is not the case. */ - if (ctx->region_type == ORT_ACC_PARALLEL) + if (ctx->region_type == ORT_ACC_PARALLEL + || ctx->region_type == ORT_ACC_SERIAL) { tree t = DECL_P (decl) ? decl : TREE_OPERAND (decl, 0); n = NULL; @@ -10277,7 +10281,8 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, decl = OMP_CLAUSE_DECL (c); /* OpenACC reductions need a present_or_copy data clause. Add one if necessary. Emit error when the reduction is private. */ - if (ctx->region_type == ORT_ACC_PARALLEL) + if (ctx->region_type == ORT_ACC_PARALLEL + || ctx->region_type == ORT_ACC_SERIAL) { n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl); if (n->value & (GOVD_PRIVATE | GOVD_FIRSTPRIVATE)) @@ -12529,6 +12534,9 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p) case OACC_PARALLEL: ort = ORT_ACC_PARALLEL; break; + case OACC_SERIAL: + ort = ORT_ACC_SERIAL; + break; case OACC_DATA: ort = ORT_ACC_DATA; break; @@ -12612,6 +12620,10 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p) stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_OACC_PARALLEL, OMP_CLAUSES (expr)); break; + case OACC_SERIAL: + stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_OACC_SERIAL, + OMP_CLAUSES (expr)); + break; case OMP_SECTIONS: stmt = gimple_build_omp_sections (body, OMP_CLAUSES (expr)); break; @@ -13870,6 +13882,7 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p, case OACC_DATA: case OACC_KERNELS: case OACC_PARALLEL: + case OACC_SERIAL: case OMP_SECTIONS: case OMP_SINGLE: case OMP_TARGET: @@ -14286,6 +14299,7 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p, && code != EH_ELSE_EXPR && code != OACC_PARALLEL && code != OACC_KERNELS + && code != OACC_SERIAL && code != OACC_DATA && code != OACC_HOST_DATA && code != OACC_DECLARE diff --git a/gcc/omp-expand.c b/gcc/omp-expand.c index eadff6e50f86..d242f4e1ae99 100644 --- a/gcc/omp-expand.c +++ b/gcc/omp-expand.c @@ -7901,19 +7901,22 @@ expand_omp_target (struct omp_region *region) gimple *stmt; edge e; bool offloaded, data_region; + int target_kind; entry_stmt = as_a (last_stmt (region->entry)); + target_kind = gimple_omp_target_kind (entry_stmt); new_bb = region->entry; offloaded = is_gimple_omp_offloaded (entry_stmt); - switch (gimple_omp_target_kind (entry_stmt)) + switch (target_kind) { case GF_OMP_TARGET_KIND_REGION: case GF_OMP_TARGET_KIND_UPDATE: case GF_OMP_TARGET_KIND_ENTER_DATA: case GF_OMP_TARGET_KIND_EXIT_DATA: - case GF_OMP_TARGET_KIND_OACC_PARALLEL: case GF_OMP_TARGET_KIND_OACC_KERNELS: + case GF_OMP_TARGET_KIND_OACC_PARALLEL: + case GF_OMP_TARGET_KIND_OACC_SERIAL: case GF_OMP_TARGET_KIND_OACC_UPDATE: case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA: case GF_OMP_TARGET_KIND_OACC_DECLARE: @@ -7944,16 +7947,28 @@ expand_omp_target (struct omp_region *region) entry_bb = region->entry; exit_bb = region->exit; - if (gimple_omp_target_kind (entry_stmt) == GF_OMP_TARGET_KIND_OACC_KERNELS) + switch (target_kind) { + case GF_OMP_TARGET_KIND_OACC_KERNELS: mark_loops_in_oacc_kernels_region (region->entry, region->exit); - /* Further down, both OpenACC kernels and OpenACC parallel constructs - will be mappted to BUILT_IN_GOACC_PARALLEL, and to distinguish the - two, there is an "oacc kernels" attribute set for OpenACC kernels. */ + /* Further down, all OpenACC compute constructs will be mapped to + BUILT_IN_GOACC_PARALLEL, and to distinguish between them, there + is an "oacc kernels" attribute set for OpenACC kernels. */ DECL_ATTRIBUTES (child_fn) = tree_cons (get_identifier ("oacc kernels"), NULL_TREE, DECL_ATTRIBUTES (child_fn)); + break; + case GF_OMP_TARGET_KIND_OACC_SERIAL: + /* Further down, all OpenACC compute constructs will be mapped to + BUILT_IN_GOACC_PARALLEL, and to distinguish between them, there + is an "oacc serial" attribute set for OpenACC serial. */ + DECL_ATTRIBUTES (child_fn) + = tree_cons (get_identifier ("oacc serial"), + NULL_TREE, DECL_ATTRIBUTES (child_fn)); + break; + default: + break; } if (offloaded) @@ -8158,6 +8173,7 @@ expand_omp_target (struct omp_region *region) break; case GF_OMP_TARGET_KIND_OACC_KERNELS: case GF_OMP_TARGET_KIND_OACC_PARALLEL: + case GF_OMP_TARGET_KIND_OACC_SERIAL: start_ix = BUILT_IN_GOACC_PARALLEL; break; case GF_OMP_TARGET_KIND_OACC_DATA: @@ -8352,7 +8368,18 @@ expand_omp_target (struct omp_region *region) args.quick_push (get_target_arguments (&gsi, entry_stmt)); break; case BUILT_IN_GOACC_PARALLEL: - oacc_set_fn_attrib (child_fn, clauses, &args); + if (lookup_attribute ("oacc serial", DECL_ATTRIBUTES (child_fn)) != NULL) + { + tree dims = NULL_TREE; + unsigned int ix; + + /* For serial constructs we set all dimensions to 1. */ + for (ix = GOMP_DIM_MAX; ix--;) + dims = tree_cons (NULL_TREE, integer_one_node, dims); + oacc_replace_fn_attrib (child_fn, dims); + } + else + oacc_set_fn_attrib (child_fn, clauses, &args); tagging = true; /* FALLTHRU */ case BUILT_IN_GOACC_ENTER_EXIT_DATA: @@ -8911,8 +8938,9 @@ build_omp_regions_1 (basic_block bb, struct omp_region *parent, { case GF_OMP_TARGET_KIND_REGION: case GF_OMP_TARGET_KIND_DATA: - case GF_OMP_TARGET_KIND_OACC_PARALLEL: case GF_OMP_TARGET_KIND_OACC_KERNELS: + case GF_OMP_TARGET_KIND_OACC_PARALLEL: + case GF_OMP_TARGET_KIND_OACC_SERIAL: case GF_OMP_TARGET_KIND_OACC_DATA: case GF_OMP_TARGET_KIND_OACC_HOST_DATA: break; @@ -9165,8 +9193,9 @@ omp_make_gimple_edges (basic_block bb, struct omp_region **region, { case GF_OMP_TARGET_KIND_REGION: case GF_OMP_TARGET_KIND_DATA: - case GF_OMP_TARGET_KIND_OACC_PARALLEL: case GF_OMP_TARGET_KIND_OACC_KERNELS: + case GF_OMP_TARGET_KIND_OACC_PARALLEL: + case GF_OMP_TARGET_KIND_OACC_SERIAL: case GF_OMP_TARGET_KIND_OACC_DATA: case GF_OMP_TARGET_KIND_OACC_HOST_DATA: break; diff --git a/gcc/omp-low.c b/gcc/omp-low.c index fa76ceba33c6..fb2ddc5f354a 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -185,15 +185,17 @@ static tree scan_omp_1_op (tree *, int *, void *); *handled_ops_p = false; \ break; -/* Return true if CTX corresponds to an oacc parallel region. */ +/* Return true if CTX corresponds to an oacc parallel or serial region. */ static bool -is_oacc_parallel (omp_context *ctx) +is_oacc_parallel_or_serial (omp_context *ctx) { enum gimple_code outer_type = gimple_code (ctx->stmt); return ((outer_type == GIMPLE_OMP_TARGET) - && (gimple_omp_target_kind (ctx->stmt) - == GF_OMP_TARGET_KIND_OACC_PARALLEL)); + && ((gimple_omp_target_kind (ctx->stmt) + == GF_OMP_TARGET_KIND_OACC_PARALLEL) + || (gimple_omp_target_kind (ctx->stmt) + == GF_OMP_TARGET_KIND_OACC_SERIAL))); } /* Return true if CTX corresponds to an oacc kernels region. */ @@ -1149,7 +1151,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) goto do_private; case OMP_CLAUSE_REDUCTION: - if (is_oacc_parallel (ctx) || is_oacc_kernels (ctx)) + if (is_oacc_parallel_or_serial (ctx) || is_oacc_kernels (ctx)) ctx->local_reduction_clauses = tree_cons (NULL, c, ctx->local_reduction_clauses); /* FALLTHRU */ @@ -2391,7 +2393,7 @@ scan_omp_for (gomp_for *stmt, omp_context *outer_ctx) { omp_context *tgt = enclosing_target_ctx (outer_ctx); - if (!tgt || is_oacc_parallel (tgt)) + if (!tgt || is_oacc_parallel_or_serial (tgt)) for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c)) { char const *check = NULL; @@ -2945,6 +2947,7 @@ check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx) { case GF_OMP_TARGET_KIND_OACC_PARALLEL: case GF_OMP_TARGET_KIND_OACC_KERNELS: + case GF_OMP_TARGET_KIND_OACC_SERIAL: ok = true; break; @@ -3393,6 +3396,7 @@ check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx) stmt_name = "target exit data"; break; case GF_OMP_TARGET_KIND_OACC_PARALLEL: stmt_name = "parallel"; break; case GF_OMP_TARGET_KIND_OACC_KERNELS: stmt_name = "kernels"; break; + case GF_OMP_TARGET_KIND_OACC_SERIAL: stmt_name = "serial"; break; case GF_OMP_TARGET_KIND_OACC_DATA: stmt_name = "data"; break; case GF_OMP_TARGET_KIND_OACC_UPDATE: stmt_name = "update"; break; case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA: @@ -3410,6 +3414,8 @@ check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx) ctx_stmt_name = "parallel"; break; case GF_OMP_TARGET_KIND_OACC_KERNELS: ctx_stmt_name = "kernels"; break; + case GF_OMP_TARGET_KIND_OACC_SERIAL: + ctx_stmt_name = "serial"; break; case GF_OMP_TARGET_KIND_OACC_DATA: ctx_stmt_name = "data"; break; case GF_OMP_TARGET_KIND_OACC_HOST_DATA: ctx_stmt_name = "host_data"; break; @@ -6711,8 +6717,10 @@ lower_oacc_reductions (location_t loc, tree clauses, tree level, bool inner, break; case GIMPLE_OMP_TARGET: - if (gimple_omp_target_kind (probe->stmt) - != GF_OMP_TARGET_KIND_OACC_PARALLEL) + if ((gimple_omp_target_kind (probe->stmt) + != GF_OMP_TARGET_KIND_OACC_PARALLEL) + && (gimple_omp_target_kind (probe->stmt) + != GF_OMP_TARGET_KIND_OACC_SERIAL)) goto do_lookup; cls = gimple_omp_target_clauses (probe->stmt); @@ -7518,7 +7526,7 @@ lower_oacc_head_mark (location_t loc, tree ddvar, tree clauses, /* In a parallel region, loops are implicitly INDEPENDENT. */ omp_context *tgt = enclosing_target_ctx (ctx); - if (!tgt || is_oacc_parallel (tgt)) + if (!tgt || is_oacc_parallel_or_serial (tgt)) tag |= OLF_INDEPENDENT; if (tag & OLF_TILE) @@ -11357,6 +11365,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) case GF_OMP_TARGET_KIND_EXIT_DATA: case GF_OMP_TARGET_KIND_OACC_PARALLEL: case GF_OMP_TARGET_KIND_OACC_KERNELS: + case GF_OMP_TARGET_KIND_OACC_SERIAL: case GF_OMP_TARGET_KIND_OACC_UPDATE: case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA: case GF_OMP_TARGET_KIND_OACC_DECLARE: @@ -11531,7 +11540,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) break; case OMP_CLAUSE_FIRSTPRIVATE: - if (is_oacc_parallel (ctx)) + if (is_oacc_parallel_or_serial (ctx)) goto oacc_firstprivate; map_cnt++; var = OMP_CLAUSE_DECL (c); @@ -11905,7 +11914,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) break; case OMP_CLAUSE_FIRSTPRIVATE: - if (is_oacc_parallel (ctx)) + if (is_oacc_parallel_or_serial (ctx)) goto oacc_firstprivate_map; ovar = OMP_CLAUSE_DECL (c); if (omp_is_reference (ovar)) @@ -12439,7 +12448,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) gimple_seq fork_seq = NULL; gimple_seq join_seq = NULL; - if (is_oacc_parallel (ctx)) + if (is_oacc_parallel_or_serial (ctx)) { /* If there are reductions on the offloaded region itself, treat them as a dummy GANG loop. */ diff --git a/gcc/testsuite/c-c++-common/goacc/serial-dims.c b/gcc/testsuite/c-c++-common/goacc/serial-dims.c new file mode 100644 index 000000000000..41698d279c98 --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/serial-dims.c @@ -0,0 +1,12 @@ +/* Invalid use of OpenACC parallelism dimensions clauses: num_gangs, + num_workers, vector_length with the serial construct. */ + +void f(void) +{ +#pragma acc serial num_gangs (1) /* { dg-error "'num_gangs' is not valid for '#pragma acc serial'" } */ + ; +#pragma acc serial num_workers (1) /* { dg-error "'num_workers' is not valid for '#pragma acc serial'" } */ + ; +#pragma acc serial vector_length (1) /* { dg-error "'vector_length' is not valid for '#pragma acc serial'" } */ + ; +} diff --git a/gcc/testsuite/gfortran.dg/goacc/serial-dims.f90 b/gcc/testsuite/gfortran.dg/goacc/serial-dims.f90 new file mode 100644 index 000000000000..72b4a8361776 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/serial-dims.f90 @@ -0,0 +1,40 @@ +! Invalid use of OpenACC parallelism dimensions clauses: num_gangs, +! num_workers, vector_length with the serial construct. + +subroutine s() + integer :: i + !$acc parallel + !$acc end parallel + + !$acc parallel loop + do i = 1, 5 + end do + + !$acc parallel loop + do i = 1, 5 + end do + !$acc end parallel loop + + !$acc serial loop + do i = 1, 5 + end do + + !$acc serial loop + do i = 1, 5 + end do + !$acc end serial loop + + !$acc serial + !$acc end serial +end subroutine s + +subroutine f() +!$acc serial num_gangs (1) ! { dg-error "Failed to match clause at" } +!$acc end serial ! { dg-error "Unexpected !.ACC END SERIAL statement" } + +!$acc serial num_workers (1) ! { dg-error "Failed to match clause at" } +!$acc end serial ! { dg-error "Unexpected !.ACC END SERIAL statement" } + +!$acc serial vector_length (1) ! { dg-error "Failed to match clause at" } +!$acc end serial ! { dg-error "Unexpected !.ACC END SERIAL statement" } +end subroutine f diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c index 53b3f55a3e6a..1cf7a9121336 100644 --- a/gcc/tree-pretty-print.c +++ b/gcc/tree-pretty-print.c @@ -3223,6 +3223,10 @@ dump_generic_node (pretty_printer *pp, tree node, int spc, dump_flags_t flags, pp_string (pp, "#pragma acc kernels"); goto dump_omp_clauses_body; + case OACC_SERIAL: + pp_string (pp, "#pragma acc serial"); + goto dump_omp_clauses_body; + case OACC_DATA: pp_string (pp, "#pragma acc data"); dump_omp_clauses (pp, OACC_DATA_CLAUSES (node), spc, flags); diff --git a/gcc/tree.def b/gcc/tree.def index fb6e7344fa6b..e8bb4f37f802 100644 --- a/gcc/tree.def +++ b/gcc/tree.def @@ -1095,6 +1095,12 @@ DEFTREECODE (OACC_PARALLEL, "oacc_parallel", tcc_statement, 2) DEFTREECODE (OACC_KERNELS, "oacc_kernels", tcc_statement, 2) +/* OpenACC - #pragma acc serial [clause1 ... clauseN] + Operand 0: OMP_BODY: Code to be executed sequentially. + Operand 1: OMP_CLAUSES: List of clauses. */ + +DEFTREECODE (OACC_SERIAL, "oacc_serial", tcc_statement, 2) + /* OpenACC - #pragma acc data [clause1 ... clauseN] Operand 0: OACC_DATA_BODY: Data construct body. Operand 1: OACC_DATA_CLAUSES: List of clauses. */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/serial-dims.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/serial-dims.c new file mode 100644 index 000000000000..bb91c9221f89 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/serial-dims.c @@ -0,0 +1,92 @@ +/* OpenACC dimensions with the serial construct. */ + +#include +#include +#include + +/* TODO: "(int) acc_device_*" casts because of the C++ acc_on_device wrapper + not behaving as expected for -O0. */ +#pragma acc routine seq +static unsigned int __attribute__ ((optimize ("O2"))) acc_gang () +{ + if (acc_on_device ((int) acc_device_host)) + return 0; + else if (acc_on_device ((int) acc_device_nvidia)) + return __builtin_goacc_parlevel_id (GOMP_DIM_GANG); + else + __builtin_abort (); +} + +#pragma acc routine seq +static unsigned int __attribute__ ((optimize ("O2"))) acc_worker () +{ + if (acc_on_device ((int) acc_device_host)) + return 0; + else if (acc_on_device ((int) acc_device_nvidia)) + return __builtin_goacc_parlevel_id (GOMP_DIM_WORKER); + else + __builtin_abort (); +} + +#pragma acc routine seq +static unsigned int __attribute__ ((optimize ("O2"))) acc_vector () +{ + if (acc_on_device ((int) acc_device_host)) + return 0; + else if (acc_on_device ((int) acc_device_nvidia)) + return __builtin_goacc_parlevel_id (GOMP_DIM_VECTOR); + else + __builtin_abort (); +} + + +int main () +{ + acc_init (acc_device_default); + + /* Serial OpenACC constructs must get launched as 1 x 1 x 1. */ + { + int gangs_min, gangs_max; + int workers_min, workers_max; + int vectors_min, vectors_max; + int gangs_actual, workers_actual, vectors_actual; + int i, j, k; + + gangs_min = workers_min = vectors_min = INT_MAX; + gangs_max = workers_max = vectors_max = INT_MIN; + gangs_actual = workers_actual = vectors_actual = 1; +#pragma acc serial /* { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } } */ + { + if (acc_on_device (acc_device_nvidia)) + { + /* The GCC nvptx back end enforces vector_length (32). */ + vectors_actual = 32; + } + else if (!acc_on_device (acc_device_host)) + __builtin_abort (); +#pragma acc loop gang \ + reduction (min: gangs_min, workers_min, vectors_min) \ + reduction (max: gangs_max, workers_max, vectors_max) + for (i = 100 * gangs_actual; i > -100 * gangs_actual; i--) +#pragma acc loop worker \ + reduction (min: gangs_min, workers_min, vectors_min) \ + reduction (max: gangs_max, workers_max, vectors_max) + for (j = 100 * workers_actual; j > -100 * workers_actual; j--) +#pragma acc loop vector \ + reduction (min: gangs_min, workers_min, vectors_min) \ + reduction (max: gangs_max, workers_max, vectors_max) + for (k = 100 * vectors_actual; k > -100 * vectors_actual; k--) + { + gangs_min = gangs_max = acc_gang (); + workers_min = workers_max = acc_worker (); + vectors_min = vectors_max = acc_vector (); + } + if (gangs_min != 0 || gangs_max != gangs_actual - 1 + || workers_min != 0 || workers_max != workers_actual - 1 + || vectors_min != 0 || vectors_max != vectors_actual - 1) + __builtin_abort (); + } + } + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-fortran/serial-dims-aux.c b/libgomp/testsuite/libgomp.oacc-fortran/serial-dims-aux.c new file mode 100644 index 000000000000..45c260510c29 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/serial-dims-aux.c @@ -0,0 +1,41 @@ +/* OpenACC dimensions with the serial construct. */ +/* Used by serial-dims.f90. */ + +#include +#include +#include + +/* TODO: "(int) acc_device_*" casts because of the C++ acc_on_device wrapper + not behaving as expected for -O0. */ +#pragma acc routine seq +static unsigned int __attribute__ ((optimize ("O2"))) acc_gang () +{ + if (acc_on_device ((int) acc_device_host)) + return 0; + else if (acc_on_device ((int) acc_device_nvidia)) + return __builtin_goacc_parlevel_id (GOMP_DIM_GANG); + else + __builtin_abort (); +} + +#pragma acc routine seq +static unsigned int __attribute__ ((optimize ("O2"))) acc_worker () +{ + if (acc_on_device ((int) acc_device_host)) + return 0; + else if (acc_on_device ((int) acc_device_nvidia)) + return __builtin_goacc_parlevel_id (GOMP_DIM_WORKER); + else + __builtin_abort (); +} + +#pragma acc routine seq +static unsigned int __attribute__ ((optimize ("O2"))) acc_vector () +{ + if (acc_on_device ((int) acc_device_host)) + return 0; + else if (acc_on_device ((int) acc_device_nvidia)) + return __builtin_goacc_parlevel_id (GOMP_DIM_VECTOR); + else + __builtin_abort (); +} diff --git a/libgomp/testsuite/libgomp.oacc-fortran/serial-dims.f90 b/libgomp/testsuite/libgomp.oacc-fortran/serial-dims.f90 new file mode 100644 index 000000000000..25c933629045 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/serial-dims.f90 @@ -0,0 +1,89 @@ +! OpenACC dimensions with the serial construct. + +! { dg-additional-sources serial-dims-aux.c } +! { dg-warning "command line option '-fintrinsic-modules-path=.*' is valid for Fortran but not for C" } + +module acc_routines + implicit none (type, external) + + interface + integer function acc_gang() bind(C) + !$acc routine seq + end function acc_gang + + integer function acc_worker() bind(C) + !$acc routine seq + end function acc_worker + + integer function acc_vector() bind(C) + !$acc routine seq + end function acc_vector + end interface +end module acc_routines + +program main + use iso_c_binding + use openacc + use acc_routines + implicit none (type, external) + + integer :: gangs_min, gangs_max + integer :: workers_min, workers_max + integer :: vectors_min, vectors_max + integer :: gangs_actual, workers_actual, vectors_actual + integer :: i, j, k + + call acc_init (acc_device_default) + + ! Serial OpenACC constructs must get launched as 1 x 1 x 1. + gangs_min = huge(gangs_min) + workers_min = huge(workers_min) + vectors_min = huge(vectors_min) + gangs_max = -huge(gangs_max) - 1 ! INT_MIN + workers_max = -huge(gangs_max) - 1 + vectors_max = -huge(gangs_max) - 1 + gangs_actual = 1 + workers_actual = 1 + vectors_actual = 1 + + !$acc serial ! { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } } + if (acc_on_device (acc_device_nvidia)) then + ! The GCC nvptx back end enforces vector_length (32). + vectors_actual = 32 + elseif (acc_on_device (acc_device_gcn)) then + ! AMD GCN relies on the autovectorizer for the vector dimension: + ! the loop below isn't likely to be vectorized, so vectors_actual + ! is effectively 1. + vectors_actual = 1 + elseif (.not. acc_on_device (acc_device_host)) then + stop 1 + end if + +!$acc loop gang & +!$acc & reduction (min: gangs_min, workers_min, vectors_min) & +!$acc & reduction (max: gangs_max, workers_max, vectors_max) + do i = 100 * gangs_actual, -99 * gangs_actual, -1 +!$acc loop worker & +!$acc & reduction (min: gangs_min, workers_min, vectors_min) & +!$acc & reduction (max: gangs_max, workers_max, vectors_max) + do j = 100 * workers_actual, -99 * workers_actual, -1 +!$acc loop vector & +!$acc & reduction (min: gangs_min, workers_min, vectors_min) & +!$acc & reduction (max: gangs_max, workers_max, vectors_max) + do k = 100 * vectors_actual, -99 * vectors_actual, -1 + gangs_min = acc_gang (); + gangs_max = acc_gang (); + workers_min = acc_worker (); + workers_max = acc_worker (); + vectors_min = acc_vector (); + vectors_max = acc_vector (); + end do + end do + end do + if (gangs_min /= 0 .or. gangs_max /= gangs_actual - 1 & + .or. workers_min /= 0 .or. workers_max /= workers_actual - 1 & + .or. vectors_min /= 0 .or. vectors_max /= vectors_actual - 1) & + stop 2 +!$acc end serial + +end program main