/* { dg-require-effective-target offload_device_nonshared_as } */

#include <stdlib.h>
#include <assert.h>

#define N 32

void test_array_section (int *p)
{
  #pragma omp target data map(alloc: p[0:N])
    {
      int ok = 1;
      for (int i = 10; i < 10 + 4; i++)
	p[i] = 997 * i;

      #pragma omp target map(always to:p[10:4]) map(tofrom: ok)
	for (int i = 10; i < 10 + 4; i++)
	  if (p[i] != 997 * i)
	    ok = 0;

      assert (ok);

      #pragma omp target map(always from:p[7:9])
	for (int i = 0; i < N; i++)
	  p[i] = i;
    }
}

int main ()
{
  int aa = 0, bb = 0, cc = 0, dd = 0;

  #pragma omp target data map(tofrom: aa) map(to: bb) map(from: cc, dd)
    {
      int ok;
      aa = bb = cc = 1;

      /* Set dd on target to 0 for the further check.  */
      #pragma omp target map(always to: dd)
	;

      dd = 1;
      #pragma omp target map(tofrom: aa) map(always to: bb) \
	map(always from: cc) map(to: dd) map(from: ok)
	{
	  /* bb is always to, aa and dd are not.  */
	  ok = (aa == 0) && (bb == 1) && (dd == 0);
	  aa = bb = cc = dd = 2;
	}

      assert (ok);
      assert (aa == 1);
      assert (bb == 1);
      assert (cc == 2); /* cc is always from.  */
      assert (dd == 1);

      dd = 3;
      #pragma omp target map(from: cc) map(always to: dd) map(from: ok)
	{
	  ok = (dd == 3); /* dd is always to.  */
	  cc = dd = 4;
	}

      assert (ok);
      assert (cc == 2);
      assert (dd == 3);
    }

  assert (aa == 2);
  assert (bb == 1);
  assert (cc == 4);
  assert (dd == 4);

  int *array = calloc (N, sizeof (int));
  test_array_section (array);

  for (int i = 0; i < 7; i++)
    assert (array[i] == 0);
  for (int i = 7; i < 7 + 9; i++)
    assert (array[i] == i);
  for (int i = 7 + 9; i < N; i++)
    assert (array[i] == 0);

  free (array);
  return 0;
}