+2018-06-29 Cesar Philippidis <cesar@codesourcery.com>
+ James Norris <jnorris@codesourcery.com>
+
+ * gimplify.cc (enum gimplify_omp_var_data): Add GOVD_DEVICETPR.
+ (omp_notice_variable): Add GOVD_DEVICEPTR attribute when appropriate.
+ (gimplify_scan_omp_clauses): Add GOVD_DEVICEPTR attribute when
+ appropriate.
+ (gimplify_adjust_omp_clauses_1): Set GOMP_MAP_FORCE_DEVICEPTR for
+ implicit deviceptr mappings.
+
2020-04-19 Chung-Lin Tang <cltang@codesourcery.com>
PR other/76739
+2018-06-29 Cesar Philippidis <cesar@codesourcery.com>
+ James Norris <jnorris@codesourcery.com>
+
+ * openmp.cc (resolve_positive_int_expr): Promote the warning to an
+ error.
+
2020-04-19 Chung-Lin Tang <cltang@codesourcery.com>
PR other/76739
if (expr->expr_type == EXPR_CONSTANT
&& expr->ts.type == BT_INTEGER
&& mpz_sgn (expr->value.integer) <= 0)
- gfc_warning (0, "INTEGER expression of %s clause at %L must be positive",
- clause, &expr->where);
+ gfc_error ("INTEGER expression of %s clause at %L must be positive",
+ clause, &expr->where);
}
static void
return;
}
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FORCE_DEVICEPTR)
+ return;
+
tree c2 = NULL_TREE, c3 = NULL_TREE, c4 = NULL_TREE;
tree present = gfc_omp_check_optional_argument (decl, true);
if (POINTER_TYPE_P (TREE_TYPE (decl)))
OMP_CLAUSE_SIZE (node3) = size_int (0);
goto finalize_map_clause;
}
+ else if (POINTER_TYPE_P (TREE_TYPE (decl))
+ && n->u.map_op == OMP_MAP_FORCE_DEVICEPTR)
+ {
+ OMP_CLAUSE_DECL (node) = decl;
+ goto finalize_map_clause;
+ }
else if (POINTER_TYPE_P (TREE_TYPE (decl))
&& (gfc_omp_privatize_by_reference (decl)
|| GFC_DECL_GET_SCALAR_POINTER (decl)
/* Flag for GOVD_FIRSTPRIVATE: OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT. */
GOVD_FIRSTPRIVATE_IMPLICIT = 0x8000000,
+ /* Flag for OpenACC deviceptrs. */
+ GOVD_DEVICEPTR = (1<<24),
+
GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE
| GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR
| GOVD_LOCAL)
error ("variable %qE declared in enclosing "
"%<host_data%> region", DECL_NAME (decl));
nflags |= GOVD_MAP;
+ nflags |= (n2->value & GOVD_DEVICEPTR);
if (octx->region_type == ORT_ACC_DATA
&& (n2->value & GOVD_MAP_0LEN_ARRAY))
nflags |= GOVD_MAP_0LEN_ARRAY;
if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TO
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TOFROM)
flags |= GOVD_MAP_ALWAYS_TO;
+ else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FORCE_DEVICEPTR)
+ flags |= GOVD_DEVICEPTR;
if ((code == OMP_TARGET
|| code == OMP_TARGET_DATA
| GOVD_MAP_FORCE
| GOVD_MAP_FORCE_PRESENT
| GOVD_MAP_ALLOC_ONLY
- | GOVD_MAP_FROM_ONLY))
+ | GOVD_MAP_FROM_ONLY
+ | GOVD_DEVICEPTR))
{
case 0:
kind = GOMP_MAP_TOFROM;
case GOVD_MAP_FORCE_PRESENT:
kind = GOMP_MAP_FORCE_PRESENT;
break;
+ case GOVD_DEVICEPTR:
+ kind = GOMP_MAP_FORCE_DEVICEPTR;
+ break;
default:
gcc_unreachable ();
}
+2018-06-29 Cesar Philippidis <cesar@codesourcery.com>
+ James Norris <jnorris@codesourcery.com>
+
+ * c-c++-common/goacc/deviceptr-4.c: Update.
+ * gfortran.dg/goacc/loop-2-kernels-tile.f95: Update.
+ * gfortran.dg/goacc/loop-2-parallel-tile.f95: Update.
+ * gfortran.dg/goacc/sie.f95: Update.
+ * gfortran.dg/goacc/tile-1.f90: Update.
+ * gfortran.dg/gomp/pr77516.f90: Update.
+
2020-04-19 Chung-Lin Tang <cltang@codesourcery.com>
PR other/76739
a[0] += 1.0;
}
-/* { dg-final { scan-tree-dump-times "#pragma omp target oacc_parallel.*map\\(tofrom:a" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "#pragma omp target oacc_parallel.*map\\(force_deviceptr:a" 1 "gimple" } } */
DO j = 1,10
ENDDO
ENDDO
- !$acc loop tile(-1) ! { dg-warning "must be positive" }
+ !$acc loop tile(-1) ! { dg-error "must be positive" }
do i = 1,10
enddo
!$acc loop tile(i) ! { dg-error "constant expression" }
DO j = 1,10
ENDDO
ENDDO
- !$acc kernels loop tile(-1) ! { dg-warning "must be positive" }
+ !$acc kernels loop tile(-1) ! { dg-error "must be positive" }
do i = 1,10
enddo
!$acc kernels loop tile(i) ! { dg-error "constant expression" }
DO j = 1,10
ENDDO
ENDDO
- !$acc loop tile(-1) ! { dg-warning "must be positive" }
+ !$acc loop tile(-1) ! { dg-error "must be positive" }
do i = 1,10
enddo
!$acc loop tile(i) ! { dg-error "constant expression" }
DO j = 1,10
ENDDO
ENDDO
- !$acc parallel loop tile(-1) ! { dg-warning "must be positive" }
+ !$acc parallel loop tile(-1) ! { dg-error "must be positive" }
do i = 1,10
enddo
!$acc parallel loop tile(i) ! { dg-error "constant expression" }
!$acc parallel num_gangs(i+1)
!$acc end parallel
- !$acc parallel num_gangs(-1) ! { dg-warning "must be positive" }
+ !$acc parallel num_gangs(-1) ! { dg-error "must be positive" }
!$acc end parallel
- !$acc parallel num_gangs(0) ! { dg-warning "must be positive" }
+ !$acc parallel num_gangs(0) ! { dg-error "must be positive" }
!$acc end parallel
!$acc parallel num_gangs() ! { dg-error "Invalid character in name" }
!$acc kernels num_gangs(i+1)
!$acc end kernels
- !$acc kernels num_gangs(-1) ! { dg-warning "must be positive" }
+ !$acc kernels num_gangs(-1) ! { dg-error "must be positive" }
!$acc end kernels
- !$acc kernels num_gangs(0) ! { dg-warning "must be positive" }
+ !$acc kernels num_gangs(0) ! { dg-error "must be positive" }
!$acc end kernels
!$acc kernels num_gangs() ! { dg-error "Invalid character in name" }
!$acc parallel num_workers(i+1)
!$acc end parallel
- !$acc parallel num_workers(-1) ! { dg-warning "must be positive" }
+ !$acc parallel num_workers(-1) ! { dg-error "must be positive" }
!$acc end parallel
- !$acc parallel num_workers(0) ! { dg-warning "must be positive" }
+ !$acc parallel num_workers(0) ! { dg-error "must be positive" }
!$acc end parallel
!$acc parallel num_workers() ! { dg-error "Invalid expression after 'num_workers\\('" }
!$acc kernels num_workers(i+1)
!$acc end kernels
- !$acc kernels num_workers(-1) ! { dg-warning "must be positive" }
+ !$acc kernels num_workers(-1) ! { dg-error "must be positive" }
!$acc end kernels
- !$acc kernels num_workers(0) ! { dg-warning "must be positive" }
+ !$acc kernels num_workers(0) ! { dg-error "must be positive" }
!$acc end kernels
!$acc kernels num_workers() ! { dg-error "Invalid expression after 'num_workers\\('" }
!$acc parallel vector_length(i+1)
!$acc end parallel
- !$acc parallel vector_length(-1) ! { dg-warning "must be positive" }
+ !$acc parallel vector_length(-1) ! { dg-error "must be positive" }
!$acc end parallel
- !$acc parallel vector_length(0) ! { dg-warning "must be positive" }
+ !$acc parallel vector_length(0) ! { dg-error "must be positive" }
!$acc end parallel
!$acc parallel vector_length() ! { dg-error "Invalid expression after 'vector_length\\('" }
!$acc kernels vector_length(i+1)
!$acc end kernels
- !$acc kernels vector_length(-1) ! { dg-warning "must be positive" }
+ !$acc kernels vector_length(-1) ! { dg-error "must be positive" }
!$acc end kernels
- !$acc kernels vector_length(0) ! { dg-warning "must be positive" }
+ !$acc kernels vector_length(0) ! { dg-error "must be positive" }
!$acc end kernels
!$acc kernels vector_length() ! { dg-error "Invalid expression after 'vector_length\\('" }
!$acc loop gang(i+1)
do i = 1,10
enddo
- !$acc loop gang(-1) ! { dg-warning "must be positive" }
+ !$acc loop gang(-1) ! { dg-error "must be positive" }
do i = 1,10
enddo
- !$acc loop gang(0) ! { dg-warning "must be positive" }
+ !$acc loop gang(0) ! { dg-error "must be positive" }
do i = 1,10
enddo
!$acc loop gang() ! { dg-error "Invalid character in name" }
!$acc loop worker(i+1)
do i = 1,10
enddo
- !$acc loop worker(-1) ! { dg-warning "must be positive" }
+ !$acc loop worker(-1) ! { dg-error "must be positive" }
do i = 1,10
enddo
- !$acc loop worker(0) ! { dg-warning "must be positive" }
+ !$acc loop worker(0) ! { dg-error "must be positive" }
do i = 1,10
enddo
!$acc loop worker() ! { dg-error "Invalid character in name" }
!$acc loop vector(i+1)
do i = 1,10
enddo
- !$acc loop vector(-1) ! { dg-warning "must be positive" }
+ !$acc loop vector(-1) ! { dg-error "must be positive" }
do i = 1,10
enddo
- !$acc loop vector(0) ! { dg-warning "must be positive" }
+ !$acc loop vector(0) ! { dg-error "must be positive" }
do i = 1,10
enddo
!$acc loop vector() ! { dg-error "Invalid character in name" }
do i = 1, n
end do
- !$acc parallel loop tile(-3) ! { dg-warning "must be positive" }
+ !$acc parallel loop tile(-3) ! { dg-error "must be positive" }
do i = 1, n
end do
- !$acc parallel loop tile(10, -3) ! { dg-warning "must be positive" }
+ !$acc parallel loop tile(10, -3) ! { dg-error "must be positive" }
do i = 1, n
do j = 1, n
end do
end do
- !$acc parallel loop tile(-100, 10, 5) ! { dg-warning "must be positive" }
+ !$acc parallel loop tile(-100, 10, 5) ! { dg-error "must be positive" }
do i = 1, n
do j = 1, n
do k = 1, n
end do
end do
- !$acc loop tile(-2) ! { dg-warning "must be positive" }
+ !$acc loop tile(-2) ! { dg-error "must be positive" }
do i = 1, n
end do
end do
end do
- !$acc loop tile(-2) ! { dg-warning "must be positive" }
+ !$acc loop tile(-2) ! { dg-error "must be positive" }
do i = 1, n
end do
do i = 1, n
end do
- !$acc kernels loop tile(-3) ! { dg-warning "must be positive" }
+ !$acc kernels loop tile(-3) ! { dg-error "must be positive" }
do i = 1, n
end do
- !$acc kernels loop tile(10, -3) ! { dg-warning "must be positive" }
+ !$acc kernels loop tile(10, -3) ! { dg-error "must be positive" }
do i = 1, n
do j = 1, n
end do
end do
- !$acc kernels loop tile(-100, 10, 5) ! { dg-warning "must be positive" }
+ !$acc kernels loop tile(-100, 10, 5) ! { dg-error "must be positive" }
do i = 1, n
do j = 1, n
do k = 1, n
program pr77516
integer :: i, x
x = 0
-!$omp simd safelen(0) reduction(+:x) ! { dg-warning "must be positive" }
+!$omp simd safelen(0) reduction(+:x) ! { dg-error "must be positive" }
do i = 1, 8
x = x + 1
end do
+2018-06-29 Cesar Philippidis <cesar@codesourcery.com>
+ James Norris <jnorris@codesourcery.com>
+
+ * oacc-parallel.c (GOACC_parallel_keyed): Handle Fortran deviceptr
+ clause.
+ (GOACC_data_start): Likewise.
+ * testsuite/libgomp.oacc-fortran/deviceptr-1.f90: New test.
+
2019-02-12 Julian Brown <julian@codesourcery.com>
* oacc-cuda.c (acc_set_cuda_stream): Return 0 on error/invalid
if (profiling_p)
goacc_profiling_dispatch (&prof_info, &enter_data_event_info, &api_info);
+ handle_ftn_pointers (mapnum, hostaddrs, sizes, kinds);
+
/* Host fallback or 'do nothing'. */
if ((acc_dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
|| (flags & GOACC_FLAG_HOST_FALLBACK))
--- /dev/null
+! { dg-do run }
+
+! Test the deviceptr clause with various directives
+! and in combination with other directives where
+! the deviceptr variable is implied.
+
+subroutine subr1 (a, b)
+ implicit none
+ integer, parameter :: N = 8
+ integer :: a(N)
+ integer :: b(N)
+ integer :: i = 0
+
+ !$acc data deviceptr (a)
+
+ !$acc parallel copy (b)
+ do i = 1, N
+ a(i) = i * 2
+ b(i) = a(i)
+ end do
+ !$acc end parallel
+
+ !$acc end data
+
+end subroutine
+
+subroutine subr2 (a, b)
+ implicit none
+ integer, parameter :: N = 8
+ integer :: a(N)
+ !$acc declare deviceptr (a)
+ integer :: b(N)
+ integer :: i = 0
+
+ !$acc parallel copy (b)
+ do i = 1, N
+ a(i) = i * 4
+ b(i) = a(i)
+ end do
+ !$acc end parallel
+
+end subroutine
+
+subroutine subr3 (a, b)
+ implicit none
+ integer, parameter :: N = 8
+ integer :: a(N)
+ !$acc declare deviceptr (a)
+ integer :: b(N)
+ integer :: i = 0
+
+ !$acc kernels copy (b)
+ do i = 1, N
+ a(i) = i * 8
+ b(i) = a(i)
+ end do
+ !$acc end kernels
+
+end subroutine
+
+subroutine subr4 (a, b)
+ implicit none
+ integer, parameter :: N = 8
+ integer :: a(N)
+ integer :: b(N)
+ integer :: i = 0
+
+ !$acc parallel deviceptr (a) copy (b)
+ do i = 1, N
+ a(i) = i * 16
+ b(i) = a(i)
+ end do
+ !$acc end parallel
+
+end subroutine
+
+subroutine subr5 (a, b)
+ implicit none
+ integer, parameter :: N = 8
+ integer :: a(N)
+ integer :: b(N)
+ integer :: i = 0
+
+ !$acc kernels deviceptr (a) copy (b)
+ do i = 1, N
+ a(i) = i * 32
+ b(i) = a(i)
+ end do
+ !$acc end kernels
+
+end subroutine
+
+subroutine subr6 (a, b)
+ implicit none
+ integer, parameter :: N = 8
+ integer :: a(N)
+ integer :: b(N)
+ integer :: i = 0
+
+ !$acc parallel deviceptr (a) copy (b)
+ do i = 1, N
+ b(i) = i
+ end do
+ !$acc end parallel
+
+end subroutine
+
+subroutine subr7 (a, b)
+ implicit none
+ integer, parameter :: N = 8
+ integer :: a(N)
+ integer :: b(N)
+ integer :: i = 0
+
+ !$acc data deviceptr (a)
+
+ !$acc parallel copy (b)
+ do i = 1, N
+ a(i) = i * 2
+ b(i) = a(i)
+ end do
+ !$acc end parallel
+
+ !$acc parallel copy (b)
+ do i = 1, N
+ a(i) = b(i) * 2
+ b(i) = a(i)
+ end do
+ !$acc end parallel
+
+ !$acc end data
+
+end subroutine
+
+program main
+ use iso_c_binding, only: c_ptr, c_f_pointer
+ implicit none
+ type (c_ptr) :: cp
+ integer, parameter :: N = 8
+ integer, pointer :: fp(:)
+ integer :: i = 0
+ integer :: b(N)
+
+ interface
+ function acc_malloc (s) bind (C)
+ use iso_c_binding, only: c_ptr, c_size_t
+ integer (c_size_t), value :: s
+ type (c_ptr) :: acc_malloc
+ end function
+ end interface
+
+ cp = acc_malloc (N * sizeof (fp(N)))
+ call c_f_pointer (cp, fp, [N])
+
+ call subr1 (fp, b)
+
+ do i = 1, N
+ if (b(i) .ne. i * 2) call abort
+ end do
+
+ call subr2 (fp, b)
+
+ do i = 1, N
+ if (b(i) .ne. i * 4) call abort
+ end do
+
+ call subr3 (fp, b)
+
+ do i = 1, N
+ if (b(i) .ne. i * 8) call abort
+ end do
+
+ call subr4 (fp, b)
+
+ do i = 1, N
+ if (b(i) .ne. i * 16) call abort
+ end do
+
+ call subr5 (fp, b)
+
+ do i = 1, N
+ if (b(i) .ne. i * 32) call abort
+ end do
+
+ call subr6 (fp, b)
+
+ do i = 1, N
+ if (b(i) .ne. i) call abort
+ end do
+
+ call subr7 (fp, b)
+
+ do i = 1, N
+ if (b(i) .ne. i * 4) call abort
+ end do
+
+end program main
projects / thirdparty / gcc.git / commitdiff
