On Fri, 31 Aug 2018 16:20:08 +0200
Jakub Jelinek <ja...@redhat.com> wrote:
> On Fri, Aug 31, 2018 at 10:04:07AM -0400, Nathan Sidwell wrote:
> > On 08/30/2018 04:27 PM, Jason Merrill wrote:
> >
> > > On Thu, Aug 30, 2018 at 3:31 PM, Julian Brown
> > > <jul...@codesourcery.com> wrote:
> > > > "Apart from parsing, it's necessary to prevent the "cannot take
> > > > the address of 'this', which is an rvalue expression" error
> > > > from appearing
> >
> > Breaking a rather fundamental language attribute does not seem wise.
> >
> > > Why does referring to this[0:1] require making 'this' addressable?
> > > Surely what we're interested in is the value of 'this', not the
> > > address.
> > Yes, transferring the this pointer is very unlikely to be what the
> > user wants -- the object being referred to contains the data. It
> > might be wise to look at the DR's and changes relating to lambdas
> > and this capture. Those changes now make it much harder to simply
> > capture the pointer unintentionally.
>
> Yeah, I agree we shouldn't try to make this addressable.
> Does OpenACC try to map the base of the array section (rather than
> what e.g. OpenMP does, privatize the pointer base instead and assign
> the pointer the new value inside of the region)?
> Even if it is mapped, can't it be mapped by taking an address of a
> temporary initialized from this?
For OpenACC, two mappings are created for an array section: one for the
data (to, from, tofrom, etc.) and a firstprivate pointer with a bias to
locate the (possibly virtual) zero'th element of the array. I think
that's the same as OpenMP.
For the test case given, it's sufficient to merely allow "this" to be
used as the base pointer for an array section. That usage doesn't
require "this" to be made addressable.
The this[0:1] syntax is accepted by PGI
(https://www.pgroup.com/resources/docs/18.4/x86/openacc-gs/index.htm,
2.4 C++ Classes in OpenACC) -- in order to copy "the class itself" to
the accelerator.
Referring to class member variables in OpenACC clauses (as the example
in that section does also) is still problematic in GCC, though.
PGI also allows the user to specify just "this" in OpenACC clauses,
which presumably does the same thing as specifying this[0:1]. For PGI,
but not for OpenACC <= 2.5, that seems to follow a general case for
pointers to structs (2.3. C Structs in OpenACC), "A pointer to a scalar
struct is treated as a one-element array, and should be shaped as
r[0:1]". That's notably different from OpenMP 4.5, which treats plain
mapped pointers as zero-length array sections, and also differs from
the current behaviour of GCC (which bizarrely, IIUC, is to copy the
bits of the host pointer verbatim to the target). OpenACC 2.5 arguably
leaves the behaviour unspecified for pointers without an explicit array
section.
The attached patch allows basic class-wrapping-an-array kinds of
usages, anyway. Re-tested with offloading to nvptx and bootstrapped. OK
to apply?
Thanks,
Julian
2018-09-03 Joseph Myers <jos...@codesourcery.com>
Julian Brown <jul...@codesourcery.com>
PR C++/66053
* semantics.c (handle_omp_array_sections_1): Allow array
sections with "this" pointer for OpenACC.
commit 355411e5415f65e09a06f42d400761fff065f7c7
Author: Julian Brown <jul...@codesourcery.com>
Date: Fri Aug 31 17:30:20 2018 -0700
Allow this[:] array slices for OpenACC
2018-09-03 Joseph Myers <jos...@codesourcery.com>
Julian Brown <jul...@codesourcery.com>
* semantics.c (handle_omp_array_sections_1): Allow array sections with
"this" pointer for OpenACC.
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index 676de01..98511ed 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -4598,7 +4598,8 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
return error_mark_node;
}
- else if (TREE_CODE (t) == PARM_DECL
+ else if (ort == C_ORT_OMP
+ && TREE_CODE (t) == PARM_DECL
&& DECL_ARTIFICIAL (t)
&& DECL_NAME (t) == this_identifier)
{
diff --git a/libgomp/testsuite/libgomp.oacc-c++/this.C b/libgomp/testsuite/libgomp.oacc-c++/this.C
new file mode 100644
index 0000000..510c690
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c++/this.C
@@ -0,0 +1,43 @@
+#include <cstdlib>
+#include <iostream>
+using namespace std;
+
+class test {
+ public:
+ int a;
+
+ test ()
+ {
+ a = -1;
+#pragma acc enter data copyin (this[0:1])
+ }
+
+ ~test ()
+ {
+#pragma acc exit data delete (this[0:1])
+ }
+
+ void set (int i)
+ {
+ a = i;
+#pragma acc update device (this[0:1])
+ }
+
+ int get ()
+ {
+#pragma acc update host (this[0:1])
+ return a;
+ }
+};
+
+int
+main ()
+{
+ test t;
+
+ t.set (4);
+ if (t.get () != 4)
+ abort ();
+
+ return 0;
+}
