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