From patchwork Fri Jun 16 16:17:18 2023 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Tobias Burnus X-Patchwork-Id: 109244 Return-Path: Delivered-To: ouuuleilei@gmail.com Received: by 2002:a59:994d:0:b0:3d9:f83d:47d9 with SMTP id k13csp1459718vqr; Fri, 16 Jun 2023 09:18:10 -0700 (PDT) X-Google-Smtp-Source: ACHHUZ5YQ49kP5nOGOIfq88fPHIGUpFUFRd7LxOqf2f38dRhU7NG/McVBlfEFbncHo9u6ky9INbT X-Received: by 2002:a05:6402:5112:b0:514:387c:930 with SMTP id m18-20020a056402511200b00514387c0930mr8635746edd.5.1686932289794; Fri, 16 Jun 2023 09:18:09 -0700 (PDT) ARC-Seal: i=1; a=rsa-sha256; t=1686932289; cv=none; d=google.com; s=arc-20160816; b=Xq6cooxKv0RpVhO8X0ez9NCEU/kG4SDue2qUVj8+TMxnKNPXwh5u7fx7nxo5DiWuJs WrTARrIib3Czurj6Us3tpIu4FKDJ/vwnST9ymkUcVuc0BIFpW8seTxjPluQchzyj25Pi jW98IT+Ao8Ife7wKPvQE4dVHU7im0NeE65oFtsSwxeVVVp/cF8hTE9ZT/XefnvoENfBn Moyr0CZQHWiciJlVE1gebuQzgd8DdrSQ+Gui7xCE3RPL2O8HFUMabQXxNSjXw1kNOYdT LZVL445eHm9PqoYSK6V40Vru8YKqOPSVXceYbfxOacumVarsMxGrv1/By0MwaCYa0ohp wUvg== 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=Pu5wr0u/GuCR5SiEsfeIaSeOCGOXZek8Xvxe+FmorFw=; b=oRvl783BegHi72HEx5CPFWTnriwGJRMBPyo+XeR5i6G/B41BygxEWK00K0EGOnIg3B aqZJGf4HmZ93lrck1neBohT7F/Xc10NE9gpTdbvQL89fUu9YBTAtws83bfar9T13+FoR tE3SIpLp+bsFLJVwOsBsVAZ8EkME+Kmd3J5vOrB+n6CbLAWl+5/D7rfEiXSLZ/jCGCo/ kuEkXFxGkSDDNWWCCYrRclDYy4OXTxj7S4DMclRZmf4Ve6q61y02qarRxKlxqh+vDoKy H3aLDymEuu/fUsLucZ3zaLJmYnHCn0UtW6S6LHJXkPotxBWiniLxEN3+KEYZTqG4zTua mAbQ== 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 b14-20020aa7c6ce000000b0051a3281db25si1134949eds.680.2023.06.16.09.18.09 for (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Fri, 16 Jun 2023 09:18:09 -0700 (PDT) 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 DB545385701E for ; Fri, 16 Jun 2023 16:17:55 +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 3A1CF3858D35 for ; Fri, 16 Jun 2023 16:17:28 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 3A1CF3858D35 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="6.00,247,1681200000"; d="diff'?scan'208";a="8998340" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa3.mentor.iphmx.com with ESMTP; 16 Jun 2023 08:17:24 -0800 IronPort-SDR: znPjk3Fpw67TRhA8Aaf5wryW6uujZgwfac+Ah9jzOlE8HkxVbVcEPci656zvf+og1+PKkGF8l8 wvoMEa6eaGgKpxkGvC+zqKKLBFm/l4KhilpORW/leAYZ1IyOhGFKc1TNHP2pK1GdldKFsR0DcN JrbaUbXmq344qz0Hb9dTw+DjjqrEFTYal/7rAP6sYKnjINP0BwCDVZSdifDLhTnYWfW+4umPv8 kdhL3NL5e1C4hVa8lVxot241tdmv9KuMXklZq94zmkiLcRHpG0oawUc1RPgHA+WrJrtCqRk3yg cy8= Message-ID: Date: Fri, 16 Jun 2023 18:17:18 +0200 MIME-Version: 1.0 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:102.0) Gecko/20100101 Thunderbird/102.12.0 Content-Language: en-US To: gcc-patches , Jakub Jelinek From: Tobias Burnus Subject: [patch] OpenMP (C/C++): Keep pointer value of unmapped ptr with default mapping [PR110270] X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-10.mgc.mentorg.com (139.181.222.10) To svr-ies-mbx-12.mgc.mentorg.com (139.181.222.12) X-Spam-Status: No, score=-11.3 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.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?1768876712723770812?= X-GMAIL-MSGID: =?utf-8?q?1768876712723770812?= This fixes an issue related to OpenMP C/C++'s default mapping of pointer variables. (That's 'defaultmap(default:pointer)' – which is possibly surprisingly *not* the same as 'defaultmap(firstprivate:pointer)'). Namely, OpenMP supports the following: int *ptr = malloc(sizeof(int)*5); #pragma omp target enter data map(ptr[:5]) #pragma omp target p[2] = 5; which matches 'firstprivate(p)' + attaching the device address of 'p[:0]' (a zero-sized array), the latter making it possible to use 'p' automatically without the need to add any map clauses at least as long as *p has been mapped before. However, for int *ptr = omp_target_alloc (sizeof(int)*5, dev_num); #pragma omp target p[2] = 5; or for #pragma omp requires unified_shared_memory int pa = &A[0]; #pragma omp target pa[0] = 6; it failed before because neither 'ptr' nor 'pa' were mapped. Solution as a user was either to add a (default)map clause (with map type than 'default'), a firstprivate or an is_device_ptr clause. The problem was that with default mapping, p and pa had the value NULL in the example above on the device. (As required by OpenMP 5.0/5.1). With the commit, they retain the original value avoiding surprises for the code above. (See PR for the reference to the relevant sections of the OpenMP 5.{0,1,2} specifications.) I would love if someone would give a review it; albeit the actual code change in libgomp/target.c is just a changing a single enum value. If there are no comments, I intent to push it next week ... Tobias ----------------- 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 (C/C++): Keep pointer value of unmapped ptr with default mapping [PR110270] For C/C++ pointers, default implicit mapping firstprivatizes the pointer but if the memory it points to is mapped, the it is updated to point to the device memory (by attaching a zero sized array section of the pointed-to storage). However, if the pointed-to storage wasn't mapped, the pointer was set to NULL on the device side (OpenMP 5.0/5.1 semantic). With this commit, the pointer retains the on-host address in that case (OpenMP 5.2 semantic). The new semantic avoids an explicit map/firstprivate/is_device_ptr in the following sensible cases: Special values (e.g. pointer or 0x1, 0x2 etc.), explicitly device allocated memory (e.g. omp_target_alloc), and with (unified) shared memory. (Note: With (U)SM, mappings still must be tracked, at least when omp_target_associate_ptr does not fail when passing in two destinct pointers.) libgomp/ PR middle-end/110270 * target.c (gomp_map_vars_internal): Copy host value instead of NULL for GOMP_MAP_ZERO_LEN_ARRAY_SECTION if not mapped. * libgomp.texi (OpenMP 5.2 Impl.): Mark as 'Y'. * testsuite/libgomp.c/target-19.c: Update expected value. * testsuite/libgomp.c++/target-18.C: Likewise. * testsuite/libgomp.c++/target-19.C: Likewise. * testsuite/libgomp.c-c++-common/requires-unified-addr-2.c: New test. * testsuite/libgomp.c-c++-common/target-implicit-map-3.c: New test. * testsuite/libgomp.c-c++-common/target-implicit-map-4.c: New test. libgomp/libgomp.texi | 2 +- libgomp/target.c | 2 +- libgomp/testsuite/libgomp.c++/target-18.C | 21 ++- libgomp/testsuite/libgomp.c++/target-19.C | 13 +- .../libgomp.c-c++-common/requires-unified-addr-2.c | 85 +++++++++++ .../libgomp.c-c++-common/target-implicit-map-3.c | 105 ++++++++++++++ .../libgomp.c-c++-common/target-implicit-map-4.c | 159 +++++++++++++++++++++ libgomp/testsuite/libgomp.c/target-19.c | 21 ++- 8 files changed, 392 insertions(+), 16 deletions(-) diff --git a/libgomp/libgomp.texi b/libgomp/libgomp.texi index 1c57f5aa261..db8b1f1427e 100644 --- a/libgomp/libgomp.texi +++ b/libgomp/libgomp.texi @@ -384,7 +384,7 @@ to address of matching mapped list item per 5.1, Sect. 2.21.7.2 @tab N @tab @item @code{declare mapper} with iterator and @code{present} modifiers @tab N @tab @item If a matching mapped list item is not found in the data environment, the - pointer retains its original value @tab N @tab + pointer retains its original value @tab Y @tab @item New @code{enter} clause as alias for @code{to} on declare target directive @tab Y @tab @item Deprecation of @code{to} clause on declare target directive @tab N @tab diff --git a/libgomp/target.c b/libgomp/target.c index e39ef8f6e82..aa2410c0f16 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -1149,7 +1149,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, if (!n) { tgt->list[i].key = NULL; - tgt->list[i].offset = OFFSET_POINTER; + tgt->list[i].offset = OFFSET_INLINED; continue; } } diff --git a/libgomp/testsuite/libgomp.c++/target-18.C b/libgomp/testsuite/libgomp.c++/target-18.C index f1085b14022..a21ed4e81f9 100644 --- a/libgomp/testsuite/libgomp.c++/target-18.C +++ b/libgomp/testsuite/libgomp.c++/target-18.C @@ -20,7 +20,9 @@ foo (int *&p, int *&q, int *&r, int n, int m) err = 1; if (sep) { - if (q != (int *) 0 || r != (int *) 0) + /* Since OpenMP 5.2, 'q'/'r' are no longer set to NULL if pointing to + unmapped storage. */ + if (q == (int *) 0 || r == (int *) 0) err = 1; } else if (p + 8 != q || r != s) @@ -37,7 +39,9 @@ foo (int *&p, int *&q, int *&r, int n, int m) err = 1; if (sep) { - if (q != (int *) 0 || r != (int *) 0) + /* Since OpenMP 5.2, 'q'/'r' are no longer set to NULL if pointing to + unmapped storage. */ + if (q == (int *) 0 || r == (int *) 0) err = 1; } else if (p + 8 != q || r != s) @@ -55,7 +59,9 @@ foo (int *&p, int *&q, int *&r, int n, int m) err = 1; if (sep) { - if (q != (int *) 0 || r != (int *) 0) + /* Since OpenMP 5.2, 'q'/'r' are no longer set to NULL if pointing to + unmapped storage. */ + if (q == (int *) 0 || r == (int *) 0) err = 1; } else if (p + 8 != q || r != s) @@ -91,7 +97,8 @@ foo (int *&p, int *&q, int *&r, int n, int m) err = 1; else if (sep) { - if (r != (int *) 0) + /* Since OpenMP 5.2, 'r' is no longer set to NULL if *r is unmapped.*/ + if (r == (int *) 0) err = 1; } else if (r != q + 1) @@ -110,7 +117,8 @@ foo (int *&p, int *&q, int *&r, int n, int m) err = 1; else if (sep) { - if (r != (int *) 0) + /* Since OpenMP 5.2, 'r' is no longer set to NULL if *r is unmapped.*/ + if (r == (int *) 0) err = 1; } else if (r != q + 1) @@ -130,7 +138,8 @@ foo (int *&p, int *&q, int *&r, int n, int m) err = 1; else if (sep) { - if (r != (int *) 0) + /* Since OpenMP 5.2, 'r' is no longer set to NULL if *r is unmapped.*/ + if (r == (int *) 0) err = 1; } else if (r != q + 1) diff --git a/libgomp/testsuite/libgomp.c++/target-19.C b/libgomp/testsuite/libgomp.c++/target-19.C index afa6e68d5cc..7bae31d2734 100644 --- a/libgomp/testsuite/libgomp.c++/target-19.C +++ b/libgomp/testsuite/libgomp.c++/target-19.C @@ -1,3 +1,8 @@ +/* { dg-additional-options "-O0" } */ +/* Disable optimization to ensure that the compiler does not exploit that + S::r + t will never be NULL due to int (&r) and (&t). */ + + extern "C" void abort (); struct S { char a[64]; int (&r)[2]; char b[64]; }; @@ -19,7 +24,9 @@ foo (S s, int (&t)[3], int z) #pragma omp target map(from: err) map(tofrom: s.r[:0], t[:0]) { if (sep) - err = s.r != (int *) 0 || t != (int *) 0; + /* Since OpenMP 5.2, if no matching mapped list it has been found, + pointers retain their original value. */ + err = s.r == (int *) 0 || t == (int *) 0; else err = t[0] != 1 || t[1] != 2 || t[2] != 3 || s.r[0] != 6 || s.r[1] != 7; } @@ -28,7 +35,9 @@ foo (S s, int (&t)[3], int z) #pragma omp target map(from: err) map(tofrom: s.r[:z], t[:z]) { if (sep) - err = s.r != (int *) 0 || t != (int *) 0; + /* Since OpenMP 5.2, if no matching mapped list it has been found, + pointers retain their original value. */ + err = s.r == (int *) 0 || t == (int *) 0; else err = t[0] != 1 || t[1] != 2 || t[2] != 3 || s.r[0] != 6 || s.r[1] != 7; } diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-unified-addr-2.c b/libgomp/testsuite/libgomp.c-c++-common/requires-unified-addr-2.c new file mode 100644 index 00000000000..3b5dcd38c1a --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/requires-unified-addr-2.c @@ -0,0 +1,85 @@ +/* PR middle-end/110270 */ + +/* OpenMP 5.2's 'defaultmap(default : pointer) for C/C++ pointers retains the + pointer value instead of setting it to NULL if the pointer cannot be found. + Contrary to requires-unified-addr-1.c which is valid OpenMP 5.0/5.1/5.2, + this testcase is only valid since OpenMP 5.2. */ + +/* This is kind of a follow-up to the requires-unified-addr-1.c testcase + and PR libgomp/109837 */ + + +#include +#include +#include +#include + +#pragma omp requires unified_address + +#define N 15 + +void +test_device (int dev) +{ + struct st { + int *ptr; + int n; + }; + struct st s; + + s.n = 10; + s.ptr = (int *) omp_target_alloc (sizeof (int)*s.n, dev); + int *ptr1 = (int *) omp_target_alloc (sizeof (int)*N, dev); + assert (s.ptr != NULL); + assert (ptr1 != NULL); + + int q[4] = {1,2,3,4}; + int *qptr; + #pragma omp target enter data map(q) device(device_num: dev) + #pragma omp target data use_device_addr(q) device(device_num: dev) + qptr = q; + + #pragma omp target map(to:s) device(device_num: dev) + for (int i = 0; i < s.n; i++) + s.ptr[i] = 23*i; + + int *ptr2 = &s.ptr[3]; + + /* s.ptr is not mapped (but omp_target_alloc'ed) thus ptr2 shall retain its value. */ + #pragma omp target device(device_num: dev) /* implied: defaultmap(default : pointer) */ + for (int i = 0; i < 4; i++) + *(qptr++) = ptr2[i]; + + #pragma omp target exit data map(q) device(device_num: dev) + for (int i = 0; i < 4; i++) + q[i] = 23 * (i+3); + + /* ptr1 retains the value as it is not mapped (but it is omp_target_alloc'ed). */ + #pragma omp target defaultmap(default : pointer) device(device_num: dev) + for (int i = 0; i < N; i++) + ptr1[i] = 11*i; + + int *ptr3 = (int *) malloc (sizeof (int)*N); + assert (0 == omp_target_memcpy(ptr3, ptr1, N * sizeof(int), 0, 0, + omp_get_initial_device(), dev)); + for (int i = 0; i < N; i++) + assert (ptr3[i] == 11*i); + + free (ptr3); + omp_target_free (ptr1, dev); + omp_target_free (s.ptr, dev); +} + +int +main() +{ + int ntgts = omp_get_num_devices(); + if (ntgts) + fprintf (stderr, "Offloading devices exist\n"); /* { dg-output "Offloading devices exist(\n|\r\n|\r)" { target offload_device } } */ + else + fprintf (stderr, "Only host fallback\n"); /* { dg-output "Only host fallback(\n|\r\n|\r)" { target { ! offload_device } } } */ + + for (int i = 0; i <= ntgts; i++) + test_device (i); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-3.c b/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-3.c new file mode 100644 index 00000000000..e69f6d0db2f --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-3.c @@ -0,0 +1,105 @@ +/* PR middle-end/110270 */ + +/* Ensure that defaultmap(default : pointer) uses correct OpenMP 5.2 + semantic, i.e. keeping the pointer value even if not mapped; + before OpenMP 5.0/5.1 required that it is NULL, causing issues + especially with unified-shared memory but also the code below + shows why that's not a good idea. */ + +#include +#include +#include + +/* 'unified_address' is required by the OpenMP spec as only then + 'is_device_ptr' can be left out. All our devices support this + while remote offloading would not. However, in practice it is + sufficient that the host and device pointer size is the same + (or the device pointer is smaller) - and then a device pointer is + representable and omp_target_alloc can return a bare device pointer. + + We here assume that this weaker condition holds and do not + require: #pragma omp requires unified_address */ + +void +test_device (int dev) +{ + int *p1 = (int*) 0x12345; + int *p1a = (int*) 0x67890; + int *p2 = (int*) omp_target_alloc (sizeof (int) * 5, dev); + int *p2a = (int*) omp_target_alloc (sizeof (int) * 10, dev); + intptr_t ip = (intptr_t) p2; + intptr_t ipa = (intptr_t) p2a; + + int A[3] = {1,2,3}; + int B[5] = {4,5,6,7,8}; + int *p3 = &A[0]; + int *p3a = &B[0]; + + #pragma omp target enter data map(to:A) device(dev) + + #pragma omp target device(dev) /* defaultmap(default:pointer) */ + { + /* The pointees aren't mapped. */ + /* OpenMP 5.2 -> same value as before the target region. */ + if ((intptr_t) p1 != 0x12345) __builtin_abort (); + if ((intptr_t) p2 != ip) __builtin_abort (); + for (int i = 0; i < 5; i++) + p2[i] = 13*i; + + for (int i = 0; i < 10; i++) + ((int *)ipa)[i] = 7*i; + + /* OpenMP: Mapped => must point to the corresponding device storage of 'A' */ + if (p3[0] != 1 || p3[1] != 2 || p3[2] != 3) + __builtin_abort (); + p3[0] = -11; p3[1] = -22; p3[2] = -33; + } + #pragma omp target exit data map(from:A) device(dev) + + if (p3[0] != -11 || p3[1] != -22 || p3[2] != -33) + __builtin_abort (); + + // With defaultmap: + + #pragma omp target enter data map(to:B) device(dev) + + #pragma omp target device(dev) defaultmap(default:pointer) + { + /* The pointees aren't mapped. */ + /* OpenMP 5.2 -> same value as before the target region. */ + if ((intptr_t) p1a != 0x67890) __builtin_abort (); + if ((intptr_t) p2a != ipa) __builtin_abort (); + + for (int i = 0; i < 5; i++) + ((int *)ip)[i] = 13*i; + + for (int i = 0; i < 10; i++) + p2a[i] = 7*i; + + /* OpenMP: Mapped => must point to the corresponding device storage of 'B' */ + if (p3a[0] != 4 || p3a[1] != 5 || p3a[2] != 6 || p3a[3] != 7 || p3a[4] != 8) + __builtin_abort (); + p3a[0] = -44; p3a[1] = -55; p3a[2] = -66; p3a[3] = -77; p3a[4] = -88; + } + #pragma omp target exit data map(from:B) device(dev) + + if (p3a[0] != -44 || p3a[1] != -55 || p3a[2] != -66 || p3a[3] != -77 || p3a[4] != -88) + __builtin_abort (); + + omp_target_free (p2, dev); + omp_target_free (p2a, dev); +} + +int +main() +{ + int ntgts = omp_get_num_devices(); + if (ntgts) + fprintf (stderr, "Offloading devices exist\n"); /* { dg-output "Offloading devices exist(\n|\r\n|\r)" { target offload_device } } */ + else + fprintf (stderr, "Only host fallback\n"); /* { dg-output "Only host fallback(\n|\r\n|\r)" { target { ! offload_device } } } */ + + for (int i = 0; i <= ntgts; i++) + test_device (i); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-4.c b/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-4.c new file mode 100644 index 00000000000..f92abfdb30d --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-4.c @@ -0,0 +1,159 @@ +/* PR middle-end/110270 */ + +/* Same as target-implicit-map-3.c but uses the following requiement + and for not mapping the stack variables 'A' and 'B' (not mapped + but accessible -> USM makes this tested feature even more important.) */ + +#pragma omp requires unified_shared_memory + +/* Ensure that defaultmap(default : pointer) uses correct OpenMP 5.2 + semantic, i.e. keeping the pointer value even if not mapped; + before OpenMP 5.0/5.1 required that it is NULL. */ + +#include +#include +#include +#include + +void +test_device (int dev) +{ + int *p1 = (int*) 0x12345; + int *p1a = (int*) 0x67890; + int *p2 = (int*) omp_target_alloc (sizeof (int) * 5, dev); + int *p2a = (int*) omp_target_alloc (sizeof (int) * 10, dev); + intptr_t ip = (intptr_t) p2; + intptr_t ipa = (intptr_t) p2a; + + int A[3] = {1,2,3}; + int B[5] = {4,5,6,7,8}; + int *p3 = &A[0]; + int *p3a = &B[0]; + + const omp_alloctrait_t traits[] + = { { omp_atk_alignment, 128 }, + { omp_atk_pool_size, 1024 }}; + omp_allocator_handle_t a = omp_init_allocator (omp_default_mem_space, 2, traits); + + int *p4 = (int*) malloc (sizeof (int) * 5); + int *p4a = (int*) omp_alloc (sizeof (int) * 10, a); + intptr_t ip4 = (intptr_t) p4; + intptr_t ip4a = (intptr_t) p4a; + + for (int i = 0; i < 5; i++) + p4[i] = -31*i; + + for (int i = 0; i < 10; i++) + p4a[i] = -43*i; + + /* Note: 'A' is not mapped but USM accessible. */ + #pragma omp target device(dev) /* defaultmap(default:pointer) */ + { + /* The pointees aren't mapped. */ + /* OpenMP 5.2 -> same value as before the target region. */ + if ((intptr_t) p1 != 0x12345) abort (); + if ((intptr_t) p2 != ip) abort (); + for (int i = 0; i < 5; i++) + p2[i] = 13*i; + + for (int i = 0; i < 10; i++) + ((int *)ipa)[i] = 7*i; + + /* OpenMP: Points to 'A'. */ + if (p3[0] != 1 || p3[1] != 2 || p3[2] != 3) + abort (); + p3[0] = -11; p3[1] = -22; p3[2] = -33; + + /* USM accesible allocated host memory. */ + if ((intptr_t) p4 != ip4) + abort (); + for (int i = 0; i < 5; i++) + if (p4[i] != -31*i) + abort (); + for (int i = 0; i < 10; i++) + if (((int *)ip4a)[i] != -43*i) + abort (); + for (int i = 0; i < 5; i++) + p4[i] = 9*i; + for (int i = 0; i < 10; i++) + ((int *)ip4a)[i] = 18*i; + } + + if (p3[0] != -11 || p3[1] != -22 || p3[2] != -33) + abort (); + + for (int i = 0; i < 5; i++) + if (p4[i] != 9*i) + abort (); + for (int i = 0; i < 10; i++) + if (p4a[i] != 18*i) + abort (); + for (int i = 0; i < 5; i++) + p4[i] = -77*i; + for (int i = 0; i < 10; i++) + p4a[i] = -65*i; + + // With defaultmap: + + /* Note: 'B' is not mapped but USM accessible. */ + #pragma omp target device(dev) defaultmap(default:pointer) + { + /* The pointees aren't mapped. */ + /* OpenMP 5.2 -> same value as before the target region. */ + if ((intptr_t) p1a != 0x67890) abort (); + if ((intptr_t) p2a != ipa) abort (); + + for (int i = 0; i < 5; i++) + ((int *)ip)[i] = 13*i; + + for (int i = 0; i < 10; i++) + p2a[i] = 7*i; + + /* USM accesible allocated host memory. */ + if ((intptr_t) p4a != ip4a) abort (); + + /* OpenMP: Points to 'B'. */ + if (p3a[0] != 4 || p3a[1] != 5 || p3a[2] != 6 || p3a[3] != 7 || p3a[4] != 8) + abort (); + p3a[0] = -44; p3a[1] = -55; p3a[2] = -66; p3a[3] = -77; p3a[4] = -88; + + /* USM accesible allocated host memory. */ + if ((intptr_t) p4a != ip4a) + abort (); + for (int i = 0; i < 5; i++) + if (((int *)ip4)[i] != -77*i) + abort (); + for (int i = 0; i < 10; i++) + if (p4a[i] != -65*i) + abort (); + for (int i = 0; i < 5; i++) + p4[i] = 36*i; + for (int i = 0; i < 10; i++) + ((int *)ip4a)[i] = 4*i; + } + + if (p3a[0] != -44 || p3a[1] != -55 || p3a[2] != -66 || p3a[3] != -77 || p3a[4] != -88) + abort (); + + for (int i = 0; i < 5; i++) + if (p4[i] != 36*i) + abort (); + for (int i = 0; i < 10; i++) + if (p4a[i] != 4*i) + abort (); + + omp_target_free (p2, dev); + omp_target_free (p2a, dev); + free (p4); + omp_free (p4a, a); + omp_destroy_allocator (a); +} + +int +main() +{ + int ntgts = omp_get_num_devices(); + for (int i = 0; i <= ntgts; i++) + test_device (i); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/target-19.c b/libgomp/testsuite/libgomp.c/target-19.c index 2505cafca9f..dac7c56cbde 100644 --- a/libgomp/testsuite/libgomp.c/target-19.c +++ b/libgomp/testsuite/libgomp.c/target-19.c @@ -20,7 +20,9 @@ foo (int *p, int *q, int *r, int n, int m) err = 1; if (sep) { - if (q != (int *) 0 || r != (int *) 0) + /* Since OpenMP 5.2, 'q'/'r' are no longer set to NULL if pointing to + unmapped storage. */ + if (q == (int *) 0 || r == (int *) 0) err = 1; } else if (p + 8 != q || r != s) @@ -37,7 +39,9 @@ foo (int *p, int *q, int *r, int n, int m) err = 1; if (sep) { - if (q != (int *) 0 || r != (int *) 0) + /* Since OpenMP 5.2, 'q'/'r' are no longer set to NULL if pointing to + unmapped storage. */ + if (q == (int *) 0 || r == (int *) 0) err = 1; } else if (p + 8 != q || r != s) @@ -55,7 +59,9 @@ foo (int *p, int *q, int *r, int n, int m) err = 1; if (sep) { - if (q != (int *) 0 || r != (int *) 0) + /* Since OpenMP 5.2, 'q'/'r' are no longer set to NULL if pointing to + unmapped storage. */ + if (q == (int *) 0 || r == (int *) 0) err = 1; } else if (p + 8 != q || r != s) @@ -91,7 +97,8 @@ foo (int *p, int *q, int *r, int n, int m) err = 1; else if (sep) { - if (r != (int *) 0) + /* Since OpenMP 5.2, 'r' is no longer set to NULL if *r is unmapped.*/ + if (r == (int *) 0) err = 1; } else if (r != q + 1) @@ -110,7 +117,8 @@ foo (int *p, int *q, int *r, int n, int m) err = 1; else if (sep) { - if (r != (int *) 0) + /* Since OpenMP 5.2, 'r' is no longer set to NULL if *r is unmapped.*/ + if (r == (int *) 0) err = 1; } else if (r != q + 1) @@ -130,7 +138,8 @@ foo (int *p, int *q, int *r, int n, int m) err = 1; else if (sep) { - if (r != (int *) 0) + /* Since OpenMP 5.2, 'r' is no longer set to NULL if *r is unmapped.*/ + if (r == (int *) 0) err = 1; } else if (r != q + 1)