aboutsummaryrefslogtreecommitdiff
path: root/libgomp/testsuite/libgomp.c-c++-common/requires-4a.c
blob: 0e0db927c2c18fabb9a11e258a9becb261e11ea7 (plain)
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
/* { 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;
}