From patchwork Wed Mar 7 13:20:26 2018 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Tom de Vries X-Patchwork-Id: 882601 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (mailfrom) smtp.mailfrom=gcc.gnu.org (client-ip=209.132.180.131; helo=sourceware.org; envelope-from=gcc-patches-return-474384-incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=mentor.com Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b="rdgWOed2"; dkim-atps=neutral Received: from sourceware.org (server1.sourceware.org [209.132.180.131]) (using TLSv1.2 with cipher ECDHE-RSA-AES256-GCM-SHA384 (256/256 bits)) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 3zxDmF4nNZz9s32 for ; Thu, 8 Mar 2018 00:20:43 +1100 (AEDT) DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:to:cc :from:subject:message-id:date:mime-version:content-type; q=dns; s=default; b=W6XwpOLxsadEX5Yp4hzVcmhwdPu/xPJ1+AcLxgu1nbgboVoIzH +A8I6Le0mKx7OFVCeDTr69rKEUeY6J9mFyduXbV1qXgy9k5Vat1yEjvESzJA+pO/ u5nPWqPfQlRGsTOQDOw9cMXW/ii9SmbmT2+zjri1PnZJ+1LUjQpQmfQAM= DKIM-Signature: v=1; a=rsa-sha1; c=relaxed; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:to:cc :from:subject:message-id:date:mime-version:content-type; s= default; bh=TQuJwEQo+0ynlrO0iaNUuTTPX0U=; b=rdgWOed2/wa0x8eRRGCL R5WG8AkJsH/RdtyRf3AkYkIPp8sXCfPzVq/croGgXUW0N/gdcKgo0iAPOlHul3lr ErmWzet7pLI6ZhoZgDz6JrN2+GkJN5NsH654779fMwBT0ymlaMxW/Q6yziLPD89w VK0fbZETZxslV4C54EbPwNY= 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-Unsubscribe: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Delivered-To: mailing list gcc-patches@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, 7 Mar 2018 14:20:26 +0100 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:52.0) Gecko/20100101 Thunderbird/52.6.0 MIME-Version: 1.0 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) 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 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; +}