mirror of
https://github.com/gcc-mirror/gcc.git
synced 2024-11-21 13:40:47 +00:00
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.
This commit is contained in:
Tobias Burnus
parent
cca8d9e5be
commit
d5c58ad1eb
@ -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;
|
||||
}
|
||||
111
libgomp/testsuite/libgomp.fortran/requires-unified-addr-1.f90
Normal file
111
libgomp/testsuite/libgomp.fortran/requires-unified-addr-1.f90
Normal file
@ -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
|
||||
Loading…
Reference in New Issue
Block a user