On 10/08/2014 02:38 AM, Thomas Schwinge wrote:
> On Fri, 3 Oct 2014 09:22:52 -0700, Cesar Philippidis <ce...@codesourcery.com>
> wrote:
>> There is a reduction bug [...]
>
> Thanks for looking into this!
>
>> This is a problem because initialize_reduction_data originally expected
>> a GIMPLE_BIND at the very beginning of a parallel block. [...]
>
> Does your patch also fix the issue that Jim saw in his C++ work?
Yes it does.
>> This patch also includes a runtime test case. I won't apply it to
>> gomp-4_0-branch just yet. But I wanted to demonstrate a test case
>> nonetheless.
>
> You can add it as a compile test, and I'll toggle to a run test as part
> of the merge into our internal development branch.
OK. I've added a compile test to gcc.dg/goaccc/. Note that this test
depends on -std=c99, so I couldn't put it in the c-c++-common directory.
>> Also, note that part of this patch also changes a comment.
>> I found some typos in the original comment, so I took the opportunity to
>> fix them, I hope.
>
> Sure, and it's also always fine to separately apply such patches under
> the obvious rule.
>
>> Is this OK for gomp-4_0-branch?
>
> Yes, with the following addressed:
>
>> --- a/gcc/omp-low.c
>> +++ b/gcc/omp-low.c
>> @@ -10140,11 +10140,20 @@ process_reduction_data (gimple_seq *body,
>> gimple_seq *in_stmt_seqp,
>
> [...]/source-gcc/gcc/omp-low.c: In function 'void
> _ZL22process_reduction_dataPP21gimple_statement_baseS1_S1_P11omp_context.isra.167.constprop.180(gimple_statement_base**,
> gimple_statement_base**, gimple_statement_base**, gimple)':
> [...]/source-gcc/gcc/omp-low.c:10172:14: warning: 'inner' may be used
> uninitialized in this function [-Wmaybe-uninitialized]
> gimple_seq inner;
> ^
>
>> --- /dev/null
>> +++ b/libgomp/testsuite/libgomp.oacc-c/collapse-4.c
>> @@ -0,0 +1,34 @@
>> +/* { dg-do run } */
>> +/* { dg-options "-O2 -std=c99" } */
>> +
>> +#include <string.h>
>> +#include <stdlib.h>
>> +#include <stdio.h>
>> +
>> +int
>> +main (void)
>> +{
>> + int i2, l = 0, r = 0;
>> + int a[3][3][3];
>> + int b[3][3];
>
> A lot of these variables are unused.
>
>> +printf("&a %p\n", &a[0][0][0]);
>> +printf("&i2 %p\n", &i2);
>> +printf("&l %p\n", &l);
>> +printf("&r %p\n", &r);
>
> Please remove <stdio.h> include and the printfs, or at least put those
> into DEBUG conditionals.
I removed those thanks. This updated patch has been committed.
Thanks,
Cesar
2014-10-08 Cesar Philippidis <ce...@codesourcery.com>
gcc/
* omp-low.c (lower_reduction_clauses): Clarify comment.
(process_reduction_data): Scan for nonempty bind statements at
the beginning of parallel blocks.
gcc/testsuite/
* gcc.dg/goacc/collapse.c: New test.
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 3bb6a24..5c452c6 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -4434,10 +4434,10 @@ lower_reduction_clauses (tree clauses, gimple_seq *stmt_seqp, omp_context *ctx)
else
{
/* The atomic add at the end of the sum creates unnecessary
- write contention on accelerators. To work around that,
- create an array or vector_length and assign an element to
- each thread. Later, in lower_omp_for (for openacc), the
- values of array will be combined. */
+ write contention on accelerators. To work around this,
+ create an array to store the partial reductions. Later, in
+ lower_omp_for (for openacc), the values of array will be
+ combined. */
tree t = NULL_TREE, array, nthreads;
tree type = get_base_type (var);
@@ -10140,11 +10140,20 @@ process_reduction_data (gimple_seq *body, gimple_seq *in_stmt_seqp,
gimple stmt;
/* A collapse clause may have inserted a new bind block. */
- stmt = gimple_seq_first (*body);
- if (stmt && gimple_code (stmt) == GIMPLE_BIND)
+ gsi = gsi_start (*body);
+ while (!gsi_end_p (gsi))
{
- inner = gimple_bind_body (gimple_seq_first (*body));
- body = &inner;
+ stmt = gsi_stmt (gsi);
+ if (gimple_code (stmt) == GIMPLE_BIND)
+ {
+ inner = gimple_bind_body (stmt);
+ body = &inner;
+ gsi = gsi_start (*body);
+ }
+ else if (gimple_code (stmt) == GIMPLE_OMP_FOR)
+ break;
+ else
+ gsi_next (&gsi);
}
for (gsi = gsi_start (*body); !gsi_end_p (gsi); gsi_next (&gsi))
diff --git a/gcc/testsuite/gcc.dg/goacc/collapse.c b/gcc/testsuite/gcc.dg/goacc/collapse.c
new file mode 100644
index 0000000..1ec20a4
--- /dev/null
+++ b/gcc/testsuite/gcc.dg/goacc/collapse.c
@@ -0,0 +1,24 @@
+/* { dg-do compile } */
+/* { dg-options "-O2 -std=c99" } */
+
+#include <string.h>
+
+int
+main (void)
+{
+ int l = 0;
+ int b[3][3];
+
+ memset (b, '\0', sizeof (b));
+
+#pragma acc parallel copy(b[0:3][0:3]) copy(l)
+ {
+#pragma acc loop collapse(2) reduction(+:l)
+ for (int i = 0; i < 2; i++)
+ for (int j = 0; j < 2; j++)
+ if (b[i][j] != 16)
+ l += 1;
+ }
+
+ return 0;
+}
