On 08/29/2016 12:46 AM, Chung-Lin Tang wrote:
> Index: oacc-parallel.c
> ===================================================================
> --- oacc-parallel.c (revision 239814)
> +++ oacc-parallel.c (working copy)
> @@ -38,15 +38,23 @@
> #include <stdarg.h>
> #include <assert.h>
>
> +/* Returns the number of mappings associated with the pointer or pset. PSET
> + have three mappings, whereas pointer have two. */
> +
> static int
> -find_pset (int pos, size_t mapnum, unsigned short *kinds)
> +find_pointer (int pos, size_t mapnum, unsigned short *kinds)
> {
> if (pos + 1 >= mapnum)
> return 0;
>
> unsigned char kind = kinds[pos+1] & 0xff;
>
> - return kind == GOMP_MAP_TO_PSET;
> + if (kind == GOMP_MAP_TO_PSET)
> + return 3;
> + else if (kind == GOMP_MAP_POINTER)
> + return 2;
> +
> + return 0;
> }
Is this still necessary with the firstprivatization of subarrays
pointers? Well, it might be for fortran. Conceptually, the gimplifier
should prune out those unnecessary firstprivate pointer clauses for
executable constructs such as enter/exit data and update.
Actually, this is one area in the spec where the intent of enter/exit
data conflicts with what it describes. If you look at the runtime
documentation for, say, acc_create, it states that
acc_create (pvar, n*sizeof(var))
is equivalent to
acc enter data create (pvar[n])
And to free acc_create, you use acc_delete. So in theory, you should be
able to
#pragma acc enter data create (pvar[n])
acc_free (pvar)
but this may result in a memory leak if the pointer mapping isn't freed.
Fortran is somewhat special because of the pointer sets. I'm not sure if
its possible to make the OpenACC runtime API compatible with enter/exit
data.
> static void goacc_wait (int async, int num_waits, va_list *ap);
> @@ -298,7 +306,9 @@ GOACC_enter_exit_data (int device, size_t mapnum,
>
> if (kind == GOMP_MAP_FORCE_ALLOC
> || kind == GOMP_MAP_FORCE_PRESENT
> - || kind == GOMP_MAP_FORCE_TO)
> + || kind == GOMP_MAP_FORCE_TO
> + || kind == GOMP_MAP_TO
> + || kind == GOMP_MAP_ALLOC)
> {
> data_enter = true;
> break;
> @@ -312,31 +322,39 @@ GOACC_enter_exit_data (int device, size_t mapnum,
> kind);
> }
>
> + /* In c, non-pointers and arrays are represented by a single data clause.
> + Dynamically allocated arrays and subarrays are represented by a data
> + clause followed by an internal GOMP_MAP_POINTER.
> +
> + In fortran, scalars and not allocated arrays are represented by a
> + single data clause. Allocated arrays and subarrays have three mappings:
> + 1) the original data clause, 2) a PSET 3) a pointer to the array data.
> + */
> +
> if (data_enter)
> {
> for (i = 0; i < mapnum; i++)
> {
> unsigned char kind = kinds[i] & 0xff;
>
> - /* Scan for PSETs. */
> - int psets = find_pset (i, mapnum, kinds);
> + /* Scan for pointers and PSETs. */
> + int pointer = find_pointer (i, mapnum, kinds);
>
> - if (!psets)
> + if (!pointer)
> {
> switch (kind)
> {
> - case GOMP_MAP_POINTER:
> - gomp_acc_insert_pointer (1, &hostaddrs[i], &sizes[i],
> - &kinds[i]);
> + case GOMP_MAP_ALLOC:
> + acc_present_or_create (hostaddrs[i], sizes[i]);
> break;
> case GOMP_MAP_FORCE_ALLOC:
> acc_create (hostaddrs[i], sizes[i]);
> break;
> - case GOMP_MAP_FORCE_PRESENT:
> + case GOMP_MAP_TO:
> acc_present_or_copyin (hostaddrs[i], sizes[i]);
> break;
> case GOMP_MAP_FORCE_TO:
> - acc_present_or_copyin (hostaddrs[i], sizes[i]);
> + acc_copyin (hostaddrs[i], sizes[i]);
> break;
Thanks for correcting that. I had some of those data mappings wrong.
> default:
> gomp_fatal (">>>> GOACC_enter_exit_data UNHANDLED kind
> 0x%.2x",
> @@ -346,12 +364,16 @@ GOACC_enter_exit_data (int device, size_t mapnum,
> }
> else
> {
> - gomp_acc_insert_pointer (3, &hostaddrs[i], &sizes[i], &kinds[i]);
> + if (!acc_is_present (hostaddrs[i], sizes[i]))
> + {
> + gomp_acc_insert_pointer (pointer, &hostaddrs[i],
> + &sizes[i], &kinds[i]);
> + }
> /* Increment 'i' by two because OpenACC requires fortran
> arrays to be contiguous, so each PSET is associated with
> one of MAP_FORCE_ALLOC/MAP_FORCE_PRESET/MAP_FORCE_TO, and
> one MAP_POINTER. */
> - i += 2;
> + i += pointer - 1;
> }
> }
> }
> @@ -360,19 +382,15 @@ GOACC_enter_exit_data (int device, size_t mapnum,
> {
> unsigned char kind = kinds[i] & 0xff;
>
> - int psets = find_pset (i, mapnum, kinds);
> + int pointer = find_pointer (i, mapnum, kinds);
>
> - if (!psets)
> + if (!pointer)
> {
> switch (kind)
> {
> - case GOMP_MAP_POINTER:
> - gomp_acc_remove_pointer (hostaddrs[i], (kinds[i] & 0xff)
> - == GOMP_MAP_FORCE_FROM,
> - async, 1);
> - break;
> case GOMP_MAP_DELETE:
> - acc_delete (hostaddrs[i], sizes[i]);
> + if (acc_is_present (hostaddrs[i], sizes[i]))
> + acc_delete (hostaddrs[i], sizes[i]);
> break;
> case GOMP_MAP_FORCE_FROM:
> acc_copyout (hostaddrs[i], sizes[i]);
> @@ -385,10 +403,14 @@ GOACC_enter_exit_data (int device, size_t mapnum,
> }
> else
> {
> - gomp_acc_remove_pointer (hostaddrs[i], (kinds[i] & 0xff)
> - == GOMP_MAP_FORCE_FROM, async, 3);
> - /* See the above comment. */
> - i += 2;
> + if (acc_is_present (hostaddrs[i], sizes[i]))
> + {
> + gomp_acc_remove_pointer (hostaddrs[i], (kinds[i] & 0xff)
> + == GOMP_MAP_FORCE_FROM, async,
> + pointer);
> + /* See the above comment. */
> + }
> + i += pointer - 1;
> }
> }
>
>
>
> libgomp-enter-exit-testsuite.patch
>
>
> Index: testsuite/libgomp.oacc-c-c++-common/data-2.c
> ===================================================================
> --- testsuite/libgomp.oacc-c-c++-common/data-2.c (revision 239814)
> +++ testsuite/libgomp.oacc-c-c++-common/data-2.c (working copy)
> @@ -3,6 +3,7 @@
> /* { dg-do run } */
>
> #include <stdlib.h>
> +#include <openacc.h>
>
> int
> main (int argc, char **argv)
> @@ -32,7 +33,7 @@ main (int argc, char **argv)
> for (i = 0; i < N; i++)
> b[i] = a[i];
>
> -#pragma acc exit data copyout (a[0:N]) copyout (b[0:N]) wait async
> +#pragma acc exit data copyout (a[0:N], b[0:N]) delete (N) wait async
> #pragma acc wait
One note about these tests in general. I wonder if we should also be
testing subarrays with non-zero base offsets. We already hit one bug
with local arrays.
Cesar
