From patchwork Wed Oct 21 15:59:03 2015 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Nathan Sidwell X-Patchwork-Id: 533917 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 1D9C7141358 for ; Thu, 22 Oct 2015 02:59:17 +1100 (AEDT) Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b=PNzSrKmk; 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=DzRQ8Y7tEipq0tBavtX3XUXWTAVSyg4bKpr7YQYwA6g8vEqgKu E76ydgu6wI66cuonPU/6G287xvrCsOgirQEH4i8yvcpCM3yCxfIAKPEkFJuSqK7A RMJVlcTLX9zuKS4aK31CPprFJzx8L0IcjqAApMsbsW8rRWu8TkkG/w3M8= 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=FsBcvxcd0ozmB73agZnfYz6w4Cw=; b=PNzSrKmkFfN+WAV6z+C/ dpz+C+CXMzDdr3/lWIyOlSKXL2qijSdQCCtxAqh1YOGYsuHdG7AdKvS/5ab4bOdX CnwFBJlEzFeU1f+ya/Npa0KVClYgt5O3xM+G2CvADxYiNDjf3ZXdtaQ61ITNn8YI YSkQi9J/GKaecVi1P6Po91w= Received: (qmail 13479 invoked by alias); 21 Oct 2015 15:59:09 -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 13467 invoked by uid 89); 21 Oct 2015 15:59:08 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.8 required=5.0 tests=BAYES_00, FREEMAIL_FROM, KAM_ASCII_DIVIDERS, RCVD_IN_DNSWL_LOW, SPF_PASS autolearn=no version=3.3.2 X-HELO: mail-qk0-f169.google.com Received: from mail-qk0-f169.google.com (HELO mail-qk0-f169.google.com) (209.85.220.169) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (AES128-GCM-SHA256 encrypted) ESMTPS; Wed, 21 Oct 2015 15:59:06 +0000 Received: by qkca6 with SMTP id a6so40524692qkc.3 for ; Wed, 21 Oct 2015 08:59:04 -0700 (PDT) X-Received: by 10.140.28.11 with SMTP id 11mr12153864qgy.94.1445443144474; Wed, 21 Oct 2015 08:59:04 -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 o76sm3471350qhb.25.2015.10.21.08.59.03 (version=TLSv1.2 cipher=ECDHE-RSA-AES128-GCM-SHA256 bits=128/128); Wed, 21 Oct 2015 08:59:03 -0700 (PDT) To: GCC Patches Cc: Jakub Jelinek , Bernd Schmidt From: Nathan Sidwell Subject: [OpenACC] loop nesting check Message-ID: <5627B647.1030407@acm.org> Date: Wed, 21 Oct 2015 11:59:03 -0400 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:38.0) Gecko/20100101 Thunderbird/38.3.0 MIME-Version: 1.0 Hi, in preparing a patch set of the OpenACC execution model for trunk, I discovered some broken tests in the testsuite. We were failing to diagnose some incorrectly nested loops, and that led to ICEs in my patch set. This patch implements a check for openacc loop directives, checking that they are immediately inside one of: 1) another openacc loop 2) an openacc offload region 3) an openacc routine The broken tests are amended with the now expected diagnostic. tested on x86_64-linux-gnu with nvptx accelerator. ok for trunk? nathan 2015-10-21 Nathan Sidwell gcc/ * omp-low.c (check_omp_nesting_restrictions): Check OpenACC loop nesting. testsuite/ * 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/loop-1.c: Adjust errors. * c-c++-common/goacc/nesting-1.c: Adjust errors. * c-c++-common/goacc-gomp/nesting-fail-1.c: Adjust errors. * c-c++-common/goacc-gomp/nesting-1.c: Adjust errors. Index: gcc/omp-low.c =================================================================== --- gcc/omp-low.c (revision 229101) +++ gcc/omp-low.c (working copy) @@ -3178,6 +3200,43 @@ 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 (gimple_omp_for_kind (stmt) == GF_OMP_FOR_KIND_OACC_LOOP) + { + 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: if (is_gimple_call (stmt) Index: gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c =================================================================== --- gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c (revision 229101) +++ gcc/testsuite/c-c++-common/goacc-gomp/nesting-1.c (working copy) @@ -5,7 +5,7 @@ f_omp_parallel (void) { int i; -#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/c-c++-common/goacc-gomp/nesting-fail-1.c =================================================================== --- gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c (revision 229101) +++ 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) ; } @@ -379,6 +379,7 @@ f_acc_data (void) void f_acc_loop (void) { +#pragma acc parallel #pragma acc loop for (i = 0; i < 2; ++i) { @@ -386,6 +387,7 @@ f_acc_loop (void) ; } +#pragma acc parallel #pragma acc loop for (i = 0; i < 2; ++i) { @@ -394,6 +396,7 @@ f_acc_loop (void) ; } +#pragma acc parallel #pragma acc loop for (i = 0; i < 2; ++i) { @@ -403,6 +406,7 @@ f_acc_loop (void) } } +#pragma acc parallel #pragma acc loop for (i = 0; i < 2; ++i) { @@ -410,6 +414,7 @@ f_acc_loop (void) ; } +#pragma acc parallel #pragma acc loop for (i = 0; i < 2; ++i) { @@ -417,6 +422,7 @@ f_acc_loop (void) ; } +#pragma acc parallel #pragma acc loop for (i = 0; i < 2; ++i) { @@ -424,6 +430,7 @@ f_acc_loop (void) ; } +#pragma acc parallel #pragma acc loop for (i = 0; i < 2; ++i) { @@ -431,6 +438,7 @@ f_acc_loop (void) ; } +#pragma acc parallel #pragma acc loop for (i = 0; i < 2; ++i) { @@ -438,6 +446,7 @@ f_acc_loop (void) i = 0; /* { dg-error "non-OpenACC construct inside of OpenACC region" } */ } +#pragma acc parallel #pragma acc loop for (i = 0; i < 2; ++i) { @@ -445,6 +454,7 @@ f_acc_loop (void) ; } +#pragma acc parallel #pragma acc loop for (i = 0; i < 2; ++i) { Index: gcc/testsuite/c-c++-common/goacc/nesting-1.c =================================================================== --- gcc/testsuite/c-c++-common/goacc/nesting-1.c (revision 229101) +++ gcc/testsuite/c-c++-common/goacc/nesting-1.c (working copy) @@ -58,7 +58,7 @@ f_acc_data (void) #pragma acc exit data delete(i) -#pragma acc loop +#pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */ for (i = 0; i < 2; ++i) ; @@ -93,7 +93,7 @@ f_acc_data (void) #pragma acc exit data delete(i) -#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/c-c++-common/goacc/loop-1.c =================================================================== --- gcc/testsuite/c-c++-common/goacc/loop-1.c (revision 229101) +++ gcc/testsuite/c-c++-common/goacc/loop-1.c (working copy) @@ -38,16 +38,16 @@ int test1() i = d; a[i] = 1; } - #pragma acc loop + #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 + #pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */ for (i = 1; i < 10; i++) { } - #pragma acc loop + #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 229101) +++ gcc/testsuite/c-c++-common/goacc/clauses-fail.c (working copy) @@ -12,6 +12,7 @@ f (void) #pragma acc data two /* { dg-error "expected '#pragma acc' clause before 'two'" } */ ; +#pragma acc parallel #pragma acc loop deux /* { dg-error "expected '#pragma acc' clause before 'deux'" } */ for (i = 0; i < 2; ++i) ; Index: gcc/testsuite/c-c++-common/goacc/sb-1.c =================================================================== --- gcc/testsuite/c-c++-common/goacc/sb-1.c (revision 229101) +++ 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 + #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 + #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 + #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 229101) +++ gcc/testsuite/c-c++-common/goacc/sb-3.c (working copy) @@ -3,11 +3,11 @@ void f (void) { int i, j; -#pragma acc loop +#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" } -#pragma acc loop // { dg-error "work-sharing region may not be closely nested inside of work-sharing, critical, ordered, master or explicit task region" } +#pragma acc loop for(j = 5; j < 10; j++) { if (i == 6 && j == 7) goto out; // { dg-error "invalid branch to/from OpenACC structured block" }