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