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
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(format: "update: num_teams=%d\n", numTeams);
28 printf(format: "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(s: xCaptured, c: 0, n: 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(format: "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(format: "capture: xCapturedCount=%d\n", xCapturedCount);
49 return 0;
50}
51

source code of offload/test/offloading/target-teams-atomic.c