{"id":810188,"url":"http://patchwork.ozlabs.org/api/1.2/patches/810188/?format=json","web_url":"http://patchwork.ozlabs.org/project/gcc/patch/4103cee4-003f-bfb1-a811-fd321a8669ad@mentor.com/","project":{"id":17,"url":"http://patchwork.ozlabs.org/api/1.2/projects/17/?format=json","name":"GNU Compiler Collection","link_name":"gcc","list_id":"gcc-patches.gcc.gnu.org","list_email":"gcc-patches@gcc.gnu.org","web_url":null,"scm_url":null,"webscm_url":null,"list_archive_url":"","list_archive_url_format":"","commit_url_format":""},"msgid":"<4103cee4-003f-bfb1-a811-fd321a8669ad@mentor.com>","list_archive_url":null,"date":"2017-09-05T14:32:22","name":"[openacc,og7,committed] Make reduction copy clauses 'private'","commit_ref":null,"pull_url":null,"state":"new","archived":false,"hash":"e791356fa5628dca56b287c6f00f887a774ea448","submitter":{"id":63304,"url":"http://patchwork.ozlabs.org/api/1.2/people/63304/?format=json","name":"Chung-Lin Tang","email":"chunglin_tang@mentor.com"},"delegate":null,"mbox":"http://patchwork.ozlabs.org/project/gcc/patch/4103cee4-003f-bfb1-a811-fd321a8669ad@mentor.com/mbox/","series":[{"id":1591,"url":"http://patchwork.ozlabs.org/api/1.2/series/1591/?format=json","web_url":"http://patchwork.ozlabs.org/project/gcc/list/?series=1591","date":"2017-09-05T14:32:22","name":"[openacc,og7,committed] Make reduction copy clauses 'private'","version":1,"mbox":"http://patchwork.ozlabs.org/series/1591/mbox/"}],"comments":"http://patchwork.ozlabs.org/api/patches/810188/comments/","check":"pending","checks":"http://patchwork.ozlabs.org/api/patches/810188/checks/","tags":{},"related":[],"headers":{"Return-Path":"","X-Original-To":"incoming@patchwork.ozlabs.org","Delivered-To":["patchwork-incoming@bilbo.ozlabs.org","mailing list gcc-patches@gcc.gnu.org"],"Authentication-Results":["ozlabs.org;\n\tspf=pass (mailfrom) smtp.mailfrom=gcc.gnu.org\n\t(client-ip=209.132.180.131; helo=sourceware.org;\n\tenvelope-from=gcc-patches-return-461512-incoming=patchwork.ozlabs.org@gcc.gnu.org;\n\treceiver=)","ozlabs.org; dkim=pass (1024-bit key;\n\tunprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org\n\theader.b=\"wpizXtKo\"; dkim-atps=neutral","sourceware.org; auth=none"],"Received":["from sourceware.org (server1.sourceware.org [209.132.180.131])\n\t(using TLSv1.2 with cipher ECDHE-RSA-AES256-GCM-SHA384 (256/256\n\tbits)) (No client certificate requested)\n\tby ozlabs.org (Postfix) with ESMTPS id 3xmq220zJ2z9t2W\n\tfor ;\n\tWed, 6 Sep 2017 00:32:56 +1000 (AEST)","(qmail 72153 invoked by alias); 5 Sep 2017 14:32:46 -0000","(qmail 69234 invoked by uid 89); 5 Sep 2017 14:32:40 -0000","from relay1.mentorg.com (HELO relay1.mentorg.com) (192.94.38.131)\n\tby sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with\n\tESMTP; Tue, 05 Sep 2017 14:32:28 +0000","from svr-orw-mbx-06.mgc.mentorg.com ([147.34.90.206])\tby\n\trelay1.mentorg.com with esmtp id 1dpEtt-0003J1-F1 from\n\tChungLin_Tang@mentor.com for gcc-patches@gcc.gnu.org;\n\tTue, 05 Sep 2017 07:32:25 -0700","from svr-orw-mbx-08.mgc.mentorg.com (147.34.90.208) by\n\tSVR-ORW-MBX-06.mgc.mentorg.com (147.34.90.206) with Microsoft\n\tSMTP Server (TLS) id 15.0.1263.5; Tue, 5 Sep 2017 07:32:22 -0700","from [0.0.0.0] (147.34.91.1) by svr-orw-mbx-08.mgc.mentorg.com\n\t(147.34.90.208) with Microsoft SMTP Server (TLS) id\n\t15.0.1263.5 via Frontend Transport; Tue, 5 Sep 2017 07:32:21 -0700"],"DomainKey-Signature":"a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id\n\t:list-unsubscribe:list-archive:list-post:list-help:sender:from\n\t:subject:to:cc:message-id:date:mime-version:content-type; q=dns;\n\ts=default; b=XfYBNd/lMoabD/rCkolkP1Jno5jifne33G9uYHAV3wF1QTFzWu\n\txf48KZUaLTnAJSsAOHWKaa4G2q+0kWL3kL0d6ugn4Lxuku3b9yKQHHyYACl9jrKJ\n\tQSQL9YuwhafotrJEslStShbM+EesCFaekpuSDG5WzRp07JUFk0w6B9bLQ=","DKIM-Signature":"v=1; a=rsa-sha1; c=relaxed; d=gcc.gnu.org; h=list-id\n\t:list-unsubscribe:list-archive:list-post:list-help:sender:from\n\t:subject:to:cc:message-id:date:mime-version:content-type; s=\n\tdefault; bh=ZESHlsySfobc0k/BH+gabjVc580=; b=wpizXtKoJTI+GnFoTh45\n\tRb+WyuAAlciafqkAqb49rUZtoZ/xIPIdKS6TMLXivHfKQOB25oqqzzEfPowS34Yl\n\tKI3kY2O2CXhJHA2co5ERqFcShPsUarkDAtdNts2HCE+GotyImC8Ikhpig8OQh/HK\n\t/yH4DPie8WXwCjyrP820Tvg=","Mailing-List":"contact gcc-patches-help@gcc.gnu.org; run by ezmlm","Precedence":"bulk","List-Id":"","List-Unsubscribe":"","List-Archive":"","List-Post":"","List-Help":"","Sender":"gcc-patches-owner@gcc.gnu.org","X-Virus-Found":"No","X-Spam-SWARE-Status":"No, score=-24.6 required=5.0 tests=AWL, BAYES_00,\n\tGIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3,\n\tRCVD_IN_DNSWL_NONE, SPF_PASS,\n\tURIBL_RED autolearn=ham version=3.3.2 spammy=backs, ACC, gangs,\n\t88949","X-HELO":"relay1.mentorg.com","From":"Chung-Lin Tang ","Subject":"[PATCH, openacc, og7,\n\tcommitted] Make reduction copy clauses 'private'","To":"gcc-patches ","CC":"Cesar Philippidis ,\n\tThomas Schwinge\t","Message-ID":"<4103cee4-003f-bfb1-a811-fd321a8669ad@mentor.com>","Date":"Tue, 5 Sep 2017 22:32:22 +0800","User-Agent":"Mozilla/5.0 (Macintosh; Intel Mac OS X 10.11;\n\trv:52.0) Gecko/20100101 Thunderbird/52.3.0","MIME-Version":"1.0","Content-Type":"multipart/mixed;\n\tboundary=\"------------6E4B456EDF5987015451DFD2\""},"content":"As we discussed, we are to support a behavior where within individual gangs,\nworker/vector level reductions will correctly work with results immediately available.\nThis is on top of the implicit 'copy' clause added when we have loop reductions.\n\nThis patch adds a capability to mark map clauses additionally as 'private' (we may\nbe overloading this word a little too much :P), such that within offloaded regions\nand wrt to our reduction lowering, the variable is (first)private, with additional\ncopy back appended at end of the offloaded region.\n\nCare is taken to make sure this behavior is not applied when potential loop gang\nreductions may happen (which this will not work). In other cases, for gang-redundant\ncode, supposedly the multiple copy backs should all be the same, so the behavior\nis same.\n\nThis is sort of a refinement of the implicit copy clause for reductions in PR70895.\nA libgomp testcase is added to test the multiple worker-level reduction result case\nacross multiple gangs. Patch was tested and pushed to openacc-gcc-7-branch.\n\nChung-Lin\nFrom 2dc21f336368889c1ebf031801a7613f65899ef1 Mon Sep 17 00:00:00 2001\nFrom: Chung-Lin Tang \nDate: Tue, 5 Sep 2017 22:09:34 +0800\nSubject: [PATCH] Add support for making maps 'private' inside offloaded\n regions.\n\n2017-09-05 Chung-Lin Tang \n\n\tgcc/\n\t* tree.h (OMP_CLAUSE_MAP_PRIVATE): Define macro.\n\t* gimplify.c (enum gimplify_omp_var_data): Add GOVD_MAP_PRIVATE enum value.\n\t(omp_add_variable): Add GOVD_MAP_PRIVATE to reduction clause flags if\n\tnot a gang-partitioned loop directive.\n\t(gimplify_adjust_omp_clauses_1): Set OMP_CLAUSE_MAP_PRIVATE of new map\n\tclause to 1 if GOVD_MAP_PRIVATE flag is present.\n\t* omp-low.c (lower_oacc_reductions): Handle map clauses with\n\tOMP_CLAUSE_MAP_PRIVATE set in same matter as firstprivate/private.\n\t(lower_omp_target): Likewise. Add copy back code for map clauses with\n\tOMP_CLAUSE_MAP_PRIVATE set.\n\n\tlibgomp/\n\t* testsuite/libgomp.oacc-c-c++-common/reduction-9.c: New test.\n---\n gcc/ChangeLog.openacc | 14 ++++++++\n gcc/gimplify.c | 34 ++++++++++++++++--\n gcc/omp-low.c | 28 +++++++++++++--\n gcc/tree.h | 3 ++\n libgomp/ChangeLog.openacc | 4 +++\n .../libgomp.oacc-c-c++-common/reduction-9.c | 41 ++++++++++++++++++++++\n 6 files changed, 119 insertions(+), 5 deletions(-)\n create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-9.c","diff":"diff --git a/gcc/ChangeLog.openacc b/gcc/ChangeLog.openacc\nindex 4b1ce0b..23e19d9 100644\n--- a/gcc/ChangeLog.openacc\n+++ b/gcc/ChangeLog.openacc\n@@ -1,3 +1,17 @@\n+2017-09-05 Chung-Lin Tang \n+\n+\t* tree.h (OMP_CLAUSE_MAP_PRIVATE): Define macro.\n+\t* gimplify.c (enum gimplify_omp_var_data): Add GOVD_MAP_PRIVATE enum value.\n+\t(omp_add_variable): Add GOVD_MAP_PRIVATE to reduction clause flags if\n+\tnot a gang-partitioned loop directive.\n+\t(gimplify_adjust_omp_clauses_1): Set OMP_CLAUSE_MAP_PRIVATE of new map\n+\tclause to 1 if GOVD_MAP_PRIVATE flag is present.\n+\t* omp-low.c (lower_oacc_reductions): Handle map clauses with\n+\tOMP_CLAUSE_MAP_PRIVATE set in same matter as firstprivate/private.\n+\t(lower_omp_target): Likewise. Add copy back code for map clauses with\n+\tOMP_CLAUSE_MAP_PRIVATE set.\n+\t* tree.h (OMP_CLAUSE_MAP_PRIVATE): Define macro.\n+\n 2017-08-11 Cesar Philippidis \n \n \t* config/nvptx/nvptx.c (PTX_GANG_DEFAULT): Delete define.\ndiff --git a/gcc/gimplify.c b/gcc/gimplify.c\nindex e481a72..2c10c64 100644\n--- a/gcc/gimplify.c\n+++ b/gcc/gimplify.c\n@@ -102,6 +102,9 @@ enum gimplify_omp_var_data\n /* Flag for GOVD_MAP: must be present already. */\n GOVD_MAP_FORCE_PRESENT = 524288,\n \n+ /* Flag for GOVD_MAP, copy to/from private storage inside offloaded region. */\n+ GOVD_MAP_PRIVATE = 1048576,\n+\n GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE\n \t\t\t | GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR\n \t\t\t | GOVD_LOCAL)\n@@ -6717,6 +6720,21 @@ omp_add_variable (struct gimplify_omp_ctx *ctx, tree decl, unsigned int flags)\n if (ctx->region_type == ORT_ACC && (flags & GOVD_REDUCTION))\n {\n struct gimplify_omp_ctx *outer_ctx = ctx->outer_context;\n+\n+ bool gang = false, worker = false, vector = false;\n+ for (tree c = ctx->clauses; c; c = OMP_CLAUSE_CHAIN (c))\n+\t{\n+\t if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_GANG)\n+\t gang = true;\n+\t else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_WORKER)\n+\t worker = true;\n+\t else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_VECTOR)\n+\t vector = true;\n+\t}\n+\n+ /* Set new copy map as 'private' if sure we're not gang-partitioning. */\n+ bool map_private = !gang && (worker || vector);\n+\n while (outer_ctx)\n \t{\n \t n = splay_tree_lookup (outer_ctx->variables, (splay_tree_key)decl);\n@@ -6738,12 +6756,21 @@ omp_add_variable (struct gimplify_omp_ctx *ctx, tree decl, unsigned int flags)\n \t\t /* Remove firstprivate and make it a copy map. */\n \t\t n->value &= ~GOVD_FIRSTPRIVATE;\n \t\t n->value |= GOVD_MAP;\n+\n+\t\t /* If not gang-partitioned, add MAP_PRIVATE on the map\n+\t\t clause. */\n+\t\t if (map_private)\n+\t\t n->value |= GOVD_MAP_PRIVATE;\n \t\t}\n \t }\n \t else if (outer_ctx->region_type == ORT_ACC_PARALLEL)\n \t {\n-\t splay_tree_insert (outer_ctx->variables, (splay_tree_key)decl,\n-\t\t\t\t GOVD_MAP | GOVD_SEEN);\n+\t unsigned f = GOVD_MAP | GOVD_SEEN;\n+\n+\t /* If not gang-partitioned, add MAP_PRIVATE on the map clause. */\n+\t if (map_private)\n+\t\tf |= GOVD_MAP_PRIVATE;\n+\t splay_tree_insert (outer_ctx->variables, (splay_tree_key)decl, f);\n \t break;\n \t }\n \t outer_ctx = outer_ctx->outer_context;\n@@ -8867,6 +8894,9 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)\n \t gcc_unreachable ();\n \t}\n OMP_CLAUSE_SET_MAP_KIND (clause, kind);\n+ if ((flags & GOVD_MAP_PRIVATE)\n+\t && TREE_CODE (OMP_CLAUSE_DECL (clause)) == VAR_DECL)\n+\tOMP_CLAUSE_MAP_PRIVATE (clause) = 1;\n tree c2 = gomp_needs_data_present (decl);\n /* Handle OpenACC pointers that were declared inside acc data\n \t regions. */\ndiff --git a/gcc/omp-low.c b/gcc/omp-low.c\nindex f45c5c3..e790f0f 100644\n--- a/gcc/omp-low.c\n+++ b/gcc/omp-low.c\n@@ -5220,7 +5220,9 @@ lower_oacc_reductions (location_t loc, tree clauses, tree level, bool inner,\n \t\t goto has_outer_reduction;\n \t\t }\n \t\t else if ((OMP_CLAUSE_CODE (cls) == OMP_CLAUSE_FIRSTPRIVATE\n-\t\t\t || OMP_CLAUSE_CODE (cls) == OMP_CLAUSE_PRIVATE)\n+\t\t\t || OMP_CLAUSE_CODE (cls) == OMP_CLAUSE_PRIVATE\n+\t\t\t || (OMP_CLAUSE_CODE (cls) == OMP_CLAUSE_MAP\n+\t\t\t\t&& OMP_CLAUSE_MAP_PRIVATE (cls)))\n \t\t\t && orig == OMP_CLAUSE_DECL (cls))\n \t\t {\n \t\t is_private = true;\n@@ -8120,7 +8122,10 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)\n \t\t&& TREE_CODE (var_type) == ARRAY_TYPE\n \t\t&& !oacc_firstprivate_int)\n \t x = build_simple_mem_ref (x);\n-\t if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE)\n+\t if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE\n+\t\t|| (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP\n+\t\t && (OMP_CLAUSE_MAP_KIND (c) & GOMP_MAP_TO)\n+\t\t && OMP_CLAUSE_MAP_PRIVATE (c)))\n \t {\n \t\tgcc_assert (is_gimple_omp_oacc (ctx->stmt));\n \t\tif (oacc_firstprivate_int)\n@@ -9054,7 +9059,24 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)\n gimple_seq_add_seq (&new_body, join_seq);\n \n if (offloaded)\n-\tnew_body = maybe_catch_exception (new_body);\n+\t{\n+\t /* For OMP_CLAUSE_MAP_PRIVATE maps, add a copy back from private\n+\t storage to receiver ref, for copying back to host. */\n+\t for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))\n+\t if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP\n+\t\t&& (OMP_CLAUSE_MAP_KIND (c) & GOMP_MAP_FROM)\n+\t\t&& OMP_CLAUSE_MAP_PRIVATE (c))\n+\t {\n+\t\ttree var = OMP_CLAUSE_DECL (c);\n+\t\ttree new_var = lookup_decl (var, ctx);\n+\t\ttree x = build_receiver_ref (var, true, ctx);\n+\t\tgimple_seq seq = NULL;\n+\t\tgimplify_assign (x, new_var, &seq);\n+\t\tgimple_seq_add_seq (&new_body, seq);\n+\t }\n+\n+\t new_body = maybe_catch_exception (new_body);\n+\t}\n \n gimple_seq_add_stmt (&new_body, gimple_build_omp_return (false));\n gimple_omp_set_body (stmt, new_body);\ndiff --git a/gcc/tree.h b/gcc/tree.h\nindex a92ea11..cfe0ee2 100644\n--- a/gcc/tree.h\n+++ b/gcc/tree.h\n@@ -1554,6 +1554,9 @@ extern void protected_set_expr_location (tree, location_t);\n /* Nonzero if this map clause is for an ACC parallel reduction variable. */\n #define OMP_CLAUSE_MAP_IN_REDUCTION(NODE) \\\n TREE_PRIVATE (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP))\n+/* Nozero if this map is loaded to private storage inside offloaded region. */\n+#define OMP_CLAUSE_MAP_PRIVATE(NODE) \\\n+ TREE_NOTHROW (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP))\n \n #define OMP_CLAUSE_PROC_BIND_KIND(NODE) \\\n (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_PROC_BIND)->omp_clause.subcode.proc_bind_kind)\ndiff --git a/libgomp/ChangeLog.openacc b/libgomp/ChangeLog.openacc\nindex 74681f2..cd1f3ab 100644\n--- a/libgomp/ChangeLog.openacc\n+++ b/libgomp/ChangeLog.openacc\n@@ -1,3 +1,7 @@\n+2017-09-05 Chung-Lin Tang \n+\n+\t* testsuite/libgomp.oacc-c-c++-common/reduction-9.c: New test.\n+\n 2017-08-11 Cesar Philippidis \n \n \t* plugin/plugin-nvptx.c (nvptx_exec): Dynamically allocate\ndiff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-9.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-9.c\nnew file mode 100644\nindex 0000000..d6e02fc\n--- /dev/null\n+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-9.c\n@@ -0,0 +1,41 @@\n+#include \n+#include \n+\n+int\n+main (int argc, char *argv[])\n+{\n+#define N 100\n+ int n = N;\n+ int i, j, tmp;\n+ int input[N*N], output[N], houtput[N];\n+\n+ for (i = 0; i < n * n; i++)\n+ input[i] = i;\n+\n+ for (i = 0; i < n; i++)\n+ {\n+ tmp = 0;\n+ for (j = 0; j < n; j++)\n+\ttmp += input[i * n + j];\n+ houtput[i] = tmp;\n+ }\n+ \n+ #pragma acc parallel loop gang\n+ for (i = 0; i < n; i++)\n+ {\n+ tmp = 0;\n+\n+ #pragma acc loop worker reduction(+:tmp)\n+ for (j = 0; j < n; j++)\n+\ttmp += input[i * n + j];\n+\n+ output[i] = tmp;\n+ }\n+\n+ /* Test if every worker-level reduction had correct private result. */\n+ for (i = 0; i < n; i++)\n+ if (houtput[i] != output[i])\n+ abort ();\n+\n+ return 0;\n+}\n","prefixes":["openacc","og7","committed"]}