From patchwork Wed Oct 21 17:00:56 2015 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Nathan Sidwell X-Patchwork-Id: 533963 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 B74F71409F8 for ; Thu, 22 Oct 2015 04:01:10 +1100 (AEDT) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b=cbgWM3M0; 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 :from:subject:message-id:date:mime-version:content-type; q=dns; s=default; b=oYuRUrGutxCCDLHpR/+yIgx/mv9F67qOhUQ7b4T4+f+UPaqiCZ h8dhiaDzmJMJf40BHdmLShJO8tQfK2vA5FjR0PHzxrU9y8oy6ORGXAmchvCjl67K GNzL5XDQ00mkuNdGNWLMbyLxxpzJYZTTg+3yF1xPAgqKhwcS06pn/ooDw= 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=+2CKN8jJksxt09fxnDClYJ3quUQ=; b=cbgWM3M00HMGWrXMSs0y ssaHNfgAKB/LzaPUyMDvF9orDl0TmWMhCa5xt6AkAwSq7pexKdZ9ZrzlQ52uvKnk lcfyBTXWD0om6K+GqMLxXkc9jfEiqBTCHFL65kdyhDN6TxuDF6NPbTnTGKtckigO zNHU8XBk1wiS4JwFxB5+3Vc= Received: (qmail 17084 invoked by alias); 21 Oct 2015 17:01:03 -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 17073 invoked by uid 89); 21 Oct 2015 17:01:02 -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 X-HELO: mail-qg0-f51.google.com Received: from mail-qg0-f51.google.com (HELO mail-qg0-f51.google.com) (209.85.192.51) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (AES128-GCM-SHA256 encrypted) ESMTPS; Wed, 21 Oct 2015 17:01:00 +0000 Received: by qgem9 with SMTP id m9so32810564qge.1 for ; Wed, 21 Oct 2015 10:00:58 -0700 (PDT) X-Received: by 10.140.93.1 with SMTP id c1mr12396929qge.59.1445446857782; Wed, 21 Oct 2015 10:00:57 -0700 (PDT) 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 r75sm3597827qkl.35.2015.10.21.10.00.57 (version=TLSv1.2 cipher=ECDHE-RSA-AES128-GCM-SHA256 bits=128/128); Wed, 21 Oct 2015 10:00:57 -0700 (PDT) To: GCC Patches From: Nathan Sidwell Subject: [gomp4] loop nesting check Message-ID: <5627C4C8.1060308@acm.org> Date: Wed, 21 Oct 2015 13:00:56 -0400 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:38.0) Gecko/20100101 Thunderbird/38.3.0 MIME-Version: 1.0 This is the gomp4-branch variant of the loop nesting patch I just committed to trunk. The gomp4 branch had some checking, but a) it didn't catch all erroreous cases b) gave an ambiguous error, by not mentioning 'OpenACC' committed to gomp4 nathan 2015-10-21 Nathan Sidwell gcc/ * omp-low.c (check_omp_nesting_restrictions): Correctly check OpenACC loop nesting. gcc/testsuite/g * c-c++-common/goacc/non-routine.c: Adjust errors. * c-c++-common/goacc/loop-1.c: Adjust errors. * c-c++-common/goacc/clauses-fail.c: Adjust errors. * c-c++-common/goacc/sb-1.c: Adjust errors. * c-c++-common/goacc/sb-3.c: Adjust errors. * c-c++-common/goacc-gomp/nesting-fail-1.c: Adjust errors. * gfortran.dg/goacc/loop-4.f95 : Adjust errors. * gcc.dg/goacc/sb-3.c: Adjust errors. * gcc.dg/goacc/sb-1.c: Adjust errors. Index: gcc/omp-low.c =================================================================== --- gcc/omp-low.c (revision 229094) +++ gcc/omp-low.c (working copy) @@ -3106,13 +3106,42 @@ check_omp_nesting_restrictions (gimple * /* We split taskloop into task and nested taskloop in it. */ if (gimple_omp_for_kind (stmt) == GF_OMP_FOR_KIND_TASKLOOP) return true; - if (is_gimple_omp_oacc (stmt) && ctx == NULL - && get_oacc_fn_attrib (current_function_decl) == NULL) + if (gimple_omp_for_kind (stmt) == GF_OMP_FOR_KIND_OACC_LOOP) { - error_at (gimple_location (stmt), - "loop directive must be associated with a compute " - "region"); - return false; + bool ok = false; + + if (ctx) + switch (gimple_code (ctx->stmt)) + { + case GIMPLE_OMP_FOR: + ok = (gimple_omp_for_kind (ctx->stmt) + == GF_OMP_FOR_KIND_OACC_LOOP); + break; + + case GIMPLE_OMP_TARGET: + switch (gimple_omp_target_kind (ctx->stmt)) + { + case GF_OMP_TARGET_KIND_OACC_PARALLEL: + case GF_OMP_TARGET_KIND_OACC_KERNELS: + ok = true; + break; + + default: + break; + } + + default: + break; + } + else if (get_oacc_fn_attrib (current_function_decl)) + ok = true; + if (!ok) + { + error_at (gimple_location (stmt), + "OpenACC loop directive must be associated with" + " an OpenACC compute region"); + return false; + } } /* FALLTHRU */ case GIMPLE_CALL: Index: gcc/testsuite/c-c++-common/goacc/non-routine.c =================================================================== --- gcc/testsuite/c-c++-common/goacc/non-routine.c (revision 229094) +++ gcc/testsuite/c-c++-common/goacc/non-routine.c (working copy) @@ -8,7 +8,7 @@ main () { int i, v = 0; -#pragma acc loop gang reduction (+:v) /* { dg-error "loop directive must be associated with a compute region" } */ +#pragma acc loop gang reduction (+:v) /* { dg-error "loop directive must be associated with an OpenACC compute region" } */ for (i = 0; i < 10; i++) v++; Index: gcc/testsuite/c-c++-common/goacc/loop-1.c =================================================================== --- gcc/testsuite/c-c++-common/goacc/loop-1.c (revision 229094) +++ gcc/testsuite/c-c++-common/goacc/loop-1.c (working copy) @@ -36,16 +36,16 @@ int test1() i = d; a[i] = 1; } - #pragma acc loop /* { dg-error "loop directive must be associated with a compute region" } */ + #pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */ for (i = 1; i < 30; i++ ) if (i == 16) break; /* { dg-error "break statement used" } */ /* different types of for loop are allowed */ - #pragma acc loop /* { dg-error "loop directive must be associated with a compute region" } */ + #pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */ for (i = 1; i < 10; i++) { } - #pragma acc loop /* { dg-error "loop directive must be associated with a compute region" } */ + #pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */ for (i = 1; i < 10; i+=2) { a[i] = i; Index: gcc/testsuite/c-c++-common/goacc/clauses-fail.c =================================================================== --- gcc/testsuite/c-c++-common/goacc/clauses-fail.c (revision 229094) +++ gcc/testsuite/c-c++-common/goacc/clauses-fail.c (working copy) @@ -17,4 +17,4 @@ f (void) ; } -/* { dg-error "loop directive must be associated with a compute region" "" { target *-*-* } 15 } */ +/* { dg-error "loop directive must be associated with an OpenACC compute region" "" { target *-*-* } 15 } */ Index: gcc/testsuite/c-c++-common/goacc/sb-1.c =================================================================== --- gcc/testsuite/c-c++-common/goacc/sb-1.c (revision 229094) +++ gcc/testsuite/c-c++-common/goacc/sb-1.c (working copy) @@ -11,7 +11,7 @@ void foo() goto bad1; // { dg-error "invalid branch to/from OpenACC structured block" } #pragma acc data goto bad1; // { dg-error "invalid branch to/from OpenACC structured block" } - #pragma acc loop /* { dg-error "loop directive must be associated with a compute region" } */ + #pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */ for (l = 0; l < 2; ++l) goto bad1; // { dg-error "invalid branch to/from OpenACC structured block" } @@ -34,7 +34,7 @@ void foo() } goto bad2_loop; // { dg-error "invalid entry to OpenACC structured block" } - #pragma acc loop /* { dg-error "loop directive must be associated with a compute region" } */ + #pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */ for (l = 0; l < 2; ++l) { bad2_loop: ; @@ -64,7 +64,7 @@ void foo() { ok1_data: break; } } - #pragma acc loop /* { dg-error "loop directive must be associated with a compute region" } */ + #pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */ for (l = 0; l < 2; ++l) { int i; Index: gcc/testsuite/c-c++-common/goacc/sb-3.c =================================================================== --- gcc/testsuite/c-c++-common/goacc/sb-3.c (revision 229094) +++ gcc/testsuite/c-c++-common/goacc/sb-3.c (working copy) @@ -3,7 +3,7 @@ void f (void) { int i, j; -#pragma acc loop /* { dg-error "loop directive must be associated with a compute region" } */ +#pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */ for(i = 1; i < 30; i++) { if (i == 7) goto out; // { dg-error "invalid branch to/from OpenACC structured block" } Index: gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c =================================================================== --- gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c (revision 229094) +++ gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c (working copy) @@ -28,7 +28,7 @@ f_omp (void) #pragma acc update host(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */ #pragma acc enter data copyin(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */ #pragma acc exit data delete(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */ -#pragma acc loop /* { dg-error "may not be closely nested" } */ +#pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */ for (i = 0; i < 2; ++i) ; } @@ -63,7 +63,7 @@ f_omp (void) } #pragma omp section { -#pragma acc loop /* { dg-error "may not be closely nested" } */ +#pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */ for (i = 0; i < 2; ++i) ; } @@ -80,7 +80,7 @@ f_omp (void) #pragma acc update host(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */ #pragma acc enter data copyin(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */ #pragma acc exit data delete(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */ -#pragma acc loop /* { dg-error "may not be closely nested" } */ +#pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */ for (i = 0; i < 2; ++i) ; } @@ -96,7 +96,7 @@ f_omp (void) #pragma acc update host(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */ #pragma acc enter data copyin(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */ #pragma acc exit data delete(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */ -#pragma acc loop /* { dg-error "may not be closely nested" } */ +#pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */ for (i = 0; i < 2; ++i) ; } @@ -112,7 +112,7 @@ f_omp (void) #pragma acc update host(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */ #pragma acc enter data copyin(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */ #pragma acc exit data delete(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */ -#pragma acc loop /* { dg-error "may not be closely nested" } */ +#pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */ for (i = 0; i < 2; ++i) ; } @@ -128,7 +128,7 @@ f_omp (void) #pragma acc update host(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */ #pragma acc enter data copyin(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */ #pragma acc exit data delete(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */ -#pragma acc loop /* { dg-error "may not be closely nested" } */ +#pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */ for (i = 0; i < 2; ++i) ; } @@ -144,7 +144,7 @@ f_omp (void) #pragma acc update host(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */ #pragma acc enter data copyin(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */ #pragma acc exit data delete(i) /* { dg-error "OpenACC construct inside of non-OpenACC region" } */ -#pragma acc loop /* { dg-error "may not be closely nested" } */ +#pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */ for (i = 0; i < 2; ++i) ; } @@ -160,7 +160,7 @@ f_omp (void) #pragma acc update host(i) /* { dg-error "OpenACC update construct inside of OpenMP target region" } */ #pragma acc enter data copyin(i) /* { dg-error "OpenACC enter/exit data construct inside of OpenMP target region" } */ #pragma acc exit data delete(i) /* { dg-error "OpenACC enter/exit data construct inside of OpenMP target region" } */ -#pragma acc loop +#pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */ for (i = 0; i < 2; ++i) ; } Index: gcc/testsuite/gfortran.dg/goacc/loop-4.f95 =================================================================== --- gcc/testsuite/gfortran.dg/goacc/loop-4.f95 (revision 229094) +++ gcc/testsuite/gfortran.dg/goacc/loop-4.f95 (working copy) @@ -1,7 +1,7 @@ ! Ensure that loops not affiliated with acc compute regions cause an error. subroutine test1 - !$acc loop gang ! { dg-error "loop directive must be associated with a compute region" } + !$acc loop gang ! { dg-error "loop directive must be associated with an OpenACC compute region" } DO i = 1,10 ENDDO end subroutine test1 Index: gcc/testsuite/gcc.dg/goacc/sb-3.c =================================================================== --- gcc/testsuite/gcc.dg/goacc/sb-3.c (revision 229094) +++ gcc/testsuite/gcc.dg/goacc/sb-3.c (working copy) @@ -1,7 +1,7 @@ void f (void) { int i, j; -#pragma acc loop /* { dg-error "loop directive must be associated with a compute region" } */ +#pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */ for(i = 1; i < 30; i++) { if (i == 7) goto out; // { dg-error "invalid branch to/from OpenACC structured block" } Index: gcc/testsuite/gcc.dg/goacc/sb-1.c =================================================================== --- gcc/testsuite/gcc.dg/goacc/sb-1.c (revision 229094) +++ gcc/testsuite/gcc.dg/goacc/sb-1.c (working copy) @@ -9,7 +9,7 @@ void foo() goto bad1; // { dg-error "invalid branch to/from OpenACC structured block" } #pragma acc data goto bad1; // { dg-error "invalid branch to/from OpenACC structured block" } - #pragma acc loop /* { dg-error "loop directive must be associated with a compute region" } */ + #pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */ for (l = 0; l < 2; ++l) goto bad1; // { dg-error "invalid branch to/from OpenACC structured block" } @@ -32,7 +32,7 @@ void foo() } goto bad2_loop; // { dg-error "invalid entry to OpenACC structured block" } - #pragma acc loop /* { dg-error "loop directive must be associated with a compute region" } */ + #pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */ for (l = 0; l < 2; ++l) { bad2_loop: ; @@ -62,7 +62,7 @@ void foo() { ok1_data: break; } } - #pragma acc loop /* { dg-error "loop directive must be associated with a compute region" } */ + #pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */ for (l = 0; l < 2; ++l) { int i;