1 // expected-no-diagnostics 2 #ifndef HEADER 3 #define HEADER 4 5 ///==========================================================================/// 6 // RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK1 --check-prefix CK1-64 7 // RUN: %clang_cc1 -DCK1 -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 %s --check-prefix CK1 --check-prefix CK1-64 9 // RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK1 --check-prefix CK1-32 10 // RUN: %clang_cc1 -DCK1 -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 %s --check-prefix CK1 --check-prefix CK1-32 12 #ifdef CK1 13 14 // CK1: [[ST:%.+]] = type { i32, double* } 15 template <typename T> 16 struct ST { 17 T a; 18 double *b; 19 }; 20 21 ST<int> gb; 22 double gc[100]; 23 24 // CK1: [[SIZE00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 800] 25 // CK1: [[MTYPE00:@.+]] = {{.+}}constant [1 x i32] [i32 34] 26 27 // CK1: [[SIZE02:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] 4] 28 // CK1: [[MTYPE02:@.+]] = {{.+}}constant [1 x i32] [i32 33] 29 30 // CK1: [[MTYPE03:@.+]] = {{.+}}constant [1 x i32] [i32 37] 31 32 // CK1: [[SIZE04:@.+]] = {{.+}}constant [2 x i[[sz]]] [i[[sz]] {{8|4}}, i[[sz]] 24] 33 // CK1: [[MTYPE04:@.+]] = {{.+}}constant [2 x i32] [i32 33, i32 17] 34 35 // CK1-LABEL: _Z3fooi 36 void foo(int arg) { 37 int la; 38 float lb[arg]; 39 40 // Region 00 41 // CK1-DAG: call void @__tgt_target_data_begin(i32 [[DEV:%[^,]+]], i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}) 42 // CK1-DAG: [[DEV]] = load i32, i32* %{{[^,]+}}, 43 // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 44 // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 45 46 // CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 47 // CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 48 // CK1-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [100 x double]** 49 // CK1-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [100 x double]** 50 // CK1-DAG: store [100 x double]* @gc, [100 x double]** [[CBP0]] 51 // CK1-DAG: store [100 x double]* @gc, [100 x double]** [[CP0]] 52 53 // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1 54 55 // CK1-DAG: call void @__tgt_target_data_end(i32 [[DEV:%[^,]+]], i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}) 56 // CK1-DAG: [[DEV]] = load i32, i32* %{{[^,]+}}, 57 // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]] 58 // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P]] 59 #pragma omp target data if(1+3-5) device(arg) map(from: gc) 60 {++arg;} 61 62 // Region 01 63 // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1 64 #pragma omp target data map(la) if(1+3-4) 65 {++arg;} 66 67 // Region 02 68 // CK1: br i1 %{{[^,]+}}, label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]] 69 // CK1: [[IFTHEN]] 70 // CK1-DAG: call void @__tgt_target_data_begin(i32 4, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE02]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE02]]{{.+}}) 71 // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 72 // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 73 74 // CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 75 // CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 76 // CK1-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32** 77 // CK1-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** 78 // CK1-DAG: store i32* [[VAR0:%.+]], i32** [[CBP0]] 79 // CK1-DAG: store i32* [[VAR0]], i32** [[CP0]] 80 // CK1: br label %[[IFEND:[^,]+]] 81 82 // CK1: [[IFELSE]] 83 // CK1: br label %[[IFEND]] 84 // CK1: [[IFEND]] 85 // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1 86 // CK1: br i1 %{{[^,]+}}, label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]] 87 88 // CK1: [[IFTHEN]] 89 // CK1-DAG: call void @__tgt_target_data_end(i32 4, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE02]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE02]]{{.+}}) 90 // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]] 91 // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P]] 92 // CK1: br label %[[IFEND:[^,]+]] 93 // CK1: [[IFELSE]] 94 // CK1: br label %[[IFEND]] 95 // CK1: [[IFEND]] 96 #pragma omp target data map(to: arg) if(arg) device(4) 97 {++arg;} 98 99 // Region 03 100 // CK1-DAG: call void @__tgt_target_data_begin(i32 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE03]]{{.+}}) 101 // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 102 // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 103 // CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] 104 105 // CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 106 // CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 107 // CK1-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 108 // CK1-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to float** 109 // CK1-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to float** 110 // CK1-DAG: store float* [[VAR0:%.+]], float** [[CBP0]] 111 // CK1-DAG: store float* [[VAR0]], float** [[CP0]] 112 // CK1-DAG: store i[[sz]] [[CSVAL0:%[^,]+]], i[[sz]]* [[S0]] 113 // CK1-DAG: [[CSVAL0]] = mul nuw i[[sz]] %{{[^,]+}}, 4 114 // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1 115 116 // CK1-DAG: call void @__tgt_target_data_end(i32 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE03]]{{.+}}) 117 // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]] 118 // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P]] 119 // CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S]] 120 #pragma omp target data map(always, to: lb) 121 {++arg;} 122 123 // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1 124 {++arg;} 125 126 // Region 04 127 // CK1-DAG: call void @__tgt_target_data_begin(i32 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE04]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE04]]{{.+}}) 128 // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 129 // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 130 131 // CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 132 // CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 133 // CK1-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]** 134 // CK1-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to double*** 135 // CK1-DAG: store [[ST]]* @gb, [[ST]]** [[CBP0]] 136 // CK1-DAG: store double** getelementptr inbounds ([[ST]], [[ST]]* @gb, i32 0, i32 1), double*** [[CP0]] 137 138 139 // CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 140 // CK1-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 141 // CK1-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to double*** 142 // CK1-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to double** 143 // CK1-DAG: store double** getelementptr inbounds ([[ST]], [[ST]]* @gb, i32 0, i32 1), double*** [[CBP1]] 144 // CK1-DAG: store double* [[SEC1:%.+]], double** [[CP1]] 145 // CK1-DAG: [[SEC1]] = getelementptr inbounds {{.+}}double* [[SEC11:%[^,]+]], i{{.+}} 0 146 // CK1-DAG: [[SEC11]] = load double*, double** getelementptr inbounds ([[ST]], [[ST]]* @gb, i32 0, i32 1), 147 148 // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1 149 150 // CK1-DAG: call void @__tgt_target_data_end(i32 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE04]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE04]]{{.+}}) 151 // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]] 152 // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P]] 153 #pragma omp target data map(to: gb.b[:3]) 154 {++arg;} 155 } 156 #endif 157 ///==========================================================================/// 158 // RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK2 --check-prefix CK2-64 159 // RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 160 // 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 %s --check-prefix CK2 --check-prefix CK2-64 161 // RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK2 --check-prefix CK2-32 162 // RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 163 // 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 %s --check-prefix CK2 --check-prefix CK2-32 164 #ifdef CK2 165 166 // CK2: [[ST:%.+]] = type { i32, double* } 167 template <typename T> 168 struct ST { 169 T a; 170 double *b; 171 172 T foo(T arg) { 173 // Region 00 174 #pragma omp target data map(always, to: b[1:3]) if(a>123) device(arg) 175 {arg++;} 176 return arg; 177 } 178 }; 179 180 // CK2: [[SIZE00:@.+]] = {{.+}}constant [2 x i[[sz:64|32]]] [i{{64|32}} {{8|4}}, i{{64|32}} 24] 181 // CK2: [[MTYPE00:@.+]] = {{.+}}constant [2 x i32] [i32 37, i32 21] 182 183 // CK2-LABEL: _Z3bari 184 int bar(int arg){ 185 ST<int> A; 186 return A.foo(arg); 187 } 188 189 // Region 00 190 // CK2: br i1 %{{[^,]+}}, label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]] 191 // CK2: [[IFTHEN]] 192 // CK2-DAG: call void @__tgt_target_data_begin(i32 [[DEV:%[^,]+]], i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}}) 193 // CK2-DAG: [[DEV]] = load i32, i32* %{{[^,]+}}, 194 // CK2-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 195 // CK2-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 196 197 // CK2-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 198 // CK2-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 199 // CK2-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]** 200 // CK2-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to double*** 201 // CK2-DAG: store [[ST]]* [[VAR0:%.+]], [[ST]]** [[CBP0]] 202 // CK2-DAG: store double** [[SEC0:%.+]], double*** [[CP0]] 203 // CK2-DAG: [[SEC0]] = getelementptr inbounds {{.*}}[[ST]]* [[VAR0]], i32 0, i32 1 204 205 206 // CK2-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 207 // CK2-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 208 // CK2-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to double*** 209 // CK2-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to double** 210 // CK2-DAG: store double** [[SEC0]], double*** [[CBP1]] 211 // CK2-DAG: store double* [[SEC1:%.+]], double** [[CP1]] 212 // CK2-DAG: [[SEC1]] = getelementptr inbounds {{.*}}double* [[SEC11:%[^,]+]], i{{.+}} 1 213 // CK2-DAG: [[SEC11]] = load double*, double** [[SEC111:%[^,]+]], 214 // CK2-DAG: [[SEC111]] = getelementptr inbounds {{.*}}[[ST]]* [[VAR0]], i32 0, i32 1 215 216 // CK2: br label %[[IFEND:[^,]+]] 217 218 // CK2: [[IFELSE]] 219 // CK2: br label %[[IFEND]] 220 // CK2: [[IFEND]] 221 // CK2: %{{.+}} = add nsw i32 %{{[^,]+}}, 1 222 // CK2: br i1 %{{[^,]+}}, label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]] 223 224 // CK2: [[IFTHEN]] 225 // CK2-DAG: call void @__tgt_target_data_end(i32 [[DEV:%[^,]+]], i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}}) 226 // CK2-DAG: [[DEV]] = load i32, i32* %{{[^,]+}}, 227 // CK2-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]] 228 // CK2-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P]] 229 // CK2: br label %[[IFEND:[^,]+]] 230 // CK2: [[IFELSE]] 231 // CK2: br label %[[IFEND]] 232 // CK2: [[IFEND]] 233 #endif 234 ///==========================================================================/// 235 // RUN: %clang_cc1 -DCK3 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK3 --check-prefix CK3-64 236 // RUN: %clang_cc1 -DCK3 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 237 // RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK3 --check-prefix CK3-64 238 // RUN: %clang_cc1 -DCK3 -verify -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK3 --check-prefix CK3-32 239 // RUN: %clang_cc1 -DCK3 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 240 // RUN: %clang_cc1 -fopenmp -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK3 --check-prefix CK3-32 241 #ifdef CK3 242 243 // CK3-LABEL: no_target_devices 244 void no_target_devices(int arg) { 245 // CK3-NOT: tgt_target_data_begin 246 // CK3: %{{.+}} = add nsw i32 %{{[^,]+}}, 1 247 // CK3-NOT: tgt_target_data_end 248 // CK3: ret 249 #pragma omp target data map(to: arg) if(arg) device(4) 250 {++arg;} 251 } 252 #endif 253 #endif 254