1  /* { dg-do run } */
       2  /* { dg-require-effective-target offload_device_nonshared_as } */
       3  
       4  #include <stdlib.h>
       5  #include <omp.h>
       6  
       7  #define EPS 0.000001
       8  #define THRESHOLD 1000
       9  
      10  const int MAX = 1800;
      11  
      12  void check (float *a, float *b, int N)
      13  {
      14    int i;
      15    for (i = 0; i < N; i++)
      16      if (a[i] - b[i] > EPS || b[i] - a[i] > EPS)
      17        abort ();
      18  }
      19  
      20  void init (float *a1, float *a2, int N)
      21  {
      22    float s = -1;
      23    int i;
      24    for (i = 0; i < N; i++)
      25      {
      26        a1[i] = s;
      27        a2[i] = i;
      28        s = -s;
      29      }
      30  }
      31  
      32  void init_again (float *a1, float *a2, int N)
      33  {
      34    float s = -1;
      35    int i;
      36    for (i = 0; i < N; i++)
      37      {
      38        a1[i] = s * 10;
      39        a2[i] = i;
      40        s = -s;
      41      }
      42  }
      43  
      44  void vec_mult_ref (float *p, float *v1, float *v2, int N)
      45  {
      46    int i;
      47  
      48    init (v1, v2, N);
      49  
      50    for (i = 0; i < N; i++)
      51      p[i] = v1[i] * v2[i];
      52  
      53    init_again (v1, v2, N);
      54  
      55    for (i = 0; i < N; i++)
      56      p[i] = p[i] + (v1[i] * v2[i]);
      57  }
      58  
      59  void vec_mult (float *p, float *v1, float *v2, int N)
      60  {
      61    int i;
      62  
      63    init (v1, v2, N);
      64  
      65    #pragma omp target data if(N > THRESHOLD) map(from: p[0:N])
      66      {
      67        #pragma omp target if (N > THRESHOLD) map(to: v1[:N], v2[:N])
      68  	{
      69  	  if (omp_is_initial_device ())
      70  	    abort;
      71  
      72  	  #pragma omp parallel for
      73  	    for (i = 0; i < N; i++)
      74  	      p[i] = v1[i] * v2[i];
      75  	}
      76  
      77        init_again (v1, v2, N);
      78  
      79        #pragma omp target if (N > THRESHOLD) map(to: v1[:N], v2[:N])
      80  	{
      81  	  if (omp_is_initial_device ())
      82  	    abort ();
      83  
      84  	  #pragma omp parallel for
      85  	    for (i = 0; i < N; i++)
      86  	      p[i] = p[i] + (v1[i] * v2[i]);
      87  	}
      88      }
      89  }
      90  
      91  int main ()
      92  {
      93    float *p1 = (float *) malloc (MAX * sizeof (float));
      94    float *p2 = (float *) malloc (MAX * sizeof (float));
      95    float *v1 = (float *) malloc (MAX * sizeof (float));
      96    float *v2 = (float *) malloc (MAX * sizeof (float));
      97  
      98    vec_mult_ref (p1, v1, v2, MAX);
      99    vec_mult (p2, v1, v2, MAX);
     100  
     101    check (p1, p2, MAX);
     102  
     103    free (p1);
     104    free (p2);
     105    free (v1);
     106    free (v2);
     107  
     108    return 0;
     109  }