From patchwork Tue Sep 3 17:06:06 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: 1980229 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=t8bBuLpB; 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 4WysWT73Qcz1yZ9 for ; Wed, 4 Sep 2024 03:07:49 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 1D88B384F01B for ; Tue, 3 Sep 2024 17:07:48 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from mail-wm1-x335.google.com (mail-wm1-x335.google.com [IPv6:2a00:1450:4864:20::335]) by sourceware.org (Postfix) with ESMTPS id 5F7D0385E027 for ; Tue, 3 Sep 2024 17:07:23 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 5F7D0385E027 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 5F7D0385E027 Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=2a00:1450:4864:20::335 ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1725383247; cv=none; b=bwtqR8cinY10uhC5eiCqMBvUgJ4JRblD+6mvX0/L4caoUN2Q+f8c5Pqmg6GgnzmOudIshtWcKVKDWFNb+SWND/LJMWoRPzYzUcB6sZ6YYHUeUclE+Y/9ujsAnT2g5u1otfFyQx+lSMUy6UcbdzNkYO4mfxojj5pQGHFqkig26nU= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1725383247; c=relaxed/simple; bh=Mv9LA0pDjMaTIDurBNFE8M9vuiSw0VU+Di6vOCWI4wM=; h=DKIM-Signature:Message-ID:Date:MIME-Version:To:Subject:From; b=X1V30eSmx+VhFKvHVJRYphtcJxgD0P99fsYaSOu7OHplc6t+mAwtUYxxYATKqspM0+XjrNAuv+GWntGBTp5UoE9f9uJBcJSraicdLhudhoNFCd6LBLFU52z9TDoUBoobqs+5zUxZB39Gh1kaPR1ERR+cJqy5wsBT/LyKs4E78PY= ARC-Authentication-Results: i=1; server2.sourceware.org Received: by mail-wm1-x335.google.com with SMTP id 5b1f17b1804b1-42c7b5b2d01so35979845e9.3 for ; Tue, 03 Sep 2024 10:07:23 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre-com.20230601.gappssmtp.com; s=20230601; t=1725383242; x=1725988042; 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=Mv9LA0pDjMaTIDurBNFE8M9vuiSw0VU+Di6vOCWI4wM=; b=t8bBuLpBX53VPolG00DaAdABKXWStG6iUgdK4IZH6zttcqormkVBF9lLQOXxFar+is jKNpFp7P6TkzQsbVtW2+B/69fIqi7OniT7gabdGwkdoJQnsms5QWOF+LQ8PO97+P+d8G rrIXO5xPHtZGe+SP8FXw7/Ezr5io3VsTX6KgmXy/ElNpJiQp4z6IrilG6yOsmu0LrN+Y p+aFLcgLNevL7sDualoAwc+6v8KCmABjSJYAcO98ie8O0SymOiS0DafoPGqJMGmWP4Lk sMRUVDxP8XbxhoXYkGYcx2a+n0G5zMdN4ueq3rBZEPeZbvAk2TOtGycC29QifwvqB3yQ 2Iow== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1725383242; x=1725988042; 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=Mv9LA0pDjMaTIDurBNFE8M9vuiSw0VU+Di6vOCWI4wM=; b=I6WbkjrpgW17dO4lDFZ7a+oAuS59qnJgRk74EQ4uSus72jEPuUdmDf9b/djUKVzLmL S9MDIXjeRnM5lB4Qw/0P4YAqvyCoNAsFcjmYL2hIXBK1YWN2XajvXXzlz/+Q9LH8Q8wt gydYAwtQ5TMoCePH1nEGPxjgFCV0RymW2PvkcBf3uDJrUc+Ft2STKrTxuRn7V36SOj6v 9kKFBYkpDDx1w70dDzCL335JPqZJOqyz0t3COnQiySCA1hMrAwkWZSnaX1ObJnk/kwlB cPTOTrd7ROm6gqP6AUBMQ7Zevfbj+Re76DOPiI6shd6cW9NZIs2BUhegIQAO5OqQXdW1 e0Sg== X-Gm-Message-State: AOJu0Yxp5pI/1pQ2QleUM64v0PbiBWCIlwE8sqeVo08gpfmyvEl3YrHn Ayj4ssstCaER5mpYGxPe1vl7VX6GP6xKEvmpRPMetlq6pEDDnZpzFbOoZGv3+BHCA9zN9Z/4YrA c X-Google-Smtp-Source: AGHT+IHOKIkAxmDN81MgaxBM2oZ/spVJPDkITgGR0ybJhnBszuB1ImyJo8SyZRESYrXhgDEcCXGXKg== X-Received: by 2002:adf:f550:0:b0:374:bd01:707c with SMTP id ffacd0b85a97d-374bd0170famr8216408f8f.48.1725383241297; Tue, 03 Sep 2024 10:07:21 -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-a8988feae4dsm708932866b.31.2024.09.03.10.07.03 (version=TLS1_3 cipher=TLS_AES_128_GCM_SHA256 bits=128/128); Tue, 03 Sep 2024 10:07:11 -0700 (PDT) Message-ID: <86875a51-8c1b-4d28-a75e-4e0a080bb1c2@baylibre.com> Date: Tue, 3 Sep 2024 18:06:06 +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 1/5] openmp: Refactor handling of iterators Content-Language: en-GB From: Kwok Cheung Yeung In-Reply-To: <77f9ccb8-6f5e-4462-aa32-71f74fd7ff26@baylibre.com> X-Spam-Status: No, score=-12.2 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 factors out the code to calculate the number of iterations required and to generate the iteration loop into separate functions from gimplify_omp_depend for reuse later. I have also replaced the 'TREE_CODE (*tp) == TREE_LIST && ...' checks used for detecting an iterator clause with a macro OMP_ITERATOR_DECL_P, as it needs to be done frequently. From d2cf47a312d9decc14d0cf37fa57ad358a96743d Mon Sep 17 00:00:00 2001 From: Kwok Cheung Yeung Date: Mon, 2 Sep 2024 19:33:08 +0100 Subject: [PATCH 1/5] openmp: Refactor handling of iterators Move code to calculate the iteration size and to generate the iterator expansion loop into separate functions. Use OMP_ITERATOR_DECL_P to check for iterators in clause declarations. 2024-09-02 Kwok Cheung Yeung gcc/c-family/ * c-omp.cc (c_finish_omp_depobj): Use OMP_ITERATOR_DECL_P. gcc/c/ * c-typeck.cc (handle_omp_array_sections): Use OMP_ITERATOR_DECL_P. (c_finish_omp_clauses): Likewise. gcc/cp/ * pt.cc (tsubst_omp_clause_decl): Use OMP_ITERATOR_DECL_P. * semantics.cc (handle_omp_array_sections): Likewise. (finish_omp_clauses): Likewise. gcc/ * gimplify.cc (gimplify_omp_affinity): Use OMP_ITERATOR_DECL_P. (compute_iterator_count): New. (build_iterator_loop): New. (gimplify_omp_depend): Use OMP_ITERATOR_DECL_P, compute_iterator_count and build_iterator_loop. * tree-inline.cc (copy_tree_body_r): Use OMP_ITERATOR_DECL_P. * tree-pretty-print.cc (dump_omp_clause): Likewise. * tree.h (OMP_ITERATOR_DECL_P): New macro. --- gcc/c-family/c-omp.cc | 4 +- gcc/c/c-typeck.cc | 13 +- gcc/cp/pt.cc | 4 +- gcc/cp/semantics.cc | 8 +- gcc/gimplify.cc | 326 +++++++++++++++++++-------------------- gcc/tree-inline.cc | 5 +- gcc/tree-pretty-print.cc | 8 +- gcc/tree.h | 6 + 8 files changed, 175 insertions(+), 199 deletions(-) diff --git a/gcc/c-family/c-omp.cc b/gcc/c-family/c-omp.cc index b5ce1466e5d..5e469a4ee4d 100644 --- a/gcc/c-family/c-omp.cc +++ b/gcc/c-family/c-omp.cc @@ -744,9 +744,7 @@ c_finish_omp_depobj (location_t loc, tree depobj, kind = OMP_CLAUSE_DEPEND_KIND (clause); t = OMP_CLAUSE_DECL (clause); gcc_assert (t); - if (TREE_CODE (t) == TREE_LIST - && TREE_PURPOSE (t) - && TREE_CODE (TREE_PURPOSE (t)) == TREE_VEC) + if (OMP_ITERATOR_DECL_P (t)) { error_at (OMP_CLAUSE_LOCATION (clause), "% modifier may not be specified on " diff --git a/gcc/c/c-typeck.cc b/gcc/c/c-typeck.cc index 58b2724b39e..521c0e85605 100644 --- a/gcc/c/c-typeck.cc +++ b/gcc/c/c-typeck.cc @@ -14501,9 +14501,7 @@ handle_omp_array_sections (tree &c, enum c_omp_region_type ort) tree *tp = &OMP_CLAUSE_DECL (c); if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEPEND || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_AFFINITY) - && TREE_CODE (*tp) == TREE_LIST - && TREE_PURPOSE (*tp) - && TREE_CODE (TREE_PURPOSE (*tp)) == TREE_VEC) + && OMP_ITERATOR_DECL_P (*tp)) tp = &TREE_VALUE (*tp); tree first = handle_omp_array_sections_1 (c, *tp, types, maybe_zero_len, first_non_one, @@ -15694,9 +15692,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) case OMP_CLAUSE_DEPEND: case OMP_CLAUSE_AFFINITY: t = OMP_CLAUSE_DECL (c); - if (TREE_CODE (t) == TREE_LIST - && TREE_PURPOSE (t) - && TREE_CODE (TREE_PURPOSE (t)) == TREE_VEC) + if (OMP_ITERATOR_DECL_P (t)) { if (TREE_PURPOSE (t) != last_iterators) last_iterators_remove @@ -15796,10 +15792,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) break; } } - if (TREE_CODE (OMP_CLAUSE_DECL (c)) == TREE_LIST - && TREE_PURPOSE (OMP_CLAUSE_DECL (c)) - && (TREE_CODE (TREE_PURPOSE (OMP_CLAUSE_DECL (c))) - == TREE_VEC)) + if (OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (c))) TREE_VALUE (OMP_CLAUSE_DECL (c)) = t; else OMP_CLAUSE_DECL (c) = t; diff --git a/gcc/cp/pt.cc b/gcc/cp/pt.cc index 024fa8a5529..6f344665fbd 100644 --- a/gcc/cp/pt.cc +++ b/gcc/cp/pt.cc @@ -17562,9 +17562,7 @@ tsubst_omp_clause_decl (tree decl, tree args, tsubst_flags_t complain, return decl; /* Handle OpenMP iterators. */ - if (TREE_CODE (decl) == TREE_LIST - && TREE_PURPOSE (decl) - && TREE_CODE (TREE_PURPOSE (decl)) == TREE_VEC) + if (OMP_ITERATOR_DECL_P (decl)) { tree ret; if (iterator_cache[0] == TREE_PURPOSE (decl)) diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc index 5ab2076b673..7ecad569900 100644 --- a/gcc/cp/semantics.cc +++ b/gcc/cp/semantics.cc @@ -5898,9 +5898,7 @@ handle_omp_array_sections (tree &c, enum c_omp_region_type ort) tree *tp = &OMP_CLAUSE_DECL (c); if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEPEND || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_AFFINITY) - && TREE_CODE (*tp) == TREE_LIST - && TREE_PURPOSE (*tp) - && TREE_CODE (TREE_PURPOSE (*tp)) == TREE_VEC) + && OMP_ITERATOR_DECL_P (*tp)) tp = &TREE_VALUE (*tp); tree first = handle_omp_array_sections_1 (c, *tp, types, maybe_zero_len, first_non_one, @@ -8204,9 +8202,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) case OMP_CLAUSE_DEPEND: case OMP_CLAUSE_AFFINITY: t = OMP_CLAUSE_DECL (c); - if (TREE_CODE (t) == TREE_LIST - && TREE_PURPOSE (t) - && TREE_CODE (TREE_PURPOSE (t)) == TREE_VEC) + if (OMP_ITERATOR_DECL_P (t)) { if (TREE_PURPOSE (t) != last_iterators) last_iterators_remove diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index 9300138aa0c..8519095adef 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -8717,9 +8717,7 @@ gimplify_omp_affinity (tree *list_p, gimple_seq *pre_p) if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_AFFINITY) { tree t = OMP_CLAUSE_DECL (c); - if (TREE_CODE (t) == TREE_LIST - && TREE_PURPOSE (t) - && TREE_CODE (TREE_PURPOSE (t)) == TREE_VEC) + if (OMP_ITERATOR_DECL_P (t)) { if (TREE_VALUE (t) == null_pointer_node) continue; @@ -8824,6 +8822,159 @@ gimplify_omp_affinity (tree *list_p, gimple_seq *pre_p) return; } +/* Returns a tree expression containing the total iteration count of the + iterator clause decl T. */ + +static tree +compute_iterator_count (tree t, gimple_seq *pre_p) +{ + tree tcnt = size_one_node; + for (tree it = TREE_PURPOSE (t); it; it = TREE_CHAIN (it)) + { + if (gimplify_expr (&TREE_VEC_ELT (it, 1), pre_p, NULL, + is_gimple_val, fb_rvalue) == GS_ERROR + || gimplify_expr (&TREE_VEC_ELT (it, 2), pre_p, NULL, + is_gimple_val, fb_rvalue) == GS_ERROR + || gimplify_expr (&TREE_VEC_ELT (it, 3), pre_p, NULL, + is_gimple_val, fb_rvalue) == GS_ERROR + || (gimplify_expr (&TREE_VEC_ELT (it, 4), pre_p, NULL, + is_gimple_val, fb_rvalue) == GS_ERROR)) + return NULL_TREE; + tree var = TREE_VEC_ELT (it, 0); + tree begin = TREE_VEC_ELT (it, 1); + tree end = TREE_VEC_ELT (it, 2); + tree step = TREE_VEC_ELT (it, 3); + tree orig_step = TREE_VEC_ELT (it, 4); + tree type = TREE_TYPE (var); + tree stype = TREE_TYPE (step); + location_t loc = DECL_SOURCE_LOCATION (var); + tree endmbegin; + /* Compute count for this iterator as + orig_step > 0 + ? (begin < end ? (end - begin + (step - 1)) / step : 0) + : (begin > end ? (end - begin + (step + 1)) / step : 0) + and compute product of those for the entire clause. */ + if (POINTER_TYPE_P (type)) + endmbegin = fold_build2_loc (loc, POINTER_DIFF_EXPR, stype, end, begin); + else + endmbegin = fold_build2_loc (loc, MINUS_EXPR, type, end, begin); + tree stepm1 = fold_build2_loc (loc, MINUS_EXPR, stype, step, + build_int_cst (stype, 1)); + tree stepp1 = fold_build2_loc (loc, PLUS_EXPR, stype, step, + build_int_cst (stype, 1)); + tree pos = fold_build2_loc (loc, PLUS_EXPR, stype, + unshare_expr (endmbegin), stepm1); + pos = fold_build2_loc (loc, TRUNC_DIV_EXPR, stype, pos, step); + tree neg = fold_build2_loc (loc, PLUS_EXPR, stype, endmbegin, stepp1); + if (TYPE_UNSIGNED (stype)) + { + neg = fold_build1_loc (loc, NEGATE_EXPR, stype, neg); + step = fold_build1_loc (loc, NEGATE_EXPR, stype, step); + } + neg = fold_build2_loc (loc, TRUNC_DIV_EXPR, stype, neg, step); + step = NULL_TREE; + tree cond = fold_build2_loc (loc, LT_EXPR, boolean_type_node, begin, end); + pos = fold_build3_loc (loc, COND_EXPR, stype, cond, pos, + build_int_cst (stype, 0)); + cond = fold_build2_loc (loc, LT_EXPR, boolean_type_node, end, begin); + neg = fold_build3_loc (loc, COND_EXPR, stype, cond, neg, + build_int_cst (stype, 0)); + tree osteptype = TREE_TYPE (orig_step); + cond = fold_build2_loc (loc, GT_EXPR, boolean_type_node, orig_step, + build_int_cst (osteptype, 0)); + tree cnt = fold_build3_loc (loc, COND_EXPR, stype, cond, pos, neg); + cnt = fold_convert_loc (loc, sizetype, cnt); + if (gimplify_expr (&cnt, pre_p, NULL, is_gimple_val, + fb_rvalue) == GS_ERROR) + return NULL_TREE; + tcnt = size_binop_loc (loc, MULT_EXPR, tcnt, cnt); + } + if (gimplify_expr (&tcnt, pre_p, NULL, is_gimple_val, fb_rvalue) == GS_ERROR) + return NULL_TREE; + + return tcnt; +} + +/* Build loops iterating over the space defined by the iterators in clause C. + 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 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); + *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)) + { + tree var = TREE_VEC_ELT (it, 0); + tree begin = TREE_VEC_ELT (it, 1); + tree end = TREE_VEC_ELT (it, 2); + tree step = TREE_VEC_ELT (it, 3); + tree orig_step = TREE_VEC_ELT (it, 4); + tree type = TREE_TYPE (var); + location_t loc = DECL_SOURCE_LOCATION (var); + /* Emit: + var = begin; + goto cond_label; + beg_label: + ... + var = var + step; + cond_label: + if (orig_step > 0) { + if (var < end) goto beg_label; + } else { + if (var > end) goto beg_label; + } + for each iterator, with inner iterators added to + the ... above. */ + tree beg_label = create_artificial_label (loc); + tree cond_label = NULL_TREE; + tree tem = build2_loc (loc, MODIFY_EXPR, void_type_node, var, begin); + append_to_statement_list_force (tem, p); + tem = build_and_jump (&cond_label); + append_to_statement_list_force (tem, p); + tem = build1 (LABEL_EXPR, void_type_node, beg_label); + append_to_statement_list (tem, p); + tree bind = build3 (BIND_EXPR, void_type_node, NULL_TREE, + NULL_TREE, NULL_TREE); + TREE_SIDE_EFFECTS (bind) = 1; + SET_EXPR_LOCATION (bind, loc); + append_to_statement_list_force (bind, p); + if (POINTER_TYPE_P (type)) + tem = build2_loc (loc, POINTER_PLUS_EXPR, type, + var, fold_convert_loc (loc, sizetype, step)); + else + tem = build2_loc (loc, PLUS_EXPR, type, var, step); + tem = build2_loc (loc, MODIFY_EXPR, void_type_node, var, tem); + append_to_statement_list_force (tem, p); + tem = build1 (LABEL_EXPR, void_type_node, cond_label); + append_to_statement_list (tem, p); + tree cond = fold_build2_loc (loc, LT_EXPR, boolean_type_node, var, end); + tree pos = fold_build3_loc (loc, COND_EXPR, void_type_node, cond, + build_and_jump (&beg_label), void_node); + cond = fold_build2_loc (loc, GT_EXPR, boolean_type_node, var, end); + tree neg = fold_build3_loc (loc, COND_EXPR, void_type_node, cond, + build_and_jump (&beg_label), void_node); + tree osteptype = TREE_TYPE (orig_step); + cond = fold_build2_loc (loc, GT_EXPR, boolean_type_node, orig_step, + build_int_cst (osteptype, 0)); + tem = fold_build3_loc (loc, COND_EXPR, void_type_node, cond, pos, neg); + append_to_statement_list_force (tem, p); + p = &BIND_EXPR_BODY (bind); + } + + return p; +} + /* If *LIST_P contains any OpenMP depend clauses with iterators, lower all the depend clauses by populating corresponding depend array. Returns 0 if there are no such depend clauses, or @@ -8868,89 +9019,12 @@ gimplify_omp_depend (tree *list_p, gimple_seq *pre_p) tree t = OMP_CLAUSE_DECL (c); if (first_loc == UNKNOWN_LOCATION) first_loc = OMP_CLAUSE_LOCATION (c); - if (TREE_CODE (t) == TREE_LIST - && TREE_PURPOSE (t) - && TREE_CODE (TREE_PURPOSE (t)) == TREE_VEC) + if (OMP_ITERATOR_DECL_P (t)) { if (TREE_PURPOSE (t) != last_iter) { - tree tcnt = size_one_node; - for (tree it = TREE_PURPOSE (t); it; it = TREE_CHAIN (it)) - { - if (gimplify_expr (&TREE_VEC_ELT (it, 1), pre_p, NULL, - is_gimple_val, fb_rvalue) == GS_ERROR - || gimplify_expr (&TREE_VEC_ELT (it, 2), pre_p, NULL, - is_gimple_val, fb_rvalue) == GS_ERROR - || gimplify_expr (&TREE_VEC_ELT (it, 3), pre_p, NULL, - is_gimple_val, fb_rvalue) == GS_ERROR - || (gimplify_expr (&TREE_VEC_ELT (it, 4), pre_p, NULL, - is_gimple_val, fb_rvalue) - == GS_ERROR)) - return 2; - tree var = TREE_VEC_ELT (it, 0); - tree begin = TREE_VEC_ELT (it, 1); - tree end = TREE_VEC_ELT (it, 2); - tree step = TREE_VEC_ELT (it, 3); - tree orig_step = TREE_VEC_ELT (it, 4); - tree type = TREE_TYPE (var); - tree stype = TREE_TYPE (step); - location_t loc = DECL_SOURCE_LOCATION (var); - tree endmbegin; - /* Compute count for this iterator as - orig_step > 0 - ? (begin < end ? (end - begin + (step - 1)) / step : 0) - : (begin > end ? (end - begin + (step + 1)) / step : 0) - and compute product of those for the entire depend - clause. */ - if (POINTER_TYPE_P (type)) - endmbegin = fold_build2_loc (loc, POINTER_DIFF_EXPR, - stype, end, begin); - else - endmbegin = fold_build2_loc (loc, MINUS_EXPR, type, - end, begin); - tree stepm1 = fold_build2_loc (loc, MINUS_EXPR, stype, - step, - build_int_cst (stype, 1)); - tree stepp1 = fold_build2_loc (loc, PLUS_EXPR, stype, step, - build_int_cst (stype, 1)); - tree pos = fold_build2_loc (loc, PLUS_EXPR, stype, - unshare_expr (endmbegin), - stepm1); - pos = fold_build2_loc (loc, TRUNC_DIV_EXPR, stype, - pos, step); - tree neg = fold_build2_loc (loc, PLUS_EXPR, stype, - endmbegin, stepp1); - if (TYPE_UNSIGNED (stype)) - { - neg = fold_build1_loc (loc, NEGATE_EXPR, stype, neg); - step = fold_build1_loc (loc, NEGATE_EXPR, stype, step); - } - neg = fold_build2_loc (loc, TRUNC_DIV_EXPR, stype, - neg, step); - step = NULL_TREE; - tree cond = fold_build2_loc (loc, LT_EXPR, - boolean_type_node, - begin, end); - pos = fold_build3_loc (loc, COND_EXPR, stype, cond, pos, - build_int_cst (stype, 0)); - cond = fold_build2_loc (loc, LT_EXPR, boolean_type_node, - end, begin); - neg = fold_build3_loc (loc, COND_EXPR, stype, cond, neg, - build_int_cst (stype, 0)); - tree osteptype = TREE_TYPE (orig_step); - cond = fold_build2_loc (loc, GT_EXPR, boolean_type_node, - orig_step, - build_int_cst (osteptype, 0)); - tree cnt = fold_build3_loc (loc, COND_EXPR, stype, - cond, pos, neg); - cnt = fold_convert_loc (loc, sizetype, cnt); - if (gimplify_expr (&cnt, pre_p, NULL, is_gimple_val, - fb_rvalue) == GS_ERROR) - return 2; - tcnt = size_binop_loc (loc, MULT_EXPR, tcnt, cnt); - } - if (gimplify_expr (&tcnt, pre_p, NULL, is_gimple_val, - fb_rvalue) == GS_ERROR) + tree tcnt = compute_iterator_count (t, pre_p); + if (!tcnt) return 2; last_iter = TREE_PURPOSE (t); last_count = tcnt; @@ -9104,92 +9178,10 @@ gimplify_omp_depend (tree *list_p, gimple_seq *pre_p) gcc_unreachable (); } tree t = OMP_CLAUSE_DECL (c); - if (TREE_CODE (t) == TREE_LIST - && TREE_PURPOSE (t) - && TREE_CODE (TREE_PURPOSE (t)) == TREE_VEC) + if (OMP_ITERATOR_DECL_P (t)) { if (TREE_PURPOSE (t) != last_iter) - { - if (last_bind) - gimplify_and_add (last_bind, pre_p); - tree block = TREE_VEC_ELT (TREE_PURPOSE (t), 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)) - { - tree var = TREE_VEC_ELT (it, 0); - tree begin = TREE_VEC_ELT (it, 1); - tree end = TREE_VEC_ELT (it, 2); - tree step = TREE_VEC_ELT (it, 3); - tree orig_step = TREE_VEC_ELT (it, 4); - tree type = TREE_TYPE (var); - location_t loc = DECL_SOURCE_LOCATION (var); - /* Emit: - var = begin; - goto cond_label; - beg_label: - ... - var = var + step; - cond_label: - if (orig_step > 0) { - if (var < end) goto beg_label; - } else { - if (var > end) goto beg_label; - } - for each iterator, with inner iterators added to - the ... above. */ - tree beg_label = create_artificial_label (loc); - tree cond_label = NULL_TREE; - tem = build2_loc (loc, MODIFY_EXPR, void_type_node, - var, begin); - append_to_statement_list_force (tem, p); - tem = build_and_jump (&cond_label); - append_to_statement_list_force (tem, p); - tem = build1 (LABEL_EXPR, void_type_node, beg_label); - append_to_statement_list (tem, p); - tree bind = build3 (BIND_EXPR, void_type_node, NULL_TREE, - NULL_TREE, NULL_TREE); - TREE_SIDE_EFFECTS (bind) = 1; - SET_EXPR_LOCATION (bind, loc); - append_to_statement_list_force (bind, p); - if (POINTER_TYPE_P (type)) - tem = build2_loc (loc, POINTER_PLUS_EXPR, type, - var, fold_convert_loc (loc, sizetype, - step)); - else - tem = build2_loc (loc, PLUS_EXPR, type, var, step); - tem = build2_loc (loc, MODIFY_EXPR, void_type_node, - var, tem); - append_to_statement_list_force (tem, p); - tem = build1 (LABEL_EXPR, void_type_node, cond_label); - append_to_statement_list (tem, p); - tree cond = fold_build2_loc (loc, LT_EXPR, - boolean_type_node, - var, end); - tree pos - = fold_build3_loc (loc, COND_EXPR, void_type_node, - cond, build_and_jump (&beg_label), - void_node); - cond = fold_build2_loc (loc, GT_EXPR, boolean_type_node, - var, end); - tree neg - = fold_build3_loc (loc, COND_EXPR, void_type_node, - cond, build_and_jump (&beg_label), - void_node); - tree osteptype = TREE_TYPE (orig_step); - cond = fold_build2_loc (loc, GT_EXPR, boolean_type_node, - orig_step, - build_int_cst (osteptype, 0)); - tem = fold_build3_loc (loc, COND_EXPR, void_type_node, - cond, pos, neg); - append_to_statement_list_force (tem, p); - p = &BIND_EXPR_BODY (bind); - } - last_body = p; - } + last_body = build_iterator_loop (c, pre_p, &last_bind); last_iter = TREE_PURPOSE (t); if (TREE_CODE (TREE_VALUE (t)) == COMPOUND_EXPR) { diff --git a/gcc/tree-inline.cc b/gcc/tree-inline.cc index f31a34ac410..05dea9473a0 100644 --- a/gcc/tree-inline.cc +++ b/gcc/tree-inline.cc @@ -1453,10 +1453,7 @@ copy_tree_body_r (tree *tp, int *walk_subtrees, void *data) || OMP_CLAUSE_CODE (*tp) == OMP_CLAUSE_DEPEND)) { tree t = OMP_CLAUSE_DECL (*tp); - if (t - && TREE_CODE (t) == TREE_LIST - && TREE_PURPOSE (t) - && TREE_CODE (TREE_PURPOSE (t)) == TREE_VEC) + if (t && OMP_ITERATOR_DECL_P (t)) { *walk_subtrees = 0; OMP_CLAUSE_DECL (*tp) = copy_node (t); diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc index 4bb946bb0e8..ab7ecbfd1ef 100644 --- a/gcc/tree-pretty-print.cc +++ b/gcc/tree-pretty-print.cc @@ -805,9 +805,7 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) pp_string (pp, "affinity("); { tree t = OMP_CLAUSE_DECL (clause); - if (TREE_CODE (t) == TREE_LIST - && TREE_PURPOSE (t) - && TREE_CODE (TREE_PURPOSE (t)) == TREE_VEC) + if (OMP_ITERATOR_DECL_P (t)) { dump_omp_iterators (pp, TREE_PURPOSE (t), spc, flags); pp_colon (pp); @@ -847,9 +845,7 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) } { tree t = OMP_CLAUSE_DECL (clause); - if (TREE_CODE (t) == TREE_LIST - && TREE_PURPOSE (t) - && TREE_CODE (TREE_PURPOSE (t)) == TREE_VEC) + if (OMP_ITERATOR_DECL_P (t)) { dump_omp_iterators (pp, TREE_PURPOSE (t), spc, flags); pp_colon (pp); diff --git a/gcc/tree.h b/gcc/tree.h index 75efc760a16..83075b82cc7 100644 --- a/gcc/tree.h +++ b/gcc/tree.h @@ -2156,6 +2156,12 @@ class auto_suppress_location_wrappers #define OMP_CLAUSE_OPERAND(NODE, I) \ OMP_CLAUSE_ELT_CHECK (NODE, I) +/* True if the clause decl NODE contains an iterator. */ +#define OMP_ITERATOR_DECL_P(NODE) \ + (TREE_CODE (NODE) == TREE_LIST \ + && TREE_PURPOSE (NODE) \ + && TREE_CODE (TREE_PURPOSE (NODE)) == TREE_VEC) + /* In a BLOCK (scope) node: Variables declared in the scope NODE. */ #define BLOCK_VARS(NODE) (BLOCK_CHECK (NODE)->block.vars) 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; +} From patchwork Tue Sep 3 17:08:54 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: 1980235 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=uloxNtPv; 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 4WysZF4bmxz1yg9 for ; Wed, 4 Sep 2024 03:10:13 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id DA84F386074B for ; Tue, 3 Sep 2024 17:10:11 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from mail-ej1-x632.google.com (mail-ej1-x632.google.com [IPv6:2a00:1450:4864:20::632]) by sourceware.org (Postfix) with ESMTPS id 55B563858402 for ; Tue, 3 Sep 2024 17:09:47 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 55B563858402 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 55B563858402 Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=2a00:1450:4864:20::632 ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1725383391; cv=none; b=RVeSObir8nO1MGNl7xu+O2cWuVwbm8NCdImZNRZdSpgWEsCVrqcYeVZfGA9SUtr4WROf58UsTSQqHW9LAZH9AU0xe1KKJAA2J18huM6QknXEN54TekyjRdJTzS2evO55ZfZZzjtrehjg8++7wkGddoWNnQAEyhhlBKaFz5NDtUc= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1725383391; c=relaxed/simple; bh=7iWNw6uuFFdsNsYSLk3L/oJ7nq6MdQTFnClyUGLvuvc=; h=DKIM-Signature:Message-ID:Date:MIME-Version:To:Subject:From; b=hTk59sJQhihevHuP8J91T4EKu+Y/dKwa5IycYipI8jlPtddCL6OIpKpNfKM0qQ90dW1zUz3NbGHjNbacPti6Yf9lcdSwxlFTV9kRakHaHNDj/c9H8rUQs/C2DthJQlp90B+Z2updpc5G/s3aRf52qowguBHBqf77u3/j1jMMGys= ARC-Authentication-Results: i=1; server2.sourceware.org Received: by mail-ej1-x632.google.com with SMTP id a640c23a62f3a-a7a9cf7d3f3so626306666b.1 for ; Tue, 03 Sep 2024 10:09:47 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre-com.20230601.gappssmtp.com; s=20230601; t=1725383386; x=1725988186; 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=7iWNw6uuFFdsNsYSLk3L/oJ7nq6MdQTFnClyUGLvuvc=; b=uloxNtPvX7K2MFps4I8+gvxqGCye1wYjUILPtoIDmpNO7OsIDMlEVuOBkV7nj+l3j3 qilD/h79w8jfbn5F5z0GbOW3bZq1nwOEVAg9Fh7F6QNbQ4JbdMZNvpK+Hi+Bqyob/hz4 S0Mc3W8qrVhgZ4sm00K0CgexAxhiUOPMS74A6TNqnbconLHIxOQ371r0Y4T2vBEJV0kF Xlkghe87d/kclIdBLL8Cx76id8Jj0Hd893LX8jK3cCjKcxQBt9piyvzdZRsSKAHTDOVx EvgSYMKiwEYEVhr6n0dbRNuu8UNQB4lu2b5Ygp+e1CcV7N9kOEU3nJ63fB7ffdSfwzD7 Jx6g== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1725383386; x=1725988186; 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=7iWNw6uuFFdsNsYSLk3L/oJ7nq6MdQTFnClyUGLvuvc=; b=MpW7iccbGNauuQ7gvFQa8zDSkcnQCl+DfXmCSdSw85cTWtr53TVU1XlHGqTQaur9C3 I9eLwis0KiMreNUdnueSFZzNW5UNLviYpKy5PerAV1BetZwfxdwDfb0tqMHXmKwVl0Vk B/TC1Rnw/NbJQxzd3h3Xs4OgM3AnAfDP9Q/avT64rafWNpCRY2ylu+dI+WGU6ewTfqVy wM/VT1luP9f3sjuEdhP1u3MolTipyYNABSmJASdKb4uED9pr9ElLXbiOBv1pWg7ODMfz nl/i1Iix9bc3sHqnNVexE2oRvpzJWFJ1scNiVX7Mclk8M09WFT25BzWcBnMFdvI+4TzN wbgA== X-Gm-Message-State: AOJu0YzKoQccwwp2aX+miDOC09EBn9TDu1D+5n7xaKaks8dD/2VZyAEG 3IDKwIrL0gV1SCrCI5NL536dwxyIQFMHTcRxZbTHG7DVjCJBwExFxF8RQDcde9XXpjrfXFcA4Hv d X-Google-Smtp-Source: AGHT+IGRKlt21Gb1M0PtGCpIdhtPTB6qK0o5r4+mA8GifyuZeCoQQxMBao9vgz9E++qhBMrb8R5erA== X-Received: by 2002:a17:907:26c4:b0:a7a:acae:3415 with SMTP id a640c23a62f3a-a89fad7c599mr563992966b.10.1725383385317; Tue, 03 Sep 2024 10:09:45 -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-a89892333a0sm699321566b.220.2024.09.03.10.09.44 (version=TLS1_3 cipher=TLS_AES_128_GCM_SHA256 bits=128/128); Tue, 03 Sep 2024 10:09:44 -0700 (PDT) Message-ID: <548eeb7f-729a-4845-8db7-316d8cbc03a0@baylibre.com> Date: Tue, 3 Sep 2024 18:08:54 +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 3/5] openmp: Add support for iterators in 'target update' 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.7 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 extends the previous patch to cover to/from clauses in 'target update'. From c3dfc4a792610530a4ab729c3f250917b828e469 Mon Sep 17 00:00:00 2001 From: Kwok Cheung Yeung Date: Mon, 2 Sep 2024 19:34:09 +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-09-02 Kwok Cheung Yeung gcc/c/ * c-parser.cc (c_parser_omp_clause_from_to): Parse 'iterator' modifier. gcc/cp/ * parser.cc (cp_parser_omp_clause_from_to): Parse 'iterator' modifier. gcc/ * gimplify.cc (gimplify_omp_map_iterators): Gimplify iterators in to/from clauses. (gimplify_scan_omp_clauses): Skip gimplification of clause decl and size for clauses with iterators. * tree-pretty-print.cc (dump_omp_clause): Call dump_omp_map_iterators for to/from clauses with iterators. gcc/testsuite/ * c-c++-common/gomp/target-update-iterator-1.c: New. * c-c++-common/gomp/target-update-iterator-2.c: New. * c-c++-common/gomp/target-update-iterator-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/cp/parser.cc | 111 ++++++++++++++++-- gcc/gimplify.cc | 24 ++-- .../gomp/target-update-iterator-1.c | 20 ++++ .../gomp/target-update-iterator-2.c | 17 +++ .../gomp/target-update-iterator-3.c | 17 +++ gcc/tree-pretty-print.cc | 20 +++- libgomp/target.c | 12 ++ .../target-update-iterators-1.c | 65 ++++++++++ .../target-update-iterators-2.c | 57 +++++++++ .../target-update-iterators-3.c | 66 +++++++++++ 11 files changed, 487 insertions(+), 27 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/gomp/target-update-iterator-1.c create mode 100644 gcc/testsuite/c-c++-common/gomp/target-update-iterator-2.c create mode 100644 gcc/testsuite/c-c++-common/gomp/target-update-iterator-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 f72fca1a711..37c419eb326 100644 --- a/gcc/c/c-parser.cc +++ b/gcc/c/c-parser.cc @@ -19305,8 +19305,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, @@ -19317,15 +19320,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); } @@ -19336,6 +19412,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_DECL (c) = build_tree_list (iterators, OMP_CLAUSE_DECL (c)); + return nl; } diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc index 29947177415..cda0ec6efeb 100644 --- a/gcc/cp/parser.cc +++ b/gcc/cp/parser.cc @@ -41659,8 +41659,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, @@ -41669,15 +41672,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); } @@ -41686,6 +41768,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_DECL (c) = build_tree_list (iterators, OMP_CLAUSE_DECL (c)); + return nl; } diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index 549acf4dfbb..6e938296245 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -9387,7 +9387,9 @@ gimplify_omp_map_iterators (tree *list_p, gimple_seq *pre_p) /* 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 + if ((OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO + || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FROM) && OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (c))) all_iterators.add (TREE_PURPOSE (OMP_CLAUSE_DECL (c))); @@ -9397,7 +9399,9 @@ gimplify_omp_map_iterators (tree *list_p, gimple_seq *pre_p) { for (tree c = *list_p; c; c = OMP_CLAUSE_CHAIN (c)) { - if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP + if ((OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP + && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_TO + && OMP_CLAUSE_CODE (c) != OMP_CLAUSE_FROM) || !OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (c))) continue; @@ -13004,6 +13008,10 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, case OMP_CLAUSE_FROM: case OMP_CLAUSE__CACHE_: 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)) { remove = true; @@ -13012,17 +13020,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 (!last_iterators + && 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 (!last_iterators + && gimplify_expr (&OMP_CLAUSE_DECL (c), pre_p, + NULL, is_gimple_lvalue, fb_lvalue) + == GS_ERROR) { remove = true; break; diff --git a/gcc/testsuite/c-c++-common/gomp/target-update-iterator-1.c b/gcc/testsuite/c-c++-common/gomp/target-update-iterator-1.c new file mode 100644 index 00000000000..3a64f511da4 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/target-update-iterator-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-iterator-2.c b/gcc/testsuite/c-c++-common/gomp/target-update-iterator-2.c new file mode 100644 index 00000000000..3789a559b6f --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/target-update-iterator-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-iterator-3.c b/gcc/testsuite/c-c++-common/gomp/target-update-iterator-3.c new file mode 100644 index 00000000000..d8672b3a242 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/target-update-iterator-3.c @@ -0,0 +1,17 @@ +/* { 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 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 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "if \\(j <= 19\\) goto ; else goto ;" 1 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "to\\(iterator\\(int i=0:10:1, int j=0:20:1\\):iterator_array=D\.\[0-9\]+" 2 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "from\\(iterator\\(int i=0:10:1\\):iterator_array=D\.\[0-9\]+" 1 "gimple" } } */ diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc index 21ecf94ada5..e43f30818d0 100644 --- a/gcc/tree-pretty-print.cc +++ b/gcc/tree-pretty-print.cc @@ -1099,16 +1099,28 @@ 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:"); - dump_generic_node (pp, OMP_CLAUSE_DECL (clause), - spc, flags, false); + 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); + } + dump_generic_node (pp, decl, spc, flags, false); goto print_clause_size; case OMP_CLAUSE_TO: pp_string (pp, "to("); if (OMP_CLAUSE_MOTION_PRESENT (clause)) pp_string (pp, "present:"); - dump_generic_node (pp, OMP_CLAUSE_DECL (clause), - spc, flags, false); + 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); + } + dump_generic_node (pp, decl, spc, flags, false); goto print_clause_size; case OMP_CLAUSE__CACHE_: diff --git a/libgomp/target.c b/libgomp/target.c index 37a4c539647..c69418f0b78 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -2224,6 +2224,7 @@ 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; if (!devicep) return; @@ -2231,6 +2232,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); + gomp_mutex_lock (&devicep->lock); if (devicep->state == GOMP_DEVICE_FINALIZED) { @@ -2324,6 +2329,13 @@ 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); + } } 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..949cc266d84 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-2.c @@ -0,0 +1,57 @@ +/* { dg-do run } */ + +/* 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..852635e50f6 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-3.c @@ -0,0 +1,66 @@ +/* { dg-do run } */ + +/* 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; +} From patchwork Tue Sep 3 17:10:09 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: 1980239 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=bGztyvVl; 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 4Wysbr107vz1yg9 for ; Wed, 4 Sep 2024 03:11:36 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id B58C63864819 for ; Tue, 3 Sep 2024 17:11:33 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from mail-lf1-x132.google.com (mail-lf1-x132.google.com [IPv6:2a00:1450:4864:20::132]) by sourceware.org (Postfix) with ESMTPS id 9A76C3858402 for ; Tue, 3 Sep 2024 17:11:03 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 9A76C3858402 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 9A76C3858402 Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=2a00:1450:4864:20::132 ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1725383469; cv=none; b=mqShKcUpsBOHi3bNNNuz2h02sK9fJH2Qabsg5m4o8By1syp9HswVMAxkHg2ny1y9lot9K8M0SJZB1hFzhkKz7sgV7mdvdoHjAcyRFTk/ANxMRphbcmWKkSmLSKo0ff1r80Mkkp8uZjByYfKgjql2WIitAN675nSGPAcAsQWehuY= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1725383469; c=relaxed/simple; bh=iHAWk7bf/OuBTZNJBBcuHIMwkOgfd9Uz7fEXX2RfE4I=; h=DKIM-Signature:Message-ID:Date:MIME-Version:To:Subject:From; b=k9rb63QL5QyLhaTZPrAG+8g8SivuHNr3OtHKW6gsIapDwGusvPOxw6GxYQgxWMDvRMj8PmyXIuw9FzxeJu0CzCa4N8XfwHRe0YCyRBvcVMD1YTO6yfW+XPBl1l0OgifCA/aA39ckZH70vIWN6+tLayCkhlcu/pdLBM3U5HZ/k1I= ARC-Authentication-Results: i=1; server2.sourceware.org Received: by mail-lf1-x132.google.com with SMTP id 2adb3069b0e04-534366c194fso5144297e87.0 for ; Tue, 03 Sep 2024 10:11:03 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre-com.20230601.gappssmtp.com; s=20230601; t=1725383462; x=1725988262; 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=ICsZtNDLeBZaJG7/EbftGOAx4eUaBIHwtUS4MyLOZLE=; b=bGztyvVl1cyS4+UG9lkMGY/wZXZoySWpKlrxd8jktzWHaWzPzfpH1TH8hQWP0Arswp fKkygU/pwndyy9cAXX2h7kZVDNjV2dLyROEISroDizfY+d9pJ+kOd8wIJBdlinQoOjN9 Y0BIVu0jwp0n3FI3noyVKD90L2zk1FgYSocjMR9nPC9ld+SpPncy+7wwqy/g2ebscY32 XRjRoFYk+hPN5qkZbOBuELQO8LpDN5A45yBR9/wj+LztAiFdtqRYqLHFshRKi2a6z4HC f0JwnzGYjQaKegzrav2ULU5Wp+w48KW5EyaRQ2QL2UBBTCZxyv4qi9c2JgjP+k3uTvOD UakA== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1725383462; x=1725988262; 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=ICsZtNDLeBZaJG7/EbftGOAx4eUaBIHwtUS4MyLOZLE=; b=F73mFOB3L/xQMWl3FzxnVejBU6IQVMu5mgaZR4NCz/K9Pyd3u5GSNfYzdeSJ6ihG3a A02mUEvdke7qkIEx5lgHI09ojM9UqwvHaZLXul8D1+Nrm9FzslF6TZQIW1ehD+PTelFR /kA/d2oO//BunipIUE7x3twf4aFSuHDW4m4YhbNaR2n8kTtyTyY82AuVRpAU85lgnSU0 Xu4JFSeO7YfbDg4iriKC8Lw9f5T98zBFy62y60NjburoRdGlkbjpPRDF9b3E4qpwDr/U dcGrni7+jYeU7RxrM2BT2p3djZNsY9CDMdL0kNSkqO/+VTsxq5G02vPYzWvubpE1RG/Z jmZg== X-Gm-Message-State: AOJu0YzrL+VLJ+1lyeGU29eGmdTTd6MCi/pcPN8pZ7Eg25+iNjUqqDKK DYACxZbZcWZ4llztlN9ij/lhekU7ozC4zOBS0cTkb71q4C29Tfyaq3zmYtb7CqgsUXynVAfkg9k 7 X-Google-Smtp-Source: AGHT+IH0Q3fYY3pjpaqKriO7TKfC3cqV+oS7vgFPs6W15FVIsBb2NSZPAxs4pShUOgQL41ubvkN53g== X-Received: by 2002:a05:6512:2804:b0:52c:d905:9645 with SMTP id 2adb3069b0e04-53546b32d3dmr9100566e87.13.1725383460843; Tue, 03 Sep 2024 10:11:00 -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-a8988feaf1esm703912566b.7.2024.09.03.10.11.00 (version=TLS1_3 cipher=TLS_AES_128_GCM_SHA256 bits=128/128); Tue, 03 Sep 2024 10:11:00 -0700 (PDT) Message-ID: <1728c2ce-3a61-4ad8-beef-21b361e9a0d0@baylibre.com> Date: Tue, 3 Sep 2024 18:10:09 +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 4/5] openmp, fortran: Add support for map iterators in OpenMP target construct (Fortran) Content-Language: en-GB From: Kwok Cheung Yeung In-Reply-To: <77f9ccb8-6f5e-4462-aa32-71f74fd7ff26@baylibre.com> X-Spam-Status: No, score=-12.8 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 adds support for iterators in the map clause of OpenMP target constructs. The parsing and translation of iterators in the front-end works the same as for the affinity and depend clauses. The iterator gimplification needed to be modified slightly to handle Fortran. The difference in how ranges work in loops (i.e. the condition on the upper bound is <=, rather than < as in C/C++) needs to be compensated for when calculating the iteration count and in the iteration loop itself. During Fortran translation of iterators, statements for the side-effects of any translated expressions are placed into BLOCK_SUBBLOCKS of the block containing the iterator variables (this also occurs with the other clauses supporting iterators). However, the previous lowering of iterators into Gimple does not appear to do anything with these statements, which causes issues if anything in the loop body references these side-effects (typically calculation of array boundaries and strides). This appears to be a bug that was simply not triggered by existing testcases. These statements are now gimplified into the innermost loop body. The libgomp runtime was modified to handle GOMP_MAP_STRUCTs in iterators, which can result from the use of derived types (which I used in test cases to implement arrays of pointers). libgomp expects a GOMP_MAP_STRUCT map to be followed immediately by a number of maps corresponding to the fields of the struct, so an iterator GOMP_MAP_STRUCT and its fields need to be expanded in a breadth-first order, rather than the usual depth-first manner (which would result in multiple GOMP_MAP_STRUCTS, followed by multiple instances of the first field, then multiples of the second etc.). When filling in the .omp_data_t data structure for the target, only the address associated with the first map generated by an iterator is set (as only a single slot in the data structure is allocated for each iterator map). From f7cdf555e9d5c49b455a364a1eef2123c7bb76d1 Mon Sep 17 00:00:00 2001 From: Kwok Cheung Yeung Date: Mon, 2 Sep 2024 19:34:15 +0100 Subject: [PATCH 4/5] openmp, fortran: Add support for map iterators in OpenMP target construct (Fortran) This adds support for iterators in map clauses within OpenMP 'target' constructs in Fortran. Some special handling for struct field maps has been added to libgomp in order to handle arrays of derived types. 2024-09-02 Kwok Cheung Yeung gcc/fortran/ * dump-parse-tree.cc (show_omp_namelist): Add iterator support for OMP_LIST_MAP. * openmp.cc (gfc_free_omp_clauses): Free namespace in namelist for OMP_LIST_MAP. (gfc_match_omp_clauses): Parse 'iterator' modifier for 'map' clause. (resolve_omp_clauses): Resolve iterators for OMP_LIST_MAP. * trans-openmp.cc (gfc_trans_omp_clauses): Handle iterators in OMP_LIST_MAP clauses. gcc/ * gimplify.cc (compute_iterator_count): Account for difference in loop boundaries in Fortran. (build_iterator_loop): Change upper boundary condition for Fortran. Insert block statements into innermost loop. (omp_accumulate_sibling_list): Prevent structs generated by iterators from being treated as unordered. * tree-pretty-print.cc (dump_block_node): Ignore BLOCK_SUBBLOCKS containing iterator block statements. gcc/testsuite/ * gfortran.dg/gomp/target-iterator-1.f90: New. * gfortran.dg/gomp/target-iterator-2.f90: New. * gfortran.dg/gomp/target-iterator-3.f90: New. libgomp/ * target.c (kind_to_name): New. (gomp_add_map): New. (gomp_merge_iterator_maps): Return array indicating the iteration that a map originated from. Expand fields of a struct mapping breadth-first. (gomp_map_vars_internal): Add extra argument in call to gomp_merge_iterator_maps and free it at the end. Only add address of first iteration for field maps to target variables. (gomp_update): Add extra argument in call to gomp_merge_iterator_maps. Free it at the end of the function. * testsuite/libgomp.fortran/target-map-iterators-1.f90: New. * testsuite/libgomp.fortran/target-map-iterators-2.f90: New. * testsuite/libgomp.fortran/target-map-iterators-3.f90: New. --- gcc/fortran/dump-parse-tree.cc | 9 +- gcc/fortran/openmp.cc | 35 ++++- gcc/fortran/trans-openmp.cc | 73 ++++++++-- gcc/gimplify.cc | 36 +++-- .../gfortran.dg/gomp/target-iterator-1.f90 | 26 ++++ .../gfortran.dg/gomp/target-iterator-2.f90 | 27 ++++ .../gfortran.dg/gomp/target-iterator-3.f90 | 24 ++++ gcc/tree-pretty-print.cc | 4 +- libgomp/target.c | 132 ++++++++++++++---- .../target-map-iterators-1.f90 | 45 ++++++ .../target-map-iterators-2.f90 | 45 ++++++ .../target-map-iterators-3.f90 | 57 ++++++++ 12 files changed, 451 insertions(+), 62 deletions(-) create mode 100644 gcc/testsuite/gfortran.dg/gomp/target-iterator-1.f90 create mode 100644 gcc/testsuite/gfortran.dg/gomp/target-iterator-2.f90 create mode 100644 gcc/testsuite/gfortran.dg/gomp/target-iterator-3.f90 create mode 100644 libgomp/testsuite/libgomp.fortran/target-map-iterators-1.f90 create mode 100644 libgomp/testsuite/libgomp.fortran/target-map-iterators-2.f90 create mode 100644 libgomp/testsuite/libgomp.fortran/target-map-iterators-3.f90 diff --git a/gcc/fortran/dump-parse-tree.cc b/gcc/fortran/dump-parse-tree.cc index 80aa8ef84e7..0272a443f65 100644 --- a/gcc/fortran/dump-parse-tree.cc +++ b/gcc/fortran/dump-parse-tree.cc @@ -1349,7 +1349,8 @@ show_omp_namelist (int list_type, gfc_omp_namelist *n) for (; n; n = n->next) { gfc_current_ns = ns_curr; - if (list_type == OMP_LIST_AFFINITY || list_type == OMP_LIST_DEPEND) + if (list_type == OMP_LIST_AFFINITY || list_type == OMP_LIST_DEPEND + || list_type == OMP_LIST_MAP) { gfc_current_ns = n->u2.ns ? n->u2.ns : ns_curr; if (n->u2.ns != ns_iter) @@ -1361,8 +1362,12 @@ show_omp_namelist (int list_type, gfc_omp_namelist *n) fputs ("AFFINITY (", dumpfile); else if (n->u.depend_doacross_op == OMP_DOACROSS_SINK_FIRST) fputs ("DOACROSS (", dumpfile); - else + else if (list_type == OMP_LIST_DEPEND) fputs ("DEPEND (", dumpfile); + else if (list_type == OMP_LIST_MAP) + fputs ("MAP (", dumpfile); + else + gcc_unreachable (); } if (n->u2.ns) { diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc index 333f0c7fe7f..996126e6e7f 100644 --- a/gcc/fortran/openmp.cc +++ b/gcc/fortran/openmp.cc @@ -191,7 +191,8 @@ gfc_free_omp_clauses (gfc_omp_clauses *c) gfc_free_expr (c->vector_length_expr); for (i = 0; i < OMP_LIST_NUM; i++) gfc_free_omp_namelist (c->lists[i], - i == OMP_LIST_AFFINITY || i == OMP_LIST_DEPEND, + i == OMP_LIST_AFFINITY || i == OMP_LIST_DEPEND + || i == OMP_LIST_MAP, i == OMP_LIST_ALLOCATE, i == OMP_LIST_USES_ALLOCATORS); gfc_free_expr_list (c->wait_list); @@ -3079,9 +3080,12 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, int always_modifier = 0; int close_modifier = 0; int present_modifier = 0; + int iterator_modifier = 0; + gfc_namespace *ns_iter = NULL, *ns_curr = gfc_current_ns; locus second_always_locus = old_loc2; locus second_close_locus = old_loc2; locus second_present_locus = old_loc2; + locus second_iterator_locus = old_loc2; for (;;) { @@ -3101,6 +3105,11 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, if (present_modifier++ == 1) second_present_locus = current_locus; } + else if (gfc_match_iterator (&ns_iter, true) == MATCH_YES) + { + if (iterator_modifier++ == 1) + second_iterator_locus = current_locus; + } else break; gfc_match (", "); @@ -3157,15 +3166,30 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, &second_present_locus); break; } + if (iterator_modifier > 1) + { + gfc_error ("too many % modifiers at %L", + &second_iterator_locus); + break; + } head = NULL; - if (gfc_match_omp_variable_list ("", &c->lists[OMP_LIST_MAP], + if (ns_iter) + gfc_current_ns = ns_iter; + m = gfc_match_omp_variable_list ("", &c->lists[OMP_LIST_MAP], false, NULL, &head, - true, true) == MATCH_YES) + true, true); + gfc_current_ns = ns_curr; + if (m == MATCH_YES) { gfc_omp_namelist *n; for (n = *head; n; n = n->next) - n->u.map.op = map_op; + { + n->u.map.op = map_op; + n->u2.ns = ns_iter; + if (ns_iter) + ns_iter->refs++; + } continue; } gfc_current_locus = old_loc; @@ -8411,7 +8435,8 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, case OMP_LIST_CACHE: for (; n != NULL; n = n->next) { - if ((list == OMP_LIST_DEPEND || list == OMP_LIST_AFFINITY) + if ((list == OMP_LIST_DEPEND || list == OMP_LIST_AFFINITY + || list == OMP_LIST_MAP) && n->u2.ns && !n->u2.ns->resolved) { n->u2.ns->resolved = 1; diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc index df1bf144e23..a9929430e53 100644 --- a/gcc/fortran/trans-openmp.cc +++ b/gcc/fortran/trans-openmp.cc @@ -2694,7 +2694,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, locus where, bool declare_simd = false, bool openacc = false, gfc_exec_op op = EXEC_NOP) { - tree omp_clauses = NULL_TREE, prev_clauses, chunk_size, c; + tree omp_clauses = NULL_TREE, prev_clauses = NULL_TREE, chunk_size, c; tree iterator = NULL_TREE; tree tree_block = NULL_TREE; stmtblock_t iter_block; @@ -3129,11 +3129,40 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, } break; case OMP_LIST_MAP: + iterator = NULL_TREE; + prev = NULL; + prev_clauses = omp_clauses; for (; n != NULL; n = n->next) { if (!n->sym->attr.referenced) continue; + if (iterator && prev->u2.ns != n->u2.ns) + { + /* Finish previous iterator group. */ + BLOCK_SUBBLOCKS (tree_block) = gfc_finish_block (&iter_block); + TREE_VEC_ELT (iterator, 5) = tree_block; + for (tree c = omp_clauses; c != prev_clauses; + c = OMP_CLAUSE_CHAIN (c)) + OMP_CLAUSE_DECL (c) = build_tree_list (iterator, + OMP_CLAUSE_DECL (c)); + prev_clauses = omp_clauses; + iterator = NULL_TREE; + } + if (n->u2.ns && (!prev || prev->u2.ns != n->u2.ns)) + { + /* Start a new iterator group. */ + gfc_init_block (&iter_block); + tree_block = make_node (BLOCK); + TREE_USED (tree_block) = 1; + BLOCK_VARS (tree_block) = NULL_TREE; + prev_clauses = omp_clauses; + iterator = handle_iterator (n->u2.ns, block, tree_block); + } + if (!iterator) + gfc_init_block (&iter_block); + prev = n; + bool always_modifier = false; tree node = build_omp_clause (input_location, OMP_CLAUSE_MAP); tree node2 = NULL_TREE; @@ -3332,7 +3361,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, TRUTH_NOT_EXPR, boolean_type_node, present); - gfc_add_expr_to_block (block, + gfc_add_expr_to_block (&iter_block, build3_loc (input_location, COND_EXPR, void_type_node, @@ -3392,7 +3421,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, tree type = TREE_TYPE (decl); tree ptr = gfc_conv_descriptor_data_get (decl); if (present) - ptr = gfc_build_cond_assign_expr (block, present, ptr, + ptr = gfc_build_cond_assign_expr (&iter_block, + present, ptr, null_pointer_node); gcc_assert (POINTER_TYPE_P (TREE_TYPE (ptr))); ptr = build_fold_indirect_ref (ptr); @@ -3420,7 +3450,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, ptr = gfc_conv_descriptor_data_get (decl); ptr = gfc_build_addr_expr (NULL, ptr); ptr = gfc_build_cond_assign_expr ( - block, present, ptr, null_pointer_node); + &iter_block, present, ptr, null_pointer_node); ptr = build_fold_indirect_ref (ptr); OMP_CLAUSE_DECL (node3) = ptr; } @@ -3509,7 +3539,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, TRUTH_ANDIF_EXPR, boolean_type_node, present, cond); - gfc_add_expr_to_block (block, + gfc_add_expr_to_block (&iter_block, build3_loc (input_location, COND_EXPR, void_type_node, @@ -3538,12 +3568,12 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, tree cond = build3_loc (input_location, COND_EXPR, void_type_node, present, cond_body, NULL_TREE); - gfc_add_expr_to_block (block, cond); + gfc_add_expr_to_block (&iter_block, cond); OMP_CLAUSE_SIZE (node) = var; } else { - gfc_add_block_to_block (block, &cond_block); + gfc_add_block_to_block (&iter_block, &cond_block); OMP_CLAUSE_SIZE (node) = size; } } @@ -3555,7 +3585,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, /* A single indirectref is handled by the middle end. */ gcc_assert (!POINTER_TYPE_P (TREE_TYPE (decl))); decl = TREE_OPERAND (decl, 0); - decl = gfc_build_cond_assign_expr (block, present, decl, + decl = gfc_build_cond_assign_expr (&iter_block, + present, decl, null_pointer_node); OMP_CLAUSE_DECL (node) = build_fold_indirect_ref (decl); } @@ -3589,7 +3620,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, size_type_node, cond, size, size_zero_node); - size = gfc_evaluate_now (size, block); + size = gfc_evaluate_now (size, &iter_block); OMP_CLAUSE_SIZE (node) = size; } } @@ -3608,7 +3639,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, && !(POINTER_TYPE_P (type) && GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (type)))) k = GOMP_MAP_FIRSTPRIVATE_POINTER; - gfc_trans_omp_array_section (block, op, n, decl, element, + gfc_trans_omp_array_section (&iter_block, + op, n, decl, element, !openacc, k, node, node2, node3, node4); } @@ -3626,12 +3658,12 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, gfc_init_se (&se, NULL); gfc_conv_expr (&se, n->expr); - gfc_add_block_to_block (block, &se.pre); + gfc_add_block_to_block (&iter_block, &se.pre); /* For BT_CHARACTER a pointer is returned. */ OMP_CLAUSE_DECL (node) = POINTER_TYPE_P (TREE_TYPE (se.expr)) ? build_fold_indirect_ref (se.expr) : se.expr; - gfc_add_block_to_block (block, &se.post); + gfc_add_block_to_block (&iter_block, &se.post); if (pointer || allocatable) { /* If it's a bare attach/detach clause, we just want @@ -3843,7 +3875,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, OMP_CLAUSE_DECL (node) = ptr; int rank = GFC_TYPE_ARRAY_RANK (type); OMP_CLAUSE_SIZE (node) - = gfc_full_array_size (block, inner, rank); + = gfc_full_array_size (&iter_block, inner, rank); tree elemsz = TYPE_SIZE_UNIT (gfc_get_element_type (type)); map_kind = OMP_CLAUSE_MAP_KIND (node); @@ -3981,7 +4013,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, /* An array element or section. */ bool element = lastref->u.ar.type == AR_ELEMENT; gomp_map_kind kind = GOMP_MAP_ATTACH_DETACH; - gfc_trans_omp_array_section (block, op, n, inner, element, + gfc_trans_omp_array_section (&iter_block, + op, n, inner, element, !openacc, kind, node, node2, node3, node4); } @@ -3993,6 +4026,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, finalize_map_clause: + if (!iterator) + gfc_add_block_to_block (block, &iter_block); omp_clauses = gfc_trans_add_clause (node, omp_clauses); if (node2) omp_clauses = gfc_trans_add_clause (node2, omp_clauses); @@ -4003,6 +4038,16 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, if (node5) omp_clauses = gfc_trans_add_clause (node5, omp_clauses); } + if (iterator) + { + /* Finish last iterator group. */ + BLOCK_SUBBLOCKS (tree_block) = gfc_finish_block (&iter_block); + TREE_VEC_ELT (iterator, 5) = tree_block; + for (tree c = omp_clauses; c != prev_clauses; + c = OMP_CLAUSE_CHAIN (c)) + OMP_CLAUSE_DECL (c) = build_tree_list (iterator, + OMP_CLAUSE_DECL (c)); + } break; case OMP_LIST_TO: case OMP_LIST_FROM: diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index 6e938296245..09e6b927d72 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -8858,10 +8858,17 @@ compute_iterator_count (tree t, gimple_seq *pre_p) endmbegin = fold_build2_loc (loc, POINTER_DIFF_EXPR, stype, end, begin); else endmbegin = fold_build2_loc (loc, MINUS_EXPR, type, end, begin); - tree stepm1 = fold_build2_loc (loc, MINUS_EXPR, stype, step, - build_int_cst (stype, 1)); - tree stepp1 = fold_build2_loc (loc, PLUS_EXPR, stype, step, - build_int_cst (stype, 1)); + /* Account for iteration stopping on the end value in Fortran rather + than before it. */ + tree stepm1 = step; + tree stepp1 = step; + if (!lang_GNU_Fortran ()) + { + stepm1 = fold_build2_loc (loc, MINUS_EXPR, stype, step, + build_int_cst (stype, 1)); + stepp1 = fold_build2_loc (loc, PLUS_EXPR, stype, step, + build_int_cst (stype, 1)); + } tree pos = fold_build2_loc (loc, PLUS_EXPR, stype, unshare_expr (endmbegin), stepm1); pos = fold_build2_loc (loc, TRUNC_DIV_EXPR, stype, pos, step); @@ -8913,6 +8920,7 @@ build_iterator_loop (tree c, gimple_seq *pre_p, tree *last_bind) gimplify_ctxp->into_ssa = saved_into_ssa; } tree block = TREE_VEC_ELT (TREE_PURPOSE (t), 5); + tree block_stmts = lang_GNU_Fortran () ? BLOCK_SUBBLOCKS (block) : NULL_TREE; *last_bind = build3 (BIND_EXPR, void_type_node, BLOCK_VARS (block), NULL, block); TREE_SIDE_EFFECTS (*last_bind) = 1; @@ -8925,6 +8933,7 @@ build_iterator_loop (tree c, gimple_seq *pre_p, tree *last_bind) tree end = TREE_VEC_ELT (it, 2); tree step = TREE_VEC_ELT (it, 3); tree orig_step = TREE_VEC_ELT (it, 4); + block = TREE_VEC_ELT (it, 5); tree type = TREE_TYPE (var); location_t loc = DECL_SOURCE_LOCATION (var); /* Emit: @@ -8935,9 +8944,9 @@ build_iterator_loop (tree c, gimple_seq *pre_p, tree *last_bind) var = var + step; cond_label: if (orig_step > 0) { - if (var < end) goto beg_label; + if (var < end) goto beg_label; // <= for Fortran } else { - if (var > end) goto beg_label; + if (var > end) goto beg_label; // >= for Fortran } for each iterator, with inner iterators added to the ... above. */ @@ -8963,10 +8972,12 @@ build_iterator_loop (tree c, gimple_seq *pre_p, tree *last_bind) append_to_statement_list_force (tem, p); tem = build1 (LABEL_EXPR, void_type_node, cond_label); append_to_statement_list (tem, p); - tree cond = fold_build2_loc (loc, LT_EXPR, boolean_type_node, var, end); + tree cond = fold_build2_loc (loc, lang_GNU_Fortran () ? LE_EXPR : LT_EXPR, + boolean_type_node, var, end); tree pos = fold_build3_loc (loc, COND_EXPR, void_type_node, cond, build_and_jump (&beg_label), void_node); - cond = fold_build2_loc (loc, GT_EXPR, boolean_type_node, var, end); + cond = fold_build2_loc (loc, lang_GNU_Fortran () ? GE_EXPR : GT_EXPR, + boolean_type_node, var, end); tree neg = fold_build3_loc (loc, COND_EXPR, void_type_node, cond, build_and_jump (&beg_label), void_node); tree osteptype = TREE_TYPE (orig_step); @@ -8975,6 +8986,11 @@ build_iterator_loop (tree c, gimple_seq *pre_p, tree *last_bind) tem = fold_build3_loc (loc, COND_EXPR, void_type_node, cond, pos, neg); append_to_statement_list_force (tem, p); p = &BIND_EXPR_BODY (bind); + /* The Fortran front-end stashes statements into the BLOCK_SUBBLOCKS + of the last element of the first iterator. These should go into the + body of the innermost loop. */ + if (!TREE_CHAIN (it)) + append_to_statement_list_force (block_stmts, p); } return p; @@ -11398,6 +11414,8 @@ omp_accumulate_sibling_list (enum omp_region_type region_type, poly_offset_int coffset; poly_int64 cbitpos; tree ocd = OMP_ITERATOR_CLAUSE_DECL (grp_end); + tree iterator = OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (grp_end)) + ? TREE_PURPOSE (OMP_CLAUSE_DECL (grp_end)) : NULL_TREE; bool openmp = !(region_type & ORT_ACC); bool target = (region_type & ORT_TARGET) != 0; tree *continue_at = NULL; @@ -11476,7 +11494,7 @@ omp_accumulate_sibling_list (enum omp_region_type region_type, if (struct_map_to_clause == NULL) struct_map_to_clause = new hash_map; - if (variable_offset) + if (variable_offset && !iterator) str_kind = GOMP_MAP_STRUCT_UNORD; tree l = build_omp_clause (OMP_CLAUSE_LOCATION (grp_end), OMP_CLAUSE_MAP); diff --git a/gcc/testsuite/gfortran.dg/gomp/target-iterator-1.f90 b/gcc/testsuite/gfortran.dg/gomp/target-iterator-1.f90 new file mode 100644 index 00000000000..25abbaf741e --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/target-iterator-1.f90 @@ -0,0 +1,26 @@ +! { dg-do compile } +! { dg-options "-fopenmp" } + +program main + implicit none + + integer, parameter :: DIM1 = 17 + integer, parameter :: DIM2 = 39 + type :: array_ptr + integer, pointer :: ptr(:) + end type + + type (array_ptr) :: x(DIM1), y(DIM1) + + !$omp target map (iterator(i=1:DIM1), to: x(i)%ptr(:)) + !$omp end target + + !$omp target map (iterator(i=1:DIM1), to: x(i)%ptr(:), y(i)%ptr(:)) + !$omp end target + + !$omp target map (iterator(i=1:DIM1), to: x(i)%ptr(:) + 3) ! { dg-error "Syntax error in OpenMP variable list at .1." } + !$omp end target ! { dg-error "Unexpected \\\!\\\$OMP END TARGET statement at .1." } + + !$omp target map(iterator(i=1:DIM1), iterator(j=1:DIM2), to: x(i)%ptr(j)) ! { dg-error "too many 'iterator' modifiers at .1." } + !$omp end target ! { dg-error "Unexpected \\\!\\\$OMP END TARGET statement at .1." } +end program diff --git a/gcc/testsuite/gfortran.dg/gomp/target-iterator-2.f90 b/gcc/testsuite/gfortran.dg/gomp/target-iterator-2.f90 new file mode 100644 index 00000000000..b7d7501cf63 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/target-iterator-2.f90 @@ -0,0 +1,27 @@ +! { dg-do compile } +! { dg-options "-fopenmp" } + +program main + implicit none + + integer, parameter :: DIM = 40 + type :: array_ptr + integer, pointer :: ptr(:) + end type + + type (array_ptr) :: x(DIM), y(DIM), z(DIM) + + !$omp target map(iterator(i=1: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(1)%ptr(1) = 0 + !$omp end target + + !$omp target map(iterator(i=1:10, j=1:20), to: x(i)) ! { dg-error "iterator variable .j. not used in clause expression" } + !$omp end target + + !$omp target map(iterator(i=1:10, j=1:20, k=1:30), to: x(i), y(j), z(k)) + !$omp end target + ! { dg-error "iterator variable .i. not used in clause expression" "" { target *-*-* } .-2 } + ! { dg-error "iterator variable .j. not used in clause expression" "" { target *-*-* } .-3 } + ! { dg-error "iterator variable .k. not used in clause expression" "" { target *-*-* } .-4 } +end program diff --git a/gcc/testsuite/gfortran.dg/gomp/target-iterator-3.f90 b/gcc/testsuite/gfortran.dg/gomp/target-iterator-3.f90 new file mode 100644 index 00000000000..3cff65ab072 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/target-iterator-3.f90 @@ -0,0 +1,24 @@ +! { dg-do compile } +! { dg-options "-fopenmp -fdump-tree-gimple" } + +program main + implicit none + + integer, parameter :: DIM1 = 17 + integer, parameter :: DIM2 = 27 + type :: ptr_t + integer, pointer :: ptr(:) + end type + + type (ptr_t) :: x(DIM1), y(DIM2) + + !$omp target map(iterator(i=1:DIM1), to: x(i)%ptr(:)) map(iterator(i=1:DIM2), from: y(i)%ptr(:)) + !$omp end target +end program + +! { dg-final { scan-tree-dump-times "if \\(i <= 17\\) goto ; else goto ;" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "if \\(i <= 27\\) goto ; else goto ;" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "map\\(iterator\\(integer\\(kind=4\\) i=1:17:1\\):iterator_array=D\.\[0-9\]+:to:" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "map\\(iterator\\(integer\\(kind=4\\) i=1:27:1\\):iterator_array=D\.\[0-9\]+:from:" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "map\\(iterator\\(integer\\(kind=4\\) i=1:17:1\\):iterator_array=D\.\[0-9\]+:attach:" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "map\\(iterator\\(integer\\(kind=4\\) i=1:27:1\\):iterator_array=D\.\[0-9\]+:attach:" 1 "gimple" } } diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc index e43f30818d0..8af83b934f0 100644 --- a/gcc/tree-pretty-print.cc +++ b/gcc/tree-pretty-print.cc @@ -1688,7 +1688,9 @@ dump_block_node (pretty_printer *pp, tree block, int spc, dump_flags_t flags) newline_and_indent (pp, spc + 2); } - if (BLOCK_SUBBLOCKS (block)) + if (BLOCK_SUBBLOCKS (block) + && (!lang_GNU_Fortran () + || TREE_CODE (BLOCK_SUBBLOCKS (block)) != STATEMENT_LIST)) { pp_string (pp, "SUBBLOCKS: "); for (t = BLOCK_SUBBLOCKS (block); t; t = BLOCK_CHAIN (t)) diff --git a/libgomp/target.c b/libgomp/target.c index c69418f0b78..dabe88bc900 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -972,14 +972,74 @@ 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_ATTACH: return "GOMP_MAP_ATTACH"; + case GOMP_MAP_DETACH: return "GOMP_MAP_DETACH"; + case GOMP_MAP_STRUCT: return "GOMP_MAP_STRUCT"; + case GOMP_MAP_STRUCT_UNORD: return "GOMP_MAP_STRUCT_UNORD"; + default: return "unknown"; + } +} + +static void +gomp_add_map (size_t idx, size_t *new_idx, + void ***hostaddrs, size_t **sizes, unsigned short **skinds, + void ***new_hostaddrs, size_t **new_sizes, + unsigned short **new_kinds, size_t *iterator_count) +{ + if ((*sizes)[idx] == SIZE_MAX) + { + uintptr_t *iterator_array = (*hostaddrs)[idx]; + size_t count = *iterator_array++; + for (size_t i = 0; i < count; i++) + { + (*new_hostaddrs)[*new_idx] = (void *) *iterator_array++; + (*new_sizes)[*new_idx] = *iterator_array++; + (*new_kinds)[*new_idx] = (*skinds)[idx]; + iterator_count[*new_idx] = i + 1; + gomp_debug (1, + "Expanding map %ld <%s>: " + "hostaddrs[%ld] = %p, sizes[%ld] = %ld\n", + idx, kind_to_name ((*new_kinds)[*new_idx]), + *new_idx, (*new_hostaddrs)[*new_idx], + *new_idx, (*new_sizes)[*new_idx]); + (*new_idx)++; + } + } + else + { + (*new_hostaddrs)[*new_idx] = (*hostaddrs)[idx]; + (*new_sizes)[*new_idx] = (*sizes)[idx]; + (*new_kinds)[*new_idx] = (*skinds)[idx]; + iterator_count[*new_idx] = 0; + (*new_idx)++; + } +} + /* 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. */ + 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) + void **kinds, size_t **iterator_count) { bool iterator_p = false; size_t map_count = 0; @@ -1006,33 +1066,36 @@ gomp_merge_iterator_maps (size_t *mapnum, void ***hostaddrs, size_t **sizes, 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) + int map_type = get_kind (true, *skinds, i) & 0xff; + if (map_type == GOMP_MAP_STRUCT || map_type == GOMP_MAP_STRUCT_UNORD) { - uintptr_t *iterator_array = (*hostaddrs)[i]; - size_t count = iterator_array[0]; - for (int j = 1; j < count * 2 + 1; j += 2) + size_t field_count = (*sizes)[i]; + + gomp_add_map (i, &new_idx, hostaddrs, sizes, skinds, + &new_hostaddrs, &new_sizes, &new_kinds, *iterator_count); + + for (size_t j = i + 1; j <= i + field_count; j++) { - 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++; + if ((*sizes)[j] == SIZE_MAX) + { + uintptr_t *iterator_array = (*hostaddrs)[j]; + size_t count = iterator_array[0]; + new_sizes[i] += count - 1; + } + gomp_add_map (j, &new_idx, hostaddrs, sizes, skinds, + &new_hostaddrs, &new_sizes, &new_kinds, + *iterator_count); } + gomp_debug (1, "Map %ld new field count = %ld\n", i, new_sizes[i]); + i += field_count; } else - { - new_hostaddrs[new_idx] = (*hostaddrs)[i]; - new_sizes[new_idx] = (*sizes)[i]; - new_kinds[new_idx] = (*skinds)[i]; - new_idx++; - } + gomp_add_map (i, &new_idx, hostaddrs, sizes, skinds, + &new_hostaddrs, &new_sizes, &new_kinds, *iterator_count); } *mapnum = map_count; @@ -1060,9 +1123,10 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, 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); + &kinds, &iterator_count); struct target_mem_desc *tgt = gomp_malloc (sizeof (*tgt) + sizeof (tgt->list[0]) * mapnum); tgt->list_count = mapnum; @@ -1912,14 +1976,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) @@ -1957,6 +2024,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, free (hostaddrs); free (sizes); free (kinds); + free (iterator_count); } return tgt; @@ -2225,6 +2293,7 @@ gomp_update (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs, 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; @@ -2234,7 +2303,7 @@ gomp_update (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs, if (short_mapkind) iterators_p = gomp_merge_iterator_maps (&mapnum, &hostaddrs, &sizes, - &kinds); + &kinds, &iterator_count); gomp_mutex_lock (&devicep->lock); if (devicep->state == GOMP_DEVICE_FINALIZED) @@ -2335,6 +2404,7 @@ gomp_update (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs, free (hostaddrs); free (sizes); free (kinds); + free (iterator_count); } } diff --git a/libgomp/testsuite/libgomp.fortran/target-map-iterators-1.f90 b/libgomp/testsuite/libgomp.fortran/target-map-iterators-1.f90 new file mode 100644 index 00000000000..80e077e69fd --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/target-map-iterators-1.f90 @@ -0,0 +1,45 @@ +! { dg-do run } + +! Test transfer of dynamically-allocated arrays to target using map +! iterators. + +program test + implicit none + + integer, parameter :: DIM1 = 8 + integer, parameter :: DIM2 = 15 + + type :: array_ptr + integer, pointer :: arr(:) + end type + + type (array_ptr) :: x(DIM1) + integer :: expected, sum, i, j + + expected = mkarray () + + !$omp target map(iterator(i=1:DIM1), to: x(i)%arr(:)) map(from: sum) + sum = 0 + do i = 1, DIM1 + do j = 1, DIM2 + sum = sum + x(i)%arr(j) + end do + end do + !$omp end target + + if (sum .ne. expected) stop 1 +contains + integer function mkarray () + integer :: exp = 0 + + do i = 1, DIM1 + allocate (x(i)%arr(DIM2)) + do j = 1, DIM2 + x(i)%arr(j) = i * j + exp = exp + x(i)%arr(j) + end do + end do + + mkarray = exp + end function +end program diff --git a/libgomp/testsuite/libgomp.fortran/target-map-iterators-2.f90 b/libgomp/testsuite/libgomp.fortran/target-map-iterators-2.f90 new file mode 100644 index 00000000000..cf0e7fbd9b3 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/target-map-iterators-2.f90 @@ -0,0 +1,45 @@ +! { dg-do run } + +! Test transfer of dynamically-allocated arrays from target using map +! iterators. + +program test + implicit none + + integer, parameter :: DIM1 = 8 + integer, parameter :: DIM2 = 15 + + type :: array_ptr + integer, pointer :: arr(:) + end type + + type (array_ptr) :: x(DIM1) + integer :: expected, sum, i, j + + call mkarray + + !$omp target map(iterator(i=1:DIM1), from: x(i)%arr(:)) map(from: expected) + expected = 0 + do i = 1, DIM1 + do j = 1, DIM2 + x(i)%arr(j) = (i+1) * (j+1) + expected = expected + x(i)%arr(j) + end do + end do + !$omp end target + + sum = 0 + do i = 1, DIM1 + do j = 1, DIM2 + sum = sum + x(i)%arr(j) + end do + end do + + if (sum .ne. expected) stop 1 +contains + subroutine mkarray + do i = 1, DIM1 + allocate (x(i)%arr(DIM2)) + end do + end subroutine +end program diff --git a/libgomp/testsuite/libgomp.fortran/target-map-iterators-3.f90 b/libgomp/testsuite/libgomp.fortran/target-map-iterators-3.f90 new file mode 100644 index 00000000000..8072c074557 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/target-map-iterators-3.f90 @@ -0,0 +1,57 @@ +! { dg-do run } + +! Test transfer of dynamically-allocated arrays to target using map +! iterators, with multiple iterators and function calls in the iterator +! expression. + +program test + implicit none + + integer, parameter :: DIM1 = 16 + integer, parameter :: DIM2 = 4 + + type :: array_ptr + integer, pointer :: arr(:) + end type + + type (array_ptr) :: x(DIM1), y(DIM1) + integer :: expected, sum, i, j + + expected = mkarrays () + + !$omp target map(iterator(i=0:DIM1/4-1, j=0:3), to: x(f (i, j))%arr(:)) & + !$omp map(iterator(i=1:DIM1), to: y(i)%arr(:)) & + !$omp map(from: sum) + sum = 0 + do i = 1, DIM1 + do j = 1, DIM2 + sum = sum + x(i)%arr(j) * y(i)%arr(j) + end do + end do + !$omp end target + + print *, sum, expected + if (sum .ne. expected) stop 1 +contains + integer function mkarrays () + integer :: exp = 0 + + do i = 1, DIM1 + allocate (x(i)%arr(DIM2)) + allocate (y(i)%arr(DIM2)) + do j = 1, DIM2 + x(i)%arr(j) = i * j + y(i)%arr(j) = i + j + exp = exp + x(i)%arr(j) * y(i)%arr(j) + end do + end do + + mkarrays = exp + end function + + integer function f (i, j) + integer, intent(in) :: i, j + + f = i * 4 + j + 1 + end function +end program From patchwork Tue Sep 3 17:11:53 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: 1980240 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=brldXTrL; 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 4Wysdf5StWz1yg9 for ; Wed, 4 Sep 2024 03:13:10 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id F1C1F385020E for ; Tue, 3 Sep 2024 17:13:08 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from mail-ej1-x634.google.com (mail-ej1-x634.google.com [IPv6:2a00:1450:4864:20::634]) by sourceware.org (Postfix) with ESMTPS id 48D62385EC2F for ; Tue, 3 Sep 2024 17:12:46 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 48D62385EC2F 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 48D62385EC2F Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=2a00:1450:4864:20::634 ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1725383569; cv=none; b=fT6vpEHda4Na6FbsvQLrKJpBo2ofUOA6c6WntbQG9zTaNm+yRUbi2yuvvFpjU0jhtxo3Ytd2dHAufdX6jhybpDJXr0sR7G8cAn9lm1ggOVkPqPmmHAk/kkc4tL9qeLpkSKkH2NvqftOH4wLFNxVf9zPd/hcMVdTCuqXzt5EN5vU= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1725383569; c=relaxed/simple; bh=8IlS9q67XioC8nYg6vp8bs68dU98OWxghQ2zgHXX9eg=; h=DKIM-Signature:Message-ID:Date:MIME-Version:To:Subject:From; b=tz/SeUXVl7Yl/5kPX18nFozCwvZxwgCdIuyTtHjJQsanCjeiZyZpKQU6OKg+GdOC/Z6mHVJiUguQ2qjcQWTLTnKmxzWYAUVYpnaUcEaeSIymrniFj+orKVP7qL/Fk1ypGrN/13U1oj+RWKNh2L52mf3UJbEz2yx2t0NqRE9WxwU= ARC-Authentication-Results: i=1; server2.sourceware.org Received: by mail-ej1-x634.google.com with SMTP id a640c23a62f3a-a8a1acb51a7so180386166b.2 for ; Tue, 03 Sep 2024 10:12:46 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre-com.20230601.gappssmtp.com; s=20230601; t=1725383565; x=1725988365; darn=gcc.gnu.org; h=in-reply-to:from:cc:content-language:subject:references:to :user-agent:mime-version:date:message-id:from:to:cc:subject:date :message-id:reply-to; bh=9ArZnYRXSFzLKSeuoNDQvUfAeP7WhEIiwyk/my3zOGA=; b=brldXTrL89boqJAXbEAACyiTkpkbtTxGv2tJMa5JZPba7Y7n2jJARhUNnKakNpMJaE NRyL+o3nAEHn7a3SvR0P66Mw4MazcAdFds4KENceIQfd4TgL1zIRpo+LfDH+jrlGpl3F tQoUsbmoaS07QshhUlGrBHkxie2DX1kojNxdM3Kri7ETKbBxNXEtn9aeNkmaMsTfONO3 8AuRk754cNP2rrp9Yq1mz8/qPzEj0JIqfwLKzwG/HG3liD/fyIAl21HMEdRVQB6XADPI 2483HozkfMqhXv83O7gr4M1XR8El44W1etZ9SrqxT0YFOmm2k4+hqtR2gzAet7VjrAa2 JuKQ== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1725383565; x=1725988365; h=in-reply-to:from:cc: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=9ArZnYRXSFzLKSeuoNDQvUfAeP7WhEIiwyk/my3zOGA=; b=Vzqjb/ks9AyvepqtVeO7qp/5sBoResRd2D1fsUn452FTi5abzskUXXJH8UXRP4CuB5 3nzaeNQLu+yr6e4QXHMypFQi+0qzmuaZjUEJCgPop/e+DdOCAKAM4qA6K+5wM1gqCVr1 8QZyKb7i2hfOCSLJ5a5FT0yjgNfR1+2oB+UyXGVp4lwuVWkWiFj0jciC0wuow1FbFw+x KBkES2IYcdNoYdDUIx+OKSs6OGXL0CtlKmyYSwhK0XhTRNDquwOXqzIjAI4aoCbrkfHk 5Jc9fdW018JwrDc0ek+ybpY8FyFpEZwwZK13c5uv1DtX3FJ+nufI2XzekrPQggB5poL1 kmZw== X-Gm-Message-State: AOJu0YzTAnus/Jdta6BLuRQVqsHoTrRaFOmqYaR1F4hTrG21CndvWP7M RKNaALMRAqXxJ1s1i9elr+sN2Tire94tYFulrh5qvIYBg2DWVy6+D4zKvMawT6tsmo3vjsS14Kh m X-Google-Smtp-Source: AGHT+IEc6CO0393AzSo8T9OdlvJQ1nDM8YvW5/aZMmsp1A2kRNgfxcI/QKhV461DAR4Is2R1tFTltA== X-Received: by 2002:a17:907:7292:b0:a88:a49a:59e2 with SMTP id a640c23a62f3a-a8a1d4c336bmr311972066b.44.1725383564131; Tue, 03 Sep 2024 10:12:44 -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-a8989196980sm703573066b.126.2024.09.03.10.12.43 (version=TLS1_3 cipher=TLS_AES_128_GCM_SHA256 bits=128/128); Tue, 03 Sep 2024 10:12:43 -0700 (PDT) Message-ID: <6cae3adf-eeb9-4731-8011-2c878e8de9df@baylibre.com> Date: Tue, 3 Sep 2024 18:11:53 +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 5/5] openmp, fortran: Add support for iterators in OpenMP 'target update' constructs (Fortran) Content-Language: en-GB Cc: fortran@gcc.gnu.org From: Kwok Cheung Yeung In-Reply-To: <77f9ccb8-6f5e-4462-aa32-71f74fd7ff26@baylibre.com> X-Spam-Status: No, score=-13.6 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 adds parsing and translation of the 'to' and 'from' clauses for the 'target update' construct in Fortran. From cfb6b76da5bba038d854d510a4fd44ddf4fa8f1f Mon Sep 17 00:00:00 2001 From: Kwok Cheung Yeung Date: Mon, 2 Sep 2024 19:34:29 +0100 Subject: [PATCH 5/5] openmp, fortran: Add support for iterators in OpenMP 'target update' constructs (Fortran) This adds Fortran support for iterators in 'to' and 'from' clauses in the 'target update' OpenMP directive. 2024-09-02 Kwok Cheung Yeung gcc/fortran/ * dump-parse-tree.cc (show_omp_namelist): Add iterator support for OMP_LIST_TO and OMP_LIST_FROM. * openmp.cc (gfc_free_omp_clauses): Free namespace for OMP_LIST_TO and OMP_LIST_FROM. (gfc_match_motion_var_list): Parse 'iterator' modifier. (resolve_omp_clauses): Resolve iterators for OMP_LIST_TO and OMP_LIST_FROM. * trans-openmp.cc (gfc_trans_omp_clauses): Handle iterators in OMP_LIST_TO and OMP_LIST_FROM clauses. gcc/testsuite/ * gfortran.dg/gomp/target-update-iterator-1.f90: New. * gfortran.dg/gomp/target-update-iterator-2.f90: New. * gfortran.dg/gomp/target-update-iterator-3.f90: New. libgomp/ * testsuite/libgomp.fortran/target-update-iterators-1.f90: New. * testsuite/libgomp.fortran/target-update-iterators-2.f90: New. * testsuite/libgomp.fortran/target-update-iterators-3.f90: New. --- gcc/fortran/dump-parse-tree.cc | 7 +- gcc/fortran/openmp.cc | 62 +++++++++++++-- gcc/fortran/trans-openmp.cc | 52 +++++++++++-- .../gomp/target-update-iterator-1.f90 | 25 ++++++ .../gomp/target-update-iterator-2.f90 | 22 ++++++ .../gomp/target-update-iterator-3.f90 | 23 ++++++ .../target-update-iterators-1.f90 | 68 ++++++++++++++++ .../target-update-iterators-2.f90 | 62 +++++++++++++++ .../target-update-iterators-3.f90 | 77 +++++++++++++++++++ 9 files changed, 386 insertions(+), 12 deletions(-) create mode 100644 gcc/testsuite/gfortran.dg/gomp/target-update-iterator-1.f90 create mode 100644 gcc/testsuite/gfortran.dg/gomp/target-update-iterator-2.f90 create mode 100644 gcc/testsuite/gfortran.dg/gomp/target-update-iterator-3.f90 create mode 100644 libgomp/testsuite/libgomp.fortran/target-update-iterators-1.f90 create mode 100644 libgomp/testsuite/libgomp.fortran/target-update-iterators-2.f90 create mode 100644 libgomp/testsuite/libgomp.fortran/target-update-iterators-3.f90 diff --git a/gcc/fortran/dump-parse-tree.cc b/gcc/fortran/dump-parse-tree.cc index 0272a443f65..1a602fb953c 100644 --- a/gcc/fortran/dump-parse-tree.cc +++ b/gcc/fortran/dump-parse-tree.cc @@ -1350,7 +1350,8 @@ show_omp_namelist (int list_type, gfc_omp_namelist *n) { gfc_current_ns = ns_curr; if (list_type == OMP_LIST_AFFINITY || list_type == OMP_LIST_DEPEND - || list_type == OMP_LIST_MAP) + || list_type == OMP_LIST_MAP + || list_type == OMP_LIST_TO || list_type == OMP_LIST_FROM) { gfc_current_ns = n->u2.ns ? n->u2.ns : ns_curr; if (n->u2.ns != ns_iter) @@ -1366,6 +1367,10 @@ show_omp_namelist (int list_type, gfc_omp_namelist *n) fputs ("DEPEND (", dumpfile); else if (list_type == OMP_LIST_MAP) fputs ("MAP (", dumpfile); + else if (list_type == OMP_LIST_TO) + fputs ("TO (", dumpfile); + else if (list_type == OMP_LIST_FROM) + fputs ("FROM (", dumpfile); else gcc_unreachable (); } diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc index 996126e6e7f..4eb4a8e53e2 100644 --- a/gcc/fortran/openmp.cc +++ b/gcc/fortran/openmp.cc @@ -192,7 +192,8 @@ gfc_free_omp_clauses (gfc_omp_clauses *c) for (i = 0; i < OMP_LIST_NUM; i++) gfc_free_omp_namelist (c->lists[i], i == OMP_LIST_AFFINITY || i == OMP_LIST_DEPEND - || i == OMP_LIST_MAP, + || i == OMP_LIST_MAP + || i == OMP_LIST_TO || i == OMP_LIST_FROM, i == OMP_LIST_ALLOCATE, i == OMP_LIST_USES_ALLOCATORS); gfc_free_expr_list (c->wait_list); @@ -1362,17 +1363,65 @@ gfc_match_motion_var_list (const char *str, gfc_omp_namelist **list, if (m != MATCH_YES) return m; - match m_present = gfc_match (" present : "); + gfc_namespace *ns_iter = NULL, *ns_curr = gfc_current_ns; + int present_modifier = 0, iterator_modifier = 0; + locus present_locus = gfc_current_locus, iterator_locus = gfc_current_locus; - m = gfc_match_omp_variable_list ("", list, false, NULL, headp, true, true); + for (;;) + { + locus current_locus = gfc_current_locus; + if (gfc_match ("present ") == MATCH_YES) + { + if (present_modifier++ == 1) + present_locus = current_locus; + } + else if (gfc_match_iterator (&ns_iter, true) == MATCH_YES) + { + if (iterator_modifier++ == 1) + iterator_locus = current_locus; + } + else + break; + gfc_match (", "); + } + + if (present_modifier > 1) + { + gfc_error ("too many % modifiers at %L", + &present_locus); + return MATCH_ERROR; + } + if (iterator_modifier > 1) + { + gfc_error ("too many % modifiers at %L", + &iterator_locus); + return MATCH_ERROR; + } + + if (ns_iter) + gfc_current_ns = ns_iter; + + const char *exp = (present_modifier || iterator_modifier) ? " :" : ""; + m = gfc_match_omp_variable_list (exp, list, false, NULL, headp, true, true); + gfc_current_ns = ns_curr; if (m != MATCH_YES) return m; - if (m_present == MATCH_YES) + + if (present_modifier || iterator_modifier) { gfc_omp_namelist *n; for (n = **headp; n; n = n->next) - n->u.present_modifier = true; + { + if (present_modifier) + n->u.present_modifier = true; + if (iterator_modifier) + { + n->u2.ns = ns_iter; + ns_iter->refs++; + } + } } + return MATCH_YES; } @@ -8436,7 +8485,8 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, for (; n != NULL; n = n->next) { if ((list == OMP_LIST_DEPEND || list == OMP_LIST_AFFINITY - || list == OMP_LIST_MAP) + || list == OMP_LIST_MAP + || list == OMP_LIST_TO || list == OMP_LIST_FROM) && n->u2.ns && !n->u2.ns->resolved) { n->u2.ns->resolved = 1; diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc index a9929430e53..1be8f2ad806 100644 --- a/gcc/fortran/trans-openmp.cc +++ b/gcc/fortran/trans-openmp.cc @@ -4052,11 +4052,40 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, case OMP_LIST_TO: case OMP_LIST_FROM: case OMP_LIST_CACHE: + iterator = NULL_TREE; + prev = NULL; + prev_clauses = omp_clauses; for (; n != NULL; n = n->next) { if (!n->sym->attr.referenced) continue; + if (iterator && prev->u2.ns != n->u2.ns) + { + /* Finish previous iterator group. */ + BLOCK_SUBBLOCKS (tree_block) = gfc_finish_block (&iter_block); + TREE_VEC_ELT (iterator, 5) = tree_block; + for (tree c = omp_clauses; c != prev_clauses; + c = OMP_CLAUSE_CHAIN (c)) + OMP_CLAUSE_DECL (c) = build_tree_list (iterator, + OMP_CLAUSE_DECL (c)); + prev_clauses = omp_clauses; + iterator = NULL_TREE; + } + if (n->u2.ns && (!prev || prev->u2.ns != n->u2.ns)) + { + /* Start a new iterator group. */ + gfc_init_block (&iter_block); + tree_block = make_node (BLOCK); + TREE_USED (tree_block) = 1; + BLOCK_VARS (tree_block) = NULL_TREE; + prev_clauses = omp_clauses; + iterator = handle_iterator (n->u2.ns, block, tree_block); + } + if (!iterator) + gfc_init_block (&iter_block); + prev = n; + switch (list) { case OMP_LIST_TO: @@ -4094,7 +4123,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, ptr = build_fold_indirect_ref (ptr); OMP_CLAUSE_DECL (node) = ptr; OMP_CLAUSE_SIZE (node) - = gfc_full_array_size (block, decl, + = gfc_full_array_size (&iter_block, decl, GFC_TYPE_ARRAY_RANK (type)); tree elemsz = TYPE_SIZE_UNIT (gfc_get_element_type (type)); @@ -4119,7 +4148,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, { gfc_conv_expr_reference (&se, n->expr); ptr = se.expr; - gfc_add_block_to_block (block, &se.pre); + gfc_add_block_to_block (&iter_block, &se.pre); OMP_CLAUSE_SIZE (node) = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (ptr))); } @@ -4128,9 +4157,9 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, gfc_conv_expr_descriptor (&se, n->expr); ptr = gfc_conv_array_data (se.expr); tree type = TREE_TYPE (se.expr); - gfc_add_block_to_block (block, &se.pre); + gfc_add_block_to_block (&iter_block, &se.pre); OMP_CLAUSE_SIZE (node) - = gfc_full_array_size (block, se.expr, + = gfc_full_array_size (&iter_block, se.expr, GFC_TYPE_ARRAY_RANK (type)); tree elemsz = TYPE_SIZE_UNIT (gfc_get_element_type (type)); @@ -4139,7 +4168,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, = fold_build2 (MULT_EXPR, gfc_array_index_type, OMP_CLAUSE_SIZE (node), elemsz); } - gfc_add_block_to_block (block, &se.post); + gfc_add_block_to_block (&iter_block, &se.post); gcc_assert (POINTER_TYPE_P (TREE_TYPE (ptr))); OMP_CLAUSE_DECL (node) = build_fold_indirect_ref (ptr); } @@ -4147,8 +4176,21 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, OMP_CLAUSE_MOTION_PRESENT (node) = 1; if (list == OMP_LIST_CACHE && n->u.map.readonly) OMP_CLAUSE__CACHE__READONLY (node) = 1; + + if (!iterator) + gfc_add_block_to_block (block, &iter_block); omp_clauses = gfc_trans_add_clause (node, omp_clauses); } + if (iterator) + { + /* Finish last iterator group. */ + BLOCK_SUBBLOCKS (tree_block) = gfc_finish_block (&iter_block); + TREE_VEC_ELT (iterator, 5) = tree_block; + for (tree c = omp_clauses; c != prev_clauses; + c = OMP_CLAUSE_CHAIN (c)) + OMP_CLAUSE_DECL (c) = build_tree_list (iterator, + OMP_CLAUSE_DECL (c)); + } break; case OMP_LIST_USES_ALLOCATORS: /* Ignore pre-defined allocators as no special treatment is needed. */ diff --git a/gcc/testsuite/gfortran.dg/gomp/target-update-iterator-1.f90 b/gcc/testsuite/gfortran.dg/gomp/target-update-iterator-1.f90 new file mode 100644 index 00000000000..08dc3d79911 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/target-update-iterator-1.f90 @@ -0,0 +1,25 @@ +! { dg-do compile } +! { dg-options "-fopenmp" } + +program test + implicit none + + integer, parameter :: DIM1 = 17 + integer, parameter :: DIM2 = 39 + + type :: array_ptr + integer, pointer :: ptr(:) + end type + + type (array_ptr) :: x(DIM1), y(DIM1) + + !$omp target update to (iterator(i=1:DIM1): x(i)%ptr(:)) + + !$omp target update to (iterator(i=1:DIM1): x(i)%ptr(:DIM2), y(i)%ptr(:)) + + !$omp target update to (iterator(i=1:DIM1), present: x(i)%ptr(:)) + + !$omp target update to (iterator(i=1:DIM1), iterator(j=i:DIM2): x(i)%ptr(j)) ! { dg-error "too many 'iterator' modifiers at .1." } + + !$omp target update to (iterator(i=1:DIM1), something: x(i, j)) ! { dg-error "Failed to match clause at .1." } +end program diff --git a/gcc/testsuite/gfortran.dg/gomp/target-update-iterator-2.f90 b/gcc/testsuite/gfortran.dg/gomp/target-update-iterator-2.f90 new file mode 100644 index 00000000000..89f645bda23 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/target-update-iterator-2.f90 @@ -0,0 +1,22 @@ +! { dg-do compile } +! { dg-options "-fopenmp" } + +program test + implicit none + + integer, parameter :: DIM1 = 17 + integer, parameter :: DIM2 = 39 + + type :: array_ptr + integer, pointer :: ptr(:) + end type + + type (array_ptr) :: x(DIM1), y(DIM1), z(DIM1) + + !$omp target update to(iterator(i=1:10): x) ! { dg-error "iterator variable .i. not used in clause expression" } + !$omp target update from(iterator(i=1:10, j=1:20): x(i)) ! { dg-error "iterator variable .j. not used in clause expression" } + !$omp target update to(iterator(i=1:10, j=1:20, k=1: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 } +end program diff --git a/gcc/testsuite/gfortran.dg/gomp/target-update-iterator-3.f90 b/gcc/testsuite/gfortran.dg/gomp/target-update-iterator-3.f90 new file mode 100644 index 00000000000..a8b8ce8f171 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/target-update-iterator-3.f90 @@ -0,0 +1,23 @@ +! { dg-do compile } +! { dg-options "-fopenmp -fdump-tree-gimple" } + +program test + implicit none + + integer, parameter :: DIM1 = 17 + integer, parameter :: DIM2 = 39 + + type :: array_ptr + integer, pointer :: ptr(:) + end type + + type (array_ptr) :: x(DIM1, DIM2), y(DIM1, DIM2), z(DIM1) + + !$omp target update to (iterator(i=1:DIM1, j=1:DIM2): x(i, j)%ptr(:), y(i, j)%ptr(:)) + !$omp target update from (iterator(i=1:DIM1): z(i)%ptr(:)) +end program + +! { dg-final { scan-tree-dump-times "if \\(i <= 17\\) goto ; else goto ;" 2 "gimple" } } +! { dg-final { scan-tree-dump-times "if \\(j <= 39\\) goto ; else goto ;" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "to\\(iterator\\(integer\\(kind=4\\) j=1:39:1, integer\\(kind=4\\) i=1:17:1\\):iterator_array=D\.\[0-9\]+" 2 "gimple" } } +! { dg-final { scan-tree-dump-times "from\\(iterator\\(integer\\(kind=4\\) i=1:17:1\\):iterator_array=D\.\[0-9\]+" 1 "gimple" } } diff --git a/libgomp/testsuite/libgomp.fortran/target-update-iterators-1.f90 b/libgomp/testsuite/libgomp.fortran/target-update-iterators-1.f90 new file mode 100644 index 00000000000..e9a13a3c737 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/target-update-iterators-1.f90 @@ -0,0 +1,68 @@ +! { dg-do run } + +! Test target enter data and target update to the target using map +! iterators. + +program test + integer, parameter :: DIM1 = 8 + integer, parameter :: DIM2 = 15 + + type :: array_ptr + integer, pointer :: arr(:) + end type + + type (array_ptr) :: x(DIM1) + integer :: expected, sum, i, j + + expected = mkarray (x) + + !$omp target enter data map(to: x) + !$omp target enter data map(iterator(i=1:DIM1), to: x(i)%arr(:)) + !$omp target map(from: sum) + sum = 0 + do i = 1, DIM1 + do j = 1, DIM2 + sum = sum + x(i)%arr(j) + end do + end do + !$omp end target + + print *, sum, expected + if (sum .ne. expected) stop 1 + + expected = 0 + do i = 1, DIM1 + do j = 1, DIM2 + x(i)%arr(j) = x(i)%arr(j) * i * j + expected = expected + x(i)%arr(j) + end do + end do + + !$omp target update to(iterator(i=1:DIM1): x(i)%arr(:)) + + !$omp target map(from: sum) + sum = 0 + do i = 1, DIM1 + do j = 1, DIM2 + sum = sum + x(i)%arr(j) + end do + end do + !$omp end target + + if (sum .ne. expected) stop 2 +contains + integer function mkarray (x) + type (array_ptr), intent(inout) :: x(DIM1) + integer :: exp = 0 + + do i = 1, DIM1 + allocate (x(i)%arr(DIM2)) + do j = 1, DIM2 + x(i)%arr(j) = i * j + exp = exp + x(i)%arr(j) + end do + end do + + mkarray = exp + end function +end program diff --git a/libgomp/testsuite/libgomp.fortran/target-update-iterators-2.f90 b/libgomp/testsuite/libgomp.fortran/target-update-iterators-2.f90 new file mode 100644 index 00000000000..6c1a8a7b5dd --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/target-update-iterators-2.f90 @@ -0,0 +1,62 @@ +! { dg-do run } + +! Test target enter data and target update from the target using map +! iterators. + +program test + integer, parameter :: DIM1 = 8 + integer, parameter :: DIM2 = 15 + + type :: array_ptr + integer, pointer :: arr(:) + end type + + type (array_ptr) :: x(DIM1) + integer :: sum, expected + + call mkarray (x) + + !$omp target enter data map(to: x(:DIM1)) + !$omp target enter data map(iterator(i=1:DIM1), to: x(i)%arr(:)) + !$omp target map(from: expected) + expected = 0 + do i = 1, DIM1 + do j = 1, DIM2 + x(i)%arr(j) = (i + 1) * (j + 2) + expected = expected + x(i)%arr(j) + end do + end do + !$omp end target + + ! Host copy of x should remain unchanged. + sum = 0 + do i = 1, DIM1 + do j = 1, DIM2 + sum = sum + x(i)%arr(j) + end do + end do + if (sum .ne. 0) stop 1 + + !$omp target update from(iterator(i=1:DIM1): x(i)%arr(:)) + + ! Host copy should now be updated. + sum = 0 + do i = 1, DIM1 + do j = 1, DIM2 + sum = sum + x(i)%arr(j) + end do + end do + + if (sum .ne. expected) stop 2 +contains + subroutine mkarray (x) + type (array_ptr), intent(inout) :: x(DIM1) + + do i = 1, DIM1 + allocate (x(i)%arr(DIM2)) + do j = 1, DIM2 + x(i)%arr(j) = 0 + end do + end do + end subroutine +end program diff --git a/libgomp/testsuite/libgomp.fortran/target-update-iterators-3.f90 b/libgomp/testsuite/libgomp.fortran/target-update-iterators-3.f90 new file mode 100644 index 00000000000..8dbfb45fe8d --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/target-update-iterators-3.f90 @@ -0,0 +1,77 @@ +! { dg-do run } + +! Test target enter data and target update to the target using map +! iterators with a function. + +program test + implicit none + + integer, parameter :: DIM1 = 8 + integer, parameter :: DIM2 = 15 + + type :: array_ptr + integer, pointer :: arr(:) + end type + + type (array_ptr) :: x(DIM1) + integer :: x_new(DIM1, DIM2) + integer :: expected, sum, i, j + + call mkarray (x) + + !$omp target enter data map(to: x(:DIM1)) + !$omp target enter data map(iterator(i=1:DIM1), to: x(i)%arr(:)) + + ! Update x on host. + do i = 1, DIM1 + do j = 1, DIM2 + x_new(i, j) = x(i)%arr(j) + x(i)%arr(j) = (i + 1) * (j + 2); + end do + end do + + ! Update a subset of x on target. + !$omp target update to(iterator(i=1:DIM1/2): x(f (i))%arr(:)) + + !$omp target map(from: sum) + sum = 0 + do i = 1, DIM1 + do j = 1, DIM2 + sum = sum + x(i)%arr(j) + end do + end do + !$omp end target + + ! Calculate expected value on host. + do i = 1, DIM1/2 + do j = 1, DIM2 + x_new(f (i), j) = x(f (i))%arr(j) + end do + end do + + expected = 0 + do i = 1, DIM1 + do j = 1, DIM2 + expected = expected + x_new(i, j) + end do + end do + + if (sum .ne. expected) stop 1 +contains + subroutine mkarray (x) + type (array_ptr), intent(inout) :: x(DIM1) + + do i = 1, DIM1 + allocate (x(i)%arr(DIM2)) + do j = 1, DIM2 + x(i)%arr(j) = i * j + end do + end do + end subroutine + + integer function f (i) + integer, intent(in) :: i + + f = i * 2 + end function +end program