offload/test/offloading/target-teams-atomic.c

   1 // Check that omp atomic is permitted and behaves when strictly nested within
   2 // omp target teams.  This is an extension to OpenMP 5.2 and is enabled by
   3 // default.
   4
   5 // RUN: %libomptarget-compile-run-and-check-generic
   6
   7 #include <omp.h>
   8 #include <stdbool.h>
   9 #include <stdio.h>
  10 #include <string.h>
  11
  12 // High parallelism increases our chances of detecting a lack of atomicity.
  13 #define NUM_TEAMS_TRY 256
  14
  15 int main() {
  16   //      CHECK: update: num_teams=[[#NUM_TEAMS:]]{{$}}
  17   // CHECK-NEXT: update: x=[[#NUM_TEAMS]]{{$}}
  18   int x = 0;
  19   int numTeams;
  20 #pragma omp target teams num_teams(NUM_TEAMS_TRY) map(tofrom : x, numTeams)
  21   {
  22 #pragma omp atomic update
  23     ++x;
  24     if (omp_get_team_num() == 0)
  25       numTeams = omp_get_num_teams();
  26   }
  27   printf("update: num_teams=%d\n", numTeams);
  28   printf("update: x=%d\n", x);
  29
  30   // CHECK-NEXT: capture: x=[[#NUM_TEAMS]]{{$}}
  31   // CHECK-NEXT: capture: xCapturedCount=[[#NUM_TEAMS]]{{$}}
  32   bool xCaptured[numTeams];
  33   memset(xCaptured, 0, sizeof xCaptured);
  34   x = 0;
  35 #pragma omp target teams num_teams(NUM_TEAMS_TRY) map(tofrom : x, numTeams)
  36   {
  37     int v;
  38 #pragma omp atomic capture
  39     v = x++;
  40     xCaptured[v] = true;
  41   }
  42   printf("capture: x=%d\n", x);
  43   int xCapturedCount = 0;
  44   for (int i = 0; i < numTeams; ++i) {
  45     if (xCaptured[i])
  46       ++xCapturedCount;
  47   }
  48   printf("capture: xCapturedCount=%d\n", xCapturedCount);
  49   return 0;
  50 }