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