clang/test/OpenMP/target_map_codegen_14.cpp

   1 // expected-no-diagnostics
   2 #ifndef HEADER
   3 #define HEADER
   4
   5 ///==========================================================================///
   6 // RUN: %clang_cc1 -DCK15 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  %s --check-prefix CK15 --check-prefix CK15-64
   7 // RUN: %clang_cc1 -DCK15 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
   8 // RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap  %s  --check-prefix CK15 --check-prefix CK15-64
   9 // RUN: %clang_cc1 -DCK15 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  %s  --check-prefix CK15 --check-prefix CK15-32
  10 // RUN: %clang_cc1 -DCK15 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
  11 // RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap  %s  --check-prefix CK15 --check-prefix CK15-32
  12
  13 // RUN: %clang_cc1 -DCK15 -verify -fopenmp -fopenmp-version=45 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  %s --check-prefix CK15 --check-prefix CK15-64
  14 // RUN: %clang_cc1 -DCK15 -fopenmp -fopenmp-version=45 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
  15 // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap  %s  --check-prefix CK15 --check-prefix CK15-64
  16 // RUN: %clang_cc1 -DCK15 -verify -fopenmp -fopenmp-version=45 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  %s  --check-prefix CK15 --check-prefix CK15-32
  17 // RUN: %clang_cc1 -DCK15 -fopenmp -fopenmp-version=45 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
  18 // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap  %s  --check-prefix CK15 --check-prefix CK15-32
  19
  20 // RUN: %clang_cc1 -DCK15 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  %s --check-prefix CK15 --check-prefix CK15-64
  21 // RUN: %clang_cc1 -DCK15 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
  22 // RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap  %s  --check-prefix CK15 --check-prefix CK15-64
  23 // RUN: %clang_cc1 -DCK15 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  %s  --check-prefix CK15 --check-prefix CK15-32
  24 // RUN: %clang_cc1 -DCK15 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
  25 // RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap  %s  --check-prefix CK15 --check-prefix CK15-32
  26
  27 // RUN: %clang_cc1 -DCK15 -verify -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY14 %s
  28 // RUN: %clang_cc1 -DCK15 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
  29 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY14 %s
  30 // RUN: %clang_cc1 -DCK15 -verify -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY14 %s
  31 // RUN: %clang_cc1 -DCK15 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
  32 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY14 %s
  33 // SIMD-ONLY14-NOT: {{__kmpc|__tgt}}
  34 #ifdef CK15
  35
  36 // CK15: [[ST:%.+]] = type { i32, double }
  37 // CK15: [[SIZES:@.+]] = {{.+}}constant [4 x i64] [i64 0, i64 4, i64 8, i64 4]
  38 // Map types:
  39 // - OMP_MAP_TARGET_PARAM = 32
  40 // - OMP_MAP_TO + OMP_MAP_FROM | OMP_MAP_IMPLICIT | OMP_MAP_MEMBER_OF = 281474976711171
  41 // - OMP_MAP_PRIVATE_VAL + OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 800
  42 // CK15: [[TYPES:@.+]] = {{.+}}constant [4 x i64] [i64 32, i64 281474976711171, i64 281474976711171, i64 800]
  43
  44 // CK15: [[SIZES2:@.+]] = {{.+}}constant [4 x i64] [i64 0, i64 4, i64 8, i64 4]
  45 // Map types:
  46 // - OMP_MAP_TARGET_PARAM = 32
  47 // - OMP_MAP_TO + OMP_MAP_FROM | OMP_MAP_IMPLICIT | OMP_MAP_MEMBER_OF = 281474976711171
  48 // - OMP_MAP_PRIVATE_VAL + OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 800
  49 // CK15: [[TYPES2:@.+]] = {{.+}}constant [4 x i64] [i64 32, i64 281474976711171, i64 281474976711171, i64 800]
  50
  51 template<int x>
  52 class SSST {
  53 public:
  54   int a;
  55   double b;
  56
  57   void foo(int c) {
  58     #pragma omp target
  59     {
  60       a += c + x;
  61       b += (double)(c + x);
  62     }
  63   }
  64   template<int y>
  65   void bar(int c) {
  66     #pragma omp target
  67     {
  68       a += c + x + y;
  69       b += (double)(c + x + y);
  70     }
  71   }
  72
  73   SSST(int a, double b) : a(a), b(b) {}
  74 };
  75
  76 // CK15-LABEL: implicit_maps_templated_class{{.*}}(
  77 void implicit_maps_templated_class (int a){
  78   SSST<123> ssst(a, (double)a);
  79
  80   // CK15: define {{.*}}void @{{.+}}foo{{.+}}(ptr {{[^,]+}}, i32 {{[^,]+}})
  81   // CK15-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
  82   // CK15-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
  83   // CK15-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
  84   // CK15-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
  85   // CK15-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
  86   // CK15-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
  87   // CK15-DAG: store ptr [[SIZES:%.+]], ptr [[SARG]]
  88   // CK15-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
  89   // CK15-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
  90   // CK15-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]], i32 0, i32 0
  91
  92   // CK15-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
  93   // CK15-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
  94   // CK15-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i32 0, i32 0
  95   // CK15-DAG: store ptr [[DECL:%.+]], ptr [[BP0]]
  96   // CK15-DAG: store ptr [[A:%.+]], ptr [[P0]]
  97   // CK15-DAG: store i64 %{{.+}}, ptr [[S0]]
  98
  99   // CK15-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 1
 100   // CK15-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 1
 101   // CK15-DAG: store ptr [[DECL]], ptr [[BP1]]
 102   // CK15-DAG: store ptr [[A]], ptr [[P1]]
 103
 104   // CK15-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 2
 105   // CK15-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 2
 106   // CK15-DAG: store ptr [[DECL]], ptr [[BP2]]
 107   // CK15-DAG: store ptr %{{.+}}, ptr [[P2]]
 108
 109   // CK15-DAG: [[BP3:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 3
 110   // CK15-DAG: [[P3:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 3
 111   // CK15-DAG: store i[[sz:64|32]] [[VAL:%.+]], ptr [[BP3]]
 112   // CK15-DAG: store i[[sz]] [[VAL]], ptr [[P3]]
 113   // CK15-DAG: [[VAL]] = load i[[sz]], ptr [[ADDR:%.+]],
 114   // CK15-64-DAG: store i32 {{.+}}, ptr [[ADDR]],
 115
 116   // CK15: call void [[KERNEL:@.+]](ptr [[DECL]], i[[sz]] {{.+}})
 117   ssst.foo(456);
 118
 119   // CK15: define {{.*}}void @{{.+}}bar{{.+}}(ptr {{[^,]+}}, i32 {{[^,]+}})
 120   // CK15-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
 121   // CK15-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
 122   // CK15-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
 123   // CK15-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 124   // CK15-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
 125   // CK15-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
 126   // CK15-DAG: store ptr [[SIZES:%.+]], ptr [[SARG]]
 127   // CK15-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
 128   // CK15-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
 129   // CK15-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]], i32 0, i32 0
 130
 131   // CK15-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
 132   // CK15-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
 133   // CK15-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i32 0, i32 0
 134   // CK15-DAG: store ptr [[DECL:%.+]], ptr [[BP0]]
 135   // CK15-DAG: store ptr [[A:%.+]], ptr [[P0]]
 136   // CK15-DAG: store i64 %{{.+}}, ptr [[S0]]
 137
 138   // CK15-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 1
 139   // CK15-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 1
 140   // CK15-DAG: store ptr [[DECL]], ptr [[BP1]]
 141   // CK15-DAG: store ptr [[A]], ptr [[P1]]
 142
 143   // CK15-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 2
 144   // CK15-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 2
 145   // CK15-DAG: store ptr [[DECL]], ptr [[BP2]]
 146   // CK15-DAG: store ptr %{{.+}}, ptr [[P2]]
 147
 148   // CK15-DAG: [[BP3:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 3
 149   // CK15-DAG: [[P3:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 3
 150   // CK15-DAG: store i[[sz]] [[VAL:%.+]], ptr [[BP3]]
 151   // CK15-DAG: store i[[sz]] [[VAL]], ptr [[P3]]
 152   // CK15-DAG: [[VAL]] = load i[[sz]], ptr [[ADDR:%.+]],
 153   // CK15-64-DAG: store i32 {{.+}}, ptr [[ADDR]],
 154
 155   // CK15: call void [[KERNEL2:@.+]](ptr [[DECL]], i[[sz]] {{.+}})
 156   ssst.bar<210>(789);
 157 }
 158
 159 // CK15: define internal void [[KERNEL]](ptr noundef [[THIS:%.+]], i[[sz]] noundef [[ARG:%.+]])
 160 // CK15: [[ADDR0:%.+]] = alloca ptr,
 161 // CK15: [[ADDR1:%.+]] = alloca i[[sz]],
 162 // CK15: store ptr [[THIS]], ptr [[ADDR0]],
 163 // CK15: store i[[sz]] [[ARG]], ptr [[ADDR1]],
 164 // CK15: [[REF0:%.+]] = load ptr, ptr [[ADDR0]],
 165 // CK15-64: {{.+}} = load i32,  ptr [[ADDR1]],
 166 // CK15-32: {{.+}} = load i32, ptr [[ADDR1]],
 167 // CK15: {{.+}} = getelementptr inbounds [[ST]], ptr [[REF0]], i32 0, i32 0
 168
 169 // CK15: define internal void [[KERNEL2]](ptr noundef [[THIS:%.+]], i[[sz]] noundef [[ARG:%.+]])
 170 // CK15: [[ADDR0:%.+]] = alloca ptr,
 171 // CK15: [[ADDR1:%.+]] = alloca i[[sz]],
 172 // CK15: store ptr [[THIS]], ptr [[ADDR0]],
 173 // CK15: store i[[sz]] [[ARG]], ptr [[ADDR1]],
 174 // CK15: [[REF0:%.+]] = load ptr, ptr [[ADDR0]],
 175 // CK15-64: {{.+}} = load i32,  ptr [[ADDR1]],
 176 // CK15-32: {{.+}} = load i32, ptr [[ADDR1]],
 177 // CK15: {{.+}} = getelementptr inbounds [[ST]], ptr [[REF0]], i32 0, i32 0
 178
 179 #endif // CK15
 180 #endif