From patchwork Wed Jul 29 12:57:25 2020 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Tobias Burnus X-Patchwork-Id: 1338321 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=2620:52:3:1:0:246e:9693:128c; helo=sourceware.org; envelope-from=gcc-patches-bounces@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Received: from sourceware.org (server2.sourceware.org [IPv6:2620:52:3:1:0:246e:9693:128c]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature RSA-PSS (4096 bits) server-digest SHA256) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 4BGtqn4z5hz9sRW for ; Wed, 29 Jul 2020 22:57:41 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 133D938708D3; Wed, 29 Jul 2020 12:57:37 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa1.mentor.iphmx.com (esa1.mentor.iphmx.com [68.232.129.153]) by sourceware.org (Postfix) with ESMTPS id 930123842408; Wed, 29 Jul 2020 12:57:33 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org 930123842408 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=Tobias_Burnus@mentor.com IronPort-SDR: sSqGsOLz1AjH5x9e2BWiYWfnRNfU6ZCXXcBud9nETbNEaq7wXxjamkQp9+Bx9kS07lL0gsMNur KlAfrqMN+9DyxXdCZkEmw0zvh8eWHXBHy20pVzKlsnuRuIr2vUdja0RMmnV2xjPW2LpFXszfse 85tP3dqih4Dab7X171bCn0m0e9zQra6MGBg6vL+RDFtrMv2x7mHyOVEZV513+TPjstM8W9OP3S IgEVkZdWJHQGGtumfN9RRxq3sFeCs687WELDCy0KNbcklFkQAA7+j3qlHi4Rt28kj9zeMHxBt/ KbI= X-IronPort-AV: E=Sophos;i="5.75,410,1589270400"; d="diff'?scan'208";a="53568027" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa1.mentor.iphmx.com with ESMTP; 29 Jul 2020 04:57:32 -0800 IronPort-SDR: UyxvDoQkLmUID3908Y5zh2CSBNdVfNhZSOetyo9EzD2JBkFApQK0taLdvvlmPK4k4aQ7dKIwHf vOGhqrwu6y7kvybVOAgK5s+y0e+OsjgIYl5TGxACM5sg9+hCz7aup1QhINxEareERNhgNzIboE UOef1vHV3AwMqf8s97wwNt9vHCWBRaPtVsWNCCKsshWbQLjus6JRSSKTaer/FXymHv3GhrISrs Mn51LxHnpJ0vTlxy00/8IGGCFDc+6WHH9krGyJ6V6aVM3Hos5SJCMOuktN6XSOXO5SYYapROvv mVo= To: gcc-patches , fortran , Jakub Jelinek From: Tobias Burnus Subject: [Patch] OpenMP: Permit in Fortran omp target data without map Message-ID: <24e69703-ad51-3c06-66d1-cacea7728abc@codesourcery.com> Date: Wed, 29 Jul 2020 14:57:25 +0200 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:68.0) Gecko/20100101 Thunderbird/68.10.0 MIME-Version: 1.0 Content-Language: en-US X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-02.mgc.mentorg.com (139.181.222.2) To svr-ies-mbx-01.mgc.mentorg.com (139.181.222.1) X-Spam-Status: No, score=-13.0 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, SPF_HELO_PASS, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.2 X-Spam-Checker-Version: SpamAssassin 3.4.2 (2018-09-13) 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@gcc.gnu.org Sender: "Gcc-patches" I have missed this OpenMP 5 feature already a couple of times myself ... (The other related OpenMP 5 feature is that 'map(x) use_device_ptr(x)' is permitted on the same directive; but that already works and is tested for in the large use_device*.f90 libgomp tests.) OK? Tobias ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter OpenMP: Permit in Fortran omp target data without map gcc/fortran/ChangeLog: * openmp.c (resolve_omp_clauses): Permit 'omp target data' without map if use_device_{addr,ptr} is present. gcc/testsuite/ChangeLog: * gfortran.dg/gomp/map-3.f90: New test. * gfortran.dg/gomp/map-4.f90: New test. gcc/fortran/openmp.c | 12 +++++++--- gcc/testsuite/gfortran.dg/gomp/map-3.f90 | 38 ++++++++++++++++++++++++++++++++ gcc/testsuite/gfortran.dg/gomp/map-4.f90 | 7 ++++++ 3 files changed, 54 insertions(+), 3 deletions(-) diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c index 0fd998839b2..16f39a4e086 100644 --- a/gcc/fortran/openmp.c +++ b/gcc/fortran/openmp.c @@ -5330,17 +5330,23 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, if (omp_clauses->depend_source && code->op != EXEC_OMP_ORDERED) gfc_error ("SOURCE dependence type only allowed " "on ORDERED directive at %L", &code->loc); - if (!openacc && code && omp_clauses->lists[OMP_LIST_MAP] == NULL) + if (!openacc + && code + && omp_clauses->lists[OMP_LIST_MAP] == NULL + && omp_clauses->lists[OMP_LIST_USE_DEVICE_PTR] == NULL + && omp_clauses->lists[OMP_LIST_USE_DEVICE_ADDR] == NULL) { const char *p = NULL; switch (code->op) { - case EXEC_OMP_TARGET_DATA: p = "TARGET DATA"; break; case EXEC_OMP_TARGET_ENTER_DATA: p = "TARGET ENTER DATA"; break; case EXEC_OMP_TARGET_EXIT_DATA: p = "TARGET EXIT DATA"; break; default: break; } - if (p) + if (code->op == EXEC_OMP_TARGET_DATA) + gfc_error ("TARGET DATA must contain at least one MAP, USE_DEVICE_PTR, " + "or USE_DEVICE_ADDR clause at %L", &code->loc); + else if (p) gfc_error ("%s must contain at least one MAP clause at %L", p, &code->loc); } diff --git a/gcc/testsuite/gfortran.dg/gomp/map-3.f90 b/gcc/testsuite/gfortran.dg/gomp/map-3.f90 new file mode 100644 index 00000000000..13f63647bda --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/map-3.f90 @@ -0,0 +1,38 @@ +! { dg-additional-options "-fdump-tree-original" } + +subroutine bar +integer, target :: x +integer, allocatable, target :: y(:,:), z(:,:) +x = 7 +!$omp target enter data map(to:x) + +x = 8 +!$omp target data map(always, to: x) +call foo(x) +!$omp end target data + +!$omp target data use_device_ptr(x) +call foo2(x) +!$omp end target data + +!$omp target data use_device_addr(x) +call foo2(x) +!$omp end target data +!$omp target exit data map(release:x) + +!$omp target data map(y) use_device_addr(y) +call foo3(y) +!$omp end target data + +!$omp target data map(z) use_device_ptr(z) +call foo3(z) +!$omp end target data +end + +! { dg-final { scan-tree-dump-times "#pragma omp target enter data map\\(to:x\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(always,to:x\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "#pragma omp target data use_device_ptr\\(x\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "#pragma omp target data use_device_addr\\(x\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "#pragma omp target exit data map\\(release:x\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(tofrom:\\*\\(c_char \\*\\) y.data \\\[len: .*\\) map\\(to:y \\\[pointer set, len: .*\\) map\\(alloc:.*y.data \\\[pointer assign, bias: 0\\\]\\) use_device_addr\\(y\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(tofrom:\\*\\(c_char \\*\\) z.data \\\[len: .*\\) map\\(to:z \\\[pointer set, len: .*\\) map\\(alloc:.*z.data \\\[pointer assign, bias: 0\\\]\\) use_device_ptr\\(z\\)" 1 "original" } } diff --git a/gcc/testsuite/gfortran.dg/gomp/map-4.f90 b/gcc/testsuite/gfortran.dg/gomp/map-4.f90 new file mode 100644 index 00000000000..3aa1c8096d7 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/map-4.f90 @@ -0,0 +1,7 @@ +!$omp target enter data device(1) if (.true.) nowait ! { dg-error "TARGET ENTER DATA must contain at least one MAP clause" } + +!$omp target data device(1) ! { dg-error "TARGET DATA must contain at least one MAP, USE_DEVICE_PTR, or USE_DEVICE_ADDR clause" } +!$omp endtarget data + +!$omp target exit data device(1) if (.true.) nowait ! { dg-error "TARGET EXIT DATA must contain at least one MAP clause" } +end