[OpenACC] declare directive

[Jakub Jelinek]

On Mon, Nov 09, 2015 at 05:11:44PM -0600, James Norris wrote:
> diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h
> index 953c4e3..c6a2981 100644
> --- a/gcc/c-family/c-pragma.h
> +++ b/gcc/c-family/c-pragma.h
> @@ -30,6 +30,7 @@ enum pragma_kind {
>    PRAGMA_OACC_ATOMIC,
>    PRAGMA_OACC_CACHE,
>    PRAGMA_OACC_DATA,
> +  PRAGMA_OACC_DECLARE,
>    PRAGMA_OACC_ENTER_DATA,
>    PRAGMA_OACC_EXIT_DATA,
>    PRAGMA_OACC_KERNELS,

This change will make PR68271 even worse, so would be really nice to
get that fixed first.

> +	  case GOMP_MAP_ALLOC:
> +	    if (!acc_is_present (hostaddrs[i], sizes[i]))
> +	      {
> +		GOACC_enter_exit_data (device, 1, &hostaddrs[i], &sizes[i],
> +				       &kinds[i], 0, 0);
> +	      }

No {}s around single statement body.

> +	  case GOMP_MAP_FORCE_PRESENT:
> +	    if (!acc_is_present (hostaddrs[i], sizes[i]))
> +	      gomp_fatal ("[%p,%zd] is not mapped", hostaddrs[i], sizes[i]);

This isn't portable unfortunately to all targets that build libgomp.
Looking around, we use various ways to print sizes in gomp_fatal:
1) use %ld and cast to unsigned long
2) use %d and cast to int
3)
#ifdef HAVE_INTTYPES_H
                      gomp_fatal ("present clause: !acc_is_present (%p, "
                                  "%"PRIu64" (0x%"PRIx64"))",
                                  (void *) k->host_start,
                                  (uint64_t) size, (uint64_t) size);
#else
                      gomp_fatal ("present clause: !acc_is_present (%p, "
                                  "%lu (0x%lx))", (void *) k->host_start,
                                  (unsigned long) size, (unsigned long) size);
#endif

I'd say use any of those for now, and if you or one of your collegues could
clean all this up, it would be greatly appreciated.
The best might be to handle this somewhere in libgomp.h by testing something
like:
#if defined(__GLIBC__) // or some configure check whether %zd works
# define GOMP_PRIuSIZE_T "zu"
# define GOMP_PRIxSIZE_T "zx"
typedef size_t gomp_prisize_t;
#elif __SIZEOF_SIZE_T__ == __SIZEOF_INT__
# define GOMP_PRIuSIZE_T "u"
# define GOMP_PRIxSIZE_T "x"
typedef unsigned int gomp_prisize_t;
#elif __SIZEOF_SIZE_T__ == __SIZEOF_LONG__
# define GOMP_PRIuSIZE_T "lu"
# define GOMP_PRIxSIZE_T "lx"
typedef unsigned long gomp_prisize_t;
#elif defined (HAVE_INTTYPES_H) && __SIZEOF_SIZE_T__ == 8 && __CHAR_BIT__ == 8
# define GOMP_PRIuSIZE_T PRIu64
# define GOMP_PRIxSIZE_T PRIx64
typedef uint64_t gomp_prisize_t;
#else
# define GOMP_PRIuSIZE_T "lu"
# define GOMP_PRIxSIZE_T "lx"
typedef unsigned long gomp_prisize_t;
#endif
and then use those macros and always cast size_t arguments to
gomp_prisize_t in the various gomp_fatal or other printing calls.

> +int
> +main (int argc, char **argv)
> +{
> +  int a[8] __attribute__((unused));
> +
> +  __builtin_printf ("CheCKpOInT\n");
> +#pragma acc declare present (a)
> +}
> +
> +/* { dg-output "CheCKpOInT" } */
> +/* { dg-shouldfail "" } */

Are you sure printf will have the buffers flushed before abort on all
targets?

	Jakub