From patchwork Fri Oct 16 09:17:25 2015 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Tom de Vries X-Patchwork-Id: 531156 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Received: from sourceware.org (unknown [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 196351402B4 for ; Fri, 16 Oct 2015 20:20:53 +1100 (AEDT) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b=j65r0hPb; 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=v+EkqCzD2krWFepobGNFdB6pnaFY1H7EftzRyBlpCUsJWSCM+l tgfZq5OYMMyyJ/FpC9OJtJnoCQSAdAgEw+3IizUrCvYP4JgU4/GRpR6ir0gTucQw 4PY35O7AsmaNzXfuh3Us5qYVTMEa/Q0l0x4m2O3P3FsTBLYNv0Un6LXac= 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=if7wdnDWSebdVZrGvRzDDZKsI1g=; b=j65r0hPb3dBXba7Dg7SU v8hGqgaIEHJOkg6IXgUoQkSeSYX3Ti1hu4emLBiujlcSPJk+vm+wa59hcF6LPM92 fHWpeu4+3r5SUmFNL6x20wLfl/U1gXVY6I31tKkX1/M4qT9cytfyUHV63iSyrHJd M24LJVmygIUSvwVn9+89HHQ= Received: (qmail 39353 invoked by alias); 16 Oct 2015 09:20:45 -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 39341 invoked by uid 89); 16 Oct 2015 09:20:44 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-2.0 required=5.0 tests=AWL, BAYES_00 autolearn=ham version=3.3.2 X-HELO: fencepost.gnu.org Received: from Unknown (HELO fencepost.gnu.org) (208.118.235.10) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (AES128-SHA encrypted) ESMTPS; Fri, 16 Oct 2015 09:18:44 +0000 Received: from eggs.gnu.org ([2001:4830:134:3::10]:37833) by fencepost.gnu.org with esmtps (TLS1.0:RSA_AES_256_CBC_SHA1:256) (Exim 4.82) (envelope-from ) id 1Zn19o-0000oT-36 for gcc-patches@gnu.org; Fri, 16 Oct 2015 05:18:36 -0400 Received: from Debian-exim by eggs.gnu.org with spam-scanned (Exim 4.71) (envelope-from ) id 1Zn19k-0002jX-MX for gcc-patches@gnu.org; Fri, 16 Oct 2015 05:18:35 -0400 Received: from relay1.mentorg.com ([192.94.38.131]:41361) by eggs.gnu.org with esmtp (Exim 4.71) (envelope-from ) id 1Zn19k-0002dD-Gu for gcc-patches@gnu.org; Fri, 16 Oct 2015 05:18:32 -0400 Received: from nat-ies.mentorg.com ([192.94.31.2] helo=SVR-IES-FEM-02.mgc.mentorg.com) by relay1.mentorg.com with esmtp id 1Zn19Y-0000pY-T0 from Tom_deVries@mentor.com ; Fri, 16 Oct 2015 02:18:21 -0700 Received: from [127.0.0.1] (137.202.0.76) by SVR-IES-FEM-02.mgc.mentorg.com (137.202.0.106) with Microsoft SMTP Server id 14.3.224.2; Fri, 16 Oct 2015 10:18:19 +0100 To: "gcc-patches@gnu.org" CC: Thomas Schwinge , Jakub Jelinek From: Tom de Vries Subject: [gomp4,committed] Handle oacc region in oacc routine Message-ID: <5620C0A5.20506@mentor.com> Date: Fri, 16 Oct 2015 11:17:25 +0200 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:38.0) Gecko/20100101 Thunderbird/38.3.0 MIME-Version: 1.0 X-detected-operating-system: by eggs.gnu.org: Windows NT kernel [generic] [fuzzy] X-Received-From: 192.94.38.131 Hi, this patch checks for occurance of oacc offload regions in oacc routines (which means nested parallelism, which is currently not supported) and gives an appropriate error message. Committed to gomp-4_0-branch. Thanks, - Tom Handle oacc region in oacc routine 2015-10-16 Tom de Vries * omp-low.c (check_omp_nesting_restrictions): Check for oacc region in oacc routine. * c-c++-common/goacc/parallel-in-routine.c: New test. --- gcc/omp-low.c | 9 +++++++++ gcc/testsuite/c-c++-common/goacc/parallel-in-routine.c | 8 ++++++++ 2 files changed, 17 insertions(+) create mode 100644 gcc/testsuite/c-c++-common/goacc/parallel-in-routine.c diff --git a/gcc/omp-low.c b/gcc/omp-low.c index f27bde7..f7e4afc 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -3204,6 +3204,15 @@ check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx) } break; case GIMPLE_OMP_TARGET: + if (is_gimple_omp_offloaded (stmt) + && get_oacc_fn_attrib (cfun->decl) != NULL) + { + error_at (gimple_location (stmt), + "OpenACC region inside of OpenACC routine, nested " + "parallelism not supported yet"); + return false; + } + for (; ctx != NULL; ctx = ctx->outer) { if (gimple_code (ctx->stmt) != GIMPLE_OMP_TARGET) diff --git a/gcc/testsuite/c-c++-common/goacc/parallel-in-routine.c b/gcc/testsuite/c-c++-common/goacc/parallel-in-routine.c new file mode 100644 index 0000000..b93d63b --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/parallel-in-routine.c @@ -0,0 +1,8 @@ +#pragma acc routine +void +foo (void) +{ +#pragma acc parallel /* { dg-error "OpenACC region inside of OpenACC routine, nested parallelism not supported yet" } */ + ; +} + -- 1.9.1