From patchwork Fri Mar 4 14:47:31 2022 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Marcel Vollweiler X-Patchwork-Id: 1601135 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+incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) 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 bilbo.ozlabs.org (Postfix) with ESMTPS id 4K99hT2vGPz9sG7 for ; Sat, 5 Mar 2022 01:48:24 +1100 (AEDT) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id CFA883857827 for ; Fri, 4 Mar 2022 14:48:21 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa3.mentor.iphmx.com (esa3.mentor.iphmx.com [68.232.137.180]) by sourceware.org (Postfix) with ESMTPS id D5BCF3858D39; Fri, 4 Mar 2022 14:47:55 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org D5BCF3858D39 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.90,155,1643702400"; d="diff'?scan'208";a="72568264" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa3.mentor.iphmx.com with ESMTP; 04 Mar 2022 06:47:54 -0800 IronPort-SDR: D+SPo4l7BdlV3qbTGQ1YSWLBr55JiWksySM3ul2Iz11guWqrbSGDlUr2D1/WI84jrJGhyBYTaU IQ67Efw9V3IOhnHY05HuIPYTSXXsOyaGtGcCtgQweHYoyUUdwji2SLc/rvtqZ9Dt+r+UmU2gmt Ucbf0dF0o7B4H/dfi7bwX8UFbkzUrpZItAcrjxYrRLkja0peFbR7i154BplnadzPEIVGscJ4KA kc8QV3Ds8xdNu3Da66pi6s1HqmFFfeRtHUrtWe+ef3NYZl/E8by7OZsi92znhOXiJ/V4yrBigj crg= Message-ID: <43ad60ea-875a-5f6f-106b-206ecd4807ba@codesourcery.com> Date: Fri, 4 Mar 2022 15:47:31 +0100 MIME-Version: 1.0 User-Agent: Mozilla/5.0 (Windows NT 10.0; Win64; x64; rv:91.0) Gecko/20100101 Thunderbird/91.6.1 From: Marcel Vollweiler Subject: [PATCH] OpenMP, libgomp: Add new runtime routine omp_get_mapped_ptr. To: X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: SVR-IES-MBX-08.mgc.mentorg.com (139.181.222.8) To SVR-IES-MBX-03.mgc.mentorg.com (139.181.222.3) X-Spam-Status: No, score=-12.8 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, SPF_HELO_PASS, SPF_PASS, TXREP, T_SCC_BODY_TEXT_LINE autolearn=ham autolearn_force=no version=3.4.4 X-Spam-Checker-Version: SpamAssassin 3.4.4 (2020-01-24) 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: , Cc: fortran@gcc.gnu.org Errors-To: gcc-patches-bounces+incoming=patchwork.ozlabs.org@gcc.gnu.org Sender: "Gcc-patches" Hi, This patch adds the OpenMP runtime routine "omp_get_mapped_ptr" which was introduced in OpenMP 5.1 (specification section 3.8.11): "The omp_get_mapped_ptr routine returns the device pointer that is associated with a host pointer for a given device." "The device_num argument must be greater than or equal to zero and less than or equal to the result of omp_get_num_devices()." "A call to this routine for a pointer that is not NULL (or C_NULL_PTR, for Fortran) and does not have an associated pointer on the given device results in a NULL pointer." "The routine returns NULL (or C_NULL_PTR, for Fortran) if unsuccessful. Otherwise it returns the device pointer, which is ptr if device_num is the value returned by omp_get_initial_device()." Implementation and tests were added for C/C++ and Fortran. There is a small inconvenience considering zero-length arrays as list items of the "target map" construct: it seems that zero-length arrays are not associated correctly there, such that omp_get_mapped_ptr returns NULL instead of the associated device pointer - in contrast to the situation where a device pointer is associated with the host pointer via omp_target_associate_ptr. However, the result for omp_get_mapped_ptr is consistent with omp_target_is_present (which returns 0, i.e. "not present") in this situation. The patch was tested on x86_64-linux with nvptx and amdgcn offloading. All with no regressions. Marcel ----------------- 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, libgomp: Add new runtime routine omp_get_mapped_ptr. libgomp/ChangeLog: * libgomp.map: Added omp_get_mapped_ptr. * libgomp.texi: Tagged omp_get_mapped_ptr as supported. * omp.h.in: Added omp_get_mapped_ptr. * omp_lib.f90.in: Added interface for omp_get_mapped_ptr. * omp_lib.h.in: Likewise. * target.c (omp_get_mapped_ptr): Added implementation of omp_get_mapped_ptr. * testsuite/libgomp.c-c++-common/get-mapped-ptr-1.c: New test. * testsuite/libgomp.c-c++-common/get-mapped-ptr-2.c: New test. * testsuite/libgomp.c-c++-common/get-mapped-ptr-3.c: New test. * testsuite/libgomp.c-c++-common/get-mapped-ptr-4.c: New test. * testsuite/libgomp.fortran/get-mapped-ptr-1.f90: New test. * testsuite/libgomp.fortran/get-mapped-ptr-2.f90: New test. * testsuite/libgomp.fortran/get-mapped-ptr-3.f90: New test. * testsuite/libgomp.fortran/get-mapped-ptr-4.f90: New test. diff --git a/libgomp/libgomp.map b/libgomp/libgomp.map index 2ac5809..00a4858 100644 --- a/libgomp/libgomp.map +++ b/libgomp/libgomp.map @@ -224,6 +224,7 @@ OMP_5.1 { omp_set_teams_thread_limit_8_; omp_get_teams_thread_limit; omp_get_teams_thread_limit_; + omp_get_mapped_ptr; } OMP_5.0.2; GOMP_1.0 { diff --git a/libgomp/libgomp.texi b/libgomp/libgomp.texi index 161a423..c163b56 100644 --- a/libgomp/libgomp.texi +++ b/libgomp/libgomp.texi @@ -314,7 +314,7 @@ The OpenMP 4.5 specification is fully supported. @item @code{omp_target_is_accessible} runtime routine @tab N @tab @item @code{omp_target_memcpy_async} and @code{omp_target_memcpy_rect_async} runtime routines @tab N @tab -@item @code{omp_get_mapped_ptr} runtime routine @tab N @tab +@item @code{omp_get_mapped_ptr} runtime routine @tab Y @tab @item @code{omp_calloc}, @code{omp_realloc}, @code{omp_aligned_alloc} and @code{omp_aligned_calloc} runtime routines @tab Y @tab @item @code{omp_alloctrait_key_t} enum: @code{omp_atv_serialized} added, diff --git a/libgomp/omp.h.in b/libgomp/omp.h.in index 89c5d65..18d0152 100644 --- a/libgomp/omp.h.in +++ b/libgomp/omp.h.in @@ -282,6 +282,7 @@ extern int omp_target_memcpy_rect (void *, const void *, __SIZE_TYPE__, int, extern int omp_target_associate_ptr (const void *, const void *, __SIZE_TYPE__, __SIZE_TYPE__, int) __GOMP_NOTHROW; extern int omp_target_disassociate_ptr (const void *, int) __GOMP_NOTHROW; +extern void *omp_get_mapped_ptr (const void *, int) __GOMP_NOTHROW; extern void omp_set_affinity_format (const char *) __GOMP_NOTHROW; extern __SIZE_TYPE__ omp_get_affinity_format (char *, __SIZE_TYPE__) diff --git a/libgomp/omp_lib.f90.in b/libgomp/omp_lib.f90.in index daf40dc..506f15c 100644 --- a/libgomp/omp_lib.f90.in +++ b/libgomp/omp_lib.f90.in @@ -835,6 +835,15 @@ end function omp_target_disassociate_ptr end interface + interface + function omp_get_mapped_ptr (ptr, device_num) bind(c) + use, intrinsic :: iso_c_binding, only : c_ptr, c_int + type(c_ptr) :: omp_get_mapped_ptr + type(c_ptr), value :: ptr + integer(c_int), value :: device_num + end function omp_get_mapped_ptr + end interface + #if _OPENMP >= 201811 !GCC$ ATTRIBUTES DEPRECATED :: omp_get_nested, omp_set_nested #endif diff --git a/libgomp/omp_lib.h.in b/libgomp/omp_lib.h.in index ff857a4..0f48510 100644 --- a/libgomp/omp_lib.h.in +++ b/libgomp/omp_lib.h.in @@ -416,3 +416,12 @@ integer(c_int), value :: device_num end function omp_target_disassociate_ptr end interface + + interface + function omp_get_mapped_ptr (ptr, device_num) bind(c) + use, intrinsic :: iso_c_binding, only : c_ptr, c_int + type(c_ptr) :: omp_get_mapped_ptr + type(c_ptr), value :: ptr + integer(c_int), value :: device_num + end function omp_get_mapped_ptr + end interface diff --git a/libgomp/target.c b/libgomp/target.c index 9017458..735d70b 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -3665,6 +3665,49 @@ omp_target_disassociate_ptr (const void *ptr, int device_num) return ret; } +void * +omp_get_mapped_ptr (const void *ptr, int device_num) +{ + if (device_num < 0 || device_num > omp_get_num_devices ()) + return NULL; + + if (device_num == omp_get_initial_device ()) + return (void*)ptr; + + struct gomp_device_descr *devicep = resolve_device (device_num); + if (devicep == NULL) + return NULL; + + if (!(devicep->capabilities & GOMP_OFFLOAD_CAP_OPENMP_400) + || devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) + return (void*)ptr; + + gomp_mutex_lock (&devicep->lock); + + struct splay_tree_s *mem_map = &devicep->mem_map; + struct splay_tree_key_s cur_node; + void *ret = NULL; + uintptr_t offset = 0; + + cur_node.host_start = (uintptr_t) ptr; + cur_node.host_end = cur_node.host_start; + splay_tree_key n = gomp_map_0len_lookup (mem_map, &cur_node); + + if (n && n->host_start == cur_node.host_start) + { + ret = (void*) n->tgt->tgt_start + n->tgt_offset; + } + else if (n) + { + offset = cur_node.host_start - n->host_start; + ret = (void*) n->tgt->tgt_start + n->tgt_offset + offset; + } + + gomp_mutex_unlock (&devicep->lock); + + return ret; +} + int omp_pause_resource (omp_pause_resource_t kind, int device_num) { diff --git a/libgomp/testsuite/libgomp.c-c++-common/get-mapped-ptr-1.c b/libgomp/testsuite/libgomp.c-c++-common/get-mapped-ptr-1.c new file mode 100644 index 0000000..28b9332 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/get-mapped-ptr-1.c @@ -0,0 +1,41 @@ +#include +#include + +int +main () +{ + int d = omp_get_default_device (); + int id = omp_get_initial_device (); + void *p , *q; + + if (d < 0 || d >= omp_get_num_devices ()) + d = id; + + p = omp_target_alloc (sizeof (int), d); + if (p == NULL) + return 0; + + if (omp_target_associate_ptr (q, p, sizeof (int), 0, d) != 0) + return 0; + + if (omp_get_mapped_ptr (q, -1) != NULL) + __builtin_abort (); + + if (omp_get_mapped_ptr (q, omp_get_num_devices () + 1) != NULL) + __builtin_abort (); + + if (omp_get_mapped_ptr (q, id) != q) + __builtin_abort (); + + if (omp_get_mapped_ptr (q, d) != p) + __builtin_abort (); + + if (omp_target_disassociate_ptr (q, d) != 0) + __builtin_abort (); + + if (omp_get_mapped_ptr (q, d) != NULL) + __builtin_abort (); + + omp_target_free (p, d); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/get-mapped-ptr-2.c b/libgomp/testsuite/libgomp.c-c++-common/get-mapped-ptr-2.c new file mode 100644 index 0000000..bbe4714 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/get-mapped-ptr-2.c @@ -0,0 +1,106 @@ +#include +#include +#include + +int +main () +{ + int d = omp_get_default_device (); + int id = omp_get_initial_device (); + int a = 42; + int b[] = { 24, 42 }; + int c[] = { 47, 11 }; + int e[128]; + int *q = &a; + void *p1 = NULL, *p2 = NULL, *p3 = NULL; + void *devptrs[128]; + + if (d < 0 || d >= omp_get_num_devices ()) + d = id; + + for (int i = 0; i < 128; i++) + e[i] = i; + + #pragma omp target data map(alloc: a, b, c[1], e[32:64]) device(d) + { + #pragma omp target map(from: p1, p2, p3, devptrs) map(alloc: a, b, c[1], e[32:64]) device(d) + { + p1 = &a; + p2 = &b; + p3 = &c[1]; + for (int i = 32; i < 96; i++) + devptrs[i] = &e[i]; + } + + if (omp_get_mapped_ptr (&a, d) != (d == id ? &a : p1) + || omp_get_mapped_ptr (q, d) != (d == id ? q : p1) + || omp_get_mapped_ptr (b, d) != (d == id ? b : p2) + || omp_get_mapped_ptr (&b[0], d) != (d == id ? &b[0] : p2) + || omp_get_mapped_ptr (&c[1], d) != (d == id ? &c[1] : p3) + || omp_get_mapped_ptr (&c[0], d) != (d == id ? &c[0] : NULL)) + __builtin_abort (); + + for (int i = 0; i < 32; i++) + if (omp_get_mapped_ptr (&e[i], d) != (d == id ? &e[i] : NULL)) + __builtin_abort (); + for (int i = 32; i < 96; i++) + if (omp_get_mapped_ptr (&e[i], d) != (d == id ? &e[i] : devptrs[i])) + __builtin_abort (); + for (int i = 96; i < 128; i++) + if (omp_get_mapped_ptr (&e[i], d) != (d == id ? &e[i] : NULL)) + __builtin_abort (); + } + + if (omp_get_mapped_ptr (&a, d) != (d == id ? &a : NULL) + || omp_get_mapped_ptr (q, d) != (d == id ? q : NULL) + || omp_get_mapped_ptr (b, d) != (d == id ? b : NULL) + || omp_get_mapped_ptr (&b[0], d) != (d == id ? &b[0] : NULL) + || omp_get_mapped_ptr (&c[1], d) != (d == id ? &c[1] : NULL) + || omp_get_mapped_ptr (&c[0], d) != (d == id ? &c[0] : NULL)) + __builtin_abort (); + for (int i = 0; i < 128; i++) + if (omp_get_mapped_ptr (&e[i], d) != (d == id ? &e[i] : NULL)) + __builtin_abort (); + + #pragma omp target enter data map (alloc: a, b, c[1], e[32:64]) device (d) + #pragma omp target map(from: p1, p2, p3, devptrs) map(alloc: a, b, c[1], e[32:64]) device(d) + { + p1 = &a; + p2 = &b; + p3 = &c[1]; + for (int i = 32; i < 96; i++) + devptrs[i] = &e[i]; + } + + if (omp_get_mapped_ptr (&a, d) != (d == id ? &a : p1) + || omp_get_mapped_ptr (q, d) != (d == id ? q : p1) + || omp_get_mapped_ptr (b, d) != (d == id ? b : p2) + || omp_get_mapped_ptr (&b[0], d) != (d == id ? &b[0] : p2) + || omp_get_mapped_ptr (&c[1], d) != (d == id ? &c[1] : p3) + || omp_get_mapped_ptr (&c[0], d) != (d == id ? &c[0] : NULL)) + __builtin_abort (); + for (int i = 0; i < 32; i++) + if (omp_get_mapped_ptr (&e[i], d) != (d == id ? &e[i] : NULL)) + __builtin_abort (); + for (int i = 32; i < 96; i++) + if (omp_get_mapped_ptr (&e[i], d) != (d == id ? &e[i] : devptrs[i])) + __builtin_abort (); + for (int i = 96; i < 128; i++) + if (omp_get_mapped_ptr (&e[i], d) != (d == id ? &e[i] : NULL)) + __builtin_abort (); + + #pragma omp target exit data map (delete: a, b, c[1], e[32:64]) device (d) + + if (omp_get_mapped_ptr (&a, d) != (d == id ? &a : NULL) + || omp_get_mapped_ptr (q, d) != (d == id ? q : NULL) + || omp_get_mapped_ptr (b, d) != (d == id ? b : NULL) + || omp_get_mapped_ptr (&b[0], d) != (d == id ? &b[0] : NULL) + || omp_get_mapped_ptr (&c[1], d) != (d == id ? &c[1] : NULL) + || omp_get_mapped_ptr (&c[0], d) != (d == id ? &c[0] : NULL)) + __builtin_abort (); + for (int i = 0; i < 128; i++) + if (omp_get_mapped_ptr (&e[i], d) != (d == id ? &e[i] : NULL)) + __builtin_abort (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/get-mapped-ptr-3.c b/libgomp/testsuite/libgomp.c-c++-common/get-mapped-ptr-3.c new file mode 100644 index 0000000..67429d7 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/get-mapped-ptr-3.c @@ -0,0 +1,51 @@ +#include +#include + +int +main () +{ + int d = omp_get_default_device (); + int id = omp_get_initial_device (); + int a[0]; + int b[] = { 24, 42 }; + void *p1 = NULL, *p2 = NULL; + + if (d < 0 || d >= omp_get_num_devices ()) + d = id; + + void *p = omp_target_alloc (sizeof (int), d); + if (p == NULL) + return 0; + + if (omp_target_associate_ptr (a, p, sizeof (int), 0, d) != 0) + return 0; + + if (omp_get_mapped_ptr (a, d) != (d == id ? a : p)) + __builtin_abort (); + + if (omp_target_disassociate_ptr (a, d) != 0) + __builtin_abort (); + + if (omp_get_mapped_ptr (a, d) != (d == id ? a : NULL)) + __builtin_abort (); + + #pragma omp target data map(alloc: a, b[1:0]) device(d) + { + #pragma omp target map(from: p1, p2) map(alloc: a, b[1:0]) device(d) + { + p1 = &a; + p2 = &b[1]; + } + + /* This is probably expected to be p1/p2 instead of NULL. Zero-length arrays + as list items of the map clause are currently not inserted into the mem + map ?! However by returning NULL, omp_get_mapped_ptr is consistent with + omp_target_is_present. */ + if (omp_get_mapped_ptr (a, d) != NULL + || omp_get_mapped_ptr (&b[1], d) != NULL) + __builtin_abort (); + } + + omp_target_free (p, d); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/get-mapped-ptr-4.c b/libgomp/testsuite/libgomp.c-c++-common/get-mapped-ptr-4.c new file mode 100644 index 0000000..c8eac06 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/get-mapped-ptr-4.c @@ -0,0 +1,49 @@ +#include +#include + +int +main () +{ + int d = omp_get_default_device (); + int id = omp_get_initial_device (); + struct s_t { int m1; char m2; } s; + void *p1 = NULL, *p2 = NULL; + + if (d < 0 || d >= omp_get_num_devices ()) + d = id; + + #pragma omp target data map(alloc: s, s.m2) device(d) + { + #pragma omp target map(from: p1, p2) map(alloc: s, s.m2) device(d) + { + p1 = &s; + p2 = &s.m2; + } + if (omp_get_mapped_ptr (&s, d) != (d == id ? &s : p1) + || omp_get_mapped_ptr (&s.m2, d) != (d == id ? &s.m2 : p2)) + __builtin_abort (); + } + + if (omp_get_mapped_ptr (&s, d) != (d == id ? &s : NULL) + || omp_get_mapped_ptr (&s.m2, d) != (d == id ? &s.m2 : NULL)) + __builtin_abort (); + + #pragma omp target enter data map(alloc: s, s.m2) device (d) + #pragma omp target map(from: p1, p2) map(alloc: s, s.m2) device(d) + { + p1 = &s; + p2 = &s.m2; + } + + if (omp_get_mapped_ptr (&s, d) != (d == id ? &s : p1) + || omp_get_mapped_ptr (&s.m2, d) != (d == id ? &s.m2 : p2)) + __builtin_abort (); + + #pragma omp target exit data map (delete: s, s.m2) device (d) + + if (omp_get_mapped_ptr (&s, d) != (d == id ? &s : NULL) + || omp_get_mapped_ptr (&s.m2, d) != (d == id ? &s.m2 : NULL)) + __builtin_abort (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.fortran/get-mapped-ptr-1.f90 b/libgomp/testsuite/libgomp.fortran/get-mapped-ptr-1.f90 new file mode 100644 index 0000000..de05179 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/get-mapped-ptr-1.f90 @@ -0,0 +1,43 @@ +program main + use omp_lib + use iso_c_binding + implicit none (external, type) + integer :: d, id + type(c_ptr) :: p + integer, target :: q + + d = omp_get_default_device () + id = omp_get_initial_device () + + if (d < 0 .or. d >= omp_get_num_devices ()) & + d = id + + p = omp_target_alloc (c_sizeof (q), d) + if (.not. c_associated (p)) & + stop 0 ! okay + + if (omp_target_associate_ptr (c_loc (q), p, c_sizeof (q), & + 0_c_size_t, d) == 0) then + + if(c_associated (omp_get_mapped_ptr (c_loc (q), -1))) & + stop 1 + + if(c_associated (omp_get_mapped_ptr (c_loc (q), & + omp_get_num_devices () + 1))) & + stop 2 + + if(.not. c_associated (omp_get_mapped_ptr (c_loc (q), id), c_loc (q))) & + stop 3 + + if(.not. c_associated (omp_get_mapped_ptr (c_loc (q), d), p)) & + stop 4 + + if (omp_target_disassociate_ptr (c_loc (q), d) /= 0) & + stop 5 + + if(c_associated (omp_get_mapped_ptr (c_loc (q), d))) & + stop 6 + end if + + call omp_target_free (p, d) +end program main diff --git a/libgomp/testsuite/libgomp.fortran/get-mapped-ptr-2.f90 b/libgomp/testsuite/libgomp.fortran/get-mapped-ptr-2.f90 new file mode 100644 index 0000000..66a0b88 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/get-mapped-ptr-2.f90 @@ -0,0 +1,175 @@ +program main + use omp_lib + use iso_c_binding + implicit none (external, type) + integer :: d, id, i, j + integer, target :: a, b(1:2), c(1:2), e(0:127) + type(c_ptr) :: p1, p2, p3, q, devptrs(0:63) + + a = 42; + q = c_loc (a); + e = [(i, i = 0, 127)] + + d = omp_get_default_device () + id = omp_get_initial_device () + + if (d < 0 .or. d >= omp_get_num_devices ()) & + d = id + + if (d /= id) then + !$omp target data map(alloc: a, b, c(2), e(32:95)) device(d) + !$omp target map(from: p1, p2, p3, devptrs) map(alloc: a, b, c(2), e(32:95)) device(d) + p1 = c_loc (a); + p2 = c_loc (b); + p3 = c_loc (c(2)) + devptrs = [(c_loc (e(i)), i = 32, 95)] + !$omp end target + + if (.not. c_associated (omp_get_mapped_ptr (c_loc (a), d), p1) & + .or. .not. c_associated (omp_get_mapped_ptr (q, d), p1) & + .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b), d), p2) & + .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b(1)), d), p2) & + .or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(2)), d), p3) & + .or. c_associated (omp_get_mapped_ptr (c_loc (c(1)), d))) & + stop 0 + + do j = 0, 31 + if (c_associated (omp_get_mapped_ptr (c_loc (e(j)), d))) & + stop 1 + end do + do j = 32, 95 + if (.not. c_associated (omp_get_mapped_ptr (c_loc (e(j)), d), devptrs(j-32))) & + stop 2 + end do + do j = 96, 128 + if (c_associated (omp_get_mapped_ptr (c_loc (e(j)), d))) & + stop 3 + end do + !$omp end target data + + if (c_associated (omp_get_mapped_ptr (c_loc (a), d)) & + .or. c_associated (omp_get_mapped_ptr (q, d)) & + .or. c_associated (omp_get_mapped_ptr (c_loc (b), d)) & + .or. c_associated (omp_get_mapped_ptr (c_loc (b(1)), d)) & + .or. c_associated (omp_get_mapped_ptr (c_loc (c(2)), d)) & + .or. c_associated (omp_get_mapped_ptr (c_loc (c(1)), d))) & + stop 4 + do j = 0, 127 + if (c_associated (omp_get_mapped_ptr (c_loc (e(j)), d))) & + stop 5 + end do + + !$omp target enter data map (alloc: a, b, c(2), e(32:95)) device (d) + !$omp target map(from: p1, p2, p3, devptrs) map(alloc: a, b, c(2), e(32:95)) device(d) + p1 = c_loc (a); + p2 = c_loc (b); + p3 = c_loc (c(2)) + devptrs = [(c_loc (e(i)), i = 32, 95)] + !$omp end target + + if (.not. c_associated (omp_get_mapped_ptr (c_loc (a), d), p1) & + .or. .not. c_associated (omp_get_mapped_ptr (q, d), p1) & + .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b), d), p2) & + .or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(2)), d), p3) & + .or. c_associated (omp_get_mapped_ptr (c_loc (c(1)), d))) & + stop 6 + + do j = 0, 31 + if (c_associated (omp_get_mapped_ptr (c_loc (e(j)), d))) & + stop 7 + end do + do j = 32, 95 + if (.not. c_associated (omp_get_mapped_ptr (c_loc (e(j)), d), devptrs(j-32))) & + stop 8 + end do + do j = 96, 128 + if (c_associated (omp_get_mapped_ptr (c_loc (e(j)), d))) & + stop 9 + end do + !$omp target exit data map (delete: a, b, c(2), e(32:95)) device (d) + + if (c_associated (omp_get_mapped_ptr (c_loc (a), d)) & + .or. c_associated (omp_get_mapped_ptr (q, d)) & + .or. c_associated (omp_get_mapped_ptr (c_loc (b), d)) & + .or. c_associated (omp_get_mapped_ptr (c_loc (b(1)), d)) & + .or. c_associated (omp_get_mapped_ptr (c_loc (c(1)), d)) & + .or. c_associated (omp_get_mapped_ptr (c_loc (c(2)), d))) & + stop 10 + do j = 0, 127 + if (c_associated (omp_get_mapped_ptr (c_loc (e(j)), d))) & + stop 11 + end do + + else ! d == id + + !$omp target data map(alloc: a, b, c(2), e(32:95)) device(d) + !$omp target map(from: p1, p2, p3, devptrs) map(alloc: a, b, c(2), e(32:95)) device(d) + p1 = c_loc (a); + p2 = c_loc (b); + p3 = c_loc (c(2)) + devptrs = [(c_loc (e(i)), i = 32, 95)] + !$omp end target + + if (.not. c_associated (omp_get_mapped_ptr (c_loc (a), d), c_loc (a)) & + .or. .not. c_associated (omp_get_mapped_ptr (q, d), q) & + .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b), d), c_loc (b)) & + .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b(1)), d), c_loc (b(1))) & + .or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(2)), d), c_loc (c(2))) & + .or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(1)), d), c_loc (c(1)))) & + stop 12 + + do j = 0, 127 + if (.not. c_associated (omp_get_mapped_ptr (c_loc (e(j)), d), c_loc (e(j)))) & + stop 13 + end do + !$omp end target data + + if (.not. c_associated (omp_get_mapped_ptr (c_loc (a), d), c_loc (a)) & + .or. .not. c_associated (omp_get_mapped_ptr (q, d), q) & + .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b), d), c_loc (b)) & + .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b(1)), d), c_loc (b(1))) & + .or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(2)), d), c_loc (c(2))) & + .or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(1)), d), c_loc (c(1)))) & + stop 14 + do j = 0, 127 + if (.not. c_associated (omp_get_mapped_ptr (c_loc (e(j)), d))) & + stop 15 + end do + + !$omp target enter data map (alloc: a, b, c(2), e(32:95)) device (d) + !$omp target map(from: p1, p2, p3, devptrs) map(alloc: a, b, c(2), e(32:95)) device(d) + p1 = c_loc (a); + p2 = c_loc (b); + p3 = c_loc (c(2)) + devptrs = [(c_loc (e(i)), i = 32, 95)] + !$omp end target + + if (.not. c_associated (omp_get_mapped_ptr (c_loc (a), d), c_loc (a)) & + .or. .not. c_associated (omp_get_mapped_ptr (q, d), q) & + .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b), d), c_loc (b)) & + .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b(1)), d), c_loc (b(1))) & + .or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(2)), d), c_loc (c(2))) & + .or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(1)), d), c_loc (c(1)))) & + stop 16 + + do j = 0, 127 + if (.not. c_associated (omp_get_mapped_ptr (c_loc (e(j)), d), c_loc (e(j)))) & + stop 17 + end do + !$omp target exit data map (delete: a, b, c(2), e(32:95)) device (d) + + if (.not. c_associated (omp_get_mapped_ptr (c_loc (a), d), c_loc (a)) & + .or. .not. c_associated (omp_get_mapped_ptr (q, d), q) & + .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b), d), c_loc (b)) & + .or. .not. c_associated (omp_get_mapped_ptr (c_loc (b(1)), d), c_loc (b(1))) & + .or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(2)), d), c_loc (c(2))) & + .or. .not. c_associated (omp_get_mapped_ptr (c_loc (c(1)), d), c_loc (c(1)))) & + stop 18 + + do j = 0, 127 + if (.not. c_associated (omp_get_mapped_ptr (c_loc (e(j)), d), c_loc (e(j)))) & + stop 19 + end do + end if + +end program main diff --git a/libgomp/testsuite/libgomp.fortran/get-mapped-ptr-3.f90 b/libgomp/testsuite/libgomp.fortran/get-mapped-ptr-3.f90 new file mode 100644 index 0000000..8e7ccac --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/get-mapped-ptr-3.f90 @@ -0,0 +1,48 @@ +program main + use omp_lib + use iso_c_binding + implicit none (external, type) + integer :: d, id + type(c_ptr) :: p, p1, p2 + integer, target :: a(1:0), b(1:2) + + d = omp_get_default_device () + id = omp_get_initial_device () + + if (d < 0 .or. d >= omp_get_num_devices ()) & + d = id + + p = omp_target_alloc (c_sizeof (c_int), d) + if (.not. c_associated (p)) & + stop 0 ! okay + + if (omp_target_associate_ptr (c_loc (a), p, c_sizeof (c_int), & + 0_c_size_t, d) == 0) then + + if(.not. c_associated (omp_get_mapped_ptr (c_loc (a), d), p)) & + stop 1 + + if (omp_target_disassociate_ptr (c_loc (a), d) /= 0) & + stop 2 + + if(c_associated (omp_get_mapped_ptr (c_loc (a), d))) & + stop 3 + + !$omp target data map(alloc: a) device(d) + !$omp target map(from: p1) map(alloc: a) device(d) + p1 = c_loc (a); + !$omp end target + if (c_associated (omp_get_mapped_ptr (c_loc (a), d))) & + stop 4 + !$omp end target data + + !$omp target data map(alloc: b(1:0)) device(d) + !$omp target map(from: p2) map(alloc: b(1:0)) device(d) + p2 = c_loc (b(1)); + !$omp end target + if (c_associated (omp_get_mapped_ptr (c_loc (b(1)), d))) & + stop 5 + !$omp end target data + end if + call omp_target_free (p, d) +end program main diff --git a/libgomp/testsuite/libgomp.fortran/get-mapped-ptr-4.f90 b/libgomp/testsuite/libgomp.fortran/get-mapped-ptr-4.f90 new file mode 100644 index 0000000..4300a55 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/get-mapped-ptr-4.f90 @@ -0,0 +1,84 @@ +program main + use omp_lib + use iso_c_binding + implicit none (external, type) + integer :: d, id + type(c_ptr) :: p1, p2 + + type t + integer :: m1, m2 + end type t + type(t), target :: s + + d = omp_get_default_device () + id = omp_get_initial_device () + + if (d < 0 .or. d >= omp_get_num_devices ()) & + d = id + + if (d /= id) then + !$omp target data map(alloc: s, s%m2) device(d) + !$omp target map(from: p1, p2) map(alloc: s, s%m2) device(d) + p1 = c_loc (s); + p2 = c_loc (s%m2); + !$omp end target + + if (.not. c_associated (omp_get_mapped_ptr (c_loc (s), d), p1) & + .or. .not. c_associated (omp_get_mapped_ptr (c_loc (s%m2), d), p2)) & + stop 0 + !$omp end target data + + if (c_associated (omp_get_mapped_ptr (c_loc (s), d)) & + .or. c_associated (omp_get_mapped_ptr (c_loc (s%m2), d))) & + stop 1 + + !$omp target enter data map (alloc: s, s%m2) device (d) + !$omp target map(from: p1, p2) map(alloc: s, s%m2) device(d) + p1 = c_loc (s); + p2 = c_loc (s%m2); + !$omp end target + + if (.not. c_associated (omp_get_mapped_ptr (c_loc (s), d), p1) & + .or. .not. c_associated (omp_get_mapped_ptr (c_loc (s%m2), d), p2)) & + stop 2 + !$omp target exit data map (delete: s, s%m2) device (d) + + if (c_associated (omp_get_mapped_ptr (c_loc (s), d)) & + .or. c_associated (omp_get_mapped_ptr (c_loc (s%m2), d))) & + stop 3 + + else ! d == id + + !$omp target data map(alloc: s, s%m2) device(d) + !$omp target map(from: p1, p2) map(alloc: s, s%m2) device(d) + p1 = c_loc (s); + p2 = c_loc (s%m2); + !$omp end target + + if (.not. c_associated (omp_get_mapped_ptr (c_loc (s), d), c_loc (s)) & + .or. .not. c_associated (omp_get_mapped_ptr (c_loc (s%m2), d), c_loc (s%m2))) & + stop 4 + !$omp end target data + + if (.not. c_associated (omp_get_mapped_ptr (c_loc (s), d), c_loc (s)) & + .or. .not. c_associated (omp_get_mapped_ptr (c_loc (s%m2), d), c_loc (s%m2))) & + stop 5 + + !$omp target enter data map (alloc: s, s%m2) device (d) + !$omp target map(from: p1, p2) map(alloc: s, s%m2) device(d) + p1 = c_loc (s); + p2 = c_loc (s%m2); + !$omp end target + + if (.not. c_associated (omp_get_mapped_ptr (c_loc (s), d), c_loc (s)) & + .or. .not. c_associated (omp_get_mapped_ptr (c_loc (s%m2), d), c_loc (s%m2))) & + stop 6 + + !$omp target exit data map (delete: s, s%m2) device (d) + + if (.not. c_associated (omp_get_mapped_ptr (c_loc (s), d), c_loc (s)) & + .or. .not. c_associated (omp_get_mapped_ptr (c_loc (s%m2), d), c_loc (s%m2))) & + stop 7 + end if + +end program main