[gomp4] Questions about "declare target" and "target update" pragmas

[gomp4] Questions about "declare target" and "target update" pragmas

[Ilya Verbin]

Hi Jakub,

I have 2 questions concerning OpenMP 4.0 specification.


1.  Do I understand correctly that every "declare target" directive should be
closed with "end declare target"?  E.g. in this example GCC marks both foo1 and
foo2 with "omp declare target" attribute:

#pragma omp declare target
int foo1 () { return 1; }
int foo2 () { return 2; }
/* EOF */

Shouldn't the frontend issue an error message that there is "declare target"
without corresponding "end declare target"?


2.  Do I understand correctly that the "target update" directive can be used
outside the "target" or "target data" regions?  E.g. this example should
print '2' (and it prints '2' while building with icc):

#pragma omp declare target
int G = 1;
#pragma omp end declare target

int main ()
{
  G = 2;
  #pragma omp target update to(G)
  G = 3;
  int x = 0;
  #pragma omp target
    {
      x = G;
    }
  printf ("%d\n", x);
}

If it is acceptable, then GOMP_target_update should also map variables that
wasn't mapped.

Thanks,
    -- Ilya

Re: [gomp4] Questions about "declare target" and "target update" pragmas

[Jakub Jelinek]

On Wed, Jan 22, 2014 at 07:51:51PM +0400, Ilya Verbin wrote:
> I have 2 questions concerning OpenMP 4.0 specification.
> 
> 
> 1.  Do I understand correctly that every "declare target" directive should be
> closed with "end declare target"?  E.g. in this example GCC marks both foo1 and
> foo2 with "omp declare target" attribute:
> 
> #pragma omp declare target
> int foo1 () { return 1; }
> int foo2 () { return 2; }
> /* EOF */

Yes, this is invalid.

> Shouldn't the frontend issue an error message that there is "declare target"
> without corresponding "end declare target"?

I guess so, e.g. the C FE has current_omp_declare_target_attribute,
supposedly it could check it after parsing the whole file and if it is
non-zero at the end, complain.

> 2.  Do I understand correctly that the "target update" directive can be used
> outside the "target" or "target data" regions?  E.g. this example should
> print '2' (and it prints '2' while building with icc):
> 
> #pragma omp declare target
> int G = 1;
> #pragma omp end declare target
> 
> int main ()
> {
>   G = 2;
>   #pragma omp target update to(G)
>   G = 3;
>   int x = 0;
>   #pragma omp target
>     {
>       x = G;
>     }
>   printf ("%d\n", x);
> }

This can print 3 (if doing host fallback or device shares address space
with host), or 2 (otherwise).  It shouldn't print 1 ever, and yes,
the target update is then well defined.  All variables from omp declare
target are allocated on the device sometime before
the first target data/target update/target region; given that they will
be allocated in the data section of the target DSO, they actually just need
to be registered with the mapping data structure when the DSO is loaded.

> If it is acceptable, then GOMP_target_update should also map variables that
> wasn't mapped.

No, the target DSO initialization should use the tables we've talked about
to initialize the mapping.

	Jakub

Re: [gomp4] Questions about "declare target" and "target update" pragmas

[Ilya Verbin]

2014/1/22 Jakub Jelinek <jakub@redhat.com>:
> This can print 3 (if doing host fallback or device shares address space
> with host), or 2 (otherwise).  It shouldn't print 1 ever, and yes,
> the target update is then well defined.  All variables from omp declare
> target are allocated on the device sometime before
> the first target data/target update/target region; given that they will
> be allocated in the data section of the target DSO, they actually just need
> to be registered with the mapping data structure when the DSO is loaded.
>
> No, the target DSO initialization should use the tables we've talked about
> to initialize the mapping.
>
>         Jakub

Yes, when G is global variable marked with 'declare target', everything works
fine.  But this testcase crashes at runtime in GOMP_target_update:

