From patchwork Fri Sep 20 15:47:59 2019 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Tobias Burnus X-Patchwork-Id: 1165306 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-509367-incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: ozlabs.org; dkim=pass (1024-bit key; unprotected) header.d=gcc.gnu.org header.i=@gcc.gnu.org header.b="opAOXAnB"; 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 46ZdSH1lnZz9s00 for ; Sat, 21 Sep 2019 01:49:17 +1000 (AEST) 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=tZMrQL6kNe53EkESER7M0wRw7hFuO9ggoHuzkuPEGpwVyuYkcI tQ+Wg/Kqk5ZnER5IUke8pbMeELwOlLDa95JtcKWdmpP5HIWK/aMojWmlMZEZVsUa kEe9ikEXlSC2rW9ys4X+T1+Q0Jf7M5gusiMelW07sfgA7DcAL5QEZ9EKk= 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=L8xspSutoV/h00j4tH2vKegCoDQ=; b=opAOXAnBbGJPzYOd1/Ts zkEBCH5EEFSAKMkynjnXAlmWj8mShbl4Y9kJpD5OiL+ZxkZBkUfAv7XN3tXMcsYB 0zzMlUp7rTT6wOuUG2KUKxZCxbhE39FT76QSIKUgj7VjkYqvnvADLkh4rf4aHhc9 qIKFqHUtZXdI219E5py8a5A= Received: (qmail 15776 invoked by alias); 20 Sep 2019 15:49:02 -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 15493 invoked by uid 89); 20 Sep 2019 15:48:41 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-24.3 required=5.0 tests=AWL, BAYES_00, GIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3, RCVD_IN_DNSWL_LOW, SPF_PASS autolearn=ham version=3.3.1 spammy=wrongly, sk:fdumpt, sk:fdump-t, noted X-HELO: esa1.mentor.iphmx.com Received: from esa1.mentor.iphmx.com (HELO esa1.mentor.iphmx.com) (68.232.129.153) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Fri, 20 Sep 2019 15:48:21 +0000 IronPort-SDR: oLlEr/N1WbtZp/71ygR3mIalFyMzPr3VLnCF72pcnJsHXYdrmzzuFCaLsSbYUT2HbdGuEJWA6J ke6nWkiyxSJwAOCdIKMQdtIUmI7tFfnTIaLuSbdGlaAVvYMRFpOkNVrvKo70idSbn3Pq1jfxqg UGcZDwQHjWYaHsF5HS42ALzbw2l0M1Lbd+kM8Bo3hujEd+901TIjW2VikdWB3wxGJM9VAjHrNa kd7L7iG1fzvAUaV5zauNEeaA+2BuTvTDH2+raPEXnmHr0nl/Yiyh2Sg/5ELLbJ4wwtTl/XApDW lb8= Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa1.mentor.iphmx.com with ESMTP; 20 Sep 2019 07:48:12 -0800 IronPort-SDR: hVDVH8BpZDMUaOFYoJO0HTV0sSRP29LILPVFMq+JIj0pRlPHlY0NdHZpF1s0nK/4/b6zfOzkWI v3afnAy4vDMC4BFpDhDN9sphj/yTPIyc/jhrGfv9Z01Y4OuZknmByCqzUpfIpR+KA6iBRDBKsf WEelrQhTjri7g9rEX8ButcbcSVfHmDW1gwxKPHalb2TPyrNnAHc3A5qIZZMrvsA2S74kKtNzKr UdQafr4OMXn/FwsqDl+EXiCgHYingEE8rMSevT3OOIlDIi4A2oh5bd+Xrmi/1f6Pbn+58oZRQ0 nRI= To: Jakub Jelinek , gcc-patches , fortran CC: Thomas Schwinge From: Tobias Burnus Subject: [Patch] PR fortran/78260 - OpenACC + OpenMP target fixes - esp. with function-result variables Message-ID: <64ef5bde-759c-9a41-985f-6f1ccb02529d@codesourcery.com> Date: Fri, 20 Sep 2019 17:47:59 +0200 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:68.0) Gecko/20100101 Thunderbird/68.1.0 MIME-Version: 1.0 X-IsSubscribed: yes Hi all, This patch does two things: (A) For OpenACC, only, it fixes the is-variable check. That check missed to reject module names (as noted in the PR) but as my testing showed, it also wrongly rejected function-result variables. (i.e. where the return-value variable has the same name as the function). - For the invalid input of the PR, gfortran gave an ICE in the gimplifier. (B) Using such function-result variables did not work properly. OpenACC used in both cases (see pr78260-2.f90) the function name – and at least one variant failed with an ICE. OpenMP used the result variable for "target data map" but not for "target update". Additionally "task depend" had the same issue. Bootstrapped and regtested on x86_64-gnu-linux w/o accelerator. I intent to build/regtest it also applied to the OG9 (openacc-gnu-9) branch and run the test case with actual nvptx+AMDGCN offloading, but I have not done so, yet. OK for the trunk? Tobias PS: Regtesting fails for continuation_6.f but that's PR fortran/91253 (fails to show a warning when a newer GLIBC is used). And for gfortran.dg/vect/vect-8.f90 fails because 23 instead of 22 loops get vectorized - probably someone (richi?) didn't update the expected value. 2019-09-20 Tobias Burnus PR fortran/78260 * openmp.c (gfc_resolve_oacc_declare): Reject all non variables but accept function result variables. * trans-openmp.c (gfc_trans_omp_clauses): Handle function-result variables for remaing cases. 2019-09-20 Tobias Burnus PR fortran/78260 * gfortran.dg/goacc/parameter.f95: Change dg-error as it is now detected earlier. * gfortran.dg/goacc/pr85701.f90: Modify to use a separate result variable. * gfortran.dg/goacc/pr78260.f90: New. * gfortran.dg/goacc/pr78260-2.f90: New. * gfortran.dg/gomp/pr78260.f90: New. * gfortran.dg/gomp/pr78260-2.f90: New. * gfortran.dg/gomp/pr78260-3.f90: New. diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c index 44fcb9db8c6..bda7f288989 100644 --- a/gcc/fortran/openmp.c +++ b/gcc/fortran/openmp.c @@ -6048,18 +6048,14 @@ gfc_resolve_oacc_declare (gfc_namespace *ns) for (n = oc->clauses->lists[list]; n; n = n->next) { n->sym->mark = 0; - if (n->sym->attr.function || n->sym->attr.subroutine) + if (n->sym->attr.flavor != FL_VARIABLE + && (n->sym->attr.flavor != FL_PROCEDURE + || n->sym->result != n->sym)) { gfc_error ("Object %qs is not a variable at %L", n->sym->name, &oc->loc); continue; } - if (n->sym->attr.flavor == FL_PARAMETER) - { - gfc_error ("PARAMETER object %qs is not allowed at %L", - n->sym->name, &oc->loc); - continue; - } if (n->expr && n->expr->ref->type == REF_ARRAY) { diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c index 8eae7bc0a52..b4c77aebf4d 100644 --- a/gcc/fortran/trans-openmp.c +++ b/gcc/fortran/trans-openmp.c @@ -2075,7 +2075,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, tree node = build_omp_clause (input_location, OMP_CLAUSE_DEPEND); if (n->expr == NULL || n->expr->ref->u.ar.type == AR_FULL) { - tree decl = gfc_get_symbol_decl (n->sym); + tree decl = gfc_trans_omp_variable (n->sym, false); if (gfc_omp_privatize_by_reference (decl)) decl = build_fold_indirect_ref (decl); if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl))) @@ -2136,7 +2136,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, tree node2 = NULL_TREE; tree node3 = NULL_TREE; tree node4 = NULL_TREE; - tree decl = gfc_get_symbol_decl (n->sym); + tree decl = gfc_trans_omp_variable (n->sym, false); if (DECL_P (decl)) TREE_ADDRESSABLE (decl) = 1; if (n->expr == NULL || n->expr->ref->u.ar.type == AR_FULL) @@ -2398,7 +2398,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, tree node = build_omp_clause (input_location, clause_code); if (n->expr == NULL || n->expr->ref->u.ar.type == AR_FULL) { - tree decl = gfc_get_symbol_decl (n->sym); + tree decl = gfc_trans_omp_variable (n->sym, false); if (gfc_omp_privatize_by_reference (decl)) decl = build_fold_indirect_ref (decl); else if (DECL_P (decl)) diff --git a/gcc/testsuite/gfortran.dg/goacc/parameter.f95 b/gcc/testsuite/gfortran.dg/goacc/parameter.f95 index 84274611915..cbe67dba788 100644 --- a/gcc/testsuite/gfortran.dg/goacc/parameter.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/parameter.f95 @@ -6,7 +6,7 @@ contains implicit none integer :: i integer, parameter :: a = 1 - !$acc declare device_resident (a) ! { dg-error "PARAMETER" } + !$acc declare device_resident (a) ! { dg-error "is not a variable" } !$acc data copy (a) ! { dg-error "not a variable" } !$acc end data !$acc data deviceptr (a) ! { dg-error "not a variable" } diff --git a/gcc/testsuite/gfortran.dg/goacc/pr78260-2.f90 b/gcc/testsuite/gfortran.dg/goacc/pr78260-2.f90 new file mode 100644 index 00000000000..e28564d6f70 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/pr78260-2.f90 @@ -0,0 +1,20 @@ +! { dg-do compile } +! { dg-options "-fopenacc -fdump-tree-original" } +! { dg-require-effective-target fopenacc } + +! PR fortran/78260 + +module m + implicit none + integer :: n = 0 +contains + integer function f1() + !$acc declare present(f1) + !$acc kernels copyin(f1) + f1 = 5 + !$acc end kernels + end function f1 +end module m +! { dg-final { scan-tree-dump-times "#pragma acc data map\\(force_present:__result_f1\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "#pragma acc data map\\(force_present:__result_f1\\)" 1 "original" } } + diff --git a/gcc/testsuite/gfortran.dg/goacc/pr78260.f90 b/gcc/testsuite/gfortran.dg/goacc/pr78260.f90 new file mode 100644 index 00000000000..21bde854919 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/pr78260.f90 @@ -0,0 +1,36 @@ +! { dg-do compile } +! { dg-options "-fopenacc" } +! { dg-require-effective-target fopenacc } + +! PR fortran/78260 +! Contributed by Gerhard Steinmetz + +module m + implicit none + integer :: n = 0 +contains + subroutine s + !$acc declare present(m) ! { dg-error "Object .m. is not a variable" } + !$acc kernels copyin(m) ! { dg-error "Object .m. is not a variable" } + n = n + 1 + !$acc end kernels + end subroutine s + subroutine s2 + !$acc declare present(s2) ! { dg-error "Object .s2. is not a variable" } + !$acc kernels copyin(s2) ! { dg-error "Object .s2. is not a variable" } + n = n + 1 + !$acc end kernels + end subroutine s2 + integer function f1() + !$acc declare present(f1) ! OK, f1 is also the result variable + !$acc kernels copyin(f1) ! OK, f1 is also the result variable + f1 = 5 + !$acc end kernels + end function f1 + integer function f2() result(res) + !$acc declare present(f2) ! { dg-error "Object .f2. is not a variable" } + !$acc kernels copyin(f2) ! { dg-error "Object .f2. is not a variable" } + res = 5 + !$acc end kernels + end function f2 +end module m diff --git a/gcc/testsuite/gfortran.dg/goacc/pr85701.f90 b/gcc/testsuite/gfortran.dg/goacc/pr85701.f90 index 9c201b865b2..bae09de90ac 100644 --- a/gcc/testsuite/gfortran.dg/goacc/pr85701.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/pr85701.f90 @@ -9,11 +9,11 @@ subroutine s2 !$acc declare present(s2) ! { dg-error "is not a variable" } end -function f1 () +function f1 () result(res) !$acc declare copy(f1) ! { dg-error "is not a variable" } end -function f2 () +function f2 () result(res) !$acc declare present(f2) ! { dg-error "is not a variable" } end diff --git a/gcc/testsuite/gfortran.dg/gomp/pr78260-2.f90 b/gcc/testsuite/gfortran.dg/gomp/pr78260-2.f90 new file mode 100644 index 00000000000..c58ad93471c --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/pr78260-2.f90 @@ -0,0 +1,59 @@ +! { dg-do compile } +! { dg-options "-fopenmp -fdump-tree-original" } + +! PR fortran/78260 + +module m + implicit none + integer :: n = 0 +contains + integer function f1() + !$omp target data map(f1) + !$omp target update to(f1) + f1 = 5 + !$omp end target data + end function f1 + + integer function f2() + dimension :: f2(1) + !$omp target data map(f2) + !$omp target update to(f2) + f2(1) = 5 + !$omp end target data + end function f2 + + integer function f3() result(res) + dimension :: res(1) + !$omp target data map(res) + !$omp target update to(res) + res(1) = 5 + !$omp end target data + end function f3 + + integer function f4() result(res) + allocatable :: res + dimension :: res(:) + !$omp target data map(res) + !$omp target update to(res) + res = [5] + !$omp end target data + end function f4 + + subroutine sub() + integer, allocatable :: arr(:) + !$omp target data map(arr) + !$omp target update to(arr) + arr = [5] + !$omp end target data + end subroutine sub +end module m + +! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(tofrom:\\*\\(c_char \\*\\) arr.data \\\[len: D.\[0-9\]+ \\* 4\\\]\\) map\\(to:arr \\\[pointer set, len: ..\\\]\\) map\\(alloc:\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) arr.data \\\[pointer assign, bias: 0\\\]\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "#pragma omp target update to\\(\\*\\(c_char \\*\\) arr.data \\\[len: D.\[0-9\]+ \\* 4\\\]\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(tofrom:\\*\\(c_char \\*\\) __result->data \\\[len: D.\[0-9\]+ \\* 4\\\]\\) map\\(to:\\*__result \\\[pointer set, len: ..\\\]\\) map\\(alloc:\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) __result->data \\\[pointer assign, bias: 0\\\]\\) map\\(alloc:__result \\\[pointer assign, bias: 0\\\]\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "#pragma omp target update to\\(\\*\\(c_char \\*\\) __result->data \\\[len: D.\[0-9\]+ \\* 4\\\]\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(tofrom:\\*__result.0\\) map\\(alloc:__result.0 \\\[pointer assign, bias: 0\\\]\\)" 2 "original" } } +! { dg-final { scan-tree-dump-times "#pragma omp target update to\\(\\*__result.0\\)" 2 "original" } } +! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(tofrom:__result_f1\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "#pragma omp target update to\\(__result_f1\\)" 1 "original" } } + diff --git a/gcc/testsuite/gfortran.dg/gomp/pr78260-3.f90 b/gcc/testsuite/gfortran.dg/gomp/pr78260-3.f90 new file mode 100644 index 00000000000..4ca3e361a59 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/pr78260-3.f90 @@ -0,0 +1,74 @@ +! { dg-do compile } +! { dg-options "-fopenmp -fdump-tree-original" } + +! PR fortran/78260 + +integer function f1() + implicit none + + f1 = 0 + + !$omp task depend(inout:f1) + !$omp end task + + !$omp task depend(inout:f1) + !$omp end task +end function f1 + +integer function f2() + implicit none + dimension :: f2(1) + + f2(1) = 0 + + !$omp task depend(inout:f2) + !$omp end task + + !$omp task depend(inout:f2) + !$omp end task +end function f2 + +integer function f3() result(res) + implicit none + dimension :: res(1) + + res(1) = 0 + + !$omp task depend(inout:res) + !$omp end task + + !$omp task depend(inout:res) + !$omp end task +end function f3 + +integer function f4() result(res) + implicit none + allocatable :: res + dimension :: res(:) + + res = [0] + + !$omp task depend(inout:res) + !$omp end task + + !$omp task depend(inout:res) + !$omp end task +end function f4 + +subroutine sub() + implicit none + integer, allocatable :: arr(:) + + arr = [3] + + !$omp task depend(inout:arr) + !$omp end task + + !$omp task depend(inout:arr) + !$omp end task +end subroutine sub + +! { dg-final { scan-tree-dump-times "#pragma omp task depend\\(inout:__result_f1\\)" 2 "original" } } +! { dg-final { scan-tree-dump-times "#pragma omp task depend\\(inout:\\*__result.0\\)" 4 "original" } } +! { dg-final { scan-tree-dump-times "#pragma omp task depend\\(inout:\\*\\(c_char \\*\\) __result->data\\)" 2 "original" } } +! { dg-final { scan-tree-dump-times "#pragma omp task depend\\(inout:\\*\\(c_char \\*\\) arr.data\\)" 2 "original" } } diff --git a/gcc/testsuite/gfortran.dg/gomp/pr78260.f90 b/gcc/testsuite/gfortran.dg/gomp/pr78260.f90 new file mode 100644 index 00000000000..23acd4c1bf9 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/pr78260.f90 @@ -0,0 +1,33 @@ +! { dg-do compile } + +! PR fortran/78260 + +module m + implicit none + integer :: n = 0 +contains + subroutine s + !$omp target data map(m) ! { dg-error "Object .m. is not a variable" } + !$omp target update to(m) ! { dg-error "Object .m. is not a variable" } + n = n + 1 + !$omp end target data + end subroutine s + subroutine s2 + !$omp target data map(s2) ! { dg-error "Object .s2. is not a variable" } + !$omp target update to(s2) ! { dg-error "Object .s2. is not a variable" } + n = n + 1 + !$omp end target data + end subroutine s2 + integer function f1() + !$omp target data map(f1) ! OK, f1 is also the result variable + !$omp target update to(f1) ! OK, f1 is also the result variable + f1 = 5 + !$omp end target data + end function f1 + integer function f2() result(res) + !$omp target data map(f2) ! { dg-error "Object .f2. is not a variable" } + !$omp target update to(f2) ! { dg-error "Object .f2. is not a variable" } + res = 5 + !$omp end target data + end function f2 +end module m