mirrors/gcc
1
0
Fork 0
mirror of https://gcc.gnu.org/git/gcc.git synced 2024-11-24 03:14:08 +08:00
Code Issues Actions Packages Projects Releases Wiki Activity

OpenMP/Fortran: Fix has_device_addr clause splitting [PR108558]

Browse Source 
gcc/fortran/ChangeLog:

	PR fortran/108558
	* trans-openmp.cc (gfc_split_omp_clauses): Handle has_device_addr.

libgomp/ChangeLog:

	PR fortran/108558
	* testsuite/libgomp.fortran/has_device_addr.f90: New test.

(cherry picked from commit 2325c8920b)
This commit is contained in:
Tobias Burnus 2023-01-27 11:32:19 +01:00
parent 355fe4fa4c
commit 591ec4820a
2 changed files with 61 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
gcc/fortran/trans-openmp.cc
View File

@ -6070,6 +6070,8 @@ gfc_split_omp_clauses (gfc_code *code,
= code->ext.omp_clauses->lists[OMP_LIST_MAP];
clausesa[GFC_OMP_SPLIT_TARGET].lists[OMP_LIST_IS_DEVICE_PTR]
= code->ext.omp_clauses->lists[OMP_LIST_IS_DEVICE_PTR];
clausesa[GFC_OMP_SPLIT_TARGET].lists[OMP_LIST_HAS_DEVICE_ADDR]
= code->ext.omp_clauses->lists[OMP_LIST_HAS_DEVICE_ADDR];
clausesa[GFC_OMP_SPLIT_TARGET].device
= code->ext.omp_clauses->device;
clausesa[GFC_OMP_SPLIT_TARGET].thread_limit

59
libgomp/testsuite/libgomp.fortran/has_device_addr.f90 Normal file
View File

@ -0,0 +1,59 @@
! { dg-additional-options "-fdump-tree-original" }
!
! PR fortran/108558
!
! { dg-final { scan-tree-dump-times "#pragma omp target has_device_addr\\(x\\) has_device_addr\\(y\\)" 2 "original" } }
! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(tofrom:x\\) map\\(tofrom:y\\)" 2 "original" } }
! { dg-final { scan-tree-dump-times "#pragma omp target data use_device_addr\\(x\\) use_device_addr\\(y\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "#pragma omp target update from\\(y\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(tofrom:x\\) map\\(tofrom:y\\) use_device_addr\\(x\\) use_device_addr\\(y\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "#pragma omp teams" 2 "original" } }
! { dg-final { scan-tree-dump-times "#pragma omp distribute" 2 "original" } }
! { dg-final { scan-tree-dump-times "#pragma omp parallel" 2 "original" } }
! { dg-final { scan-tree-dump-times "#pragma omp for nowait" 2 "original" } }
module m
contains
subroutine vectorAdd(x, y, N)
implicit none
integer :: N
integer(4) :: x(N), y(N)
integer :: i
!$omp target teams distribute parallel do has_device_addr(x, y)
do i = 1, N
y(i) = x(i) + y(i)
end do
end subroutine vectorAdd
end module m
program main
use m
implicit none
integer, parameter :: N = 9876
integer(4) :: x(N), y(N)
integer :: i
x(:) = 1
y(:) = 2
!$omp target data map(x, y)
!$omp target data use_device_addr(x, y)
call vectorAdd(x, y, N)
!$omp end target data
!$omp target update from(y)
if (any (y /= 3)) error stop
!$omp end target data
x = 1
y = 2
!$omp target data map(x, y) use_device_addr(x, y)
!$omp target teams distribute parallel do has_device_addr(x, y)
do i = 1, N
y(i) = x(i) + y(i)
end do
!$omp end target data
if (any (y /= 3)) error stop
end program