mirror/gcc
2
0
Fork 0
mirror of git://gcc.gnu.org/git/gcc.git synced 2025-04-10 05:50:28 +08:00
Code Issues Packages Projects Releases Wiki Activity

OpenMP/Fortran: Fixes for {use,is}_device_ptr

Browse Source 
gcc/fortran/ChangeLog:

	PR fortran/98476
	* openmp.c (resolve_omp_clauses): Change use_device_ptr
	to use_device_addr for unless type(c_ptr); check all
	list item for is_device_ptr.

gcc/ChangeLog:

	PR fortran/98476
	* omp-low.c (lower_omp_target): Handle nonpointer is_device_ptr.

libgomp/ChangeLog:

	PR fortran/98476
	* testsuite/libgomp.fortran/is_device_ptr-1.f90: New test.

gcc/testsuite/ChangeLog:

	PR fortran/98476
	* gfortran.dg/gomp/map-3.f90: Update expected scan-dump-tree.
	* gfortran.dg/gomp/is_device_ptr-2.f90: New test.
	* gfortran.dg/gomp/use_device_ptr-1.f90: New test.
This commit is contained in:
Tobias Burnus 2021-01-19 11:57:34 +01:00
parent 9b8741c98f
commit 049bfd186f
6 changed files with 160 additions and 23 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

67
gcc/fortran/openmp.c
View File

@ -5345,22 +5345,25 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
}
break;
case OMP_LIST_IS_DEVICE_PTR:
if (!n->sym->attr.dummy)
gfc_error ("Non-dummy object %qs in %s clause at %L",
n->sym->name, name, &n->where);
if (n->sym->attr.allocatable
|| (n->sym->ts.type == BT_CLASS
&& CLASS_DATA (n->sym)->attr.allocatable))
gfc_error ("ALLOCATABLE object %qs in %s clause at %L",
n->sym->name, name, &n->where);
if (n->sym->attr.pointer
|| (n->sym->ts.type == BT_CLASS
&& CLASS_DATA (n->sym)->attr.pointer))
gfc_error ("POINTER object %qs in %s clause at %L",
n->sym->name, name, &n->where);
if (n->sym->attr.value)
gfc_error ("VALUE object %qs in %s clause at %L",
n->sym->name, name, &n->where);
for (n = omp_clauses->lists[list]; n != NULL; n = n->next)
{
if (!n->sym->attr.dummy)
gfc_error ("Non-dummy object %qs in %s clause at %L",
n->sym->name, name, &n->where);
if (n->sym->attr.allocatable
|| (n->sym->ts.type == BT_CLASS
&& CLASS_DATA (n->sym)->attr.allocatable))
gfc_error ("ALLOCATABLE object %qs in %s clause at %L",
n->sym->name, name, &n->where);
if (n->sym->attr.pointer
|| (n->sym->ts.type == BT_CLASS
&& CLASS_DATA (n->sym)->attr.pointer))
gfc_error ("POINTER object %qs in %s clause at %L",
n->sym->name, name, &n->where);
if (n->sym->attr.value)
gfc_error ("VALUE object %qs in %s clause at %L",
n->sym->name, name, &n->where);
}
break;
case OMP_LIST_USE_DEVICE_PTR:
case OMP_LIST_USE_DEVICE_ADDR:
@ -5657,6 +5660,38 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
break;
}
}
/* OpenMP 5.1: use_device_ptr acts like use_device_addr, except for
type(c_ptr). */
if (omp_clauses->lists[OMP_LIST_USE_DEVICE_PTR])
{
gfc_omp_namelist *n_prev, *n_next, *n_addr;
n_addr = omp_clauses->lists[OMP_LIST_USE_DEVICE_ADDR];
for (; n_addr && n_addr->next; n_addr = n_addr->next)
;
n_prev = NULL;
n = omp_clauses->lists[OMP_LIST_USE_DEVICE_PTR];
while (n)
{
n_next = n->next;
if (n->sym->ts.type != BT_DERIVED
|| n->sym->ts.u.derived->ts.f90_type != BT_VOID)
{
n->next = NULL;
if (n_addr)
n_addr->next = n;
else
omp_clauses->lists[OMP_LIST_USE_DEVICE_ADDR] = n;
n_addr = n;
if (n_prev)
n_prev->next = n_next;
else
omp_clauses->lists[OMP_LIST_USE_DEVICE_PTR] = n_next;
}
else
n_prev = n;
n = n_next;
}
}
if (omp_clauses->safelen_expr)
resolve_positive_int_expr (omp_clauses->safelen_expr, "SAFELEN");
if (omp_clauses->simdlen_expr)

6
gcc/omp-low.c
View File

