From patchwork Thu Mar 20 14:34:37 2014 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 332192 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 9182E2C009A for ; Fri, 21 Mar 2014 01:35:00 +1100 (EST) DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:subject:in-reply-to:references:date:message-id:mime-version :content-type; q=dns; s=default; b=kql/88I/Cf9Ostw0iwqBeDeRt8bKo Ep0hxdtayVg14OdRZCqleLmxfPkEZn2d/gLu78l0SuxmA2G7Lneu9vp7pG262Swa KkXv+TfyVAqZ37pXL0nCLolU6fVNXDMM3Eq5J7tK8V+sfW7xJtiJXByhgrn3Hlu/ Hg/KudxqVHPPoA= 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:from :to:subject:in-reply-to:references:date:message-id:mime-version :content-type; s=default; bh=P39g3NG54W8Q5veh2TklEgPd+8E=; b=lIm GPJKXay86fXAMjCYM3EGIa7tVnKblu9tZgh2hQ5/FgmCfoqxqPh/ioKyDXcvm4Dr 5fomLZSeX4xSKtIQPiYINxzIjbOh4wXER8LyMk+vAI48hHRNI4mtoDUAibyRskod OuYX/W2Phnl/rTiXZOjKG3SFCiGQiFsmCz733vok= Received: (qmail 25059 invoked by alias); 20 Mar 2014 14:34:50 -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 24922 invoked by uid 89); 20 Mar 2014 14:34:49 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-2.1 required=5.0 tests=AWL, BAYES_00 autolearn=ham version=3.3.2 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; Thu, 20 Mar 2014 14:34:46 +0000 Received: from svr-orw-exc-10.mgc.mentorg.com ([147.34.98.58]) by relay1.mentorg.com with esmtp id 1WQe3O-0004gX-DE from Thomas_Schwinge@mentor.com for gcc-patches@gcc.gnu.org; Thu, 20 Mar 2014 07:34:42 -0700 Received: from SVR-IES-FEM-01.mgc.mentorg.com ([137.202.0.104]) by SVR-ORW-EXC-10.mgc.mentorg.com with Microsoft SMTPSVC(6.0.3790.4675); Thu, 20 Mar 2014 07:34:42 -0700 Received: from feldtkeller.schwinge.homeip.net (137.202.0.76) by SVR-IES-FEM-01.mgc.mentorg.com (137.202.0.104) with Microsoft SMTP Server id 14.2.247.3; Thu, 20 Mar 2014 14:34:40 +0000 From: Thomas Schwinge To: Subject: [gomp4] Nesting of OpenACC constructs inside of OpenACC data constructs (was: [gomp4 2/3] OpenACC data construct implementation in terms of GF_OMP_TARGET_KIND_OACC_DATA.) In-Reply-To: <1393014736-19719-2-git-send-email-thomas@codesourcery.com> References: <877g8os0vx.fsf@kepler.schwinge.homeip.net> <1393014736-19719-1-git-send-email-thomas@codesourcery.com> <1393014736-19719-2-git-send-email-thomas@codesourcery.com> User-Agent: Notmuch/0.9-101-g81dad07 (http://notmuchmail.org) Emacs/23.4.1 (i486-pc-linux-gnu) Date: Thu, 20 Mar 2014 15:34:37 +0100 Message-ID: <874n2tlz0i.fsf@kepler.schwinge.homeip.net> MIME-Version: 1.0 Hi! Applied in r208701 to gomp-4_0-branch: commit 22dd36a31c433dcd8bcc890d245a9e4ac6ed9c7f Author: tschwinge Date: Thu Mar 20 14:33:28 2014 +0000 Nesting of OpenACC constructs inside of OpenACC data constructs. gcc/ * omp-low.c (check_omp_nesting_restrictions): Allow nesting of OpenACC constructs inside of OpenACC data constructs. gcc/testsuite/ * c-c++-common/goacc/nesting-1.c: New file. * c-c++-common/goacc/nesting-data-1.c: Likewise. * c-c++-common/goacc/nesting-fail-1.c: Update. git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@208701 138bc75d-0d04-0410-961f-82ee72b054a4 Grüße, Thomas diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp index 1aebc4d..f43452c 100644 --- gcc/ChangeLog.gomp +++ gcc/ChangeLog.gomp @@ -1,3 +1,8 @@ +2014-03-20 Thomas Schwinge + + * omp-low.c (check_omp_nesting_restrictions): Allow nesting of + OpenACC constructs inside of OpenACC data constructs. + 2014-03-18 Ilmir Usmanov * tree.def (OACC_LOOP): New tree code. diff --git gcc/omp-low.c gcc/omp-low.c index f1b0fa5..23a0dda 100644 --- gcc/omp-low.c +++ gcc/omp-low.c @@ -2416,26 +2416,31 @@ scan_omp_teams (gimple stmt, omp_context *outer_ctx) static bool check_omp_nesting_restrictions (gimple stmt, omp_context *ctx) { - omp_context *ctx_; - /* TODO: While the OpenACC specification does allow for certain kinds of - nesting, we don't support that yet. */ - /* No nesting of STMT (which is an OpenACC or OpenMP one, or a GOMP builtin) - inside any OpenACC CTX. */ - for (ctx_ = ctx; ctx_ != NULL; ctx_ = ctx_->outer) - if (is_gimple_omp (ctx_->stmt) - && is_gimple_omp_oacc_specifically (ctx_->stmt)) - { - error_at (gimple_location (stmt), - "may not be nested"); - return false; - } - /* No nesting of OpenACC STMT inside any OpenACC or OpenMP CTX. */ + nesting, we don't support many of these yet. */ if (is_gimple_omp (stmt) && is_gimple_omp_oacc_specifically (stmt)) { - for (ctx_ = ctx; ctx_ != NULL; ctx_ = ctx_->outer) - if (is_gimple_omp (ctx_->stmt)) + /* No nesting of OpenACC STMT inside any OpenACC or OpenMP CTX different + from an OpenACC data construct. */ + for (omp_context *ctx_ = ctx; ctx_ != NULL; ctx_ = ctx_->outer) + if (is_gimple_omp (ctx_->stmt) + && !(gimple_code (ctx_->stmt) == GIMPLE_OMP_TARGET + && (gimple_omp_target_kind (ctx_->stmt) + == GF_OMP_TARGET_KIND_OACC_DATA))) + { + error_at (gimple_location (stmt), + "may not be nested"); + return false; + } + } + else + { + /* No nesting of non-OpenACC STMT (that is, an OpenMP one, or a GOMP + builtin) inside any OpenACC CTX. */ + for (omp_context *ctx_ = ctx; ctx_ != NULL; ctx_ = ctx_->outer) + if (is_gimple_omp (ctx_->stmt) + && is_gimple_omp_oacc_specifically (ctx_->stmt)) { error_at (gimple_location (stmt), "may not be nested"); diff --git gcc/testsuite/ChangeLog.gomp gcc/testsuite/ChangeLog.gomp index fd38d80..13e99d5 100644 --- gcc/testsuite/ChangeLog.gomp +++ gcc/testsuite/ChangeLog.gomp @@ -1,5 +1,9 @@ 2014-03-20 Thomas Schwinge + * c-c++-common/goacc/nesting-1.c: New file. + * c-c++-common/goacc/nesting-data-1.c: Likewise. + * c-c++-common/goacc/nesting-fail-1.c: Update. + * c-c++-common/goacc/nesting-fail-1.c (f_acc_kernels): Replace OpenACC parallel with kernels directive. diff --git gcc/testsuite/c-c++-common/goacc/nesting-1.c gcc/testsuite/c-c++-common/goacc/nesting-1.c new file mode 100644 index 0000000..3a22292 --- /dev/null +++ gcc/testsuite/c-c++-common/goacc/nesting-1.c @@ -0,0 +1,13 @@ +void +f_acc_data (void) +{ +#pragma acc data + { +#pragma acc parallel + ; +#pragma acc kernels + ; +#pragma acc data + ; + } +} diff --git gcc/testsuite/c-c++-common/goacc/nesting-data-1.c gcc/testsuite/c-c++-common/goacc/nesting-data-1.c new file mode 100644 index 0000000..fefe6cd --- /dev/null +++ gcc/testsuite/c-c++-common/goacc/nesting-data-1.c @@ -0,0 +1,61 @@ +void +f (void) +{ + unsigned char c, ca[15], caa[20][30]; + +#pragma acc data copyin(c) + { + c = 5; + ca[3] = c; + caa[3][12] = ca[3] + caa[3][12]; + +#pragma acc data copyin(ca[2:4]) + { + c = 6; + ca[4] = c; + caa[3][12] = ca[3] + caa[3][12]; + } + +#pragma acc parallel copyout(ca[3:4]) + { + c = 7; + ca[5] = c; + caa[3][12] = ca[3] + caa[3][12]; + } + +#pragma acc kernels copy(ca[4:4]) + { + c = 8; + ca[6] = c; + caa[3][12] = ca[3] + caa[3][12]; + } + +#pragma acc data pcopy(ca[5:7]) + { + c = 15; + ca[7] = c; + caa[3][12] = ca[3] + caa[3][12]; + +#pragma acc data pcopyin(caa[3:7][0:30]) + { + c = 16; + ca[8] = c; + caa[3][12] = ca[3] + caa[3][12]; + } + +#pragma acc parallel pcopyout(caa[3:7][0:30]) + { + c = 17; + ca[9] = c; + caa[3][12] = ca[3] + caa[3][12]; + } + +#pragma acc kernels pcopy(caa[3:7][0:30]) + { + c = 18; + ca[10] = c; + caa[3][12] = ca[3] + caa[3][12]; + } + } + } +} diff --git gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c index ca8921f..00dc602 100644 --- gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c +++ gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c @@ -1,5 +1,5 @@ /* TODO: While the OpenACC specification does allow for certain kinds of - nesting, we don't support that yet. */ + nesting, we don't support many of these yet. */ void f_acc_parallel (void) { @@ -15,7 +15,7 @@ f_acc_parallel (void) } /* TODO: While the OpenACC specification does allow for certain kinds of - nesting, we don't support that yet. */ + nesting, we don't support many of these yet. */ void f_acc_kernels (void) { @@ -29,19 +29,3 @@ f_acc_kernels (void) ; } } - -/* TODO: While the OpenACC specification does allow for certain kinds of - nesting, we don't support that yet. */ -void -f_acc_data (void) -{ -#pragma acc data - { -#pragma acc parallel /* { dg-error "may not be nested" } */ - ; -#pragma acc kernels /* { dg-error "may not be nested" } */ - ; -#pragma acc data /* { dg-error "may not be nested" } */ - ; - } -}