1 // 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 -fopenmp-version=50 -DGPU
2 // 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 - -fopenmp-version=50 -DGPU | FileCheck %s --implicit-check-not='ret i32 {{1|81|84}}'
3 // 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 -fopenmp-version=50 -DGPU
4 // 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 - -fopenmp-version=50 -DGPU | FileCheck %s --implicit-check-not='ret i32 {{1|81|84}}'
5 
6 // 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 -fopenmp-version=50 -DNOHOST
7 // 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 - -fopenmp-version=50 -DNOHOST | FileCheck %s --implicit-check-not='ret i32 {{1|81|84}}'
8 // 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 -fopenmp-version=50 -DNOHOST
9 // 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 - -fopenmp-version=50 -DNOHOST | FileCheck %s --implicit-check-not='ret i32 {{1|81|84}}'
10 // expected-no-diagnostics
11 
12 // CHECK-NOT: ret i32 {{1|81|84}}
13 // CHECK-DAG: define {{.*}}i32 @_Z3barv()
14 // CHECK-DAG: define {{.*}}i32 @_ZN16SpecSpecialFuncs6MethodEv(%struct.SpecSpecialFuncs* %{{.+}})
15 // CHECK-DAG: define {{.*}}i32 @_ZN12SpecialFuncs6MethodEv(%struct.SpecialFuncs* %{{.+}})
16 // CHECK-DAG: define linkonce_odr {{.*}}i32 @_ZN16SpecSpecialFuncs6methodEv(%struct.SpecSpecialFuncs* %{{.+}})
17 // CHECK-DAG: define linkonce_odr {{.*}}i32 @_ZN12SpecialFuncs6methodEv(%struct.SpecialFuncs* %{{.+}})
18 // CHECK-DAG: define {{.*}}i32 @_Z5prio_v()
19 // CHECK-DAG: define internal i32 @_ZL6prio1_v()
20 // CHECK-DAG: define {{.*}}i32 @_Z4callv()
21 // CHECK-DAG: define internal i32 @_ZL9stat_usedv()
22 // CHECK-DAG: define {{.*}}i32 @fn_linkage()
23 // CHECK-DAG: define {{.*}}i32 @_Z11fn_linkage1v()
24 
25 // CHECK-DAG: ret i32 2
26 // CHECK-DAG: ret i32 3
27 // CHECK-DAG: ret i32 4
28 // CHECK-DAG: ret i32 5
29 // CHECK-DAG: ret i32 6
30 // CHECK-DAG: ret i32 7
31 // CHECK-DAG: ret i32 82
32 // CHECK-DAG: ret i32 83
33 // CHECK-DAG: ret i32 85
34 // CHECK-DAG: ret i32 86
35 // CHECK-DAG: ret i32 87
36 
37 // Outputs for function members
38 // CHECK-DAG: ret i32 6
39 // CHECK-DAG: ret i32 7
40 // CHECK-NOT: ret i32 {{1|81|84}}
41 
42 #ifndef HEADER
43 #define HEADER
44 
45 #ifdef GPU
46 #define CORRECT gpu
47 #define SUBSET nohost, gpu
48 #define WRONG cpu, gpu
49 #endif // GPU
50 #ifdef NOHOST
51 #define CORRECT nohost
52 #define SUBSET nohost, gpu
53 #define WRONG nohost, host
54 #endif // NOHOST
55 
foo()56 int foo() { return 2; }
57 int bazzz();
58 int test();
59 static int stat_unused_();
60 static int stat_used_();
61 
62 #pragma omp declare target
63 
64 #pragma omp declare variant(foo) match(device = {kind(CORRECT)})
bar()65 int bar() { return 1; }
66 
67 #pragma omp declare variant(bazzz) match(device = {kind(CORRECT)})
baz()68 int baz() { return 1; }
69 
70 #pragma omp declare variant(test) match(device = {kind(CORRECT)})
call()71 int call() { return 1; }
72 
73 #pragma omp declare variant(stat_unused_) match(device = {kind(CORRECT)})
stat_unused()74 static int stat_unused() { return 1; }
75 
76 #pragma omp declare variant(stat_used_) match(device = {kind(CORRECT)})
stat_used()77 static int stat_used() { return 1; }
78 
79 #pragma omp end declare target
80 
main()81 int main() {
82   int res;
83 #pragma omp target map(from \
84                        : res)
85   res = bar() + baz() + call();
86   return res;
87 }
88 
test()89 int test() { return 3; }
stat_unused_()90 static int stat_unused_() { return 4; }
stat_used_()91 static int stat_used_() { return 5; }
92 
93 #pragma omp declare target
94 
95 struct SpecialFuncs {
vdSpecialFuncs96   void vd() {}
97   SpecialFuncs();
98   ~SpecialFuncs();
99 
method_SpecialFuncs100   int method_() { return 6; }
101 #pragma omp declare variant(SpecialFuncs::method_) \
102     match(device = {kind(CORRECT)})
methodSpecialFuncs103   int method() { return 1; }
104 #pragma omp declare variant(SpecialFuncs::method_) \
105     match(device = {kind(CORRECT)})
106   int Method();
107 } s;
108 
Method()109 int SpecialFuncs::Method() { return 1; }
110 
111 struct SpecSpecialFuncs {
vdSpecSpecialFuncs112   void vd() {}
113   SpecSpecialFuncs();
114   ~SpecSpecialFuncs();
115 
116   int method_();
117 #pragma omp declare variant(SpecSpecialFuncs::method_) \
118     match(device = {kind(CORRECT)})
methodSpecSpecialFuncs119   int method() { return 1; }
120 #pragma omp declare variant(SpecSpecialFuncs::method_) \
121     match(device = {kind(CORRECT)})
122   int Method();
123 } s1;
124 
125 #pragma omp end declare target
126 
method_()127 int SpecSpecialFuncs::method_() { return 7; }
Method()128 int SpecSpecialFuncs::Method() { return 1; }
129 
prio()130 int prio() { return 81; }
prio1()131 int prio1() { return 82; }
prio2()132 static int prio2() { return 83; }
prio3()133 static int prio3() { return 84; }
prio4()134 static int prio4() { return 84; }
fn_linkage_variant()135 int fn_linkage_variant() { return 85; }
fn_linkage_variant1()136 extern "C" int fn_linkage_variant1() { return 86; }
fn_variant2()137 int fn_variant2() { return 1; }
138 
139 #pragma omp declare target
140 
xxx()141 void xxx() {
142   (void)s.method();
143   (void)s1.method();
144 }
145 
146 #pragma omp declare variant(prio) match(device = {kind(SUBSET)})
147 #pragma omp declare variant(prio1) match(device = {kind(CORRECT)})
prio_()148 int prio_() { return 1; }
149 
150 #pragma omp declare variant(prio4) match(device = {kind(SUBSET)})
151 #pragma omp declare variant(prio2) match(device = {kind(CORRECT)})
152 #pragma omp declare variant(prio3) match(device = {kind(SUBSET)})
prio1_()153 static int prio1_() { return 1; }
154 
int_fn()155 int int_fn() { return prio1_(); }
156 
157 extern "C" {
158 #pragma omp declare variant(fn_linkage_variant) match(device = {kind(CORRECT)})
fn_linkage()159 int fn_linkage() { return 1; }
160 }
161 
162 #pragma omp declare variant(fn_linkage_variant1) match(device = {kind(CORRECT)})
fn_linkage1()163 int fn_linkage1() { return 1; }
164 
165 #pragma omp declare variant(fn_variant2) match(device = {kind(WRONG)})
fn2()166 int fn2() { return 87; }
167 
168 #pragma omp end declare target
169 
170 #endif // HEADER
171