Hi Thomas,
I have now fixed it temporarily differently – without actually
testing that feature. – See attachment.
For the proper fix, it would help to get the semantic right in
OpenACC itself (→ ongoing discussion). I think the patch
for PR94120 is probably not completely right – and should be
fixed for GCC 10, but I think one should wait for the next two
OpenACC meetings to get converged. One presumably needs to handle
'device_resident' and 'link' differently from 'copy* etc.
I added a note to PR 94140.
Regards,
Tobias
On 4/10/20 3:20 PM, Thomas Schwinge wrote:
Hi Tobias!
I do see the new 'libgomp.oacc-c++/declare-pr94120.C' (see below, for
reference) FAIL for '-foffload=nvptx-none' execution testing. The reason
is that the 'C' array doesn't have the content it's checked to have.
Now, my question, shouldn't it be changed like the attached patch, or
similar, because we actually first need to return from function 'f' in
order for the 'copyout(C)' to be executed?
Or, otherwise, what's the use of the 'copyout' clause with OpenACC
'declare'? I suppose it's only meant to be used with function arguments,
because for every locally-defined variable (like 'C' in
'libgomp.oacc-c++/declare-pr94120.C'), that variable will be gone just
after the 'copyout' is executed, so the 'copyout' conceptually
equals/acts as 'create' in that case. However, OpenACC explicitly does
allow 'copyout' in certain scenarios. Do you think my interpretation is
correct, or what am I missing? (In case the answer isn't obvious to you,
too, please file an issue with OpenACC, linking to this email for
reference.)
However, GCC doesn't like my changes either; actually two errors are
diagnosed:
[...]/libgomp.oacc-c++/declare-pr94120.C: In function ‘void f(int*)’:
[...]/libgomp.oacc-c++/declare-pr94120.C:21:30: error: array section in
‘#pragma acc declare’
21 | #pragma acc declare copyout (C[0:N])
| ^
[...]/libgomp.oacc-c++/declare-pr94120.C:21:30: error: ‘C’ must be a
variable declared in the same scope as ‘#pragma acc declare’
Please have a look, and as necessary file GCC PRs, and XFAIL the
'libgomp.oacc-c++/declare-pr94120.C' execution testing for
'-DACC_MEM_SHARED=0'.
Grüße
Thomas
On 2020-03-11T13:28:44+0100, Tobias Burnus<tob...@codesourcery.com> wrote:
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c++/declare-pr94120.C
@@ -0,0 +1,57 @@
+#include <openacc.h>
+#include <stdlib.h>
+
+#define N 8
+
+namespace one {
+ int A[N] = { 1, 2, 3, 4, 5, 6, 7, 8 };
+ #pragma acc declare copyin (A)
+};
+
+namespace outer {
+ namespace inner {
+ int B[N];
+ #pragma acc declare create (B)
+ };
+};
+
+static void
+f (void)
+{
+ int i;
+ int C[N];
+ #pragma acc declare copyout (C)
+
+ if (!acc_is_present (&one::A, sizeof (one::A)))
+ abort ();
+
+ if (!acc_is_present (&outer::inner::B, sizeof (outer::inner::B)))
+ abort ();
+
+#pragma acc parallel
+ for (i = 0; i < N; i++)
+ {
+ outer::inner::B[i] = one::A[i];
+ C[i] = outer::inner::B[i];
+ }
+
+ for (i = 0; i < N; i++)
+ {
+ if (C[i] != i + 1)
+ abort ();
+ }
+
+#pragma acc parallel
+ for (i = 0; i < N; i++)
+ if (outer::inner::B[i] != i + 1)
+ abort ();
+}
+
+
+int
+main (int argc, char **argv)
+{
+ f ();
+
+ return 0;
+}
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander
Walter
r10-7810-g85d8c05a02bf7d1b256f806582a11e3fd8970a32
commit 85d8c05a02bf7d1b256f806582a11e3fd8970a32
Author: Tobias Burnus <tob...@codesourcery.com>
Date: Mon Apr 20 12:38:50 2020 +0200
Fix declare copyout in libgomp.oacc-c++/declare-pr94120.C
Testing on the host does not make sense for 'declare copyout' for
a same-scope stack-allocated variable. Once the copyout is done,
the variable is gone. Hence, test the variable on the device. This
can be revisit after the OpenACC semantic has been fixed; but with
that fix, the test PASSes again with devices.
PR middle-end/94120
* testsuite/libgomp.oacc-c++/declare-pr94120.C: Fix 'declare copy(out)'
test case.
diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index ce71ac6e783..b1cf297a0d7 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,9 @@
+2020-04-20 Tobias Burnus <tob...@codesourcery.com>
+
+ PR middle-end/94120
+ * testsuite/libgomp.oacc-c++/declare-pr94120.C: Fix 'declare copy(out)'
+ test case.
+
2020-04-17 Tobias Burnus <tob...@codesourcery.com>
PR middle-end/94635
diff --git a/libgomp/testsuite/libgomp.oacc-c++/declare-pr94120.C b/libgomp/testsuite/libgomp.oacc-c++/declare-pr94120.C
index 1e1254187ea..ed69359b533 100644
--- a/libgomp/testsuite/libgomp.oacc-c++/declare-pr94120.C
+++ b/libgomp/testsuite/libgomp.oacc-c++/declare-pr94120.C
@@ -35,6 +35,7 @@ f (void)
C[i] = outer::inner::B[i];
}
+#pragma acc parallel
for (i = 0; i < N; i++)
{
if (C[i] != i + 1)
