From patchwork Fri Jan 27 09:19:42 2023 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Tobias Burnus X-Patchwork-Id: 63782 Return-Path: X-Original-To: patchwork@sourceware.org Delivered-To: patchwork@sourceware.org Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 544153858C00 for ; Fri, 27 Jan 2023 09:20:24 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa2.mentor.iphmx.com (esa2.mentor.iphmx.com [68.232.141.98]) by sourceware.org (Postfix) with ESMTPS id 5B4863858409; Fri, 27 Jan 2023 09:19:53 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 5B4863858409 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com X-IronPort-AV: E=Sophos;i="5.97,250,1669104000"; d="diff'?scan'208";a="95958822" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa2.mentor.iphmx.com with ESMTP; 27 Jan 2023 01:19:52 -0800 IronPort-SDR: 1CHdokXzE6m17jLJNZk2S7XhBdvHtnMhmFvtcIBPadj0ZW54uhc+F9bppzYbk35fOjQR+XpdC0 2xlWrbKg/DraN4L82JU1B5Ru01MOCijDndno3PURyHLRuqp4jY+PI3/kjk06lrw7xwn6s2iCVR 8n8JIfi3PS9vNDkPJHGjvw89/eXQ8aEc1LyiD8xhR8tTrB3GSZG6Ed0F2fknAQDCSew/8SbWSD omVu5PTmZ7CUocd2PEvsb5GomJLXURwwt/X/9UtsUsOIBAuqr21w1qXRfooXoWwVxHS8nSCzux nJo= Message-ID: Date: Fri, 27 Jan 2023 10:19:42 +0100 MIME-Version: 1.0 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:102.0) Gecko/20100101 Thunderbird/102.6.1 Content-Language: en-US To: gcc-patches , fortran , Jakub Jelinek From: Tobias Burnus Subject: [Patch] OpenMP/Fortran: Fix has_device_addr clause splitting [PR108558] X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-14.mgc.mentorg.com (139.181.222.14) To svr-ies-mbx-12.mgc.mentorg.com (139.181.222.12) X-Spam-Status: No, score=-11.4 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, RCVD_IN_MSPIKE_H2, SPF_HELO_PASS, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces+patchwork=sourceware.org@gcc.gnu.org Sender: "Gcc-patches" Rather obvious fix. Hence, I intent to commit it later as obvious, unless there are any comments. Tobias PS: Thanks goes to Thomas for finding + reporting the issue. ----------------- Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955 OpenMP/Fortran: Fix has_device_addr clause splitting [PR108558] gcc/fortran/ChangeLog: PR fortran/108558 * trans-openmp.cc (gfc_split_omp_clauses): Handle has_device_addr. libgomp/ChangeLog: PR fortran/108558 * testsuite/libgomp.fortran/has_device_addr.f90: New test. gcc/fortran/trans-openmp.cc | 2 + .../testsuite/libgomp.fortran/has_device_addr.f90 | 59 ++++++++++++++++++++++ 2 files changed, 61 insertions(+) diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc index 87213de0918..5283d0ce5f3 100644 --- a/gcc/fortran/trans-openmp.cc +++ b/gcc/fortran/trans-openmp.cc @@ -6205,6 +6205,8 @@ gfc_split_omp_clauses (gfc_code *code, = code->ext.omp_clauses->lists[OMP_LIST_MAP]; clausesa[GFC_OMP_SPLIT_TARGET].lists[OMP_LIST_IS_DEVICE_PTR] = code->ext.omp_clauses->lists[OMP_LIST_IS_DEVICE_PTR]; + clausesa[GFC_OMP_SPLIT_TARGET].lists[OMP_LIST_HAS_DEVICE_ADDR] + = code->ext.omp_clauses->lists[OMP_LIST_HAS_DEVICE_ADDR]; clausesa[GFC_OMP_SPLIT_TARGET].device = code->ext.omp_clauses->device; clausesa[GFC_OMP_SPLIT_TARGET].thread_limit diff --git a/libgomp/testsuite/libgomp.fortran/has_device_addr.f90 b/libgomp/testsuite/libgomp.fortran/has_device_addr.f90 new file mode 100644 index 00000000000..95cc7788f2d --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/has_device_addr.f90 @@ -0,0 +1,59 @@ +! { dg-additional-options "-fdump-tree-original" } + +! +! PR fortran/108558 +! + +! { dg-final { scan-tree-dump-times "#pragma omp target has_device_addr\\(x\\) has_device_addr\\(y\\)" 2 "original" } } +! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(tofrom:x\\) map\\(tofrom:y\\)" 2 "original" } } +! { dg-final { scan-tree-dump-times "#pragma omp target data use_device_addr\\(x\\) use_device_addr\\(y\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "#pragma omp target update from\\(y\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(tofrom:x\\) map\\(tofrom:y\\) use_device_addr\\(x\\) use_device_addr\\(y\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "#pragma omp teams" 2 "original" } } +! { dg-final { scan-tree-dump-times "#pragma omp distribute" 2 "original" } } +! { dg-final { scan-tree-dump-times "#pragma omp parallel" 2 "original" } } +! { dg-final { scan-tree-dump-times "#pragma omp for nowait" 2 "original" } } + +module m +contains +subroutine vectorAdd(x, y, N) + implicit none + integer :: N + integer(4) :: x(N), y(N) + integer :: i + + !$omp target teams distribute parallel do has_device_addr(x, y) + do i = 1, N + y(i) = x(i) + y(i) + end do +end subroutine vectorAdd +end module m + +program main + use m + implicit none + integer, parameter :: N = 9876 + integer(4) :: x(N), y(N) + integer :: i + + x(:) = 1 + y(:) = 2 + + !$omp target data map(x, y) + !$omp target data use_device_addr(x, y) + call vectorAdd(x, y, N) + !$omp end target data + !$omp target update from(y) + if (any (y /= 3)) error stop + !$omp end target data + + x = 1 + y = 2 + !$omp target data map(x, y) use_device_addr(x, y) + !$omp target teams distribute parallel do has_device_addr(x, y) + do i = 1, N + y(i) = x(i) + y(i) + end do + !$omp end target data + if (any (y /= 3)) error stop +end program