From patchwork Fri Oct 18 09:27:39 2019 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Tobias Burnus X-Patchwork-Id: 1179226 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=209.132.180.131; helo=sourceware.org; envelope-from=gcc-patches-return-511274-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="OHD3h+uf"; 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 46vggN6Ng6z9sPJ for ; Fri, 18 Oct 2019 20:27:58 +1100 (AEDT) DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:to :from:subject:message-id:date:mime-version:content-type; q=dns; s=default; b=HfLZk/T9lwGflTonbHLLJAfgcClyK8R9Gz8qynahvy8kCmm4o/ 4q5qRHlo0Li16nHZ/djhxfhseH8NmqQRd9hzweScHAh3dfuIh/3x1qSB3mphcajV 7CwsnXeag6ZoqdP0cx9hbW7e1vO0BF1vOeBXB0k32k3LH21Sw+3ppXWAk= 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 :from:subject:message-id:date:mime-version:content-type; s= default; bh=PB+naB6vDMcLCCy8LaAyuEf7vQY=; b=OHD3h+uflo8EDnpne1YF h19qJMaxB31UiOOEyTzG+ByFq5jgF+BJRWEoEtPpkd+mmIbws2YU6WL3bvArCPAg ZuzQzaWzw//fpxbV+clGS4Xf7X7VHDt+88XoyH4ADWXBNyjn9kj5xivknMso+EN7 gvtaFpH7fLzIUMmL09WONCs= Received: (qmail 10146 invoked by alias); 18 Oct 2019 09:27:49 -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 10120 invoked by uid 89); 18 Oct 2019 09:27:49 -0000 Authentication-Results: sourceware.org; auth=none X-Spam-SWARE-Status: No, score=-21.5 required=5.0 tests=AWL, BAYES_00, GIT_PATCH_0, GIT_PATCH_1, GIT_PATCH_2, GIT_PATCH_3, SPF_PASS autolearn=ham version=3.3.1 spammy=POINTER_TYPE_P, pointer_type_p, !$omp X-HELO: esa4.mentor.iphmx.com Received: from esa4.mentor.iphmx.com (HELO esa4.mentor.iphmx.com) (68.232.137.252) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Fri, 18 Oct 2019 09:27:47 +0000 IronPort-SDR: qHz7IvPJ2eG0aTq0iFtB5ZGNScVWxPDPXDaV3WOsS9ciAPMMb28Iq1fLf6Xb5uWo0RoCRRVzYW ZVxTjgZq4+a/AIYLUAuisI4wKQ96swLK0679/+JD/EHeajfHHjzWCHSpBEZ3tXi8uHB8gTRpXl 66wWafPigyx0HGugd5bOiZdTEdHzEN+KKyzD1FrxghncGnUNMCqYZ2xu7x4frZuN7xWtyubec7 uJF5pauPHb0HVmCGEl4iEh2gfTY0UJ40yH9Q+M8UFaaxOpeN4RGgwZsc7VUpsJprcV309HJrnQ Bbs= Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa4.mentor.iphmx.com with ESMTP; 18 Oct 2019 01:27:45 -0800 IronPort-SDR: WahTXHlZPJny6tlnaEVYb0wxPs+vTaErYO9POohCssYfSnOWO7K7qTqIwIv+l95KnDDWnsj8jQ 1aWOcbpYcRB2zGZH4liFbFyaKo0/pMg4urH+TT4vbUGQa/SoTc5xC6Pk827EBeoBG1OzaFC71v +GW3JJT9QNjl+AEBTY9Epr/gTE8WoUCsFKcIUT1X6lVfaK7Dt5Fj/8usm49XdTdkkmyzoW6KJl JK6i36NDp36K3BKAzn6mHNjT42UwCpGREIy6g/qSAk6CvfA2+LjqV+Xr6xkBtMe62qgpgLaPQC I4E= To: gcc-patches , fortran , Jakub Jelinek From: Tobias Burnus Subject: [Patch][Fortran/OpenMP] Don't create "alloc:" for 'target exit data' Message-ID: <7a5f39e8-a33b-048a-f9c1-1355b941771e@codesourcery.com> Date: Fri, 18 Oct 2019 11:27:39 +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 Currently, one has for   !$omp target exit data map(delete:x) in the original dump:   #pragma omp target exit data map(delete:*x) map(alloc:x [pointer assign, bias: 0]) The "alloc:" not only does not make sense but also gives run-time messages like: libgomp: GOMP_target_enter_exit_data unhandled kind 0x04 [Depending on the data type, in gfc_trans_omp_clauses's OMP_LIST_MAP, add map clauses of type GOMP_MAP_POINTER and/or GOMP_MAP_TO_PSET.] That's for release:/delete:. However, for 'target exit data' (GOMP_target_enter_exit_data) the same issue occurs for "from:"/"always, from:".  But "from:" implies "alloc:". – While "alloc:" does not make sense for "target exit data" or "update", for "target" or "target data" it surely matters. Hence, I only exclude "from:" for exit data and update. See attached patch. I have additionally Fortran-fied libgomp.c/target-20.c to have at least one 'enter/exit target data' test case for Fortran. Build + regtested on x86_64-gnu-linux w/o offloading. And I have tested the new test case with nvptx. Tobias gcc/fortran/ * trans-openmp.c (gfc_trans_omp_clauses): Do not create map(alloc:) for map(delete:/release:) and for (from:/always,from:) only if new arg require_from_alloc is true, which is the default. (gfc_trans_omp_target_exit_data, gfc_trans_omp_target_update): Call it with require_from_alloc = false. libgomp/ * testsuite/libgomp.fortran/target9.f90: New. diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c index dad11a24430..f890629c73d 100644 --- a/gcc/fortran/trans-openmp.c +++ b/gcc/fortran/trans-openmp.c @@ -1852,7 +1852,8 @@ static vec *doacross_steps; static tree gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, - locus where, bool declare_simd = false) + locus where, bool declare_simd = false, + bool require_from_alloc = true) { tree omp_clauses = NULL_TREE, chunk_size, c; int list, ifc; @@ -2163,6 +2164,16 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, if (!n->sym->attr.referenced) continue; + /* map(alloc:) etc. is not needed for delete/release + For 'from:', it is needed when setting up the environment + but not for updating or copying out of the data. */ + bool no_extra_pointer = n->u.map_op == OMP_MAP_DELETE + || n->u.map_op == OMP_MAP_RELEASE + || (!require_from_alloc + && (n->u.map_op == OMP_MAP_FROM + || n->u.map_op + == OMP_MAP_ALWAYS_FROM)); + tree node = build_omp_clause (input_location, OMP_CLAUSE_MAP); tree node2 = NULL_TREE; tree node3 = NULL_TREE; @@ -2172,7 +2183,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, TREE_ADDRESSABLE (decl) = 1; if (n->expr == NULL || n->expr->ref->u.ar.type == AR_FULL) { - if (POINTER_TYPE_P (TREE_TYPE (decl)) + if (!no_extra_pointer + && POINTER_TYPE_P (TREE_TYPE (decl)) && (gfc_omp_privatize_by_reference (decl) || GFC_DECL_GET_SCALAR_POINTER (decl) || GFC_DECL_GET_SCALAR_ALLOCATABLE (decl) @@ -2208,17 +2220,20 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, ptr); ptr = build_fold_indirect_ref (ptr); OMP_CLAUSE_DECL (node) = ptr; - node2 = build_omp_clause (input_location, - OMP_CLAUSE_MAP); - OMP_CLAUSE_SET_MAP_KIND (node2, GOMP_MAP_TO_PSET); - OMP_CLAUSE_DECL (node2) = decl; - OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type); - node3 = build_omp_clause (input_location, - OMP_CLAUSE_MAP); - OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_POINTER); - OMP_CLAUSE_DECL (node3) - = gfc_conv_descriptor_data_get (decl); - OMP_CLAUSE_SIZE (node3) = size_int (0); + if (!no_extra_pointer) + { + node2 = build_omp_clause (input_location, + OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (node2, GOMP_MAP_TO_PSET); + OMP_CLAUSE_DECL (node2) = decl; + OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type); + node3 = build_omp_clause (input_location, + OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_POINTER); + OMP_CLAUSE_DECL (node3) + = gfc_conv_descriptor_data_get (decl); + OMP_CLAUSE_SIZE (node3) = size_int (0); + } /* We have to check for n->sym->attr.dimension because of scalar coarrays. */ @@ -2302,6 +2317,9 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, ptr); OMP_CLAUSE_DECL (node) = build_fold_indirect_ref (ptr); + if (no_extra_pointer) + goto skip_extra_map_pointer; + if (POINTER_TYPE_P (TREE_TYPE (decl)) && GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (TREE_TYPE (decl)))) { @@ -2346,6 +2364,9 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, OMP_CLAUSE_SIZE (node3) = fold_build2 (MINUS_EXPR, sizetype, ptr, ptr2); } + + skip_extra_map_pointer: + switch (n->u.map_op) { case OMP_MAP_ALLOC: @@ -4979,7 +5000,7 @@ gfc_trans_omp_target_exit_data (gfc_code *code) gfc_start_block (&block); omp_clauses = gfc_trans_omp_clauses (&block, code->ext.omp_clauses, - code->loc); + code->loc, false, false); stmt = build1_loc (input_location, OMP_TARGET_EXIT_DATA, void_type_node, omp_clauses); gfc_add_expr_to_block (&block, stmt); @@ -4994,7 +5015,7 @@ gfc_trans_omp_target_update (gfc_code *code) gfc_start_block (&block); omp_clauses = gfc_trans_omp_clauses (&block, code->ext.omp_clauses, - code->loc); + code->loc, false, false); stmt = build1_loc (input_location, OMP_TARGET_UPDATE, void_type_node, omp_clauses); gfc_add_expr_to_block (&block, stmt); diff --git a/libgomp/testsuite/libgomp.fortran/target9.f90 b/libgomp/testsuite/libgomp.fortran/target9.f90 new file mode 100644 index 00000000000..91d60a33307 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/target9.f90 @@ -0,0 +1,123 @@ +! { dg-require-effective-target offload_device_nonshared_as } */ + +module target_test + implicit none (type, external) + integer, parameter :: N = 40 + integer :: sum + integer :: var1 = 1 + integer :: var2 = 2 + + !$omp declare target to(D) + integer :: D(N) = 0 +contains + subroutine enter_data (X) + integer :: X(:) + !$omp target enter data map(to: var1, var2, X) map(alloc: sum) + end subroutine enter_data + + subroutine exit_data_0 (D) + integer :: D(N) + !$omp target exit data map(delete: D) + end subroutine exit_data_0 + + subroutine exit_data_1 () + !$omp target exit data map(from: var1) + end subroutine exit_data_1 + + subroutine exit_data_2 (X) + integer :: X(N) + !$omp target exit data map(from: var2) map(release: X, sum) + end subroutine exit_data_2 + + subroutine exit_data_3 (p, idx) + integer :: p(:) + integer, value :: idx + !$omp target exit data map(from: p(idx)) + end subroutine exit_data_3 + + subroutine test_nested () + integer :: X, Y, Z + X = 0 + Y = 0 + Z = 0 + + !$omp target data map(from: X, Y, Z) + !$omp target data map(from: X, Y, Z) + !$omp target map(from: X, Y, Z) + X = 1337 + Y = 1337 + Z = 1337 + !$omp end target + if (X /= 0) stop 11 + if (Y /= 0) stop 12 + if (Z /= 0) stop 13 + + !$omp target exit data map(from: X) map(release: Y) + if (X /= 0) stop 14 + if (Y /= 0) stop 15 + + !$omp target exit data map(release: Y) map(delete: Z) + if (Y /= 0) stop 16 + if (Z /= 0) stop 17 + !$omp end target data + if (X /= 1337) stop 18 + if (Y /= 0) stop 19 + if (Z /= 0) stop 20 + + !$omp target map(from: X) + X = 2448 + !$omp end target + if (X /= 2448) stop 21 + if (Y /= 0) stop 22 + if (Z /= 0) stop 23 + + X = 4896 + !$omp end target data + if (X /= 4896) stop 24 + if (Y /= 0) stop 25 + if (Z /= 0) stop 26 + end subroutine test_nested +end module target_test + +program main + use target_test + implicit none (type, external) + + integer, allocatable :: X(:) + integer, pointer, contiguous :: Y(:) + + + allocate(X(N), Y(N)) + X(10) = 10 + Y(20) = 20 + call enter_data (X) + + call exit_data_0 (D) ! This should have no effect on D. + + !$omp target map(alloc: var1, var2, X) map(to: Y) map(always, from: sum) + var1 = var1 + X(10) + var2 = var2 + Y(20) + sum = var1 + var2 + D(sum) = D(sum) + 1 + !$omp end target + + if (var1 /= 1) stop 1 + if (var2 /= 2) stop 2 + if (sum /= 33) stop 3 + + call exit_data_1 () + if (var1 /= 11) stop 4 + if (var2 /= 2) stop 5 + + ! Increase refcount of already mapped X(1:N). + !$omp target enter data map(alloc: X(16:17)) + + call exit_data_2 (X) + if (var2 /= 22) stop 6 + + call exit_data_3 (X, 5) ! Unmap X(1:N). + + deallocate (X, Y) + + call test_nested () +end program main