[committed] libgomp/testsuite: Add requires-unified-addr-1.{c,f90} [PR109837]

[committed] libgomp/testsuite: Add requires-unified-addr-1.{c,f90} [PR109837]

[Tobias Burnus]

[-- Attachment #1: Type: text/plain, Size: 1519 bytes --]

Add a testcase for "omp requires unified_address" as we hadn't one.

The feature itself worked since the beginning (hardware + implementation
wise); that the devices report 'omp requires unified_address' as
supported is newer: for nvptx since r13-3460-g131d18e928a3ea and for GCN
since r14-1584-gf1af7d65ff64fe (a week ago).

The test assumes (→ dg-output) that all offload devices support
unified_address; this implies: if an offloading device is available, it
also remains available after adding the unified-address requirement.
Goal: ensure that we don't end up with only host fallback.

Unified address implies: Pointer size is the same such that no
'is_device_ptr' is required to convert an opaque pointer, it also
permits to do device-pointer pointer arithmetic on the host. This
testcase also assumes that 'int' / 'integer' has the same size on host
and device. (If not: good luck with offloading in general!)

The test also makes the sound assumption that derived-type component
pointers are passed through with the derived-type itself such that the
pointer address remains well defined. In terms of the standard, it would
have an undefined association status.

Committed as Rev. r14-1783-gd5c58ad1ebaff9

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

[-- Attachment #2: committed.diff --]
[-- Type: text/x-patch, Size: 6516 bytes --]

commit d5c58ad1ebaff924c2546df074174cffb128feb8
Author: Tobias Burnus <tobias@codesourcery.com>
Date:   Tue Jun 13 11:27:47 2023 +0200

    libgomp/testsuite: Add requires-unified-addr-1.{c,f90} [PR109837]
    
    Add a testcase for 'omp requires unified_address' that is currently supported
    by all devices but was not tested for.
    
    libgomp/
    
            PR libgomp/109837
            * testsuite/libgomp.c-c++-common/requires-unified-addr-1.c: New test.
            * testsuite/libgomp.fortran/requires-unified-addr-1.f90: New test.
---
 .../libgomp.c-c++-common/requires-unified-addr-1.c |  74 ++++++++++++++
 .../libgomp.fortran/requires-unified-addr-1.f90    | 111 +++++++++++++++++++++
 2 files changed, 185 insertions(+)

diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-unified-addr-1.c b/libgomp/testsuite/libgomp.c-c++-common/requires-unified-addr-1.c
new file mode 100644
index 00000000000..bff0a6b31ab
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/requires-unified-addr-1.c
@@ -0,0 +1,74 @@
+/* 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];
+
+  #pragma omp target firstprivate(qptr) map(tofrom:ptr2) device(device_num: dev)
+  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);
+
+  #pragma omp target map(to: ptr1) 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.fortran/requires-unified-addr-1.f90 b/libgomp/testsuite/libgomp.fortran/requires-unified-addr-1.f90
new file mode 100644
index 00000000000..f5a5adf093b
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/requires-unified-addr-1.f90
@@ -0,0 +1,111 @@
+! PR libgomp/109837
+
+program main
+  use iso_c_binding
+  use iso_fortran_env
+  use omp_lib
+  implicit none (external, type)
+  !$omp requires unified_address
+
+  integer(c_intptr_t), parameter :: N = 15
+  integer :: i, ntgts
+
+  ntgts = omp_get_num_devices();
+  if (ntgts > 0) then
+    write (ERROR_UNIT, '(a)') "Offloading devices exist"  ! { dg-output "Offloading devices exist(\n|\r\n|\r)" { target offload_device } }
+  else
+    write (ERROR_UNIT, '(a)') "Only host fallback"      ! { dg-output "Only host fallback(\n|\r\n|\r)" { target { ! offload_device } } }
+  endif
+
+  do i = 0, ntgts
+    call test_device (i);
+  end do
+
+contains
+
+  subroutine test_device (dev)
+    integer, value, intent(in) :: dev
+
+    type t
+      integer(c_intptr_t) :: n, m
+      integer, pointer :: fptr(:)
+      type(c_ptr) :: cptr      
+    end type t
+    type(t) :: s
+    type(c_ptr) :: cptr, qptr, cptr2, cptr2a
+    integer, target :: q(4)
+    integer, pointer :: fptr(:)
+    integer(c_intptr_t) :: i
+
+    s%n = 10;
+    s%m = 23;
+    s%cptr = omp_target_alloc (s%n * NUMERIC_STORAGE_SIZE/CHARACTER_STORAGE_SIZE, dev);
+    cptr = omp_target_alloc (s%m * NUMERIC_STORAGE_SIZE/CHARACTER_STORAGE_SIZE, dev);
+    if (.not. c_associated(s%cptr)) stop 1
+    if (.not. c_associated(cptr)) stop 2
+    call c_f_pointer (cptr, s%fptr, [s%m])
+
+    cptr = omp_target_alloc (N * NUMERIC_STORAGE_SIZE/CHARACTER_STORAGE_SIZE, dev);
+    if (.not. c_associated(cptr)) stop 3
+
+    q = [1, 2, 3, 4]
+    !$omp target enter data map(q) device(device_num: dev)
+    !$omp target data use_device_addr(q) device(device_num: dev)
+       qptr = c_loc(q)
+    !$omp end target data
+
+    !$omp target map(to:s) device(device_num: dev)
+    block
+      integer, pointer :: iptr(:)
+      call c_f_pointer(s%cptr, iptr, [s%n])
+      do i = 1, s%n
+        iptr(i) = 23 * int(i)
+      end do
+      do i = 1, s%m
+        s%fptr(i) = 35 * int(i)
+      end do
+    end block
+
+    cptr2 = c_loc(s%fptr(4))
+    cptr2a = s%cptr
+
+    !$omp target firstprivate(qptr) map(tofrom: cptr2) map(to :cptr2a) device(device_num: dev)
+    block
+      integer, pointer :: iptr(:), iptr2(:), qvar(:)
+      call c_f_pointer(cptr2, iptr, [4])
+      call c_f_pointer(cptr2a, iptr2, [4])
+      call c_f_pointer(qptr, qvar, [4])
+      qvar = iptr + iptr2
+    end block
+
+    !$omp target exit data map(q) device(device_num: dev)
+    do i = 1, 4
+      if (q(i) /= 23 * int(i)  +  35 * (int(i) + 4 - 1)) stop 4
+    end do
+
+    !$omp target map(to: cptr) device(device_num: dev)
+    block
+      integer, pointer :: p(:)
+      call c_f_pointer(cptr, p, [N])
+      do i = 1, N
+        p(i) = 11 * int(i)
+      end do
+    end block
+
+    allocate(fptr(N))
+    if (0 /= omp_target_memcpy (c_loc(fptr), cptr,  &
+                                N * NUMERIC_STORAGE_SIZE/CHARACTER_STORAGE_SIZE,  &
+                                0_c_intptr_t, 0_c_intptr_t, &
+                                omp_get_initial_device(), dev))  &
+      stop 5
+
+    do i = 1, N
+      if (fptr(i) /= 11 * int(i)) stop 6
+    end do
+
+    deallocate (fptr);
+    call omp_target_free (cptr, dev);
+    call omp_target_free (s%cptr, dev);
+    call omp_target_free (c_loc(s%fptr), dev);
+  end
+end