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

source code of offload/test/offloading/barrier_fence.c