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()51int 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()60int bar() { return 3; } 61 62 #pragma omp declare variant(bazzz) match(device = {kind(CORRECT)}) baz()63int baz() { return 4; } 64 65 #pragma omp declare variant(test) match(device = {kind(CORRECT)}) call()66int call() { return 5; } 67 68 #pragma omp declare variant(stat_unused_) match(device = {kind(CORRECT)}) stat_unused()69static int stat_unused() { return 6; } 70 71 #pragma omp declare variant(stat_used_) match(device = {kind(CORRECT)}) stat_used()72static int stat_used() { return 7; } 73 74 #pragma omp end declare target 75 main()76int main() { 77 int res; 78 #pragma omp target map(from \ 79 : res) 80 res = bar() + baz() + call(); 81 return res; 82 } 83 test()84int test() { return 8; } stat_unused_()85static int stat_unused_() { return 9; } stat_used_()86static 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()103int 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_()121int SpecSpecialFuncs::method_() { return 15; } Method()122int SpecSpecialFuncs::Method() { return 16; } 123 prio()124int prio() { return 17; } prio1()125int prio1() { return 18; } prio2()126static int prio2() { return 19; } prio3()127static int prio3() { return 20; } prio4()128static int prio4() { return 21; } fn_linkage_variant()129int fn_linkage_variant() { return 22; } fn_linkage_variant1()130extern "C" int fn_linkage_variant1() { return 23; } fn_variant2()131int fn_variant2() { return 24; } 132 133 #pragma omp declare target 134 xxx()135void 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_()142int 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_()147static int prio1_() { return 26; } 148 int_fn()149int int_fn() { return prio1_(); } 150 151 extern "C" { 152 #pragma omp declare variant(fn_linkage_variant) match(device = {kind(CORRECT)}) fn_linkage()153int fn_linkage() { return 27; } 154 } 155 156 #pragma omp declare variant(fn_linkage_variant1) match(device = {kind(CORRECT)}) fn_linkage1()157int fn_linkage1() { return 28; } 158 159 #pragma omp declare variant(fn_variant2) match(device = {kind(WRONG)}) fn2()160int fn2() { return 29; } 161 162 #pragma omp end declare target 163 164 #endif // HEADER 165