Hi Tom,
this patch provides a 'bool independent' field in struct loop, which
will be switched on by an "independent" clause in a #pragma acc loop directive.
I assume you'll be wiring it to the kernels parloops pass in a followup patch.
Note: there are already a few other similar fields in struct loop, namely
'safelen' and 'can_be_parallel', used by OMP simd safelen and GRAPHITE
respectively.
The intention and/or setting of these fields are all a bit different, so I've
decided to add a new bool for OpenACC.
Tested and committed to gomp-4_0-branch.
Chung-Lin
2015-07-14 Chung-Lin Tang <clt...@codesourcery.com>
* cfgloop.h (struct loop): Add 'bool marked_independent' field.
* gimplify.c (gimplify_scan_omp_clauses): Keep OMP_CLAUSE_INDEPENDENT.
* omp-low.c (struct omp_region): Add 'int kind' and
'bool independent' fields.
(expand_omp_for): Set 'marked_independent' field for loop
corresponding to region.
(find_omp_for_region_data): New function.
(find_omp_target_region_data): Set kind field.
(build_omp_regions_1): Call find_omp_for_region_data() for
GIMPLE_OMP_FOR statements.
Index: cfgloop.h
===================================================================
--- cfgloop.h (revision 225758)
+++ cfgloop.h (working copy)
@@ -194,6 +194,10 @@ struct GTY ((chain_next ("%h.next"))) loop {
/* True if the loop is part of an oacc kernels region. */
bool in_oacc_kernels_region;
+ /* True if loop is tagged as having independent iterations by user,
+ e.g. the OpenACC independent clause. */
+ bool marked_independent;
+
/* For SIMD loops, this is a unique identifier of the loop, referenced
by IFN_GOMP_SIMD_VF, IFN_GOMP_SIMD_LANE and IFN_GOMP_SIMD_LAST_LANE
builtins. */
Index: gimplify.c
===================================================================
--- gimplify.c (revision 225758)
+++ gimplify.c (working copy)
@@ -6602,7 +6602,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_se
break;
case OMP_CLAUSE_DEVICE_RESIDENT:
- case OMP_CLAUSE_INDEPENDENT:
remove = true;
break;
@@ -6612,6 +6611,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_se
case OMP_CLAUSE_COLLAPSE:
case OMP_CLAUSE_AUTO:
case OMP_CLAUSE_SEQ:
+ case OMP_CLAUSE_INDEPENDENT:
case OMP_CLAUSE_MERGEABLE:
case OMP_CLAUSE_PROC_BIND:
case OMP_CLAUSE_SAFELEN:
Index: omp-low.c
===================================================================
--- omp-low.c (revision 225758)
+++ omp-low.c (working copy)
@@ -136,8 +136,16 @@ struct omp_region
/* True if this is nested inside an OpenACC kernels construct. */
bool inside_kernels_p;
+ /* 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;
+
+ tree broadcast_array;
};
/* Context structure. Used to store information about each parallel
@@ -8273,8 +8281,15 @@ expand_omp_for (struct omp_region *region, gimple
loops_state_set (LOOPS_NEED_FIXUP);
if (region->inside_kernels_p)
- expand_omp_for_generic (region, &fd, BUILT_IN_NONE, BUILT_IN_NONE,
- inner_stmt);
+ {
+ expand_omp_for_generic (region, &fd, BUILT_IN_NONE, BUILT_IN_NONE,
+ inner_stmt);
+ if (region->independent && region->cont->loop_father)
+ {
+ struct loop *loop = region->cont->loop_father;
+ loop->marked_independent = true;
+ }
+ }
else if (gimple_omp_for_kind (fd.for_stmt) & GF_OMP_FOR_SIMD)
expand_omp_simd (region, &fd);
else if (gimple_omp_for_kind (fd.for_stmt) == GF_OMP_FOR_KIND_CILKFOR)
@@ -9943,6 +9958,34 @@ find_omp_for_region_gwv (gimple stmt)
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)
+ {
+ struct omp_region *target_region = region->outer;
+ while (target_region
+ && target_region->type != GIMPLE_OMP_TARGET)
+ target_region = target_region->outer;
+ if (!target_region)
+ return;
+
+ 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))
+ region->independent = true;
+ }
+}
+
/* Fill in additional data for a region REGION associated with an
OMP_TARGET STMT. */
@@ -9960,6 +10003,7 @@ find_omp_target_region_data (struct omp_region *re
region->gwv_this |= OACC_LOOP_MASK (OACC_worker);
if (find_omp_clause (clauses, OMP_CLAUSE_VECTOR_LENGTH))
region->gwv_this |= OACC_LOOP_MASK (OACC_vector);
+ region->kind = gimple_omp_target_kind (stmt);
}
/* Helper for build_omp_regions. Scan the dominator tree starting at
@@ -10046,7 +10090,7 @@ build_omp_regions_1 (basic_block bb, struct omp_re
}
}
else if (code == GIMPLE_OMP_FOR)
- region->gwv_this = find_omp_for_region_gwv (stmt);
+ find_omp_for_region_data (region, as_a <gomp_for *> (stmt));
/* ..., this directive becomes the parent for a new region. */
if (region)
parent = region;
