offload/test/offloading/barrier_fence.c

   1 // RUN: %libomptarget-compile-generic -fopenmp-offload-mandatory -O3
   2 // RUN: %libomptarget-run-generic
   3 // RUN: %libomptarget-compileopt-generic -fopenmp-offload-mandatory -O3
   4 // RUN: %libomptarget-run-generic
   5
   6 // REQUIRES: gpu
   7
   8 #include <omp.h>
   9 #include <stdio.h>
  10
  11 struct IdentTy;
  12 void __kmpc_barrier_simple_spmd(struct IdentTy *Loc, int32_t TId);
  13 void __kmpc_barrier_simple_generic(struct IdentTy *Loc, int32_t TId);
  14
  15 #pragma omp begin declare target device_type(nohost)
  16 static int A[512] __attribute__((address_space(3), loader_uninitialized));
  17 static int B[512 * 32] __attribute__((loader_uninitialized));
  18 #pragma omp end declare target
  19
  20 int main() {
  21   printf("Testing simple spmd barrier\n");
  22   for (int r = 0; r < 50; r++) {
  23 #pragma omp target teams distribute thread_limit(512) num_teams(440)
  24     for (int j = 0; j < 512 * 32; ++j) {
  25 #pragma omp parallel firstprivate(j)
  26       {
  27         int TId = omp_get_thread_num();
  28         int TeamId = omp_get_team_num();
  29         int NT = omp_get_num_threads();
  30         // Sequential
  31         for (int i = 0; i < NT; ++i) {
  32           // Test shared memory globals
  33           if (TId == i)
  34             A[i] = i + j;
  35           __kmpc_barrier_simple_spmd(0, TId);
  36           if (A[i] != i + j)
  37             __builtin_trap();
  38           __kmpc_barrier_simple_spmd(0, TId);
  39           // Test generic globals
  40           if (TId == i)
  41             B[TeamId] = i;
  42           __kmpc_barrier_simple_spmd(0, TId);
  43           if (B[TeamId] != i)
  44             __builtin_trap();
  45           __kmpc_barrier_simple_spmd(0, TId);
  46         }
  47       }
  48     }
  49   }
  50
  51   printf("Testing simple generic barrier\n");
  52   for (int r = 0; r < 50; r++) {
  53 #pragma omp target teams distribute thread_limit(512) num_teams(440)
  54     for (int j = 0; j < 512 * 32; ++j) {
  55 #pragma omp parallel firstprivate(j)
  56       {
  57         int TId = omp_get_thread_num();
  58         int TeamId = omp_get_team_num();
  59         int NT = omp_get_num_threads();
  60         // Sequential
  61         for (int i = 0; i < NT; ++i) {
  62           if (TId == i)
  63             A[i] = i + j;
  64           __kmpc_barrier_simple_generic(0, TId);
  65           if (A[i] != i + j)
  66             __builtin_trap();
  67           __kmpc_barrier_simple_generic(0, TId);
  68           if (TId == i)
  69             B[TeamId] = i;
  70           __kmpc_barrier_simple_generic(0, TId);
  71           if (B[TeamId] != i)
  72             __builtin_trap();
  73           __kmpc_barrier_simple_generic(0, TId);
  74         }
  75       }
  76     }
  77   }
  78   return 0;
  79 }