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