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
11struct IdentTy;
12void __kmpc_barrier_simple_spmd(struct IdentTy *Loc, int32_t TId);
13void __kmpc_barrier_simple_generic(struct IdentTy *Loc, int32_t TId);
14
15#pragma omp begin declare target device_type(nohost)
16static int A[512] __attribute__((address_space(3), loader_uninitialized));
17static int B[512 * 32] __attribute__((loader_uninitialized));
18#pragma omp end declare target
19
20int main() {
21 printf(format: "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(Loc: 0, TId);
36 if (A[i] != i + j)
37 __builtin_trap();
38 __kmpc_barrier_simple_spmd(Loc: 0, TId);
39 // Test generic globals
40 if (TId == i)
41 B[TeamId] = i;
42 __kmpc_barrier_simple_spmd(Loc: 0, TId);
43 if (B[TeamId] != i)
44 __builtin_trap();
45 __kmpc_barrier_simple_spmd(Loc: 0, TId);
46 }
47 }
48 }
49 }
50
51 printf(format: "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(Loc: 0, TId);
65 if (A[i] != i + j)
66 __builtin_trap();
67 __kmpc_barrier_simple_generic(Loc: 0, TId);
68 if (TId == i)
69 B[TeamId] = i;
70 __kmpc_barrier_simple_generic(Loc: 0, TId);
71 if (B[TeamId] != i)
72 __builtin_trap();
73 __kmpc_barrier_simple_generic(Loc: 0, TId);
74 }
75 }
76 }
77 }
78 return 0;
79}
80

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