1  /* Test OpenACC's support for manual deep copy, including the attach
       2     and detach clauses.  */
       3  
       4  /* TODO The tree dump scanning has certain expectations.
       5     { dg-do compile { target { lp64 || llp64 } } } */
       6  /* { dg-additional-options "-fdump-tree-omplower" } */
       7  
       8  /* { dg-additional-options -Wuninitialized } */
       9  
      10  void
      11  t1 ()
      12  {
      13    struct foo {
      14      int *a, *b, c, d, *e;
      15    } s;
      16  
      17    int *a, *z;
      18    /* { dg-note {'z' was declared here} {} { target *-*-* } .-1 } */
      19  
      20  #pragma acc enter data copyin(s)
      21    {
      22  #pragma acc data copy(s.a[0:10]) copy(z[0:10])
      23      /* { dg-warning {'z' is used uninitialized} {} { target *-*-* } .-1 } */
      24      {
      25        s.e = z;
      26  #pragma acc parallel loop attach(s.e)
      27        for (int i = 0; i < 10; i++)
      28          s.a[i] = s.e[i];
      29  
      30  
      31        a = s.e;
      32  #pragma acc enter data attach(a)
      33  #pragma acc exit data detach(a)
      34      }
      35  
      36  #pragma acc enter data copyin(a)
      37  #pragma acc enter data attach(s.e)
      38  #pragma acc exit data detach(s.e)
      39  
      40  #pragma acc data attach(s.e)
      41      {
      42      }
      43  #pragma acc exit data delete(a)
      44  
      45  #pragma acc exit data detach(a) finalize
      46  #pragma acc exit data detach(s.a) finalize
      47    }
      48  }
      49  
      50  /* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_data map.to:s .len: 32.." 1 "omplower" } } */
      51  /* { dg-final { scan-tree-dump-times "pragma omp target oacc_data map.tofrom:.z .len: 40.. map.struct:s .len: 1.. map.alloc:s.a .len: 8.. map.tofrom:._1 .len: 40.. map.attach:s.a .bias: 0.." 1 "omplower" } } */
      52  /* { dg-final { scan-tree-dump-times "pragma omp target oacc_parallel map.attach:s.e .bias: 0.. map.tofrom:s .len: 32" 1 "omplower" } } */
      53  /* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_data map.attach:a .bias: 0.." 1 "omplower" } } */
      54  /* { dg-final { scan-tree-dump-times "pragma omp target oacc_exit_data map.detach:a .bias: 0.." 1 "omplower" } } */
      55  /* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_data map.to:a .len: 8.." 1 "omplower" } } */
      56  /* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_data map.attach:s.e .bias: 0.." 1 "omplower" } } */
      57  /* { dg-final { scan-tree-dump-times "pragma omp target oacc_exit_data map.detach:s.e .bias: 0.." 1 "omplower" } } */
      58  /* { dg-final { scan-tree-dump-times "pragma omp target oacc_data map.attach:s.e .bias: 0.." 1 "omplower" } } */
      59  /* { dg-final { scan-tree-dump-times "pragma omp target oacc_exit_data map.release:a .len: 8.." 1 "omplower" } } */
      60  /* { dg-final { scan-tree-dump-times "pragma omp target oacc_exit_data finalize map.force_detach:a .bias: 0.." 1 "omplower" } } */
      61  /* { dg-final { scan-tree-dump-times "pragma omp target oacc_exit_data finalize map.force_detach:s.a .bias: 0.." 1 "omplower" } } */