From patchwork Tue Feb 12 19:02:46 2019 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: =?utf-8?q?Gerg=C3=B6_Barany?= X-Patchwork-Id: 1040788 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-495952-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="dTrH/UK8"; 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 43zX9M63NNz9sCh for ; Wed, 13 Feb 2019 06:03:03 +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 :from:subject:message-id:date:mime-version:content-type; q=dns; s=default; b=QWz2BD+m6/xv+Or9Q4K7IsOit6U6FgXgvXfC2BHxzXnL27jKNJ k3sDRWNIIOqNaSISfEVK4qRkPF5tCJtcEtnV1X/P5/WvDCZ9n+Irmr8uFPfkW5IO shYjogNNvu0HZotUzl5rpP0QxdgcPQvaX3KOOOj6YHJISu2mRaH1pJLc0= 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 :from:subject:message-id:date:mime-version:content-type; s= default; bh=CLp58B9MNzCsd4s3UDNqzbXvSr4=; b=dTrH/UK8jonK4mzndqGX UhiI6sIL9QO0/yGPVLRwsONNqwq/dI85yPD5Ccly58BR0FJpKLJ2qZQNz6tUYxtt qHfO221lXvbAfIOuGzqDPirxCLcYDWwXYX2qy+tLklJPdI5yHALhcQSr9wotoqYE nbvpLG/cX4DWKnTAn9PvsFw= Received: (qmail 120047 invoked by alias); 12 Feb 2019 19:02:57 -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 120039 invoked by uid 89); 12 Feb 2019 19:02:56 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-26.9 required=5.0 tests=BAYES_00, GIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3, RCVD_IN_DNSWL_NONE, SPF_PASS 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; Tue, 12 Feb 2019 19:02:54 +0000 Received: from nat-ies.mentorg.com ([192.94.31.2] helo=SVR-IES-MBX-03.mgc.mentorg.com) by relay1.mentorg.com with esmtps (TLSv1.2:ECDHE-RSA-AES256-SHA384:256) id 1gtdKV-0005LQ-V3 from Gergo_Barany@mentor.com for gcc-patches@gcc.gnu.org; Tue, 12 Feb 2019 11:02:52 -0800 Received: from [172.30.73.203] (137.202.0.90) by SVR-IES-MBX-03.mgc.mentorg.com (139.181.222.3) with Microsoft SMTP Server (TLS) id 15.0.1320.4; Tue, 12 Feb 2019 19:02:48 +0000 To: , Thomas Schwinge From: =?utf-8?q?Gerg=C3=B6_Barany?= Subject: [PATCH, OpenACC, og8] OpenACC kernels control flow analysis bug fix Message-ID: Date: Tue, 12 Feb 2019 20:02:46 +0100 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:60.0) Gecko/20100101 Thunderbird/60.2.1 MIME-Version: 1.0 Hi all, The attached patch fixes a bug in recent work on OpenACC "kernels" regions. Jumps within nested binds or try statements were not analyzed correctly and could lead to ICEs. Tested on x86_64 with offloading to NVPTX. Thanks, Gergö Correctly handle nested bind and try statements in the OpenACC kernels conversion control-flow region analysis. gcc/ * omp-oacc-kernels.c (control_flow_regions::compute_regions): Factored out... (control_flow_regions::visit_gimple_seq): ... this new method, now also handling bind and try statements. gcc/testsuite/ * gcc/testsuite/c-c++-common/goacc/kernels-decompose-1.c: Add tests. From 6db1f381f344b4482dcca6b82fc6316d172840be Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Gerg=C3=B6=20Barany?= Date: Mon, 11 Feb 2019 08:23:31 -0800 Subject: [PATCH] OpenACC kernels control flow analysis bug fix Correctly handle nested bind and try statements in the OpenACC kernels conversion control-flow region analysis. gcc/ * omp-oacc-kernels.c (control_flow_regions::compute_regions): Factored out... (control_flow_regions::visit_gimple_seq): ... this new method, now also handling bind and try statements. gcc/testsuite/ * gcc/testsuite/c-c++-common/goacc/kernels-decompose-1.c: Add tests. --- gcc/omp-oacc-kernels.c | 94 +++++++++++++++------- .../c-c++-common/goacc/kernels-decompose-1.c | 44 ++++++++++ 2 files changed, 107 insertions(+), 31 deletions(-) diff --git a/gcc/omp-oacc-kernels.c b/gcc/omp-oacc-kernels.c index d1db492..1fa2647 100644 --- a/gcc/omp-oacc-kernels.c +++ b/gcc/omp-oacc-kernels.c @@ -935,12 +935,24 @@ class control_flow_regions control-flow regions in the statement sequence SEQ. */ void compute_regions (gimple_seq seq); + /* Helper for compute_regions, scanning a single statement sequence SEQ + starting at index IDX and returning the next index after the last + statement in the sequence. */ + size_t visit_gimple_seq (gimple_seq seq, size_t idx); + /* The mapping from statement indices to region representatives. */ vec representatives; /* A cache mapping statement indices to a flag indicating whether the statement is a top level OpenACC for loop. */ vec omp_for_loops; + + /* A mapping of control flow statements (goto, switch, cond) to their + representatives. */ + hash_map control_flow_reps; + + /* A mapping of labels to their representatives. */ + hash_map label_reps; }; control_flow_regions::control_flow_regions (gimple_seq seq) @@ -1008,41 +1020,12 @@ control_flow_regions::union_reps (size_t a, size_t b) void control_flow_regions::compute_regions (gimple_seq seq) { - hash_map control_flow_reps; - hash_map label_reps; - size_t current_region = 0, idx = 0; + size_t idx = 0; /* In a first pass, assign an initial region to each statement. Except in the case of OpenACC loops, each statement simply gets the same region representative as its predecessor. */ - for (gimple_stmt_iterator gsi = gsi_start (seq); - !gsi_end_p (gsi); - gsi_next (&gsi)) - { - gimple *stmt = gsi_stmt (gsi); - gimple *omp_for = top_level_omp_for_in_stmt (stmt); - omp_for_loops.safe_push (omp_for != NULL); - if (omp_for != NULL) - { - /* Assign a new region to this loop and to its successor. */ - current_region = idx; - representatives.safe_push (current_region); - current_region++; - } - else - { - representatives.safe_push (current_region); - /* Remember any jumps and labels for the second pass below. */ - if (gimple_code (stmt) == GIMPLE_COND - || gimple_code (stmt) == GIMPLE_SWITCH - || gimple_code (stmt) == GIMPLE_GOTO) - control_flow_reps.put (stmt, current_region); - else if (gimple_code (stmt) == GIMPLE_LABEL) - label_reps.put (gimple_label_label (as_a (stmt)), - current_region); - } - idx++; - } + visit_gimple_seq (seq, idx); gcc_assert (representatives.length () == omp_for_loops.length ()); /* Revisit all the control flow statements and union the region of each @@ -1087,6 +1070,55 @@ control_flow_regions::compute_regions (gimple_seq seq) } } +size_t +control_flow_regions::visit_gimple_seq (gimple_seq seq, size_t idx) +{ + size_t current_region = idx; + + for (gimple_stmt_iterator gsi = gsi_start (seq); + !gsi_end_p (gsi); + gsi_next (&gsi)) + { + gimple *stmt = gsi_stmt (gsi); + gimple *omp_for = top_level_omp_for_in_stmt (stmt); + omp_for_loops.safe_push (omp_for != NULL); + if (omp_for != NULL) + { + /* Assign a new region to this loop and to its successor. */ + current_region = idx; + representatives.safe_push (current_region); + current_region++; + } + else if (gimple_code (stmt) == GIMPLE_BIND) + { + representatives.safe_push (current_region); + gimple_seq body = gimple_bind_body (as_a (stmt)); + idx = visit_gimple_seq (body, idx); + } + else if (gimple_code (stmt) == GIMPLE_TRY) + { + representatives.safe_push (current_region); + idx = visit_gimple_seq (gimple_try_eval (stmt), idx); + idx = visit_gimple_seq (gimple_try_cleanup (stmt), idx); + } + else + { + representatives.safe_push (current_region); + /* Remember any jumps and labels for the second pass below. */ + if (gimple_code (stmt) == GIMPLE_COND + || gimple_code (stmt) == GIMPLE_SWITCH + || gimple_code (stmt) == GIMPLE_GOTO) + control_flow_reps.put (stmt, current_region); + else if (gimple_code (stmt) == GIMPLE_LABEL) + label_reps.put (gimple_label_label (as_a (stmt)), + current_region); + } + idx++; + } + + return idx; +} + /* Decompose the body of the KERNELS_REGION, which was originally annotated with the KERNELS_CLAUSES, into a series of regions. */ diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-1.c b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-1.c index 7255d69..6a3381e 100644 --- a/gcc/testsuite/c-c++-common/goacc/kernels-decompose-1.c +++ b/gcc/testsuite/c-c++-common/goacc/kernels-decompose-1.c @@ -118,5 +118,49 @@ main () #pragma acc kernels /* { dg-message "note: beginning .gang-single. region in OpenACC .kernels. construct" } */ ; + #pragma acc kernels + { + int a[20], b[20]; + int n = 0; + + /* The goto in this loop is placed into a bind; make sure it is + traversed correctly. */ + for (int i = 0; i < 20; i++) + goto label_43; + + #pragma acc loop independent /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" } */ + for (int m = 0; m < 20; m++) + b[m] = a[m]; + +label_43: + b[n] = a[n]; + } + +#ifdef __cplusplus + #pragma acc kernels + { + int a[20], b[20]; + int n = 0; + + for (int i = 0; i < 20; i++) + { + try + { + goto label_44; + } + catch (int) + { + } + } + + #pragma acc loop independent /* { dg-message "note: unparallelized loop nest in OpenACC .kernels. region: it's executed conditionally" "" { target c++ } } */ + for (int m = 0; m < 20; m++) + b[m] = a[m]; + +label_44: + b[n] = a[n]; + } +#endif + return 0; } -- 2.8.1