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()56int 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()65int bar() { return 1; } 66 67 #pragma omp declare variant(bazzz) match(device = {kind(CORRECT)}) baz()68int baz() { return 1; } 69 70 #pragma omp declare variant(test) match(device = {kind(CORRECT)}) call()71int call() { return 1; } 72 73 #pragma omp declare variant(stat_unused_) match(device = {kind(CORRECT)}) stat_unused()74static int stat_unused() { return 1; } 75 76 #pragma omp declare variant(stat_used_) match(device = {kind(CORRECT)}) stat_used()77static int stat_used() { return 1; } 78 79 #pragma omp end declare target 80 main()81int main() { 82 int res; 83 #pragma omp target map(from \ 84 : res) 85 res = bar() + baz() + call(); 86 return res; 87 } 88 test()89int test() { return 3; } stat_unused_()90static int stat_unused_() { return 4; } stat_used_()91static 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()109int 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_()127int SpecSpecialFuncs::method_() { return 7; } Method()128int SpecSpecialFuncs::Method() { return 1; } 129 prio()130int prio() { return 81; } prio1()131int prio1() { return 82; } prio2()132static int prio2() { return 83; } prio3()133static int prio3() { return 84; } prio4()134static int prio4() { return 84; } fn_linkage_variant()135int fn_linkage_variant() { return 85; } fn_linkage_variant1()136extern "C" int fn_linkage_variant1() { return 86; } fn_variant2()137int fn_variant2() { return 1; } 138 139 #pragma omp declare target 140 xxx()141void 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_()148int 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_()153static int prio1_() { return 1; } 154 int_fn()155int int_fn() { return prio1_(); } 156 157 extern "C" { 158 #pragma omp declare variant(fn_linkage_variant) match(device = {kind(CORRECT)}) fn_linkage()159int fn_linkage() { return 1; } 160 } 161 162 #pragma omp declare variant(fn_linkage_variant1) match(device = {kind(CORRECT)}) fn_linkage1()163int fn_linkage1() { return 1; } 164 165 #pragma omp declare variant(fn_variant2) match(device = {kind(WRONG)}) fn2()166int fn2() { return 87; } 167 168 #pragma omp end declare target 169 170 #endif // HEADER 171