Hi Ilya! On Tue, 18 Feb 2014 20:49:39 +0400, Ilya Verbin <iver...@gmail.com> wrote: > 2014-02-18 20:25 GMT+04:00 Thomas Schwinge <tho...@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
pgpdQGEfbKbKo.pgp
Description: PGP signature
