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 main()15int 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 } 51