From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 29814 invoked by alias); 7 Mar 2018 13:20:37 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Received: (qmail 29796 invoked by uid 89); 7 Mar 2018 13:20:36 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-24.8 required=5.0 tests=AWL,BAYES_00,GIT_PATCH_0,GIT_PATCH_1,GIT_PATCH_2,GIT_PATCH_3,RCVD_IN_DNSWL_NONE,SPF_PASS,URIBL_RED autolearn=ham version=3.3.2 spammy= X-HELO: relay1.mentorg.com Received: from relay1.mentorg.com (HELO relay1.mentorg.com) (192.94.38.131) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Wed, 07 Mar 2018 13:20:34 +0000 Received: from nat-ies.mentorg.com ([192.94.31.2] helo=SVR-IES-MBX-04.mgc.mentorg.com) by relay1.mentorg.com with esmtps (TLSv1.2:ECDHE-RSA-AES256-SHA384:256) id 1etYzf-0006Dr-6R from Tom_deVries@mentor.com ; Wed, 07 Mar 2018 05:20:31 -0800 Received: from [172.30.72.166] (137.202.0.87) by SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) with Microsoft SMTP Server (TLS) id 15.0.1320.4; Wed, 7 Mar 2018 13:20:27 +0000 To: Jakub Jelinek , Richard Biener CC: GCC Patches From: Tom de Vries Subject: [PATCH] Fix ICE for static vars in offloaded functions Message-ID: <2299637a-3124-8609-e68d-225703f16abe@mentor.com> Date: Wed, 07 Mar 2018 13:20:00 -0000 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:52.0) Gecko/20100101 Thunderbird/52.6.0 MIME-Version: 1.0 Content-Type: multipart/mixed; boundary="------------3E37F83377C4F29AD1B847BE" X-ClientProxiedBy: svr-ies-mbx-02.mgc.mentorg.com (139.181.222.2) To SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) X-SW-Source: 2018-03/txt/msg00320.txt.bz2 --------------3E37F83377C4F29AD1B847BE Content-Type: text/plain; charset="utf-8"; format=flowed Content-Transfer-Encoding: 7bit Content-length: 1919 Hi, if we compile the testcase pr84592-2.c from the patch: ... #include int main (void) { int n[1]; n[0] = 3; #pragma omp target { static int test[4] = { 1, 2, 3, 4 }; n[0] += test[n[0]]; } if (n[0] != 7) abort (); return 0; } ... for nvptx offloading, we run into an assert: ... lto1: internal compiler error: in input_varpool_node, at lto-cgraph.c:1424 0x959ebb input_varpool_node gcc/lto-cgraph.c:1422 0x959ebb input_cgraph_1 gcc/lto-cgraph.c:1544 0x959ebb input_symtab() gcc/lto-cgraph.c:1858 0x5aceac read_cgraph_and_symbols gcc/lto/lto.c:2891 0x5aceac lto_main() gcc/lto/lto.c:3356 ... The assert we run into is: ... 1422 gcc_assert (flag_ltrans 1423 || (!node->in_other_partition 1424 && !node->used_from_other_partition)); ... where node is: ... (gdb) call debug_generic_expr (node.decl) test ... and the reason the assert triggers is: ... (gdb) p node.in_other_partition $1 = 1 ... AFAIU, what this means is that the variable test is placed in a different partition than the offloading function main._omp_fn.0 that uses the variable. I looked at where global variables are put into offload_vars, and found that that happens in varpool_node::get_create: ... if ((flag_openacc || flag_openmp) && lookup_attribute ("omp declare target", DECL_ATTRIBUTES (decl))) { node->offloadable = 1; if (ENABLE_OFFLOADING && !DECL_EXTERNAL (decl)) { g->have_offload = true; if (!in_lto_p) vec_safe_push (offload_vars, decl); } } ... The patch fixes the ICE there by marking the varpool_node test as offloadable as well. Build and reg-tested libgomp on x86_64 with nvptx accelerator. Bootstrapped and reg-tested on x86_64. OK for stage4 trunk? Thanks, - Tom --------------3E37F83377C4F29AD1B847BE Content-Type: text/x-patch; name="0001-Fix-ICE-for-static-vars-in-offloaded-functions.patch" Content-Transfer-Encoding: 7bit Content-Disposition: inline; filename="0001-Fix-ICE-for-static-vars-in-offloaded-functions.patch" Content-length: 3559 Fix ICE for static vars in offloaded functions 2018-03-06 Tom de Vries PR lto/84592 * varpool.c (varpool_node::get_create): Mark static variables in offloaded functions as offloadable. * testsuite/libgomp.c/pr84592-2.c: New test. * testsuite/libgomp.c/pr84592.c: New test. * testsuite/libgomp.oacc-c-c++-common/pr84592-3.c: New test. --- gcc/varpool.c | 18 +++++++++--- libgomp/testsuite/libgomp.c/pr84592-2.c | 20 ++++++++++++++ libgomp/testsuite/libgomp.c/pr84592.c | 32 ++++++++++++++++++++++ .../libgomp.oacc-c-c++-common/pr84592-3.c | 32 ++++++++++++++++++++++ 4 files changed, 98 insertions(+), 4 deletions(-) diff --git a/gcc/varpool.c b/gcc/varpool.c index 418753cca2a..a4fd892ca4d 100644 --- a/gcc/varpool.c +++ b/gcc/varpool.c @@ -151,11 +151,21 @@ varpool_node::get_create (tree decl) node = varpool_node::create_empty (); node->decl = decl; - if ((flag_openacc || flag_openmp) - && lookup_attribute ("omp declare target", DECL_ATTRIBUTES (decl))) + if (flag_openacc || flag_openmp) { - node->offloadable = 1; - if (ENABLE_OFFLOADING && !DECL_EXTERNAL (decl)) + bool offload_var + = lookup_attribute ("omp declare target", DECL_ATTRIBUTES (decl)); + bool in_offload_func + = (cfun + && TREE_STATIC (decl) + && (lookup_attribute ("omp target entrypoint", + DECL_ATTRIBUTES (cfun->decl)) + || lookup_attribute ("omp declare target", + DECL_ATTRIBUTES (cfun->decl)))); + if (offload_var || in_offload_func) + node->offloadable = 1; + + if (offload_var && ENABLE_OFFLOADING && !DECL_EXTERNAL (decl)) { g->have_offload = true; if (!in_lto_p) diff --git a/libgomp/testsuite/libgomp.c/pr84592-2.c b/libgomp/testsuite/libgomp.c/pr84592-2.c new file mode 100644 index 00000000000..021497b28ff --- /dev/null +++ b/libgomp/testsuite/libgomp.c/pr84592-2.c @@ -0,0 +1,20 @@ +#include + +int +main (void) +{ + int n[1]; + + n[0] = 3; + +#pragma omp target + { + static int test[4] = { 1, 2, 3, 4 }; + n[0] += test[n[0]]; + } + + if (n[0] != 7) + abort (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/pr84592.c b/libgomp/testsuite/libgomp.c/pr84592.c new file mode 100644 index 00000000000..197fd19bacc --- /dev/null +++ b/libgomp/testsuite/libgomp.c/pr84592.c @@ -0,0 +1,32 @@ +/* { dg-additional-options "-ftree-switch-conversion" } */ + +#include + +int +main (void) +{ + int n[1]; + + n[0] = 4; + +#pragma omp target + { + int a = n[0]; + + switch (a & 3) + { + case 0: a = 4; break; + case 1: a = 3; break; + case 2: a = 2; break; + default: + a = 1; break; + } + + n[0] = a; + } + + if (n[0] != 4) + abort (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr84592-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr84592-3.c new file mode 100644 index 00000000000..afcc1de7635 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr84592-3.c @@ -0,0 +1,32 @@ +/* { dg-additional-options "-ftree-switch-conversion" } */ + +#include + +#pragma acc routine seq +static int __attribute__((noinline)) foo (int n) +{ + switch (n & 3) + { + case 0: return 4; + case 1: return 3; + case 2: return 2; + default: + return 1; + } +} + +int +main (void) +{ + int n[1]; + n[0] = 4; +#pragma acc parallel copy(n) + { + n[0] = foo (n[0]); + } + + if (n[0] != 4) + abort (); + + return 0; +} --------------3E37F83377C4F29AD1B847BE--