From patchwork Fri Oct 4 14:56:01 2024 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Kwok Cheung Yeung X-Patchwork-Id: 1992779 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@legolas.ozlabs.org Authentication-Results: legolas.ozlabs.org; dkim=pass (2048-bit key; unprotected) header.d=baylibre-com.20230601.gappssmtp.com header.i=@baylibre-com.20230601.gappssmtp.com header.a=rsa-sha256 header.s=20230601 header.b=Un0fCwJ5; dkim-atps=neutral Authentication-Results: legolas.ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=8.43.85.97; helo=server2.sourceware.org; envelope-from=gcc-patches-bounces~incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=patchwork.ozlabs.org) Received: from server2.sourceware.org (server2.sourceware.org [8.43.85.97]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature ECDSA (secp384r1) server-digest SHA384) (No client certificate requested) by legolas.ozlabs.org (Postfix) with ESMTPS id 4XKs8Q2vDkz1xt7 for ; Sat, 5 Oct 2024 00:57:10 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 971E538432ED for ; Fri, 4 Oct 2024 14:57:08 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from mail-ed1-x536.google.com (mail-ed1-x536.google.com [IPv6:2a00:1450:4864:20::536]) by sourceware.org (Postfix) with ESMTPS id 1BBBD385E011 for ; Fri, 4 Oct 2024 14:56:30 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 1BBBD385E011 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=baylibre.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=baylibre.com ARC-Filter: OpenARC Filter v1.0.0 sourceware.org 1BBBD385E011 Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=2a00:1450:4864:20::536 ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1728053796; cv=none; b=EoHVJmyF6ovJ+tRB80uH+6vbc9N98mIcVsCDuM6VAlh0OahpP1Ts89TGqtOM9ifVCbwXFpnGghuKsfDmKskauw4cscwSzY9xeTjpnktkmnxi2S5n8U7UtNmxsCuR99O1n1UMTD5EDosGMhAGojSY04qIxeCLwlblobLNFsZN4NU= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1728053796; c=relaxed/simple; bh=2hz5f0lZQdtjVkAy0812rYNF0zGjcmXmmBn9/RNmD3g=; h=DKIM-Signature:Message-ID:Date:MIME-Version:To:Subject:From; b=ckk5y0WUi2+eqifliWa/8CKtLNcwAdfl/kaY0FR4OqP80gWP9CaNVnrum2Heq4OndhLY+nKFeCUZADe9SAdhsnenD9WN4KXaZtqvAyXpmkmPxpyFIAnS3+2IIB5tn70L3pRP3jzszcYhzaQkb+gYEMQYhoLajCxHMaaEH6KpRDo= ARC-Authentication-Results: i=1; server2.sourceware.org Received: by mail-ed1-x536.google.com with SMTP id 4fb4d7f45d1cf-5c5b9d2195eso2993155a12.1 for ; Fri, 04 Oct 2024 07:56:30 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre-com.20230601.gappssmtp.com; s=20230601; t=1728053788; x=1728658588; darn=gcc.gnu.org; h=in-reply-to:from:content-language:subject:references:to:user-agent :mime-version:date:message-id:from:to:cc:subject:date:message-id :reply-to; bh=UTaJgqmXWq6YUKhFg0fhJKu6JPci6SR/kHhYDe64zF0=; b=Un0fCwJ582R6e99/KbbYAiuqbEqwZedNE41AGQQS8ALGTigVVeFTHEQaBlOwz2rVI0 tS5RBy95SOCAc5dHbaGrf+PNeCutnJCVdE485dbUgM1/OGVXnOMUil7AgKnbNMJqCYjN mLFf0NwAwRQPqSu6PM8zY5LuhvtJlKO+jOD6O6gKwFVzYburiYQkyg65dm6vSDTOrqk0 DH1AUavAwn0fh1I4OzZu6ZcBzk/9erBeKPN8O2aaOsjkwSmkCumk38E+lVSCvqSrXjRw slOMtceJFhtF/KSg4MTk939v550xkkLDq6I77EAZWF94nE08yjOXhPqRAWa8JNtEl3BA i/5Q== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1728053788; x=1728658588; h=in-reply-to:from:content-language:subject:references:to:user-agent :mime-version:date:message-id:x-gm-message-state:from:to:cc:subject :date:message-id:reply-to; bh=UTaJgqmXWq6YUKhFg0fhJKu6JPci6SR/kHhYDe64zF0=; b=UB2U5iQkyFtxPXvavNnv4tv5+E2JinPyQ69UuUU++N2zmGm4RVansxZJpWjx4VL1PI tHc+40AnJkGNfZ5gONHeHaTB5YCv6OF8nDUFL0k5JgjktQQAXG78d7ED1rD+9TF2h+VN en5Cq5Mck2PW1BtV2YpRRIeMVGu1Uyk7bp+/hZRjL8wi4X9varRp7ZpSgqrxlrDZ0uAT FQkAfsxhXxGbyFZ790BewJ5CQA5k7Lfe/vbQfbavwyH4TJOPlOk3dXxuM/nkQ0h4PGnZ bG97bNpCOZOAF9pOj5g5cVWnF7XwK+lp8z98aoS4chPBGL/BTqsxVl4O9HiF0SG7qvwD 2ppg== X-Gm-Message-State: AOJu0YwZp8tf1ud+DHDwaUr0Ortv/jP2ZQNF2BvFaRdwHDDqVYljFqOi Wd3w6TmNjmWYqD8yAV1p3vM1N/xBwEXAgRXiD9Pa2dibR8k/CKEKD3Z618oYg/dCRUrc9kBlTp2 5 X-Google-Smtp-Source: AGHT+IHWPo8jowzuWfpUoR58x7yPXZhU5PWbCqrpgxaDplLvQuNlSMDVoE2EAQzKXOBCU/1fA6Kl4A== X-Received: by 2002:a17:907:2cc7:b0:a90:13b6:3ece with SMTP id a640c23a62f3a-a991bd3fb6emr291721266b.15.1728053788341; Fri, 04 Oct 2024 07:56:28 -0700 (PDT) Received: from ?IPV6:2a00:23c6:88fe:9301:2d7d:f734:bc6:c47b? ([2a00:23c6:88fe:9301:2d7d:f734:bc6:c47b]) by smtp.gmail.com with ESMTPSA id a640c23a62f3a-a992e784a4bsm1055166b.114.2024.10.04.07.56.27 (version=TLS1_3 cipher=TLS_AES_128_GCM_SHA256 bits=128/128); Fri, 04 Oct 2024 07:56:28 -0700 (PDT) Message-ID: Date: Fri, 4 Oct 2024 15:56:01 +0100 MIME-Version: 1.0 User-Agent: Mozilla Thunderbird To: gcc-patches , Jakub Jelinek , Tobias Burnus References: <6b94b8ed-020b-47e2-b02a-4891891f2847@baylibre.com> Subject: [PATCH v3 2/5] openmp: Add support for iterators in map clauses (C/C++) Content-Language: en-GB From: Kwok Cheung Yeung In-Reply-To: <6b94b8ed-020b-47e2-b02a-4891891f2847@baylibre.com> X-Spam-Status: No, score=-12.5 required=5.0 tests=BAYES_00, DKIM_SIGNED, DKIM_VALID, GIT_PATCH_0, RCVD_IN_DNSWL_NONE, SPF_HELO_NONE, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.30 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces~incoming=patchwork.ozlabs.org@gcc.gnu.org This patch modifies the C and C++ parsers to accept an iterator as a map type modifier, storing it in the OMP_CLAUSE_ITERATOR argument of the clause. When finishing clauses, any clauses generated from a clause with iterators also has the iterator applied to them. During gimplification, check_omp_map_iterators is called to check that all iterator variables are referenced at some point with a clause. Gimplification of the clause decl and size are delayed until iterator expansion as they may reference iterator variables. In lower_target, lower_omp_map_iterators is called to construct the expansion loop for iterator clauses. Clauses using the same set of iterators reuse the loop, though with different storage allocated for them. lower_omp_map_iterator_expr is called to add the final expression that is sent as the hostaddr for libgomp to the loop, and a reference to the array generated by the iterator loop is returned to replace the original expression. lower_omp_map_iterator_size works similarly for the clause size. finish_omp_map_iterators is called later to finalise the loop. Libgomp has a new function gomp_merge_iterator_maps which identifies data coming from an iterator, and effectively creates new maps on-the-fly from the iterator info array, inserting them into the list of mappings at the point where iterator data occurred. As there are now multiple maps where one was previously, an entry is only added to the target vars for the first expanded map, otherwise it will get out of sync with the expected layout and the wrong variables will be picked up by the target function. From 50557e513ca534ba32f50d99991b056a07a6f671 Mon Sep 17 00:00:00 2001 From: Kwok Cheung Yeung Date: Fri, 4 Oct 2024 15:16:12 +0100 Subject: [PATCH 2/5] openmp: Add support for iterators in map clauses (C/C++) This adds preliminary support for iterators in map clauses within OpenMP 'target' constructs (which includes constructs such as 'target enter data'). Iterators with non-constant loop bounds are not currently supported. 2024-10-04 Kwok Cheung Yeung gcc/c/ * c-parser.cc (c_parser_omp_clause_map): Parse 'iterator' modifier. * c-typeck.cc (c_finish_omp_clauses): Finish iterators. Apply iterators to generated clauses. gcc/cp/ * parser.cc (cp_parser_omp_clause_map): Parse 'iterator' modifier. * semantics.cc (finish_omp_clauses): Finish iterators. Apply iterators to generated clauses. gcc/ * gimplify.cc (compute_iterator_count): Make non-static. Take an iterator instead of a clause for an operand. (build_iterator_loop): Likewise. (gimplify_omp_depend): Pass iterator in call to compute_iterator_count and build_iterator_loop. (find_var_decl): New. (check_omp_map_iterators): New. (gimplify_scan_omp_clauses): Call check_omp_map_iterators on clauses with iterators. (gimplify_adjust_omp_clauses): Skip gimplification of clause decl and size for clauses with iterators. * omp-low.cc (struct iterator_loop_info_t): New type. (iterator_loop_map_t): New type. (lower_omp_map_iterators): New. (lower_omp_map_iterator_expr): New. (lower_omp_map_iterator_size): New. (finish_omp_map_iterators): New. (lower_omp_target): Call lower_omp_map_iterators on clauses with iterators. Call lower_omp_map_iterator_expr before assigning to sender ref. Call lower_omp_map_iterator_size before setting the size. Call finish_omp_map_iterators. Insert statements generated during iterator expansion before the statements for the target clause. * tree-pretty-print.cc (dump_omp_clause): Call dump_omp_iterators for iterators in map clauses. * tree.cc (omp_clause_num_ops): Add operand for OMP_CLAUSE_MAP. (walk_tree_1): Do not walk last operand of OMP_CLAUSE_MAP. * tree.h (OMP_CLAUSE_HAS_ITERATORS): New. (OMP_CLAUSE_ITERATORS: New. gcc/testsuite/ * c-c++-common/gomp/map-6.c (foo): Amend expected error message. * c-c++-common/gomp/target-map-iterators-1.c: New. * c-c++-common/gomp/target-map-iterators-2.c: New. * c-c++-common/gomp/target-map-iterators-3.c: New. libgomp/ * target.c (kind_to_name): New. (gomp_merge_iterator_maps): New. (gomp_map_vars_internal): Call gomp_merge_iterator_maps. Copy address of only the first iteration to target vars. Free allocated variables. * testsuite/libgomp.c-c++-common/target-map-iterators-1.c: New. * testsuite/libgomp.c-c++-common/target-map-iterators-2.c: New. * testsuite/libgomp.c-c++-common/target-map-iterators-3.c: New. --- gcc/c/c-parser.cc | 59 +++++- gcc/c/c-typeck.cc | 22 ++- gcc/cp/parser.cc | 62 +++++- gcc/cp/semantics.cc | 22 ++- gcc/gimplify.cc | 88 +++++++-- gcc/omp-low.cc | 186 +++++++++++++++++- gcc/testsuite/c-c++-common/gomp/map-6.c | 10 +- .../gomp/target-map-iterators-1.c | 23 +++ .../gomp/target-map-iterators-2.c | 19 ++ .../gomp/target-map-iterators-3.c | 23 +++ gcc/tree-pretty-print.cc | 5 + gcc/tree.cc | 5 +- gcc/tree.h | 7 + libgomp/target.c | 130 +++++++++++- .../target-map-iterators-1.c | 47 +++++ .../target-map-iterators-2.c | 44 +++++ .../target-map-iterators-3.c | 56 ++++++ 17 files changed, 759 insertions(+), 49 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/gomp/target-map-iterators-1.c create mode 100644 gcc/testsuite/c-c++-common/gomp/target-map-iterators-2.c create mode 100644 gcc/testsuite/c-c++-common/gomp/target-map-iterators-3.c create mode 100644 libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-1.c create mode 100644 libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-2.c create mode 100644 libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-3.c diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc index a681438cbbe..184fc076388 100644 --- a/gcc/c/c-parser.cc +++ b/gcc/c/c-parser.cc @@ -18892,7 +18892,7 @@ c_parser_omp_clause_doacross (c_parser *parser, tree list) map ( [map-type-modifier[,] ...] map-kind: variable-list ) map-type-modifier: - always | close */ + always | close | present | iterator (iterators-definition) */ static tree c_parser_omp_clause_map (c_parser *parser, tree list) @@ -18907,15 +18907,35 @@ c_parser_omp_clause_map (c_parser *parser, tree list) int pos = 1; int map_kind_pos = 0; - while (c_parser_peek_nth_token_raw (parser, pos)->type == CPP_NAME) + int iterator_length = 0; + for (;;) { - if (c_parser_peek_nth_token_raw (parser, pos + 1)->type == CPP_COLON) + c_token *tok = c_parser_peek_nth_token_raw (parser, pos); + if (tok->type != CPP_NAME) + break; + + const char *p = IDENTIFIER_POINTER (tok->value); + c_token *next_tok = c_parser_peek_nth_token_raw (parser, pos + 1); + if (strcmp (p, "iterator") == 0 && next_tok->type == CPP_OPEN_PAREN) + { + unsigned n = pos + 2; + if (c_parser_check_balanced_raw_token_sequence (parser, &n) + && c_parser_peek_nth_token_raw (parser, n)->type + == CPP_CLOSE_PAREN) + { + iterator_length = n - pos + 1; + pos = n; + next_tok = c_parser_peek_nth_token_raw (parser, pos + 1); + } + } + + if (next_tok->type == CPP_COLON) { map_kind_pos = pos; break; } - if (c_parser_peek_nth_token_raw (parser, pos + 1)->type == CPP_COMMA) + if (next_tok->type == CPP_COMMA) pos++; pos++; } @@ -18923,6 +18943,7 @@ c_parser_omp_clause_map (c_parser *parser, tree list) int always_modifier = 0; int close_modifier = 0; int present_modifier = 0; + tree iterators = NULL_TREE; for (int pos = 1; pos < map_kind_pos; ++pos) { c_token *tok = c_parser_peek_token (parser); @@ -18964,10 +18985,24 @@ c_parser_omp_clause_map (c_parser *parser, tree list) } present_modifier++; } + else if (strcmp ("iterator", p) == 0 + && c_parser_peek_2nd_token (parser)->type == CPP_OPEN_PAREN) + { + if (iterators) + { + c_parser_error (parser, "too many % modifiers"); + parens.skip_until_found_close (parser); + return list; + } + iterators = c_parser_omp_iterators (parser); + pos += iterator_length - 1; + continue; + } else { c_parser_error (parser, "% clause with map-type modifier other " - "than %, % or %"); + "than %, %, % " + "or %"); parens.skip_until_found_close (parser); return list; } @@ -19016,8 +19051,20 @@ c_parser_omp_clause_map (c_parser *parser, tree list) nl = c_parser_omp_variable_list (parser, clause_loc, OMP_CLAUSE_MAP, list, true); + if (iterators) + { + tree block = pop_scope (); + if (iterators == error_mark_node) + iterators = NULL_TREE; + else + TREE_VEC_ELT (iterators, 5) = block; + } + for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c)) - OMP_CLAUSE_SET_MAP_KIND (c, kind); + { + OMP_CLAUSE_SET_MAP_KIND (c, kind); + OMP_CLAUSE_ITERATORS (c) = iterators; + } parens.skip_until_found_close (parser); return nl; diff --git a/gcc/c/c-typeck.cc b/gcc/c/c-typeck.cc index 30a03f071d8..cca9f1c000c 100644 --- a/gcc/c/c-typeck.cc +++ b/gcc/c/c-typeck.cc @@ -15058,7 +15058,15 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) /* We've reached the end of a list of expanded nodes. Reset the group start pointer. */ if (c == grp_sentinel) - grp_start_p = NULL; + { + if (grp_start_p + && OMP_CLAUSE_HAS_ITERATORS (*grp_start_p) + && OMP_CLAUSE_ITERATORS (*grp_start_p)) + for (tree gc = *grp_start_p; gc != grp_sentinel; + gc = OMP_CLAUSE_CHAIN (gc)) + OMP_CLAUSE_ITERATORS (gc) = OMP_CLAUSE_ITERATORS (*grp_start_p); + grp_start_p = NULL; + } switch (OMP_CLAUSE_CODE (c)) { @@ -15805,6 +15813,12 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) case OMP_CLAUSE_MAP: if (OMP_CLAUSE_MAP_IMPLICIT (c) && !implicit_moved) goto move_implicit; + if (OMP_CLAUSE_ITERATORS (c) + && c_omp_finish_iterators (OMP_CLAUSE_ITERATORS (c))) + { + t = error_mark_node; + break; + } /* FALLTHRU */ case OMP_CLAUSE_TO: case OMP_CLAUSE_FROM: @@ -16497,6 +16511,12 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) pc = &OMP_CLAUSE_CHAIN (c); } + if (grp_start_p + && OMP_CLAUSE_HAS_ITERATORS (*grp_start_p) + && OMP_CLAUSE_ITERATORS (*grp_start_p)) + for (tree gc = *grp_start_p; gc; gc = OMP_CLAUSE_CHAIN (gc)) + OMP_CLAUSE_ITERATORS (gc) = OMP_CLAUSE_ITERATORS (*grp_start_p); + if (simdlen && safelen && tree_int_cst_lt (OMP_CLAUSE_SAFELEN_EXPR (safelen), diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc index f50534f5f39..79d6e115d16 100644 --- a/gcc/cp/parser.cc +++ b/gcc/cp/parser.cc @@ -41776,16 +41776,34 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list) int pos = 1; int map_kind_pos = 0; - while (cp_lexer_peek_nth_token (parser->lexer, pos)->type == CPP_NAME - || cp_lexer_peek_nth_token (parser->lexer, pos)->keyword == RID_DELETE) + int iterator_length = 0; + for (;;) { - if (cp_lexer_peek_nth_token (parser->lexer, pos + 1)->type == CPP_COLON) + cp_token *tok = cp_lexer_peek_nth_token (parser->lexer, pos); + if (!(tok->type == CPP_NAME || tok->keyword == RID_DELETE)) + break; + + cp_token *next_tok = cp_lexer_peek_nth_token (parser->lexer, pos + 1); + if (tok->type == CPP_NAME + && strcmp (IDENTIFIER_POINTER (tok->u.value), "iterator") == 0 + && next_tok->type == CPP_OPEN_PAREN) + { + int n = cp_parser_skip_balanced_tokens (parser, pos + 1); + if (n != pos + 1) + { + iterator_length = n - pos; + pos = n - 1; + next_tok = cp_lexer_peek_nth_token (parser->lexer, n); + } + } + + if (next_tok->type == CPP_COLON) { map_kind_pos = pos; break; } - if (cp_lexer_peek_nth_token (parser->lexer, pos + 1)->type == CPP_COMMA) + if (next_tok->type == CPP_COMMA) pos++; pos++; } @@ -41793,6 +41811,7 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list) bool always_modifier = false; bool close_modifier = false; bool present_modifier = false; + tree iterators = NULL_TREE; for (int pos = 1; pos < map_kind_pos; ++pos) { cp_token *tok = cp_lexer_peek_token (parser->lexer); @@ -41842,10 +41861,29 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list) } present_modifier = true; } + else if (strcmp ("iterator", p) == 0 + && cp_lexer_peek_nth_token (parser->lexer, 2)->type + == CPP_OPEN_PAREN) + { + if (iterators) + { + cp_parser_error (parser, "too many % modifiers"); + cp_parser_skip_to_closing_parenthesis (parser, + /*recovering=*/true, + /*or_comma=*/false, + /*consume_paren=*/true); + return list; + } + begin_scope (sk_omp, NULL); + iterators = cp_parser_omp_iterators (parser); + pos += iterator_length - 1; + continue; + } else { cp_parser_error (parser, "% clause with map-type modifier other" - " than %, % or %"); + " than %, %, %" + " or %"); cp_parser_skip_to_closing_parenthesis (parser, /*recovering=*/true, /*or_comma=*/false, @@ -41909,8 +41947,20 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list) NULL, true); finish_scope (); + if (iterators) + { + tree block = poplevel (1, 1, 0); + if (iterators == error_mark_node) + iterators = NULL_TREE; + else + TREE_VEC_ELT (iterators, 5) = block; + } + for (c = nlist; c != list; c = OMP_CLAUSE_CHAIN (c)) - OMP_CLAUSE_SET_MAP_KIND (c, kind); + { + OMP_CLAUSE_SET_MAP_KIND (c, kind); + OMP_CLAUSE_ITERATORS (c) = iterators; + } return nlist; } diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc index 4f856a9d749..ba5657f7bc2 100644 --- a/gcc/cp/semantics.cc +++ b/gcc/cp/semantics.cc @@ -7263,7 +7263,15 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) /* We've reached the end of a list of expanded nodes. Reset the group start pointer. */ if (c == grp_sentinel) - grp_start_p = NULL; + { + if (grp_start_p + && OMP_CLAUSE_HAS_ITERATORS (*grp_start_p) + && OMP_CLAUSE_ITERATORS (*grp_start_p)) + for (tree gc = *grp_start_p; gc != grp_sentinel; + gc = OMP_CLAUSE_CHAIN (gc)) + OMP_CLAUSE_ITERATORS (gc) = OMP_CLAUSE_ITERATORS (*grp_start_p); + grp_start_p = NULL; + } switch (OMP_CLAUSE_CODE (c)) { @@ -8484,6 +8492,12 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) case OMP_CLAUSE_MAP: if (OMP_CLAUSE_MAP_IMPLICIT (c) && !implicit_moved) goto move_implicit; + if (OMP_CLAUSE_ITERATORS (c) + && cp_omp_finish_iterators (OMP_CLAUSE_ITERATORS (c))) + { + t = error_mark_node; + break; + } /* FALLTHRU */ case OMP_CLAUSE_TO: case OMP_CLAUSE_FROM: @@ -9348,6 +9362,12 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) pc = &OMP_CLAUSE_CHAIN (c); } + if (grp_start_p + && OMP_CLAUSE_HAS_ITERATORS (*grp_start_p) + && OMP_CLAUSE_ITERATORS (*grp_start_p)) + for (tree gc = *grp_start_p; gc; gc = OMP_CLAUSE_CHAIN (gc)) + OMP_CLAUSE_ITERATORS (gc) = OMP_CLAUSE_ITERATORS (*grp_start_p); + if (reduction_seen < 0 && (ordered_seen || schedule_seen)) reduction_seen = -2; diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index b5b1f83db8f..6e532d07fcf 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -8823,13 +8823,13 @@ gimplify_omp_affinity (tree *list_p, gimple_seq *pre_p) } /* Returns a tree expression containing the total iteration count of the - iterator clause decl T. */ + iterator IT. */ -static tree -compute_iterator_count (tree t, gimple_seq *pre_p) +tree +compute_iterator_count (tree it, gimple_seq *pre_p) { tree tcnt = size_one_node; - for (tree it = TREE_PURPOSE (t); it; it = TREE_CHAIN (it)) + for (; it; it = TREE_CHAIN (it)) { if (gimplify_expr (&TREE_VEC_ELT (it, 1), pre_p, NULL, is_gimple_val, fb_rvalue) == GS_ERROR @@ -8899,21 +8899,17 @@ compute_iterator_count (tree t, gimple_seq *pre_p) Returns a pointer to the BIND_EXPR_BODY in the innermost loop body. LAST_BIND is set to point to the BIND_EXPR containing the whole loop. */ -static tree * -build_iterator_loop (tree c, gimple_seq *pre_p, tree *last_bind) +tree * +build_iterator_loop (tree it, gimple_seq *pre_p, tree *last_bind) { - tree t = OMP_CLAUSE_DECL (c); - gcc_assert (OMP_ITERATOR_DECL_P (t)); - if (*last_bind) gimplify_and_add (*last_bind, pre_p); - tree block = TREE_VEC_ELT (TREE_PURPOSE (t), 5); + tree block = TREE_VEC_ELT (it, 5); *last_bind = build3 (BIND_EXPR, void_type_node, BLOCK_VARS (block), NULL, block); TREE_SIDE_EFFECTS (*last_bind) = 1; - SET_EXPR_LOCATION (*last_bind, OMP_CLAUSE_LOCATION (c)); tree *p = &BIND_EXPR_BODY (*last_bind); - for (tree it = TREE_PURPOSE (t); it; it = TREE_CHAIN (it)) + for (; it; it = TREE_CHAIN (it)) { tree var = TREE_VEC_ELT (it, 0); tree begin = TREE_VEC_ELT (it, 1); @@ -9023,7 +9019,7 @@ gimplify_omp_depend (tree *list_p, gimple_seq *pre_p) { if (TREE_PURPOSE (t) != last_iter) { - tree tcnt = compute_iterator_count (t, pre_p); + tree tcnt = compute_iterator_count (TREE_PURPOSE (t), pre_p); if (!tcnt) return 2; last_iter = TREE_PURPOSE (t); @@ -9181,7 +9177,9 @@ gimplify_omp_depend (tree *list_p, gimple_seq *pre_p) if (OMP_ITERATOR_DECL_P (t)) { if (TREE_PURPOSE (t) != last_iter) - last_body = build_iterator_loop (c, pre_p, &last_bind); + last_body = build_iterator_loop (TREE_PURPOSE (t), pre_p, + &last_bind); + SET_EXPR_LOCATION (last_bind, OMP_CLAUSE_LOCATION (c)); last_iter = TREE_PURPOSE (t); if (TREE_CODE (TREE_VALUE (t)) == COMPOUND_EXPR) { @@ -12078,6 +12076,51 @@ error_out: return success; } +/* Callback for walk_tree to find a VAR_DECL (stored in DATA) in the + tree TP. */ + +static tree +find_var_decl (tree *tp, int *, void *data) +{ + tree t = *tp; + + if (TREE_CODE (t) == VAR_DECL && t == (tree) data) + return t; + + return NULL_TREE; +} + +/* Check that the clause C uses all the iterator variables. + Return TRUE if there are no errors. */ + +static bool +check_omp_map_iterators (tree c) +{ + bool error = false; + gcc_assert (OMP_CLAUSE_ITERATORS (c)); + + /* Do not check internal map kinds. */ + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH) + return true; + + for (tree it = OMP_CLAUSE_ITERATORS (c); it; it = TREE_CHAIN (it)) + { + tree var = TREE_VEC_ELT (it, 0); + tree t = walk_tree (&OMP_CLAUSE_DECL (c), find_var_decl, var, NULL); + if (t == NULL_TREE) + t = walk_tree (&OMP_CLAUSE_SIZE (c), find_var_decl, var, NULL); + if (t == NULL_TREE) + { + error_at (OMP_CLAUSE_LOCATION (c), + "iterator variable %qD not used in clause expression", + var); + error = true; + } + } + return !error; +} + /* Scan the OMP clauses in *LIST_P, installing mappings into a new and previous omp contexts. */ @@ -12478,6 +12521,12 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, break; } + if (OMP_CLAUSE_ITERATORS (c) && !check_omp_map_iterators (c)) + { + remove = true; + break; + } + if (!omp_parse_expr (addr_tokens, decl)) { remove = true; @@ -14168,7 +14217,11 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, : TYPE_SIZE_UNIT (TREE_TYPE (decl)); } gimplify_omp_ctxp = ctx->outer_context; - if (gimplify_expr (&OMP_CLAUSE_SIZE (c), pre_p, NULL, + if (OMP_CLAUSE_ITERATORS (c)) + /* Gimplify the OMP_CLAUSE_SIZE later, when the iterator is + gimplified. */ + ; + else if (gimplify_expr (&OMP_CLAUSE_SIZE (c), pre_p, NULL, is_gimple_val, fb_rvalue) == GS_ERROR) { gimplify_omp_ctxp = ctx; @@ -14333,6 +14386,11 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, if (code == OMP_TARGET && OMP_CLAUSE_MAP_IN_REDUCTION (c)) break; + /* Do not gimplify the declaration yet for clauses with + iterators. */ + if (OMP_CLAUSE_ITERATORS (c)) + break; + gimplify_omp_ctxp = ctx->outer_context; if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue, fb_lvalue) == GS_ERROR) diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc index da2051b0279..9cf6e207d1c 100644 --- a/gcc/omp-low.cc +++ b/gcc/omp-low.cc @@ -12607,6 +12607,163 @@ lower_omp_taskreg (gimple_stmt_iterator *gsi_p, omp_context *ctx) } } +extern tree compute_iterator_count (tree it, gimple_seq *pre_p); +extern tree *build_iterator_loop (tree it, gimple_seq *pre_p, tree *last_bind); + +struct iterator_loop_info_t +{ + tree bind; + tree count; + tree index; + tree *body; + tree *iterator; + hash_map elems; +}; + +typedef hash_map iterator_loop_map_t; + +/* Builds a loop to expand any iterators in clause C, reusing any previously + built loops if they use the same set of iterators. Generated Gimple + statements are placed into PRE_P. Information on the loops is held in + LOOPS. finish_omp_map_iterators must be called before the loops are + used. */ + +static void +lower_omp_map_iterators (tree c, gimple_seq *pre_p, iterator_loop_map_t *loops) +{ + if (!OMP_CLAUSE_HAS_ITERATORS (c) || !OMP_CLAUSE_ITERATORS (c)) + return; + + bool built_p; + iterator_loop_info_t &loop = loops->get_or_insert (OMP_CLAUSE_ITERATORS (c), + &built_p); + if (!built_p) + { + loop.count = compute_iterator_count (OMP_CLAUSE_ITERATORS (c), pre_p); + if (!loop.count) + return; + + loop.body = build_iterator_loop (OMP_CLAUSE_ITERATORS (c), pre_p, + &loop.bind); + loop.index = create_tmp_var (sizetype); + SET_EXPR_LOCATION (loop.bind, OMP_CLAUSE_LOCATION (c)); + loop.iterator = &OMP_CLAUSE_ITERATORS (c); + + /* idx = -1; */ + /* This should be initialized to before the individual elements, + as idx is pre-incremented in the loop body. */ + gimple *g = gimple_build_assign (loop.index, size_int (-1)); + gimple_seq_add_stmt (pre_p, g); + + /* IN LOOP BODY: */ + /* idx += 2; */ + tree tem = build2_loc (OMP_CLAUSE_LOCATION (c), MODIFY_EXPR, + void_type_node, loop.index, + size_binop (PLUS_EXPR, loop.index, size_int (2))); + append_to_statement_list_force (tem, loop.body); + } + + /* Create array to hold expanded values. */ + tree last_count_2 = size_binop (MULT_EXPR, loop.count, size_int (2)); + tree arr_length = size_binop (PLUS_EXPR, last_count_2, size_int (1)); + tree elems = NULL_TREE; + if (TREE_CONSTANT (arr_length)) + { + tree type = build_array_type (ptr_type_node, + build_index_type (arr_length)); + elems = create_tmp_var_raw (type); + TREE_ADDRESSABLE (elems) = 1; + gimple_add_tmp_var (elems); + } + else + { + /* Handle dynamic sizes. */ + sorry ("dynamic iterator sizes not implemented yet"); + } + loop.elems.put (c, elems); + + /* elems[0] = count; */ + tree lhs = build4 (ARRAY_REF, ptr_type_node, elems, size_int (0), + NULL_TREE, NULL_TREE); + tree tem = build2_loc (OMP_CLAUSE_LOCATION (c), MODIFY_EXPR, + void_type_node, lhs, loop.count); + gimplify_and_add (tem, pre_p); +} + +/* Set EXPR as the hostaddr expression that should result from the clause C. + LOOPS holds the intermediate loop info. Returns the tree that should be + passed as the hostaddr. */ + +static tree +lower_omp_map_iterator_expr (tree expr, tree c, iterator_loop_map_t *loops) +{ + if (!OMP_CLAUSE_HAS_ITERATORS (c) || !OMP_CLAUSE_ITERATORS (c)) + return expr; + + iterator_loop_info_t *loop = loops->get (OMP_CLAUSE_ITERATORS (c)); + gcc_assert (loop); + tree *elems = loop->elems.get (c); + + /* IN LOOP BODY: */ + /* elems[idx] = ; */ + tree lhs = build4 (ARRAY_REF, ptr_type_node, *elems, loop->index, NULL_TREE, + NULL_TREE); + tree tem = build2_loc (OMP_CLAUSE_LOCATION (c), MODIFY_EXPR, void_type_node, + lhs, expr); + append_to_statement_list_force (tem, loop->body); + + return build_fold_addr_expr_with_type (*elems, ptr_type_node); +} + +/* Set SIZE as the size expression that should result from the clause C. + LOOPS holds the intermediate loop info. Returns the tree that should be + passed as the clause size. */ + +static tree +lower_omp_map_iterator_size (tree size, tree c, iterator_loop_map_t *loops) +{ + if (!OMP_CLAUSE_HAS_ITERATORS (c) || !OMP_CLAUSE_ITERATORS (c)) + return size; + + iterator_loop_info_t *loop = loops->get (OMP_CLAUSE_ITERATORS (c)); + gcc_assert (loop); + tree *elems = loop->elems.get (c); + + /* IN LOOP BODY: */ + /* elems[idx+1] = size; */ + tree lhs = build4 (ARRAY_REF, ptr_type_node, *elems, + size_binop (PLUS_EXPR, loop->index, size_int (1)), + NULL_TREE, NULL_TREE); + tree tem = build2_loc (OMP_CLAUSE_LOCATION (c), MODIFY_EXPR, + void_type_node, lhs, size); + append_to_statement_list_force (tem, loop->body); + + return size_int (SIZE_MAX); +} + +/* Finish building the iterator loops in LOOPS, with generated Gimple + statements going in PRE_P. The loops cannot be amended after this is + called. */ + +static void +finish_omp_map_iterators (iterator_loop_map_t *loops, gimple_seq *pre_p) +{ + for (iterator_loop_map_t::iterator it = loops->begin (); + it != loops->end (); ++it) + { + iterator_loop_info_t &loop = (*it).second; + gimplify_and_add (loop.bind, pre_p); + + for (hash_map::iterator it2 = loop.elems.begin (); + it2 != loop.elems.end (); ++it2) + { + tree clause = (*it2).first; + OMP_CLAUSE_DECL (clause) = (*it2).second; + OMP_CLAUSE_SIZE (clause) = size_int (SIZE_MAX); + } + } +} + /* Lower the GIMPLE_OMP_TARGET in the current statement in GSI_P. CTX holds context information for the directive. */ @@ -12617,7 +12774,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) tree child_fn, t, c; gomp_target *stmt = as_a (gsi_stmt (*gsi_p)); gbind *tgt_bind, *bind, *dep_bind = NULL; - gimple_seq tgt_body, olist, ilist, fplist, new_body; + gimple_seq tgt_body, olist, iterlist, ilist, fplist, new_body; location_t loc = gimple_location (stmt); bool offloaded, data_region; unsigned int map_cnt = 0; @@ -12628,6 +12785,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) tree deep_map_offset_data = NULL_TREE; tree deep_map_offset = NULL_TREE; + iterator_loop_map_t iterator_loops; + offloaded = is_gimple_omp_offloaded (stmt); switch (gimple_omp_target_kind (stmt)) { @@ -12706,6 +12865,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) push_gimplify_context (); fplist = NULL; + iterlist = NULL; ilist = NULL; olist = NULL; for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c)) @@ -12761,7 +12921,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) gcc_unreachable (); } #endif - /* FALLTHRU */ + lower_omp_map_iterators (c, &iterlist, &iterator_loops); + /* FALLTHRU */ case OMP_CLAUSE_TO: case OMP_CLAUSE_FROM: oacc_firstprivate: @@ -13190,6 +13351,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) *p = build_fold_indirect_ref (nd); } v = build_fold_addr_expr_with_type (v, ptr_type_node); + v = lower_omp_map_iterator_expr (v, c, &iterator_loops); gimplify_assign (x, v, &ilist); nc = NULL_TREE; } @@ -13263,12 +13425,18 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) && TREE_CODE (TREE_TYPE (ovar)) == ARRAY_TYPE) { gcc_assert (offloaded); - tree avar - = create_tmp_var (TREE_TYPE (TREE_TYPE (x))); - mark_addressable (avar); - gimplify_assign (avar, build_fold_addr_expr (var), &ilist); - talign = DECL_ALIGN_UNIT (avar); + tree avar = build_fold_addr_expr (var); + if (!OMP_CLAUSE_ITERATORS (c)) + { + tree tmp = create_tmp_var (TREE_TYPE (TREE_TYPE (x))); + mark_addressable (tmp); + gimplify_assign (tmp, avar, &ilist); + avar = tmp; + } + talign = TYPE_ALIGN_UNIT (TREE_TYPE (TREE_TYPE (x))); avar = build_fold_addr_expr (avar); + avar = lower_omp_map_iterator_expr (avar, c, + &iterator_loops); gimplify_assign (x, avar, &ilist); } else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE) @@ -13348,6 +13516,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) if (s == NULL_TREE) s = TYPE_SIZE_UNIT (TREE_TYPE (ovar)); s = fold_convert (size_type_node, s); + s = lower_omp_map_iterator_size (s, c, &iterator_loops); purpose = size_int (map_idx++); CONSTRUCTOR_APPEND_ELT (vsize, purpose, s); if (TREE_CODE (s) != INTEGER_CST) @@ -13713,6 +13882,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) DECL_INITIAL (TREE_VEC_ELT (t, 2)) = build_constructor (TREE_TYPE (TREE_VEC_ELT (t, 2)), vkind); } + finish_omp_map_iterators (&iterator_loops, &iterlist); for (int i = 1; i <= 2; i++) if (deep_map_cnt || !TREE_STATIC (TREE_VEC_ELT (t, i))) { @@ -14280,6 +14450,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) gimple_omp_set_body (stmt, new_body); } + gsi_insert_seq_before (gsi_p, iterlist, GSI_SAME_STMT); + bind = gimple_build_bind (NULL, NULL, tgt_bind ? gimple_bind_block (tgt_bind) : NULL_TREE); diff --git a/gcc/testsuite/c-c++-common/gomp/map-6.c b/gcc/testsuite/c-c++-common/gomp/map-6.c index 014ed35ab41..13e3b58cc92 100644 --- a/gcc/testsuite/c-c++-common/gomp/map-6.c +++ b/gcc/testsuite/c-c++-common/gomp/map-6.c @@ -13,19 +13,19 @@ foo (void) #pragma omp target map (to:a) ; - #pragma omp target map (a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present'" } */ + #pragma omp target map (a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close', 'iterator' or 'present'" } */ ; - #pragma omp target map (close, a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present'" } */ + #pragma omp target map (close, a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close', 'iterator' or 'present'" } */ ; - #pragma omp target enter data map(b7) map (close, a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present'" } */ + #pragma omp target enter data map(b7) map (close, a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close', 'iterator' or 'present'" } */ ; - #pragma omp target exit data map(b7) map (close, a from: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present'" } */ + #pragma omp target exit data map(b7) map (close, a from: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close', 'iterator' or 'present'" } */ ; - #pragma omp target data map(b7) map (close, a from: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present'" } */ + #pragma omp target data map(b7) map (close, a from: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close', 'iterator' or 'present'" } */ ; diff --git a/gcc/testsuite/c-c++-common/gomp/target-map-iterators-1.c b/gcc/testsuite/c-c++-common/gomp/target-map-iterators-1.c new file mode 100644 index 00000000000..7d6c8dc6255 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/target-map-iterators-1.c @@ -0,0 +1,23 @@ +/* { dg-do compile } */ +/* { dg-options "-fopenmp" } */ + +#define DIM1 17 +#define DIM2 39 + +void f (int **x, int **y) +{ + #pragma omp target map(iterator(i=0:DIM1), to: x[i][:DIM2]) + ; + + #pragma omp target map(iterator(i=0:DIM1), to: x[i][:DIM2], y[i][:DIM2]) + ; + + #pragma omp target map(iterator(i=0:DIM1), to: x[i][:DIM2] + 2) /* { dg-message "unsupported map expression" } */ + ; + + #pragma omp target map(iterator(i=0:DIM1), iterator(j=0:DIM2), to: x[i][j]) /* { dg-error "too many 'iterator' modifiers" } */ + ; + + #pragma omp target map(iterator(i=0:DIM1), to: (i % 2 == 0) ? x[i] : y[i]) /* { dg-message "unsupported map expression" } */ + ; +} diff --git a/gcc/testsuite/c-c++-common/gomp/target-map-iterators-2.c b/gcc/testsuite/c-c++-common/gomp/target-map-iterators-2.c new file mode 100644 index 00000000000..da14d068f19 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/target-map-iterators-2.c @@ -0,0 +1,19 @@ +/* { dg-do compile } */ +/* { dg-options "-fopenmp" } */ + +void f (int *x, float *y, double *z) +{ + #pragma omp target map(iterator(i=0:10), to: x) /* { dg-error "iterator variable .i. not used in clause expression" } */ + /* Add a reference to x to ensure that the 'to' clause does not get + dropped. */ + x[0] = 0; + + #pragma omp target map(iterator(i=0:10, j=0:20), to: x[i]) /* { dg-error "iterator variable .j. not used in clause expression" } */ + ; + + #pragma omp target map(iterator(i=0:10, j=0:20, k=0:30), to: x[i], y[j], z[k]) + /* { dg-error "iterator variable .i. not used in clause expression" "" { target *-*-* } .-1 } */ + /* { dg-error "iterator variable .j. not used in clause expression" "" { target *-*-* } .-2 } */ + /* { dg-error "iterator variable .k. not used in clause expression" "" { target *-*-* } .-3 } */ + ; +} diff --git a/gcc/testsuite/c-c++-common/gomp/target-map-iterators-3.c b/gcc/testsuite/c-c++-common/gomp/target-map-iterators-3.c new file mode 100644 index 00000000000..fb0c761018a --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/target-map-iterators-3.c @@ -0,0 +1,23 @@ +/* { dg-do compile } */ +/* { dg-options "-fopenmp -fdump-tree-omplower" } */ + +#define DIM1 10 +#define DIM2 20 +#define DIM3 30 + +void f (int ***x, float ***y, double **z) +{ + #pragma omp target \ + map(to: x, y) \ + map(iterator(i=0:DIM1, j=0:DIM2), to: x[i][j][:DIM3], y[i][j][:DIM3]) \ + map(from: z) \ + map(iterator(i=0:DIM1), from: z[i][:DIM2]) + ; +} + +/* { dg-final { scan-tree-dump-times "if \\(i <= 9\\) goto ; else goto ;" 2 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "if \\(j <= 19\\) goto ; else goto ;" 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "map\\(iterator\\(int i=0:10:1\\):from:D\.\[0-9\]+" 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "map\\(iterator\\(int i=0:10:1\\):attach:D\.\[0-9\]+" 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "map\\(iterator\\(int i=0:10:1, int j=0:20:1\\):to:D\.\[0-9\]+" 2 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "map\\(iterator\\(int i=0:10:1, int j=0:20:1\\):attach:D\.\[0-9\]+" 4 "omplower" } } */ diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc index 39e586c808c..be2723dcdae 100644 --- a/gcc/tree-pretty-print.cc +++ b/gcc/tree-pretty-print.cc @@ -911,6 +911,11 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) pp_string (pp, "map("); if (OMP_CLAUSE_MAP_READONLY (clause)) pp_string (pp, "readonly,"); + if (OMP_CLAUSE_ITERATORS (clause)) + { + dump_omp_iterators (pp, OMP_CLAUSE_ITERATORS (clause), spc, flags); + pp_colon (pp); + } switch (OMP_CLAUSE_MAP_KIND (clause)) { case GOMP_MAP_ALLOC: diff --git a/gcc/tree.cc b/gcc/tree.cc index bc50afca9a3..f12d7b8bb8a 100644 --- a/gcc/tree.cc +++ b/gcc/tree.cc @@ -266,7 +266,7 @@ unsigned const char omp_clause_num_ops[] = 1, /* OMP_CLAUSE_EXCLUSIVE */ 2, /* OMP_CLAUSE_FROM */ 2, /* OMP_CLAUSE_TO */ - 2, /* OMP_CLAUSE_MAP */ + 3, /* OMP_CLAUSE_MAP */ 1, /* OMP_CLAUSE_HAS_DEVICE_ADDR */ 1, /* OMP_CLAUSE_DOACROSS */ 2, /* OMP_CLAUSE__CACHE_ */ @@ -11598,6 +11598,9 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data, case OMP_CLAUSE: { int len = omp_clause_num_ops[OMP_CLAUSE_CODE (t)]; + /* Do not walk the iterator operand of OpenMP MAP clauses. */ + if (OMP_CLAUSE_HAS_ITERATORS (t)) + len--; for (int i = 0; i < len; i++) WALK_SUBTREE (OMP_CLAUSE_OPERAND (t, i)); WALK_SUBTREE_TAIL (OMP_CLAUSE_CHAIN (t)); diff --git a/gcc/tree.h b/gcc/tree.h index 83075b82cc7..384a5f1f250 100644 --- a/gcc/tree.h +++ b/gcc/tree.h @@ -1619,6 +1619,13 @@ class auto_suppress_location_wrappers != UNKNOWN_LOCATION) #define OMP_CLAUSE_LOCATION(NODE) (OMP_CLAUSE_CHECK (NODE))->omp_clause.locus +#define OMP_CLAUSE_HAS_ITERATORS(NODE) \ + (OMP_CLAUSE_CODE (NODE) == OMP_CLAUSE_MAP) +#define OMP_CLAUSE_ITERATORS(NODE) \ + OMP_CLAUSE_OPERAND (OMP_CLAUSE_RANGE_CHECK (OMP_CLAUSE_CHECK (NODE), \ + OMP_CLAUSE_MAP, \ + OMP_CLAUSE_MAP), 2) + /* True on OMP_FOR and other OpenMP/OpenACC looping constructs if the loop nest is non-rectangular. */ #define OMP_FOR_NON_RECTANGULAR(NODE) \ diff --git a/libgomp/target.c b/libgomp/target.c index cf62af61f3b..463a162879b 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -975,6 +975,105 @@ gomp_map_val (struct target_mem_desc *tgt, void **hostaddrs, size_t i) } } +static const char * +kind_to_name (unsigned short kind) +{ + if (GOMP_MAP_IMPLICIT_P (kind)) + kind &= ~GOMP_MAP_IMPLICIT; + + switch (kind & 0xff) + { + case GOMP_MAP_ALLOC: return "GOMP_MAP_ALLOC"; + case GOMP_MAP_FIRSTPRIVATE: return "GOMP_MAP_FIRSTPRIVATE"; + case GOMP_MAP_FIRSTPRIVATE_INT: return "GOMP_MAP_FIRSTPRIVATE_INT"; + case GOMP_MAP_TO: return "GOMP_MAP_TO"; + case GOMP_MAP_TO_PSET: return "GOMP_MAP_TO_PSET"; + case GOMP_MAP_FROM: return "GOMP_MAP_FROM"; + case GOMP_MAP_TOFROM: return "GOMP_MAP_TOFROM"; + case GOMP_MAP_POINTER: return "GOMP_MAP_POINTER"; + case GOMP_MAP_ATTACH: return "GOMP_MAP_ATTACH"; + case GOMP_MAP_DETACH: return "GOMP_MAP_DETACH"; + default: return "unknown"; + } +} + +/* Map entries containing expanded iterators will be flattened and merged into + HOSTADDRS, SIZES and KINDS, and MAPNUM updated. Returns true if there are + any iterators found. ITERATOR_COUNT holds the iteration count of the + iterator that generates each map (0 if not generated from an iterator). + HOSTADDRS, SIZES, KINDS and ITERATOR_COUNT must be freed afterwards if any + merging occurs. */ + +static bool +gomp_merge_iterator_maps (size_t *mapnum, void ***hostaddrs, size_t **sizes, + void **kinds, size_t **iterator_count) +{ + bool iterator_p = false; + size_t map_count = 0; + unsigned short **skinds = (unsigned short **) kinds; + + for (size_t i = 0; i < *mapnum; i++) + if ((*sizes)[i] == SIZE_MAX) + { + uintptr_t *iterator_array = (*hostaddrs)[i]; + map_count += iterator_array[0]; + iterator_p = true; + } + else + map_count++; + + if (!iterator_p) + return false; + + gomp_debug (1, + "Expanding iterator maps - number of map entries: %u -> %u\n", + (int) *mapnum, (int) map_count); + void **new_hostaddrs = (void **) gomp_malloc (map_count * sizeof (void *)); + size_t *new_sizes = (size_t *) gomp_malloc (map_count * sizeof (size_t)); + unsigned short *new_kinds + = (unsigned short *) gomp_malloc (map_count * sizeof (unsigned short)); + size_t new_idx = 0; + *iterator_count = (size_t *) gomp_malloc (map_count * sizeof (size_t)); + + for (size_t i = 0; i < *mapnum; i++) + { + if ((*sizes)[i] == SIZE_MAX) + { + uintptr_t *iterator_array = (*hostaddrs)[i]; + size_t count = *iterator_array++; + for (size_t j = 0; j < count; j++) + { + new_hostaddrs[new_idx] = (void *) *iterator_array++; + new_sizes[new_idx] = *iterator_array++; + new_kinds[new_idx] = (*skinds)[i]; + (*iterator_count)[new_idx] = j + 1; + gomp_debug (1, + "Expanding map %u <%s>: " + "hostaddrs[%u] = %p, sizes[%u] = %lu\n", + (int) i, kind_to_name (new_kinds[new_idx]), + (int) new_idx, new_hostaddrs[new_idx], + (int) new_idx, (unsigned long) new_sizes[new_idx]); + new_idx++; + } + } + else + { + new_hostaddrs[new_idx] = (*hostaddrs)[i]; + new_sizes[new_idx] = (*sizes)[i]; + new_kinds[new_idx] = (*skinds)[i]; + (*iterator_count)[new_idx] = 0; + new_idx++; + } + } + + *mapnum = map_count; + *hostaddrs = new_hostaddrs; + *sizes = new_sizes; + *kinds = new_kinds; + + return true; +} + static inline __attribute__((always_inline)) struct target_mem_desc * gomp_map_vars_internal (struct gomp_device_descr *devicep, struct goacc_asyncqueue *aq, size_t mapnum, @@ -991,6 +1090,11 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, const int typemask = short_mapkind ? 0xff : 0x7; struct splay_tree_s *mem_map = &devicep->mem_map; struct splay_tree_key_s cur_node; + bool iterators_p = false; + size_t *iterator_count = NULL; + if (short_mapkind) + iterators_p = gomp_merge_iterator_maps (&mapnum, &hostaddrs, &sizes, + &kinds, &iterator_count); struct target_mem_desc *tgt = gomp_malloc (sizeof (*tgt) + sizeof (tgt->list[0]) * mapnum); tgt->list_count = mapnum; @@ -1840,14 +1944,17 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, if (pragma_kind & GOMP_MAP_VARS_TARGET) { + size_t map_num = 0; for (i = 0; i < mapnum; i++) - { - cur_node.tgt_offset = gomp_map_val (tgt, hostaddrs, i); - gomp_copy_host2dev (devicep, aq, - (void *) (tgt->tgt_start + i * sizeof (void *)), - (void *) &cur_node.tgt_offset, sizeof (void *), - true, cbufp); - } + if (!iterator_count || iterator_count[i] <= 1) + { + cur_node.tgt_offset = gomp_map_val (tgt, hostaddrs, i); + gomp_copy_host2dev (devicep, aq, + (void *) (tgt->tgt_start + map_num * sizeof (void *)), + (void *) &cur_node.tgt_offset, sizeof (void *), + true, cbufp); + map_num++; + } } if (cbufp) @@ -1879,6 +1986,15 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, } gomp_mutex_unlock (&devicep->lock); + + if (iterators_p) + { + free (hostaddrs); + free (sizes); + free (kinds); + free (iterator_count); + } + return tgt; } diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-1.c b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-1.c new file mode 100644 index 00000000000..b3d87f231df --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-1.c @@ -0,0 +1,47 @@ +/* { dg-do run } */ +/* { dg-require-effective-target offload_device_nonshared_as } */ + +/* Test transfer of dynamically-allocated arrays to target using map + iterators. */ + +#include + +#define DIM1 8 +#define DIM2 15 + +int mkarray (int *x[]) +{ + int expected = 0; + + for (int i = 0; i < DIM1; i++) + { + x[i] = (int *) malloc (DIM2 * sizeof (int)); + for (int j = 0; j < DIM2; j++) + { + x[i][j] = rand (); + expected += x[i][j]; + } + } + + return expected; +} + +int main (void) +{ + int *x[DIM1]; + int y; + + int expected = mkarray (x); + + #pragma omp target enter data map(to: x) + #pragma omp target map(iterator(i=0:DIM1), to: x[i][:DIM2]) \ + map(from: y) + { + y = 0; + for (int i = 0; i < DIM1; i++) + for (int j = 0; j < DIM2; j++) + y += x[i][j]; + } + + return y - expected; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-2.c b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-2.c new file mode 100644 index 00000000000..8569b55ab5b --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-2.c @@ -0,0 +1,44 @@ +/* { dg-do run } */ +/* { dg-require-effective-target offload_device_nonshared_as } */ + +/* Test transfer of dynamically-allocated arrays from target using map + iterators. */ + +#include + +#define DIM1 8 +#define DIM2 15 + +void mkarray (int *x[]) +{ + for (int i = 0; i < DIM1; i++) + x[i] = (int *) malloc (DIM2 * sizeof (int)); +} + +int main (void) +{ + int *x[DIM1]; + int y, expected; + + mkarray (x); + + #pragma omp target enter data map(alloc: x) + #pragma omp target map(iterator(i=0:DIM1), from: x[i][:DIM2]) \ + map(from: expected) + { + expected = 0; + for (int i = 0; i < DIM1; i++) + for (int j = 0; j < DIM2; j++) + { + x[i][j] = (i+1) * (j+1); + expected += x[i][j]; + } + } + + y = 0; + for (int i = 0; i < DIM1; i++) + for (int j = 0; j < DIM2; j++) + y += x[i][j]; + + return y - expected; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-3.c b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-3.c new file mode 100644 index 00000000000..be30fa65d80 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-3.c @@ -0,0 +1,56 @@ +/* { dg-do run } */ +/* { dg-require-effective-target offload_device_nonshared_as } */ + +/* Test transfer of dynamically-allocated arrays to target using map + iterators, with multiple iterators and function calls in the iterator + expression. */ + +#include + +#define DIM1 16 +#define DIM2 15 + +int mkarrays (int *x[], int *y[]) +{ + int expected = 0; + + for (int i = 0; i < DIM1; i++) + { + x[i] = (int *) malloc (DIM2 * sizeof (int)); + y[i] = (int *) malloc (sizeof (int)); + *y[i] = rand (); + for (int j = 0; j < DIM2; j++) + { + x[i][j] = rand (); + expected += x[i][j] * *y[i]; + } + } + + return expected; +} + +int f (int i, int j) +{ + return i * 4 + j; +} + +int main (void) +{ + int *x[DIM1], *y[DIM1]; + int sum; + + int expected = mkarrays (x, y); + + #pragma omp target enter data map(to: x, y) + #pragma omp target map(iterator(i=0:DIM1/4, j=0:4), to: x[f(i, j)][:DIM2]) \ + map(iterator(i=0:DIM1), to: y[i][:1]) \ + map(from: sum) + { + sum = 0; + for (int i = 0; i < DIM1; i++) + for (int j = 0; j < DIM2; j++) + sum += x[i][j] * y[i][0]; + } + + return sum - expected; +}