From patchwork Fri Oct 4 14:56:49 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: 1992780 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=kqltwzNP; 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 4XKsBR4WS9z1xt7 for ; Sat, 5 Oct 2024 00:58:55 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id C4D953842FFF for ; Fri, 4 Oct 2024 14:58:53 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from mail-lj1-x230.google.com (mail-lj1-x230.google.com [IPv6:2a00:1450:4864:20::230]) by sourceware.org (Postfix) with ESMTPS id 013743842AFC for ; Fri, 4 Oct 2024 14:57:18 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 013743842AFC 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 013743842AFC Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=2a00:1450:4864:20::230 ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1728053842; cv=none; b=Wq3DvO4smRvNnlzpmkjGQJY2sT029VLUmsoBJYwi9NpvL50UmmRVzrTdVigPdqtbU4F3U6HIyaaUnn8SwoIqFLtDk2C4lk1rLdBc61AkhW1cFzbhwr1MnYZXxiTBSK6WYayQllYtULpeuwZslIhq2OPr73CPP8IVvMtROEMRAGs= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1728053842; c=relaxed/simple; bh=4hcac8RTOycRjQ/ypLGqNOHJ0PQobwI3Z9lCl0I5uyM=; h=DKIM-Signature:Message-ID:Date:MIME-Version:To:Subject:From; b=l6M0Y3UIiBjcoqH6N7xvuN8+Z2AxEy+e+lalGR72j83PYSIs84zcnha3TawrlA+gVCTb2LDYdLMY5w1Yrj5vhw/9SnOkh/Tkb6rgQIgDDRPxXVJqXF5ijbqFzfq8G2dGnKYJ5DDQ035OTcfny9cCq925oIf4d7KBoaVZINpZJZI= ARC-Authentication-Results: i=1; server2.sourceware.org Received: by mail-lj1-x230.google.com with SMTP id 38308e7fff4ca-2fad100dd9fso35405881fa.3 for ; Fri, 04 Oct 2024 07:57:17 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre-com.20230601.gappssmtp.com; s=20230601; t=1728053836; x=1728658636; 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=4hcac8RTOycRjQ/ypLGqNOHJ0PQobwI3Z9lCl0I5uyM=; b=kqltwzNPR2DZXMNC++f/IkiJ+gH6fGCGWpFI4Y1//68kLdyiOVUxxzQPkNxI+COWDQ HF77QXHRNOq/0q5bKmY/1w0yd9V2zucSmqbqw/27WwJKcfU0Jp5Bj6sXuHu2sj0oU0iN V5MIgxer//rrGhMXoCj1kKOQ18FHDFHTsmUu60IOyTfX7eDhd3hlL2V9+DyUmIWH+tCh GMnEd+135VmvzLaqnC+CzgmoE8ihorTWRSIk+ayydHs9vs8jO4b0G7oAteNiw8HDjXrV V0liB/tpYnNkTCp3xDG/CbVYIYfpAG6wsk3m0pn+tc1sLy6KoGNYMhaKqKxVH+pk3Ube bUIw== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1728053836; x=1728658636; 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=4hcac8RTOycRjQ/ypLGqNOHJ0PQobwI3Z9lCl0I5uyM=; b=IvvAmP288xqNnKXEqX8VP9DgraR9ZeyHBgR2FFuzX/VOlXqJKB74MQSCP2sSaYvy9L paHu2ktDzF+Z3Y6CYgn1oOuOKBh1WafcpbekJr4akCGX7+ReSOYszTNC2e7UkD4LP/sS 5inDqqEEbCnbAsnWCiir1galXBgF88d5vAiypVv+D0Dp8+SQrarKpFn+01BkEpZLMhoY qnnQL0gFJYXf/EKZcPPanx7T1s7CnO/hUicb/HuAFdovp9ZpTGXhbVAVawUGVlfzeiR9 TOTi4U183lCeTR8LMHb9rVR0KvTeLoT09clePn9YZjYANBDIDe7faSciYcBKKHD74F99 nqsg== X-Gm-Message-State: AOJu0YwL6B704bQvUCjlz9+4dK4OnkT0Gv+7uy0BSzurX9tlWFdZDnpD 1hYIOtLr/KZFt7GiDILgl2RINmMMtWxyAtGncVUQpcl4OxAUTQifP3OlpeuiqpM3zaC8NAKJV53 h X-Google-Smtp-Source: AGHT+IGsY08kdDCjBpDmbklHQnGPqzEI88uIF5M7n50urOAbAlf//HBmNhpV9ycm+n4FJdZIkwIWVQ== X-Received: by 2002:a2e:878a:0:b0:2f7:5a41:b0b with SMTP id 38308e7fff4ca-2faf3c44dc5mr23111241fa.26.1728053835973; Fri, 04 Oct 2024 07:57:15 -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-a992e5bbc2csm2048766b.35.2024.10.04.07.57.15 (version=TLS1_3 cipher=TLS_AES_128_GCM_SHA256 bits=128/128); Fri, 04 Oct 2024 07:57:15 -0700 (PDT) Message-ID: Date: Fri, 4 Oct 2024 15:56:49 +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 3/5] openmp: Add support for iterators in 'target update' 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.7 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 extends the previous patch to cover to/from clauses in 'target update'. From 1c8bf84ec99fe2fd371e345f012eb0d84a923153 Mon Sep 17 00:00:00 2001 From: Kwok Cheung Yeung Date: Fri, 4 Oct 2024 15:16:21 +0100 Subject: [PATCH 3/5] openmp: Add support for iterators in 'target update' clauses (C/C++) This adds support for iterators in 'to' and 'from' clauses in the 'target update' OpenMP directive. 2024-10-04 Kwok Cheung Yeung gcc/c/ * c-parser.cc (c_parser_omp_clause_from_to): Parse 'iterator' modifier. * c-typeck.cc (c_finish_omp_clauses): Finish iterators for to/from clauses. gcc/cp/ * parser.cc (cp_parser_omp_clause_from_to): Parse 'iterator' modifier. * semantics.cc (finish_omp_clauses): Finish iterators for to/from clauses. gcc/ * gimplify.cc (gimplify_scan_omp_clauses): Call check_omp_map_iterators on clauses with iterators. Skip gimplification of clause decl and size for clauses with iterators. * omp-low.cc (lower_omp_target): Call lower_omp_map_iterators on to/from clauses. * tree-pretty-print.cc (dump_omp_clause): Call dump_omp_iterators for to/from clauses with iterators. * tree.cc (omp_clause_num_ops): Add extra operand for OMP_CLAUSE_FROM and OMP_CLAUSE_TO. * tree.h (OMP_CLAUSE_HAS_ITERATORS): Add check for OMP_CLAUSE_TO and OMP_CLAUSE_FROM. (OMP_CLAUSE_ITERATORS): Likewise. gcc/testsuite/ * c-c++-common/gomp/target-update-iterators-1.c: New. * c-c++-common/gomp/target-update-iterators-2.c: New. * c-c++-common/gomp/target-update-iterators-3.c: New. libgomp/ * target.c (gomp_update): Call gomp_merge_iterator_maps. Free allocated variables. * testsuite/libgomp.c-c++-common/target-update-iterators-1.c: New. * testsuite/libgomp.c-c++-common/target-update-iterators-2.c: New. * testsuite/libgomp.c-c++-common/target-update-iterators-3.c: New. --- gcc/c/c-parser.cc | 105 +++++++++++++++-- gcc/c/c-typeck.cc | 5 +- gcc/cp/parser.cc | 111 ++++++++++++++++-- gcc/cp/semantics.cc | 5 +- gcc/gimplify.cc | 18 ++- gcc/omp-low.cc | 3 +- .../gomp/target-update-iterators-1.c | 20 ++++ .../gomp/target-update-iterators-2.c | 17 +++ .../gomp/target-update-iterators-3.c | 17 +++ gcc/tree-pretty-print.cc | 10 ++ gcc/tree.cc | 4 +- gcc/tree.h | 8 +- libgomp/target.c | 14 +++ .../target-update-iterators-1.c | 65 ++++++++++ .../target-update-iterators-2.c | 58 +++++++++ .../target-update-iterators-3.c | 67 +++++++++++ 16 files changed, 496 insertions(+), 31 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/gomp/target-update-iterators-1.c create mode 100644 gcc/testsuite/c-c++-common/gomp/target-update-iterators-2.c create mode 100644 gcc/testsuite/c-c++-common/gomp/target-update-iterators-3.c create mode 100644 libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-1.c create mode 100644 libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-2.c create mode 100644 libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-3.c diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc index 184fc076388..c2a5985c89b 100644 --- a/gcc/c/c-parser.cc +++ b/gcc/c/c-parser.cc @@ -19304,8 +19304,11 @@ c_parser_omp_clause_device_type (c_parser *parser, tree list) to ( variable-list ) OpenMP 5.1: - from ( [present :] variable-list ) - to ( [present :] variable-list ) */ + from ( [motion-modifier[,] [motion-modifier[,]...]:] variable-list ) + to ( [motion-modifier[,] [motion-modifier[,]...]:] variable-list ) + + motion-modifier: + present | iterator (iterators-definition) */ static tree c_parser_omp_clause_from_to (c_parser *parser, enum omp_clause_code kind, @@ -19316,15 +19319,88 @@ c_parser_omp_clause_from_to (c_parser *parser, enum omp_clause_code kind, if (!parens.require_open (parser)) return list; + int pos = 1, colon_pos = 0; + int iterator_length = 0; + while (c_parser_peek_nth_token_raw (parser, pos)->type == CPP_NAME) + { + if (c_parser_peek_nth_token_raw (parser, pos + 1)->type + == CPP_OPEN_PAREN) + { + unsigned int 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; + } + } + if (c_parser_peek_nth_token_raw (parser, pos + 1)->type == CPP_COMMA) + pos += 2; + else + pos++; + if (c_parser_peek_nth_token_raw (parser, pos)->type == CPP_COLON) + { + colon_pos = pos; + break; + } + } + bool present = false; - c_token *token = c_parser_peek_token (parser); + tree iterators = NULL_TREE; - if (token->type == CPP_NAME - && strcmp (IDENTIFIER_POINTER (token->value), "present") == 0 - && c_parser_peek_2nd_token (parser)->type == CPP_COLON) + for (pos = 1; pos < colon_pos; pos++) { - present = true; - c_parser_consume_token (parser); + c_token *token = c_parser_peek_token (parser); + + if (token->type == CPP_COMMA) + { + c_parser_consume_token (parser); + continue; + } + if (token->type == CPP_NAME) + { + const char *name = IDENTIFIER_POINTER (token->value); + if (strcmp (name, "present") == 0) + { + if (present) + { + c_parser_error (parser, "too many % modifiers"); + parens.skip_until_found_close (parser); + return list; + } + present = true; + c_parser_consume_token (parser); + } + else if (strcmp (name, "iterator") == 0) + { + 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; + } + else + { + if (kind == OMP_CLAUSE_TO) + c_parser_error (parser, "% clause with motion modifier " + "other than % or %"); + else + c_parser_error (parser, "% clause with motion modifier " + "other than % or %"); + parens.skip_until_found_close (parser); + return list; + } + } + } + + if (colon_pos) + { + gcc_assert (pos == colon_pos); + gcc_assert (c_parser_next_token_is (parser, CPP_COLON)); c_parser_consume_token (parser); } @@ -19335,6 +19411,19 @@ c_parser_omp_clause_from_to (c_parser *parser, enum omp_clause_code kind, for (tree c = nl; c != list; c = OMP_CLAUSE_CHAIN (c)) OMP_CLAUSE_MOTION_PRESENT (c) = 1; + if (iterators) + { + tree block = pop_scope (); + if (iterators == error_mark_node) + iterators = NULL_TREE; + else + TREE_VEC_ELT (iterators, 5) = block; + } + + if (iterators) + for (tree c = nl; c != list; c = OMP_CLAUSE_CHAIN (c)) + OMP_CLAUSE_ITERATORS (c) = iterators; + return nl; } diff --git a/gcc/c/c-typeck.cc b/gcc/c/c-typeck.cc index cca9f1c000c..5d16f749133 100644 --- a/gcc/c/c-typeck.cc +++ b/gcc/c/c-typeck.cc @@ -15813,6 +15813,9 @@ 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; + /* FALLTHRU */ + case OMP_CLAUSE_TO: + case OMP_CLAUSE_FROM: if (OMP_CLAUSE_ITERATORS (c) && c_omp_finish_iterators (OMP_CLAUSE_ITERATORS (c))) { @@ -15820,8 +15823,6 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) break; } /* FALLTHRU */ - case OMP_CLAUSE_TO: - case OMP_CLAUSE_FROM: case OMP_CLAUSE__CACHE_: { using namespace omp_addr_tokenizer; diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc index 79d6e115d16..861337803e6 100644 --- a/gcc/cp/parser.cc +++ b/gcc/cp/parser.cc @@ -41716,8 +41716,11 @@ cp_parser_omp_clause_doacross (cp_parser *parser, tree list, location_t loc) to ( variable-list ) OpenMP 5.1: - from ( [present :] variable-list ) - to ( [present :] variable-list ) */ + from ( [motion-modifier[,] [motion-modifier[,]...]:] variable-list ) + to ( [motion-modifier[,] [motion-modifier[,]...]:] variable-list ) + + motion-modifier: + present | iterator (iterators-definition) */ static tree cp_parser_omp_clause_from_to (cp_parser *parser, enum omp_clause_code kind, @@ -41726,15 +41729,94 @@ cp_parser_omp_clause_from_to (cp_parser *parser, enum omp_clause_code kind, if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN)) return list; + size_t pos = 1, colon_pos = 0; + int iterator_length = 0; + while (cp_lexer_nth_token_is (parser->lexer, pos, CPP_NAME)) + { + if (cp_lexer_nth_token_is (parser->lexer, pos + 1, CPP_OPEN_PAREN)) + { + unsigned int n = cp_parser_skip_balanced_tokens (parser, pos + 1); + if (n != pos + 1) + { + iterator_length = n - pos; + pos = n - 1; + } + } + if (cp_lexer_nth_token_is (parser->lexer, pos + 1, CPP_COMMA)) + pos += 2; + else + pos++; + if (cp_lexer_nth_token_is (parser->lexer, pos, CPP_COLON)) + { + colon_pos = pos; + break; + } + } + bool present = false; - cp_token *token = cp_lexer_peek_token (parser->lexer); + tree iterators = NULL_TREE; + for (pos = 1; pos < colon_pos; pos++) + { + cp_token *token = cp_lexer_peek_token (parser->lexer); - if (token->type == CPP_NAME - && strcmp (IDENTIFIER_POINTER (token->u.value), "present") == 0 - && cp_lexer_nth_token_is (parser->lexer, 2, CPP_COLON)) + if (token->type == CPP_COMMA) + { + cp_lexer_consume_token (parser->lexer); + continue; + } + if (token->type == CPP_NAME) + { + const char *name = IDENTIFIER_POINTER (token->u.value); + if (strcmp (name, "present") == 0) + { + if (present) + { + cp_parser_error (parser, "too many % modifiers"); + cp_parser_skip_to_closing_parenthesis (parser, + /*recovering=*/true, + /*or_comma=*/false, + /*consume_paren=*/true); + return list; + } + present = true; + cp_lexer_consume_token (parser->lexer); + } + else if (strcmp (name, "iterator") == 0) + { + 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; + } + else + { + if (kind == OMP_CLAUSE_TO) + cp_parser_error (parser, "% clause with motion modifier " + "other than % or %"); + else + cp_parser_error (parser, "% clause with motion modifier " + "other than % or %"); + cp_parser_skip_to_closing_parenthesis (parser, + /*recovering=*/true, + /*or_comma=*/false, + /*consume_paren=*/true); + return list; + } + } + } + + if (colon_pos) { - present = true; - cp_lexer_consume_token (parser->lexer); + gcc_assert (pos == colon_pos); + gcc_assert (cp_lexer_next_token_is (parser->lexer, CPP_COLON)); cp_lexer_consume_token (parser->lexer); } @@ -41743,6 +41825,19 @@ cp_parser_omp_clause_from_to (cp_parser *parser, enum omp_clause_code kind, for (tree c = nl; c != list; c = OMP_CLAUSE_CHAIN (c)) OMP_CLAUSE_MOTION_PRESENT (c) = 1; + if (iterators) + { + tree block = poplevel (1, 1, 0); + if (iterators == error_mark_node) + iterators = NULL_TREE; + else + TREE_VEC_ELT (iterators, 5) = block; + } + + if (iterators) + for (tree c = nl; c != list; c = OMP_CLAUSE_CHAIN (c)) + OMP_CLAUSE_ITERATORS (c) = iterators; + return nl; } diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc index ba5657f7bc2..fdf814d3ce6 100644 --- a/gcc/cp/semantics.cc +++ b/gcc/cp/semantics.cc @@ -8492,6 +8492,9 @@ 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; + /* FALLTHRU */ + case OMP_CLAUSE_TO: + case OMP_CLAUSE_FROM: if (OMP_CLAUSE_ITERATORS (c) && cp_omp_finish_iterators (OMP_CLAUSE_ITERATORS (c))) { @@ -8499,8 +8502,6 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) break; } /* FALLTHRU */ - case OMP_CLAUSE_TO: - case OMP_CLAUSE_FROM: case OMP_CLAUSE__CACHE_: { using namespace omp_addr_tokenizer; diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index 6e532d07fcf..ba972a2892a 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -12862,6 +12862,12 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, case OMP_CLAUSE_TO: case OMP_CLAUSE_FROM: + if (OMP_CLAUSE_ITERATORS (c) && !check_omp_map_iterators (c)) + { + remove = true; + break; + } + /* FALLTHRU */ case OMP_CLAUSE__CACHE_: decl = OMP_CLAUSE_DECL (c); if (error_operand_p (decl)) @@ -12872,17 +12878,19 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, if (OMP_CLAUSE_SIZE (c) == NULL_TREE) OMP_CLAUSE_SIZE (c) = DECL_P (decl) ? DECL_SIZE_UNIT (decl) : TYPE_SIZE_UNIT (TREE_TYPE (decl)); - if (gimplify_expr (&OMP_CLAUSE_SIZE (c), pre_p, - NULL, is_gimple_val, fb_rvalue) == GS_ERROR) + if (!(OMP_CLAUSE_HAS_ITERATORS (c) && OMP_CLAUSE_ITERATORS (c)) + && gimplify_expr (&OMP_CLAUSE_SIZE (c), pre_p, + NULL, is_gimple_val, fb_rvalue) == GS_ERROR) { remove = true; break; } if (!DECL_P (decl)) { - if (gimplify_expr (&OMP_CLAUSE_DECL (c), pre_p, - NULL, is_gimple_lvalue, fb_lvalue) - == GS_ERROR) + if (!(OMP_CLAUSE_HAS_ITERATORS (c) && OMP_CLAUSE_ITERATORS (c)) + && gimplify_expr (&OMP_CLAUSE_DECL (c), pre_p, + NULL, is_gimple_lvalue, fb_lvalue) + == GS_ERROR) { remove = true; break; diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc index 9cf6e207d1c..a8b86889c66 100644 --- a/gcc/omp-low.cc +++ b/gcc/omp-low.cc @@ -12921,10 +12921,11 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) gcc_unreachable (); } #endif - lower_omp_map_iterators (c, &iterlist, &iterator_loops); /* FALLTHRU */ case OMP_CLAUSE_TO: case OMP_CLAUSE_FROM: + lower_omp_map_iterators (c, &iterlist, &iterator_loops); + /* FALLTHRU */ oacc_firstprivate: var = OMP_CLAUSE_DECL (c); { diff --git a/gcc/testsuite/c-c++-common/gomp/target-update-iterators-1.c b/gcc/testsuite/c-c++-common/gomp/target-update-iterators-1.c new file mode 100644 index 00000000000..3a64f511da4 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/target-update-iterators-1.c @@ -0,0 +1,20 @@ +/* { dg-do compile } */ +/* { dg-options "-fopenmp" } */ + +#define DIM1 17 +#define DIM2 39 + +void f (int **x, float **y) +{ + #pragma omp target update to (iterator(i=0:DIM1): x[i][:DIM2]) + + #pragma omp target update to (iterator(i=0:DIM1): x[i][:DIM2], y[i][:DIM2]) + + #pragma omp target update to (iterator(i=0:DIM1), present: x[i][:DIM2]) + + #pragma omp target update to (iterator(i=0:DIM1), iterator(j=0:DIM2): x[i][j]) /* { dg-error "too many 'iterator' modifiers" } */ + /* { dg-error ".#pragma omp target update. must contain at least one .from. or .to. clauses" "" { target *-*-* } .-1 } */ + + #pragma omp target update to (iterator(i=0:DIM1), something: x[i][j]) /* { dg-error ".to. clause with motion modifier other than .iterator. or .present. before .something." } */ + /* { dg-error ".#pragma omp target update. must contain at least one .from. or .to. clauses" "" { target *-*-* } .-1 } */ +} diff --git a/gcc/testsuite/c-c++-common/gomp/target-update-iterators-2.c b/gcc/testsuite/c-c++-common/gomp/target-update-iterators-2.c new file mode 100644 index 00000000000..3789a559b6f --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/target-update-iterators-2.c @@ -0,0 +1,17 @@ +/* { dg-do compile } */ +/* { dg-options "-fopenmp" } */ + +void f (int *x, float *y, double *z) +{ + #pragma omp target update to(iterator(i=0:10): x) /* { dg-error "iterator variable .i. not used in clause expression" }*/ + ; + + #pragma omp target update from(iterator(i=0:10, j=0:20): x[i]) /* { dg-error "iterator variable .j. not used in clause expression" }*/ + ; + + #pragma omp target update to(iterator(i=0:10, j=0:20, k=0:30): 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-update-iterators-3.c b/gcc/testsuite/c-c++-common/gomp/target-update-iterators-3.c new file mode 100644 index 00000000000..b256674442f --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/target-update-iterators-3.c @@ -0,0 +1,17 @@ +/* { 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 update to (iterator(i=0:DIM1, j=0:DIM2): x[i][j][:DIM3], y[i][j][:DIM3]) + #pragma omp target update from (iterator(i=0:DIM1): 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 "to\\(iterator\\(int i=0:10:1, int j=0:20:1\\):D\.\[0-9\]+" 2 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "from\\(iterator\\(int i=0:10:1\\):D\.\[0-9\]+" 1 "omplower" } } */ diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc index be2723dcdae..fa1b2dce27f 100644 --- a/gcc/tree-pretty-print.cc +++ b/gcc/tree-pretty-print.cc @@ -1084,6 +1084,11 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) pp_string (pp, "from("); if (OMP_CLAUSE_MOTION_PRESENT (clause)) pp_string (pp, "present:"); + if (OMP_CLAUSE_ITERATORS (clause)) + { + dump_omp_iterators (pp, OMP_CLAUSE_ITERATORS (clause), spc, flags); + pp_colon (pp); + } dump_generic_node (pp, OMP_CLAUSE_DECL (clause), spc, flags, false); goto print_clause_size; @@ -1092,6 +1097,11 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) pp_string (pp, "to("); if (OMP_CLAUSE_MOTION_PRESENT (clause)) pp_string (pp, "present:"); + if (OMP_CLAUSE_ITERATORS (clause)) + { + dump_omp_iterators (pp, OMP_CLAUSE_ITERATORS (clause), spc, flags); + pp_colon (pp); + } dump_generic_node (pp, OMP_CLAUSE_DECL (clause), spc, flags, false); goto print_clause_size; diff --git a/gcc/tree.cc b/gcc/tree.cc index f12d7b8bb8a..5bba8df4518 100644 --- a/gcc/tree.cc +++ b/gcc/tree.cc @@ -264,8 +264,8 @@ unsigned const char omp_clause_num_ops[] = 1, /* OMP_CLAUSE_IS_DEVICE_PTR */ 1, /* OMP_CLAUSE_INCLUSIVE */ 1, /* OMP_CLAUSE_EXCLUSIVE */ - 2, /* OMP_CLAUSE_FROM */ - 2, /* OMP_CLAUSE_TO */ + 3, /* OMP_CLAUSE_FROM */ + 3, /* OMP_CLAUSE_TO */ 3, /* OMP_CLAUSE_MAP */ 1, /* OMP_CLAUSE_HAS_DEVICE_ADDR */ 1, /* OMP_CLAUSE_DOACROSS */ diff --git a/gcc/tree.h b/gcc/tree.h index 384a5f1f250..a57419447a1 100644 --- a/gcc/tree.h +++ b/gcc/tree.h @@ -1619,11 +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_HAS_ITERATORS(NODE) \ + (OMP_CLAUSE_CODE (NODE) == OMP_CLAUSE_FROM \ + || OMP_CLAUSE_CODE (NODE) == OMP_CLAUSE_TO \ + || 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_FROM, \ OMP_CLAUSE_MAP), 2) /* True on OMP_FOR and other OpenMP/OpenACC looping constructs if the loop nest diff --git a/libgomp/target.c b/libgomp/target.c index 463a162879b..60d57a19dd0 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -2260,6 +2260,8 @@ gomp_update (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs, size_t i; struct splay_tree_key_s cur_node; const int typemask = short_mapkind ? 0xff : 0x7; + bool iterators_p = false; + size_t *iterator_count = NULL; if (!devicep) return; @@ -2267,6 +2269,10 @@ gomp_update (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs, if (mapnum == 0) return; + if (short_mapkind) + iterators_p = gomp_merge_iterator_maps (&mapnum, &hostaddrs, &sizes, + &kinds, &iterator_count); + gomp_mutex_lock (&devicep->lock); if (devicep->state == GOMP_DEVICE_FINALIZED) { @@ -2360,6 +2366,14 @@ gomp_update (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs, } } gomp_mutex_unlock (&devicep->lock); + + if (iterators_p) + { + free (hostaddrs); + free (sizes); + free (kinds); + free (iterator_count); + } } static struct gomp_offload_icv_list * diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-1.c b/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-1.c new file mode 100644 index 00000000000..5a4cad5c219 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-1.c @@ -0,0 +1,65 @@ +/* { dg-do run } */ + +/* Test target enter data and target update to the 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 sum; + int expected = mkarray (x); + + #pragma omp target enter data map(to: x[:DIM1]) + #pragma omp target enter data map(iterator(i=0:DIM1), to: x[i][:DIM2]) + #pragma omp target map(from: sum) + { + sum = 0; + for (int i = 0; i < DIM1; i++) + for (int j = 0; j < DIM2; j++) + sum += x[i][j]; + } + + if (sum != expected) + return 1; + + expected = 0; + for (int i = 0; i < DIM1; i++) + for (int j = 0; j < DIM2; j++) + { + x[i][j] *= rand (); + expected += x[i][j]; + } + + #pragma omp target update to(iterator(i=0:DIM1): x[i][:DIM2]) + + #pragma omp target map(from: sum) + { + sum = 0; + for (int i = 0; i < DIM1; i++) + for (int j = 0; j < DIM2; j++) + sum += x[i][j]; + } + + return sum != expected; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-2.c b/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-2.c new file mode 100644 index 00000000000..93438d01c97 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-2.c @@ -0,0 +1,58 @@ +/* { dg-do run } */ +/* { dg-require-effective-target offload_device_nonshared_as } */ + +/* Test target enter data and target update from the 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)); + for (int j = 0; j < DIM2; j++) + x[i][j] = 0; + } +} + +int main (void) +{ + int *x[DIM1]; + int sum, expected; + + mkarray (x); + + #pragma omp target enter data map(alloc: x[:DIM1]) + #pragma omp target enter data map(iterator(i=0:DIM1), to: x[i][:DIM2]) + #pragma omp target 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 + 2); + expected += x[i][j]; + } + } + + /* Host copy of x should remain unchanged. */ + sum = 0; + for (int i = 0; i < DIM1; i++) + for (int j = 0; j < DIM2; j++) + sum += x[i][j]; + if (sum != 0) + return 1; + + #pragma omp target update from(iterator(i=0:DIM1): x[i][:DIM2]) + + /* Host copy should now be updated. */ + sum = 0; + for (int i = 0; i < DIM1; i++) + for (int j = 0; j < DIM2; j++) + sum += x[i][j]; + return sum - expected; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-3.c b/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-3.c new file mode 100644 index 00000000000..a70b21c4b75 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-3.c @@ -0,0 +1,67 @@ +/* { dg-do run } */ +/* { dg-require-effective-target offload_device_nonshared_as } */ + +/* Test target enter data and target update to the target using map + iterators with a function. */ + +#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)); + for (int j = 0; j < DIM2; j++) + x[i][j] = rand (); + } +} + +int f (int i) +{ + return i * 2; +} + +int main (void) +{ + int *x[DIM1], x_new[DIM1][DIM2]; + int sum, expected; + + mkarray (x); + + #pragma omp target enter data map(alloc: x[:DIM1]) + #pragma omp target enter data map(iterator(i=0:DIM1), to: x[i][:DIM2]) + + /* Update x on host. */ + for (int i = 0; i < DIM1; i++) + for (int j = 0; j < DIM2; j++) + { + x_new[i][j] = x[i][j]; + x[i][j] = (i + 1) * (j + 2); + } + + /* Update a subset of x on target. */ + #pragma omp target update to(iterator(i=0:DIM1/2): x[f (i)][:DIM2]) + + #pragma omp target map(from: sum) + { + sum = 0; + for (int i = 0; i < DIM1; i++) + for (int j = 0; j < DIM2; j++) + sum += x[i][j]; + } + + /* Calculate expected value on host. */ + for (int i = 0; i < DIM1/2; i++) + for (int j = 0; j < DIM2; j++) + x_new[f (i)][j] = x[f (i)][j]; + + expected = 0; + for (int i = 0; i < DIM1; i++) + for (int j = 0; j < DIM2; j++) + expected += x_new[i][j]; + + return sum - expected; +}