1  /* { dg-do run }  */
       2  /* { dg-additional-options "-foffload-options=nvptx-none=-misa=sm_35" { target { offload_target_nvptx } } } */
       3  /* { dg-additional-sources reverse-offload-1-aux.c } */
       4  
       5  /* Check that reverse offload works in particular:
       6     - no code is generated on the device side (i.e. no
       7       implicit declare target of called functions and no
       8       code gen for the target-region body)
       9       -> would otherwise fail due to 'add_3' symbol
      10     - Plus the usual (compiles, runs, produces correct result)
      11  
      12     Note: Running also the non-reverse-offload target regions
      13     on the host (host fallback) is valid and will pass.  */
      14  
      15  #pragma omp requires reverse_offload
      16  
      17  extern int add_3 (int);
      18  
      19  static int global_var = 5;
      20  
      21  void
      22  check_offload (int *x, int *y)
      23  {
      24    *x = add_3 (*x);
      25    *y = add_3 (*y);
      26  }
      27  
      28  #pragma omp declare target
      29  void
      30  tg_fn (int *x, int *y)
      31  {
      32    int x2 = *x, y2 = *y;
      33    if (x2 != 2 || y2 != 3)
      34      __builtin_abort ();
      35    x2 = x2 + 2;
      36    y2 = y2 + 7;
      37  
      38    #pragma omp target device(ancestor : 1) map(tofrom: x2)
      39      check_offload(&x2, &y2);
      40  
      41    if (x2 != 2+2+3 || y2 != 3 + 7)
      42      __builtin_abort ();
      43    *x = x2, *y = y2;
      44  }
      45  #pragma omp end declare target
      46  
      47  void
      48  my_func (int *x, int *y)
      49  {
      50    if (global_var != 5)
      51      __builtin_abort ();
      52    global_var = 242;
      53    *x = 2*add_3(*x);
      54    *y = 3*add_3(*y);
      55  }
      56  
      57  int
      58  main ()
      59  {
      60    #pragma omp target
      61    {
      62       int x = 2, y = 3;
      63       tg_fn (&x, &y);
      64    }
      65  
      66    #pragma omp target
      67    {
      68       int x = -2, y = -1;
      69       #pragma omp target device ( ancestor:1 ) firstprivate(y) map(tofrom:x)
      70       {
      71         if (x != -2 || y != -1)
      72           __builtin_abort ();
      73         my_func (&x, &y);
      74         if (x != 2*(3-2) || y != 3*(3-1))
      75           __builtin_abort ();
      76       }
      77       if (x != 2*(3-2) || y != -1)
      78         __builtin_abort ();
      79    }
      80  
      81    if (global_var != 242)
      82      __builtin_abort ();
      83    return 0;
      84  }