From patchwork Wed Feb 24 20:07:57 2016 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Nathan Sidwell X-Patchwork-Id: 587703 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org 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 77C26140C36 for ; Thu, 25 Feb 2016 07:08:13 +1100 (AEDT) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b=u+PBY/8G; dkim-atps=neutral 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=innRmr/2YURxA0IYWktRrmnLtK9b18DRHeQiXM2LoVD5Ro0yrx ITh8i/BNUxpFIIPXk3/4hYevIydxLfZ4YYtV1xX17zUSw7NBBHNSRPrU/3JzDoDP M4VrsvINGtB7kVzppP/z14LJ5cV9s/G8FyQVM5N1119r7aFbAi6DNACq0= 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=t4mqwnME7suQgnE6pTjnuVaotbI=; b=u+PBY/8GBPtUzSkCvB3L XqzmWva8Y+JGfGsLKVRqbyYXU//luDa7/XvJrDKH0z1XPxGn1Ofi9CabaH+O6U6k ynMAp3fXyFi8GfxuBoBTnNYpKozmZpfKBm3ig4/XG77wIGXuqcNmgka5nonwWbgY S7JxthYwG4Is0hfXEDpMAr0= Received: (qmail 115234 invoked by alias); 24 Feb 2016 20:08:04 -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 115223 invoked by uid 89); 24 Feb 2016 20:08:03 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=0.9 required=5.0 tests=BAYES_50, FREEMAIL_FROM, KAM_ASCII_DIVIDERS, RCVD_IN_DNSWL_LOW, SPF_PASS autolearn=no version=3.3.2 spammy=falling, sk:unsigne, ifn, abstraction X-HELO: mail-qg0-f47.google.com Received: from mail-qg0-f47.google.com (HELO mail-qg0-f47.google.com) (209.85.192.47) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (AES128-GCM-SHA256 encrypted) ESMTPS; Wed, 24 Feb 2016 20:08:01 +0000 Received: by mail-qg0-f47.google.com with SMTP id b67so23361548qgb.1 for ; Wed, 24 Feb 2016 12:08:00 -0800 (PST) X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20130820; h=x-gm-message-state:sender:to:cc:from:subject:message-id:date :user-agent:mime-version:content-type; bh=nb5LUD0tsXuQNKHiupTQ9qcZFe+2rJ0STNUyl9BdSnw=; b=B65l4+ytJHxFWQcZdUGGzlf/3sDNNH0e8bG2vEnrP8DI0DAdWAruv5l0QRy5awF2gb zm0xhhdTweaKasmZrxgKhI29b+SQs45Xhap0Zes28cZE6Nc/iq/VNetp89gmAsy2Ov5K p+nTUs304BFYh34gEoP/4Aevr939cSCN805nSwrTTL7biLVONto6UpldIvuKMTGl3dBF zAIJ6KkztMmGMVR+bWCCNRzrWiV/vlV6j76lNIU47JGFDoYQS5KageSGsG20yol2Pcyl JiMgHzjvNdkKsWw4MS25opUAlESXkXfZbBN309LUrcycb82RlY9MIXuoCnLeM5GmFEcI BYkA== X-Gm-Message-State: AG10YOReKI4fvM8um6FECDg3+bgmiC3YLdbCkF2R5KwP/gBgegBZfRKxydIt7wdJ/yoVNQ== X-Received: by 10.140.22.139 with SMTP id 11mr51917675qgn.34.1456344478882; Wed, 24 Feb 2016 12:07:58 -0800 (PST) Received: from ?IPv6:2601:181:c000:c497:a2a8:cdff:fe3e:b48? ([2601:181:c000:c497:a2a8:cdff:fe3e:b48]) by smtp.googlemail.com with ESMTPSA id r132sm1836461qhc.41.2016.02.24.12.07.57 (version=TLSv1/SSLv3 cipher=OTHER); Wed, 24 Feb 2016 12:07:58 -0800 (PST) To: Jakub Jelinek Cc: GCC Patches From: Nathan Sidwell Subject: [pr/69916] ICE with empty loops Message-ID: <56CE0D9D.8030704@acm.org> Date: Wed, 24 Feb 2016 15:07:57 -0500 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:38.0) Gecko/20100101 Thunderbird/38.6.0 MIME-Version: 1.0 Jakub, this patch fixes the ICE reported in pr69916 (https://gcc.gnu.org/bugzilla/show_bug.cgi?id=69916) The loop is lowered at omp-lowering, but subsequently determined to be dead before we get to oacc-target-lower. The loop CF is removed along with the (pure) IFN_OACC_LOOP function calls inserted during lowering. However the IFN_UNIQUE loop head & tail calls remain (because they are not pure). Thus in the oacc-target-lower pass we rediscover the loop structure. Firstly we assign a specific axis for this loop -- as it's auto. That's a pessimization, but not wrong. However, we then scan the loop to adjust the expected OACC_LOOP calls with the determined partitioning information. As they're not there, we end up falling out of the function and die with a single_succ_edge assert. (In general we might end up finding OACC_LOOP calls of an inner loop, or meeting a block with more than one successor. Either would be bad.) This patch changes the loop transformation to count OACC_LOOP calls it encounters when rediscovering the loops, and uses that count for the OACC_LOOP adjustment scan (rather than expect OACC_LOOP_BOUND to be the last one). That fixes the ICE. While there it is trivial to mark the loop as not to be partitioned, if we discover no OACC_LOOP calls, which addresses the pessimization mentioned above. As the loop is no longer partitioned, the fork and join markers, end up being deleted. ok for trunk? nathan 2016-02-24 Nathan Sidwell gcc/ PR middle-end/69916 * omp-low.c (struct oacc_loop): Add ifns. (new_oacc_loop_raw): Initialize it. (finish_oacc_loop): Clear mask & flags if no ifns. (oacc_loop_discover_walk): Count IFN_GOACC_LOOP calls. (oacc_loop_xform_loop): Add ifns arg & adjust. (oacc_loop_process): Adjust oacc_loop_xform_loop call. gcc/testsuite/ PR middle-end/69916 * c-c-++-common/goacc/pr69916.c: New. Index: omp-low.c =================================================================== --- omp-low.c (revision 233663) +++ omp-low.c (working copy) @@ -241,8 +241,9 @@ struct oacc_loop tree routine; /* Pseudo-loop enclosing a routine. */ unsigned mask; /* Partitioning mask. */ - unsigned flags; /* Partitioning flags. */ - tree chunk_size; /* Chunk size. */ + unsigned flags; /* Partitioning flags. */ + unsigned ifns; /* Contained loop abstraction functions. */ + tree chunk_size; /* Chunk size. */ gcall *head_end; /* Final marker of head sequence. */ }; @@ -20393,6 +20394,7 @@ new_oacc_loop_raw (oacc_loop *parent, lo loop->routine = NULL_TREE; loop->mask = loop->flags = 0; + loop->ifns = 0; loop->chunk_size = 0; loop->head_end = NULL; @@ -20454,6 +20456,9 @@ new_oacc_loop_routine (oacc_loop *parent static oacc_loop * finish_oacc_loop (oacc_loop *loop) { + /* If the loop has been collapsed, don't partition it. */ + if (!loop->ifns) + loop->mask = loop->flags = 0; return loop->parent; } @@ -20584,43 +20589,54 @@ oacc_loop_discover_walk (oacc_loop *loop if (!gimple_call_internal_p (call)) continue; - if (gimple_call_internal_fn (call) != IFN_UNIQUE) - continue; + switch (gimple_call_internal_fn (call)) + { + default: + break; - enum ifn_unique_kind kind - = (enum ifn_unique_kind) TREE_INT_CST_LOW (gimple_call_arg (call, 0)); - if (kind == IFN_UNIQUE_OACC_HEAD_MARK - || kind == IFN_UNIQUE_OACC_TAIL_MARK) - { - if (gimple_call_num_args (call) == 2) - { - gcc_assert (marker && !remaining); - marker = 0; - if (kind == IFN_UNIQUE_OACC_TAIL_MARK) - loop = finish_oacc_loop (loop); - else - loop->head_end = call; - } - else - { - int count = TREE_INT_CST_LOW (gimple_call_arg (call, 2)); + case IFN_GOACC_LOOP: + /* Count the goacc loop abstraction fns, to determine if the + loop was collapsed already. */ + loop->ifns++; + break; - if (!marker) + case IFN_UNIQUE: + enum ifn_unique_kind kind + = (enum ifn_unique_kind) (TREE_INT_CST_LOW + (gimple_call_arg (call, 0))); + if (kind == IFN_UNIQUE_OACC_HEAD_MARK + || kind == IFN_UNIQUE_OACC_TAIL_MARK) + { + if (gimple_call_num_args (call) == 2) { - if (kind == IFN_UNIQUE_OACC_HEAD_MARK) - loop = new_oacc_loop (loop, call); - remaining = count; + gcc_assert (marker && !remaining); + marker = 0; + if (kind == IFN_UNIQUE_OACC_TAIL_MARK) + loop = finish_oacc_loop (loop); + else + loop->head_end = call; } - gcc_assert (count == remaining); - if (remaining) + else { - remaining--; - if (kind == IFN_UNIQUE_OACC_HEAD_MARK) - loop->heads[marker] = call; - else - loop->tails[remaining] = call; + int count = TREE_INT_CST_LOW (gimple_call_arg (call, 2)); + + if (!marker) + { + if (kind == IFN_UNIQUE_OACC_HEAD_MARK) + loop = new_oacc_loop (loop, call); + remaining = count; + } + gcc_assert (count == remaining); + if (remaining) + { + remaining--; + if (kind == IFN_UNIQUE_OACC_HEAD_MARK) + loop->heads[marker] = call; + else + loop->tails[remaining] = call; + } + marker++; } - marker++; } } } @@ -20726,10 +20742,12 @@ oacc_loop_xform_head_tail (gcall *from, determined partitioning mask and chunking argument. */ static void -oacc_loop_xform_loop (gcall *end_marker, tree mask_arg, tree chunk_arg) +oacc_loop_xform_loop (gcall *end_marker, unsigned ifns, + tree mask_arg, tree chunk_arg) { gimple_stmt_iterator gsi = gsi_for_stmt (end_marker); + gcc_checking_assert (ifns); for (;;) { for (; !gsi_end_p (gsi); gsi_next (&gsi)) @@ -20749,13 +20767,13 @@ oacc_loop_xform_loop (gcall *end_marker, *gimple_call_arg_ptr (call, 5) = mask_arg; *gimple_call_arg_ptr (call, 4) = chunk_arg; - if (TREE_INT_CST_LOW (gimple_call_arg (call, 0)) - == IFN_GOACC_LOOP_BOUND) + ifns--; + if (!ifns) return; } - /* If we didn't see LOOP_BOUND, it should be in the single - successor block. */ + /* The LOOP_BOUND ifn, could be in the single successor + block. */ basic_block bb = single_succ (gsi_bb (gsi)); gsi = gsi_start_bb (bb); } @@ -20778,7 +20796,7 @@ oacc_loop_process (oacc_loop *loop) tree mask_arg = build_int_cst (unsigned_type_node, mask); tree chunk_arg = loop->chunk_size; - oacc_loop_xform_loop (loop->head_end, mask_arg, chunk_arg); + oacc_loop_xform_loop (loop->head_end, loop->ifns, mask_arg, chunk_arg); for (ix = 0; ix != GOMP_DIM_MAX && loop->heads[ix]; ix++) { Index: testsuite/c-c++-common/goacc/pr69916.c =================================================================== --- testsuite/c-c++-common/goacc/pr69916.c (nonexistent) +++ testsuite/c-c++-common/goacc/pr69916.c (working copy) @@ -0,0 +1,20 @@ +/* { dg-additional-options "-O2" } */ + +/* PR 69916, an loop determined to be empty sometime after omp-lower + and before oacc-device-lower can evaporate leading to no GOACC_LOOP + internal functions existing. */ + +int +main (void) +{ + +#pragma acc parallel + { + int j = 0; +#pragma acc loop private (j) + for (int i = 0; i < 10; i++) + j++; + } + + return 0; +}