int main ()
{
  int G = 2;
  #pragma omp target update to(G)
  G = 3;
  int x = 0;
  #pragma omp target
    {
      x = G;
    }
  printf ("%d\n", x);
}

Is it right, that such usage of 'target update' is not allowed by omp
specification?

  -- Ilya

Re: [gomp4] Questions about "declare target" and "target update" pragmas

[Jakub Jelinek]

On Tue, Jan 28, 2014 at 09:54:09PM +0400, Ilya Verbin wrote:
> Yes, when G is global variable marked with 'declare target', everything works
> fine.  But this testcase crashes at runtime in GOMP_target_update:
> 
> int main ()
> {
>   int G = 2;
>   #pragma omp target update to(G)
>   G = 3;
>   int x = 0;
>   #pragma omp target
>     {
>       x = G;
>     }
>   printf ("%d\n", x);
> }
> 
> Is it right, that such usage of 'target update' is not allowed by omp
> specification?

Yes, this testcase is invalid.
"If the corresponding list item is not present in the device data
environment, the behavior is unspecified."
Perhaps we shouldn't crash but do nothing, or complain to stderr, as QoI.

	Jakub

Fwd: [gomp4] Questions about "declare target" and "target update" pragmas

[Ilya Verbin]

One more question.  Is it valid to use arr[MAX/2..MAX] on target?

#define MAX 20
void foo ()
{
  int arr[MAX];
  #pragma omp target map(from: arr[0:MAX/2])
    {
      int i;
      for (i = 0; i < MAX; i++)
        arr[i] = i;
    }
}

In this case GOMP_target gets sizes[0]==40 as input.  Due to this,
gomp_map_vars allocates 40 bytes of memory on target for 'arr',
instead of 80 bytes.

  -- Ilya

Re: [gomp4] Questions about "declare target" and "target update" pragmas

[Ilya Verbin]

Hi Jakub,

I discovered an issue related to global variables.
In this testcase the 'omp target' child fn uses the local copy of glob_var.
But the 'omp parallel' child fn tries to use the glob_var directly and therefore
crashes.

int glob_var;

void
foo (void)
{
  glob_var = 1;
  #pragma omp target map(to: glob_var)
    {
      glob_var = 2;
      #pragma omp parallel
        {
          glob_var = 3;
        }
    }
}

In the spec I found only that if a variable is not present in the enclosing
device data environment, then it is allocated in the device data environment
associated with the construct.  Effectively this means glob_var become a
non-global variable within the omp_target construct.  Then it's not quite clear
what type of glob_var should be in case target is not present and we fallback to
host execution.

Should we forbid 'omp target' to map global vars that are not defined as target?
Or force 'omp parallel' to use local copies inside the target regions?
(ICC reports an error about mapping glob_var for this testcase)

  -- Ilya

Re: [gomp4] Questions about "declare target" and "target update" pragmas

[Ilya Verbin]

Hi Jakub,

I have one more question :)
This testcase seems to be correct... or not?

#pragma omp declare target
extern int G;
#pragma omp end declare target

int G;

int main ()
{
  #pragma omp target update to(G)

  return 0;
}

If yes, then we have a problem that the decl of G in varpool_node::get_create
doesn't have "omp declare target" attribute.

Thanks,
  -- Ilya

Re: [gomp4] Questions about "declare target" and "target update" pragmas

[Ilya Verbin]

On Tue, Mar 10, 2015 at 19:52:52 +0300, Ilya Verbin wrote:
> Hi Jakub,
> 
> I have one more question :)
> This testcase seems to be correct... or not?
> 
> #pragma omp declare target
> extern int G;
> #pragma omp end declare target
> 
> int G;
> 
> int main ()
> {
>   #pragma omp target update to(G)
> 
>   return 0;
> }
> 
> If yes, then we have a problem that the decl of G in varpool_node::get_create
> doesn't have "omp declare target" attribute.

Ping?

