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
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
|
// RUN: %clang_cc1 -verify -fopenmp -triple x86_64-unknown-linux-gnu -x c -std=c99 -ast-print %s -o - | FileCheck %s
// RUN: %clang_cc1 -verify -fopenmp-simd -triple x86_64-unknown-linux-gnu -x c -std=c99 -ast-print %s -o - | FileCheck %s
// RUN: %clang_cc1 -verify -fopenmp -triple amdgcn-amd-amdhsa -x c -std=c99 -ast-print %s -o - | FileCheck %s --check-prefix=CHECK-AMDGCN
// RUN: %clang_cc1 -verify -fopenmp-simd -triple amdgcn-amd-amdhsa -x c -std=c99 -ast-print %s -o - | FileCheck %s --check-prefix=CHECK-AMDGCN
// expected-no-diagnostics
#ifndef HEADER
#define HEADER
void bar(void);
#define N 10
void foo(void) {
#pragma omp metadirective when(device = {kind(cpu)} \
: parallel) default()
bar();
#pragma omp metadirective when(implementation = {vendor(score(0) \
: llvm)}, \
device = {kind(cpu)} \
: parallel) default(target teams)
bar();
#pragma omp metadirective when(device = {kind(gpu)} \
: target teams) when(implementation = {vendor(llvm)} \
: parallel) default()
bar();
#pragma omp metadirective default(target) when(implementation = {vendor(score(5) \
: llvm)}, \
device = {kind(cpu, host)} \
: parallel)
bar();
#pragma omp metadirective when(user = {condition(N > 10)} \
: target) when(user = {condition(N == 10)} \
: parallel)
bar();
#pragma omp metadirective when(device = {kind(host)} \
: parallel for)
for (int i = 0; i < 100; i++)
;
#pragma omp metadirective when(implementation = {extension(match_all)} \
: parallel) default(parallel for)
for (int i = 0; i < 100; i++)
;
#pragma omp metadirective when(implementation = {extension(match_any)} \
: parallel) default(parallel for)
for (int i = 0; i < 100; i++)
;
#pragma omp metadirective when(implementation = {extension(match_none)} \
: parallel) default(parallel for)
for (int i = 0; i < 100; i++)
;
// Test metadirective with nested OpenMP directive.
int array[16];
#pragma omp metadirective when(user = {condition(1)} \
: parallel for)
for (int i = 0; i < 16; i++) {
#pragma omp simd
for (int j = 0; j < 16; j++)
array[i] = i;
}
#pragma omp metadirective when(device={arch("amdgcn")}: \
teams distribute parallel for)\
default(parallel for)
for (int i = 0; i < 100; i++)
;
#pragma omp metadirective when(implementation = {extension(match_all)} \
: nothing) default(parallel for)
for (int i = 0; i < 16; i++)
;
#pragma omp metadirective when(implementation = {extension(match_any)} \
: parallel) default(nothing)
for (int i = 0; i < 16; i++)
;
}
// CHECK: void bar(void);
// CHECK: void foo(void)
// CHECK-NEXT: #pragma omp parallel
// CHECK-NEXT: bar()
// CHECK-NEXT: #pragma omp parallel
// CHECK-NEXT: bar()
// CHECK-NEXT: #pragma omp parallel
// CHECK-NEXT: bar()
// CHECK-NEXT: #pragma omp parallel
// CHECK-NEXT: bar()
// CHECK-NEXT: #pragma omp parallel
// CHECK-NEXT: bar()
// CHECK-NEXT: #pragma omp parallel for
// CHECK-NEXT: for (int i = 0; i < 100; i++)
// CHECK: #pragma omp parallel
// CHECK-NEXT: for (int i = 0; i < 100; i++)
// CHECK: #pragma omp parallel for
// CHECK-NEXT: for (int i = 0; i < 100; i++)
// CHECK: #pragma omp parallel
// CHECK-NEXT: for (int i = 0; i < 100; i++)
// CHECK: #pragma omp parallel for
// CHECK-NEXT: for (int i = 0; i < 16; i++) {
// CHECK-NEXT: #pragma omp simd
// CHECK-NEXT: for (int j = 0; j < 16; j++)
// CHECK-AMDGCN: #pragma omp teams distribute parallel for
// CHECK-AMDGCN-NEXT: for (int i = 0; i < 100; i++)
// CHECK: for (int i = 0; i < 16; i++)
// CHECK: for (int i = 0; i < 16; i++)
#endif
|
