+2018-06-29 Cesar Philippidis <cesar@codesourcery.com>
+ James Norris <jnorris@codesourcery.com>
+
+ * gimplify.c (enum gimplify_omp_var_data): Add GOVD_DEVICETPR.
+ (oacc_default_clause): Privatize fortran common blocks.
+ (omp_notice_variable): Add GOVD_DEVICEPTR attribute when appropriate.
+ Defer the expansion of DECL_VALUE_EXPR for common block decls.
+ (gimplify_scan_omp_clauses): Add GOVD_DEVICEPTR attribute when
+ appropriate.
+ (gimplify_adjust_omp_clauses_1): Set GOMP_MAP_FORCE_DEVICEPTR for
+ implicit deviceptr mappings.
+
2018-10-02 Thomas Schwinge <thomas@codesourcery.com>
Cesar Philippidis <cesar@codesourcery.com>
+2018-06-29 Cesar Philippidis <cesar@codesourcery.com>
+ James Norris <jnorris@codesourcery.com>
+
+ * openmp.c (gfc_match_omp_map_clause): Re-write handling of the
+ deviceptr clause. Add new common_blocks argument. Propagate it to
+ gfc_match_omp_variable_list.
+ (gfc_match_omp_clauses): Update calls to gfc_match_omp_map_clauses.
+ (resolve_positive_int_expr): Promote the warning to an error.
+ (check_array_not_assumed): Remove pointer check.
+ (resolve_oacc_nested_loops): Error on do concurrent loops.
+ * trans-openmp.c (gfc_omp_finish_clause): Don't create pointer data
+ mappings for deviceptr clauses.
+ (gfc_trans_omp_clauses): Likewise.
+
2018-10-02 Thomas Schwinge <thomas@codesourcery.com>
Cesar Philippidis <cesar@codesourcery.com>
static bool
gfc_match_omp_map_clause (gfc_omp_namelist **list, gfc_omp_map_op map_op,
- bool allow_derived = false)
+ bool common_blocks, bool allow_derived = false)
{
gfc_omp_namelist **head = NULL;
- if (gfc_match_omp_variable_list ("", list, false, NULL, &head, true,
+ if (gfc_match_omp_variable_list ("", list, common_blocks, NULL, &head, true,
allow_derived)
== MATCH_YES)
{
if ((mask & OMP_CLAUSE_ATTACH)
&& gfc_match ("attach ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_ATTACH, allow_derived))
+ OMP_MAP_ATTACH, false,
+ allow_derived))
continue;
break;
case 'c':
if ((mask & OMP_CLAUSE_COPY)
&& gfc_match ("copy ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_TOFROM, allow_derived))
+ OMP_MAP_TOFROM, openacc,
+ allow_derived))
continue;
if (mask & OMP_CLAUSE_COPYIN)
{
{
if (gfc_match ("copyin ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_TO, allow_derived))
+ OMP_MAP_TO, true,
+ allow_derived))
continue;
}
else if (gfc_match_omp_variable_list ("copyin (",
if ((mask & OMP_CLAUSE_COPYOUT)
&& gfc_match ("copyout ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_FROM, allow_derived))
+ OMP_MAP_FROM, true, allow_derived))
continue;
if ((mask & OMP_CLAUSE_COPYPRIVATE)
&& gfc_match_omp_variable_list ("copyprivate (",
if ((mask & OMP_CLAUSE_CREATE)
&& gfc_match ("create ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_ALLOC, allow_derived))
+ OMP_MAP_ALLOC, true, allow_derived))
continue;
break;
case 'd':
if ((mask & OMP_CLAUSE_DELETE)
&& gfc_match ("delete ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_RELEASE, allow_derived))
+ OMP_MAP_RELEASE, true,
+ allow_derived))
continue;
if ((mask & OMP_CLAUSE_DEPEND)
&& gfc_match ("depend ( ") == MATCH_YES)
if ((mask & OMP_CLAUSE_DETACH)
&& gfc_match ("detach ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_DETACH, allow_derived))
+ OMP_MAP_DETACH, false,
+ allow_derived))
continue;
if ((mask & OMP_CLAUSE_DEVICE)
&& !openacc
&& openacc
&& gfc_match ("device ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_FORCE_TO, allow_derived))
+ OMP_MAP_FORCE_TO, false,
+ allow_derived))
continue;
if ((mask & OMP_CLAUSE_DEVICEPTR)
&& gfc_match ("deviceptr ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_FORCE_DEVICEPTR,
+ OMP_MAP_FORCE_DEVICEPTR, false,
allow_derived))
continue;
if ((mask & OMP_CLAUSE_DEVICE_RESIDENT)
if ((mask & OMP_CLAUSE_HOST_SELF)
&& gfc_match ("host ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_FORCE_FROM, allow_derived))
+ OMP_MAP_FORCE_FROM, true,
+ allow_derived))
continue;
break;
case 'i':
if ((mask & OMP_CLAUSE_COPY)
&& gfc_match ("pcopy ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_TOFROM, allow_derived))
+ OMP_MAP_TOFROM, true,
+ allow_derived))
continue;
if ((mask & OMP_CLAUSE_COPYIN)
&& gfc_match ("pcopyin ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_TO, allow_derived))
+ OMP_MAP_TO, true, allow_derived))
continue;
if ((mask & OMP_CLAUSE_COPYOUT)
&& gfc_match ("pcopyout ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_FROM, allow_derived))
+ OMP_MAP_FROM, true, allow_derived))
continue;
if ((mask & OMP_CLAUSE_CREATE)
&& gfc_match ("pcreate ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_ALLOC, allow_derived))
+ OMP_MAP_ALLOC, true, allow_derived))
continue;
if ((mask & OMP_CLAUSE_PRESENT)
&& gfc_match ("present ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_FORCE_PRESENT,
+ OMP_MAP_FORCE_PRESENT, false,
allow_derived))
continue;
if ((mask & OMP_CLAUSE_COPY)
&& gfc_match ("present_or_copy ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_TOFROM, allow_derived))
+ OMP_MAP_TOFROM, true,
+ allow_derived))
continue;
if ((mask & OMP_CLAUSE_COPYIN)
&& gfc_match ("present_or_copyin ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_TO, allow_derived))
+ OMP_MAP_TO, true, allow_derived))
continue;
if ((mask & OMP_CLAUSE_COPYOUT)
&& gfc_match ("present_or_copyout ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_FROM, allow_derived))
+ OMP_MAP_FROM, true, allow_derived))
continue;
if ((mask & OMP_CLAUSE_CREATE)
&& gfc_match ("present_or_create ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_ALLOC, allow_derived))
+ OMP_MAP_ALLOC, true, allow_derived))
continue;
if ((mask & OMP_CLAUSE_PRIORITY)
&& c->priority == NULL
if ((mask & OMP_CLAUSE_HOST_SELF)
&& gfc_match ("self ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_FORCE_FROM, allow_derived))
+ OMP_MAP_FORCE_FROM, true,
+ allow_derived))
continue;
if ((mask & OMP_CLAUSE_SEQ)
&& !c->seq
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
if (sym->as && sym->as->type == AS_ASSUMED_RANK)
gfc_error ("Assumed rank array %qs in %s clause at %L",
sym->name, name, &loc);
- if (sym->as && sym->as->type == AS_DEFERRED && sym->attr.pointer
- && !sym->attr.contiguous)
- gfc_error ("Noncontiguous deferred shape array %qs in %s clause at %L",
- sym->name, name, &loc);
}
static void
}
tree c2 = NULL_TREE, c3 = NULL_TREE, c4 = NULL_TREE;
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FORCE_DEVICEPTR)
+ return;
if (POINTER_TYPE_P (TREE_TYPE (decl)))
{
if (!gfc_omp_privatize_by_reference (decl)
field = TREE_OPERAND (field, 1);
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 (field)
|| GFC_DECL_GET_SCALAR_ALLOCATABLE (field)
ptr2 = fold_convert (sizetype, ptr2);
OMP_CLAUSE_SIZE (node3)
= fold_build2 (MINUS_EXPR, sizetype, ptr, ptr2);
+ finalize_map_clause:;
}
else
gcc_unreachable ();
fields. */
GOVD_MAP_HAS_ATTACHMENTS = 8388608,
+ /* 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)
{
const char *rkind;
bool on_device = false;
+ bool is_private = false;
bool declared = is_oacc_declared (decl);
tree type = TREE_TYPE (decl);
if (lang_hooks.decls.omp_privatize_by_reference (decl))
type = TREE_TYPE (type);
+ if (RECORD_OR_UNION_TYPE_P (type))
+ is_private = lang_hooks.decls.omp_disregard_value_expr (decl, false);
+
if ((ctx->region_type & (ORT_ACC_PARALLEL | ORT_ACC_KERNELS)) != 0
&& is_global_var (decl)
- && device_resident_p (decl))
+ && device_resident_p (decl)
+ && !is_private)
{
on_device = true;
flags |= GOVD_MAP_TO_ONLY;
case ORT_ACC_KERNELS:
rkind = "kernels";
- if (AGGREGATE_TYPE_P (type))
+ if (is_private)
+ flags |= GOVD_MAP;
+ else if (AGGREGATE_TYPE_P (type))
{
/* Aggregates default to 'present_or_copy', or 'present'. */
if (ctx->default_kind != OMP_CLAUSE_DEFAULT_PRESENT)
case ORT_ACC_PARALLEL:
rkind = "parallel";
- if (on_device || declared)
+ if (is_private)
+ flags |= GOVD_FIRSTPRIVATE;
+ else if (on_device || declared)
flags |= GOVD_MAP;
else if (AGGREGATE_TYPE_P (type))
{
{
tree value = get_base_address (DECL_VALUE_EXPR (decl));
- if (value && DECL_P (value) && DECL_THREAD_LOCAL_P (value))
+ if (!(ctx->region_type & ORT_ACC)
+ && value && DECL_P (value) && DECL_THREAD_LOCAL_P (value))
return omp_notice_threadprivate_variable (ctx, decl, value);
}
n = splay_tree_lookup (ctx->variables, (splay_tree_key)decl);
if ((ctx->region_type & ORT_TARGET) != 0)
{
- ret = lang_hooks.decls.omp_disregard_value_expr (decl, true);
+ shared = !(ctx->region_type & ORT_ACC);
+ ret = lang_hooks.decls.omp_disregard_value_expr (decl, shared);
if (n == NULL)
{
unsigned nflags = flags;
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;
}
shared = ((flags | n->value) & GOVD_SHARED) != 0;
+ if (ctx->region_type & ORT_ACC)
+ shared = false;
ret = lang_hooks.decls.omp_disregard_value_expr (decl, shared);
/* If nothing changed, there's nothing left to do. */
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;
goto do_add;
case OMP_CLAUSE_DEPEND:
| 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/common-block-1.f90: New test.
+ * gfortran.dg/goacc/common-block-2.f90: New test.
+ * gfortran.dg/goacc/loop-2.f95: Update.
+ * gfortran.dg/goacc/loop-3-2.f95: Update.
+ * gfortran.dg/goacc/loop-3.f95: Update.
+ * gfortran.dg/goacc/loop-5.f95: Update.
+ * gfortran.dg/goacc/pr72715.f90: New test.
+ * gfortran.dg/goacc/sie.f95: Update.
+ * gfortran.dg/goacc/tile-1.f90: Update.
+ * gfortran.dg/gomp/pr77516.f90: Update.
+
2018-10-02 Thomas Schwinge <thomas@codesourcery.com>
Cesar Philippidis <cesar@codesourcery.com>
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" } } */
--- /dev/null
+! Test data clauses involving common blocks and common block data.
+! Specifically, validate early matching errors.
+
+subroutine subtest
+ implicit none
+ integer, parameter :: n = 10
+ integer a(n), b(n), c, d(n), e
+ real*4 x(n), y(n), z, w(n), v
+ common /blockA/ a, c, x
+ common /blockB/ b, y, z
+ !$acc declare link(/blockA/, /blockB/, e, v)
+end subroutine subtest
+
+program test
+ implicit none
+ integer, parameter :: n = 10
+ integer a(n), b(n), c, d(n), e
+ real*4 x(n), y(n), z, w(n), v
+ common /blockA/ a, c, x
+ common /blockB/ b, y, z
+ !$acc declare link(/blockA/, /blockB/, e, v)
+
+ !$acc data copy(/blockA/, /blockB/, e, v)
+ !$acc end data
+
+ !$acc data copyin(/blockA/, /blockB/, e, v)
+ !$acc end data
+
+ !$acc data copyout(/blockA/, /blockB/, e, v)
+ !$acc end data
+
+ !$acc data create(/blockA/, /blockB/, e, v)
+ !$acc end data
+
+ !$acc data copyout(/blockA/, /blockB/, e, v)
+ !$acc end data
+
+ !$acc data pcopy(/blockA/, /blockB/, e, v)
+ !$acc end data
+
+ !$acc data pcopyin(/blockA/, /blockB/, e, v)
+ !$acc end data
+
+ !$acc data pcopyout(/blockA/, /blockB/, e, v)
+ !$acc end data
+
+ !$acc data pcreate(/blockA/, /blockB/, e, v)
+ !$acc end data
+
+ !$acc data pcopyout(/blockA/, /blockB/, e, v)
+ !$acc end data
+
+ !$acc data present(/blockA/, /blockB/, e, v) ! { dg-error "" }
+ !$acc end data ! { dg-error "Unexpected ..ACC END DATA statement" }
+
+ !$acc parallel private(/blockA/, /blockB/, e, v)
+ !$acc end parallel
+
+ !$acc parallel firstprivate(/blockA/, /blockB/, e, v)
+ !$acc end parallel
+
+ !$acc exit data delete(/blockA/, /blockB/, e, v)
+
+ !$acc data deviceptr(/blockA/, /blockB/, e, v) ! { dg-error "Syntax error" }
+ !$acc end data ! { dg-error "Unexpected ..ACC END DATA statement" }
+
+ !$acc data deviceptr(/blockA/, /blockB/, e, v) ! { dg-error "Syntax error" }
+ !$acc end data ! { dg-error "Unexpected ..ACC END DATA statement" }
+end program test
--- /dev/null
+! Test data clauses involving common blocks and common block data.
+! Specifically, resolver errors such as duplicate data clauses.
+
+program test
+ implicit none
+ integer, parameter :: n = 10
+ integer a(n), b(n), c, d(n), e
+ real*4 x(n), y(n), z, w(n), v
+ common /blockA/ a, c, x
+ common /blockB/ b, y, z
+
+ !$acc data copy(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+ !$acc end data
+
+ !$acc data copyin(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+ !$acc end data
+
+ !$acc data copyout(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+ !$acc end data
+
+ !$acc data create(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+ !$acc end data
+
+ !$acc data copyout(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+ !$acc end data
+
+ !$acc data pcopy(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+ !$acc end data
+
+ !$acc data pcopyin(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+ !$acc end data
+
+ !$acc data pcopyout(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+ !$acc end data
+
+ !$acc data pcreate(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+ !$acc end data
+
+ !$acc data pcopyout(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+ !$acc end data
+
+ !$acc parallel private(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+ !$acc end parallel
+
+ !$acc parallel firstprivate(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+ !$acc end parallel
+
+ !$acc exit data delete(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+end program test
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" }
DO j = 1,10
ENDDO
ENDDO
- !$acc loop tile(-1) ! { dg-warning "must be positive" }
- do i = 1,10
- enddo
!$acc loop vector tile(*)
DO i = 1,10
ENDDO
DO j = 1,10
ENDDO
ENDDO
- !$acc loop tile(-1) ! { dg-warning "must be positive" }
- do i = 1,10
- enddo
!$acc loop vector tile(*)
DO i = 1,10
ENDDO
DO j = 1,10
ENDDO
ENDDO
- !$acc kernels loop tile(-1) ! { dg-warning "must be positive" }
- do i = 1,10
- enddo
!$acc kernels loop vector tile(*)
DO i = 1,10
ENDDO
DO j = 1,10
ENDDO
ENDDO
- !$acc parallel loop tile(-1) ! { dg-warning "must be positive" }
- do i = 1,10
- enddo
!$acc parallel loop vector tile(*)
DO i = 1,10
ENDDO
!$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 character in name" }
!$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 character in name" }
!$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 character in name" }
!$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 character in name" }
!$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/common-block-1.f90: New test.
+ * testsuite/libgomp.oacc-fortran/common-block-2.f90: New test.
+ * testsuite/libgomp.oacc-fortran/common-block-3.f90: New test.
+ * 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
struct goacc_thread *thr = goacc_thread ();
struct gomp_device_descr *acc_dev = thr->dev;
+ 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
+! Test data located inside common blocks. This test does not execrise
+! ACC DECLARE.
+
+module const
+ integer, parameter :: n = 100
+end module const
+
+subroutine check
+ use const
+
+ implicit none
+ integer i, x(n), y
+ common /BLOCK/ x, y
+
+ do i = 1, n
+ if (x(i) .ne. y) call abort
+ end do
+end subroutine check
+
+module m
+ use const
+ integer a(n), b
+ common /BLOCK/ a, b
+
+contains
+ subroutine mod_implicit_incr
+ implicit none
+ integer i
+
+ !$acc parallel loop
+ do i = 1, n
+ a(i) = b
+ end do
+ !$acc end parallel loop
+
+ call check
+ end subroutine mod_implicit_incr
+
+ subroutine mod_explicit_incr
+ implicit none
+ integer i
+
+ !$acc parallel loop copy(a(1:n)) copyin(b)
+ do i = 1, n
+ a(i) = b
+ end do
+ !$acc end parallel loop
+
+ call check
+ end subroutine mod_explicit_incr
+end module m
+
+subroutine sub_implicit_incr
+ use const
+
+ implicit none
+ integer i, x(n), y
+ common /BLOCK/ x, y
+
+ !$acc parallel loop
+ do i = 1, n
+ x(i) = y
+ end do
+ !$acc end parallel loop
+
+ call check
+end subroutine sub_implicit_incr
+
+subroutine sub_explicit_incr
+ use const
+
+ implicit none
+ integer i, x(n), y
+ common /BLOCK/ x, y
+
+ !$acc parallel loop copy(x(1:n)) copyin(y)
+ do i = 1, n
+ x(i) = y
+ end do
+ !$acc end parallel loop
+
+ call check
+end subroutine sub_explicit_incr
+
+program main
+ use m
+
+ implicit none
+
+ a(:) = -1
+ b = 5
+ call mod_implicit_incr
+
+ a(:) = -2
+ b = 6
+ call mod_explicit_incr
+
+ a(:) = -3
+ b = 7
+ call sub_implicit_incr
+
+ a(:) = -4
+ b = 8
+ call sub_explicit_incr
+end program main
--- /dev/null
+! Test data located inside common blocks. This test does not execrise
+! ACC DECLARE. All data clauses are explicit.
+
+module consts
+ integer, parameter :: n = 100
+end module consts
+
+subroutine validate
+ use consts
+
+ implicit none
+ integer i, j
+ real*4 x(n), y(n), z
+ common /BLOCK/ x, y, z, j
+
+ do i = 1, n
+ if (abs(x(i) - i - z) .ge. 0.0001) call abort
+ end do
+end subroutine validate
+
+subroutine incr
+ use consts
+
+ implicit none
+ integer i, j
+ real*4 x(n), y(n), z
+ common /BLOCK/ x, y, z, j
+
+ !$acc parallel loop pcopy(/BLOCK/)
+ do i = 1, n
+ x(i) = x(i) + z
+ end do
+ !$acc end parallel loop
+end subroutine incr
+
+program main
+ use consts
+
+ implicit none
+ integer i, j
+ real*4 a(n), b(n), c
+ common /BLOCK/ a, b, c, j
+
+ ! Test copyout, pcopy, device
+
+ !$acc data copyout(a, c)
+
+ c = 1.0
+
+ !$acc update device(c)
+
+ !$acc parallel loop pcopy(a)
+ do i = 1, n
+ a(i) = i
+ end do
+ !$acc end parallel loop
+
+ call incr
+ call incr
+ call incr
+ !$acc end data
+
+ c = 3.0
+ call validate
+
+ ! Test pcopy without copyout
+
+ c = 2.0
+ call incr
+ c = 5.0
+ call validate
+
+ ! Test create, delete, host, copyout, copyin
+
+ !$acc enter data create(b)
+
+ !$acc parallel loop pcopy(b)
+ do i = 1, n
+ b(i) = i
+ end do
+ !$acc end parallel loop
+
+ !$acc update host (b)
+
+ !$acc parallel loop pcopy(b) copyout(a) copyin(c)
+ do i = 1, n
+ a(i) = b(i) + c
+ end do
+ !$acc end parallel loop
+
+ !$acc exit data delete(b)
+
+ call validate
+
+ a(:) = b(:)
+ c = 0.0
+ call validate
+
+ ! Test copy
+
+ c = 1.0
+ !$acc parallel loop copy(/BLOCK/)
+ do i = 1, n
+ a(i) = b(i) + c
+ end do
+ !$acc end parallel loop
+
+ call validate
+
+ ! Test pcopyin, pcopyout FIXME
+
+ c = 2.0
+ !$acc data copyin(b, c) copyout(a)
+
+ !$acc parallel loop pcopyin(b, c) pcopyout(a)
+ do i = 1, n
+ a(i) = b(i) + c
+ end do
+ !$acc end parallel loop
+
+ !$acc end data
+
+ call validate
+
+ ! Test reduction, private
+
+ j = 0
+
+ !$acc parallel private(i) copy(j)
+ !$acc loop reduction(+:j)
+ do i = 1, n
+ j = j + 1
+ end do
+ !$acc end parallel
+
+ if (j .ne. n) call abort
+
+ ! Test firstprivate, copy
+
+ a(:) = 0
+ c = j
+
+ !$acc parallel loop firstprivate(c) copyout(a)
+ do i = 1, n
+ a(i) = i + c
+ end do
+ !$acc end parallel loop
+
+ call validate
+end program main
--- /dev/null
+! Test data located inside common blocks. This test does not execrise
+! ACC DECLARE. Most of the data clauses are implicit.
+
+module consts
+ integer, parameter :: n = 100
+end module consts
+
+subroutine validate
+ use consts
+
+ implicit none
+ integer i, j
+ real*4 x(n), y(n), z
+ common /BLOCK/ x, y, z, j
+
+ do i = 1, n
+ if (abs(x(i) - i - z) .ge. 0.0001) call abort
+ end do
+end subroutine validate
+
+subroutine incr_parallel
+ use consts
+
+ implicit none
+ integer i, j
+ real*4 x(n), y(n), z
+ common /BLOCK/ x, y, z, j
+
+ !$acc parallel loop
+ do i = 1, n
+ x(i) = x(i) + z
+ end do
+ !$acc end parallel loop
+end subroutine incr_parallel
+
+subroutine incr_kernels
+ use consts
+
+ implicit none
+ integer i, j
+ real*4 x(n), y(n), z
+ common /BLOCK/ x, y, z, j
+
+ !$acc kernels
+ do i = 1, n
+ x(i) = x(i) + z
+ end do
+ !$acc end kernels
+end subroutine incr_kernels
+
+program main
+ use consts
+
+ implicit none
+ integer i, j
+ real*4 a(n), b(n), c
+ common /BLOCK/ a, b, c, j
+
+ !$acc data copyout(a, c)
+
+ c = 1.0
+
+ !$acc update device(c)
+
+ !$acc parallel loop
+ do i = 1, n
+ a(i) = i
+ end do
+ !$acc end parallel loop
+
+ call incr_parallel
+ call incr_parallel
+ call incr_parallel
+ !$acc end data
+
+ c = 3.0
+ call validate
+
+ ! Test pcopy without copyout
+
+ c = 2.0
+ call incr_kernels
+ c = 5.0
+ call validate
+
+ !$acc kernels
+ do i = 1, n
+ b(i) = i
+ end do
+ !$acc end kernels
+
+ !$acc parallel loop
+ do i = 1, n
+ a(i) = b(i) + c
+ end do
+ !$acc end parallel loop
+
+ call validate
+
+ a(:) = b(:)
+ c = 0.0
+ call validate
+
+ ! Test copy
+
+ c = 1.0
+ !$acc parallel loop
+ do i = 1, n
+ a(i) = b(i) + c
+ end do
+ !$acc end parallel loop
+
+ call validate
+
+ c = 2.0
+ !$acc data copyin(b, c) copyout(a)
+
+ !$acc kernels
+ do i = 1, n
+ a(i) = b(i) + c
+ end do
+ !$acc end kernels
+
+ !$acc end data
+
+ call validate
+
+ j = 0
+
+ !$acc parallel loop reduction(+:j)
+ do i = 1, n
+ j = j + 1
+ end do
+ !$acc end parallel loop
+
+ if (j .ne. n) call abort
+end program main
--- /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
