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.
This commit is contained in:
Tobias Burnus
parent
53953b6f31
commit
b25ea7ab78
8 changed files with 392 additions and 16 deletions
|
|
@ -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
|
||||
|
|
|
|||
|
|
@ -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;
|
||||
}
|
||||
}
|
||||
|
|
|
|||
|
|
@ -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)
|
||||
|
|
|
|||
|
|
@ -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;
|
||||
}
|
||||
|
|
|
|||
|
|
@ -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 <assert.h>
|
||||
#include <omp.h>
|
||||
#include <stdio.h>
|
||||
#include <stdlib.h>
|
||||
|
||||
#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;
|
||||
}
|
||||
105
libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-3.c
Normal file
105
libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-3.c
Normal file
|
|
@ -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 <stdio.h>
|
||||
#include <stdint.h>
|
||||
#include <omp.h>
|
||||
|
||||
/* '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;
|
||||
}
|
||||
159
libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-4.c
Normal file
159
libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-4.c
Normal file
|
|
@ -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 <stdio.h>
|
||||
#include <stdlib.h>
|
||||
#include <stdint.h>
|
||||
#include <omp.h>
|
||||
|
||||
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;
|
||||
}
|
||||
|
|
@ -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)
|
||||
|
|
|
|||
Loading…
Add table
Reference in a new issue