165 lines · cpp
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 -DGPU2// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-target-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 %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -emit-pch -o %t -DGPU4// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-target-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 -DNOHOST7// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-target-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 %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -emit-pch -o %t -DNOHOST9// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-target-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 -DGPU12// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-target-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 %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -emit-pch -o %t -DGPU14// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-target-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 -DNOHOST17// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-target-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 %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -emit-pch -o %t -DNOHOST19// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-target-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-diagnostics21 22// CHECK-DAG: ret i32 223// CHECK-DAG: ret i32 324// CHECK-DAG: ret i32 425// CHECK-DAG: ret i32 526// CHECK-DAG: ret i32 827// CHECK-DAG: ret i32 1128// CHECK-DAG: ret i32 1329// CHECK-DAG: ret i32 1530// CHECK-DAG: ret i32 1631// CHECK-DAG: ret i32 1932// CHECK-DAG: ret i32 2533 34// Outputs for function members checked via implicit filecheck flag35 36 37#ifndef HEADER38#define HEADER39 40#ifdef GPU41#define SUBSET gpu42#define CORRECT nohost, gpu43#define WRONG cpu, gpu44#endif // GPU45#ifdef NOHOST46#define SUBSET nohost47#define CORRECT nohost, gpu48#define WRONG nohost, host49#endif // NOHOST50 51int foo() { return 2; }52int bazzz();53int test();54static int stat_unused_();55static int stat_used_();56 57#pragma omp declare target58 59#pragma omp declare variant(foo) match(device = {kind(CORRECT)})60int bar() { return 3; }61 62#pragma omp declare variant(bazzz) match(device = {kind(CORRECT)})63int baz() { return 4; }64 65#pragma omp declare variant(test) match(device = {kind(CORRECT)})66int call() { return 5; }67 68#pragma omp declare variant(stat_unused_) match(device = {kind(CORRECT)})69static int stat_unused() { return 6; }70 71#pragma omp declare variant(stat_used_) match(device = {kind(CORRECT)})72static int stat_used() { return 7; }73 74#pragma omp end declare target75 76int main() {77 int res;78#pragma omp target map(from \79 : res)80 res = bar() + baz() + call();81 return res;82}83 84int test() { return 8; }85static int stat_unused_() { return 9; }86static int stat_used_() { return 10; }87 88#pragma omp declare target89 90struct SpecialFuncs {91 void vd() {}92 SpecialFuncs();93 ~SpecialFuncs();94 int method_() { return 11; }95#pragma omp declare variant(SpecialFuncs::method_) \96 match(device = {kind(CORRECT)})97 int method() { return 12; }98#pragma omp declare variant(SpecialFuncs::method_) \99 match(device = {kind(CORRECT)})100 int Method();101} s;102 103int SpecialFuncs::Method() { return 13; }104 105struct SpecSpecialFuncs {106 void vd() {}107 SpecSpecialFuncs();108 ~SpecSpecialFuncs();109 110 int method_();111#pragma omp declare variant(SpecSpecialFuncs::method_) \112 match(device = {kind(CORRECT)})113 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 target120 121int SpecSpecialFuncs::method_() { return 15; }122int SpecSpecialFuncs::Method() { return 16; }123 124int prio() { return 17; }125int prio1() { return 18; }126static int prio2() { return 19; }127static int prio3() { return 20; }128static int prio4() { return 21; }129int fn_linkage_variant() { return 22; }130extern "C" int fn_linkage_variant1() { return 23; }131int fn_variant2() { return 24; }132 133#pragma omp declare target134 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)})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)})147static int prio1_() { return 26; }148 149int int_fn() { return prio1_(); }150 151extern "C" {152#pragma omp declare variant(fn_linkage_variant) match(device = {kind(CORRECT)})153int fn_linkage() { return 27; }154}155 156#pragma omp declare variant(fn_linkage_variant1) match(device = {kind(CORRECT)})157int fn_linkage1() { return 28; }158 159#pragma omp declare variant(fn_variant2) match(device = {kind(WRONG)})160int fn2() { return 29; }161 162#pragma omp end declare target163 164#endif // HEADER165