// RUN: %libomptarget-compileopt-run-and-check-generic
//
// UNSUPPORTED: x86_64-pc-linux-gnu
// UNSUPPORTED: x86_64-pc-linux-gnu-LTO
// UNSUPPORTED: aarch64-unknown-linux-gnu
// UNSUPPORTED: aarch64-unknown-linux-gnu-LTO
// UNSUPPORTED: s390x-ibm-linux-gnu
// UNSUPPORTED: s390x-ibm-linux-gnu-LTO

#include <omp.h>
#include <ompx.h>
#include <stdio.h>
#include <stdlib.h>
#include <string.h>

struct info {
  int tid, bid, tdim;
};

int main(int argc, char **argv) {
  int N = 1 << 20;
  if (argc > 1)
    N = atoi(argv[1]);

  struct info *X = (struct info *)malloc(sizeof(*X) * N);
  memset(X, '0', sizeof(*X) * N);

  int TL = 256;
  int NT = (N + TL - 1) / TL;

#pragma omp target data map(tofrom : X [0:N])
#pragma omp target teams num_teams(NT) thread_limit(TL)
  {
#pragma omp parallel
    {
      int tid = ompx_thread_id_x();
      int bid = ompx_block_id_x();
      int tdim = ompx_block_dim_x();
      int gid = tid + bid * tdim;
      if (gid < N) {
        X[gid].tid = tid;
        X[gid].bid = bid;
        X[gid].tdim = tdim;
      };
    }
  }

  int tid = 0, bid = 0, tdim = 256;
  for (int i = 0; i < N; i++) {
    if (X[i].tid != tid || X[i].bid != bid || X[i].tdim != tdim) {
      printf("%i: %i vs %i, %i vs %i, %i vs %i\n", i, X[i].tid, tid, X[i].bid,
             bid, X[i].tdim, tdim);
      return 1;
    }
    tid++;
    if (tid == tdim) {
      tid = 0;
      bid++;
    }
  }

  // CHECK: OK
  printf("OK");
  return 0;
}