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: 49122 Return-Path: Delivered-To: ouuuleilei@gmail.com Received: by 2002:adf:eb09:0:0:0:0:0 with SMTP id s9csp735091wrn; Fri, 27 Jan 2023 01:20:37 -0800 (PST) X-Google-Smtp-Source: AMrXdXtQNIbK76AW1OYNoEUWvp1UE0wKxrber7AXeVKGVZKGj4dmPU3L6wQEEApqyNSM28wMzr1Y X-Received: by 2002:a17:906:25db:b0:877:6a03:9ad1 with SMTP id n27-20020a17090625db00b008776a039ad1mr28701724ejb.7.1674811237507; Fri, 27 Jan 2023 01:20:37 -0800 (PST) ARC-Seal: i=1; a=rsa-sha256; t=1674811237; cv=none; d=google.com; s=arc-20160816; b=FRIuE2orOSyoeJQ8vQbGUSVcQ7sS29jEWz5N00Ed/8KDqGi5KFsMjOaFGrqsoVrOH7 /6sd3j5JB/XH19ERnws+ErjbFD3YEManxmNW3yPb9jpDYU9/4DdL124gwAuKqXmVXwCN y8OHrvQLDHTMy5kjzK+RZFRWrMUt+kuWin9eGEMxTmThd/UpqXwNyCjCdGtWp2bN+Juq VH9/tebc+iZHZP7RjDgKdyf/DfioY1wpeRuxnrWGSIthYNYos4F7wmYPbyJBnPfDTHbA 27PYuCr+HYmqjnp1U05ivsPlgla5AeMsq9YMoDrf+6VEmk29GtzF9ZfErWbZSO6R4e+o FECA== ARC-Message-Signature: i=1; a=rsa-sha256; c=relaxed/relaxed; d=google.com; s=arc-20160816; h=sender:errors-to:list-subscribe:list-help:list-post:list-archive :list-unsubscribe:list-id:precedence:subject:from:to :content-language:user-agent:mime-version:date:message-id :ironport-sdr:dmarc-filter:delivered-to; bh=o1iI8wO6Fs8/hMw4F3IjbpWDWitO8OxnZlEOI7M1Vz4=; b=WhyNeiaeABk8kBZPRMoFULGsV/VnkGcCWJp2Cvf+8RCQfO7WHV/9Z29VOvzdFE0RFQ MVJM9kFK8eiGmOlpUgvnZJFvIYuyIU3YSUBTsQXxgKED9++eRhJQHkKPImlqR+7QvHSN Y3hochW2JVfvOxBjEQVFgrMNwdrEPukc9FLIG/IyrP68UF0jmz9rjqLw90s5kCSRxBGo ctTy86YSabA4Tq6ahy7NKKaaXCrFmrUsR77633LNmnYT2c2PYIeeSrS4lFRtV27oPfxl ytFFJii2IE9dgRqU4peRuv1WAgxecZQt1GC5Ue372uohWBpF11R4eUnWCAt1DzxZZw0J UlOA== ARC-Authentication-Results: i=1; mx.google.com; spf=pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 2620:52:3:1:0:246e:9693:128c as permitted sender) smtp.mailfrom="gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org" Received: from sourceware.org (server2.sourceware.org. [2620:52:3:1:0:246e:9693:128c]) by mx.google.com with ESMTPS id fw32-20020a170907502000b008776f852d5fsi3776854ejc.961.2023.01.27.01.20.37 for (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Fri, 27 Jan 2023 01:20:37 -0800 (PST) Received-SPF: pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 2620:52:3:1:0:246e:9693:128c as permitted sender) client-ip=2620:52:3:1:0:246e:9693:128c; Authentication-Results: mx.google.com; spf=pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 2620:52:3:1:0:246e:9693:128c as permitted sender) smtp.mailfrom="gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org" Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 02BEE3858436 for ; Fri, 27 Jan 2023 09:20:34 +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+ouuuleilei=gmail.com@gcc.gnu.org Sender: "Gcc-patches" X-getmail-retrieved-from-mailbox: =?utf-8?q?INBOX?= X-GMAIL-THRID: =?utf-8?q?1756166868160565897?= X-GMAIL-MSGID: =?utf-8?q?1756166868160565897?= 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