[gomp4] libgomp: Don't update copy_from for the existing mappings.

Hi Thomas,

Here is the fix for the issue discussed in http://gcc.gnu.org/ml/gcc/2014-02/msg00257.html

OK to commit to gomp-4_0-branch?

---
 libgomp/ChangeLog.gomp |    5 +++++
 libgomp/target.c       |    5 -----
 2 files changed, 5 insertions(+), 5 deletions(-)

Hi Ilya!

On Mon, 17 Feb 2014 21:06:41 +0400, Ilya Verbin <iverbin@gmail.com> wrote:
> Here is the fix for the issue discussed in http://gcc.gnu.org/ml/gcc/2014-02/msg00257.html
> 
> OK to commit to gomp-4_0-branch?

> +	* target.c (gomp_map_vars_existing): Don't update copy_from for the
> +	existing mappings.

> --- a/libgomp/target.c
> +++ b/libgomp/target.c
> @@ -171,11 +171,6 @@ gomp_map_vars_existing (splay_tree_key oldn, splay_tree_key newn,
>  		"[%p..%p) is already mapped",
>  		(void *) newn->host_start, (void *) newn->host_end,
>  		(void *) oldn->host_start, (void *) oldn->host_end);
> -  if (((kind & 7) == 2 || (kind & 7) == 3)
> -      && !oldn->copy_from
> -      && oldn->host_start == newn->host_start
> -      && oldn->host_end == newn->host_end)
> -    oldn->copy_from = true;
>    oldn->refcount++;
>  }

Sure; this was the consensus, as I understand it.  Though, wouldn't it
make sense to also add a test case to a) test for this behavior, and also
b) to document the "GCC interpretation of the OpenMP standard" in this
case (including a link to the OpenMP issue tracker, or Jakub's email on
this topic)?


Grüße,
 Thomas

2014-02-18 20:25 GMT+04:00 Thomas Schwinge <thomas@codesourcery.com>:
> Sure; this was the consensus, as I understand it.  Though, wouldn't it
> make sense to also add a test case to a) test for this behavior, and also
> b) to document the "GCC interpretation of the OpenMP standard" in this
> case (including a link to the OpenMP issue tracker, or Jakub's email on
> this topic)?
>
> Grüße,
>  Thomas

This issue is not reproducible on the shared memory.  So, currently
(while we perform host fallback), it's impossible to create a
test-case.  But we're going to create a libgomp plugin, that will
emulate a target device with a private memory.  It will allow to run
tests for offloading without having accelerator device.

  -- Ilya

Hi Ilya!

On Tue, 18 Feb 2014 20:49:39 +0400, Ilya Verbin <iverbin@gmail.com> wrote:
> 2014-02-18 20:25 GMT+04:00 Thomas Schwinge <thomas@codesourcery.com>:
> > Sure; this was the consensus, as I understand it.  Though, wouldn't it
> > make sense to also add a test case to a) test for this behavior, and also
> > b) to document the "GCC interpretation of the OpenMP standard" in this
> > case (including a link to the OpenMP issue tracker, or Jakub's email on
> > this topic)?
> 
> This issue is not reproducible on the shared memory.  So, currently
> (while we perform host fallback), it's impossible to create a
> test-case.  But we're going to create a libgomp plugin, that will
> emulate a target device with a private memory.  It will allow to run
> tests for offloading without having accelerator device.

Right, that's what we also had suggested, prepared, and now posted in
<http://news.gmane.org/find-root.php?message_id=%3C87r46zt76g.fsf%40schwinge.name%3E>.


Using this, and extending your test case as follows:

@@ -6,10 +6,13 @@
 int main ()
 {
   int *a = malloc (N * sizeof (int));
+  int *b = a;
   printf ("1: %p\n", a);
 #pragma omp target data map(tofrom: a[0:N])
   {
     printf ("2: %p\n", a);
+    if (a != b)
+      abort ();
 #pragma omp target
     {
       int i;
@@ -17,8 +20,12 @@
 	a[i] = i;
     }
     printf ("3: %p\n", a);
+    if (a != b)
+      abort ();
   }
   printf ("4: %p\n", a);
+  if (a != b)
+    abort ();
   free (a);
   return 0;
 }

..., I then see the following with the current code:

    $ LIBGOMP_PLUGIN_PATH=$PWD ./a.out 
    1: 0x1737030
    libgomp plugin: ../source/libgomp/plugin-host.c:device_alloc (4015): 0x1738040
    libgomp plugin: ../source/libgomp/plugin-host.c:device_host2dev (0x1738040, 0x1737030, 4000)
    libgomp plugin: ../source/libgomp/plugin-host.c:device_host2dev (0x1738fe0, 0x7fff808f2cb8, 8)
    2: 0x1737030
    libgomp plugin: ../source/libgomp/plugin-host.c:device_alloc (39): 0x17391a0
    libgomp plugin: ../source/libgomp/plugin-host.c:device_host2dev (0x17391a0, 0x7fff808f2c38, 8)
    libgomp plugin: ../source/libgomp/plugin-host.c:device_free (0x17391a0)
    3: 0x1737030
    libgomp plugin: ../source/libgomp/plugin-host.c:device_dev2host (0x1737030, 0x1738040, 4000)
    libgomp plugin: ../source/libgomp/plugin-host.c:device_dev2host (0x7fff808f2d48, 0x1738fe0, 8)
    libgomp plugin: ../source/libgomp/plugin-host.c:device_free (0x1738040)
    4: 0x1738040
    Aborted (core dumped)

..., and with your libgomp patch applied:

    $ LIBGOMP_PLUGIN_PATH=$PWD ./a.out
    1: 0x74c030
    libgomp plugin: ../source/libgomp/plugin-host.c:device_alloc (4015): 0x74d040
    libgomp plugin: ../source/libgomp/plugin-host.c:device_host2dev (0x74d040, 0x74c030, 4000)
    libgomp plugin: ../source/libgomp/plugin-host.c:device_host2dev (0x74dfe0, 0x7fffef55f0f8, 8)
    2: 0x74c030
    libgomp plugin: ../source/libgomp/plugin-host.c:device_alloc (39): 0x74e1a0
    libgomp plugin: ../source/libgomp/plugin-host.c:device_host2dev (0x74e1a0, 0x7fffef55f078, 8)
    libgomp plugin: ../source/libgomp/plugin-host.c:device_free (0x74e1a0)
    3: 0x74c030
    libgomp plugin: ../source/libgomp/plugin-host.c:device_dev2host (0x74c030, 0x74d040, 4000)
    libgomp plugin: ../source/libgomp/plugin-host.c:device_free (0x74d040)
    4: 0x74c030


Grüße,
 Thomas