From patchwork Tue Sep 3 17:07:32 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: 1980234 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=Vf0WhWQ3; dkim-atps=neutral Authentication-Results: legolas.ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=2620:52:3:1:0:246e:9693:128c; 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 [IPv6:2620:52:3:1:0:246e:9693:128c]) (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 4WysXh74VNz1yZ9 for ; Wed, 4 Sep 2024 03:08:52 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 51E3F386182F for ; Tue, 3 Sep 2024 17:08:50 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from mail-wm1-x331.google.com (mail-wm1-x331.google.com [IPv6:2a00:1450:4864:20::331]) by sourceware.org (Postfix) with ESMTPS id B39913858402 for ; Tue, 3 Sep 2024 17:08:25 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org B39913858402 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 B39913858402 Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=2a00:1450:4864:20::331 ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1725383311; cv=none; b=qbDUOY66Ulto/flfUSOaR3NvkmObOSmQDhi9Q108v65ZYA7xo1JkT/2fK4jDfLsuO4Dr/tFRsPZni9GSEckX3yfNlxziZOSE9MU82/Ib3y6xzYZNyXIpVUhkLzDVoV2hsrTwm0ujmLXZjgVPrNyagieo4x7nkToH1tH/BAIb+S4= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1725383311; c=relaxed/simple; bh=0K8RW+2ziyweHyn6aUsnz3XBWk7zbrANm8DdB9bvqv8=; h=DKIM-Signature:Message-ID:Date:MIME-Version:To:Subject:From; b=i6o0LFa3flhCAMuIBFub0C6hjPjQxxqNUreHjHSVTDsUl011SmdUujDdhVwMbV3PwN6fCNu2jiRXoKY5XjC6qldd98lTp9s7IpC0xtnpyl88ctZYUv5pHvDWwg3YqvV4R+UbnPPGrVUezPrA36D5fK3jY49hfdyRrgLqbi5JgGI= ARC-Authentication-Results: i=1; server2.sourceware.org Received: by mail-wm1-x331.google.com with SMTP id 5b1f17b1804b1-42bb8cf8abeso44946985e9.2 for ; Tue, 03 Sep 2024 10:08:25 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre-com.20230601.gappssmtp.com; s=20230601; t=1725383304; x=1725988104; 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=MUSWmhryH7SH4uY64OHDM5xGORFtSsy22b7pH47Kafw=; b=Vf0WhWQ3sD7pOrnTPu6D8V3IPLfFK5XY+XOiVl7ILszA5vWum6SvCtiwpmrAHBH2NP a0RXxi0Dr4+cc0Tm5/EgGDELVQzI1Pq3wacCwa7nB5QNafP3JrHFo0joeYeWaSWmMSlo 6uOrukLgK4p8/qRpNdky/RYs/Q/Q/dOTfnWxUSdS3/Sh1pvFMuDk1F9nzArQBWjhK1Hb jU2zhY4sSDnDTzsvoM9M8J0crdg7VfjvR4CXfVbxmkZ1VhiCp54ueYFFcpaEBPvUS6xV sIQpb2fUGWZzwOcGoPIX1pBmG9QKnKMMaGWWNGfomd+Tx9x1ofszmqBc+p9RxZ2yEE54 s/UQ== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1725383304; x=1725988104; 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=MUSWmhryH7SH4uY64OHDM5xGORFtSsy22b7pH47Kafw=; b=Sf20pR7wghXibrnLQ2rhh5NCar5vgzXZiB+KToIPJSwJZUqLOHx8WctHrn0sWnx7BY 2U6sz1MzF2X6Ace0TffDXQxrPAle+6fW2GoRCDdGfTdELV1KbXiUiSnaBZPoeA7A1fsx eNFmUDoJoQTfx4GqoLFR/nkOHjmSa/2JIq88tIZdLugZ58+/0KBmUG/PZ1xSn+eJCeRY PmfkXPlt0f4+K3ZDzNbeLSfIuIrs/zRYVGv02QZUi/fsDe0Bjk2TM7+8OfTIjy4zanF0 ZNiLhwChgYIssSvX5rWFcFbVFosmKtZchZACfYh4rytyNqvYc0X6mEFcLPVS7X4OaMBa r8xw== X-Gm-Message-State: AOJu0YwdvGAPQCnFpPEEiAm1Cy+S7jwIm/Lj/KnsKtjFV54jwBOJD/v3 pFkpXilJ0OZM5+SV92Bt0sk5iI77m23NuJC6ulwrUf711sqJc0yO9VKQyC2w3x5ie2ER3+4HMYU y X-Google-Smtp-Source: AGHT+IGphbZSgiIWDyWOmFWk0gc1Exxsx7L6SWTW1TzgBH2mTMR3fw7Q2VM/tSRq9r4N+JDGQEoXLg== X-Received: by 2002:adf:fc0f:0:b0:374:b300:c4d5 with SMTP id ffacd0b85a97d-376dd71a859mr1007726f8f.28.1725383303682; Tue, 03 Sep 2024 10:08:23 -0700 (PDT) Received: from ?IPV6:2a00:23c6:88fe:9301:912a:9e8a:468f:40d0? ([2a00:23c6:88fe:9301:912a:9e8a:468f:40d0]) by smtp.gmail.com with ESMTPSA id a640c23a62f3a-a8989035f80sm701194066b.78.2024.09.03.10.08.23 (version=TLS1_3 cipher=TLS_AES_128_GCM_SHA256 bits=128/128); Tue, 03 Sep 2024 10:08:23 -0700 (PDT) Message-ID: <858d4273-2a1a-41a7-813b-59514e5c485a@baylibre.com> Date: Tue, 3 Sep 2024 18:07:32 +0100 MIME-Version: 1.0 User-Agent: Mozilla Thunderbird To: gcc-patches , Jakub Jelinek , Tobias Burnus References: <77f9ccb8-6f5e-4462-aa32-71f74fd7ff26@baylibre.com> Subject: [PATCH v2 2/5] openmp: Add support for iterators in map clauses (C/C++) Content-Language: en-GB From: Kwok Cheung Yeung In-Reply-To: <77f9ccb8-6f5e-4462-aa32-71f74fd7ff26@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, T_SCC_BODY_TEXT_LINE 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, encoded in the same way as the depend and affinity clauses. When finishing the clauses, clauses with iterators are treated separately from ones without to avoid clashes (e.g. iterating over x[i] will likely generate implicit clauses to map x). During gimplification, clauses with iterators are treated similarly to normal clauses, removing the iterator from the clause decl if necessary. gimplify_omp_map_iterators is called at the end of gimplify_adjust_omp_clauses. For each map clause with an iterator, gimplify_omp_map_iterators generates a loop (or multiple loops, if the iterator is multidimensional) to iterate over the iterator expression, storing the result in a new array (constant-sized for now, we could dynamically allocate the array for non-constant iteration bounds). The data array stores the total number of iterations in the first element, then the address generated by the iterator expression and the OMP_CLAUSE_SIZE (since the iteration variables may occur within the size tree) for each iteration. The clause is then rewritten to point to the new array. The original clause decl is no longer directly relevant, but is kept around for informational purposes. The original OMP_CLAUSE_SIZE is set to SIZE_MAX to indicate that the clause has an expanded iterator associated with it. Multiple clauses using the same iterator are expanded together even if they are not adjacent. When OMP lowering clauses with iterators, the data array holding the expanded iterator info is used as the variable to send. 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. From dd65c671dc9f5fb34290938a413c610eb0110ef6 Mon Sep 17 00:00:00 2001 From: Kwok Cheung Yeung Date: Mon, 2 Sep 2024 19:33:47 +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-09-02 Kwok Cheung Yeung gcc/c/ * c-parser.cc (c_parser_omp_clause_map): Parse 'iterator' modifier. * c-typeck.cc (c_finish_omp_clauses): Call recursively on iterator clauses. gcc/cp/ * parser.cc (cp_parser_omp_clause_map): Parse 'iterator' modifier. * semantics.cc (finish_omp_clauses): Call recursively on iterator clauses. gcc/ * gimplify.cc (build_iterator_loop): Do not gimplify last binding into SSA. (find_var_decl): New. (check_iterator_var_usage): New. (gimplify_omp_map_iterators): New. (omp_group_iterator): New. (omp_get_attachment): Replace OMP_CLAUSE_DECL with OMP_ITERATOR_CLAUSE_DECL. (omp_group_last): Keep decls with and without iterators in separate groups. (omp_index_mapping_groups_1): Replace OMP_CLAUSE_DECL with OMP_ITERATOR_CLAUSE_DECL. (omp_tsort_mapping_groups_1): Likewise. (omp_resolve_clause_dependencies): Likewise. Prevent removal of mapping if groups do not use the same iterators. (omp_accumulate_sibling_list): Replace OMP_CLAUSE_DECL with OMP_ITERATOR_CLAUSE_DECL. (omp_build_struct_sibling_lists): Likewise. (gimplify_scan_omp_clauses): Remove iterators from clauses before scanning clauses. Replace afterwards. (gimplify_adjust_omp_clauses): Replace OMP_CLAUSE_DECL with OMP_ITERATOR_CLAUSE_DECL. Skip gimplification of clause decl and size for clauses with iterators. Call gimplify_omp_map_iterators. * omp-low.cc (scan_sharing_clauses): Add field for iterator clauses. (lower_omp_target): Replace OMP_CLAUSE_DECL with OMP_ITERATOR_CLAUSE_DECL. Always increase map count by one for clauses with iterators. Use expanded iterator array as the output variable for iterator clauses. * tree-pretty-print.cc (dump_omp_map_iterators): New. (dump_omp_clause): Call dump_omp_map_iterators for iterators in map clauses. * tree.h (OMP_ITERATOR_CLAUSE_DECL): New. gcc/testsuite/ * c-c++-common/gomp/map-6.c (foo): Amend expected error message. * c-c++-common/gomp/target-iterator-1.c: New. * c-c++-common/gomp/target-iterator-2.c: New. * c-c++-common/gomp/target-iterator-3.c: New. libgomp/ * target.c (gomp_merge_iterator_maps): New. (gomp_map_vars_internal): Call gomp_merge_iterator_maps. 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 | 60 ++++- gcc/c/c-typeck.cc | 68 ++++++ gcc/cp/parser.cc | 63 ++++- gcc/cp/semantics.cc | 65 +++++ gcc/gimplify.cc | 229 +++++++++++++++++- gcc/omp-low.cc | 17 +- gcc/testsuite/c-c++-common/gomp/map-6.c | 10 +- .../c-c++-common/gomp/target-iterator-1.c | 23 ++ .../c-c++-common/gomp/target-iterator-2.c | 19 ++ .../c-c++-common/gomp/target-iterator-3.c | 20 ++ gcc/tree-pretty-print.cc | 24 +- gcc/tree.h | 7 + libgomp/target.c | 83 +++++++ .../target-map-iterators-1.c | 44 ++++ .../target-map-iterators-2.c | 42 ++++ .../target-map-iterators-3.c | 54 +++++ 16 files changed, 793 insertions(+), 35 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/gomp/target-iterator-1.c create mode 100644 gcc/testsuite/c-c++-common/gomp/target-iterator-2.c create mode 100644 gcc/testsuite/c-c++-common/gomp/target-iterator-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 aff5af17430..f72fca1a711 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,21 @@ 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); + if (iterators) + OMP_CLAUSE_DECL (c) = build_tree_list (iterators, OMP_CLAUSE_DECL (c)); + } parens.skip_until_found_close (parser); return nl; diff --git a/gcc/c/c-typeck.cc b/gcc/c/c-typeck.cc index 521c0e85605..d631d95f091 100644 --- a/gcc/c/c-typeck.cc +++ b/gcc/c/c-typeck.cc @@ -15811,6 +15811,74 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) auto_vec addr_tokens; t = OMP_CLAUSE_DECL (c); + if (OMP_ITERATOR_DECL_P (t)) + { + tree iterators = TREE_PURPOSE (t); + if (c_omp_finish_iterators (iterators)) + { + t = error_mark_node; + break; + } + + /* Find the end of the group of clauses that use the same + iterator.*/ + tree end_clause; + for (end_clause = c; end_clause; + end_clause = OMP_CLAUSE_CHAIN (end_clause)) + { + tree nc = OMP_CLAUSE_CHAIN (end_clause); + /* Remove iterator temporarily. */ + OMP_CLAUSE_DECL (end_clause) = + TREE_VALUE (OMP_CLAUSE_DECL (end_clause)); + if (!nc + || !OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (nc)) + || TREE_PURPOSE (OMP_CLAUSE_DECL (nc)) != iterators) + break; + } + tree next_clause = OMP_CLAUSE_CHAIN (end_clause); + + /* Temporarily split off the group of clauses with the same + iterator. */ + OMP_CLAUSE_CHAIN (end_clause) = NULL_TREE; + tree new_clauses = c_finish_omp_clauses (c, ort); + + /* Replace the iterators and splice the new clauses in. */ + tree *clause_p = &new_clauses; + while (*clause_p) + { + /* Skip unwanted clause types. + FIXME: Is this the right thing to do? */ + bool skip = false; + if (OMP_CLAUSE_CODE (*clause_p) == OMP_CLAUSE_MAP) + switch (OMP_CLAUSE_MAP_KIND (*clause_p)) + { + case GOMP_MAP_TO: + case GOMP_MAP_FROM: + case GOMP_MAP_ATTACH: + case GOMP_MAP_DETACH: + skip = false; + break; + default: + skip = true; + break; + } + if (skip) + *clause_p = OMP_CLAUSE_CHAIN (*clause_p); + else + { + OMP_CLAUSE_DECL (*clause_p) + = build_tree_list (iterators, + OMP_CLAUSE_DECL (*clause_p)); + + clause_p = &OMP_CLAUSE_CHAIN (*clause_p); + } + } + *clause_p = next_clause; + *pc = new_clauses; + pc = clause_p; + continue; + } + if (TREE_CODE (t) == OMP_ARRAY_SECTION) { grp_start_p = pc; diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc index edfa5a49440..29947177415 100644 --- a/gcc/cp/parser.cc +++ b/gcc/cp/parser.cc @@ -41719,16 +41719,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++; } @@ -41736,6 +41754,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); @@ -41785,10 +41804,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, @@ -41852,8 +41890,21 @@ 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); + if (iterators) + OMP_CLAUSE_DECL (c) = build_tree_list (iterators, OMP_CLAUSE_DECL (c)); + } return nlist; } diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc index 7ecad569900..6222cc2fe87 100644 --- a/gcc/cp/semantics.cc +++ b/gcc/cp/semantics.cc @@ -8375,6 +8375,71 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) auto_vec addr_tokens; t = OMP_CLAUSE_DECL (c); + if (OMP_ITERATOR_DECL_P (t)) + { + tree iterators = TREE_PURPOSE (t); + if (cp_omp_finish_iterators (iterators)) + { + t = error_mark_node; + break; + } + + /* Find the end of the group of clauses that use the same + iterator.*/ + tree end_clause; + for (end_clause = c; end_clause; + end_clause = OMP_CLAUSE_CHAIN (end_clause)) + { + tree nc = OMP_CLAUSE_CHAIN (end_clause); + /* Remove iterator temporarily. */ + OMP_CLAUSE_DECL (end_clause) = + TREE_VALUE (OMP_CLAUSE_DECL (end_clause)); + if (!nc + || !OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (nc)) + || TREE_PURPOSE (OMP_CLAUSE_DECL (nc)) != iterators) + break; + } + tree next_clause = OMP_CLAUSE_CHAIN (end_clause); + + /* Temporarily split off the group of clauses with the same + iterator. */ + OMP_CLAUSE_CHAIN (end_clause) = NULL_TREE; + tree new_clauses = finish_omp_clauses (c, ort); + + /* Replace the iterators and splice the new clauses in. */ + tree *clause_p = &new_clauses; + while (*clause_p) + { + OMP_CLAUSE_DECL (*clause_p) + = build_tree_list (iterators, + OMP_CLAUSE_DECL (*clause_p)); + /* Skip unwanted clause types. + FIXME: Is this the right thing to do? */ + bool skip = false; + if (OMP_CLAUSE_CODE (*clause_p) == OMP_CLAUSE_MAP) + switch (OMP_CLAUSE_MAP_KIND (*clause_p)) + { + case GOMP_MAP_TO: + case GOMP_MAP_FROM: + case GOMP_MAP_ATTACH: + case GOMP_MAP_DETACH: + skip = false; + break; + default: + skip = true; + break; + } + if (skip) + *clause_p = OMP_CLAUSE_CHAIN (*clause_p); + else + clause_p = &OMP_CLAUSE_CHAIN (*clause_p); + } + *clause_p = next_clause; + *pc = new_clauses; + pc = clause_p; + continue; + } + if (TREE_CODE (t) == OMP_ARRAY_SECTION) { grp_start_p = pc; diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index 8519095adef..549acf4dfbb 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -8906,7 +8906,12 @@ build_iterator_loop (tree c, gimple_seq *pre_p, tree *last_bind) gcc_assert (OMP_ITERATOR_DECL_P (t)); if (*last_bind) - gimplify_and_add (*last_bind, pre_p); + { + bool saved_into_ssa = gimplify_ctxp->into_ssa; + gimplify_ctxp->into_ssa = false; + gimplify_and_add (*last_bind, pre_p); + gimplify_ctxp->into_ssa = saved_into_ssa; + } tree block = TREE_VEC_ELT (TREE_PURPOSE (t), 5); *last_bind = build3 (BIND_EXPR, void_type_node, BLOCK_VARS (block), NULL, block); @@ -9330,6 +9335,166 @@ gimplify_omp_depend (tree *list_p, gimple_seq *pre_p) return 1; } +/* Callback for walk_tree to find a VAR_DECL for the given tree. */ + +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 for clause decls in iterators that do not use all the iterator + variables. */ + +static bool +check_iterator_var_usage (tree c) +{ + tree decl = OMP_CLAUSE_DECL (c); + bool error = false; + gcc_assert (OMP_ITERATOR_DECL_P (decl)); + + for (tree it = TREE_PURPOSE (decl); it; it = TREE_CHAIN (it)) + { + tree var = TREE_VEC_ELT (it, 0); + tree t = walk_tree (&TREE_VALUE (decl), 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; +} + +static void +gimplify_omp_map_iterators (tree *list_p, gimple_seq *pre_p) +{ + tree last_iter = NULL_TREE; + tree last_bind = NULL_TREE; + tree last_count = NULL_TREE; + tree last_index = NULL_TREE; + tree *last_body = NULL; + + /* Find all map iterators in use. */ + hash_set all_iterators; + for (tree c = *list_p; c; c = OMP_CLAUSE_CHAIN (c)) + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (c))) + all_iterators.add (TREE_PURPOSE (OMP_CLAUSE_DECL (c))); + + /* Expand all clauses using the same iterator together. */ + for (hash_set::iterator it = all_iterators.begin (); + it != all_iterators.end (); ++it) + { + for (tree c = *list_p; c; c = OMP_CLAUSE_CHAIN (c)) + { + if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP + || !OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (c))) + continue; + + tree t = OMP_CLAUSE_DECL (c); + if (TREE_PURPOSE (t) != *it) + continue; + + if (!check_iterator_var_usage (c)) + continue; + + if (TREE_PURPOSE (t) != last_iter) + { + tree tcnt = compute_iterator_count (t, pre_p); + if (!tcnt) + continue; + + last_iter = TREE_PURPOSE (t); + last_count = tcnt; + last_body = build_iterator_loop (c, pre_p, &last_bind); + last_index = create_tmp_var (sizetype); + SET_EXPR_LOCATION (last_bind, OMP_CLAUSE_LOCATION (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 (last_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, last_index, + size_binop (PLUS_EXPR, last_index, + size_int (2))); + append_to_statement_list_force (tem, last_body); + } + + /* Create array to hold expanded values. */ + tree last_count_2 = size_binop (MULT_EXPR, last_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."); + } + + /* 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, last_count); + gimplify_and_add (tem, pre_p); + + /* IN LOOP BODY: */ + /* elems[idx] = &; */ + lhs = build4 (ARRAY_REF, ptr_type_node, elems, last_index, NULL_TREE, + NULL_TREE); + tree rhs = build1 (ADDR_EXPR, ptr_type_node, TREE_VALUE (t)); + tem = build2_loc (OMP_CLAUSE_LOCATION (c), MODIFY_EXPR, + void_type_node, lhs, rhs); + append_to_statement_list_force (tem, last_body); + + /* elems[idx+1] = OMP_CLAUSE_SIZE (c); */ + lhs = build4 (ARRAY_REF, ptr_type_node, elems, + size_binop (PLUS_EXPR, last_index, size_int (1)), + NULL_TREE, NULL_TREE); + tem = build2_loc (OMP_CLAUSE_LOCATION (c), MODIFY_EXPR, + void_type_node, lhs, OMP_CLAUSE_SIZE (c)); + append_to_statement_list_force (tem, last_body); + + /* Replace iterator information. */ + TREE_PURPOSE (t) = make_tree_vec (2); + TREE_VEC_ELT (TREE_PURPOSE (t), 0) = last_iter; + TREE_VEC_ELT (TREE_PURPOSE (t), 1) = elems; + + OMP_CLAUSE_SIZE (c) = size_int (SIZE_MAX); + } + } + + if (last_bind) + { + bool saved_into_ssa = gimplify_ctxp->into_ssa; + gimplify_ctxp->into_ssa = false; + gimplify_and_add (last_bind, pre_p); + gimplify_ctxp->into_ssa = saved_into_ssa; + } +} + /* True if mapping node C maps, or unmaps, a (Fortran) array descriptor. */ static bool @@ -9538,6 +9703,22 @@ omp_get_base_pointer (tree expr) return NULL_TREE; } +/* Return the iterator for a mapping group, or NULL if there isn't one. */ + +static tree +omp_group_iterator (omp_mapping_group *grp) +{ + tree c = grp->grp_end; + if (!OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (c))) + return NULL_TREE; + + tree iter = TREE_PURPOSE (OMP_CLAUSE_DECL (c)); + if (TREE_VEC_LENGTH (iter) == 2) + iter = TREE_VEC_ELT (iter, 0); + + return iter; +} + /* An attach or detach operation depends directly on the address being attached/detached. Return that address, or none if there are no attachments/detachments. */ @@ -9592,7 +9773,7 @@ omp_get_attachment (omp_mapping_group *grp) case GOMP_MAP_ATTACH_DETACH: case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION: case GOMP_MAP_DETACH: - return OMP_CLAUSE_DECL (node); + return OMP_ITERATOR_CLAUSE_DECL (node); default: internal_error ("unexpected mapping node"); @@ -9604,7 +9785,7 @@ omp_get_attachment (omp_mapping_group *grp) node = OMP_CLAUSE_CHAIN (node); if (OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_ATTACH || OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_DETACH) - return OMP_CLAUSE_DECL (node); + return OMP_ITERATOR_CLAUSE_DECL (node); else internal_error ("unexpected mapping node"); return error_mark_node; @@ -9616,7 +9797,7 @@ omp_get_attachment (omp_mapping_group *grp) return OMP_CLAUSE_DECL (*grp->grp_start); if (OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_FIRSTPRIVATE_POINTER || OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_FIRSTPRIVATE_REFERENCE) - return OMP_CLAUSE_DECL (*grp->grp_start); + return OMP_ITERATOR_CLAUSE_DECL (*grp->grp_start); else internal_error ("unexpected mapping node"); return error_mark_node; @@ -9670,7 +9851,9 @@ omp_group_last (tree *start_p) == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION) || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_DETACH || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_ALWAYS_POINTER - || omp_map_clause_descriptor_p (nc))) + || omp_map_clause_descriptor_p (nc)) + && OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (c)) + == OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (nc))) { tree nc2 = OMP_CLAUSE_CHAIN (nc); if (OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_DETACH) @@ -9973,7 +10156,7 @@ omp_index_mapping_groups_1 (hash_mapgrp_start); + tree decl = OMP_ITERATOR_CLAUSE_DECL (*grp->grp_start); while (decl) { @@ -10699,7 +10882,7 @@ omp_resolve_clause_dependencies (enum tree_code code, FOR_EACH_VEC_ELT (*groups, i, grp) { tree grp_end = grp->grp_end; - tree decl = OMP_CLAUSE_DECL (grp_end); + tree decl = OMP_ITERATOR_CLAUSE_DECL (grp_end); gcc_assert (OMP_CLAUSE_CODE (grp_end) == OMP_CLAUSE_MAP); @@ -10886,7 +11069,9 @@ omp_resolve_clause_dependencies (enum tree_code code, { omp_mapping_group *struct_group; if (omp_mapped_by_containing_struct (grpmap, decl, &struct_group) - && *grp->grp_start == grp_end) + && *grp->grp_start == grp_end + && omp_group_iterator (grp) + == omp_group_iterator (struct_group)) { omp_check_mapping_compatibility (OMP_CLAUSE_LOCATION (grp_end), struct_group, grp); @@ -11208,7 +11393,7 @@ omp_accumulate_sibling_list (enum omp_region_type region_type, using namespace omp_addr_tokenizer; poly_offset_int coffset; poly_int64 cbitpos; - tree ocd = OMP_CLAUSE_DECL (grp_end); + tree ocd = OMP_ITERATOR_CLAUSE_DECL (grp_end); bool openmp = !(region_type & ORT_ACC); bool target = (region_type & ORT_TARGET) != 0; tree *continue_at = NULL; @@ -11806,7 +11991,7 @@ omp_build_struct_sibling_lists (enum tree_code code, FOR_EACH_VEC_ELT (*groups, i, grp) { tree c = grp->grp_end; - tree decl = OMP_CLAUSE_DECL (c); + tree decl = OMP_ITERATOR_CLAUSE_DECL (c); tree grp_end = grp->grp_end; auto_vec addr_tokens; tree sentinel = OMP_CLAUSE_CHAIN (grp_end); @@ -12091,6 +12276,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, tree c; tree *orig_list_p = list_p; int handled_depend_iterators = -1; + tree last_iterators = NULL_TREE; int nowait = -1; ctx = new_omp_context (region_type); @@ -12471,6 +12657,9 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, grpnum++; } decl = OMP_CLAUSE_DECL (c); + last_iterators = + OMP_ITERATOR_DECL_P (decl) ? TREE_PURPOSE (decl) : NULL_TREE; + OMP_CLAUSE_DECL (c) = decl = OMP_ITERATOR_CLAUSE_DECL (c); if (error_operand_p (decl)) { @@ -13388,6 +13577,9 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_REFERENCE)) remove = true; + if (last_iterators) + OMP_CLAUSE_DECL (c) = build_tree_list (last_iterators, + OMP_CLAUSE_DECL (c)); if (remove) *list_p = OMP_CLAUSE_CHAIN (c); else @@ -14106,7 +14298,7 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, break; case OMP_CLAUSE_MAP: - decl = OMP_CLAUSE_DECL (c); + decl = OMP_ITERATOR_CLAUSE_DECL (c); if (!grp_end) { grp_start_p = list_p; @@ -14168,7 +14360,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_ITERATOR_DECL_P (OMP_CLAUSE_DECL (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 +14529,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_ITERATOR_DECL_P (OMP_CLAUSE_DECL (c))) + break; + gimplify_omp_ctxp = ctx->outer_context; if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue, fb_lvalue) == GS_ERROR) @@ -14701,6 +14902,8 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, break; } + gimplify_omp_map_iterators (orig_list_p, pre_p); + gimplify_omp_ctxp = ctx->outer_context; delete_omp_context (ctx); } diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc index 4d003f42098..262990ecf9a 100644 --- a/gcc/omp-low.cc +++ b/gcc/omp-low.cc @@ -1518,7 +1518,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) case OMP_CLAUSE_MAP: if (ctx->outer) scan_omp_op (&OMP_CLAUSE_SIZE (c), ctx->outer); - decl = OMP_CLAUSE_DECL (c); + decl = OMP_ITERATOR_CLAUSE_DECL (c); /* If requested, make 'decl' addressable. */ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE (c)) @@ -12734,11 +12734,17 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) gcc_unreachable (); } #endif + if (OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (c)) + && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP) + { + map_cnt++; + continue; + } /* FALLTHRU */ case OMP_CLAUSE_TO: case OMP_CLAUSE_FROM: oacc_firstprivate: - var = OMP_CLAUSE_DECL (c); + var = OMP_ITERATOR_CLAUSE_DECL (c); if (!DECL_P (var)) { if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP @@ -13019,7 +13025,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) case OMP_CLAUSE_FROM: oacc_firstprivate_map: nc = c; - ovar = OMP_CLAUSE_DECL (c); + ovar = OMP_ITERATOR_CLAUSE_DECL (c); if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER || (OMP_CLAUSE_MAP_KIND (c) @@ -13039,6 +13045,11 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) { tree x = build_sender_ref (ovar, ctx); tree v = ovar; + if (OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (c))) + { + tree iterator = TREE_PURPOSE (OMP_CLAUSE_DECL (c)); + v = TREE_VEC_ELT (iterator, 1); + } if (in_reduction_clauses && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP && OMP_CLAUSE_MAP_IN_REDUCTION (c)) 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-iterator-1.c b/gcc/testsuite/c-c++-common/gomp/target-iterator-1.c new file mode 100644 index 00000000000..7d6c8dc6255 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/target-iterator-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-iterator-2.c b/gcc/testsuite/c-c++-common/gomp/target-iterator-2.c new file mode 100644 index 00000000000..da14d068f19 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/target-iterator-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-iterator-3.c b/gcc/testsuite/c-c++-common/gomp/target-iterator-3.c new file mode 100644 index 00000000000..22becdda559 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/target-iterator-3.c @@ -0,0 +1,20 @@ +/* { dg-do compile } */ +/* { dg-options "-fopenmp -fdump-tree-gimple" } */ + +#define DIM1 10 +#define DIM2 20 +#define DIM3 30 + +void f (int ***x, float ***y, double **z) +{ + #pragma omp target map(iterator(i=0:DIM1, j=0:DIM2), to: x[i][j][:DIM3], y[i][j][:DIM3]) \ + map(iterator(i=0:DIM1), from: z[i][:DIM2]) + ; +} + +/* { dg-final { scan-tree-dump-times "if \\(i <= 9\\) goto ; else goto ;" 2 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "if \\(j <= 19\\) goto ; else goto ;" 1 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "map\\(iterator\\(int i=0:10:1\\):iterator_array=D\.\[0-9\]+:from:" 1 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "map\\(iterator\\(int i=0:10:1\\):iterator_array=D\.\[0-9\]+:attach:" 1 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "map\\(iterator\\(int i=0:10:1, int j=0:20:1\\):iterator_array=D\.\[0-9\]+:to:" 2 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "map\\(iterator\\(int i=0:10:1, int j=0:20:1\\):iterator_array=D\.\[0-9\]+:attach:" 4 "gimple" } } */ diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc index ab7ecbfd1ef..21ecf94ada5 100644 --- a/gcc/tree-pretty-print.cc +++ b/gcc/tree-pretty-print.cc @@ -451,6 +451,19 @@ dump_omp_iterators (pretty_printer *pp, tree iter, int spc, dump_flags_t flags) pp_right_paren (pp); } +static void +dump_omp_map_iterators (pretty_printer *pp, tree iter, int spc, + dump_flags_t flags) +{ + if (TREE_VEC_LENGTH (iter) == 6) + dump_omp_iterators (pp, iter, spc, flags); + else + { + dump_omp_iterators (pp, TREE_VEC_ELT (iter, 0), spc, flags); + pp_string (pp, ":iterator_array="); + dump_generic_node (pp, TREE_VEC_ELT (iter, 1), spc, flags, false); + } +} /* Dump OMP clause CLAUSE, without following OMP_CLAUSE_CHAIN. @@ -461,6 +474,7 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) { const char *name; const char *modifier = NULL; + tree decl = NULL_TREE; switch (OMP_CLAUSE_CODE (clause)) { case OMP_CLAUSE_PRIVATE: @@ -911,6 +925,13 @@ 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,"); + decl = OMP_CLAUSE_DECL (clause); + if (OMP_ITERATOR_DECL_P (decl)) + { + dump_omp_map_iterators (pp, TREE_PURPOSE (decl), spc, flags); + pp_colon (pp); + decl = TREE_VALUE (decl); + } switch (OMP_CLAUSE_MAP_KIND (clause)) { case GOMP_MAP_ALLOC: @@ -1025,8 +1046,7 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) gcc_unreachable (); } pp_colon (pp); - dump_generic_node (pp, OMP_CLAUSE_DECL (clause), - spc, flags, false); + dump_generic_node (pp, decl, spc, flags, false); print_clause_size: if (OMP_CLAUSE_SIZE (clause)) { diff --git a/gcc/tree.h b/gcc/tree.h index 83075b82cc7..9fb21a95fbc 100644 --- a/gcc/tree.h +++ b/gcc/tree.h @@ -2162,6 +2162,13 @@ class auto_suppress_location_wrappers && TREE_PURPOSE (NODE) \ && TREE_CODE (TREE_PURPOSE (NODE)) == TREE_VEC) +/* Return the iterator expression if NODE contains an iterator. + Return the clause decl if NODE does not. */ +#define OMP_ITERATOR_CLAUSE_DECL(NODE) \ + (OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (NODE)) \ + ? TREE_VALUE (OMP_CLAUSE_DECL (NODE)) \ + : OMP_CLAUSE_DECL (NODE)) + /* In a BLOCK (scope) node: Variables declared in the scope NODE. */ #define BLOCK_VARS(NODE) (BLOCK_CHECK (NODE)->block.vars) diff --git a/libgomp/target.c b/libgomp/target.c index 47ec36928a6..37a4c539647 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -972,6 +972,77 @@ gomp_map_val (struct target_mem_desc *tgt, void **hostaddrs, size_t i) } } +/* 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. HOSTADDRS, SIZES and KINDS must be freed afterwards + if any merging occurs. */ + +static bool +gomp_merge_iterator_maps (size_t *mapnum, void ***hostaddrs, size_t **sizes, + void **kinds) +{ + 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: %ld -> %ld\n", + *mapnum, 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; + + for (size_t i = 0; i < *mapnum; i++) + { + if ((*sizes)[i] == SIZE_MAX) + { + uintptr_t *iterator_array = (*hostaddrs)[i]; + size_t count = iterator_array[0]; + for (int j = 1; j < count * 2 + 1; j += 2) + { + new_hostaddrs[new_idx] = (void *) iterator_array[j]; + new_sizes[new_idx] = iterator_array[j+1]; + new_kinds[new_idx] = (*skinds)[i]; + gomp_debug (1, + "Expanding map %ld: " + "hostaddrs[%ld] = %p, sizes[%ld] = %ld\n", + i, new_idx, new_hostaddrs[new_idx], + new_idx, 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]; + 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, @@ -988,6 +1059,10 @@ 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; + if (short_mapkind) + iterators_p = gomp_merge_iterator_maps (&mapnum, &hostaddrs, &sizes, + &kinds); struct target_mem_desc *tgt = gomp_malloc (sizeof (*tgt) + sizeof (tgt->list[0]) * mapnum); tgt->list_count = mapnum; @@ -1876,6 +1951,14 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, } gomp_mutex_unlock (&devicep->lock); + + if (iterators_p) + { + free (hostaddrs); + free (sizes); + free (kinds); + } + 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..900a0ba2d64 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-1.c @@ -0,0 +1,44 @@ +/* { dg-do run } */ + +/* 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 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..bad0f7f17b8 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-2.c @@ -0,0 +1,42 @@ +/* { dg-do run } */ + +/* 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 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..e3da479e6cb --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-3.c @@ -0,0 +1,54 @@ +/* { dg-do run } */ + +/* 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 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; +}