/* { dg-require-effective-target lto } */
/* { dg-additional-options "-flto" } */
/* { dg-additional-options "-foffload-options=nvptx-none=-misa=sm_35" { target { offload_target_nvptx } } } */
/* { dg-additional-sources requires-4-aux.c } */

/* Same as requires-4.c, but uses heap memory for 'a'.  */

/* Check no diagnostic by device-compiler's or host compiler's lto1.
   Other file uses: 'requires reverse_offload', but that's inactive as
   there are no declare target directives, device constructs nor device routines  */

/* Depending on offload device capabilities, it may print something like the
   following (only) if GOMP_DEBUG=1:
   "devices present but 'omp requires unified_address, unified_shared_memory, reverse_offload' cannot be fulfilled"
   and in that case does host-fallback execution.

   No offload devices support USM at present, so we may verify host-fallback
   execution by presence of separate memory spaces.  */

#pragma omp requires unified_address,unified_shared_memory

int *a;
extern void foo (void);

int
main (void)
{
  a = (int *) __builtin_calloc (sizeof (int), 10);
  #pragma omp target map(to: a)
  for (int i = 0; i < 10; i++)
    a[i] = i;

  for (int i = 0; i < 10; i++)
    if (a[i] != i)
      __builtin_abort ();

  foo ();
  __builtin_free (a);
  return 0;
}