I am investigating run-fails on some benchmark, and have found a second
questionable place, where a function argument overrides a global array.
Just to be sure, is this a bug in the test?

#pragma omp declare target
int a1[50], a2[50];
#pragma omp end declare target

void foo (int a1[])
{
  #pragma omp target
    {
      a1[10]++;
      a2[10]++;
    }
}

int main ()
{
  a1[10] = a2[10] = 10;

  #pragma omp target update to(a1, a2)
  foo (a1);
  #pragma omp target update from(a1, a2)

  if (a1[10] != a2[10])
    abort ();
  return 0;
}

Thanks,
  -- Ilya

Re: [gomp4] Questions about "declare target" and "target update" pragmas

[Jakub Jelinek]

On Mon, Mar 16, 2015 at 09:41:53PM +0300, Ilya Verbin wrote:
> On Tue, Mar 10, 2015 at 19:52:52 +0300, Ilya Verbin wrote:
> > Hi Jakub,
> > 
> > I have one more question :)
> > This testcase seems to be correct... or not?
> > 
> > #pragma omp declare target
> > extern int G;
> > #pragma omp end declare target
> > 
> > int G;
> > 
> > int main ()
> > {
> >   #pragma omp target update to(G)
> > 
> >   return 0;
> > }
> > 
> > If yes, then we have a problem that the decl of G in varpool_node::get_create
> > doesn't have "omp declare target" attribute.
> 
> Ping?

Here is untested patch.  I'm going to check it in after bootstrap/regtest.

> I am investigating run-fails on some benchmark, and have found a second
> questionable place, where a function argument overrides a global array.
> Just to be sure, is this a bug in the test?
> 
> #pragma omp declare target
> int a1[50], a2[50];
> #pragma omp end declare target
> 
> void foo (int a1[])
> {
>   #pragma omp target
>     {
>       a1[10]++;
>       a2[10]++;
>     }
> }

