On 07/11/2014 02:27 AM, Thomas Schwinge wrote:
> Hi Cesar!
>
> On Thu, 10 Jul 2014 11:47:42 -0700, Cesar Philippidis
> <ce...@codesourcery.com> wrote:
>> These patch enables the collapse clause with a value greater than one.
>
> Thanks!
>
>> Is this patch OK for gomp-4_0-branch?
>
> OK with the following addressed:
>
>> --- a/gcc/c/c-parser.c
>> +++ b/gcc/c/c-parser.c
>> @@ -11260,6 +11260,10 @@ c_parser_oacc_all_clauses (c_parser *parser,
>> omp_clause_mask mask,
>>
>> switch (c_kind)
>> {
>> + case PRAGMA_OMP_CLAUSE_COLLAPSE:
>> + clauses = c_parser_omp_clause_collapse (parser, clauses);
>> + c_name = "collapse";
>> + break;
>
> Annotate c_parser_omp_clause_collapse that it is now used for OpenACC,
> too.
>
>
>> #define OACC_LOOP_CLAUSE_MASK
>> \
>> - (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_REDUCTION)
>> + ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COLLAPSE) \
>> + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_REDUCTION))
>
> Given the syntax generally used, add a space between the two closing
> parens.
>
>
> Even though we're re-using OpenMP's tested implementation, please add
> some basic front end test coverage for OpenACC: reject invalid, accept
> valid syntax.
Thanks for the feedback. I've committed the attach patch with those changes.
Cesar
2014-07-14 Cesar Philippidis <ce...@codesourcery.com>
gcc/c/
* c-parser.c (c_parser_oacc_all_clauses): Handle
PRAGMA_OMP_CLAUSE_COLLAPSE.
(OACC_LOOP_CLAUSE_MASK): Add PRAGMA_OMP_CLAUSE_COLLAPSE.
(c_parser_omp_for_loop): Remove asserts.
gcc/
* omp-low.c (extract_omp_for_data): Likewise.
(expand_omp_for_static_nochunk): Likewise.
gcc/testsuite/
* c-c++-common/goacc/collapse-1.c: New test.
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index d638930..4e2ea33 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -10057,7 +10057,7 @@ c_parser_oacc_data_clause_deviceptr (c_parser *parser, tree list)
return list;
}
-/* OpenMP 3.0:
+/* OpenACC 2.0, OpenMP 3.0:
collapse ( constant-expression ) */
static tree
@@ -11260,6 +11260,10 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
switch (c_kind)
{
+ case PRAGMA_OMP_CLAUSE_COLLAPSE:
+ clauses = c_parser_omp_clause_collapse (parser, clauses);
+ c_name = "collapse";
+ break;
case PRAGMA_OMP_CLAUSE_COPY:
clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
c_name = "copy";
@@ -11702,7 +11706,8 @@ c_parser_oacc_kernels (location_t loc, c_parser *parser, char *p_name)
*/
#define OACC_LOOP_CLAUSE_MASK \
- (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_REDUCTION)
+ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COLLAPSE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_REDUCTION) )
static tree
c_parser_oacc_loop (location_t loc, c_parser *parser, char *p_name)
@@ -12306,10 +12311,7 @@ c_parser_omp_for_loop (location_t loc, c_parser *parser, enum tree_code code,
for (cl = clauses; cl; cl = OMP_CLAUSE_CHAIN (cl))
if (OMP_CLAUSE_CODE (cl) == OMP_CLAUSE_COLLAPSE)
- {
- gcc_assert (code != OACC_LOOP);
collapse = tree_to_shwi (OMP_CLAUSE_COLLAPSE_EXPR (cl));
- }
gcc_assert (collapse >= 1);
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index b08bc6b..6345e14 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -381,9 +381,6 @@ extract_omp_for_data (gimple for_stmt, struct omp_for_data *fd,
case OMP_CLAUSE_COLLAPSE:
if (fd->collapse > 1)
{
- if (is_gimple_omp_oacc_specifically (for_stmt))
- sorry ("collapse (>1) clause not supported yet");
-
collapse_iter = &OMP_CLAUSE_COLLAPSE_ITERVAR (t);
collapse_count = &OMP_CLAUSE_COLLAPSE_COUNT (t);
}
@@ -6684,9 +6681,6 @@ expand_omp_for_static_nochunk (struct omp_region *region,
if (fd->collapse > 1)
{
- gcc_assert (gimple_omp_for_kind (fd->for_stmt)
- != GF_OMP_FOR_KIND_OACC_LOOP);
-
int first_zero_iter = -1;
basic_block l2_dom_bb = NULL;
@@ -6697,12 +6691,7 @@ expand_omp_for_static_nochunk (struct omp_region *region,
t = NULL_TREE;
}
else if (gimple_omp_for_combined_into_p (fd->for_stmt))
- {
- gcc_assert (gimple_omp_for_kind (fd->for_stmt)
- != GF_OMP_FOR_KIND_OACC_LOOP);
-
t = integer_one_node;
- }
else
t = fold_binary (fd->loop.cond_code, boolean_type_node,
fold_convert (type, fd->loop.n1),
@@ -6736,9 +6725,6 @@ expand_omp_for_static_nochunk (struct omp_region *region,
ep->probability = REG_BR_PROB_BASE / 2000 - 1;
if (gimple_in_ssa_p (cfun))
{
- gcc_assert (gimple_omp_for_kind (fd->for_stmt)
- != GF_OMP_FOR_KIND_OACC_LOOP);
-
int dest_idx = find_edge (entry_bb, fin_bb)->dest_idx;
for (gsi = gsi_start_phis (fin_bb);
!gsi_end_p (gsi); gsi_next (&gsi))
@@ -6865,9 +6851,6 @@ expand_omp_for_static_nochunk (struct omp_region *region,
if (gimple_omp_for_combined_p (fd->for_stmt))
{
- gcc_assert (gimple_omp_for_kind (fd->for_stmt)
- != GF_OMP_FOR_KIND_OACC_LOOP);
-
tree clauses = gimple_code (inner_stmt) == GIMPLE_OMP_PARALLEL
? gimple_omp_parallel_clauses (inner_stmt)
: gimple_omp_for_clauses (inner_stmt);
@@ -6914,12 +6897,7 @@ expand_omp_for_static_nochunk (struct omp_region *region,
gsi_insert_after (&gsi, stmt, GSI_CONTINUE_LINKING);
}
if (fd->collapse > 1)
- {
- gcc_assert (gimple_omp_for_kind (fd->for_stmt)
- != GF_OMP_FOR_KIND_OACC_LOOP);
-
expand_omp_for_init_vars (fd, &gsi, counts, inner_stmt, startvar);
- }
if (!broken_loop)
{
@@ -6954,12 +6932,7 @@ expand_omp_for_static_nochunk (struct omp_region *region,
gsi_remove (&gsi, true);
if (fd->collapse > 1 && !gimple_omp_for_combined_p (fd->for_stmt))
- {
- gcc_assert (gimple_omp_for_kind (fd->for_stmt)
- != GF_OMP_FOR_KIND_OACC_LOOP);
-
collapse_bb = extract_omp_for_update_vars (fd, cont_bb, body_bb);
- }
}
/* Replace the GIMPLE_OMP_RETURN with a barrier, or nothing. */
@@ -6996,9 +6969,6 @@ expand_omp_for_static_nochunk (struct omp_region *region,
}
else if (fd->collapse > 1)
{
- gcc_assert (gimple_omp_for_kind (fd->for_stmt)
- != GF_OMP_FOR_KIND_OACC_LOOP);
-
remove_edge (ep);
ep = make_edge (cont_bb, collapse_bb, EDGE_TRUE_VALUE);
}
diff --git a/gcc/testsuite/c-c++-common/goacc/collapse-1.c b/gcc/testsuite/c-c++-common/goacc/collapse-1.c
new file mode 100644
index 0000000..b2aebb7
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/collapse-1.c
@@ -0,0 +1,97 @@
+/* { dg-do compile } */
+
+int i, j, k;
+extern int foo (void);
+
+void
+f1 (void)
+{
+ #pragma acc parallel
+ #pragma acc loop collapse (2)
+ for (i = 0; i < 5; i++)
+ ; /* { dg-error "not enough perfectly nested" } */
+ {
+ for (j = 0; j < 5; j++)
+ ;
+ }
+}
+
+void
+f2 (void)
+{
+ #pragma acc parallel
+ #pragma acc loop collapse (2)
+ for (i = 0; i < 5; i++)
+ {
+ {
+ {
+ for (j = 0; j < 5; j++)
+ {
+ }
+ }
+ }
+ }
+}
+
+void
+f3 (void)
+{
+ #pragma acc parallel
+ #pragma acc loop collapse (2)
+ for (i = 0; i < 5; i++)
+ {
+ int k = foo (); /* { dg-error "not enough perfectly nested" } */
+ {
+ {
+ for (j = 0; j < 5; j++)
+ {
+ }
+ }
+ }
+ }
+}
+
+void
+f4 (void)
+{
+ #pragma acc parallel
+ #pragma acc loop collapse (2)
+ for (i = 0; i < 5; i++)
+ {
+ {
+ for (j = 0; j < 5; j++)
+ ;
+ foo (); /* { dg-error "collapsed loops not perfectly nested before" } */
+ }
+ }
+}
+
+void
+f5 (void)
+{
+ #pragma acc parallel
+ #pragma acc loop collapse (2)
+ for (i = 0; i < 5; i++)
+ {
+ {
+ for (j = 0; j < 5; j++)
+ ;
+ }
+ foo (); /* { dg-error "collapsed loops not perfectly nested before" } */
+ }
+}
+
+void
+f6 (void)
+{
+ #pragma acc parallel
+ #pragma acc loop collapse (2)
+ for (i = 0; i < 5; i++)
+ {
+ {
+ for (j = 0; j < 5; j++)
+ ;
+ }
+ }
+ foo ();
+}