@ -12520,7 +12520,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
|| omp_is_allocatable_or_ptr (ovar))
{
type = TREE_TYPE (type);
if (TREE_CODE (type) != ARRAY_TYPE
if (POINTER_TYPE_P (type)
&& TREE_CODE (type) != ARRAY_TYPE
&& ((OMP_CLAUSE_CODE (c) != OMP_CLAUSE_USE_DEVICE_ADDR
&& !omp_is_allocatable_or_ptr (ovar))
|| (omp_is_reference (ovar)
@ -12784,7 +12785,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
if (omp_is_reference (var))
{
type = TREE_TYPE (type);
if (TREE_CODE (type) != ARRAY_TYPE
if (POINTER_TYPE_P (type)
&& TREE_CODE (type) != ARRAY_TYPE
&& (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_USE_DEVICE_ADDR
|| (omp_is_reference (var)
&& omp_is_allocatable_or_ptr (var))))

21
gcc/testsuite/gfortran.dg/gomp/is_device_ptr-2.f90 Normal file
View File

@ -0,0 +1,21 @@
! PR fortran/98476
subroutine abc(cc)
integer, target :: cc, dd
cc = 131
dd = 484
!$omp target enter data map(to: cc, dd)
!$omp target data use_device_addr(cc) use_device_ptr(dd)
!$omp target is_device_ptr(cc, dd) ! { dg-error "Non-dummy object 'cc' in IS_DEVICE_PTR clause at" }
if (cc /= 131 .or. dd /= 484) stop 1
cc = 44
dd = 45
!$omp end target
!$omp end target data
!$omp target exit data map(from:cc, dd)
if (cc /= 44 .or. dd /= 45) stop 5
end

10
gcc/testsuite/gfortran.dg/gomp/map-3.f90
View File

@ -1,10 +1,10 @@
! { dg-additional-options "-fdump-tree-original" }
subroutine bar
integer, target :: x
integer, target :: x, x2
integer, allocatable, target :: y(:,:), z(:,:)
x = 7
!$omp target enter data map(to:x)
!$omp target enter data map(to:x, x2)
x = 8
!$omp target data map(always, to: x)
@ -15,7 +15,7 @@ call foo(x)
call foo2(x)
!$omp end target data
!$omp target data use_device_addr(x)
!$omp target data use_device_addr(x2)
call foo2(x)
!$omp end target data
!$omp target exit data map(release:x)
@ -31,8 +31,8 @@ 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 data use_device_addr\\(x2\\)" 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" } }
! { 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_addr\\(z\\)" 1 "original" } }

25
gcc/testsuite/gfortran.dg/gomp/use_device_ptr-1.f90 Normal file
View File

@ -0,0 +1,25 @@
! { dg-do compile }
! { dg-additional-options "-fdump-tree-original" }
! PR fortran/98476
use iso_c_binding, only: c_ptr
implicit none (external, type)
interface
subroutine bar(x)
import
type(c_ptr), value :: x
end
end interface
type(c_ptr) :: x
!$omp target data map(alloc: x)
!$omp target data use_device_ptr(x)
call bar(x)
!$omp end target data
!$omp end target data
end
! { dg-final { scan-tree-dump-times "pragma omp target data use_device_ptr\\(x\\)" 1 "original" } }

54
libgomp/testsuite/libgomp.fortran/is_device_ptr-1.f90 Normal file
View File

@ -0,0 +1,54 @@
! { dg-additional-options "-fdump-tree-original" }
! PR fortran/98476
program abc
implicit none
integer a, b
a = 83
b = 73
call test(a, b)
contains
subroutine test(aa, bb)
use iso_c_binding, only: c_ptr, c_loc, c_f_pointer
integer :: aa, bb
integer, target :: cc, dd
type(c_ptr) :: pcc, pdd
cc = 131
dd = 484
!$omp target enter data map(to: aa, bb, cc, dd)
!$omp target data use_device_ptr(aa, cc) use_device_addr(bb, dd)
pcc = c_loc(cc)
pdd = c_loc(dd)
! TODO: has_device_addr(cc, dd)
!$omp target is_device_ptr(aa, bb)
if (aa /= 83 .or. bb /= 73) stop 1
aa = 42
bb = 43
block
integer, pointer :: c2, d2
call c_f_pointer(pcc, c2)
call c_f_pointer(pdd, d2)
if (c2 /= 131 .or. d2 /= 484) stop 2
c2 = 44
d2 = 45
end block
!$omp end target
!$omp end target data
!$omp target exit data map(from:aa, bb, cc, dd)
if (aa /= 42 .or. bb /= 43) stop 3
if (cc /= 44 .or. dd /= 45) stop 5
endsubroutine
end program
! { dg-final { scan-tree-dump-times "omp target data .*use_device_addr\\(aa\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "omp target data .*use_device_addr\\(bb\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "omp target data .*use_device_addr\\(cc\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "omp target data .*use_device_addr\\(dd\\)" 1 "original" } }