From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: by sourceware.org (Postfix, from userid 1534) id 9A2E2385841D; Mon, 19 Jun 2023 07:09:31 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 9A2E2385841D DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=gcc.gnu.org; s=default; t=1687158571; bh=IPE3+JtVY1C0xzvmY+kTSGp0mFfqyb2tUAKjOlfRAqg=; h=From:To:Subject:Date:From; b=vqWJYpm3+ikhwxM/xhJGLocCbWcNsj9HT1mQ/HC9LWQfQlyaFEljGteYsYyVlpyji xVJuukwFpVtS7raGQ00tH/2eDbSFjqO6DM+Tt2yodRTkg/H1xFOJU9RBNKnTpEsx15 TIGqbG5epg9Dj4DQ0ekEETh3n4ffyemC5pwzLa7I= MIME-Version: 1.0 Content-Transfer-Encoding: 7bit Content-Type: text/plain; charset="utf-8" From: Tobias Burnus To: gcc-cvs@gcc.gnu.org Subject: [gcc r14-1923] OpenMP (C/C++): Keep pointer value of unmapped ptr with default mapping [PR110270] X-Act-Checkin: gcc X-Git-Author: Tobias Burnus X-Git-Refname: refs/heads/master X-Git-Oldrev: 53953b6f31f18ac2e2241f0c1f3c8d7ecec78e7f X-Git-Newrev: b25ea7ab78cdb7baec694e56eadb71002726a73e Message-Id: <20230619070931.9A2E2385841D@sourceware.org> Date: Mon, 19 Jun 2023 07:09:31 +0000 (GMT) List-Id: https://gcc.gnu.org/g:b25ea7ab78cdb7baec694e56eadb71002726a73e commit r14-1923-gb25ea7ab78cdb7baec694e56eadb71002726a73e Author: Tobias Burnus Date: Mon Jun 19 09:08:51 2023 +0200 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. Diff: --- 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 b6a7214ab4f..80c25a16f1e 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -1153,7 +1153,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..863cf0e28a7 --- /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..d0b0cd178c0 --- /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)