xref: /llvm-project/offload/test/api/ompx_sync.c (revision 330d8983d25d08580fc1642fea48b2473f47a9da)
1*330d8983SJohannes Doerfert // RUN: %libomptarget-compile-run-and-check-generic
2*330d8983SJohannes Doerfert 
3*330d8983SJohannes Doerfert #include <omp.h>
4*330d8983SJohannes Doerfert #include <ompx.h>
5*330d8983SJohannes Doerfert #include <stdio.h>
6*330d8983SJohannes Doerfert 
foo(int device)7*330d8983SJohannes Doerfert void foo(int device) {
8*330d8983SJohannes Doerfert   int X;
9*330d8983SJohannes Doerfert   // clang-format off
10*330d8983SJohannes Doerfert #pragma omp target teams map(from: X) device(device) thread_limit(2) num_teams(1)
11*330d8983SJohannes Doerfert #pragma omp parallel
12*330d8983SJohannes Doerfert   // clang-format on
13*330d8983SJohannes Doerfert   {
14*330d8983SJohannes Doerfert     int tid = ompx_thread_id_x();
15*330d8983SJohannes Doerfert     int bid = ompx_block_id_x();
16*330d8983SJohannes Doerfert     if (tid == 1 && bid == 0) {
17*330d8983SJohannes Doerfert       X = 42;
18*330d8983SJohannes Doerfert       ompx_sync_block_divergent(3);
19*330d8983SJohannes Doerfert     } else {
20*330d8983SJohannes Doerfert       ompx_sync_block_divergent(1);
21*330d8983SJohannes Doerfert     }
22*330d8983SJohannes Doerfert     if (tid == 0 && bid == 0)
23*330d8983SJohannes Doerfert       X++;
24*330d8983SJohannes Doerfert     ompx_sync_block(ompx_seq_cst);
25*330d8983SJohannes Doerfert     if (tid == 1 && bid == 0)
26*330d8983SJohannes Doerfert       X++;
27*330d8983SJohannes Doerfert     ompx_sync_block_acq_rel();
28*330d8983SJohannes Doerfert     if (tid == 0 && bid == 0)
29*330d8983SJohannes Doerfert       X++;
30*330d8983SJohannes Doerfert     ompx_sync_block(ompx_release);
31*330d8983SJohannes Doerfert     if (tid == 0 && bid == 0)
32*330d8983SJohannes Doerfert       X++;
33*330d8983SJohannes Doerfert   }
34*330d8983SJohannes Doerfert   // CHECK: X: 46
35*330d8983SJohannes Doerfert   // CHECK: X: 46
36*330d8983SJohannes Doerfert   printf("X: %i\n", X);
37*330d8983SJohannes Doerfert }
38*330d8983SJohannes Doerfert 
main()39*330d8983SJohannes Doerfert int main() {
40*330d8983SJohannes Doerfert   foo(omp_get_default_device());
41*330d8983SJohannes Doerfert   foo(omp_get_initial_device());
42*330d8983SJohannes Doerfert }
43