1 // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc -DGPU
2 // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -DGPU | FileCheck %s --implicit-check-not='ret i32 {{6|7|9|10|12|14|17|18|20|21|22|23|24|26}}'
3 // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -emit-pch -o %t -DGPU
4 // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -o - -DGPU | FileCheck %s --implicit-check-not='ret i32 {{6|7|9|10|12|14|17|18|20|21|22|23|24|26}}'
5 
6 // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc -DNOHOST
7 // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -DNOHOST | FileCheck %s --implicit-check-not='ret i32 {{6|7|9|10|12|14|17|18|20|21|22|23|24|26}}'
8 // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -emit-pch -o %t -DNOHOST
9 // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -o - -DNOHOST | FileCheck %s --implicit-check-not='ret i32 {{6|7|9|10|12|14|17|18|20|21|22|23|24|26}}'
10 
11 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc -DGPU
12 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -DGPU | FileCheck %s --implicit-check-not='ret i32 {{6|7|9|10|12|14|17|18|20|21|22|23|24|26}}'
13 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -emit-pch -o %t -DGPU
14 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -o - -DGPU | FileCheck %s --implicit-check-not='ret i32 {{6|7|9|10|12|14|17|18|20|21|22|23|24|26}}'
15 
16 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc -DNOHOST
17 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -DNOHOST | FileCheck %s --implicit-check-not='ret i32 {{6|7|9|10|12|14|17|18|20|21|22|23|24|26}}'
18 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -emit-pch -o %t -DNOHOST
19 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -o - -DNOHOST | FileCheck %s --implicit-check-not='ret i32 {{6|7|9|10|12|14|17|18|20|21|22|23|24|26}}'
20 // expected-no-diagnostics
21 
22 // CHECK-DAG: ret i32 2
23 // CHECK-DAG: ret i32 3
24 // CHECK-DAG: ret i32 4
25 // CHECK-DAG: ret i32 5
26 // CHECK-DAG: ret i32 8
27 // CHECK-DAG: ret i32 11
28 // CHECK-DAG: ret i32 13
29 // CHECK-DAG: ret i32 15
30 // CHECK-DAG: ret i32 16
31 // CHECK-DAG: ret i32 19
32 // CHECK-DAG: ret i32 25
33 
34 // Outputs for function members checked via implicit filecheck flag
35 
36 
37 #ifndef HEADER
38 #define HEADER
39 
40 #ifdef GPU
41 #define SUBSET gpu
42 #define CORRECT nohost, gpu
43 #define WRONG cpu, gpu
44 #endif // GPU
45 #ifdef NOHOST
46 #define SUBSET nohost
47 #define CORRECT nohost, gpu
48 #define WRONG nohost, host
49 #endif // NOHOST
50 
foo()51 int foo() { return 2; }
52 int bazzz();
53 int test();
54 static int stat_unused_();
55 static int stat_used_();
56 
57 #pragma omp declare target
58 
59 #pragma omp declare variant(foo) match(device = {kind(CORRECT)})
bar()60 int bar() { return 3; }
61 
62 #pragma omp declare variant(bazzz) match(device = {kind(CORRECT)})
baz()63 int baz() { return 4; }
64 
65 #pragma omp declare variant(test) match(device = {kind(CORRECT)})
call()66 int call() { return 5; }
67 
68 #pragma omp declare variant(stat_unused_) match(device = {kind(CORRECT)})
stat_unused()69 static int stat_unused() { return 6; }
70 
71 #pragma omp declare variant(stat_used_) match(device = {kind(CORRECT)})
stat_used()72 static int stat_used() { return 7; }
73 
74 #pragma omp end declare target
75 
main()76 int main() {
77   int res;
78 #pragma omp target map(from \
79                        : res)
80   res = bar() + baz() + call();
81   return res;
82 }
83 
test()84 int test() { return 8; }
stat_unused_()85 static int stat_unused_() { return 9; }
stat_used_()86 static int stat_used_() { return 10; }
87 
88 #pragma omp declare target
89 
90 struct SpecialFuncs {
vdSpecialFuncs91   void vd() {}
92   SpecialFuncs();
93   ~SpecialFuncs();
method_SpecialFuncs94   int method_() { return 11; }
95 #pragma omp declare variant(SpecialFuncs::method_) \
96     match(device = {kind(CORRECT)})
methodSpecialFuncs97   int method() { return 12; }
98 #pragma omp declare variant(SpecialFuncs::method_) \
99     match(device = {kind(CORRECT)})
100   int Method();
101 } s;
102 
Method()103 int SpecialFuncs::Method() { return 13; }
104 
105 struct SpecSpecialFuncs {
vdSpecSpecialFuncs106   void vd() {}
107   SpecSpecialFuncs();
108   ~SpecSpecialFuncs();
109 
110   int method_();
111 #pragma omp declare variant(SpecSpecialFuncs::method_) \
112     match(device = {kind(CORRECT)})
methodSpecSpecialFuncs113   int method() { return 14; }
114 #pragma omp declare variant(SpecSpecialFuncs::method_) \
115     match(device = {kind(CORRECT)})
116   int Method();
117 } s1;
118 
119 #pragma omp end declare target
120 
method_()121 int SpecSpecialFuncs::method_() { return 15; }
Method()122 int SpecSpecialFuncs::Method() { return 16; }
123 
prio()124 int prio() { return 17; }
prio1()125 int prio1() { return 18; }
prio2()126 static int prio2() { return 19; }
prio3()127 static int prio3() { return 20; }
prio4()128 static int prio4() { return 21; }
fn_linkage_variant()129 int fn_linkage_variant() { return 22; }
fn_linkage_variant1()130 extern "C" int fn_linkage_variant1() { return 23; }
fn_variant2()131 int fn_variant2() { return 24; }
132 
133 #pragma omp declare target
134 
xxx()135 void xxx() {
136   (void)s.method();
137   (void)s1.method();
138 }
139 
140 #pragma omp declare variant(prio) match(device = {kind(SUBSET)})
141 #pragma omp declare variant(prio1) match(device = {kind(CORRECT)})
prio_()142 int prio_() { return 25; }
143 
144 #pragma omp declare variant(prio4) match(device = {kind(SUBSET)})
145 #pragma omp declare variant(prio2) match(device = {kind(CORRECT)})
146 #pragma omp declare variant(prio3) match(device = {kind(SUBSET)})
prio1_()147 static int prio1_() { return 26; }
148 
int_fn()149 int int_fn() { return prio1_(); }
150 
151 extern "C" {
152 #pragma omp declare variant(fn_linkage_variant) match(device = {kind(CORRECT)})
fn_linkage()153 int fn_linkage() { return 27; }
154 }
155 
156 #pragma omp declare variant(fn_linkage_variant1) match(device = {kind(CORRECT)})
fn_linkage1()157 int fn_linkage1() { return 28; }
158 
159 #pragma omp declare variant(fn_variant2) match(device = {kind(WRONG)})
fn2()160 int fn2() { return 29; }
161 
162 #pragma omp end declare target
163 
164 #endif // HEADER
165