OpenMP: Permit in Fortran omp target data without map
gcc/fortran/ChangeLog: * openmp.c (resolve_omp_clauses): Permit 'omp target data' without map if use_device_{addr,ptr} is present. gcc/testsuite/ChangeLog: * gfortran.dg/gomp/map-3.f90: New test. * gfortran.dg/gomp/map-4.f90: New test.
This commit is contained in:
parent
9f6abd2db9
commit
d6cd139c17
|
@ -5330,17 +5330,23 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
|
|||
if (omp_clauses->depend_source && code->op != EXEC_OMP_ORDERED)
|
||||
gfc_error ("SOURCE dependence type only allowed "
|
||||
"on ORDERED directive at %L", &code->loc);
|
||||
if (!openacc && code && omp_clauses->lists[OMP_LIST_MAP] == NULL)
|
||||
if (!openacc
|
||||
&& code
|
||||
&& omp_clauses->lists[OMP_LIST_MAP] == NULL
|
||||
&& omp_clauses->lists[OMP_LIST_USE_DEVICE_PTR] == NULL
|
||||
&& omp_clauses->lists[OMP_LIST_USE_DEVICE_ADDR] == NULL)
|
||||
{
|
||||
const char *p = NULL;
|
||||
switch (code->op)
|
||||
{
|
||||
case EXEC_OMP_TARGET_DATA: p = "TARGET DATA"; break;
|
||||
case EXEC_OMP_TARGET_ENTER_DATA: p = "TARGET ENTER DATA"; break;
|
||||
case EXEC_OMP_TARGET_EXIT_DATA: p = "TARGET EXIT DATA"; break;
|
||||
default: break;
|
||||
}
|
||||
if (p)
|
||||
if (code->op == EXEC_OMP_TARGET_DATA)
|
||||
gfc_error ("TARGET DATA must contain at least one MAP, USE_DEVICE_PTR, "
|
||||
"or USE_DEVICE_ADDR clause at %L", &code->loc);
|
||||
else if (p)
|
||||
gfc_error ("%s must contain at least one MAP clause at %L",
|
||||
p, &code->loc);
|
||||
}
|
||||
|
|
|
@ -0,0 +1,38 @@
|
|||
! { dg-additional-options "-fdump-tree-original" }
|
||||
|
||||
subroutine bar
|
||||
integer, target :: x
|
||||
integer, allocatable, target :: y(:,:), z(:,:)
|
||||
x = 7
|
||||
!$omp target enter data map(to:x)
|
||||
|
||||
x = 8
|
||||
!$omp target data map(always, to: x)
|
||||
call foo(x)
|
||||
!$omp end target data
|
||||
|
||||
!$omp target data use_device_ptr(x)
|
||||
call foo2(x)
|
||||
!$omp end target data
|
||||
|
||||
!$omp target data use_device_addr(x)
|
||||
call foo2(x)
|
||||
!$omp end target data
|
||||
!$omp target exit data map(release:x)
|
||||
|
||||
!$omp target data map(y) use_device_addr(y)
|
||||
call foo3(y)
|
||||
!$omp end target data
|
||||
|
||||
!$omp target data map(z) use_device_ptr(z)
|
||||
call foo3(z)
|
||||
!$omp end target data
|
||||
end
|
||||
|
||||
! { dg-final { scan-tree-dump-times "#pragma omp target enter data map\\(to:x\\)" 1 "original" } }
|
||||
! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(always,to:x\\)" 1 "original" } }
|
||||
! { dg-final { scan-tree-dump-times "#pragma omp target data use_device_ptr\\(x\\)" 1 "original" } }
|
||||
! { dg-final { scan-tree-dump-times "#pragma omp target data use_device_addr\\(x\\)" 1 "original" } }
|
||||
! { dg-final { scan-tree-dump-times "#pragma omp target exit data map\\(release:x\\)" 1 "original" } }
|
||||
! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(tofrom:\\*\\(c_char \\*\\) y.data \\\[len: .*\\) map\\(to:y \\\[pointer set, len: .*\\) map\\(alloc:.*y.data \\\[pointer assign, bias: 0\\\]\\) use_device_addr\\(y\\)" 1 "original" } }
|
||||
! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(tofrom:\\*\\(c_char \\*\\) z.data \\\[len: .*\\) map\\(to:z \\\[pointer set, len: .*\\) map\\(alloc:.*z.data \\\[pointer assign, bias: 0\\\]\\) use_device_ptr\\(z\\)" 1 "original" } }
|
|
@ -0,0 +1,7 @@
|
|||
!$omp target enter data device(1) if (.true.) nowait ! { dg-error "TARGET ENTER DATA must contain at least one MAP clause" }
|
||||
|
||||
!$omp target data device(1) ! { dg-error "TARGET DATA must contain at least one MAP, USE_DEVICE_PTR, or USE_DEVICE_ADDR clause" }
|
||||
!$omp endtarget data
|
||||
|
||||
!$omp target exit data device(1) if (.true.) nowait ! { dg-error "TARGET EXIT DATA must contain at least one MAP clause" }
|
||||
end
|
Loading…
Reference in New Issue