diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c index 0fd998839b2..16f39a4e086 100644 --- a/gcc/fortran/openmp.c +++ b/gcc/fortran/openmp.c @@ -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); } diff --git a/gcc/testsuite/gfortran.dg/gomp/map-3.f90 b/gcc/testsuite/gfortran.dg/gomp/map-3.f90 new file mode 100644 index 00000000000..13f63647bda --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/map-3.f90 @@ -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" } } diff --git a/gcc/testsuite/gfortran.dg/gomp/map-4.f90 b/gcc/testsuite/gfortran.dg/gomp/map-4.f90 new file mode 100644 index 00000000000..3aa1c8096d7 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/map-4.f90 @@ -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