1 // expected-no-diagnostics 2 #ifndef HEADER 3 #define HEADER 4 5 ///==========================================================================/// 6 // RUN: %clang_cc1 -DCK1 -verify -Wno-vla -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 -Wno-vla %s -emit-llvm -o - | FileCheck %s --check-prefix CK1 --check-prefix CK1-64 9 // RUN: %clang_cc1 -DCK1 -verify -Wno-vla -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 -Wno-vla %s -emit-llvm -o - | FileCheck %s --check-prefix CK1 --check-prefix CK1-32 12 13 // RUN: %clang_cc1 -DCK1 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s 14 // RUN: %clang_cc1 -DCK1 -fopenmp-simd -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-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s 16 // RUN: %clang_cc1 -DCK1 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s 17 // RUN: %clang_cc1 -DCK1 -fopenmp-simd -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-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s 19 // SIMD-ONLY0-NOT: {{__kmpc|__tgt}} 20 #ifdef CK1 21 22 // CK1: [[ST:%.+]] = type { i32, ptr } 23 template <typename T> 24 struct ST { 25 T a; 26 double *b; 27 }; 28 29 ST<int> gb; 30 double gc[100]; 31 32 // CK1: [[IDENT_T:%.+]] = type { i32, i32, i32, i32, ptr } 33 // CK1: [[KMP_TASK_T_WITH_PRIVATES:%.+]] = type { [[KMP_TASK_T:%[^,]+]], [[KMP_PRIVATES_T:%.+]] } 34 // CK1: [[KMP_TASK_T]] = type { ptr, ptr, i32, %{{[^,]+}}, %{{[^,]+}} } 35 // CK1-32: [[KMP_PRIVATES_T]] = type { [1 x i64], [1 x ptr], [1 x ptr] } 36 // CK1-64: [[KMP_PRIVATES_T]] = type { [1 x ptr], [1 x ptr], [1 x i64] } 37 38 // CK1: [[SIZE00:@.+]] = {{.+}}constant [1 x i64] [i64 800] 39 // CK1: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 2] 40 41 // CK1: [[SIZE02:@.+]] = {{.+}}constant [1 x i64] [i64 4] 42 // CK1: [[MTYPE02:@.+]] = {{.+}}constant [1 x i64] zeroinitializer 43 44 // CK1: [[MTYPE03:@.+]] = {{.+}}constant [1 x i64] [i64 6] 45 46 // CK1: [[SIZE04:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 24] 47 // CK1: [[MTYPE04:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 281474976710672] 48 49 // CK1: [[MTYPE05:@.+]] = {{.+}}constant [1 x i64] [i64 1026] 50 51 // CK1: [[MTYPE06:@.+]] = {{.+}}constant [1 x i64] [i64 1030] 52 53 // CK1-LABEL: _Z3fooi 54 void foo(int arg) { 55 int la; 56 float lb[arg]; 57 58 // Region 00 59 // CK1-NOT: __tgt_target_data_begin 60 // CK1-DAG: call i32 @__kmpc_omp_task(ptr @{{[^,]+}}, i32 %{{[^,]+}}, ptr [[TASK:%.+]]) 61 // CK1-DAG: [[TASK]] = call ptr @__kmpc_omp_target_task_alloc(ptr @{{[^,]+}}, i32 %{{[^,]+}}, i32 1, i[[sz:32|64]] {{36|64}}, i{{32|64}} 4, ptr [[OMP_TASK_ENTRY:@[^,]+]], i64 [[DEV:%.+]]) 62 // CK1-DAG: [[DEV]] = sext i32 [[DEV32:%.+]] to i64 63 // CK1-DAG: [[PRIVATES:%.+]] = getelementptr inbounds nuw [[KMP_TASK_T_WITH_PRIVATES]], ptr [[TASK]], i32 0, i32 1 64 // CK1-32-DAG: [[FPBPGEP:%.+]] = getelementptr inbounds nuw [[KMP_PRIVATES_T]], ptr [[PRIVATES]], i32 0, i32 1 65 // CK1-64-DAG: [[FPBPGEP:%.+]] = getelementptr inbounds nuw [[KMP_PRIVATES_T]], ptr [[PRIVATES]], i32 0, i32 0 66 // CK1-DAG: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{4|8}} [[FPBPGEP]], ptr align {{4|8}} [[BPGEP:%.+]], i[[sz]] {{4|8}}, i1 false) 67 // CK1-DAG: [[BPGEP]] = getelementptr inbounds [1 x ptr], ptr [[BP:%.+]], i32 0, i32 0 68 // CK1-DAG: [[BPGEP:%.+]] = getelementptr inbounds [1 x ptr], ptr [[BP]], i32 0, i32 0 69 // CK1-DAG: store ptr [[GC:@[^,]+]], ptr [[BPGEP]], align 70 // CK1-32-DAG: [[FPPGEP:%.+]] = getelementptr inbounds nuw [[KMP_PRIVATES_T]], ptr [[PRIVATES]], i32 0, i32 2 71 // CK1-64-DAG: [[FPPGEP:%.+]] = getelementptr inbounds nuw [[KMP_PRIVATES_T]], ptr [[PRIVATES]], i32 0, i32 1 72 // CK1-DAG: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{4|8}} [[FPPGEP]], ptr align {{4|8}} [[PGEP:%.+]], i[[sz]] {{4|8}}, i1 false) 73 // CK1-DAG: [[PGEP]] = getelementptr inbounds [1 x ptr], ptr [[P:%.+]], i32 0, i32 0 74 // CK1-DAG: [[PGEP:%.+]] = getelementptr inbounds [1 x ptr], ptr [[P]], i32 0, i32 0 75 // CK1-DAG: store ptr [[GC]], ptr [[PGEP]], align 76 // CK1-32-DAG: [[FPSZGEP:%.+]] = getelementptr inbounds nuw [[KMP_PRIVATES_T]], ptr [[PRIVATES]], i32 0, i32 0 77 // CK1-64-DAG: [[FPSZGEP:%.+]] = getelementptr inbounds nuw [[KMP_PRIVATES_T]], ptr [[PRIVATES]], i32 0, i32 2 78 // CK1-DAG: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{4|8}} [[FPSZGEP]], ptr align {{4|8}} [[SIZE00]], i[[sz]] {{4|8}}, i1 false) 79 #pragma omp target exit data if(1+3-5) device(arg) map(from: gc) nowait 80 {++arg;} 81 82 // Region 01 83 // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1 84 #pragma omp target exit data map(release: la) if(1+3-4) 85 {++arg;} 86 87 // Region 02 88 // CK1-NOT: __tgt_target_data_begin 89 // CK1: br i1 %{{[^,]+}}, label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]] 90 // CK1: [[IFTHEN]] 91 // CK1-DAG: call void @__tgt_target_data_end_mapper(ptr @{{.+}}, i64 4, i32 1, ptr [[GEPBP:%.+]], ptr [[GEPP:%.+]], ptr [[SIZE02]], ptr [[MTYPE02]]{{.+}}, ptr null) 92 // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 93 // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 94 95 // CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 96 // CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 97 // CK1-DAG: store ptr [[VAL0:%[^,]+]], ptr [[BP0]] 98 // CK1-DAG: store ptr [[VAL0]], ptr [[P0]] 99 // CK1: br label %[[IFEND:[^,]+]] 100 101 // CK1: [[IFELSE]] 102 // CK1: br label %[[IFEND]] 103 // CK1: [[IFEND]] 104 // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1 105 #pragma omp target exit data map(release: arg) if(arg) device(4) 106 {++arg;} 107 108 // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1 109 {++arg;} 110 111 // Region 03 112 // CK1-NOT: __tgt_target_data_begin 113 // CK1-DAG: call void @__tgt_target_data_end_mapper(ptr @{{.+}}, i64 -1, i32 1, ptr [[GEPBP:%.+]], ptr [[GEPP:%.+]], ptr [[GEPS:%.+]], ptr [[MTYPE03]]{{.+}}, ptr null) 114 // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 115 // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 116 // CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] 117 118 // CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 119 // CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 120 // CK1-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 121 // CK1-DAG: store ptr [[VAL0:%[^,]+]], ptr [[BP0]] 122 // CK1-DAG: store ptr [[VAL0]], ptr [[P0]] 123 // CK1-DAG: store i64 [[CSVAL0:%[^,]+]], ptr [[S0]] 124 // CK1-64-DAG: [[CSVAL0]] = mul nuw i64 %{{[^,]+}}, 4 125 // CK1-32-DAG: [[CSVAL0]] = sext i32 [[CSVAL032:%.+]] to i64 126 // CK1-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4 127 // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1 128 #pragma omp target exit data map(always, from: lb) 129 {++arg;} 130 131 // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1 132 {++arg;} 133 134 // Region 04 135 // CK1-NOT: __tgt_target_data_begin 136 // CK1-DAG: call void @__tgt_target_data_end_mapper(ptr @{{.+}}, i64 -1, i32 2, ptr [[GEPBP:%.+]], ptr [[GEPP:%.+]], ptr [[GEPS:%.+]], ptr [[MTYPE04]]{{.+}}, ptr null) 137 // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 138 // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 139 // CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]] 140 141 // CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 142 // CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 143 // CK1-DAG: [[PS0:%.+]] = getelementptr inbounds {{.+}}[[PS]], i{{.+}} 0, i{{.+}} 0 144 // CK1-DAG: store ptr @gb, ptr [[BP0]] 145 // CK1-DAG: store ptr getelementptr inbounds nuw ([[ST]], ptr @gb, i32 0, i32 1), ptr [[P0]] 146 // CK1-DAG: [[DIV:%.+]] = sdiv exact i64 sub (i64 ptrtoint (ptr getelementptr (ptr, ptr getelementptr inbounds nuw (%struct.ST, ptr @gb, i32 0, i32 1), i32 1) to i64), i64 ptrtoint (ptr getelementptr inbounds nuw (%struct.ST, ptr @gb, i32 0, i32 1) to i64)), ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64) 147 // CK1-DAG: store i64 [[DIV]], ptr [[PS0]], 148 149 150 // CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 151 // CK1-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 152 // CK1-DAG: store ptr getelementptr inbounds nuw ([[ST]], ptr @gb, i32 0, i32 1), ptr [[BP1]] 153 // CK1-DAG: store ptr [[SEC1:%[^,]+]], ptr [[P1]] 154 // CK1-DAG: [[SEC1]] = getelementptr inbounds {{.+}}ptr [[SEC11:%[^,]+]], i{{.+}} 0 155 // CK1-DAG: [[SEC11]] = load ptr, ptr getelementptr inbounds nuw ([[ST]], ptr @gb, i32 0, i32 1), 156 157 // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1 158 #pragma omp target exit data map(release: gb.b[:3]) 159 {++arg;} 160 161 // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1 162 {++arg;} 163 164 // Region 05 165 // CK1-NOT: __tgt_target_data_begin 166 // CK1-DAG: call void @__tgt_target_data_end_mapper(ptr @{{.+}}, i64 -1, i32 1, ptr [[GEPBP:%.+]], ptr [[GEPP:%.+]], ptr [[GEPS:%.+]], ptr [[MTYPE05]]{{.+}}, ptr null) 167 // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 168 // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 169 // CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] 170 171 // CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 172 // CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 173 // CK1-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 174 // CK1-DAG: store ptr [[VAL0:%[^,]+]], ptr [[BP0]] 175 // CK1-DAG: store ptr [[VAL0]], ptr [[P0]] 176 // CK1-DAG: store i64 [[CSVAL0:%[^,]+]], ptr [[S0]] 177 // CK1-64-DAG: [[CSVAL0]] = mul nuw i64 %{{[^,]+}}, 4 178 // CK1-32-DAG: [[CSVAL0]] = sext i32 [[CSVAL032:%.+]] to i64 179 // CK1-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4 180 // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1 181 #pragma omp target exit data map(close, from: lb) 182 {++arg;} 183 184 // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1 185 {++arg;} 186 187 // Region 06 188 // CK1-NOT: __tgt_target_data_begin 189 // CK1-DAG: call void @__tgt_target_data_end_mapper(ptr @{{.+}}, i64 -1, i32 1, ptr [[GEPBP:%.+]], ptr [[GEPP:%.+]], ptr [[GEPS:%.+]], ptr [[MTYPE06]]{{.+}}, ptr null) 190 // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 191 // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 192 // CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] 193 194 // CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 195 // CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 196 // CK1-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 197 // CK1-DAG: store ptr [[VAL0:%[^,]+]], ptr [[BP0]] 198 // CK1-DAG: store ptr [[VAL0]], ptr [[P0]] 199 // CK1-DAG: store i64 [[CSVAL0:%[^,]+]], ptr [[S0]] 200 // CK1-64-DAG: [[CSVAL0]] = mul nuw i64 %{{[^,]+}}, 4 201 // CK1-32-DAG: [[CSVAL0]] = sext i32 [[CSVAL032:%.+]] to i64 202 // CK1-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4 203 // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1 204 #pragma omp target exit data map(always close, from: lb) 205 {++arg;} 206 } 207 208 // CK1: define internal {{.*}}i32 [[OMP_TASK_ENTRY]](i32 {{.*}}%{{[^,]+}}, ptr noalias noundef %{{[^,]+}}) 209 // CK1-DAG: call void @__tgt_target_data_end_nowait_mapper(ptr @{{.+}}, i64 %{{[^,]+}}, i32 1, ptr [[BP:%[^,]+]], ptr [[P:%[^,]+]], ptr [[SZ:%[^,]+]], ptr [[MTYPE00]], ptr null, ptr null, i32 0, ptr null, i32 0, ptr null) 210 // CK1-DAG: [[BP]] = load ptr, ptr [[FPBPADDR:%[^,]+]], align 211 // CK1-DAG: [[P]] = load ptr, ptr [[FPPADDR:%[^,]+]], align 212 // CK1-DAG: [[SZ]] = load ptr, ptr [[FPSZADDR:%[^,]+]], align 213 // CK1-DAG: call void {{%.*}}(ptr %{{[^,]+}}, ptr [[FPBPADDR]], ptr [[FPPADDR]], ptr [[FPSZADDR]]) 214 // CK1: ret i32 0 215 // CK1: } 216 217 #endif 218 ///==========================================================================/// 219 // 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 220 // 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 221 // 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 222 // 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 223 // 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 224 // 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 225 226 // RUN: %clang_cc1 -DCK2 -verify -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY1 %s 227 // RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 228 // 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 --check-prefix SIMD-ONLY1 %s 229 // RUN: %clang_cc1 -DCK2 -verify -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY1 %s 230 // RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 231 // 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 --check-prefix SIMD-ONLY1 %s 232 // SIMD-ONLY1-NOT: {{__kmpc|__tgt}} 233 #ifdef CK2 234 235 // CK2: [[ST:%.+]] = type { i32, ptr } 236 template <typename T> 237 struct ST { 238 T a; 239 double *b; 240 241 T foo(T arg) { 242 // Region 00 243 #pragma omp target exit data map(always, release: b[1:3]) if(a>123) device(arg) 244 {arg++;} 245 return arg; 246 } 247 }; 248 249 // CK2: [[SIZES:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 24] 250 // CK2: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 281474976710676] 251 252 // CK2-LABEL: _Z3bari 253 int bar(int arg){ 254 ST<int> A; 255 return A.foo(arg); 256 } 257 258 // Region 00 259 // CK2-NOT: __tgt_target_data_begin 260 // CK2: br i1 %{{[^,]+}}, label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]] 261 // CK2: [[IFTHEN]] 262 // CK2-DAG: call void @__tgt_target_data_end_mapper(ptr @{{.+}}, i64 [[DEV:%[^,]+]], i32 2, ptr [[GEPBP:%.+]], ptr [[GEPP:%.+]], ptr [[GEPS:%.+]], ptr [[MTYPE00]]{{.+}}, ptr null) 263 // CK2-DAG: [[DEV]] = sext i32 [[DEVi32:%[^,]+]] to i64 264 // CK2-DAG: [[DEVi32]] = load i32, ptr %{{[^,]+}}, 265 // CK2-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 266 // CK2-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 267 // CK2-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]] 268 269 // CK2-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 270 // CK2-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 271 // CK2-DAG: [[PS0:%.+]] = getelementptr inbounds {{.+}}[[PS]], i{{.+}} 0, i{{.+}} 0 272 // CK2-DAG: store ptr [[VAR0:%[^,]+]], ptr [[BP0]] 273 // CK2-DAG: store ptr [[SEC0:%[^,]+]], ptr [[P0]] 274 // CK2-DAG: store i64 [[CSVAL0:%[^,]+]], ptr [[PS0]], 275 // CK2-DAG: [[SEC0]] = getelementptr inbounds {{.*}}ptr [[VAR0]], i32 0, i32 1 276 277 // CK2-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 278 // CK2-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 279 // CK2-DAG: store ptr [[SEC0]], ptr [[BP1]] 280 // CK2-DAG: store ptr [[SEC1:%[^,]+]], ptr [[P1]] 281 // CK2-DAG: [[SEC1]] = getelementptr inbounds {{.*}}ptr [[SEC11:%[^,]+]], i{{.+}} 1 282 // CK2-DAG: [[SEC11]] = load ptr, ptr [[SEC111:%[^,]+]], 283 // CK2-DAG: [[SEC111]] = getelementptr inbounds {{.*}}ptr [[VAR0]], i32 0, i32 1 284 285 // CK2: br label %[[IFEND:[^,]+]] 286 287 // CK2: [[IFELSE]] 288 // CK2: br label %[[IFEND]] 289 // CK2: [[IFEND]] 290 // CK2: %{{.+}} = add nsw i32 %{{[^,]+}}, 1 291 #endif 292 ///==========================================================================/// 293 // 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 294 // RUN: %clang_cc1 -DCK3 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 295 // 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 296 // RUN: %clang_cc1 -DCK3 -verify -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK3 --check-prefix CK3-32 297 // RUN: %clang_cc1 -DCK3 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 298 // 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 299 300 // RUN: %clang_cc1 -DCK3 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY2 %s 301 // RUN: %clang_cc1 -DCK3 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 302 // RUN: %clang_cc1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY2 %s 303 // RUN: %clang_cc1 -DCK3 -verify -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY2 %s 304 // RUN: %clang_cc1 -DCK3 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 305 // RUN: %clang_cc1 -fopenmp-simd -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY2 %s 306 // SIMD-ONLY2-NOT: {{__kmpc|__tgt}} 307 #ifdef CK3 308 309 // CK3-LABEL: no_target_devices 310 void no_target_devices(int arg) { 311 // CK3-NOT: tgt_target_data_begin 312 // CK3: %{{.+}} = add nsw i32 %{{[^,]+}}, 1 313 // CK3-NOT: tgt_target_data_end 314 // CK3: ret 315 #pragma omp target exit data map(from: arg) if(arg) device(4) 316 {++arg;} 317 } 318 #endif 319 ///==========================================================================/// 320 // RUN: %clang_cc1 -DCK4 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK4 --check-prefix CK4-64 321 // RUN: %clang_cc1 -DCK4 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 322 // 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 CK4 --check-prefix CK4-64 323 // RUN: %clang_cc1 -DCK4 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK4 --check-prefix CK4-32 324 // RUN: %clang_cc1 -DCK4 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 325 // 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 CK4 --check-prefix CK4-32 326 327 // RUN: %clang_cc1 -DCK4 -verify -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY1 %s 328 // RUN: %clang_cc1 -DCK4 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 329 // 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 --check-prefix SIMD-ONLY1 %s 330 // RUN: %clang_cc1 -DCK4 -verify -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY1 %s 331 // RUN: %clang_cc1 -DCK4 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 332 // 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 --check-prefix SIMD-ONLY1 %s 333 // SIMD-ONLY1-NOT: {{__kmpc|__tgt}} 334 #ifdef CK4 335 336 // CK4: [[STT:%.+]] = type { i32, ptr } 337 template <typename T> 338 struct STT { 339 T a; 340 double *b; 341 342 T foo(T arg) { 343 // Region 00 344 #pragma omp target exit data map(always close, release: b[1:3]) if(a>123) device(arg) 345 {arg++;} 346 return arg; 347 } 348 }; 349 350 // CK4: [[SIZES:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 24] 351 // CK4: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 281474976711700] 352 353 // CK4-LABEL: _Z3bari 354 int bar(int arg){ 355 STT<int> A; 356 return A.foo(arg); 357 } 358 359 // Region 00 360 // CK4-NOT: __tgt_target_data_begin 361 // CK4: br i1 %{{[^,]+}}, label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]] 362 // CK4: [[IFTHEN]] 363 // CK4-DAG: call void @__tgt_target_data_end_mapper(ptr @{{.+}}, i64 [[DEV:%[^,]+]], i32 2, ptr [[GEPBP:%.+]], ptr [[GEPP:%.+]], ptr [[GEPS:%.+]], ptr [[MTYPE00]]{{.+}}, ptr null) 364 // CK4-DAG: [[DEV]] = sext i32 [[DEVi32:%[^,]+]] to i64 365 // CK4-DAG: [[DEVi32]] = load i32, ptr %{{[^,]+}}, 366 // CK4-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 367 // CK4-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 368 // CK4-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]] 369 370 // CK4-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 371 // CK4-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 372 // CK4-DAG: [[PS0:%.+]] = getelementptr inbounds {{.+}}[[PS]], i{{.+}} 0, i{{.+}} 0 373 // CK4-DAG: store ptr [[VAR0:%[^,]+]], ptr [[BP0]] 374 // CK4-DAG: store ptr [[SEC0:%[^,]+]], ptr [[P0]] 375 // CK4-DAG: store i64 [[CSVAL0:%[^,]+]], ptr [[PS0]], 376 // CK4-DAG: [[SEC0]] = getelementptr inbounds {{.*}}ptr [[VAR0]], i32 0, i32 1 377 378 // CK4-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 379 // CK4-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 380 // CK4-DAG: store ptr [[SEC0]], ptr [[BP1]] 381 // CK4-DAG: store ptr [[SEC1:%[^,]+]], ptr [[P1]] 382 // CK4-DAG: [[SEC1]] = getelementptr inbounds {{.*}}ptr [[SEC11:%[^,]+]], i{{.+}} 1 383 // CK4-DAG: [[SEC11]] = load ptr, ptr [[SEC111:%[^,]+]], 384 // CK4-DAG: [[SEC111]] = getelementptr inbounds {{.*}}ptr [[VAR0]], i32 0, i32 1 385 386 // CK4: br label %[[IFEND:[^,]+]] 387 388 // CK4: [[IFELSE]] 389 // CK4: br label %[[IFEND]] 390 // CK4: [[IFEND]] 391 // CK4: %{{.+}} = add nsw i32 %{{[^,]+}}, 1 392 #endif 393 #endif 394