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