brintos

brintos / llvm-project-archived public Read only

0
0
Text · 1.2 KiB · 7cca784 Raw
61 lines · c
1// RUN: %libomptarget-compileopt-run-and-check-generic2//3// REQUIRES: gpu4 5#include <omp.h>6#include <ompx.h>7#include <stdio.h>8#include <stdlib.h>9#include <string.h>10 11struct info {12  int tid, bid, tdim;13};14 15int main(int argc, char **argv) {16  int N = 1 << 20;17  if (argc > 1)18    N = atoi(argv[1]);19 20  struct info *X = (struct info *)malloc(sizeof(*X) * N);21  memset(X, '0', sizeof(*X) * N);22 23  int TL = 256;24  int NT = (N + TL - 1) / TL;25 26#pragma omp target data map(tofrom : X [0:N])27#pragma omp target teams num_teams(NT) thread_limit(TL)28  {29#pragma omp parallel30    {31      int tid = ompx_thread_id_x();32      int bid = ompx_block_id_x();33      int tdim = ompx_block_dim_x();34      int gid = tid + bid * tdim;35      if (gid < N) {36        X[gid].tid = tid;37        X[gid].bid = bid;38        X[gid].tdim = tdim;39      };40    }41  }42 43  int tid = 0, bid = 0, tdim = 256;44  for (int i = 0; i < N; i++) {45    if (X[i].tid != tid || X[i].bid != bid || X[i].tdim != tdim) {46      printf("%i: %i vs %i, %i vs %i, %i vs %i\n", i, X[i].tid, tid, X[i].bid,47             bid, X[i].tdim, tdim);48      return 1;49    }50    tid++;51    if (tid == tdim) {52      tid = 0;53      bid++;54    }55  }56 57  // CHECK: OK58  printf("OK");59  return 0;60}61