That is a buggy test.  int a1[] function argument is changed
into int *a1, so it is actually
#pragma omp target map(tofrom:a1, a2)
{
  a1[10]++;
  a2[10]++;
}
which copies the a1 pointer to the device by value (no pointer
transformation).
Perhaps the testcase writer meant to use #pragma omp target map(a1[10])
instead (or map(a1[0:50])?

2015-03-19  Jakub Jelinek  <jakub@redhat.com>

	* c-decl.c (c_decl_attributes): Also add "omp declare target"
	attribute for DECL_EXTERNAL VAR_DECLs.

	* decl2.c (cplus_decl_attributes): Also add "omp declare target"
	attribute for DECL_EXTERNAL VAR_DECLs.

	* testsuite/libgomp.c/target-10.c: New test.
	* testsuite/libgomp.c++/target-4.C: New test.

--- gcc/c/c-decl.c.jj	2015-03-09 19:24:34.000000000 +0100
+++ gcc/c/c-decl.c	2015-03-19 13:01:15.423381262 +0100
@@ -4407,7 +4407,8 @@ c_decl_attributes (tree *node, tree attr
 {
   /* Add implicit "omp declare target" attribute if requested.  */
   if (current_omp_declare_target_attribute
-      && ((TREE_CODE (*node) == VAR_DECL && TREE_STATIC (*node))
+      && ((TREE_CODE (*node) == VAR_DECL
+	   && (TREE_STATIC (*node) || DECL_EXTERNAL (*node)))
 	  || TREE_CODE (*node) == FUNCTION_DECL))
     {
       if (TREE_CODE (*node) == VAR_DECL
--- gcc/cp/decl2.c.jj	2015-03-18 11:53:13.000000000 +0100
+++ gcc/cp/decl2.c	2015-03-19 13:04:30.739176009 +0100
@@ -1440,7 +1440,8 @@ cplus_decl_attributes (tree *decl, tree
 
   /* Add implicit "omp declare target" attribute if requested.  */
   if (scope_chain->omp_declare_target_attribute
-      && ((TREE_CODE (*decl) == VAR_DECL && TREE_STATIC (*decl))
+      && ((TREE_CODE (*decl) == VAR_DECL
+	   && (TREE_STATIC (*decl) || DECL_EXTERNAL (*decl)))
 	  || TREE_CODE (*decl) == FUNCTION_DECL))
     {
       if (TREE_CODE (*decl) == VAR_DECL
--- libgomp/testsuite/libgomp.c/target-10.c.jj	2015-03-19 13:06:56.812778618 +0100
+++ libgomp/testsuite/libgomp.c/target-10.c	2015-03-19 13:07:03.857662996 +0100
@@ -0,0 +1,14 @@
+/* { dg-do run } */
+
+#pragma omp declare target
+extern int v;
+#pragma omp end declare target
+
+int v;
+
+int
+main ()
+{
+  #pragma omp target update to(v)
+  return 0;
+}
--- libgomp/testsuite/libgomp.c++/target-4.C.jj	2015-03-19 13:07:19.286409775 +0100
+++ libgomp/testsuite/libgomp.c++/target-4.C	2015-03-19 13:07:30.322228651 +0100
@@ -0,0 +1,3 @@
+// { dg-do run }
+
+#include "../libgomp.c/target-10.c"

	Jakub

Re: [gomp4] Questions about "declare target" and "target update" pragmas

[Ilya Verbin]

On Thu, Mar 19, 2015 at 14:47:44 +0100, Jakub Jelinek wrote:
> Here is untested patch.  I'm going to check it in after bootstrap/regtest.

Thanks.

> > I am investigating run-fails on some benchmark, and have found a second
> > questionable place, where a function argument overrides a global array.
> > Just to be sure, is this a bug in the test?
> > 
> > #pragma omp declare target
> > int a1[50], a2[50];
> > #pragma omp end declare target
> > 
> > void foo (int a1[])
> > {
> >   #pragma omp target
> >     {
> >       a1[10]++;
> >       a2[10]++;
> >     }
> > }
> 
> That is a buggy test.  int a1[] function argument is changed
> into int *a1, so it is actually
> #pragma omp target map(tofrom:a1, a2)

Actually, it copies only a1 pointer, since a2 points to the global array.

> {
>   a1[10]++;
>   a2[10]++;
> }
> which copies the a1 pointer to the device by value (no pointer
> transformation).
> Perhaps the testcase writer meant to use #pragma omp target map(a1[10])
> instead (or map(a1[0:50])?

If I understand correctly, it's not allowed to map global target arrays this
way, since it's already present in the initial device data environment:

2.9.4 declare target Directive
If a list item is a variable then the original variable is mapped to a corresponding
variable in the initial device data environment for all devices.

2.14.5 map Clause
If a corresponding list item of the original list item is in the enclosing device data
environment, the new device data environment uses the corresponding list item from the
enclosing device data environment. No additional storage is allocated in the new device
data environment and neither initialization nor assignment is performed, regardless of
the map-type that is specified.

So, to fix this testcase I can just remove the "int a1[]" function argument, and
add some "#pragma omp target update" where needed.

  -- Ilya

Re: [gomp4] Questions about "declare target" and "target update" pragmas

[Jakub Jelinek]

On Thu, Mar 19, 2015 at 05:49:47PM +0300, Ilya Verbin wrote:
> > > void foo (int a1[])
> > > {
> > >   #pragma omp target
> > >     {
> > >       a1[10]++;
> > >       a2[10]++;
> > >     }
> > > }
> > 
> > That is a buggy test.  int a1[] function argument is changed
> > into int *a1, so it is actually
> > #pragma omp target map(tofrom:a1, a2)
> 
> Actually, it copies only a1 pointer, since a2 points to the global array.

Sure, the implicit map(tofrom:a2) doesn't really do anything, since a2
is clearly already mapped and will stay to be mapped, and is a global var.
> 
> > {
> >   a1[10]++;
> >   a2[10]++;
> > }
> > which copies the a1 pointer to the device by value (no pointer
> > transformation).
> > Perhaps the testcase writer meant to use #pragma omp target map(a1[10])
> > instead (or map(a1[0:50])?
> 
> If I understand correctly, it's not allowed to map global target arrays this
> way, since it's already present in the initial device data environment:

It of course is allowed.  It just means that it doesn't allocate new memory
(sizeof(int) large in the map(a1[10]) case, sizeof(int)*50 large in the map(a1[0:50])
case), nor copy the bytes around, all it will do is allocate memory for the
target copy of the a1 pointer, and do pointer transformation such that the
a1 pointer on the target will point to the global target a1 array.
Without the map(a1[10]) or map(a1[0:50]) clauses (i.e. implicit map(tofrom:a1))
it does similar thing, except it copies the pointer value to the target (and
back at the end of the region) instead, which is not what you want...

> So, to fix this testcase I can just remove the "int a1[]" function argument, and
> add some "#pragma omp target update" where needed.

Well, supposedly the test was meant to test it with a pointer parameter,
otherwise why would there be both a1 and a2 when a2 would be enough?

	Jakub

Re: [gomp4] Questions about "declare target" and "target update" pragmas

[Ilya Verbin]

On Thu, Mar 19, 2015 at 15:57:10 +0100, Jakub Jelinek wrote:
> On Thu, Mar 19, 2015 at 05:49:47PM +0300, Ilya Verbin wrote:
> > If I understand correctly, it's not allowed to map global target arrays this
> > way, since it's already present in the initial device data environment:
> 
> It of course is allowed.  It just means that it doesn't allocate new memory
> (sizeof(int) large in the map(a1[10]) case, sizeof(int)*50 large in the map(a1[0:50])
> case), nor copy the bytes around, all it will do is allocate memory for the
> target copy of the a1 pointer, and do pointer transformation such that the
> a1 pointer on the target will point to the global target a1 array.
> Without the map(a1[10]) or map(a1[0:50]) clauses (i.e. implicit map(tofrom:a1))
> it does similar thing, except it copies the pointer value to the target (and
> back at the end of the region) instead, which is not what you want...

Ok, got it.

And what about global allocatable fortran arrays?  I didn't find any
restrictions in the specification.  Here is a reduced testcase:

module test
  integer, allocatable, target :: x(:)
  !$omp declare target(x)
end module test
  use test
  integer :: n = 1000
  allocate (x(n))
  !$omp target map(x(1:n))
     x(123) = 456
  !$omp end target
  deallocate (x)
end

It crashes on target with NULL-pointer access, however the memory for x(1:n) is
allocated on target.  Looks like there's something wrong with pointer
transformation.  Is this a wrong testcase or a bug in gcc?

Thanks,
  -- Ilya

Re: [gomp4] Questions about "declare target" and "target update" pragmas

[Jakub Jelinek]

On Thu, Mar 19, 2015 at 09:42:58PM +0300, Ilya Verbin wrote:
> Ok, got it.
> 
> And what about global allocatable fortran arrays?  I didn't find any
> restrictions in the specification.  Here is a reduced testcase:

This really can't be supported.  If you have global allocatables, you really
should allocate them in the target IMNSHO.  The allocation on the host
doesn't make it allocated on the target.  Otherwise, how you could e.g.
deallocate it on the target or allocate again?
> 
> module test
>   integer, allocatable, target :: x(:)
>   !$omp declare target(x)
> end module test
>   use test
>   integer :: n = 1000
>   allocate (x(n))
>   !$omp target map(x(1:n))
>      x(123) = 456
>   !$omp end target
>   deallocate (x)
> end
> 
> It crashes on target with NULL-pointer access, however the memory for x(1:n) is
> allocated on target.  Looks like there's something wrong with pointer
> transformation.  Is this a wrong testcase or a bug in gcc?

	Jakub