From patchwork Fri Oct 4 14:54: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: 1992778 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=u/PqYHZc; 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 4XKs6t0Dk6z1xt7 for ; Sat, 5 Oct 2024 00:55:50 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id D87CE3844054 for ; Fri, 4 Oct 2024 14:55:47 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from mail-ed1-x531.google.com (mail-ed1-x531.google.com [IPv6:2a00:1450:4864:20::531]) by sourceware.org (Postfix) with ESMTPS id 4B12C385E011 for ; Fri, 4 Oct 2024 14:55:22 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 4B12C385E011 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 4B12C385E011 Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=2a00:1450:4864:20::531 ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1728053726; cv=none; b=tZ++8iACVXN9iB845bzhTkFCBeudaXozp70TCLAFgbQIz9rE/56HSzbzAp2te8eMojzUukDOI8eOal7koG17ELVmQpYP27s6NrGj8ezgB+EB9H21Zl8mYCcWpo92n8t2zLQoIOHmJ5MW0X9Nv+g/4cS/VIDHoyDanWauj+Nares= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1728053726; c=relaxed/simple; bh=Bvxq21jPdbXh4ZcqP2aXfYUMkOqElQJ72lDOBmiHa5U=; h=DKIM-Signature:Message-ID:Date:MIME-Version:To:Subject:From; b=VQ6BqeuppHWGqXKa4Ahrych6Ow0GpZfgf3iaHIo9atE8iXwi0ISuC28TspMK4KtgZaXEsGRf/L3+1paRw0Q25GaNo1wWyKWrFqt2x2Q3xpob8skO2NWEL8CpXLlmMymqVzK1hqTnVQJUi44xPnHIkPFnxgRETI5wWUnd3MAKmzg= ARC-Authentication-Results: i=1; server2.sourceware.org Received: by mail-ed1-x531.google.com with SMTP id 4fb4d7f45d1cf-5c88c9e45c2so5825577a12.0 for ; Fri, 04 Oct 2024 07:55:22 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre-com.20230601.gappssmtp.com; s=20230601; t=1728053720; x=1728658520; 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=Bvxq21jPdbXh4ZcqP2aXfYUMkOqElQJ72lDOBmiHa5U=; b=u/PqYHZcyGn78AJwxCPL/5Wo6mbtPKYOEMpE7lQn7DU1gLAh2unvE+WB0addKYPyOs mglQEC19xpCFKCD4CrgB7HcEfd2/jyt/TVOKt5DUh/DJCOsa2SLd6AdYfIiTZ4qewj+g 6Xe6JSKggUbfP0iSzcBS5/gWJntSdVj9MuaaC7btyBF8CTKPI6KaVcvPavL0eId4pn90 9AgFK9+3NisLCILMIdLsyoN/5c3uS5IfbBYttnlrCRXK2mrZC1oICFo6ZIAB0boXMicn r+1YWmqzxrCKR0r4wmZMhIbpU+GCI3dXvpHYj+zIBuXb1dl1sR30UYT4IQ0CVIV1i100 7ThQ== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1728053720; x=1728658520; 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=Bvxq21jPdbXh4ZcqP2aXfYUMkOqElQJ72lDOBmiHa5U=; b=DK/hHzHQWzld83TiUOOVbE5qBEhuF3aFibJqGZy2SN+DK888G0QgNrQdDDZURA18PL OiQK1xLcElqrixUmB41+KSeR3Sr9vld/Nh/ktNqjvZIbzA/dSK/jNepio3WjDI3AyZ/+ /waeEKpOkuH9+rY4GzOdFBE7LE5EjQAVTaB12wpRpg5/Z6U3o/yxpjQnBJOIAMSX09v2 SDVoaZSZZJbfEpCX8VaTRbll7nKqGGG1jciuQws901gEKwzXivJVw6+in1yIMDvvwbO5 7l9afLCyZosGTsmuHSuFwaWmMOY6qp3qiHn3qPX33B0QvwkFXyx6cSBfD1qRvLNV9gJx o9Mg== X-Gm-Message-State: AOJu0YzcTTMVxW/WS1SLRj01BNaBvlPp69Wg7rDH5BL3YHV7hzJ75bFr ChqINz9j01x0f9+x/OaQ+4YjuhHVXuLXD7x5Am7x23RsUgoaheydPVqU9uhw9k5sayEkZyV/UuE B X-Google-Smtp-Source: AGHT+IECk1/jAo/n/CkODQl4MpK26YbZ44VJD5pY+kAkBQDK7lbgFsMyAPQ6H0SeY+ENPgxzp2k/+A== X-Received: by 2002:a17:907:9303:b0:a91:1699:f8eb with SMTP id a640c23a62f3a-a991d003859mr276088766b.28.1728053720322; Fri, 04 Oct 2024 07:55:20 -0700 (PDT) Received: from ?IPV6:2a00:23c6:88fe:9301:2d7d:f734:bc6:c47b? ([2a00:23c6:88fe:9301:2d7d:f734:bc6:c47b]) by smtp.gmail.com with ESMTPSA id a640c23a62f3a-a992e61e716sm1633566b.67.2024.10.04.07.55.19 (version=TLS1_3 cipher=TLS_AES_128_GCM_SHA256 bits=128/128); Fri, 04 Oct 2024 07:55:20 -0700 (PDT) Message-ID: <6970b7ca-0c18-40d7-8662-a013a82eaf21@baylibre.com> Date: Fri, 4 Oct 2024 15:54:53 +0100 MIME-Version: 1.0 User-Agent: Mozilla Thunderbird To: gcc-patches , Jakub Jelinek , Tobias Burnus References: <6b94b8ed-020b-47e2-b02a-4891891f2847@baylibre.com> Subject: [PATCH v3 1/5] openmp: Refactor handling of iterators Content-Language: en-GB From: Kwok Cheung Yeung In-Reply-To: <6b94b8ed-020b-47e2-b02a-4891891f2847@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 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 34bf780b1e0395028ecdacfa1385238a8da13be6 Mon Sep 17 00:00:00 2001 From: Kwok Cheung Yeung Date: Fri, 4 Oct 2024 15:15:42 +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-10-04 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 620a3c1353a..24c8a801255 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 ba6d96d26b2..30a03f071d8 100644 --- a/gcc/c/c-typeck.cc +++ b/gcc/c/c-typeck.cc @@ -14504,9 +14504,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, @@ -15697,9 +15695,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 @@ -15799,10 +15795,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 43468e5f62e..5a72402ba1f 100644 --- a/gcc/cp/pt.cc +++ b/gcc/cp/pt.cc @@ -17604,9 +17604,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 0cb46c1986c..4f856a9d749 100644 --- a/gcc/cp/semantics.cc +++ b/gcc/cp/semantics.cc @@ -6015,9 +6015,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, @@ -8322,9 +8320,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 ceb53e5d5bb..b5b1f83db8f 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 b378ffbfb4c..39e586c808c 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 Fri Oct 4 14:56:01 2024 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Kwok Cheung Yeung X-Patchwork-Id: 1992779 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@legolas.ozlabs.org Authentication-Results: legolas.ozlabs.org; dkim=pass (2048-bit key; unprotected) header.d=baylibre-com.20230601.gappssmtp.com header.i=@baylibre-com.20230601.gappssmtp.com header.a=rsa-sha256 header.s=20230601 header.b=Un0fCwJ5; dkim-atps=neutral Authentication-Results: legolas.ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=8.43.85.97; helo=server2.sourceware.org; envelope-from=gcc-patches-bounces~incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=patchwork.ozlabs.org) Received: from server2.sourceware.org (server2.sourceware.org [8.43.85.97]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature ECDSA (secp384r1) server-digest SHA384) (No client certificate requested) by legolas.ozlabs.org (Postfix) with ESMTPS id 4XKs8Q2vDkz1xt7 for ; Sat, 5 Oct 2024 00:57:10 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 971E538432ED for ; Fri, 4 Oct 2024 14:57:08 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from mail-ed1-x536.google.com (mail-ed1-x536.google.com [IPv6:2a00:1450:4864:20::536]) by sourceware.org (Postfix) with ESMTPS id 1BBBD385E011 for ; Fri, 4 Oct 2024 14:56:30 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 1BBBD385E011 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=baylibre.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=baylibre.com ARC-Filter: OpenARC Filter v1.0.0 sourceware.org 1BBBD385E011 Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=2a00:1450:4864:20::536 ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1728053796; cv=none; b=EoHVJmyF6ovJ+tRB80uH+6vbc9N98mIcVsCDuM6VAlh0OahpP1Ts89TGqtOM9ifVCbwXFpnGghuKsfDmKskauw4cscwSzY9xeTjpnktkmnxi2S5n8U7UtNmxsCuR99O1n1UMTD5EDosGMhAGojSY04qIxeCLwlblobLNFsZN4NU= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1728053796; c=relaxed/simple; bh=2hz5f0lZQdtjVkAy0812rYNF0zGjcmXmmBn9/RNmD3g=; h=DKIM-Signature:Message-ID:Date:MIME-Version:To:Subject:From; b=ckk5y0WUi2+eqifliWa/8CKtLNcwAdfl/kaY0FR4OqP80gWP9CaNVnrum2Heq4OndhLY+nKFeCUZADe9SAdhsnenD9WN4KXaZtqvAyXpmkmPxpyFIAnS3+2IIB5tn70L3pRP3jzszcYhzaQkb+gYEMQYhoLajCxHMaaEH6KpRDo= ARC-Authentication-Results: i=1; server2.sourceware.org Received: by mail-ed1-x536.google.com with SMTP id 4fb4d7f45d1cf-5c5b9d2195eso2993155a12.1 for ; Fri, 04 Oct 2024 07:56:30 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre-com.20230601.gappssmtp.com; s=20230601; t=1728053788; x=1728658588; darn=gcc.gnu.org; h=in-reply-to:from:content-language:subject:references:to:user-agent :mime-version:date:message-id:from:to:cc:subject:date:message-id :reply-to; bh=UTaJgqmXWq6YUKhFg0fhJKu6JPci6SR/kHhYDe64zF0=; b=Un0fCwJ582R6e99/KbbYAiuqbEqwZedNE41AGQQS8ALGTigVVeFTHEQaBlOwz2rVI0 tS5RBy95SOCAc5dHbaGrf+PNeCutnJCVdE485dbUgM1/OGVXnOMUil7AgKnbNMJqCYjN mLFf0NwAwRQPqSu6PM8zY5LuhvtJlKO+jOD6O6gKwFVzYburiYQkyg65dm6vSDTOrqk0 DH1AUavAwn0fh1I4OzZu6ZcBzk/9erBeKPN8O2aaOsjkwSmkCumk38E+lVSCvqSrXjRw slOMtceJFhtF/KSg4MTk939v550xkkLDq6I77EAZWF94nE08yjOXhPqRAWa8JNtEl3BA i/5Q== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1728053788; x=1728658588; h=in-reply-to:from:content-language:subject:references:to:user-agent :mime-version:date:message-id:x-gm-message-state:from:to:cc:subject :date:message-id:reply-to; bh=UTaJgqmXWq6YUKhFg0fhJKu6JPci6SR/kHhYDe64zF0=; b=UB2U5iQkyFtxPXvavNnv4tv5+E2JinPyQ69UuUU++N2zmGm4RVansxZJpWjx4VL1PI tHc+40AnJkGNfZ5gONHeHaTB5YCv6OF8nDUFL0k5JgjktQQAXG78d7ED1rD+9TF2h+VN en5Cq5Mck2PW1BtV2YpRRIeMVGu1Uyk7bp+/hZRjL8wi4X9varRp7ZpSgqrxlrDZ0uAT FQkAfsxhXxGbyFZ790BewJ5CQA5k7Lfe/vbQfbavwyH4TJOPlOk3dXxuM/nkQ0h4PGnZ bG97bNpCOZOAF9pOj5g5cVWnF7XwK+lp8z98aoS4chPBGL/BTqsxVl4O9HiF0SG7qvwD 2ppg== X-Gm-Message-State: AOJu0YwZp8tf1ud+DHDwaUr0Ortv/jP2ZQNF2BvFaRdwHDDqVYljFqOi Wd3w6TmNjmWYqD8yAV1p3vM1N/xBwEXAgRXiD9Pa2dibR8k/CKEKD3Z618oYg/dCRUrc9kBlTp2 5 X-Google-Smtp-Source: AGHT+IHWPo8jowzuWfpUoR58x7yPXZhU5PWbCqrpgxaDplLvQuNlSMDVoE2EAQzKXOBCU/1fA6Kl4A== X-Received: by 2002:a17:907:2cc7:b0:a90:13b6:3ece with SMTP id a640c23a62f3a-a991bd3fb6emr291721266b.15.1728053788341; Fri, 04 Oct 2024 07:56:28 -0700 (PDT) Received: from ?IPV6:2a00:23c6:88fe:9301:2d7d:f734:bc6:c47b? ([2a00:23c6:88fe:9301:2d7d:f734:bc6:c47b]) by smtp.gmail.com with ESMTPSA id a640c23a62f3a-a992e784a4bsm1055166b.114.2024.10.04.07.56.27 (version=TLS1_3 cipher=TLS_AES_128_GCM_SHA256 bits=128/128); Fri, 04 Oct 2024 07:56:28 -0700 (PDT) Message-ID: Date: Fri, 4 Oct 2024 15:56:01 +0100 MIME-Version: 1.0 User-Agent: Mozilla Thunderbird To: gcc-patches , Jakub Jelinek , Tobias Burnus References: <6b94b8ed-020b-47e2-b02a-4891891f2847@baylibre.com> Subject: [PATCH v3 2/5] openmp: Add support for iterators in map clauses (C/C++) Content-Language: en-GB From: Kwok Cheung Yeung In-Reply-To: <6b94b8ed-020b-47e2-b02a-4891891f2847@baylibre.com> X-Spam-Status: No, score=-12.5 required=5.0 tests=BAYES_00, DKIM_SIGNED, DKIM_VALID, GIT_PATCH_0, RCVD_IN_DNSWL_NONE, SPF_HELO_NONE, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.30 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces~incoming=patchwork.ozlabs.org@gcc.gnu.org This patch modifies the C and C++ parsers to accept an iterator as a map type modifier, storing it in the OMP_CLAUSE_ITERATOR argument of the clause. When finishing clauses, any clauses generated from a clause with iterators also has the iterator applied to them. During gimplification, check_omp_map_iterators is called to check that all iterator variables are referenced at some point with a clause. Gimplification of the clause decl and size are delayed until iterator expansion as they may reference iterator variables. In lower_target, lower_omp_map_iterators is called to construct the expansion loop for iterator clauses. Clauses using the same set of iterators reuse the loop, though with different storage allocated for them. lower_omp_map_iterator_expr is called to add the final expression that is sent as the hostaddr for libgomp to the loop, and a reference to the array generated by the iterator loop is returned to replace the original expression. lower_omp_map_iterator_size works similarly for the clause size. finish_omp_map_iterators is called later to finalise the loop. Libgomp has a new function gomp_merge_iterator_maps which identifies data coming from an iterator, and effectively creates new maps on-the-fly from the iterator info array, inserting them into the list of mappings at the point where iterator data occurred. As there are now multiple maps where one was previously, an entry is only added to the target vars for the first expanded map, otherwise it will get out of sync with the expected layout and the wrong variables will be picked up by the target function. From 50557e513ca534ba32f50d99991b056a07a6f671 Mon Sep 17 00:00:00 2001 From: Kwok Cheung Yeung Date: Fri, 4 Oct 2024 15:16:12 +0100 Subject: [PATCH 2/5] openmp: Add support for iterators in map clauses (C/C++) This adds preliminary support for iterators in map clauses within OpenMP 'target' constructs (which includes constructs such as 'target enter data'). Iterators with non-constant loop bounds are not currently supported. 2024-10-04 Kwok Cheung Yeung gcc/c/ * c-parser.cc (c_parser_omp_clause_map): Parse 'iterator' modifier. * c-typeck.cc (c_finish_omp_clauses): Finish iterators. Apply iterators to generated clauses. gcc/cp/ * parser.cc (cp_parser_omp_clause_map): Parse 'iterator' modifier. * semantics.cc (finish_omp_clauses): Finish iterators. Apply iterators to generated clauses. gcc/ * gimplify.cc (compute_iterator_count): Make non-static. Take an iterator instead of a clause for an operand. (build_iterator_loop): Likewise. (gimplify_omp_depend): Pass iterator in call to compute_iterator_count and build_iterator_loop. (find_var_decl): New. (check_omp_map_iterators): New. (gimplify_scan_omp_clauses): Call check_omp_map_iterators on clauses with iterators. (gimplify_adjust_omp_clauses): Skip gimplification of clause decl and size for clauses with iterators. * omp-low.cc (struct iterator_loop_info_t): New type. (iterator_loop_map_t): New type. (lower_omp_map_iterators): New. (lower_omp_map_iterator_expr): New. (lower_omp_map_iterator_size): New. (finish_omp_map_iterators): New. (lower_omp_target): Call lower_omp_map_iterators on clauses with iterators. Call lower_omp_map_iterator_expr before assigning to sender ref. Call lower_omp_map_iterator_size before setting the size. Call finish_omp_map_iterators. Insert statements generated during iterator expansion before the statements for the target clause. * tree-pretty-print.cc (dump_omp_clause): Call dump_omp_iterators for iterators in map clauses. * tree.cc (omp_clause_num_ops): Add operand for OMP_CLAUSE_MAP. (walk_tree_1): Do not walk last operand of OMP_CLAUSE_MAP. * tree.h (OMP_CLAUSE_HAS_ITERATORS): New. (OMP_CLAUSE_ITERATORS: New. gcc/testsuite/ * c-c++-common/gomp/map-6.c (foo): Amend expected error message. * c-c++-common/gomp/target-map-iterators-1.c: New. * c-c++-common/gomp/target-map-iterators-2.c: New. * c-c++-common/gomp/target-map-iterators-3.c: New. libgomp/ * target.c (kind_to_name): New. (gomp_merge_iterator_maps): New. (gomp_map_vars_internal): Call gomp_merge_iterator_maps. Copy address of only the first iteration to target vars. Free allocated variables. * testsuite/libgomp.c-c++-common/target-map-iterators-1.c: New. * testsuite/libgomp.c-c++-common/target-map-iterators-2.c: New. * testsuite/libgomp.c-c++-common/target-map-iterators-3.c: New. --- gcc/c/c-parser.cc | 59 +++++- gcc/c/c-typeck.cc | 22 ++- gcc/cp/parser.cc | 62 +++++- gcc/cp/semantics.cc | 22 ++- gcc/gimplify.cc | 88 +++++++-- gcc/omp-low.cc | 186 +++++++++++++++++- gcc/testsuite/c-c++-common/gomp/map-6.c | 10 +- .../gomp/target-map-iterators-1.c | 23 +++ .../gomp/target-map-iterators-2.c | 19 ++ .../gomp/target-map-iterators-3.c | 23 +++ gcc/tree-pretty-print.cc | 5 + gcc/tree.cc | 5 +- gcc/tree.h | 7 + libgomp/target.c | 130 +++++++++++- .../target-map-iterators-1.c | 47 +++++ .../target-map-iterators-2.c | 44 +++++ .../target-map-iterators-3.c | 56 ++++++ 17 files changed, 759 insertions(+), 49 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/gomp/target-map-iterators-1.c create mode 100644 gcc/testsuite/c-c++-common/gomp/target-map-iterators-2.c create mode 100644 gcc/testsuite/c-c++-common/gomp/target-map-iterators-3.c create mode 100644 libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-1.c create mode 100644 libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-2.c create mode 100644 libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-3.c diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc index a681438cbbe..184fc076388 100644 --- a/gcc/c/c-parser.cc +++ b/gcc/c/c-parser.cc @@ -18892,7 +18892,7 @@ c_parser_omp_clause_doacross (c_parser *parser, tree list) map ( [map-type-modifier[,] ...] map-kind: variable-list ) map-type-modifier: - always | close */ + always | close | present | iterator (iterators-definition) */ static tree c_parser_omp_clause_map (c_parser *parser, tree list) @@ -18907,15 +18907,35 @@ c_parser_omp_clause_map (c_parser *parser, tree list) int pos = 1; int map_kind_pos = 0; - while (c_parser_peek_nth_token_raw (parser, pos)->type == CPP_NAME) + int iterator_length = 0; + for (;;) { - if (c_parser_peek_nth_token_raw (parser, pos + 1)->type == CPP_COLON) + c_token *tok = c_parser_peek_nth_token_raw (parser, pos); + if (tok->type != CPP_NAME) + break; + + const char *p = IDENTIFIER_POINTER (tok->value); + c_token *next_tok = c_parser_peek_nth_token_raw (parser, pos + 1); + if (strcmp (p, "iterator") == 0 && next_tok->type == CPP_OPEN_PAREN) + { + unsigned n = pos + 2; + if (c_parser_check_balanced_raw_token_sequence (parser, &n) + && c_parser_peek_nth_token_raw (parser, n)->type + == CPP_CLOSE_PAREN) + { + iterator_length = n - pos + 1; + pos = n; + next_tok = c_parser_peek_nth_token_raw (parser, pos + 1); + } + } + + if (next_tok->type == CPP_COLON) { map_kind_pos = pos; break; } - if (c_parser_peek_nth_token_raw (parser, pos + 1)->type == CPP_COMMA) + if (next_tok->type == CPP_COMMA) pos++; pos++; } @@ -18923,6 +18943,7 @@ c_parser_omp_clause_map (c_parser *parser, tree list) int always_modifier = 0; int close_modifier = 0; int present_modifier = 0; + tree iterators = NULL_TREE; for (int pos = 1; pos < map_kind_pos; ++pos) { c_token *tok = c_parser_peek_token (parser); @@ -18964,10 +18985,24 @@ c_parser_omp_clause_map (c_parser *parser, tree list) } present_modifier++; } + else if (strcmp ("iterator", p) == 0 + && c_parser_peek_2nd_token (parser)->type == CPP_OPEN_PAREN) + { + if (iterators) + { + c_parser_error (parser, "too many % modifiers"); + parens.skip_until_found_close (parser); + return list; + } + iterators = c_parser_omp_iterators (parser); + pos += iterator_length - 1; + continue; + } else { c_parser_error (parser, "% clause with map-type modifier other " - "than %, % or %"); + "than %, %, % " + "or %"); parens.skip_until_found_close (parser); return list; } @@ -19016,8 +19051,20 @@ c_parser_omp_clause_map (c_parser *parser, tree list) nl = c_parser_omp_variable_list (parser, clause_loc, OMP_CLAUSE_MAP, list, true); + if (iterators) + { + tree block = pop_scope (); + if (iterators == error_mark_node) + iterators = NULL_TREE; + else + TREE_VEC_ELT (iterators, 5) = block; + } + for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c)) - OMP_CLAUSE_SET_MAP_KIND (c, kind); + { + OMP_CLAUSE_SET_MAP_KIND (c, kind); + OMP_CLAUSE_ITERATORS (c) = iterators; + } parens.skip_until_found_close (parser); return nl; diff --git a/gcc/c/c-typeck.cc b/gcc/c/c-typeck.cc index 30a03f071d8..cca9f1c000c 100644 --- a/gcc/c/c-typeck.cc +++ b/gcc/c/c-typeck.cc @@ -15058,7 +15058,15 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) /* We've reached the end of a list of expanded nodes. Reset the group start pointer. */ if (c == grp_sentinel) - grp_start_p = NULL; + { + if (grp_start_p + && OMP_CLAUSE_HAS_ITERATORS (*grp_start_p) + && OMP_CLAUSE_ITERATORS (*grp_start_p)) + for (tree gc = *grp_start_p; gc != grp_sentinel; + gc = OMP_CLAUSE_CHAIN (gc)) + OMP_CLAUSE_ITERATORS (gc) = OMP_CLAUSE_ITERATORS (*grp_start_p); + grp_start_p = NULL; + } switch (OMP_CLAUSE_CODE (c)) { @@ -15805,6 +15813,12 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) case OMP_CLAUSE_MAP: if (OMP_CLAUSE_MAP_IMPLICIT (c) && !implicit_moved) goto move_implicit; + if (OMP_CLAUSE_ITERATORS (c) + && c_omp_finish_iterators (OMP_CLAUSE_ITERATORS (c))) + { + t = error_mark_node; + break; + } /* FALLTHRU */ case OMP_CLAUSE_TO: case OMP_CLAUSE_FROM: @@ -16497,6 +16511,12 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) pc = &OMP_CLAUSE_CHAIN (c); } + if (grp_start_p + && OMP_CLAUSE_HAS_ITERATORS (*grp_start_p) + && OMP_CLAUSE_ITERATORS (*grp_start_p)) + for (tree gc = *grp_start_p; gc; gc = OMP_CLAUSE_CHAIN (gc)) + OMP_CLAUSE_ITERATORS (gc) = OMP_CLAUSE_ITERATORS (*grp_start_p); + if (simdlen && safelen && tree_int_cst_lt (OMP_CLAUSE_SAFELEN_EXPR (safelen), diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc index f50534f5f39..79d6e115d16 100644 --- a/gcc/cp/parser.cc +++ b/gcc/cp/parser.cc @@ -41776,16 +41776,34 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list) int pos = 1; int map_kind_pos = 0; - while (cp_lexer_peek_nth_token (parser->lexer, pos)->type == CPP_NAME - || cp_lexer_peek_nth_token (parser->lexer, pos)->keyword == RID_DELETE) + int iterator_length = 0; + for (;;) { - if (cp_lexer_peek_nth_token (parser->lexer, pos + 1)->type == CPP_COLON) + cp_token *tok = cp_lexer_peek_nth_token (parser->lexer, pos); + if (!(tok->type == CPP_NAME || tok->keyword == RID_DELETE)) + break; + + cp_token *next_tok = cp_lexer_peek_nth_token (parser->lexer, pos + 1); + if (tok->type == CPP_NAME + && strcmp (IDENTIFIER_POINTER (tok->u.value), "iterator") == 0 + && next_tok->type == CPP_OPEN_PAREN) + { + int n = cp_parser_skip_balanced_tokens (parser, pos + 1); + if (n != pos + 1) + { + iterator_length = n - pos; + pos = n - 1; + next_tok = cp_lexer_peek_nth_token (parser->lexer, n); + } + } + + if (next_tok->type == CPP_COLON) { map_kind_pos = pos; break; } - if (cp_lexer_peek_nth_token (parser->lexer, pos + 1)->type == CPP_COMMA) + if (next_tok->type == CPP_COMMA) pos++; pos++; } @@ -41793,6 +41811,7 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list) bool always_modifier = false; bool close_modifier = false; bool present_modifier = false; + tree iterators = NULL_TREE; for (int pos = 1; pos < map_kind_pos; ++pos) { cp_token *tok = cp_lexer_peek_token (parser->lexer); @@ -41842,10 +41861,29 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list) } present_modifier = true; } + else if (strcmp ("iterator", p) == 0 + && cp_lexer_peek_nth_token (parser->lexer, 2)->type + == CPP_OPEN_PAREN) + { + if (iterators) + { + cp_parser_error (parser, "too many % modifiers"); + cp_parser_skip_to_closing_parenthesis (parser, + /*recovering=*/true, + /*or_comma=*/false, + /*consume_paren=*/true); + return list; + } + begin_scope (sk_omp, NULL); + iterators = cp_parser_omp_iterators (parser); + pos += iterator_length - 1; + continue; + } else { cp_parser_error (parser, "% clause with map-type modifier other" - " than %, % or %"); + " than %, %, %" + " or %"); cp_parser_skip_to_closing_parenthesis (parser, /*recovering=*/true, /*or_comma=*/false, @@ -41909,8 +41947,20 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list) NULL, true); finish_scope (); + if (iterators) + { + tree block = poplevel (1, 1, 0); + if (iterators == error_mark_node) + iterators = NULL_TREE; + else + TREE_VEC_ELT (iterators, 5) = block; + } + for (c = nlist; c != list; c = OMP_CLAUSE_CHAIN (c)) - OMP_CLAUSE_SET_MAP_KIND (c, kind); + { + OMP_CLAUSE_SET_MAP_KIND (c, kind); + OMP_CLAUSE_ITERATORS (c) = iterators; + } return nlist; } diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc index 4f856a9d749..ba5657f7bc2 100644 --- a/gcc/cp/semantics.cc +++ b/gcc/cp/semantics.cc @@ -7263,7 +7263,15 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) /* We've reached the end of a list of expanded nodes. Reset the group start pointer. */ if (c == grp_sentinel) - grp_start_p = NULL; + { + if (grp_start_p + && OMP_CLAUSE_HAS_ITERATORS (*grp_start_p) + && OMP_CLAUSE_ITERATORS (*grp_start_p)) + for (tree gc = *grp_start_p; gc != grp_sentinel; + gc = OMP_CLAUSE_CHAIN (gc)) + OMP_CLAUSE_ITERATORS (gc) = OMP_CLAUSE_ITERATORS (*grp_start_p); + grp_start_p = NULL; + } switch (OMP_CLAUSE_CODE (c)) { @@ -8484,6 +8492,12 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) case OMP_CLAUSE_MAP: if (OMP_CLAUSE_MAP_IMPLICIT (c) && !implicit_moved) goto move_implicit; + if (OMP_CLAUSE_ITERATORS (c) + && cp_omp_finish_iterators (OMP_CLAUSE_ITERATORS (c))) + { + t = error_mark_node; + break; + } /* FALLTHRU */ case OMP_CLAUSE_TO: case OMP_CLAUSE_FROM: @@ -9348,6 +9362,12 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) pc = &OMP_CLAUSE_CHAIN (c); } + if (grp_start_p + && OMP_CLAUSE_HAS_ITERATORS (*grp_start_p) + && OMP_CLAUSE_ITERATORS (*grp_start_p)) + for (tree gc = *grp_start_p; gc; gc = OMP_CLAUSE_CHAIN (gc)) + OMP_CLAUSE_ITERATORS (gc) = OMP_CLAUSE_ITERATORS (*grp_start_p); + if (reduction_seen < 0 && (ordered_seen || schedule_seen)) reduction_seen = -2; diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index b5b1f83db8f..6e532d07fcf 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -8823,13 +8823,13 @@ gimplify_omp_affinity (tree *list_p, gimple_seq *pre_p) } /* Returns a tree expression containing the total iteration count of the - iterator clause decl T. */ + iterator IT. */ -static tree -compute_iterator_count (tree t, gimple_seq *pre_p) +tree +compute_iterator_count (tree it, gimple_seq *pre_p) { tree tcnt = size_one_node; - for (tree it = TREE_PURPOSE (t); it; it = TREE_CHAIN (it)) + for (; it; it = TREE_CHAIN (it)) { if (gimplify_expr (&TREE_VEC_ELT (it, 1), pre_p, NULL, is_gimple_val, fb_rvalue) == GS_ERROR @@ -8899,21 +8899,17 @@ compute_iterator_count (tree t, gimple_seq *pre_p) Returns a pointer to the BIND_EXPR_BODY in the innermost loop body. LAST_BIND is set to point to the BIND_EXPR containing the whole loop. */ -static tree * -build_iterator_loop (tree c, gimple_seq *pre_p, tree *last_bind) +tree * +build_iterator_loop (tree it, gimple_seq *pre_p, tree *last_bind) { - tree t = OMP_CLAUSE_DECL (c); - gcc_assert (OMP_ITERATOR_DECL_P (t)); - if (*last_bind) gimplify_and_add (*last_bind, pre_p); - tree block = TREE_VEC_ELT (TREE_PURPOSE (t), 5); + tree block = TREE_VEC_ELT (it, 5); *last_bind = build3 (BIND_EXPR, void_type_node, BLOCK_VARS (block), NULL, block); TREE_SIDE_EFFECTS (*last_bind) = 1; - SET_EXPR_LOCATION (*last_bind, OMP_CLAUSE_LOCATION (c)); tree *p = &BIND_EXPR_BODY (*last_bind); - for (tree it = TREE_PURPOSE (t); it; it = TREE_CHAIN (it)) + for (; it; it = TREE_CHAIN (it)) { tree var = TREE_VEC_ELT (it, 0); tree begin = TREE_VEC_ELT (it, 1); @@ -9023,7 +9019,7 @@ gimplify_omp_depend (tree *list_p, gimple_seq *pre_p) { if (TREE_PURPOSE (t) != last_iter) { - tree tcnt = compute_iterator_count (t, pre_p); + tree tcnt = compute_iterator_count (TREE_PURPOSE (t), pre_p); if (!tcnt) return 2; last_iter = TREE_PURPOSE (t); @@ -9181,7 +9177,9 @@ gimplify_omp_depend (tree *list_p, gimple_seq *pre_p) if (OMP_ITERATOR_DECL_P (t)) { if (TREE_PURPOSE (t) != last_iter) - last_body = build_iterator_loop (c, pre_p, &last_bind); + last_body = build_iterator_loop (TREE_PURPOSE (t), pre_p, + &last_bind); + SET_EXPR_LOCATION (last_bind, OMP_CLAUSE_LOCATION (c)); last_iter = TREE_PURPOSE (t); if (TREE_CODE (TREE_VALUE (t)) == COMPOUND_EXPR) { @@ -12078,6 +12076,51 @@ error_out: return success; } +/* Callback for walk_tree to find a VAR_DECL (stored in DATA) in the + tree TP. */ + +static tree +find_var_decl (tree *tp, int *, void *data) +{ + tree t = *tp; + + if (TREE_CODE (t) == VAR_DECL && t == (tree) data) + return t; + + return NULL_TREE; +} + +/* Check that the clause C uses all the iterator variables. + Return TRUE if there are no errors. */ + +static bool +check_omp_map_iterators (tree c) +{ + bool error = false; + gcc_assert (OMP_CLAUSE_ITERATORS (c)); + + /* Do not check internal map kinds. */ + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH) + return true; + + for (tree it = OMP_CLAUSE_ITERATORS (c); it; it = TREE_CHAIN (it)) + { + tree var = TREE_VEC_ELT (it, 0); + tree t = walk_tree (&OMP_CLAUSE_DECL (c), find_var_decl, var, NULL); + if (t == NULL_TREE) + t = walk_tree (&OMP_CLAUSE_SIZE (c), find_var_decl, var, NULL); + if (t == NULL_TREE) + { + error_at (OMP_CLAUSE_LOCATION (c), + "iterator variable %qD not used in clause expression", + var); + error = true; + } + } + return !error; +} + /* Scan the OMP clauses in *LIST_P, installing mappings into a new and previous omp contexts. */ @@ -12478,6 +12521,12 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, break; } + if (OMP_CLAUSE_ITERATORS (c) && !check_omp_map_iterators (c)) + { + remove = true; + break; + } + if (!omp_parse_expr (addr_tokens, decl)) { remove = true; @@ -14168,7 +14217,11 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, : TYPE_SIZE_UNIT (TREE_TYPE (decl)); } gimplify_omp_ctxp = ctx->outer_context; - if (gimplify_expr (&OMP_CLAUSE_SIZE (c), pre_p, NULL, + if (OMP_CLAUSE_ITERATORS (c)) + /* Gimplify the OMP_CLAUSE_SIZE later, when the iterator is + gimplified. */ + ; + else if (gimplify_expr (&OMP_CLAUSE_SIZE (c), pre_p, NULL, is_gimple_val, fb_rvalue) == GS_ERROR) { gimplify_omp_ctxp = ctx; @@ -14333,6 +14386,11 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, if (code == OMP_TARGET && OMP_CLAUSE_MAP_IN_REDUCTION (c)) break; + /* Do not gimplify the declaration yet for clauses with + iterators. */ + if (OMP_CLAUSE_ITERATORS (c)) + break; + gimplify_omp_ctxp = ctx->outer_context; if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue, fb_lvalue) == GS_ERROR) diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc index da2051b0279..9cf6e207d1c 100644 --- a/gcc/omp-low.cc +++ b/gcc/omp-low.cc @@ -12607,6 +12607,163 @@ lower_omp_taskreg (gimple_stmt_iterator *gsi_p, omp_context *ctx) } } +extern tree compute_iterator_count (tree it, gimple_seq *pre_p); +extern tree *build_iterator_loop (tree it, gimple_seq *pre_p, tree *last_bind); + +struct iterator_loop_info_t +{ + tree bind; + tree count; + tree index; + tree *body; + tree *iterator; + hash_map elems; +}; + +typedef hash_map iterator_loop_map_t; + +/* Builds a loop to expand any iterators in clause C, reusing any previously + built loops if they use the same set of iterators. Generated Gimple + statements are placed into PRE_P. Information on the loops is held in + LOOPS. finish_omp_map_iterators must be called before the loops are + used. */ + +static void +lower_omp_map_iterators (tree c, gimple_seq *pre_p, iterator_loop_map_t *loops) +{ + if (!OMP_CLAUSE_HAS_ITERATORS (c) || !OMP_CLAUSE_ITERATORS (c)) + return; + + bool built_p; + iterator_loop_info_t &loop = loops->get_or_insert (OMP_CLAUSE_ITERATORS (c), + &built_p); + if (!built_p) + { + loop.count = compute_iterator_count (OMP_CLAUSE_ITERATORS (c), pre_p); + if (!loop.count) + return; + + loop.body = build_iterator_loop (OMP_CLAUSE_ITERATORS (c), pre_p, + &loop.bind); + loop.index = create_tmp_var (sizetype); + SET_EXPR_LOCATION (loop.bind, OMP_CLAUSE_LOCATION (c)); + loop.iterator = &OMP_CLAUSE_ITERATORS (c); + + /* idx = -1; */ + /* This should be initialized to before the individual elements, + as idx is pre-incremented in the loop body. */ + gimple *g = gimple_build_assign (loop.index, size_int (-1)); + gimple_seq_add_stmt (pre_p, g); + + /* IN LOOP BODY: */ + /* idx += 2; */ + tree tem = build2_loc (OMP_CLAUSE_LOCATION (c), MODIFY_EXPR, + void_type_node, loop.index, + size_binop (PLUS_EXPR, loop.index, size_int (2))); + append_to_statement_list_force (tem, loop.body); + } + + /* Create array to hold expanded values. */ + tree last_count_2 = size_binop (MULT_EXPR, loop.count, size_int (2)); + tree arr_length = size_binop (PLUS_EXPR, last_count_2, size_int (1)); + tree elems = NULL_TREE; + if (TREE_CONSTANT (arr_length)) + { + tree type = build_array_type (ptr_type_node, + build_index_type (arr_length)); + elems = create_tmp_var_raw (type); + TREE_ADDRESSABLE (elems) = 1; + gimple_add_tmp_var (elems); + } + else + { + /* Handle dynamic sizes. */ + sorry ("dynamic iterator sizes not implemented yet"); + } + loop.elems.put (c, elems); + + /* elems[0] = count; */ + tree lhs = build4 (ARRAY_REF, ptr_type_node, elems, size_int (0), + NULL_TREE, NULL_TREE); + tree tem = build2_loc (OMP_CLAUSE_LOCATION (c), MODIFY_EXPR, + void_type_node, lhs, loop.count); + gimplify_and_add (tem, pre_p); +} + +/* Set EXPR as the hostaddr expression that should result from the clause C. + LOOPS holds the intermediate loop info. Returns the tree that should be + passed as the hostaddr. */ + +static tree +lower_omp_map_iterator_expr (tree expr, tree c, iterator_loop_map_t *loops) +{ + if (!OMP_CLAUSE_HAS_ITERATORS (c) || !OMP_CLAUSE_ITERATORS (c)) + return expr; + + iterator_loop_info_t *loop = loops->get (OMP_CLAUSE_ITERATORS (c)); + gcc_assert (loop); + tree *elems = loop->elems.get (c); + + /* IN LOOP BODY: */ + /* elems[idx] = ; */ + tree lhs = build4 (ARRAY_REF, ptr_type_node, *elems, loop->index, NULL_TREE, + NULL_TREE); + tree tem = build2_loc (OMP_CLAUSE_LOCATION (c), MODIFY_EXPR, void_type_node, + lhs, expr); + append_to_statement_list_force (tem, loop->body); + + return build_fold_addr_expr_with_type (*elems, ptr_type_node); +} + +/* Set SIZE as the size expression that should result from the clause C. + LOOPS holds the intermediate loop info. Returns the tree that should be + passed as the clause size. */ + +static tree +lower_omp_map_iterator_size (tree size, tree c, iterator_loop_map_t *loops) +{ + if (!OMP_CLAUSE_HAS_ITERATORS (c) || !OMP_CLAUSE_ITERATORS (c)) + return size; + + iterator_loop_info_t *loop = loops->get (OMP_CLAUSE_ITERATORS (c)); + gcc_assert (loop); + tree *elems = loop->elems.get (c); + + /* IN LOOP BODY: */ + /* elems[idx+1] = size; */ + tree lhs = build4 (ARRAY_REF, ptr_type_node, *elems, + size_binop (PLUS_EXPR, loop->index, size_int (1)), + NULL_TREE, NULL_TREE); + tree tem = build2_loc (OMP_CLAUSE_LOCATION (c), MODIFY_EXPR, + void_type_node, lhs, size); + append_to_statement_list_force (tem, loop->body); + + return size_int (SIZE_MAX); +} + +/* Finish building the iterator loops in LOOPS, with generated Gimple + statements going in PRE_P. The loops cannot be amended after this is + called. */ + +static void +finish_omp_map_iterators (iterator_loop_map_t *loops, gimple_seq *pre_p) +{ + for (iterator_loop_map_t::iterator it = loops->begin (); + it != loops->end (); ++it) + { + iterator_loop_info_t &loop = (*it).second; + gimplify_and_add (loop.bind, pre_p); + + for (hash_map::iterator it2 = loop.elems.begin (); + it2 != loop.elems.end (); ++it2) + { + tree clause = (*it2).first; + OMP_CLAUSE_DECL (clause) = (*it2).second; + OMP_CLAUSE_SIZE (clause) = size_int (SIZE_MAX); + } + } +} + /* Lower the GIMPLE_OMP_TARGET in the current statement in GSI_P. CTX holds context information for the directive. */ @@ -12617,7 +12774,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) tree child_fn, t, c; gomp_target *stmt = as_a (gsi_stmt (*gsi_p)); gbind *tgt_bind, *bind, *dep_bind = NULL; - gimple_seq tgt_body, olist, ilist, fplist, new_body; + gimple_seq tgt_body, olist, iterlist, ilist, fplist, new_body; location_t loc = gimple_location (stmt); bool offloaded, data_region; unsigned int map_cnt = 0; @@ -12628,6 +12785,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) tree deep_map_offset_data = NULL_TREE; tree deep_map_offset = NULL_TREE; + iterator_loop_map_t iterator_loops; + offloaded = is_gimple_omp_offloaded (stmt); switch (gimple_omp_target_kind (stmt)) { @@ -12706,6 +12865,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) push_gimplify_context (); fplist = NULL; + iterlist = NULL; ilist = NULL; olist = NULL; for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c)) @@ -12761,7 +12921,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) gcc_unreachable (); } #endif - /* FALLTHRU */ + lower_omp_map_iterators (c, &iterlist, &iterator_loops); + /* FALLTHRU */ case OMP_CLAUSE_TO: case OMP_CLAUSE_FROM: oacc_firstprivate: @@ -13190,6 +13351,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) *p = build_fold_indirect_ref (nd); } v = build_fold_addr_expr_with_type (v, ptr_type_node); + v = lower_omp_map_iterator_expr (v, c, &iterator_loops); gimplify_assign (x, v, &ilist); nc = NULL_TREE; } @@ -13263,12 +13425,18 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) && TREE_CODE (TREE_TYPE (ovar)) == ARRAY_TYPE) { gcc_assert (offloaded); - tree avar - = create_tmp_var (TREE_TYPE (TREE_TYPE (x))); - mark_addressable (avar); - gimplify_assign (avar, build_fold_addr_expr (var), &ilist); - talign = DECL_ALIGN_UNIT (avar); + tree avar = build_fold_addr_expr (var); + if (!OMP_CLAUSE_ITERATORS (c)) + { + tree tmp = create_tmp_var (TREE_TYPE (TREE_TYPE (x))); + mark_addressable (tmp); + gimplify_assign (tmp, avar, &ilist); + avar = tmp; + } + talign = TYPE_ALIGN_UNIT (TREE_TYPE (TREE_TYPE (x))); avar = build_fold_addr_expr (avar); + avar = lower_omp_map_iterator_expr (avar, c, + &iterator_loops); gimplify_assign (x, avar, &ilist); } else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE) @@ -13348,6 +13516,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) if (s == NULL_TREE) s = TYPE_SIZE_UNIT (TREE_TYPE (ovar)); s = fold_convert (size_type_node, s); + s = lower_omp_map_iterator_size (s, c, &iterator_loops); purpose = size_int (map_idx++); CONSTRUCTOR_APPEND_ELT (vsize, purpose, s); if (TREE_CODE (s) != INTEGER_CST) @@ -13713,6 +13882,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) DECL_INITIAL (TREE_VEC_ELT (t, 2)) = build_constructor (TREE_TYPE (TREE_VEC_ELT (t, 2)), vkind); } + finish_omp_map_iterators (&iterator_loops, &iterlist); for (int i = 1; i <= 2; i++) if (deep_map_cnt || !TREE_STATIC (TREE_VEC_ELT (t, i))) { @@ -14280,6 +14450,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) gimple_omp_set_body (stmt, new_body); } + gsi_insert_seq_before (gsi_p, iterlist, GSI_SAME_STMT); + bind = gimple_build_bind (NULL, NULL, tgt_bind ? gimple_bind_block (tgt_bind) : NULL_TREE); diff --git a/gcc/testsuite/c-c++-common/gomp/map-6.c b/gcc/testsuite/c-c++-common/gomp/map-6.c index 014ed35ab41..13e3b58cc92 100644 --- a/gcc/testsuite/c-c++-common/gomp/map-6.c +++ b/gcc/testsuite/c-c++-common/gomp/map-6.c @@ -13,19 +13,19 @@ foo (void) #pragma omp target map (to:a) ; - #pragma omp target map (a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present'" } */ + #pragma omp target map (a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close', 'iterator' or 'present'" } */ ; - #pragma omp target map (close, a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present'" } */ + #pragma omp target map (close, a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close', 'iterator' or 'present'" } */ ; - #pragma omp target enter data map(b7) map (close, a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present'" } */ + #pragma omp target enter data map(b7) map (close, a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close', 'iterator' or 'present'" } */ ; - #pragma omp target exit data map(b7) map (close, a from: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present'" } */ + #pragma omp target exit data map(b7) map (close, a from: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close', 'iterator' or 'present'" } */ ; - #pragma omp target data map(b7) map (close, a from: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present'" } */ + #pragma omp target data map(b7) map (close, a from: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close', 'iterator' or 'present'" } */ ; diff --git a/gcc/testsuite/c-c++-common/gomp/target-map-iterators-1.c b/gcc/testsuite/c-c++-common/gomp/target-map-iterators-1.c new file mode 100644 index 00000000000..7d6c8dc6255 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/target-map-iterators-1.c @@ -0,0 +1,23 @@ +/* { dg-do compile } */ +/* { dg-options "-fopenmp" } */ + +#define DIM1 17 +#define DIM2 39 + +void f (int **x, int **y) +{ + #pragma omp target map(iterator(i=0:DIM1), to: x[i][:DIM2]) + ; + + #pragma omp target map(iterator(i=0:DIM1), to: x[i][:DIM2], y[i][:DIM2]) + ; + + #pragma omp target map(iterator(i=0:DIM1), to: x[i][:DIM2] + 2) /* { dg-message "unsupported map expression" } */ + ; + + #pragma omp target map(iterator(i=0:DIM1), iterator(j=0:DIM2), to: x[i][j]) /* { dg-error "too many 'iterator' modifiers" } */ + ; + + #pragma omp target map(iterator(i=0:DIM1), to: (i % 2 == 0) ? x[i] : y[i]) /* { dg-message "unsupported map expression" } */ + ; +} diff --git a/gcc/testsuite/c-c++-common/gomp/target-map-iterators-2.c b/gcc/testsuite/c-c++-common/gomp/target-map-iterators-2.c new file mode 100644 index 00000000000..da14d068f19 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/target-map-iterators-2.c @@ -0,0 +1,19 @@ +/* { dg-do compile } */ +/* { dg-options "-fopenmp" } */ + +void f (int *x, float *y, double *z) +{ + #pragma omp target map(iterator(i=0:10), to: x) /* { dg-error "iterator variable .i. not used in clause expression" } */ + /* Add a reference to x to ensure that the 'to' clause does not get + dropped. */ + x[0] = 0; + + #pragma omp target map(iterator(i=0:10, j=0:20), to: x[i]) /* { dg-error "iterator variable .j. not used in clause expression" } */ + ; + + #pragma omp target map(iterator(i=0:10, j=0:20, k=0:30), to: x[i], y[j], z[k]) + /* { dg-error "iterator variable .i. not used in clause expression" "" { target *-*-* } .-1 } */ + /* { dg-error "iterator variable .j. not used in clause expression" "" { target *-*-* } .-2 } */ + /* { dg-error "iterator variable .k. not used in clause expression" "" { target *-*-* } .-3 } */ + ; +} diff --git a/gcc/testsuite/c-c++-common/gomp/target-map-iterators-3.c b/gcc/testsuite/c-c++-common/gomp/target-map-iterators-3.c new file mode 100644 index 00000000000..fb0c761018a --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/target-map-iterators-3.c @@ -0,0 +1,23 @@ +/* { dg-do compile } */ +/* { dg-options "-fopenmp -fdump-tree-omplower" } */ + +#define DIM1 10 +#define DIM2 20 +#define DIM3 30 + +void f (int ***x, float ***y, double **z) +{ + #pragma omp target \ + map(to: x, y) \ + map(iterator(i=0:DIM1, j=0:DIM2), to: x[i][j][:DIM3], y[i][j][:DIM3]) \ + map(from: z) \ + map(iterator(i=0:DIM1), from: z[i][:DIM2]) + ; +} + +/* { dg-final { scan-tree-dump-times "if \\(i <= 9\\) goto ; else goto ;" 2 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "if \\(j <= 19\\) goto ; else goto ;" 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "map\\(iterator\\(int i=0:10:1\\):from:D\.\[0-9\]+" 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "map\\(iterator\\(int i=0:10:1\\):attach:D\.\[0-9\]+" 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "map\\(iterator\\(int i=0:10:1, int j=0:20:1\\):to:D\.\[0-9\]+" 2 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "map\\(iterator\\(int i=0:10:1, int j=0:20:1\\):attach:D\.\[0-9\]+" 4 "omplower" } } */ diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc index 39e586c808c..be2723dcdae 100644 --- a/gcc/tree-pretty-print.cc +++ b/gcc/tree-pretty-print.cc @@ -911,6 +911,11 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) pp_string (pp, "map("); if (OMP_CLAUSE_MAP_READONLY (clause)) pp_string (pp, "readonly,"); + if (OMP_CLAUSE_ITERATORS (clause)) + { + dump_omp_iterators (pp, OMP_CLAUSE_ITERATORS (clause), spc, flags); + pp_colon (pp); + } switch (OMP_CLAUSE_MAP_KIND (clause)) { case GOMP_MAP_ALLOC: diff --git a/gcc/tree.cc b/gcc/tree.cc index bc50afca9a3..f12d7b8bb8a 100644 --- a/gcc/tree.cc +++ b/gcc/tree.cc @@ -266,7 +266,7 @@ unsigned const char omp_clause_num_ops[] = 1, /* OMP_CLAUSE_EXCLUSIVE */ 2, /* OMP_CLAUSE_FROM */ 2, /* OMP_CLAUSE_TO */ - 2, /* OMP_CLAUSE_MAP */ + 3, /* OMP_CLAUSE_MAP */ 1, /* OMP_CLAUSE_HAS_DEVICE_ADDR */ 1, /* OMP_CLAUSE_DOACROSS */ 2, /* OMP_CLAUSE__CACHE_ */ @@ -11598,6 +11598,9 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data, case OMP_CLAUSE: { int len = omp_clause_num_ops[OMP_CLAUSE_CODE (t)]; + /* Do not walk the iterator operand of OpenMP MAP clauses. */ + if (OMP_CLAUSE_HAS_ITERATORS (t)) + len--; for (int i = 0; i < len; i++) WALK_SUBTREE (OMP_CLAUSE_OPERAND (t, i)); WALK_SUBTREE_TAIL (OMP_CLAUSE_CHAIN (t)); diff --git a/gcc/tree.h b/gcc/tree.h index 83075b82cc7..384a5f1f250 100644 --- a/gcc/tree.h +++ b/gcc/tree.h @@ -1619,6 +1619,13 @@ class auto_suppress_location_wrappers != UNKNOWN_LOCATION) #define OMP_CLAUSE_LOCATION(NODE) (OMP_CLAUSE_CHECK (NODE))->omp_clause.locus +#define OMP_CLAUSE_HAS_ITERATORS(NODE) \ + (OMP_CLAUSE_CODE (NODE) == OMP_CLAUSE_MAP) +#define OMP_CLAUSE_ITERATORS(NODE) \ + OMP_CLAUSE_OPERAND (OMP_CLAUSE_RANGE_CHECK (OMP_CLAUSE_CHECK (NODE), \ + OMP_CLAUSE_MAP, \ + OMP_CLAUSE_MAP), 2) + /* True on OMP_FOR and other OpenMP/OpenACC looping constructs if the loop nest is non-rectangular. */ #define OMP_FOR_NON_RECTANGULAR(NODE) \ diff --git a/libgomp/target.c b/libgomp/target.c index cf62af61f3b..463a162879b 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -975,6 +975,105 @@ gomp_map_val (struct target_mem_desc *tgt, void **hostaddrs, size_t i) } } +static const char * +kind_to_name (unsigned short kind) +{ + if (GOMP_MAP_IMPLICIT_P (kind)) + kind &= ~GOMP_MAP_IMPLICIT; + + switch (kind & 0xff) + { + case GOMP_MAP_ALLOC: return "GOMP_MAP_ALLOC"; + case GOMP_MAP_FIRSTPRIVATE: return "GOMP_MAP_FIRSTPRIVATE"; + case GOMP_MAP_FIRSTPRIVATE_INT: return "GOMP_MAP_FIRSTPRIVATE_INT"; + case GOMP_MAP_TO: return "GOMP_MAP_TO"; + case GOMP_MAP_TO_PSET: return "GOMP_MAP_TO_PSET"; + case GOMP_MAP_FROM: return "GOMP_MAP_FROM"; + case GOMP_MAP_TOFROM: return "GOMP_MAP_TOFROM"; + case GOMP_MAP_POINTER: return "GOMP_MAP_POINTER"; + case GOMP_MAP_ATTACH: return "GOMP_MAP_ATTACH"; + case GOMP_MAP_DETACH: return "GOMP_MAP_DETACH"; + default: return "unknown"; + } +} + +/* Map entries containing expanded iterators will be flattened and merged into + HOSTADDRS, SIZES and KINDS, and MAPNUM updated. Returns true if there are + any iterators found. ITERATOR_COUNT holds the iteration count of the + iterator that generates each map (0 if not generated from an iterator). + HOSTADDRS, SIZES, KINDS and ITERATOR_COUNT must be freed afterwards if any + merging occurs. */ + +static bool +gomp_merge_iterator_maps (size_t *mapnum, void ***hostaddrs, size_t **sizes, + void **kinds, size_t **iterator_count) +{ + bool iterator_p = false; + size_t map_count = 0; + unsigned short **skinds = (unsigned short **) kinds; + + for (size_t i = 0; i < *mapnum; i++) + if ((*sizes)[i] == SIZE_MAX) + { + uintptr_t *iterator_array = (*hostaddrs)[i]; + map_count += iterator_array[0]; + iterator_p = true; + } + else + map_count++; + + if (!iterator_p) + return false; + + gomp_debug (1, + "Expanding iterator maps - number of map entries: %u -> %u\n", + (int) *mapnum, (int) map_count); + void **new_hostaddrs = (void **) gomp_malloc (map_count * sizeof (void *)); + size_t *new_sizes = (size_t *) gomp_malloc (map_count * sizeof (size_t)); + unsigned short *new_kinds + = (unsigned short *) gomp_malloc (map_count * sizeof (unsigned short)); + size_t new_idx = 0; + *iterator_count = (size_t *) gomp_malloc (map_count * sizeof (size_t)); + + for (size_t i = 0; i < *mapnum; i++) + { + if ((*sizes)[i] == SIZE_MAX) + { + uintptr_t *iterator_array = (*hostaddrs)[i]; + size_t count = *iterator_array++; + for (size_t j = 0; j < count; j++) + { + new_hostaddrs[new_idx] = (void *) *iterator_array++; + new_sizes[new_idx] = *iterator_array++; + new_kinds[new_idx] = (*skinds)[i]; + (*iterator_count)[new_idx] = j + 1; + gomp_debug (1, + "Expanding map %u <%s>: " + "hostaddrs[%u] = %p, sizes[%u] = %lu\n", + (int) i, kind_to_name (new_kinds[new_idx]), + (int) new_idx, new_hostaddrs[new_idx], + (int) new_idx, (unsigned long) new_sizes[new_idx]); + new_idx++; + } + } + else + { + new_hostaddrs[new_idx] = (*hostaddrs)[i]; + new_sizes[new_idx] = (*sizes)[i]; + new_kinds[new_idx] = (*skinds)[i]; + (*iterator_count)[new_idx] = 0; + new_idx++; + } + } + + *mapnum = map_count; + *hostaddrs = new_hostaddrs; + *sizes = new_sizes; + *kinds = new_kinds; + + return true; +} + static inline __attribute__((always_inline)) struct target_mem_desc * gomp_map_vars_internal (struct gomp_device_descr *devicep, struct goacc_asyncqueue *aq, size_t mapnum, @@ -991,6 +1090,11 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, const int typemask = short_mapkind ? 0xff : 0x7; struct splay_tree_s *mem_map = &devicep->mem_map; struct splay_tree_key_s cur_node; + bool iterators_p = false; + size_t *iterator_count = NULL; + if (short_mapkind) + iterators_p = gomp_merge_iterator_maps (&mapnum, &hostaddrs, &sizes, + &kinds, &iterator_count); struct target_mem_desc *tgt = gomp_malloc (sizeof (*tgt) + sizeof (tgt->list[0]) * mapnum); tgt->list_count = mapnum; @@ -1840,14 +1944,17 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, if (pragma_kind & GOMP_MAP_VARS_TARGET) { + size_t map_num = 0; for (i = 0; i < mapnum; i++) - { - cur_node.tgt_offset = gomp_map_val (tgt, hostaddrs, i); - gomp_copy_host2dev (devicep, aq, - (void *) (tgt->tgt_start + i * sizeof (void *)), - (void *) &cur_node.tgt_offset, sizeof (void *), - true, cbufp); - } + if (!iterator_count || iterator_count[i] <= 1) + { + cur_node.tgt_offset = gomp_map_val (tgt, hostaddrs, i); + gomp_copy_host2dev (devicep, aq, + (void *) (tgt->tgt_start + map_num * sizeof (void *)), + (void *) &cur_node.tgt_offset, sizeof (void *), + true, cbufp); + map_num++; + } } if (cbufp) @@ -1879,6 +1986,15 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, } gomp_mutex_unlock (&devicep->lock); + + if (iterators_p) + { + free (hostaddrs); + free (sizes); + free (kinds); + free (iterator_count); + } + return tgt; } diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-1.c b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-1.c new file mode 100644 index 00000000000..b3d87f231df --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-1.c @@ -0,0 +1,47 @@ +/* { dg-do run } */ +/* { dg-require-effective-target offload_device_nonshared_as } */ + +/* Test transfer of dynamically-allocated arrays to target using map + iterators. */ + +#include + +#define DIM1 8 +#define DIM2 15 + +int mkarray (int *x[]) +{ + int expected = 0; + + for (int i = 0; i < DIM1; i++) + { + x[i] = (int *) malloc (DIM2 * sizeof (int)); + for (int j = 0; j < DIM2; j++) + { + x[i][j] = rand (); + expected += x[i][j]; + } + } + + return expected; +} + +int main (void) +{ + int *x[DIM1]; + int y; + + int expected = mkarray (x); + + #pragma omp target enter data map(to: x) + #pragma omp target map(iterator(i=0:DIM1), to: x[i][:DIM2]) \ + map(from: y) + { + y = 0; + for (int i = 0; i < DIM1; i++) + for (int j = 0; j < DIM2; j++) + y += x[i][j]; + } + + return y - expected; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-2.c b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-2.c new file mode 100644 index 00000000000..8569b55ab5b --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-2.c @@ -0,0 +1,44 @@ +/* { dg-do run } */ +/* { dg-require-effective-target offload_device_nonshared_as } */ + +/* Test transfer of dynamically-allocated arrays from target using map + iterators. */ + +#include + +#define DIM1 8 +#define DIM2 15 + +void mkarray (int *x[]) +{ + for (int i = 0; i < DIM1; i++) + x[i] = (int *) malloc (DIM2 * sizeof (int)); +} + +int main (void) +{ + int *x[DIM1]; + int y, expected; + + mkarray (x); + + #pragma omp target enter data map(alloc: x) + #pragma omp target map(iterator(i=0:DIM1), from: x[i][:DIM2]) \ + map(from: expected) + { + expected = 0; + for (int i = 0; i < DIM1; i++) + for (int j = 0; j < DIM2; j++) + { + x[i][j] = (i+1) * (j+1); + expected += x[i][j]; + } + } + + y = 0; + for (int i = 0; i < DIM1; i++) + for (int j = 0; j < DIM2; j++) + y += x[i][j]; + + return y - expected; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-3.c b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-3.c new file mode 100644 index 00000000000..be30fa65d80 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-3.c @@ -0,0 +1,56 @@ +/* { dg-do run } */ +/* { dg-require-effective-target offload_device_nonshared_as } */ + +/* Test transfer of dynamically-allocated arrays to target using map + iterators, with multiple iterators and function calls in the iterator + expression. */ + +#include + +#define DIM1 16 +#define DIM2 15 + +int mkarrays (int *x[], int *y[]) +{ + int expected = 0; + + for (int i = 0; i < DIM1; i++) + { + x[i] = (int *) malloc (DIM2 * sizeof (int)); + y[i] = (int *) malloc (sizeof (int)); + *y[i] = rand (); + for (int j = 0; j < DIM2; j++) + { + x[i][j] = rand (); + expected += x[i][j] * *y[i]; + } + } + + return expected; +} + +int f (int i, int j) +{ + return i * 4 + j; +} + +int main (void) +{ + int *x[DIM1], *y[DIM1]; + int sum; + + int expected = mkarrays (x, y); + + #pragma omp target enter data map(to: x, y) + #pragma omp target map(iterator(i=0:DIM1/4, j=0:4), to: x[f(i, j)][:DIM2]) \ + map(iterator(i=0:DIM1), to: y[i][:1]) \ + map(from: sum) + { + sum = 0; + for (int i = 0; i < DIM1; i++) + for (int j = 0; j < DIM2; j++) + sum += x[i][j] * y[i][0]; + } + + return sum - expected; +} From patchwork Fri Oct 4 14:56:49 2024 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Kwok Cheung Yeung X-Patchwork-Id: 1992780 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@legolas.ozlabs.org Authentication-Results: legolas.ozlabs.org; dkim=pass (2048-bit key; unprotected) header.d=baylibre-com.20230601.gappssmtp.com header.i=@baylibre-com.20230601.gappssmtp.com header.a=rsa-sha256 header.s=20230601 header.b=kqltwzNP; dkim-atps=neutral Authentication-Results: legolas.ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=8.43.85.97; helo=server2.sourceware.org; envelope-from=gcc-patches-bounces~incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=patchwork.ozlabs.org) Received: from server2.sourceware.org (server2.sourceware.org [8.43.85.97]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature ECDSA (secp384r1) server-digest SHA384) (No client certificate requested) by legolas.ozlabs.org (Postfix) with ESMTPS id 4XKsBR4WS9z1xt7 for ; Sat, 5 Oct 2024 00:58:55 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id C4D953842FFF for ; Fri, 4 Oct 2024 14:58:53 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from mail-lj1-x230.google.com (mail-lj1-x230.google.com [IPv6:2a00:1450:4864:20::230]) by sourceware.org (Postfix) with ESMTPS id 013743842AFC for ; Fri, 4 Oct 2024 14:57:18 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 013743842AFC Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=baylibre.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=baylibre.com ARC-Filter: OpenARC Filter v1.0.0 sourceware.org 013743842AFC Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=2a00:1450:4864:20::230 ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1728053842; cv=none; b=Wq3DvO4smRvNnlzpmkjGQJY2sT029VLUmsoBJYwi9NpvL50UmmRVzrTdVigPdqtbU4F3U6HIyaaUnn8SwoIqFLtDk2C4lk1rLdBc61AkhW1cFzbhwr1MnYZXxiTBSK6WYayQllYtULpeuwZslIhq2OPr73CPP8IVvMtROEMRAGs= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1728053842; c=relaxed/simple; bh=4hcac8RTOycRjQ/ypLGqNOHJ0PQobwI3Z9lCl0I5uyM=; h=DKIM-Signature:Message-ID:Date:MIME-Version:To:Subject:From; b=l6M0Y3UIiBjcoqH6N7xvuN8+Z2AxEy+e+lalGR72j83PYSIs84zcnha3TawrlA+gVCTb2LDYdLMY5w1Yrj5vhw/9SnOkh/Tkb6rgQIgDDRPxXVJqXF5ijbqFzfq8G2dGnKYJ5DDQ035OTcfny9cCq925oIf4d7KBoaVZINpZJZI= ARC-Authentication-Results: i=1; server2.sourceware.org Received: by mail-lj1-x230.google.com with SMTP id 38308e7fff4ca-2fad100dd9fso35405881fa.3 for ; Fri, 04 Oct 2024 07:57:17 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre-com.20230601.gappssmtp.com; s=20230601; t=1728053836; x=1728658636; darn=gcc.gnu.org; h=in-reply-to:from:content-language:subject:references:to:user-agent :mime-version:date:message-id:from:to:cc:subject:date:message-id :reply-to; bh=4hcac8RTOycRjQ/ypLGqNOHJ0PQobwI3Z9lCl0I5uyM=; b=kqltwzNPR2DZXMNC++f/IkiJ+gH6fGCGWpFI4Y1//68kLdyiOVUxxzQPkNxI+COWDQ HF77QXHRNOq/0q5bKmY/1w0yd9V2zucSmqbqw/27WwJKcfU0Jp5Bj6sXuHu2sj0oU0iN V5MIgxer//rrGhMXoCj1kKOQ18FHDFHTsmUu60IOyTfX7eDhd3hlL2V9+DyUmIWH+tCh GMnEd+135VmvzLaqnC+CzgmoE8ihorTWRSIk+ayydHs9vs8jO4b0G7oAteNiw8HDjXrV V0liB/tpYnNkTCp3xDG/CbVYIYfpAG6wsk3m0pn+tc1sLy6KoGNYMhaKqKxVH+pk3Ube bUIw== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1728053836; x=1728658636; h=in-reply-to:from:content-language:subject:references:to:user-agent :mime-version:date:message-id:x-gm-message-state:from:to:cc:subject :date:message-id:reply-to; bh=4hcac8RTOycRjQ/ypLGqNOHJ0PQobwI3Z9lCl0I5uyM=; b=IvvAmP288xqNnKXEqX8VP9DgraR9ZeyHBgR2FFuzX/VOlXqJKB74MQSCP2sSaYvy9L paHu2ktDzF+Z3Y6CYgn1oOuOKBh1WafcpbekJr4akCGX7+ReSOYszTNC2e7UkD4LP/sS 5inDqqEEbCnbAsnWCiir1galXBgF88d5vAiypVv+D0Dp8+SQrarKpFn+01BkEpZLMhoY qnnQL0gFJYXf/EKZcPPanx7T1s7CnO/hUicb/HuAFdovp9ZpTGXhbVAVawUGVlfzeiR9 TOTi4U183lCeTR8LMHb9rVR0KvTeLoT09clePn9YZjYANBDIDe7faSciYcBKKHD74F99 nqsg== X-Gm-Message-State: AOJu0YwL6B704bQvUCjlz9+4dK4OnkT0Gv+7uy0BSzurX9tlWFdZDnpD 1hYIOtLr/KZFt7GiDILgl2RINmMMtWxyAtGncVUQpcl4OxAUTQifP3OlpeuiqpM3zaC8NAKJV53 h X-Google-Smtp-Source: AGHT+IGsY08kdDCjBpDmbklHQnGPqzEI88uIF5M7n50urOAbAlf//HBmNhpV9ycm+n4FJdZIkwIWVQ== X-Received: by 2002:a2e:878a:0:b0:2f7:5a41:b0b with SMTP id 38308e7fff4ca-2faf3c44dc5mr23111241fa.26.1728053835973; Fri, 04 Oct 2024 07:57:15 -0700 (PDT) Received: from ?IPV6:2a00:23c6:88fe:9301:2d7d:f734:bc6:c47b? ([2a00:23c6:88fe:9301:2d7d:f734:bc6:c47b]) by smtp.gmail.com with ESMTPSA id a640c23a62f3a-a992e5bbc2csm2048766b.35.2024.10.04.07.57.15 (version=TLS1_3 cipher=TLS_AES_128_GCM_SHA256 bits=128/128); Fri, 04 Oct 2024 07:57:15 -0700 (PDT) Message-ID: Date: Fri, 4 Oct 2024 15:56:49 +0100 MIME-Version: 1.0 User-Agent: Mozilla Thunderbird To: gcc-patches , Jakub Jelinek , Tobias Burnus References: <6b94b8ed-020b-47e2-b02a-4891891f2847@baylibre.com> Subject: [PATCH v3 3/5] openmp: Add support for iterators in 'target update' clauses (C/C++) Content-Language: en-GB From: Kwok Cheung Yeung In-Reply-To: <6b94b8ed-020b-47e2-b02a-4891891f2847@baylibre.com> X-Spam-Status: No, score=-12.7 required=5.0 tests=BAYES_00, DKIM_SIGNED, DKIM_VALID, GIT_PATCH_0, RCVD_IN_DNSWL_NONE, SPF_HELO_NONE, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.30 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces~incoming=patchwork.ozlabs.org@gcc.gnu.org This patch extends the previous patch to cover to/from clauses in 'target update'. From 1c8bf84ec99fe2fd371e345f012eb0d84a923153 Mon Sep 17 00:00:00 2001 From: Kwok Cheung Yeung Date: Fri, 4 Oct 2024 15:16:21 +0100 Subject: [PATCH 3/5] openmp: Add support for iterators in 'target update' clauses (C/C++) This adds support for iterators in 'to' and 'from' clauses in the 'target update' OpenMP directive. 2024-10-04 Kwok Cheung Yeung gcc/c/ * c-parser.cc (c_parser_omp_clause_from_to): Parse 'iterator' modifier. * c-typeck.cc (c_finish_omp_clauses): Finish iterators for to/from clauses. gcc/cp/ * parser.cc (cp_parser_omp_clause_from_to): Parse 'iterator' modifier. * semantics.cc (finish_omp_clauses): Finish iterators for to/from clauses. gcc/ * gimplify.cc (gimplify_scan_omp_clauses): Call check_omp_map_iterators on clauses with iterators. Skip gimplification of clause decl and size for clauses with iterators. * omp-low.cc (lower_omp_target): Call lower_omp_map_iterators on to/from clauses. * tree-pretty-print.cc (dump_omp_clause): Call dump_omp_iterators for to/from clauses with iterators. * tree.cc (omp_clause_num_ops): Add extra operand for OMP_CLAUSE_FROM and OMP_CLAUSE_TO. * tree.h (OMP_CLAUSE_HAS_ITERATORS): Add check for OMP_CLAUSE_TO and OMP_CLAUSE_FROM. (OMP_CLAUSE_ITERATORS): Likewise. gcc/testsuite/ * c-c++-common/gomp/target-update-iterators-1.c: New. * c-c++-common/gomp/target-update-iterators-2.c: New. * c-c++-common/gomp/target-update-iterators-3.c: New. libgomp/ * target.c (gomp_update): Call gomp_merge_iterator_maps. Free allocated variables. * testsuite/libgomp.c-c++-common/target-update-iterators-1.c: New. * testsuite/libgomp.c-c++-common/target-update-iterators-2.c: New. * testsuite/libgomp.c-c++-common/target-update-iterators-3.c: New. --- gcc/c/c-parser.cc | 105 +++++++++++++++-- gcc/c/c-typeck.cc | 5 +- gcc/cp/parser.cc | 111 ++++++++++++++++-- gcc/cp/semantics.cc | 5 +- gcc/gimplify.cc | 18 ++- gcc/omp-low.cc | 3 +- .../gomp/target-update-iterators-1.c | 20 ++++ .../gomp/target-update-iterators-2.c | 17 +++ .../gomp/target-update-iterators-3.c | 17 +++ gcc/tree-pretty-print.cc | 10 ++ gcc/tree.cc | 4 +- gcc/tree.h | 8 +- libgomp/target.c | 14 +++ .../target-update-iterators-1.c | 65 ++++++++++ .../target-update-iterators-2.c | 58 +++++++++ .../target-update-iterators-3.c | 67 +++++++++++ 16 files changed, 496 insertions(+), 31 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/gomp/target-update-iterators-1.c create mode 100644 gcc/testsuite/c-c++-common/gomp/target-update-iterators-2.c create mode 100644 gcc/testsuite/c-c++-common/gomp/target-update-iterators-3.c create mode 100644 libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-1.c create mode 100644 libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-2.c create mode 100644 libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-3.c diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc index 184fc076388..c2a5985c89b 100644 --- a/gcc/c/c-parser.cc +++ b/gcc/c/c-parser.cc @@ -19304,8 +19304,11 @@ c_parser_omp_clause_device_type (c_parser *parser, tree list) to ( variable-list ) OpenMP 5.1: - from ( [present :] variable-list ) - to ( [present :] variable-list ) */ + from ( [motion-modifier[,] [motion-modifier[,]...]:] variable-list ) + to ( [motion-modifier[,] [motion-modifier[,]...]:] variable-list ) + + motion-modifier: + present | iterator (iterators-definition) */ static tree c_parser_omp_clause_from_to (c_parser *parser, enum omp_clause_code kind, @@ -19316,15 +19319,88 @@ c_parser_omp_clause_from_to (c_parser *parser, enum omp_clause_code kind, if (!parens.require_open (parser)) return list; + int pos = 1, colon_pos = 0; + int iterator_length = 0; + while (c_parser_peek_nth_token_raw (parser, pos)->type == CPP_NAME) + { + if (c_parser_peek_nth_token_raw (parser, pos + 1)->type + == CPP_OPEN_PAREN) + { + unsigned int n = pos + 2; + if (c_parser_check_balanced_raw_token_sequence (parser, &n) + && (c_parser_peek_nth_token_raw (parser, n)->type + == CPP_CLOSE_PAREN)) + { + iterator_length = n - pos + 1; + pos = n; + } + } + if (c_parser_peek_nth_token_raw (parser, pos + 1)->type == CPP_COMMA) + pos += 2; + else + pos++; + if (c_parser_peek_nth_token_raw (parser, pos)->type == CPP_COLON) + { + colon_pos = pos; + break; + } + } + bool present = false; - c_token *token = c_parser_peek_token (parser); + tree iterators = NULL_TREE; - if (token->type == CPP_NAME - && strcmp (IDENTIFIER_POINTER (token->value), "present") == 0 - && c_parser_peek_2nd_token (parser)->type == CPP_COLON) + for (pos = 1; pos < colon_pos; pos++) { - present = true; - c_parser_consume_token (parser); + c_token *token = c_parser_peek_token (parser); + + if (token->type == CPP_COMMA) + { + c_parser_consume_token (parser); + continue; + } + if (token->type == CPP_NAME) + { + const char *name = IDENTIFIER_POINTER (token->value); + if (strcmp (name, "present") == 0) + { + if (present) + { + c_parser_error (parser, "too many % modifiers"); + parens.skip_until_found_close (parser); + return list; + } + present = true; + c_parser_consume_token (parser); + } + else if (strcmp (name, "iterator") == 0) + { + if (iterators) + { + c_parser_error (parser, "too many % modifiers"); + parens.skip_until_found_close (parser); + return list; + } + iterators = c_parser_omp_iterators (parser); + pos += iterator_length - 1; + } + else + { + if (kind == OMP_CLAUSE_TO) + c_parser_error (parser, "% clause with motion modifier " + "other than % or %"); + else + c_parser_error (parser, "% clause with motion modifier " + "other than % or %"); + parens.skip_until_found_close (parser); + return list; + } + } + } + + if (colon_pos) + { + gcc_assert (pos == colon_pos); + gcc_assert (c_parser_next_token_is (parser, CPP_COLON)); c_parser_consume_token (parser); } @@ -19335,6 +19411,19 @@ c_parser_omp_clause_from_to (c_parser *parser, enum omp_clause_code kind, for (tree c = nl; c != list; c = OMP_CLAUSE_CHAIN (c)) OMP_CLAUSE_MOTION_PRESENT (c) = 1; + if (iterators) + { + tree block = pop_scope (); + if (iterators == error_mark_node) + iterators = NULL_TREE; + else + TREE_VEC_ELT (iterators, 5) = block; + } + + if (iterators) + for (tree c = nl; c != list; c = OMP_CLAUSE_CHAIN (c)) + OMP_CLAUSE_ITERATORS (c) = iterators; + return nl; } diff --git a/gcc/c/c-typeck.cc b/gcc/c/c-typeck.cc index cca9f1c000c..5d16f749133 100644 --- a/gcc/c/c-typeck.cc +++ b/gcc/c/c-typeck.cc @@ -15813,6 +15813,9 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) case OMP_CLAUSE_MAP: if (OMP_CLAUSE_MAP_IMPLICIT (c) && !implicit_moved) goto move_implicit; + /* FALLTHRU */ + case OMP_CLAUSE_TO: + case OMP_CLAUSE_FROM: if (OMP_CLAUSE_ITERATORS (c) && c_omp_finish_iterators (OMP_CLAUSE_ITERATORS (c))) { @@ -15820,8 +15823,6 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) break; } /* FALLTHRU */ - case OMP_CLAUSE_TO: - case OMP_CLAUSE_FROM: case OMP_CLAUSE__CACHE_: { using namespace omp_addr_tokenizer; diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc index 79d6e115d16..861337803e6 100644 --- a/gcc/cp/parser.cc +++ b/gcc/cp/parser.cc @@ -41716,8 +41716,11 @@ cp_parser_omp_clause_doacross (cp_parser *parser, tree list, location_t loc) to ( variable-list ) OpenMP 5.1: - from ( [present :] variable-list ) - to ( [present :] variable-list ) */ + from ( [motion-modifier[,] [motion-modifier[,]...]:] variable-list ) + to ( [motion-modifier[,] [motion-modifier[,]...]:] variable-list ) + + motion-modifier: + present | iterator (iterators-definition) */ static tree cp_parser_omp_clause_from_to (cp_parser *parser, enum omp_clause_code kind, @@ -41726,15 +41729,94 @@ cp_parser_omp_clause_from_to (cp_parser *parser, enum omp_clause_code kind, if (!cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN)) return list; + size_t pos = 1, colon_pos = 0; + int iterator_length = 0; + while (cp_lexer_nth_token_is (parser->lexer, pos, CPP_NAME)) + { + if (cp_lexer_nth_token_is (parser->lexer, pos + 1, CPP_OPEN_PAREN)) + { + unsigned int n = cp_parser_skip_balanced_tokens (parser, pos + 1); + if (n != pos + 1) + { + iterator_length = n - pos; + pos = n - 1; + } + } + if (cp_lexer_nth_token_is (parser->lexer, pos + 1, CPP_COMMA)) + pos += 2; + else + pos++; + if (cp_lexer_nth_token_is (parser->lexer, pos, CPP_COLON)) + { + colon_pos = pos; + break; + } + } + bool present = false; - cp_token *token = cp_lexer_peek_token (parser->lexer); + tree iterators = NULL_TREE; + for (pos = 1; pos < colon_pos; pos++) + { + cp_token *token = cp_lexer_peek_token (parser->lexer); - if (token->type == CPP_NAME - && strcmp (IDENTIFIER_POINTER (token->u.value), "present") == 0 - && cp_lexer_nth_token_is (parser->lexer, 2, CPP_COLON)) + if (token->type == CPP_COMMA) + { + cp_lexer_consume_token (parser->lexer); + continue; + } + if (token->type == CPP_NAME) + { + const char *name = IDENTIFIER_POINTER (token->u.value); + if (strcmp (name, "present") == 0) + { + if (present) + { + cp_parser_error (parser, "too many % modifiers"); + cp_parser_skip_to_closing_parenthesis (parser, + /*recovering=*/true, + /*or_comma=*/false, + /*consume_paren=*/true); + return list; + } + present = true; + cp_lexer_consume_token (parser->lexer); + } + else if (strcmp (name, "iterator") == 0) + { + if (iterators) + { + cp_parser_error (parser, "too many % modifiers"); + cp_parser_skip_to_closing_parenthesis (parser, + /*recovering=*/true, + /*or_comma=*/false, + /*consume_paren=*/true); + return list; + } + begin_scope (sk_omp, NULL); + iterators = cp_parser_omp_iterators (parser); + pos += iterator_length - 1; + } + else + { + if (kind == OMP_CLAUSE_TO) + cp_parser_error (parser, "% clause with motion modifier " + "other than % or %"); + else + cp_parser_error (parser, "% clause with motion modifier " + "other than % or %"); + cp_parser_skip_to_closing_parenthesis (parser, + /*recovering=*/true, + /*or_comma=*/false, + /*consume_paren=*/true); + return list; + } + } + } + + if (colon_pos) { - present = true; - cp_lexer_consume_token (parser->lexer); + gcc_assert (pos == colon_pos); + gcc_assert (cp_lexer_next_token_is (parser->lexer, CPP_COLON)); cp_lexer_consume_token (parser->lexer); } @@ -41743,6 +41825,19 @@ cp_parser_omp_clause_from_to (cp_parser *parser, enum omp_clause_code kind, for (tree c = nl; c != list; c = OMP_CLAUSE_CHAIN (c)) OMP_CLAUSE_MOTION_PRESENT (c) = 1; + if (iterators) + { + tree block = poplevel (1, 1, 0); + if (iterators == error_mark_node) + iterators = NULL_TREE; + else + TREE_VEC_ELT (iterators, 5) = block; + } + + if (iterators) + for (tree c = nl; c != list; c = OMP_CLAUSE_CHAIN (c)) + OMP_CLAUSE_ITERATORS (c) = iterators; + return nl; } diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc index ba5657f7bc2..fdf814d3ce6 100644 --- a/gcc/cp/semantics.cc +++ b/gcc/cp/semantics.cc @@ -8492,6 +8492,9 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) case OMP_CLAUSE_MAP: if (OMP_CLAUSE_MAP_IMPLICIT (c) && !implicit_moved) goto move_implicit; + /* FALLTHRU */ + case OMP_CLAUSE_TO: + case OMP_CLAUSE_FROM: if (OMP_CLAUSE_ITERATORS (c) && cp_omp_finish_iterators (OMP_CLAUSE_ITERATORS (c))) { @@ -8499,8 +8502,6 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) break; } /* FALLTHRU */ - case OMP_CLAUSE_TO: - case OMP_CLAUSE_FROM: case OMP_CLAUSE__CACHE_: { using namespace omp_addr_tokenizer; diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index 6e532d07fcf..ba972a2892a 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -12862,6 +12862,12 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, case OMP_CLAUSE_TO: case OMP_CLAUSE_FROM: + if (OMP_CLAUSE_ITERATORS (c) && !check_omp_map_iterators (c)) + { + remove = true; + break; + } + /* FALLTHRU */ case OMP_CLAUSE__CACHE_: decl = OMP_CLAUSE_DECL (c); if (error_operand_p (decl)) @@ -12872,17 +12878,19 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, if (OMP_CLAUSE_SIZE (c) == NULL_TREE) OMP_CLAUSE_SIZE (c) = DECL_P (decl) ? DECL_SIZE_UNIT (decl) : TYPE_SIZE_UNIT (TREE_TYPE (decl)); - if (gimplify_expr (&OMP_CLAUSE_SIZE (c), pre_p, - NULL, is_gimple_val, fb_rvalue) == GS_ERROR) + if (!(OMP_CLAUSE_HAS_ITERATORS (c) && OMP_CLAUSE_ITERATORS (c)) + && gimplify_expr (&OMP_CLAUSE_SIZE (c), pre_p, + NULL, is_gimple_val, fb_rvalue) == GS_ERROR) { remove = true; break; } if (!DECL_P (decl)) { - if (gimplify_expr (&OMP_CLAUSE_DECL (c), pre_p, - NULL, is_gimple_lvalue, fb_lvalue) - == GS_ERROR) + if (!(OMP_CLAUSE_HAS_ITERATORS (c) && OMP_CLAUSE_ITERATORS (c)) + && gimplify_expr (&OMP_CLAUSE_DECL (c), pre_p, + NULL, is_gimple_lvalue, fb_lvalue) + == GS_ERROR) { remove = true; break; diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc index 9cf6e207d1c..a8b86889c66 100644 --- a/gcc/omp-low.cc +++ b/gcc/omp-low.cc @@ -12921,10 +12921,11 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) gcc_unreachable (); } #endif - lower_omp_map_iterators (c, &iterlist, &iterator_loops); /* FALLTHRU */ case OMP_CLAUSE_TO: case OMP_CLAUSE_FROM: + lower_omp_map_iterators (c, &iterlist, &iterator_loops); + /* FALLTHRU */ oacc_firstprivate: var = OMP_CLAUSE_DECL (c); { diff --git a/gcc/testsuite/c-c++-common/gomp/target-update-iterators-1.c b/gcc/testsuite/c-c++-common/gomp/target-update-iterators-1.c new file mode 100644 index 00000000000..3a64f511da4 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/target-update-iterators-1.c @@ -0,0 +1,20 @@ +/* { dg-do compile } */ +/* { dg-options "-fopenmp" } */ + +#define DIM1 17 +#define DIM2 39 + +void f (int **x, float **y) +{ + #pragma omp target update to (iterator(i=0:DIM1): x[i][:DIM2]) + + #pragma omp target update to (iterator(i=0:DIM1): x[i][:DIM2], y[i][:DIM2]) + + #pragma omp target update to (iterator(i=0:DIM1), present: x[i][:DIM2]) + + #pragma omp target update to (iterator(i=0:DIM1), iterator(j=0:DIM2): x[i][j]) /* { dg-error "too many 'iterator' modifiers" } */ + /* { dg-error ".#pragma omp target update. must contain at least one .from. or .to. clauses" "" { target *-*-* } .-1 } */ + + #pragma omp target update to (iterator(i=0:DIM1), something: x[i][j]) /* { dg-error ".to. clause with motion modifier other than .iterator. or .present. before .something." } */ + /* { dg-error ".#pragma omp target update. must contain at least one .from. or .to. clauses" "" { target *-*-* } .-1 } */ +} diff --git a/gcc/testsuite/c-c++-common/gomp/target-update-iterators-2.c b/gcc/testsuite/c-c++-common/gomp/target-update-iterators-2.c new file mode 100644 index 00000000000..3789a559b6f --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/target-update-iterators-2.c @@ -0,0 +1,17 @@ +/* { dg-do compile } */ +/* { dg-options "-fopenmp" } */ + +void f (int *x, float *y, double *z) +{ + #pragma omp target update to(iterator(i=0:10): x) /* { dg-error "iterator variable .i. not used in clause expression" }*/ + ; + + #pragma omp target update from(iterator(i=0:10, j=0:20): x[i]) /* { dg-error "iterator variable .j. not used in clause expression" }*/ + ; + + #pragma omp target update to(iterator(i=0:10, j=0:20, k=0:30): x[i], y[j], z[k]) + /* { dg-error "iterator variable .i. not used in clause expression" "" { target *-*-* } .-1 } */ + /* { dg-error "iterator variable .j. not used in clause expression" "" { target *-*-* } .-2 } */ + /* { dg-error "iterator variable .k. not used in clause expression" "" { target *-*-* } .-3 } */ + ; +} diff --git a/gcc/testsuite/c-c++-common/gomp/target-update-iterators-3.c b/gcc/testsuite/c-c++-common/gomp/target-update-iterators-3.c new file mode 100644 index 00000000000..b256674442f --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/target-update-iterators-3.c @@ -0,0 +1,17 @@ +/* { dg-do compile } */ +/* { dg-options "-fopenmp -fdump-tree-omplower" } */ + +#define DIM1 10 +#define DIM2 20 +#define DIM3 30 + +void f (int ***x, float ***y, double **z) +{ + #pragma omp target update to (iterator(i=0:DIM1, j=0:DIM2): x[i][j][:DIM3], y[i][j][:DIM3]) + #pragma omp target update from (iterator(i=0:DIM1): z[i][:DIM2]) +} + +/* { dg-final { scan-tree-dump-times "if \\(i <= 9\\) goto ; else goto ;" 2 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "if \\(j <= 19\\) goto ; else goto ;" 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "to\\(iterator\\(int i=0:10:1, int j=0:20:1\\):D\.\[0-9\]+" 2 "omplower" } } */ +/* { dg-final { scan-tree-dump-times "from\\(iterator\\(int i=0:10:1\\):D\.\[0-9\]+" 1 "omplower" } } */ diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc index be2723dcdae..fa1b2dce27f 100644 --- a/gcc/tree-pretty-print.cc +++ b/gcc/tree-pretty-print.cc @@ -1084,6 +1084,11 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) pp_string (pp, "from("); if (OMP_CLAUSE_MOTION_PRESENT (clause)) pp_string (pp, "present:"); + if (OMP_CLAUSE_ITERATORS (clause)) + { + dump_omp_iterators (pp, OMP_CLAUSE_ITERATORS (clause), spc, flags); + pp_colon (pp); + } dump_generic_node (pp, OMP_CLAUSE_DECL (clause), spc, flags, false); goto print_clause_size; @@ -1092,6 +1097,11 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) pp_string (pp, "to("); if (OMP_CLAUSE_MOTION_PRESENT (clause)) pp_string (pp, "present:"); + if (OMP_CLAUSE_ITERATORS (clause)) + { + dump_omp_iterators (pp, OMP_CLAUSE_ITERATORS (clause), spc, flags); + pp_colon (pp); + } dump_generic_node (pp, OMP_CLAUSE_DECL (clause), spc, flags, false); goto print_clause_size; diff --git a/gcc/tree.cc b/gcc/tree.cc index f12d7b8bb8a..5bba8df4518 100644 --- a/gcc/tree.cc +++ b/gcc/tree.cc @@ -264,8 +264,8 @@ unsigned const char omp_clause_num_ops[] = 1, /* OMP_CLAUSE_IS_DEVICE_PTR */ 1, /* OMP_CLAUSE_INCLUSIVE */ 1, /* OMP_CLAUSE_EXCLUSIVE */ - 2, /* OMP_CLAUSE_FROM */ - 2, /* OMP_CLAUSE_TO */ + 3, /* OMP_CLAUSE_FROM */ + 3, /* OMP_CLAUSE_TO */ 3, /* OMP_CLAUSE_MAP */ 1, /* OMP_CLAUSE_HAS_DEVICE_ADDR */ 1, /* OMP_CLAUSE_DOACROSS */ diff --git a/gcc/tree.h b/gcc/tree.h index 384a5f1f250..a57419447a1 100644 --- a/gcc/tree.h +++ b/gcc/tree.h @@ -1619,11 +1619,13 @@ class auto_suppress_location_wrappers != UNKNOWN_LOCATION) #define OMP_CLAUSE_LOCATION(NODE) (OMP_CLAUSE_CHECK (NODE))->omp_clause.locus -#define OMP_CLAUSE_HAS_ITERATORS(NODE) \ - (OMP_CLAUSE_CODE (NODE) == OMP_CLAUSE_MAP) +#define OMP_CLAUSE_HAS_ITERATORS(NODE) \ + (OMP_CLAUSE_CODE (NODE) == OMP_CLAUSE_FROM \ + || OMP_CLAUSE_CODE (NODE) == OMP_CLAUSE_TO \ + || OMP_CLAUSE_CODE (NODE) == OMP_CLAUSE_MAP) #define OMP_CLAUSE_ITERATORS(NODE) \ OMP_CLAUSE_OPERAND (OMP_CLAUSE_RANGE_CHECK (OMP_CLAUSE_CHECK (NODE), \ - OMP_CLAUSE_MAP, \ + OMP_CLAUSE_FROM, \ OMP_CLAUSE_MAP), 2) /* True on OMP_FOR and other OpenMP/OpenACC looping constructs if the loop nest diff --git a/libgomp/target.c b/libgomp/target.c index 463a162879b..60d57a19dd0 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -2260,6 +2260,8 @@ gomp_update (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs, size_t i; struct splay_tree_key_s cur_node; const int typemask = short_mapkind ? 0xff : 0x7; + bool iterators_p = false; + size_t *iterator_count = NULL; if (!devicep) return; @@ -2267,6 +2269,10 @@ gomp_update (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs, if (mapnum == 0) return; + if (short_mapkind) + iterators_p = gomp_merge_iterator_maps (&mapnum, &hostaddrs, &sizes, + &kinds, &iterator_count); + gomp_mutex_lock (&devicep->lock); if (devicep->state == GOMP_DEVICE_FINALIZED) { @@ -2360,6 +2366,14 @@ gomp_update (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs, } } gomp_mutex_unlock (&devicep->lock); + + if (iterators_p) + { + free (hostaddrs); + free (sizes); + free (kinds); + free (iterator_count); + } } static struct gomp_offload_icv_list * diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-1.c b/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-1.c new file mode 100644 index 00000000000..5a4cad5c219 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-1.c @@ -0,0 +1,65 @@ +/* { dg-do run } */ + +/* Test target enter data and target update to the target using map + iterators. */ + +#include + +#define DIM1 8 +#define DIM2 15 + +int mkarray (int *x[]) +{ + int expected = 0; + for (int i = 0; i < DIM1; i++) + { + x[i] = (int *) malloc (DIM2 * sizeof (int)); + for (int j = 0; j < DIM2; j++) + { + x[i][j] = rand (); + expected += x[i][j]; + } + } + + return expected; +} + +int main (void) +{ + int *x[DIM1]; + int sum; + int expected = mkarray (x); + + #pragma omp target enter data map(to: x[:DIM1]) + #pragma omp target enter data map(iterator(i=0:DIM1), to: x[i][:DIM2]) + #pragma omp target map(from: sum) + { + sum = 0; + for (int i = 0; i < DIM1; i++) + for (int j = 0; j < DIM2; j++) + sum += x[i][j]; + } + + if (sum != expected) + return 1; + + expected = 0; + for (int i = 0; i < DIM1; i++) + for (int j = 0; j < DIM2; j++) + { + x[i][j] *= rand (); + expected += x[i][j]; + } + + #pragma omp target update to(iterator(i=0:DIM1): x[i][:DIM2]) + + #pragma omp target map(from: sum) + { + sum = 0; + for (int i = 0; i < DIM1; i++) + for (int j = 0; j < DIM2; j++) + sum += x[i][j]; + } + + return sum != expected; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-2.c b/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-2.c new file mode 100644 index 00000000000..93438d01c97 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-2.c @@ -0,0 +1,58 @@ +/* { dg-do run } */ +/* { dg-require-effective-target offload_device_nonshared_as } */ + +/* Test target enter data and target update from the target using map + iterators. */ + +#include + +#define DIM1 8 +#define DIM2 15 + +void mkarray (int *x[]) +{ + for (int i = 0; i < DIM1; i++) + { + x[i] = (int *) malloc (DIM2 * sizeof (int)); + for (int j = 0; j < DIM2; j++) + x[i][j] = 0; + } +} + +int main (void) +{ + int *x[DIM1]; + int sum, expected; + + mkarray (x); + + #pragma omp target enter data map(alloc: x[:DIM1]) + #pragma omp target enter data map(iterator(i=0:DIM1), to: x[i][:DIM2]) + #pragma omp target map(from: expected) + { + expected = 0; + for (int i = 0; i < DIM1; i++) + for (int j = 0; j < DIM2; j++) + { + x[i][j] = (i + 1) * (j + 2); + expected += x[i][j]; + } + } + + /* Host copy of x should remain unchanged. */ + sum = 0; + for (int i = 0; i < DIM1; i++) + for (int j = 0; j < DIM2; j++) + sum += x[i][j]; + if (sum != 0) + return 1; + + #pragma omp target update from(iterator(i=0:DIM1): x[i][:DIM2]) + + /* Host copy should now be updated. */ + sum = 0; + for (int i = 0; i < DIM1; i++) + for (int j = 0; j < DIM2; j++) + sum += x[i][j]; + return sum - expected; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-3.c b/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-3.c new file mode 100644 index 00000000000..a70b21c4b75 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/target-update-iterators-3.c @@ -0,0 +1,67 @@ +/* { dg-do run } */ +/* { dg-require-effective-target offload_device_nonshared_as } */ + +/* Test target enter data and target update to the target using map + iterators with a function. */ + +#include + +#define DIM1 8 +#define DIM2 15 + +void mkarray (int *x[]) +{ + for (int i = 0; i < DIM1; i++) + { + x[i] = (int *) malloc (DIM2 * sizeof (int)); + for (int j = 0; j < DIM2; j++) + x[i][j] = rand (); + } +} + +int f (int i) +{ + return i * 2; +} + +int main (void) +{ + int *x[DIM1], x_new[DIM1][DIM2]; + int sum, expected; + + mkarray (x); + + #pragma omp target enter data map(alloc: x[:DIM1]) + #pragma omp target enter data map(iterator(i=0:DIM1), to: x[i][:DIM2]) + + /* Update x on host. */ + for (int i = 0; i < DIM1; i++) + for (int j = 0; j < DIM2; j++) + { + x_new[i][j] = x[i][j]; + x[i][j] = (i + 1) * (j + 2); + } + + /* Update a subset of x on target. */ + #pragma omp target update to(iterator(i=0:DIM1/2): x[f (i)][:DIM2]) + + #pragma omp target map(from: sum) + { + sum = 0; + for (int i = 0; i < DIM1; i++) + for (int j = 0; j < DIM2; j++) + sum += x[i][j]; + } + + /* Calculate expected value on host. */ + for (int i = 0; i < DIM1/2; i++) + for (int j = 0; j < DIM2; j++) + x_new[f (i)][j] = x[f (i)][j]; + + expected = 0; + for (int i = 0; i < DIM1; i++) + for (int j = 0; j < DIM2; j++) + expected += x_new[i][j]; + + return sum - expected; +} From patchwork Fri Oct 4 14:58:23 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: 1992782 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=RjIQh8sF; 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 4XKsC559kbz1xv2 for ; Sat, 5 Oct 2024 00:59:29 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id BC880386F471 for ; Fri, 4 Oct 2024 14:59:27 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from mail-ej1-x62a.google.com (mail-ej1-x62a.google.com [IPv6:2a00:1450:4864:20::62a]) by sourceware.org (Postfix) with ESMTPS id 55A62386D620 for ; Fri, 4 Oct 2024 14:58:52 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 55A62386D620 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 55A62386D620 Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=2a00:1450:4864:20::62a ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1728053937; cv=none; b=sGx7IYaImeG15rE/wHKr6k6R29cQnA+pAYGJwixFUJh5CgDlpcuvjO3fFLg0DQz40mX/L++IpJTMaGPJU089HkmvlUbyos7RIE19qKMgbM5SCjzi0NKeG7KrO03YlBFYmEyNMJ3wi1WmFw3J+QOiyaaPNvJ6c7GbO37ZQ7N0C7U= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1728053937; c=relaxed/simple; bh=7uwNx11SgGpTtfVPeYUW07/4eeGE3GEBkKA/70tmU2A=; h=DKIM-Signature:Message-ID:Date:MIME-Version:To:Subject:From; b=t7EdLnyK+oAHwZ9uu7ciTeRu8F/EG1PmSAiJ2oHNllOwKMlAK1Mv0dVo9WPICVK1c25GGtd1tNUTmcjUSiYQattU6vw8P6gC4qQL6jah3B4/3dmFMBbrk/YA0Qq/esC7mvxGJ4Otav44fkzDyWU9dad7fBzJBuIzTqP6u/kRNOs= ARC-Authentication-Results: i=1; server2.sourceware.org Received: by mail-ej1-x62a.google.com with SMTP id a640c23a62f3a-a8a6d1766a7so335057666b.3 for ; Fri, 04 Oct 2024 07:58:52 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre-com.20230601.gappssmtp.com; s=20230601; t=1728053931; x=1728658731; 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=Jl03BM/sW+Gj646eByPO/EWuXr6PtdhMWa/cduvFhVQ=; b=RjIQh8sFsiH7fmNPFaPCgJDzBJVBDGSWjNLOpoUWnynMeoEWUW1O1slOBaa01wSg0u m3/91uglxa4VoG74sTLkXeSAecmxjQcZ27NP/W8BQ41m+6beuCsBu5ZNLzY1Wu4Tlen5 zysnejMKoZWDf+B1vFaZkYZPp1olTd4PKT2F4kKMZpL0CkiHSFvo8P9+lvg0A0G9nyNc /Ac+s+CQrOKi+a9srz5jnVt7Jo0sK1RXjZjs3bTAjxtbtCPdFS/6aVqGnF+IIoKsk4DV Iiq1MMhCjGUb/KRr75BzRAQ69Y5JvHbBYXWzYypaaH5HaOpKkPo5U4PTQOhoxWr53DxD H5dQ== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1728053931; x=1728658731; 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=Jl03BM/sW+Gj646eByPO/EWuXr6PtdhMWa/cduvFhVQ=; b=olXgS6g48gKMJSX/8Mr055sKnzI7ELQbaxi3VImDpg1MWTV77j4pEHmII+GBKUOHV8 oBJz9i6nh02HFZTH9HYKUGtaKVLSaxGw5ljrZdQJaCxY9C2owRgmAF2Rh2uB2Z7zhOBI P4PjXJZJ1Pgdmt7V/P+68QT3m4+o0vzrxid+1jA6w4GZtbxqQPbFgkTPeZnRYWtgp6lP R1y+B44SXv1LnhjEdmv+hX00uj4GXn0pyYf7Fvhb6Ha7zgOEKJPQvVfOqGHppxyAGbcl pFGEseRrMmz84yegxkHEcIjulrEjcPG/qHs86+naw5ZGRM6XH/p6Tu6xgMQmnAb8usqv GZFg== X-Gm-Message-State: AOJu0YxlqGyq/O7fmtAqqfc1GFzLr1hO9JgEoMT8uL+0V0Jc4DIJmn1A 9hNQH+a2Vlekct1dIGJEyFa5BnzMv9IGNFPWnaP8kkiF7NQvrjIZ0tf+i9G3RWL23aefk8yqxMY B X-Google-Smtp-Source: AGHT+IF8bU55qx0EX7dOoRYNwhXeeAwwntdkTOIbbmQBgzolOauQBgwKz7psdiIGNM3QHkgtQ1nIRA== X-Received: by 2002:a17:906:730f:b0:a86:8917:fcd6 with SMTP id a640c23a62f3a-a991c022d25mr282140466b.60.1728053930488; Fri, 04 Oct 2024 07:58:50 -0700 (PDT) Received: from ?IPV6:2a00:23c6:88fe:9301:2d7d:f734:bc6:c47b? ([2a00:23c6:88fe:9301:2d7d:f734:bc6:c47b]) by smtp.gmail.com with ESMTPSA id a640c23a62f3a-a992e784a84sm1310966b.117.2024.10.04.07.58.49 (version=TLS1_3 cipher=TLS_AES_128_GCM_SHA256 bits=128/128); Fri, 04 Oct 2024 07:58:50 -0700 (PDT) Message-ID: <300ddbe2-4ad9-49ef-9f24-a02649a64bb9@baylibre.com> Date: Fri, 4 Oct 2024 15:58:23 +0100 MIME-Version: 1.0 User-Agent: Mozilla Thunderbird To: gcc-patches , Jakub Jelinek , Tobias Burnus , fortran@gcc.gnu.org References: <6b94b8ed-020b-47e2-b02a-4891891f2847@baylibre.com> Subject: [PATCH v3 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: <6b94b8ed-020b-47e2-b02a-4891891f2847@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 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 way as for the affinity and depend clauses, except for putting the iterator into the OMP_CLAUSE_ITERATOR of the clause. 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.). The presence of variables in the field offset triggers the unwanted creation of GOMP_MAP_STRUCT_UNORD for variable offsets. The offset tree is now walked over and if it only contains iterator variables, then the offset is treated as constant again (which it is, within the context of each iteration of the iterator). From a24aa032c2e23577d4fbc61df6da79345bae8292 Mon Sep 17 00:00:00 2001 From: Kwok Cheung Yeung Date: Fri, 4 Oct 2024 15:16:29 +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-10-04 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. Add expressions to iter_block rather than block. 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. (contains_only_iterator_vars_1): New. (contains_only_iterator_vars): New. (extract_base_bit_offset): Add iterator argument. Do not set variable_offset if contains_only_iterator_vars is true. (omp_accumulate_sibling_list): Add iterator argument to extract_base_bit_offset. * omp-low.cc (lower_omp_target): Add sorry if iterators used with deep mapping. * tree-pretty-print.cc (dump_block_node): Ignore BLOCK_SUBBLOCKS containing iterator block statements. gcc/testsuite/ * gfortran.dg/gomp/target-map-iterators-1.f90: New. * gfortran.dg/gomp/target-map-iterators-2.f90: New. * gfortran.dg/gomp/target-map-iterators-3.f90: New. libgomp/ * target.c (kind_to_name): Handle GOMP_MAP_STRUCT and GOMP_MAP_STRUCT_UNORD. (gomp_add_map): New. (gomp_merge_iterator_maps): Expand fields of a struct mapping breadth-first. * 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 | 71 ++++++++++++---- gcc/gimplify.cc | 76 ++++++++++++++--- gcc/omp-low.cc | 5 ++ .../gomp/target-map-iterators-1.f90 | 26 ++++++ .../gomp/target-map-iterators-2.f90 | 27 ++++++ .../gomp/target-map-iterators-3.f90 | 24 ++++++ gcc/tree-pretty-print.cc | 4 +- libgomp/target.c | 83 ++++++++++++++----- .../target-map-iterators-1.f90 | 45 ++++++++++ .../target-map-iterators-2.f90 | 45 ++++++++++ .../target-map-iterators-3.f90 | 57 +++++++++++++ 13 files changed, 452 insertions(+), 55 deletions(-) create mode 100644 gcc/testsuite/gfortran.dg/gomp/target-map-iterators-1.f90 create mode 100644 gcc/testsuite/gfortran.dg/gomp/target-map-iterators-2.f90 create mode 100644 gcc/testsuite/gfortran.dg/gomp/target-map-iterators-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 3547d7f8aca..3ee6ed1ea7f 100644 --- a/gcc/fortran/dump-parse-tree.cc +++ b/gcc/fortran/dump-parse-tree.cc @@ -1359,7 +1359,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) @@ -1371,8 +1372,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 2d5c4305d2a..3003ba605cf 100644 --- a/gcc/fortran/openmp.cc +++ b/gcc/fortran/openmp.cc @@ -193,7 +193,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, i == OMP_LIST_INIT); @@ -3477,9 +3478,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 (;;) { @@ -3499,6 +3503,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 (", "); @@ -3555,15 +3564,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; @@ -8856,7 +8880,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 3a335ade0f7..c154975fb0b 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,39 @@ 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_ITERATORS (c) = iterator; + 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 +3360,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 +3420,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 +3449,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 +3538,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 +3567,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 +3584,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 +3619,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 +3638,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 +3657,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 +3874,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 +4012,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 +4025,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 +4037,15 @@ 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_ITERATORS (c) = iterator; + } break; case OMP_LIST_TO: case OMP_LIST_FROM: diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index ba972a2892a..4e30d335324 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -8858,10 +8858,17 @@ compute_iterator_count (tree it, 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); @@ -8905,6 +8912,7 @@ build_iterator_loop (tree it, gimple_seq *pre_p, tree *last_bind) if (*last_bind) gimplify_and_add (*last_bind, pre_p); tree block = TREE_VEC_ELT (it, 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; @@ -8916,6 +8924,7 @@ build_iterator_loop (tree it, 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: @@ -8926,9 +8935,9 @@ build_iterator_loop (tree it, 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. */ @@ -8954,10 +8963,12 @@ build_iterator_loop (tree it, 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); @@ -8966,6 +8977,11 @@ build_iterator_loop (tree it, 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; @@ -9406,6 +9422,34 @@ build_omp_struct_comp_nodes (enum tree_code code, tree grp_start, tree grp_end, return c2; } +/* Callback for walk_tree. Return any VAR_DECLS that are not found in the + iterators stored in DATA. */ + +static tree +contains_only_iterator_vars_1 (tree* tp, int *, void *data) +{ + tree iterators = (tree) data; + tree t = *tp; + + if (TREE_CODE (t) != VAR_DECL) + return NULL_TREE; + + for (tree it = iterators; it; it = TREE_CHAIN (it)) + if (t == TREE_VEC_ELT (it, 0)) + return NULL_TREE; + + return t; +} + +/* Return true if the only variables present in EXPR are iterator variables in + ITERATORS. */ + +static bool +contains_only_iterator_vars (tree expr, tree iterators) +{ + return !walk_tree (&expr, contains_only_iterator_vars_1, iterators, NULL); +} + /* Strip ARRAY_REFS or an indirect ref off BASE, find the containing object, and set *BITPOSP and *POFFSETP to the bit offset of the access. If BASE_REF is non-NULL and the containing object is a reference, set @@ -9416,7 +9460,8 @@ build_omp_struct_comp_nodes (enum tree_code code, tree grp_start, tree grp_end, static tree extract_base_bit_offset (tree base, poly_int64 *bitposp, poly_offset_int *poffsetp, - bool *variable_offset) + bool *variable_offset, + tree iterator) { tree offset; poly_int64 bitsize, bitpos; @@ -9440,6 +9485,8 @@ extract_base_bit_offset (tree base, poly_int64 *bitposp, { poffset = 0; *variable_offset = (offset != NULL_TREE); + if (iterator && *variable_offset) + *variable_offset = !contains_only_iterator_vars (offset, iterator); } if (maybe_ne (bitpos, 0)) @@ -11245,8 +11292,11 @@ omp_accumulate_sibling_list (enum omp_region_type region_type, } bool variable_offset; + tree iterators = OMP_CLAUSE_HAS_ITERATORS (grp_end) + ? OMP_CLAUSE_ITERATORS (grp_end) : NULL_TREE; tree base - = extract_base_bit_offset (ocd, &cbitpos, &coffset, &variable_offset); + = extract_base_bit_offset (ocd, &cbitpos, &coffset, &variable_offset, + iterators); int base_token; for (base_token = addr_tokens.length () - 1; base_token >= 0; base_token--) @@ -11579,8 +11629,12 @@ omp_accumulate_sibling_list (enum omp_region_type region_type, sc_decl = TREE_OPERAND (sc_decl, 0); bool variable_offset2; + tree iterators2 = OMP_CLAUSE_HAS_ITERATORS (*sc) + ? OMP_CLAUSE_ITERATORS (*sc) : NULL_TREE; + tree base2 = extract_base_bit_offset (sc_decl, &bitpos, &offset, - &variable_offset2); + &variable_offset2, + iterators2); if (!base2 || !operand_equal_p (base2, base, 0)) break; if (scp) diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc index a8b86889c66..46f40a14646 100644 --- a/gcc/omp-low.cc +++ b/gcc/omp-low.cc @@ -12938,6 +12938,11 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) deep_map_cnt = extra; } + if (deep_map_cnt + && OMP_CLAUSE_HAS_ITERATORS (c) && OMP_CLAUSE_ITERATORS (c)) + sorry ("iterators used together with deep mapping are not " + "supported yet"); + if (!DECL_P (var)) { if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP diff --git a/gcc/testsuite/gfortran.dg/gomp/target-map-iterators-1.f90 b/gcc/testsuite/gfortran.dg/gomp/target-map-iterators-1.f90 new file mode 100644 index 00000000000..25abbaf741e --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/target-map-iterators-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-map-iterators-2.f90 b/gcc/testsuite/gfortran.dg/gomp/target-map-iterators-2.f90 new file mode 100644 index 00000000000..b7d7501cf63 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/target-map-iterators-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-map-iterators-3.f90 b/gcc/testsuite/gfortran.dg/gomp/target-map-iterators-3.f90 new file mode 100644 index 00000000000..785f149c0d8 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/target-map-iterators-3.f90 @@ -0,0 +1,24 @@ +! { dg-do compile } +! { dg-options "-fopenmp -fdump-tree-omplower" } + +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 "omplower" } } +! { dg-final { scan-tree-dump-times "if \\(i <= 27\\) goto ; else goto ;" 1 "omplower" } } +! { dg-final { scan-tree-dump-times "map\\(iterator\\(integer\\(kind=4\\) i=1:17:1\\):to:D\.\[0-9\]+" 1 "omplower" } } +! { dg-final { scan-tree-dump-times "map\\(iterator\\(integer\\(kind=4\\) i=1:27:1\\):from:D\.\[0-9\]+" 1 "omplower" } } +! { dg-final { scan-tree-dump-times "map\\(iterator\\(integer\\(kind=4\\) i=1:17:1\\):attach:D\.\[0-9\]+" 1 "omplower" } } +! { dg-final { scan-tree-dump-times "map\\(iterator\\(integer\\(kind=4\\) i=1:27:1\\):attach:D\.\[0-9\]+" 1 "omplower" } } diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc index fa1b2dce27f..da6b757e212 100644 --- a/gcc/tree-pretty-print.cc +++ b/gcc/tree-pretty-print.cc @@ -1671,7 +1671,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 60d57a19dd0..e8205f6c309 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -993,10 +993,48 @@ kind_to_name (unsigned short kind) case GOMP_MAP_POINTER: return "GOMP_MAP_POINTER"; 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 %u <%s>: " + "hostaddrs[%u] = %p, sizes[%u] = %lu\n", + (int) idx, kind_to_name ((*new_kinds)[*new_idx]), + (int) *new_idx, (*new_hostaddrs)[*new_idx], + (int) *new_idx, (unsigned long) (*new_sizes)[*new_idx]); + (*new_idx)++; + } + } + else + { + (*new_hostaddrs)[*new_idx] = (*hostaddrs)[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. ITERATOR_COUNT holds the iteration count of the @@ -1037,33 +1075,34 @@ gomp_merge_iterator_maps (size_t *mapnum, void ***hostaddrs, size_t **sizes, 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++; - for (size_t j = 0; j < count; j++) + 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++; - new_sizes[new_idx] = *iterator_array++; - new_kinds[new_idx] = (*skinds)[i]; - (*iterator_count)[new_idx] = j + 1; - gomp_debug (1, - "Expanding map %u <%s>: " - "hostaddrs[%u] = %p, sizes[%u] = %lu\n", - (int) i, kind_to_name (new_kinds[new_idx]), - (int) new_idx, new_hostaddrs[new_idx], - (int) new_idx, (unsigned long) new_sizes[new_idx]); - new_idx++; + 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 %u: new field count = %lu\n", + (int) i, (unsigned long) 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]; - (*iterator_count)[new_idx] = 0; - new_idx++; - } + gomp_add_map (i, &new_idx, hostaddrs, sizes, skinds, + &new_hostaddrs, &new_sizes, &new_kinds, *iterator_count); } *mapnum = map_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 Fri Oct 4 14:59:13 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: 1992783 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=aH5ax+px; 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 4XKsDH4Qcwz1xtH for ; Sat, 5 Oct 2024 01:00:31 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id D453B38425B9 for ; Fri, 4 Oct 2024 15:00:28 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from mail-lj1-x234.google.com (mail-lj1-x234.google.com [IPv6:2a00:1450:4864:20::234]) by sourceware.org (Postfix) with ESMTPS id 974A8386F473 for ; Fri, 4 Oct 2024 14:59:41 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 974A8386F473 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 974A8386F473 Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=2a00:1450:4864:20::234 ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1728053984; cv=none; b=nRO1raIw9eSA7MjWOl7DasUvILerKbLNilkxbtaBJ8627cPbPpZ9jMEn/Jy7Kti3eH4tGRhjCye+bmFLT54K+i0fy2NGQaUjkZI/TcKFBJvvbbFVpGep0DT1ySCxpuFr28L8q7G0kQ/c3AbBh7E0b6ZKM/LGyDPHo6l4e2SD1ks= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1728053984; c=relaxed/simple; bh=KIpe9HUSgyQS7JQ1B29lgzm/XFiCh6/q/vbDpQ9oOuY=; h=DKIM-Signature:Message-ID:Date:MIME-Version:To:Subject:From; b=apGk7JRPkFsa+WQQ3O06VWdyKaRq92Dc0XC9QV47HZFS77hd1xPvoeD49ydD/E415+z8s/t4dd8F4rEKm2XSdOnCGAT5KwAa+NUmtR61DtC57/WfBCITFemJM/B9q2V8XXBvcqgn2aBac/lMWqb4n+0bAGclgA1HshQer/AwCzg= ARC-Authentication-Results: i=1; server2.sourceware.org Received: by mail-lj1-x234.google.com with SMTP id 38308e7fff4ca-2f75c56f16aso24114771fa.0 for ; Fri, 04 Oct 2024 07:59:41 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=baylibre-com.20230601.gappssmtp.com; s=20230601; t=1728053980; x=1728658780; 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=S/fLFKnnfk5FjznT0yuhRNrBrzI3f7gFs94r85zNedU=; b=aH5ax+pxrPN0pr7fj2B/OsXGrRkMDG20ofjCdtn4YtyfDTnBuqKh9fxOgMNzy4Wcjg 3TsaAx/DPi9u1cvlD1Gg9OA6N57dj/MKUH5sdnADCHwpbce2cPssZFQpe/raBR5I1xeQ vVCCd58wPuROnfavnQjy/6nxo1iAtkSjHag1UFqpr2I0Cc2iquRFntEEDVmzC/c34Fh4 jAlojeF0rEf6bcaMU/sOfFHj9eIu8ongI/145jbTw+cYwdvXxPirAZCJkjzsJ0D0x7xk O/wccLe6PmfZqWhWBWz466FhzRuXyStTTb6PKrKxCTs1+gqUTVmPckUj+uSnNYu3A4Cu TQkw== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1728053980; x=1728658780; 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=S/fLFKnnfk5FjznT0yuhRNrBrzI3f7gFs94r85zNedU=; b=TbzR30Vo4AQqQCnECseWhBMdh3XWmGjApHb34PrXvjeYx+rH6vsTAbcTBLRPFyyfVl fSrNeHw7eDLFaJCOIy4TNhBAWaQ5f35G6rYbfugRsJfBGrYfUAs3uaFDgEYppmsffpIp ltDmavd/x0OFIq8MToRiAq+yIMNhJP3p3gKEUM26txwnVDJytgKQ3lz4tuNNKsKlP0w4 7mU1H9T9/tzRqB19tYj5M/XGgE2iemzqR9F99hkgEDxOpHp2tlPxEJlPHIoeujZ0KEwI LbqhcN6HyTtPEZtyTJiGKxzhx2aZl9rkL18nW3F/rfWPcqDJqOkxlSLTZ8U0vQDgUNWt g7eg== X-Gm-Message-State: AOJu0YyIrQPZsaskFCbu75Pa3J6Rg9nYJpgKEl1Qk2nh0qLLUqu3c8yO xdNsVV27Np78vfm0iPrcmM7S6aW1lFP3YMEy+VhKr+qNanpE6jNVTP2iOPB7u9RYXrO0JkHcDG1 0 X-Google-Smtp-Source: AGHT+IFaO0lCZSsc0o+EZ5sAQ/EhdyCPC3mWMNwS+k2+gafFU0ptzPiPHBLxnW6SI2pV3UhIeKsMiA== X-Received: by 2002:a05:6512:3b85:b0:536:54df:bffc with SMTP id 2adb3069b0e04-539ab9dc722mr1687814e87.42.1728053979738; Fri, 04 Oct 2024 07:59:39 -0700 (PDT) Received: from ?IPV6:2a00:23c6:88fe:9301:2d7d:f734:bc6:c47b? ([2a00:23c6:88fe:9301:2d7d:f734:bc6:c47b]) by smtp.gmail.com with ESMTPSA id a640c23a62f3a-a992e5d0b41sm2280966b.9.2024.10.04.07.59.39 (version=TLS1_3 cipher=TLS_AES_128_GCM_SHA256 bits=128/128); Fri, 04 Oct 2024 07:59:39 -0700 (PDT) Message-ID: Date: Fri, 4 Oct 2024 15:59:13 +0100 MIME-Version: 1.0 User-Agent: Mozilla Thunderbird To: gcc-patches , fortran@gcc.gnu.org, Jakub Jelinek , Tobias Burnus References: <6b94b8ed-020b-47e2-b02a-4891891f2847@baylibre.com> Subject: [PATCH v3 5/5] openmp, fortran: Add support for iterators in OpenMP 'target update' constructs (Fortran) Content-Language: en-GB From: Kwok Cheung Yeung In-Reply-To: <6b94b8ed-020b-47e2-b02a-4891891f2847@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 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 da8ab0cb38d2bc347cf902ec417b0397c28e24e2 Mon Sep 17 00:00:00 2001 From: Kwok Cheung Yeung Date: Fri, 4 Oct 2024 15:16:38 +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-10-04 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. Add expressions to iter_block rather than block. gcc/testsuite/ * gfortran.dg/gomp/target-update-iterators-1.f90: New. * gfortran.dg/gomp/target-update-iterators-2.f90: New. * gfortran.dg/gomp/target-update-iterators-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 | 50 ++++++++++-- .../gomp/target-update-iterators-1.f90 | 25 ++++++ .../gomp/target-update-iterators-2.f90 | 22 ++++++ .../gomp/target-update-iterators-3.f90 | 23 ++++++ .../target-update-iterators-1.f90 | 68 ++++++++++++++++ .../target-update-iterators-2.f90 | 63 +++++++++++++++ .../target-update-iterators-3.f90 | 78 +++++++++++++++++++ 9 files changed, 386 insertions(+), 12 deletions(-) create mode 100644 gcc/testsuite/gfortran.dg/gomp/target-update-iterators-1.f90 create mode 100644 gcc/testsuite/gfortran.dg/gomp/target-update-iterators-2.f90 create mode 100644 gcc/testsuite/gfortran.dg/gomp/target-update-iterators-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 3ee6ed1ea7f..0a2d546d3fe 100644 --- a/gcc/fortran/dump-parse-tree.cc +++ b/gcc/fortran/dump-parse-tree.cc @@ -1360,7 +1360,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) @@ -1376,6 +1377,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 3003ba605cf..c765d5814a7 100644 --- a/gcc/fortran/openmp.cc +++ b/gcc/fortran/openmp.cc @@ -194,7 +194,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, i == OMP_LIST_INIT); @@ -1368,17 +1369,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; } @@ -8881,7 +8930,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 c154975fb0b..c83445d5885 100644 --- a/gcc/fortran/trans-openmp.cc +++ b/gcc/fortran/trans-openmp.cc @@ -4050,11 +4050,39 @@ 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_ITERATORS (c) = iterator; + 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: @@ -4092,7 +4120,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)); @@ -4117,7 +4145,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))); } @@ -4126,9 +4154,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)); @@ -4137,7 +4165,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); } @@ -4145,8 +4173,20 @@ 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_ITERATORS (c) = iterator; + } 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-iterators-1.f90 b/gcc/testsuite/gfortran.dg/gomp/target-update-iterators-1.f90 new file mode 100644 index 00000000000..08dc3d79911 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/target-update-iterators-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-iterators-2.f90 b/gcc/testsuite/gfortran.dg/gomp/target-update-iterators-2.f90 new file mode 100644 index 00000000000..89f645bda23 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/target-update-iterators-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-iterators-3.f90 b/gcc/testsuite/gfortran.dg/gomp/target-update-iterators-3.f90 new file mode 100644 index 00000000000..753811384ae --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/target-update-iterators-3.f90 @@ -0,0 +1,23 @@ +! { dg-do compile } +! { dg-options "-fopenmp -fdump-tree-omplower" } + +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 "omplower" } } +! { dg-final { scan-tree-dump-times "if \\(j <= 39\\) goto ; else goto ;" 1 "omplower" } } +! { dg-final { scan-tree-dump-times "to\\(iterator\\(integer\\(kind=4\\) j=1:39:1, integer\\(kind=4\\) i=1:17:1\\):D\.\[0-9\]+" 2 "omplower" } } +! { dg-final { scan-tree-dump-times "from\\(iterator\\(integer\\(kind=4\\) i=1:17:1\\):D\.\[0-9\]+" 1 "omplower" } } 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..2e982bc032c --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/target-update-iterators-2.f90 @@ -0,0 +1,63 @@ +! { dg-do run } +! { dg-require-effective-target offload_device_nonshared_as } + +! 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..54b2a6c37c1 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/target-update-iterators-3.f90 @@ -0,0 +1,78 @@ +! { dg-do run } +! { dg-require-effective-target offload_device_nonshared_as } + +! Test target enter data and target update to the target using map +! iterators with a function. + +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