2015-10-19 Nathan Sidwell <nathan@codesourcery.com>
gcc/
* omp-low.c (struct omp_region): Remove gwv_this field.
(struct omp_context): Remove gwv_this and gwv_belop fields.
(struct oacc_loop): Add marker field.
(enum oacc_loop): Adjust OLF_DIM_BASE initializer.
(extract_oacc_loop_mask): Delete.
(extrac_oacc_routine_gwv): Delete.
(oacc_loop_or_target_p): Delete.
(check_oacc_kernel_gwv): New.
(scan_omp_for): Remove OpenACC gwv mask handling. Check gang,
worker, vector argments.
(scan_omp_1_stmt): Remove OpenACC gwv mask checking.
(lower_oacc_head_mask): Set OLF_AUTO whenever possible. Ensure 1
level of headers.
(find_omp_for_region_gwv): Delete.
(find_omp_for_region_data): Remove gwv setting. Only set
independent field for kernels.
(find_omp_target_region_data): Delete.
(build_omp_regions_1): Set region kind here.
(new_oacc_loop_raw): Initialize marker field.
(new_oacc_loop): Likewise. Don't set mask here.
(new_oacc_loop_routine): Set marker field.
(dump_oacc_loop): Dump marker.
(oacc_loop_walk): Rename to ...
(oacc_loop_discover_walk): ... here. Adjust head & tail
recording.
(oacc_loop_sibling_nreverse): New.
(oacc_loop_discovery): Reverse siblings. Don't dump loops here.
(oacc_loop_process): Adjust & remove asserts.
(oacc_loop_fixed_partitions): New.
(oacc_loop_partition): New.
(execute_oacc_device_lower): Partition loops, neuter unused loop
heads & tails.
gcc/testsuite/
* gcc/testsuite/c-c++-common/goacc/routine-7.c: Adjust diagnostics.
* gcc/testsuite/c-c++-common/goacc/loop-3.c: Adjust diagnostics.
* gcc/testsuite/c-c++-common/goacc/routine-6.c: Adjust diagnostics.
* gcc/testsuite/c-c++-common/goacc/loop-2.c: Adjust diagnostics.
* gcc/testsuite/c-c++-common/goacc/loop-4.c: Adjust diagnostics.
* gcc/testsuite/gfortran.dg/goacc/loop-6.f95: Adjust diagnostics.
* gcc/testsuite/gfortran.dg/goacc/routine-4.f90: Adjust diagnostics.
* gcc/testsuite/gfortran.dg/goacc/routine-5.f90: Adjust diagnostics.
===================================================================
@@ -141,9 +141,6 @@ struct omp_region
/* Records a generic kind field. */
int kind;
- /* For an OpenACC loop, the level of parallelism requested. */
- int gwv_this;
-
/* For an OpenACC loop directive, true if has the 'independent' clause. */
bool independent;
};
@@ -203,14 +200,6 @@ struct omp_context
/* The number of reductions in a loop. */
int reductions;
-
- /* For OpenACC loops, a mask of gang, worker and vector used at
- levels below this one. */
- int gwv_below;
- /* For OpenACC loops, a mask of gang, worker and vector used at
- this level and above. For parallel and kernels clauses, a mask
- indicating which of num_gangs/num_workers/num_vectors was used. */
- int gwv_this;
};
/* A structure holding the elements of:
@@ -249,7 +238,8 @@ struct oacc_loop
location_t loc; /* Location of the loop start. */
- /* Start of head and tail. */
+ gcall *marker; /* Initial head marker. */
+
gcall *heads[GOMP_DIM_MAX]; /* Head marker functions. */
gcall *tails[GOMP_DIM_MAX]; /* Tail marker functions. */
@@ -271,7 +261,7 @@ enum oacc_loop_flags
OLF_GANG_STATIC = 1u << 3, /* Gang partitioning is static (has op). */
/* Explicitly specified loop axes. */
- OLF_DIM_BASE = 4 - GOMP_DIM_GANG,
+ OLF_DIM_BASE = 4,
OLF_DIM_GANG = 1u << (OLF_DIM_BASE + GOMP_DIM_GANG),
OLF_DIM_WORKER = 1u << (OLF_DIM_BASE + GOMP_DIM_WORKER),
OLF_DIM_VECTOR = 1u << (OLF_DIM_BASE + GOMP_DIM_VECTOR),
@@ -300,34 +290,6 @@ static gphi *find_phi_with_arg_on_edge (
*handled_ops_p = false; \
break;
-/* Extract the gang, worker and vector clauses associated with CTX.
-
- GWV_THIS contains the current level of parallelism the loop nest.
- I.e. if the loop above contains a gang clause, and the current loop
- contains a vector clause, gwv_this will have the GOM_DIM_GANG and
- GOMP_DIM_VECTOR bits set. This function extracts the level of
- parallelism only associated with the current loop, e.g.
- GOMP_DIM_VECTOR. */
-
-static int
-extract_oacc_loop_mask (omp_context *ctx)
-{
- int loop_flags = 0;
-
- if (is_gimple_omp_oacc (ctx->stmt))
- {
- omp_context *outer = ctx->outer;
-
- if (outer && gimple_code (outer->stmt) != GIMPLE_OMP_FOR)
- outer = NULL;
-
- loop_flags = outer ? ctx->gwv_this & (~outer->gwv_this)
- : ctx->gwv_this;
- }
-
- return loop_flags;
-}
-
static bool
is_oacc_parallel (omp_context *ctx)
{
@@ -464,36 +426,6 @@ is_combined_parallel (struct omp_region
return region->is_combined_parallel;
}
-/* Return the gang, worker and vector attributes from associated with
- FNDECL. Returns a GOMP_DIM for the lowest level of parallelism beginning
- with GOMP_DIM_GANG, or -1 if the routine is a SEQ. Otherwise, return 0 if
- the FNDECL is not an acc routine.
-*/
-
-static int
-extract_oacc_routine_gwv (tree fndecl)
-{
- tree attrs = get_oacc_fn_attrib (fndecl);
- tree pos;
- unsigned gwv = 0;
- int i;
- int ret = 0;
-
- if (attrs != NULL_TREE)
- {
- for (i = 0, pos = TREE_VALUE (attrs);
- gwv == 0 && i != GOMP_DIM_MAX;
- i++, pos = TREE_CHAIN (pos))
- if (TREE_PURPOSE (pos) != boolean_false_node)
- return 1 << i;
-
- ret = -1;
- }
-
- return ret;
-}
-
-
/* Extract the header elements of parallel loop FOR_STMT and store
them into *FD. */
@@ -2594,19 +2526,6 @@ enclosing_target_ctx (omp_context *ctx)
return ctx;
}
-static bool
-oacc_loop_or_target_p (gimple *stmt)
-{
- enum gimple_code outer_type = gimple_code (stmt);
- return ((outer_type == GIMPLE_OMP_TARGET
- && ((gimple_omp_target_kind (stmt)
- == GF_OMP_TARGET_KIND_OACC_PARALLEL)
- || (gimple_omp_target_kind (stmt)
- == GF_OMP_TARGET_KIND_OACC_KERNELS)))
- || (outer_type == GIMPLE_OMP_FOR
- && gimple_omp_for_kind (stmt) == GF_OMP_FOR_KIND_OACC_LOOP));
-}
-
/* Return true if ctx is part of an oacc kernels region. */
static bool
@@ -2623,127 +2542,134 @@ ctx_in_oacc_kernels_region (omp_context
return false;
}
+/* Check the parallelism clauses inside a kernels regions.
+ Until kernels handling moves to use the same loop indirection
+ scheme as parallel, we need to do this checking early. */
+
+static unsigned
+check_oacc_kernel_gwv (gomp_for *stmt, omp_context *ctx)
+{
+ bool checking = true;
+ unsigned outer_mask = 0;
+ unsigned this_mask = 0;
+ bool has_seq = false, has_auto = false;
+
+ if (ctx->outer)
+ outer_mask = check_oacc_kernel_gwv (NULL, ctx->outer);
+ if (!stmt)
+ {
+ checking = false;
+ if (gimple_code (ctx->stmt) != GIMPLE_OMP_FOR)
+ return outer_mask;
+ stmt = as_a <gomp_for *> (ctx->stmt);
+ }
+
+ for (tree c = gimple_omp_for_clauses (stmt); c; c = OMP_CLAUSE_CHAIN (c))
+ {
+ switch (OMP_CLAUSE_CODE (c))
+ {
+ case OMP_CLAUSE_GANG:
+ this_mask |= GOMP_DIM_MASK (GOMP_DIM_GANG);
+ break;
+ case OMP_CLAUSE_WORKER:
+ this_mask |= GOMP_DIM_MASK (GOMP_DIM_WORKER);
+ break;
+ case OMP_CLAUSE_VECTOR:
+ this_mask |= GOMP_DIM_MASK (GOMP_DIM_VECTOR);
+ break;
+ case OMP_CLAUSE_SEQ:
+ has_seq = true;
+ break;
+ case OMP_CLAUSE_AUTO:
+ has_auto = true;
+ break;
+ default:
+ break;
+ }
+ }
+
+ if (checking)
+ {
+ if (has_seq && (this_mask || has_auto))
+ error_at (gimple_location (stmt), "%<seq%> overrides other OpenACC loop specifiers");
+ else if (has_auto && this_mask)
+ error_at (gimple_location (stmt), "%<auto%> conflicts with other OpenACC loop specifiers");
+
+ if (this_mask & outer_mask)
+ error_at (gimple_location (stmt), "inner loop uses same OpenACC parallelism as containing loop");
+ }
+
+ return outer_mask | this_mask;
+}
+
/* Scan a GIMPLE_OMP_FOR. */
static void
scan_omp_for (gomp_for *stmt, omp_context *outer_ctx)
{
- enum gimple_code outer_type = GIMPLE_ERROR_MARK;
omp_context *ctx;
size_t i;
tree clauses = gimple_omp_for_clauses (stmt);
- bool gwv_clause = false;
- bool auto_clause = false;
- bool seq_clause = false;
- int gwv_routine = 0;
- bool in_oacc_kernels_region = ctx_in_oacc_kernels_region (outer_ctx);
-
- if (outer_ctx)
- outer_type = gimple_code (outer_ctx->stmt);
- else
- {
- gwv_routine = extract_oacc_routine_gwv (current_function_decl);
- if (gwv_routine > 0)
- gwv_routine = gwv_routine >> 1;
- }
ctx = new_omp_context (stmt, outer_ctx);
if (is_gimple_omp_oacc (stmt))
{
- if (outer_ctx && outer_type == GIMPLE_OMP_FOR)
- ctx->gwv_this = outer_ctx->gwv_this;
- for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
- {
- int val;
- if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_GANG)
- {
- val = GOMP_DIM_MASK (GOMP_DIM_GANG);
- gwv_clause = true;
- }
- else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_WORKER)
- {
- val = GOMP_DIM_MASK (GOMP_DIM_WORKER);
- gwv_clause = true;
- }
- else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_VECTOR)
- {
- val = GOMP_DIM_MASK (GOMP_DIM_VECTOR);
- gwv_clause = true;
- }
- else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_SEQ)
- {
- seq_clause = true;
- continue;
- }
- else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_AUTO)
- {
- auto_clause = true;
- continue;
- }
- else
- continue;
- ctx->gwv_this |= val;
- if (!outer_ctx)
- {
- /* Skip; not nested inside a region. */
- continue;
- }
- if (!oacc_loop_or_target_p (outer_ctx->stmt))
- {
- /* Skip; not nested inside an OpenACC region. */
- continue;
- }
- if (outer_type == GIMPLE_OMP_FOR)
- outer_ctx->gwv_below |= val;
- if (OMP_CLAUSE_OPERAND (c, 0) != NULL_TREE)
- {
- omp_context *enclosing = enclosing_target_ctx (outer_ctx);
- /* Enclosing may be null if we are inside an acc routine. If
- that's the case, treat this loop as a parallel. */
- if (enclosing == NULL || gimple_omp_target_kind (enclosing->stmt)
- == GF_OMP_TARGET_KIND_OACC_PARALLEL)
- error_at (gimple_location (stmt),
- "no arguments allowed to gang, worker and vector clauses inside parallel");
- }
- }
+ omp_context *tgt = enclosing_target_ctx (outer_ctx);
+
+ if (!tgt || is_oacc_parallel (tgt))
+ for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
+ {
+ char const *check = NULL;
+
+ switch (OMP_CLAUSE_CODE (c))
+ {
+ case OMP_CLAUSE_GANG:
+ check = "gang";
+ break;
+
+ case OMP_CLAUSE_WORKER:
+ check = "worker";
+ break;
+
+ case OMP_CLAUSE_VECTOR:
+ check = "vector";
+ break;
+
+ default:
+ break;
+ }
+
+ if (check && OMP_CLAUSE_OPERAND (c, 0))
+ error_at (gimple_location (stmt),
+ "argument not permitted on %<%s%> clause in"
+ " OpenACC %<parallel%>", check);
+ }
+ }
- /* Filter out any OpenACC clauses which aren't associated with
- gangs, workers or vectors. Such reductions are no-ops. */
- if (extract_oacc_loop_mask (ctx) == 0
- || in_oacc_kernels_region)
+ if (is_gimple_omp_oacc (stmt))
+ {
+ omp_context *tgt = enclosing_target_ctx (ctx);
+ if (tgt && is_oacc_kernels (tgt))
{
- /* First filter out the clauses at the beginning of the chain. */
- while (clauses && OMP_CLAUSE_CODE (clauses) == OMP_CLAUSE_REDUCTION)
- {
- clauses = OMP_CLAUSE_CHAIN (clauses);
- }
+ /* Strip out reductions, as they are not handled yet. */
+ tree *prev_ptr = &clauses;
- if (clauses != NULL)
+ while (tree probe = *prev_ptr)
{
- /* Filter out the remaining clauses. */
- for (tree c = OMP_CLAUSE_CHAIN (clauses), prev = clauses;
- c; c = OMP_CLAUSE_CHAIN (c))
- {
- if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION)
- {
- tree t = OMP_CLAUSE_CHAIN (c);
- OMP_CLAUSE_CHAIN (prev) = t;
- }
- else
- prev = c;
- }
+ tree *next_ptr = &OMP_CLAUSE_CHAIN (probe);
+
+ if (OMP_CLAUSE_CODE (probe) == OMP_CLAUSE_REDUCTION)
+ *prev_ptr = *next_ptr;
+ else
+ prev_ptr = next_ptr;
}
gimple_omp_for_set_clauses (stmt, clauses);
+ check_oacc_kernel_gwv (stmt, ctx);
}
}
- if ((gwv_clause && auto_clause) || (auto_clause && seq_clause))
- error_at (gimple_location (stmt), "incompatible use of clause auto");
- else if (gwv_clause && seq_clause)
- error_at (gimple_location (stmt), "incompatible use of clause seq");
-
scan_sharing_clauses (clauses, ctx);
scan_omp (gimple_omp_for_pre_body_ptr (stmt), ctx);
@@ -2755,25 +2681,6 @@ scan_omp_for (gomp_for *stmt, omp_contex
scan_omp_op (gimple_omp_for_incr_ptr (stmt, i), ctx);
}
scan_omp (gimple_omp_body_ptr (stmt), ctx);
-
- if (is_gimple_omp_oacc (stmt))
- {
- if (ctx->gwv_this & ctx->gwv_below)
- error_at (gimple_location (stmt),
- "gang, worker and vector may occur only once in a loop nest");
- else if (ctx->gwv_below != 0
- && ctx->gwv_this > ctx->gwv_below)
- error_at (gimple_location (stmt),
- "gang, worker and vector must occur in this order in a loop nest");
- else if (!outer_ctx && ctx->gwv_this != 0 && gwv_routine != 0
- && ((ffs (ctx->gwv_this) <= gwv_routine)
- || gwv_routine < 0))
- error_at (gimple_location (stmt),
- "invalid parallelism inside acc routine");
-
- if (outer_ctx && outer_type == GIMPLE_OMP_FOR)
- outer_ctx->gwv_below |= ctx->gwv_below;
- }
}
/* Scan an OpenMP sections directive. */
@@ -2840,19 +2747,6 @@ scan_omp_target (gomp_target *stmt, omp_
gimple_omp_target_set_child_fn (stmt, ctx->cb.dst_fn);
}
- if (is_gimple_omp_oacc (stmt))
- {
- for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
- {
- if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_NUM_GANGS)
- ctx->gwv_this |= GOMP_DIM_MASK (GOMP_DIM_GANG);
- else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_NUM_WORKERS)
- ctx->gwv_this |= GOMP_DIM_MASK (GOMP_DIM_WORKER);
- else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_VECTOR_LENGTH)
- ctx->gwv_this |= GOMP_DIM_MASK (GOMP_DIM_VECTOR);
- }
- }
-
scan_sharing_clauses (clauses, ctx);
scan_omp (gimple_omp_body_ptr (stmt), ctx);
@@ -3376,16 +3270,6 @@ scan_omp_1_stmt (gimple_stmt_iterator *g
default:
break;
}
- else if (ctx && is_gimple_omp_oacc (ctx->stmt)
- && !is_oacc_parallel (ctx))
- {
- /* Is this a call to an acc routine? */
- int gwv = extract_oacc_routine_gwv (fndecl);
-
- if (gwv > 0 && ffs (ctx->gwv_this) >= ffs (gwv))
- error_at (gimple_location (stmt),
- "incompatible parallelism with acc routine");
- }
}
}
if (remove)
@@ -4975,28 +4859,18 @@ lower_oacc_head_mark (location_t loc, tr
}
/* In a parallel region, loops are implicitly INDEPENDENT. */
- if (is_oacc_parallel (ctx))
+ omp_context *tgt = enclosing_target_ctx (ctx);
+ if (!tgt || is_oacc_parallel (tgt))
tag |= OLF_INDEPENDENT;
- /* In a kernels region, a loop lacking SEQ, GANG, WORKER and/or
- VECTOR is implicitly AUTO. */
- if (is_oacc_kernels (ctx)
- && !(tag & (((GOMP_DIM_MASK (GOMP_DIM_MAX) - 1) << OLF_DIM_BASE)
- | OLF_SEQ)))
+ /* A loop lacking SEQ, GANG, WORKER and/or VECTOR is implicitly AUTO. */
+ if (!(tag & (((GOMP_DIM_MASK (GOMP_DIM_MAX) - 1) << OLF_DIM_BASE)
+ | OLF_SEQ)))
tag |= OLF_AUTO;
- {
- /* Check we didn't discover any different partitioning from the
- existing scheme. */
- unsigned mask = ctx->gwv_this;
- if (ctx->outer && gimple_code (ctx->outer->stmt) == GIMPLE_OMP_FOR)
- mask &= ~ctx->outer->gwv_this;
-
- gcc_assert (mask == ((tag >> OLF_DIM_BASE)
- & (GOMP_DIM_MASK (GOMP_DIM_MAX) - 1)));
- }
-
- /* TODO: allocate at least one level, for auto allocation. */
+ /* Ensure at least one level. */
+ if (!levels)
+ levels++;
args.safe_push (build_int_cst (integer_type_node, levels));
args.safe_push (build_int_cst (integer_type_node, tag));
@@ -11121,32 +10995,9 @@ expand_omp (struct omp_region *region)
/* Map each basic block to an omp_region. */
static hash_map<basic_block, omp_region *> *bb_region_map;
-/* Return a mask of GWV bits for region REGION associated with an
- OMP_FOR STMT. */
-
-static int
-find_omp_for_region_gwv (gimple *stmt)
-{
- int tmp = 0;
-
- if (!is_gimple_omp_oacc (stmt))
- return 0;
-
- tree clauses = gimple_omp_for_clauses (stmt);
- if (find_omp_clause (clauses, OMP_CLAUSE_GANG))
- tmp |= GOMP_DIM_MASK (GOMP_DIM_GANG);
- if (find_omp_clause (clauses, OMP_CLAUSE_WORKER))
- tmp |= GOMP_DIM_MASK (GOMP_DIM_WORKER);
- if (find_omp_clause (clauses, OMP_CLAUSE_VECTOR))
- tmp |= GOMP_DIM_MASK (GOMP_DIM_VECTOR);
-
- return tmp;
-}
-
static void
find_omp_for_region_data (struct omp_region *region, gomp_for *stmt)
{
- region->gwv_this = find_omp_for_region_gwv (stmt);
region->kind = gimple_omp_for_kind (stmt);
if (region->kind == GF_OMP_FOR_KIND_OACC_LOOP)
@@ -11160,37 +11011,12 @@ find_omp_for_region_data (struct omp_reg
tree clauses = gimple_omp_for_clauses (stmt);
- if (target_region->kind == GF_OMP_TARGET_KIND_OACC_PARALLEL
- && !find_omp_clause (clauses, OMP_CLAUSE_SEQ))
- /* In OpenACC parallel constructs, 'independent' is implied on all
- loop directives without a 'seq' clause. */
- region->independent = true;
- else if (target_region->kind == GF_OMP_TARGET_KIND_OACC_KERNELS
- && find_omp_clause (clauses, OMP_CLAUSE_INDEPENDENT))
+ if (target_region->kind == GF_OMP_TARGET_KIND_OACC_KERNELS
+ && find_omp_clause (clauses, OMP_CLAUSE_INDEPENDENT))
region->independent = true;
}
}
-/* Fill in additional data for a region REGION associated with an
- OMP_TARGET STMT. */
-
-static void
-find_omp_target_region_data (struct omp_region *region,
- gomp_target *stmt)
-{
- if (!is_gimple_omp_oacc (stmt))
- return;
-
- tree clauses = gimple_omp_target_clauses (stmt);
- if (find_omp_clause (clauses, OMP_CLAUSE_NUM_GANGS))
- region->gwv_this |= GOMP_DIM_MASK (GOMP_DIM_GANG);
- if (find_omp_clause (clauses, OMP_CLAUSE_NUM_WORKERS))
- region->gwv_this |= GOMP_DIM_MASK (GOMP_DIM_WORKER);
- if (find_omp_clause (clauses, OMP_CLAUSE_VECTOR_LENGTH))
- region->gwv_this |= GOMP_DIM_MASK (GOMP_DIM_VECTOR);
- region->kind = gimple_omp_target_kind (stmt);
-}
-
/* Helper for build_omp_regions. Scan the dominator tree starting at
block BB. PARENT is the region that contains BB. If SINGLE_TREE is
true, the function ends once a single tree is built (otherwise, whole
@@ -11260,8 +11086,8 @@ build_omp_regions_1 (basic_block bb, str
case GF_OMP_TARGET_KIND_OACC_PARALLEL:
case GF_OMP_TARGET_KIND_OACC_KERNELS:
case GF_OMP_TARGET_KIND_OACC_DATA:
- find_omp_target_region_data (region,
- as_a <gomp_target *> (stmt));
+ if (is_gimple_omp_oacc (stmt))
+ region->kind = gimple_omp_target_kind (stmt);
break;
case GF_OMP_TARGET_KIND_UPDATE:
case GF_OMP_TARGET_KIND_OACC_UPDATE:
@@ -15894,6 +15720,7 @@ new_oacc_loop_raw (oacc_loop *parent, lo
}
loop->loc = loc;
+ loop->marker = NULL;
memset (loop->heads, 0, sizeof (loop->heads));
memset (loop->tails, 0, sizeof (loop->tails));
loop->routine = NULL_TREE;
@@ -15918,22 +15745,22 @@ new_oacc_loop_outer (tree decl)
Link into PARENT loop. Return the new loop. */
static oacc_loop *
-new_oacc_loop (oacc_loop *parent, gcall *head)
+new_oacc_loop (oacc_loop *parent, gcall *marker)
{
- oacc_loop *loop = new_oacc_loop_raw (parent, gimple_location (head));
+ oacc_loop *loop = new_oacc_loop_raw (parent, gimple_location (marker));
+
+ loop->marker = marker;
+
+ /* TODO: This is where device_type flattening would occur for the loop
+ flags. */
- loop->flags = TREE_INT_CST_LOW (gimple_call_arg (head, 2));
+ loop->flags = TREE_INT_CST_LOW (gimple_call_arg (marker, 2));
tree chunk_size = integer_zero_node;
if (loop->flags & OLF_GANG_STATIC)
- chunk_size = gimple_call_arg (head,3);
+ chunk_size = gimple_call_arg (marker, 3);
loop->chunk_size = chunk_size;
- /* Set the mask from the incoming flags.
- TODO: Be smarter and more flexible. */
- loop->mask = ((loop->flags >> OLF_DIM_BASE)
- & (GOMP_DIM_MASK (GOMP_DIM_MAX) - 1));
-
return loop;
}
@@ -15948,7 +15775,8 @@ new_oacc_loop_routine (oacc_loop *parent
int level = oacc_validate_dims (decl, attrs, dims);
gcc_assert (level >= 0);
-
+
+ loop->marker = call;
loop->routine = decl;
loop->mask = ((GOMP_DIM_MASK (GOMP_DIM_MAX) - 1)
^ (GOMP_DIM_MASK (level) - 1));
@@ -16015,6 +15843,9 @@ dump_oacc_loop (FILE *file, oacc_loop *l
loop->flags, loop->mask,
LOCATION_FILE (loop->loc), LOCATION_LINE (loop->loc));
+ if (loop->marker)
+ print_gimple_stmt (file, loop->marker, depth * 2, 0);
+
if (loop->routine)
fprintf (file, "%*sRoutine %s:%u:%s\n",
depth * 2, "", DECL_SOURCE_FILE (loop->routine),
@@ -16049,7 +15880,7 @@ debug_oacc_loop (oacc_loop *loop)
nested. */
static void
-oacc_loop_walk (oacc_loop *loop, basic_block bb)
+oacc_loop_discover_walk (oacc_loop *loop, basic_block bb)
{
if (bb->flags & BB_VISITED)
return;
@@ -16105,19 +15936,16 @@ oacc_loop_walk (oacc_loop *loop, basic_b
if (code == IFN_UNIQUE_OACC_HEAD_MARK)
loop = new_oacc_loop (loop, call);
remaining = count;
- if (remaining)
- remaining--;
}
- else
+ gcc_assert (count == remaining);
+ if (remaining)
{
- gcc_assert (count == remaining);
remaining--;
+ if (code == IFN_UNIQUE_OACC_HEAD_MARK)
+ loop->heads[marker] = call;
+ else
+ loop->tails[remaining] = call;
}
-
- if (code == IFN_UNIQUE_OACC_HEAD_MARK)
- loop->heads[marker] = call;
- else
- loop->tails[remaining] = call;
marker++;
}
}
@@ -16129,7 +15957,29 @@ oacc_loop_walk (oacc_loop *loop, basic_b
edge_iterator ei;
FOR_EACH_EDGE (e, ei, bb->succs)
- oacc_loop_walk (loop, e->dest);
+ oacc_loop_discover_walk (loop, e->dest);
+}
+
+/* LOOP is the first sibling. Reverse the order in place and return
+ the new first sibling. Recurse to child loops. */
+
+static oacc_loop *
+oacc_loop_sibling_nreverse (oacc_loop *loop)
+{
+ oacc_loop *last = NULL;
+ do
+ {
+ if (loop->child)
+ loop->child = oacc_loop_sibling_nreverse (loop->child);
+
+ oacc_loop *next = loop->sibling;
+ loop->sibling = last;
+ last = loop;
+ loop = next;
+ }
+ while (loop);
+
+ return last;
}
/* Discover the OpenACC loops marked up by HEAD and TAIL markers for
@@ -16141,19 +15991,16 @@ oacc_loop_discovery ()
basic_block bb;
oacc_loop *top = new_oacc_loop_outer (current_function_decl);
- oacc_loop_walk (top, ENTRY_BLOCK_PTR_FOR_FN (cfun));
+ oacc_loop_discover_walk (top, ENTRY_BLOCK_PTR_FOR_FN (cfun));
+
+ /* The siblings were constructed in reverse order, reverse them so
+ that diagnostics come out in an unsurprising order. */
+ top = oacc_loop_sibling_nreverse (top);
/* Reset the visited flags. */
FOR_ALL_BB_FN (bb, cfun)
bb->flags &= ~BB_VISITED;
- if (dump_file)
- {
- fprintf (dump_file, "OpenACC loops\n");
- dump_oacc_loop (dump_file, top, 0);
- fprintf (dump_file, "\n");
- }
-
return top;
}
@@ -16260,12 +16107,11 @@ oacc_loop_process (oacc_loop *loop)
if (loop->child)
oacc_loop_process (loop->child);
- int ix;
- unsigned mask = loop->mask;
- unsigned dim = GOMP_DIM_GANG;
-
- if (mask && !loop->routine)
+ if (loop->mask && !loop->routine)
{
+ int ix;
+ unsigned mask = loop->mask;
+ unsigned dim = GOMP_DIM_GANG;
tree mask_arg = build_int_cst (unsigned_type_node, mask);
tree chunk_arg = loop->chunk_size;
@@ -16284,17 +16130,134 @@ oacc_loop_process (oacc_loop *loop)
mask ^= GOMP_DIM_MASK (dim);
}
}
- else
- gcc_assert (!loop->heads[1] && !loop->tails[1]
- && (loop->routine || !loop->parent
- || integer_zerop (gimple_call_arg (loop->heads[0], 1))));
-
- gcc_assert (loop->routine || !mask);
if (loop->sibling)
oacc_loop_process (loop->sibling);
}
+/* Walk the OpenACC loop heirarchy checking and assigning the
+ programmer-specified partitionings. OUTER_MASK is the partitioning
+ this loop is contained within. Return partitiong mask used within
+ this loop nest. */
+
+static unsigned
+oacc_loop_fixed_partitions (oacc_loop *loop, unsigned outer_mask)
+{
+ unsigned this_mask = loop->mask;
+ bool has_auto = false;
+ bool noisy = true;
+
+#ifdef ACCEL_COMPILER
+ /* When device_type is supported, we want the device compiler to be
+ noisy, if the loop parameters are device_type-specific. */
+ noisy = false;
+#endif
+
+ if (!loop->routine)
+ {
+ bool auto_par = (loop->flags & OLF_AUTO) != 0;
+ bool seq_par = (loop->flags & OLF_SEQ) != 0;
+
+ this_mask = ((loop->flags >> OLF_DIM_BASE)
+ & (GOMP_DIM_MASK (GOMP_DIM_MAX) - 1));
+
+ if ((this_mask != 0) + auto_par + seq_par > 1)
+ {
+ if (noisy)
+ error_at (loop->loc,
+ seq_par
+ ? "%<seq%> overrides other OpenACC loop specifiers"
+ : "%<auto%> conflicts with other OpenACC loop specifiers");
+ auto_par = false;
+ loop->flags &= ~OLF_AUTO;
+ if (seq_par)
+ {
+ loop->flags &=
+ ~((GOMP_DIM_MASK (GOMP_DIM_MAX) - 1) << OLF_DIM_BASE);
+ this_mask = 0;
+ }
+ }
+ if (auto_par && (loop->flags & OLF_INDEPENDENT))
+ has_auto = true;
+ }
+
+ if (this_mask & outer_mask)
+ {
+ const oacc_loop *outer;
+ for (outer = loop->parent; outer; outer = outer->parent)
+ if (outer->mask & this_mask)
+ break;
+
+ if (noisy)
+ {
+ if (outer)
+ {
+ error_at (loop->loc,
+ "%s uses same OpenACC parallelism as containing loop",
+ loop->routine ? "routine call" : "inner loop");
+ inform (outer->loc, "containing loop here");
+ }
+ else
+ error_at (loop->loc,
+ "%s uses OpenACC parallelism disallowed by containing routine",
+ loop->routine ? "routine call" : "loop");
+
+ if (loop->routine)
+ inform (DECL_SOURCE_LOCATION (loop->routine),
+ "routine %qD declared here", loop->routine);
+ }
+ this_mask &= ~outer_mask;
+ }
+ else
+ {
+ unsigned outermost = this_mask & -this_mask;
+
+ if (outermost && outermost <= outer_mask)
+ {
+ if (noisy)
+ {
+ error_at (loop->loc,
+ "incorrectly nested OpenACC loop parallelism");
+
+ const oacc_loop *outer;
+ for (outer = loop->parent;
+ outer->flags && outer->flags < outermost;
+ outer = outer->parent)
+ continue;
+ inform (outer->loc, "containing loop here");
+ }
+
+ this_mask &= ~outermost;
+ }
+ }
+
+ loop->mask = this_mask;
+
+ if (loop->child
+ && oacc_loop_fixed_partitions (loop->child, outer_mask | this_mask))
+ has_auto = true;
+
+ if (loop->sibling
+ && oacc_loop_fixed_partitions (loop->sibling, outer_mask))
+ has_auto = true;
+
+ return has_auto;
+}
+
+/* Walk the OpenACC loop heirarchy to check and assign partitioning
+ axes. */
+
+static void
+oacc_loop_partition (oacc_loop *loop, int fn_level)
+{
+ unsigned outer_mask = 0;
+
+ if (fn_level >= 0)
+ outer_mask = GOMP_DIM_MASK (fn_level) - 1;
+
+ oacc_loop_fixed_partitions (loop, outer_mask);
+}
+
/* Default launch dimension validator. Force everything to 1. A
backend that wants to provide larger dimensions must override this
hook. */
@@ -16415,11 +16378,18 @@ execute_oacc_device_lower ()
/* Not an offloaded function. */
return 0;
- oacc_validate_dims (current_function_decl, attrs, dims);
+ int fn_level = oacc_validate_dims (current_function_decl, attrs, dims);
- /* Discover and process the loops. */
+ /* Discover, partition and process the loops. */
oacc_loop *loops = oacc_loop_discovery ();
+ oacc_loop_partition (loops, fn_level);
oacc_loop_process (loops);
+ if (dump_file)
+ {
+ fprintf (dump_file, "OpenACC loops\n");
+ dump_oacc_loop (dump_file, loops, 0);
+ fprintf (dump_file, "\n");
+ }
/* Offloaded targets may introduce new basic blocks, which require
dominance information to update SSA. */
@@ -16473,7 +16443,13 @@ execute_oacc_device_lower ()
case IFN_GOACC_REDUCTION_TEARDOWN:
/* Mark the function for SSA renaming. */
mark_virtual_operands_for_renaming (cfun);
- targetm.goacc.reduction (call);
+
+ /* If the level is -1, this ended up being an unused
+ axis. Handle as a default. */
+ if (integer_minus_onep (gimple_call_arg (call, 2)))
+ default_goacc_reduction (call);
+ else
+ targetm.goacc.reduction (call);
rescan = 1;
break;
@@ -16481,14 +16457,22 @@ execute_oacc_device_lower ()
{
unsigned code = TREE_INT_CST_LOW (gimple_call_arg (call, 0));
- if ((code == IFN_UNIQUE_OACC_FORK
- || code == IFN_UNIQUE_OACC_JOIN)
- && (targetm.goacc.fork_join
- (call, dims, code == IFN_UNIQUE_OACC_FORK)))
- rescan = -1;
- else if (code == IFN_UNIQUE_OACC_HEAD_MARK
- || code == IFN_UNIQUE_OACC_TAIL_MARK)
- rescan = -1;
+ switch (code)
+ {
+ case IFN_UNIQUE_OACC_FORK:
+ case IFN_UNIQUE_OACC_JOIN:
+ if (integer_minus_onep (gimple_call_arg (call, 1)))
+ rescan = -1;
+ else if (targetm.goacc.fork_join
+ (call, dims, code == IFN_UNIQUE_OACC_FORK))
+ rescan = -1;
+ break;
+
+ case IFN_UNIQUE_OACC_HEAD_MARK:
+ case IFN_UNIQUE_OACC_TAIL_MARK:
+ rescan = -1;
+ break;
+ }
break;
}
}
===================================================================
@@ -32,7 +32,7 @@ worker (int red)
for (int i = 0; i < 10; i++)
red ++;
-#pragma acc loop gang reduction (+:red) // { dg-error "invalid parallelism inside acc routine" }
+#pragma acc loop gang reduction (+:red) // { dg-error "disallowed by containing routine" }
for (int i = 0; i < 10; i++)
red ++;
@@ -55,11 +55,11 @@ vector (int red)
for (int i = 0; i < 10; i++)
red ++;
-#pragma acc loop gang reduction (+:red) // { dg-error "invalid parallelism inside acc routine" }
+#pragma acc loop gang reduction (+:red) // { dg-error "disallowed by containing routine" }
for (int i = 0; i < 10; i++)
red ++;
-#pragma acc loop worker reduction (+:red) // { dg-error "invalid parallelism inside acc routine" }
+#pragma acc loop worker reduction (+:red) // { dg-error "disallowed by containing routine" }
for (int i = 0; i < 10; i++)
red ++;
@@ -78,15 +78,15 @@ seq (int red)
for (int i = 0; i < 10; i++)
red ++;
-#pragma acc loop gang reduction (+:red) // { dg-error "invalid parallelism inside acc routine" }
+#pragma acc loop gang reduction (+:red) // { dg-error "disallowed by containing routine" }
for (int i = 0; i < 10; i++)
red ++;
-#pragma acc loop worker reduction (+:red) // { dg-error "invalid parallelism inside acc routine" }
+#pragma acc loop worker reduction (+:red) // { dg-error "disallowed by containing routine" }
for (int i = 0; i < 10; i++)
red ++;
-#pragma acc loop vector reduction (+:red) // { dg-error "invalid parallelism inside acc routine" }
+#pragma acc loop vector reduction (+:red) // { dg-error "disallowed by containing routine" }
for (int i = 0; i < 10; i++)
red ++;
===================================================================
@@ -8,27 +8,27 @@ void par1 (void)
#pragma acc parallel
{
-#pragma acc loop gang(5) // { dg-error "no arguments allowed to gang" }
+#pragma acc loop gang(5) // { dg-error "argument not permitted" }
for (i = 0; i < 10; i++)
{ }
-#pragma acc loop gang(num:5) // { dg-error "no arguments allowed to gang" }
+#pragma acc loop gang(num:5) // { dg-error "argument not permitted" }
for (i = 0; i < 10; i++)
{ }
-#pragma acc loop worker(5) // { dg-error "no arguments allowed to gang" }
+#pragma acc loop worker(5) // { dg-error "argument not permitted" }
for (i = 0; i < 10; i++)
{ }
-#pragma acc loop worker(num:5) // { dg-error "no arguments allowed to gang" }
+#pragma acc loop worker(num:5) // { dg-error "argument not permitted" }
for (i = 0; i < 10; i++)
{ }
-#pragma acc loop vector(5) // { dg-error "no arguments allowed to gang" }
+#pragma acc loop vector(5) // { dg-error "argument not permitted" }
for (i = 0; i < 10; i++)
{ }
-#pragma acc loop vector(length:5) // { dg-error "no arguments allowed to gang" }
+#pragma acc loop vector(length:5) // { dg-error "argument not permitted" }
for (i = 0; i < 10; i++)
{ }
@@ -77,26 +77,26 @@ void p2 (void)
{
int i, j;
-#pragma acc parallel loop gang(5) // { dg-error "no arguments allowed to gang" "" { target c } }
- for (i = 0; i < 10; i++) // { dg-error "no arguments allowed to gang" "" { target c++ } }
+#pragma acc parallel loop gang(5) // { dg-error "argument not permitted" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "argument not permitted" "" { target c++ } }
{ }
-#pragma acc parallel loop gang(num:5) // { dg-error "no arguments allowed to gang" "" { target c } }
- for (i = 0; i < 10; i++) // { dg-error "no arguments allowed to gang" "" { target c++ } }
+#pragma acc parallel loop gang(num:5) // { dg-error "argument not permitted" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "argument not permitted" "" { target c++ } }
{ }
#pragma acc parallel loop gang
for (i = 0; i < 10; i++)
{
#pragma acc parallel loop gang // { dg-error "OpenACC construct inside of non-OpenACC region" }
- for (j = 1; j < 10; j++)
+ for (j = 1; j < 10; j++)
{ }
}
-#pragma acc parallel loop worker(5) // { dg-error "no arguments allowed to gang" "" { target c } }
- for (i = 0; i < 10; i++) // { dg-error "no arguments allowed to gang" "" { target c++ } }
+#pragma acc parallel loop worker(5) // { dg-error "argument not permitted" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "argument not permitted" "" { target c++ } }
{ }
-#pragma acc parallel loop worker(num:5) // { dg-error "no arguments allowed to gang" "" { target c } }
- for (i = 0; i < 10; i++) // { dg-error "no arguments allowed to gang" "" { target c++ } }
+#pragma acc parallel loop worker(num:5) // { dg-error "argument not permitted" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "argument not permitted" "" { target c++ } }
{ }
#pragma acc parallel loop worker
for (i = 0; i < 10; i++)
@@ -109,11 +109,11 @@ void p2 (void)
{ }
}
-#pragma acc parallel loop vector(5) // { dg-error "no arguments allowed to gang" "" { target c } }
- for (i = 0; i < 10; i++) // { dg-error "no arguments allowed to gang" "" { target c++ } }
+#pragma acc parallel loop vector(5) // { dg-error "argument not permitted" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "argument not permitted" "" { target c++ } }
{ }
-#pragma acc parallel loop vector(length:5) // { dg-error "no arguments allowed to gang" "" { target c } }
- for (i = 0; i < 10; i++) // { dg-error "no arguments allowed to gang" "" { target c++ } }
+#pragma acc parallel loop vector(length:5) // { dg-error "argument not permitted" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "argument not permitted" "" { target c++ } }
{ }
#pragma acc parallel loop vector
for (i = 0; i < 10; i++)
===================================================================
@@ -3,21 +3,21 @@
#pragma acc routine gang
int
-gang ()
+gang () /* { dg-message "declared here" 3 } */
{
return 1;
}
#pragma acc routine worker
int
-worker ()
+worker () /* { dg-message "declared here" 2 } */
{
return 1;
}
#pragma acc routine vector
int
-vector ()
+vector () /* { dg-message "declared here" } */
{
return 1;
}
@@ -49,30 +49,30 @@ main ()
red += vector ();
/* Gang routine tests. */
-#pragma acc loop gang reduction (+:red)
+#pragma acc loop gang reduction (+:red) /* { dg-message "containing loop" } */
for (int i = 0; i < 10; i++)
- red += gang (); // { dg-error "incompatible parallelism with acc routine" }
+ red += gang (); // { dg-error "routine call uses same" }
-#pragma acc loop worker reduction (+:red)
+#pragma acc loop worker reduction (+:red) /* { dg-message "containing loop" } */
for (int i = 0; i < 10; i++)
- red += gang (); // { dg-error "incompatible parallelism with acc routine" }
+ red += gang (); // { dg-error "routine call uses same" }
-#pragma acc loop vector reduction (+:red)
+#pragma acc loop vector reduction (+:red) /* { dg-message "containing loop" } */
for (int i = 0; i < 10; i++)
- red += gang (); // { dg-error "incompatible parallelism with acc routine" }
+ red += gang (); // { dg-error "routine call uses same" }
/* Worker routine tests. */
#pragma acc loop gang reduction (+:red)
for (int i = 0; i < 10; i++)
red += worker ();
-#pragma acc loop worker reduction (+:red)
+#pragma acc loop worker reduction (+:red) /* { dg-message "containing loop" } */
for (int i = 0; i < 10; i++)
- red += worker (); // { dg-error "incompatible parallelism with acc routine" }
+ red += worker (); // { dg-error "routine call uses same" }
-#pragma acc loop vector reduction (+:red)
+#pragma acc loop vector reduction (+:red) /* { dg-message "containing loop" } */
for (int i = 0; i < 10; i++)
- red += worker (); // { dg-error "incompatible parallelism with acc routine" }
+ red += worker (); // { dg-error "routine call uses same" }
/* Vector routine tests. */
#pragma acc loop gang reduction (+:red)
@@ -83,9 +83,9 @@ main ()
for (int i = 0; i < 10; i++)
red += vector ();
-#pragma acc loop vector reduction (+:red)
+#pragma acc loop vector reduction (+:red) /* { dg-message "containing loop" } */
for (int i = 0; i < 10; i++)
- red += vector (); // { dg-error "incompatible parallelism with acc routine" }
+ red += vector (); // { dg-error "routine call uses same" }
/* Seq routine tests. */
#pragma acc loop gang reduction (+:red)
===================================================================
@@ -20,7 +20,7 @@ main ()
#pragma acc loop gang(static:*)
for (i = 0; i < 10; i++)
{ }
-#pragma acc loop gang // { dg-error "gang, worker and vector may occur only once in a loop nest" }
+#pragma acc loop gang // { dg-message "containing loop" }
for (i = 0; i < 10; i++)
{
#pragma acc loop vector
@@ -29,31 +29,31 @@ main ()
#pragma acc loop worker
for (j = 1; j < 10; j++)
{ }
-#pragma acc loop gang
+#pragma acc loop gang // { dg-error "inner loop uses same" }
for (j = 1; j < 10; j++)
{ }
}
-#pragma acc loop seq gang // { dg-error "incompatible use of clause" }
+#pragma acc loop seq gang // { dg-error "'seq' overrides" }
for (i = 0; i < 10; i++)
{ }
#pragma acc loop worker
for (i = 0; i < 10; i++)
{ }
-#pragma acc loop worker // { dg-error "gang, worker and vector may occur only once in a loop nest" }
+#pragma acc loop worker // { dg-message "containing loop" 2 }
for (i = 0; i < 10; i++)
{
#pragma acc loop vector
for (j = 1; j < 10; j++)
{ }
-#pragma acc loop worker
+#pragma acc loop worker // { dg-error "inner loop uses same" }
for (j = 1; j < 10; j++)
{ }
-#pragma acc loop gang
+#pragma acc loop gang // { dg-error "incorrectly nested" }
for (j = 1; j < 10; j++)
{ }
}
-#pragma acc loop seq worker // { dg-error "incompatible use of clause" }
+#pragma acc loop seq worker // { dg-error "'seq' overrides" }
for (i = 0; i < 10; i++)
{ }
#pragma acc loop gang worker
@@ -65,20 +65,20 @@ main ()
{ }
for (i = 0; i < 10; i++)
{ }
-#pragma acc loop vector // { dg-error "gang, worker and vector may occur only once in a loop nest" }
+#pragma acc loop vector // { dg-message "containing loop" 3 }
for (i = 0; i < 10; i++)
{
-#pragma acc loop vector
+#pragma acc loop vector // { dg-error "inner loop uses same" }
for (j = 1; j < 10; j++)
{ }
-#pragma acc loop worker
+#pragma acc loop worker // { dg-error "incorrectly nested" }
for (j = 1; j < 10; j++)
{ }
-#pragma acc loop gang
+#pragma acc loop gang // { dg-error "incorrectly nested" }
for (j = 1; j < 10; j++)
{ }
}
-#pragma acc loop seq vector // { dg-error "incompatible use of clause" }
+#pragma acc loop seq vector // { dg-error "'seq' overrides" }
for (i = 0; i < 10; i++)
{ }
#pragma acc loop gang vector
@@ -91,16 +91,16 @@ main ()
#pragma acc loop auto
for (i = 0; i < 10; i++)
{ }
-#pragma acc loop seq auto // { dg-error "incompatible use of clause" }
+#pragma acc loop seq auto // { dg-error "'seq' overrides" }
for (i = 0; i < 10; i++)
{ }
-#pragma acc loop gang auto // { dg-error "incompatible use of clause" }
+#pragma acc loop gang auto // { dg-error "'auto' conflicts" }
for (i = 0; i < 10; i++)
{ }
-#pragma acc loop worker auto // { dg-error "incompatible use of clause" }
+#pragma acc loop worker auto // { dg-error "'auto' conflicts" }
for (i = 0; i < 10; i++)
{ }
-#pragma acc loop vector auto // { dg-error "incompatible use of clause" }
+#pragma acc loop vector auto // { dg-error "'auto' conflicts" }
for (i = 0; i < 10; i++)
{ }
@@ -119,16 +119,16 @@ main ()
for (i = 0; i < 10; i++)
{ }
-#pragma acc parallel loop seq gang // { dg-error "incompatible use of clause" "" { target c } }
- for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" "" { target c++ } }
+#pragma acc parallel loop seq gang // { dg-error "'seq' overrides" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "'seq' overrides" "" { target c++ } }
{ }
#pragma acc parallel loop worker
for (i = 0; i < 10; i++)
{ }
-#pragma acc parallel loop seq worker // { dg-error "incompatible use of clause" "" { target c } }
- for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" "" { target c++ } }
+#pragma acc parallel loop seq worker // { dg-error "'seq' overrides" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "'seq' overrides" "" { target c++ } }
{ }
#pragma acc parallel loop gang worker
for (i = 0; i < 10; i++)
@@ -138,8 +138,8 @@ main ()
for (i = 0; i < 10; i++)
{ }
-#pragma acc parallel loop seq vector // { dg-error "incompatible use of clause" "" { target c } }
- for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" "" { target c++ } }
+#pragma acc parallel loop seq vector // { dg-error "'seq' overrides" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "'seq' overrides" "" { target c++ } }
{ }
#pragma acc parallel loop gang vector
for (i = 0; i < 10; i++)
@@ -151,19 +151,20 @@ main ()
#pragma acc parallel loop auto
for (i = 0; i < 10; i++)
{ }
-#pragma acc parallel loop seq auto // { dg-error "incompatible use of clause" "" { target c } }
- for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" "" { target c++ } }
+#pragma acc parallel loop seq auto // { dg-error "'seq' overrides" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "'seq' overrides" "" { target c++ } }
{ }
-#pragma acc parallel loop gang auto // { dg-error "incompatible use of clause" "" { target c } }
- for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" "" { target c++ } }
+#pragma acc parallel loop gang auto // { dg-error "'auto' conflicts" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "'auto' conflicts" "" { target c++ } }
{ }
-#pragma acc parallel loop worker auto // { dg-error "incompatible use of clause" "" { target c } }
- for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" "" { target c++ } }
+#pragma acc parallel loop worker auto // { dg-error "'auto' conflicts" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "'auto' conflicts" "" { target c++ } }
{ }
-#pragma acc parallel loop vector auto // { dg-error "incompatible use of clause" "" { target c } }
- for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" "" { target c++ } }
+#pragma acc parallel loop vector auto // { dg-error "'auto' conflicts" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "'auto' conflicts" "" { target c++ } }
{ }
+
return 0;
}
===================================================================
@@ -25,7 +25,7 @@ main ()
#pragma acc loop gang(static:*)
for (i = 0; i < 10; i++)
{ }
-#pragma acc loop gang // { dg-error "gang, worker and vector may occur only once in a loop nest" }
+#pragma acc loop gang
for (i = 0; i < 10; i++)
{
#pragma acc loop vector
@@ -34,11 +34,11 @@ main ()
#pragma acc loop worker
for (j = 0; j < 10; j++)
{ }
-#pragma acc loop gang
+#pragma acc loop gang // { dg-error "inner loop uses same" }
for (j = 0; j < 10; j++)
{ }
}
-#pragma acc loop seq gang // { dg-error "incompatible use of clause" }
+#pragma acc loop seq gang // { dg-error "'seq' overrides" }
for (i = 0; i < 10; i++)
{ }
@@ -51,20 +51,20 @@ main ()
#pragma acc loop worker(num:5)
for (i = 0; i < 10; i++)
{ }
-#pragma acc loop worker // { dg-error "gang, worker and vector may occur only once in a loop nest" }
+#pragma acc loop worker
for (i = 0; i < 10; i++)
{
#pragma acc loop vector
for (j = 0; j < 10; j++)
{ }
-#pragma acc loop worker
+#pragma acc loop worker // { dg-error "inner loop uses same" }
for (j = 0; j < 10; j++)
{ }
#pragma acc loop gang
for (j = 0; j < 10; j++)
{ }
}
-#pragma acc loop seq worker // { dg-error "incompatible use of clause" }
+#pragma acc loop seq worker // { dg-error "'seq' overrides" }
for (i = 0; i < 10; i++)
{ }
#pragma acc loop gang worker
@@ -80,10 +80,10 @@ main ()
#pragma acc loop vector(length:5)
for (i = 0; i < 10; i++)
{ }
-#pragma acc loop vector // { dg-error "gang, worker and vector may occur only once in a loop nest" }
+#pragma acc loop vector
for (i = 0; i < 10; i++)
{
-#pragma acc loop vector
+#pragma acc loop vector // { dg-error "inner loop uses same" }
for (j = 1; j < 10; j++)
{ }
#pragma acc loop worker
@@ -93,7 +93,7 @@ main ()
for (j = 1; j < 10; j++)
{ }
}
-#pragma acc loop seq vector // { dg-error "incompatible use of clause" }
+#pragma acc loop seq vector // { dg-error "'seq' overrides" }
for (i = 0; i < 10; i++)
{ }
#pragma acc loop gang vector
@@ -106,19 +106,18 @@ main ()
#pragma acc loop auto
for (i = 0; i < 10; i++)
{ }
-#pragma acc loop seq auto // { dg-error "incompatible use of clause" }
+#pragma acc loop seq auto // { dg-error "'seq' overrides" }
for (i = 0; i < 10; i++)
{ }
-#pragma acc loop gang auto // { dg-error "incompatible use of clause" }
+#pragma acc loop gang auto // { dg-error "'auto' conflicts" }
for (i = 0; i < 10; i++)
{ }
-#pragma acc loop worker auto // { dg-error "incompatible use of clause" }
+#pragma acc loop worker auto // { dg-error "'auto' conflicts" }
for (i = 0; i < 10; i++)
{ }
-#pragma acc loop vector auto // { dg-error "incompatible use of clause" }
+#pragma acc loop vector auto // { dg-error "'auto' conflicts" }
for (i = 0; i < 10; i++)
{ }
-
}
@@ -150,8 +149,8 @@ main ()
#pragma acc kernels loop worker(num:5)
for (i = 0; i < 10; i++)
{ }
-#pragma acc kernels loop seq worker // { dg-error "incompatible use of clause" "" { target c } }
- for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" "" { target c++ } }
+#pragma acc kernels loop seq worker // { dg-error "'seq' overrides" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "'seq' overrides" "" { target c++ } }
{ }
#pragma acc kernels loop gang worker
for (i = 0; i < 10; i++)
@@ -166,8 +165,8 @@ main ()
#pragma acc kernels loop vector(length:5)
for (i = 0; i < 10; i++)
{ }
-#pragma acc kernels loop seq vector // { dg-error "incompatible use of clause" "" { target c } }
- for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" "" { target c++ } }
+#pragma acc kernels loop seq vector // { dg-error "'seq' overrides" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "'seq' overrides" "" { target c++ } }
{ }
#pragma acc kernels loop gang vector
for (i = 0; i < 10; i++)
@@ -179,17 +178,17 @@ main ()
#pragma acc kernels loop auto
for (i = 0; i < 10; i++)
{ }
-#pragma acc kernels loop seq auto // { dg-error "incompatible use of clause" "" { target c } }
- for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" "" { target c++ } }
+#pragma acc kernels loop seq auto // { dg-error "'seq' overrides" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "'seq' overrides" "" { target c++ } }
{ }
-#pragma acc kernels loop gang auto // { dg-error "incompatible use of clause" "" { target c } }
- for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" "" { target c++ } }
+#pragma acc kernels loop gang auto // { dg-error "'auto' conflicts" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "'auto' conflicts" "" { target c++ } }
{ }
-#pragma acc kernels loop worker auto // { dg-error "incompatible use of clause" "" { target c } }
- for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" "" { target c++ } }
+#pragma acc kernels loop worker auto // { dg-error "'auto' conflicts" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "'auto' conflicts" "" { target c++ } }
{ }
-#pragma acc kernels loop vector auto // { dg-error "incompatible use of clause" "" { target c } }
- for (i = 0; i < 10; i++) // { dg-error "incompatible use of clause" "" { target c++ } }
+#pragma acc kernels loop vector auto // { dg-error "'auto' conflicts" "" { target c } }
+ for (i = 0; i < 10; i++) // { dg-error "'auto' conflicts" "" { target c++ } }
{ }
return 0;
===================================================================
@@ -46,10 +46,10 @@ program test
!$acc loop vector
DO i = 1,10
ENDDO
- !$acc loop vector(5) ! { dg-error "no arguments allowed to gang" }
+ !$acc loop vector(5) ! { dg-error "argument not permitted" }
DO i = 1,10
ENDDO
- !$acc loop vector(length:5) ! { dg-error "no arguments allowed to gang" }
+ !$acc loop vector(length:5) ! { dg-error "argument not permitted" }
DO i = 1,10
ENDDO
!$acc loop vector
@@ -70,10 +70,10 @@ program test
!$acc parallel loop vector
DO i = 1,10
ENDDO
- !$acc parallel loop vector(5) ! { dg-error "no arguments allowed to gang" }
+ !$acc parallel loop vector(5) ! { dg-error "argument not permitted" }
DO i = 1,10
ENDDO
- !$acc parallel loop vector(length:5) ! { dg-error "no arguments allowed to gang" }
+ !$acc parallel loop vector(length:5) ! { dg-error "argument not permitted" }
DO i = 1,10
ENDDO
end
===================================================================
@@ -49,19 +49,19 @@ program main
call gang (a)
end do
- !$acc loop gang
+ !$acc loop gang ! { dg-message "containing loop" }
do i = 1, N
- call gang (a) ! { dg-error "incompatible parallelism with acc routine" }
+ call gang (a) ! { dg-error "routine call uses same" }
end do
- !$acc loop worker
+ !$acc loop worker ! { dg-message "containing loop" }
do i = 1, N
- call gang (a) ! { dg-error "incompatible parallelism with acc routine" }
+ call gang (a) ! { dg-error "routine call uses same" }
end do
- !$acc loop vector
+ !$acc loop vector ! { dg-message "containing loop" }
do i = 1, N
- call gang (a) ! { dg-error "incompatible parallelism with acc routine" }
+ call gang (a) ! { dg-error "routine call uses same" }
end do
!$acc end parallel
@@ -80,14 +80,14 @@ program main
call worker (a)
end do
- !$acc loop worker
+ !$acc loop worker ! { dg-message "containing loop" }
do i = 1, N
- call worker (a) ! { dg-error "incompatible parallelism with acc routine" }
+ call worker (a) ! { dg-error "routine call uses same" }
end do
- !$acc loop vector
+ !$acc loop vector ! { dg-message "containing loop" }
do i = 1, N
- call worker (a) ! { dg-error "incompatible parallelism with acc routine" }
+ call worker (a) ! { dg-error "routine call uses same" }
end do
!$acc end parallel
@@ -111,14 +111,14 @@ program main
call vector (a)
end do
- !$acc loop vector
+ !$acc loop vector ! { dg-message "containing loop" }
do i = 1, N
- call vector (a) ! { dg-error "incompatible parallelism with acc routine" }
+ call vector (a) ! { dg-error "routine call uses same" }
end do
!$acc end parallel
contains
- subroutine gang (a)
+ subroutine gang (a) ! { dg-message "declared here" 3 }
!$acc routine gang
integer, intent (inout) :: a(N)
integer :: i
@@ -128,7 +128,7 @@ contains
end do
end subroutine gang
- subroutine worker (a)
+ subroutine worker (a) ! { dg-message "declared here" 2 }
!$acc routine worker
integer, intent (inout) :: a(N)
integer :: i
@@ -138,7 +138,7 @@ contains
end do
end subroutine worker
- subroutine vector (a)
+ subroutine vector (a) ! { dg-message "declared here" }
!$acc routine vector
integer, intent (inout) :: a(N)
integer :: i
===================================================================
@@ -40,7 +40,7 @@ subroutine worker (a)
a(i) = a(i) - a(i)
end do
- !$acc loop gang ! { dg-error "invalid parallelism inside acc routine" }
+ !$acc loop gang ! { dg-error "disallowed by containing routine" }
do i = 1, N
a(i) = a(i) - a(i)
end do
@@ -66,12 +66,12 @@ subroutine vector (a)
a(i) = a(i) - a(i)
end do
- !$acc loop gang ! { dg-error "invalid parallelism inside acc routine" }
+ !$acc loop gang ! { dg-error "disallowed by containing routine" }
do i = 1, N
a(i) = a(i) - a(i)
end do
- !$acc loop worker ! { dg-error "invalid parallelism inside acc routine" }
+ !$acc loop worker ! { dg-error "disallowed by containing routine" }
do i = 1, N
a(i) = a(i) - a(i)
end do
@@ -92,17 +92,17 @@ subroutine seq (a)
a(i) = a(i) - a(i)
end do
- !$acc loop gang ! { dg-error "invalid parallelism inside acc routine" }
+ !$acc loop gang ! { dg-error "disallowed by containing routine" }
do i = 1, N
a(i) = a(i) - a(i)
end do
- !$acc loop worker ! { dg-error "invalid parallelism inside acc routine" }
+ !$acc loop worker ! { dg-error "disallowed by containing routine" }
do i = 1, N
a(i) = a(i) - a(i)
end do
- !$acc loop vector ! { dg-error "invalid parallelism inside acc routine" }
+ !$acc loop vector ! { dg-error "disallowed by containing routine" }
do i = 1, N
a(i) = a(i) - a(i)
end do