// RUN: %compile-run-and-check

#include <omp.h>
#include <stdio.h>

const int MaxThreads = 1024;
const int NumThreads = 64;

int main(int argc, char *argv[]) {
  int level = -1, activeLevel = -1;
  // The expected value is -1, initialize to different value.
  int ancestorTNumNeg = 1, teamSizeNeg = 1;
  int ancestorTNum0 = -1, teamSize0 = -1;
  // The expected value is -1, initialize to different value.
  int ancestorTNum1 = 1, teamSize1 = 1;
  int check1[MaxThreads];
  int check2[MaxThreads];
  int check3[MaxThreads];
  int check4[MaxThreads];
  for (int i = 0; i < MaxThreads; i++) {
    check1[i] = check2[i] = check3[i] = check4[i] = 0;
  }

  #pragma omp target map(level, activeLevel, ancestorTNumNeg, teamSizeNeg) \
                     map(ancestorTNum0, teamSize0, ancestorTNum1, teamSize1) \
                     map(check1[:], check2[:], check3[:], check4[:])
  {
    level = omp_get_level();
    activeLevel = omp_get_active_level();

    // Expected to return -1.
    ancestorTNumNeg = omp_get_ancestor_thread_num(-1);
    teamSizeNeg = omp_get_team_size(-1);

    // Expected to return 0 and 1.
    ancestorTNum0 = omp_get_ancestor_thread_num(0);
    teamSize0 = omp_get_team_size(0);

    // Expected to return -1 because the requested level is larger than
    // the nest level.
    ancestorTNum1 = omp_get_ancestor_thread_num(1);
    teamSize1 = omp_get_team_size(1);

    // Expecting active parallel region.
    #pragma omp parallel num_threads(NumThreads)
    {
      int id = omp_get_thread_num();
      // Multiply return value of omp_get_level by 5 to avoid that this test
      // passes if both API calls return wrong values.
      check1[id] += omp_get_level() * 5 + omp_get_active_level();

      // Expected to return 0 and 1.
      check2[id] += omp_get_ancestor_thread_num(0) + 5 * omp_get_team_size(0);
      // Expected to return the current thread num.
      check2[id] += (omp_get_ancestor_thread_num(1) - id);
      // Exepcted to return the current number of threads.
      check2[id] += 3 * omp_get_team_size(1);
      // Expected to return -1, see above.
      check2[id] += omp_get_ancestor_thread_num(2) + omp_get_team_size(2);

      // Expecting serialized parallel region.
      #pragma omp parallel
      {
        #pragma omp atomic
        check3[id] += omp_get_level() * 5 + omp_get_active_level();

        // Expected to return 0 and 1.
        int check4Inc = omp_get_ancestor_thread_num(0) + 5 * omp_get_team_size(0);
        // Expected to return the parent thread num.
        check4Inc += (omp_get_ancestor_thread_num(1) - id);
        // Exepcted to return the number of threads in the active parallel region.
        check4Inc += 3 * omp_get_team_size(1);
        // Exptected to return 0 and 1.
        check4Inc += omp_get_ancestor_thread_num(2) + 3 * omp_get_team_size(2);
        // Expected to return -1, see above.
        check4Inc += omp_get_ancestor_thread_num(3) + omp_get_team_size(3);

        #pragma omp atomic
        check4[id] += check4Inc;
      }
    }
  }

  // CHECK: target: level = 0, activeLevel = 0
  printf("target: level = %d, activeLevel = %d\n", level, activeLevel);
  // CHECK: level = -1: ancestorTNum = -1, teamSize = -1
  printf("level = -1: ancestorTNum = %d, teamSize = %d\n", ancestorTNumNeg, teamSizeNeg);
  // CHECK: level = 0: ancestorTNum = 0, teamSize = 1
  printf("level = 0: ancestorTNum = %d, teamSize = %d\n", ancestorTNum0, teamSize0);
  // CHECK: level = 1: ancestorTNum = -1, teamSize = -1
  printf("level = 1: ancestorTNum = %d, teamSize = %d\n", ancestorTNum1, teamSize1);

  // CHECK-NOT: invalid
  for (int i = 0; i < MaxThreads; i++) {
    // Check active parallel region:
    // omp_get_level() = 1, omp_get_active_level() = 1
    const int Expected1 = 6;
    if (i < NumThreads) {
      if (check1[i] != Expected1) {
        printf("invalid: check1[%d] should be %d, is %d\n", i, Expected1, check1[i]);
      }
    } else if (check1[i] != 0) {
      printf("invalid: check1[%d] should be 0, is %d\n", i, check1[i]);
    }

    // 5 * 1 + 3 * 64 - 1 - 1 (see above)
    const int Expected2 = 195;
    if (i < NumThreads) {
      if (check2[i] != Expected2) {
        printf("invalid: check2[%d] should be %d, is %d\n", i, Expected2, check2[i]);
      }
    } else if (check2[i] != 0) {
      printf("invalid: check2[%d] should be 0, is %d\n", i, check2[i]);
    }

    // Check serialized parallel region:
    // omp_get_level() = 2, omp_get_active_level() = 1
    const int Expected3 = 11;
    if (i < NumThreads) {
      if (check3[i] != Expected3) {
        printf("invalid: check3[%d] should be %d, is %d\n", i, Expected3, check3[i]);
      }
    } else if (check3[i] != 0) {
      printf("invalid: check3[%d] should be 0, is %d\n", i, check3[i]);
    }

    // 5 * 1 + 3 * 64 + 3 * 1 - 1 - 1 (see above)
    const int Expected4 = 198;
    if (i < NumThreads) {
      if (check4[i] != Expected4) {
        printf("invalid: check4[%d] should be %d, is %d\n", i, Expected4, check4[i]);
      }
    } else if (check4[i] != 0) {
      printf("invalid: check4[%d] should be 0, is %d\n", i, check4[i]);
    }
  }

  // Check for paraller level in non-SPMD kernels.
  level = 0;
  #pragma omp target teams distribute num_teams(1) thread_limit(32) reduction(+:level)
  for (int i=0; i<5032; i+=32) {
    int ub = (i+32 > 5032) ? 5032 : i+32;
    #pragma omp parallel for schedule(dynamic)
    for (int j=i ; j < ub; j++) ;
    level += omp_get_level();
  }
  // CHECK: Integral level = 0.
  printf("Integral level = %d.\n", level);

  return 0;
}