1 // expected-no-diagnostics 2 #ifndef HEADER 3 #define HEADER 4 5 /// 6 /// Implicit maps. 7 /// 8 9 ///==========================================================================/// 10 // RUN: %clang_cc1 -DCK1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK1 --check-prefix CK1-64 11 // RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 12 // 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 CK1 --check-prefix CK1-64 13 // RUN: %clang_cc1 -DCK1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK1 --check-prefix CK1-32 14 // RUN: %clang_cc1 -DCK1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 15 // 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 CK1 --check-prefix CK1-32 16 #ifdef CK1 17 18 // CK1-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4] 19 // Map types: OMP_MAP_BYCOPY = 128 20 // CK1-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 128] 21 22 // CK1-LABEL: implicit_maps_integer 23 void implicit_maps_integer (int a){ 24 int i = a; 25 26 // CK1-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}}) 27 // CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 28 // CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 29 // CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 30 // CK1-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 31 // CK1-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]], 32 // CK1-DAG: store i8* [[VALP:%.+]], i8** [[P1]], 33 // CK1-DAG: [[VALBP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8* 34 // CK1-DAG: [[VALP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8* 35 // CK1-DAG: [[VAL]] = load i[[sz]], i[[sz]]* [[ADDR:%.+]], 36 // CK1-64-DAG: [[CADDR:%.+]] = bitcast i[[sz]]* [[ADDR]] to i32* 37 // CK1-64-DAG: store i32 {{.+}}, i32* [[CADDR]], 38 39 // CK1: call void [[KERNEL:@.+]](i[[sz]] [[VAL]]) 40 #pragma omp target 41 { 42 ++i; 43 } 44 } 45 46 // CK1: define internal void [[KERNEL]](i[[sz]] [[ARG:%.+]]) 47 // CK1: [[ADDR:%.+]] = alloca i[[sz]], 48 // CK1: store i[[sz]] [[ARG]], i[[sz]]* [[ADDR]], 49 // CK1-64: [[CADDR:%.+]] = bitcast i64* [[ADDR]] to i32* 50 // CK1-64: {{.+}} = load i32, i32* [[CADDR]], 51 // CK1-32: {{.+}} = load i32, i32* [[ADDR]], 52 53 #endif 54 ///==========================================================================/// 55 // RUN: %clang_cc1 -DCK2 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK2 --check-prefix CK2-64 56 // RUN: %clang_cc1 -DCK2 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 57 // 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 CK2 --check-prefix CK2-64 58 // RUN: %clang_cc1 -DCK2 -verify -fopenmp -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK2 --check-prefix CK2-32 59 // RUN: %clang_cc1 -DCK2 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 60 // 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 CK2 --check-prefix CK2-32 61 #ifdef CK2 62 63 // CK2-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4] 64 // Map types: OMP_MAP_BYCOPY = 128 65 // CK2-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 128] 66 67 // CK2-LABEL: implicit_maps_integer_reference 68 void implicit_maps_integer_reference (int a){ 69 int &i = a; 70 // CK2-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}}) 71 // CK2-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 72 // CK2-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 73 // CK2-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 74 // CK2-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 75 // CK2-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]], 76 // CK2-DAG: store i8* [[VALP:%.+]], i8** [[P1]], 77 // CK2-DAG: [[VALBP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8* 78 // CK2-DAG: [[VALP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8* 79 // CK2-DAG: [[VAL]] = load i[[sz]], i[[sz]]* [[ADDR:%.+]], 80 // CK2-64-DAG: [[CADDR:%.+]] = bitcast i[[sz]]* [[ADDR]] to i32* 81 // CK2-64-DAG: store i32 {{.+}}, i32* [[CADDR]], 82 83 // CK2: call void [[KERNEL:@.+]](i[[sz]] [[VAL]]) 84 #pragma omp target 85 { 86 ++i; 87 } 88 } 89 90 // CK2: define internal void [[KERNEL]](i[[sz]] [[ARG:%.+]]) 91 // CK2: [[ADDR:%.+]] = alloca i[[sz]], 92 // CK2: [[REF:%.+]] = alloca i32*, 93 // CK2: store i[[sz]] [[ARG]], i[[sz]]* [[ADDR]], 94 // CK2-64: [[CADDR:%.+]] = bitcast i[[sz]]* [[ADDR]] to i32* 95 // CK2-64: store i32* [[CADDR]], i32** [[REF]], 96 // CK2-64: [[RVAL:%.+]] = load i32*, i32** [[REF]], 97 // CK2-64: {{.+}} = load i32, i32* [[RVAL]], 98 // CK2-32: store i32* [[ADDR]], i32** [[REF]], 99 // CK2-32: [[RVAL:%.+]] = load i32*, i32** [[REF]], 100 // CK2-32: {{.+}} = load i32, i32* [[RVAL]], 101 102 #endif 103 ///==========================================================================/// 104 // 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 105 // RUN: %clang_cc1 -DCK3 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 106 // 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 107 // RUN: %clang_cc1 -DCK3 -verify -fopenmp -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK3 --check-prefix CK3-32 108 // RUN: %clang_cc1 -DCK3 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 109 // 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 110 #ifdef CK3 111 112 // CK3-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4] 113 // Map types: OMP_MAP_BYCOPY = 128 114 // CK3-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 128] 115 116 // CK3-LABEL: implicit_maps_parameter 117 void implicit_maps_parameter (int a){ 118 119 // CK3-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}}) 120 // CK3-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 121 // CK3-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 122 // CK3-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 123 // CK3-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 124 // CK3-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]], 125 // CK3-DAG: store i8* [[VALP:%.+]], i8** [[P1]], 126 // CK3-DAG: [[VALBP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8* 127 // CK3-DAG: [[VALP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8* 128 // CK3-DAG: [[VAL]] = load i[[sz]], i[[sz]]* [[ADDR:%.+]], 129 // CK3-64-DAG: [[CADDR:%.+]] = bitcast i[[sz]]* [[ADDR]] to i32* 130 // CK3-64-DAG: store i32 {{.+}}, i32* [[CADDR]], 131 132 // CK3: call void [[KERNEL:@.+]](i[[sz]] [[VAL]]) 133 #pragma omp target 134 { 135 ++a; 136 } 137 } 138 139 // CK3: define internal void [[KERNEL]](i[[sz]] [[ARG:%.+]]) 140 // CK3: [[ADDR:%.+]] = alloca i[[sz]], 141 // CK3: store i[[sz]] [[ARG]], i[[sz]]* [[ADDR]], 142 // CK3-64: [[CADDR:%.+]] = bitcast i64* [[ADDR]] to i32* 143 // CK3-64: {{.+}} = load i32, i32* [[CADDR]], 144 // CK3-32: {{.+}} = load i32, i32* [[ADDR]], 145 146 #endif 147 ///==========================================================================/// 148 // RUN: %clang_cc1 -DCK4 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK4 --check-prefix CK4-64 149 // RUN: %clang_cc1 -DCK4 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 150 // 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 CK4 --check-prefix CK4-64 151 // RUN: %clang_cc1 -DCK4 -verify -fopenmp -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK4 --check-prefix CK4-32 152 // RUN: %clang_cc1 -DCK4 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 153 // 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 CK4 --check-prefix CK4-32 154 #ifdef CK4 155 156 // CK4-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4] 157 // Map types: OMP_MAP_BYCOPY = 128 158 // CK4-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 128] 159 160 // CK4-LABEL: implicit_maps_nested_integer 161 void implicit_maps_nested_integer (int a){ 162 int i = a; 163 164 // The captures in parallel are by reference. Only the capture in target is by 165 // copy. 166 167 // CK4: call void {{.+}}@__kmpc_fork_call({{.+}} [[KERNELP1:@.+]] to void (i32*, i32*, ...)*), i32* {{.+}}) 168 // CK4: define internal void [[KERNELP1]](i32* {{[^,]+}}, i32* {{[^,]+}}, i32* {{[^,]+}}) 169 #pragma omp parallel 170 { 171 // CK4-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}}) 172 // CK4-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 173 // CK4-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 174 // CK4-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 175 // CK4-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 176 // CK4-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]], 177 // CK4-DAG: store i8* [[VALP:%.+]], i8** [[P1]], 178 // CK4-DAG: [[VALBP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8* 179 // CK4-DAG: [[VALP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8* 180 // CK4-DAG: [[VAL]] = load i[[sz]], i[[sz]]* [[ADDR:%.+]], 181 // CK4-64-DAG: [[CADDR:%.+]] = bitcast i[[sz]]* [[ADDR]] to i32* 182 // CK4-64-DAG: store i32 {{.+}}, i32* [[CADDR]], 183 184 // CK4: call void [[KERNEL:@.+]](i[[sz]] [[VAL]]) 185 #pragma omp target 186 { 187 #pragma omp parallel 188 { 189 ++i; 190 } 191 } 192 } 193 } 194 195 // CK4: define internal void [[KERNEL]](i[[sz]] [[ARG:%.+]]) 196 // CK4: [[ADDR:%.+]] = alloca i[[sz]], 197 // CK4: store i[[sz]] [[ARG]], i[[sz]]* [[ADDR]], 198 // CK4-64: [[CADDR:%.+]] = bitcast i64* [[ADDR]] to i32* 199 // CK4-64: call void {{.+}}@__kmpc_fork_call({{.+}} [[KERNELP2:@.+]] to void (i32*, i32*, ...)*), i32* [[CADDR]]) 200 // CK4-32: call void {{.+}}@__kmpc_fork_call({{.+}} [[KERNELP2:@.+]] to void (i32*, i32*, ...)*), i32* [[ADDR]]) 201 // CK4: define internal void [[KERNELP2]](i32* {{[^,]+}}, i32* {{[^,]+}}, i32* {{[^,]+}}) 202 #endif 203 ///==========================================================================/// 204 // RUN: %clang_cc1 -DCK5 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK5 --check-prefix CK5-64 205 // RUN: %clang_cc1 -DCK5 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 206 // 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 CK5 --check-prefix CK5-64 207 // RUN: %clang_cc1 -DCK5 -verify -fopenmp -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK5 --check-prefix CK5-32 208 // RUN: %clang_cc1 -DCK5 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 209 // 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 CK5 --check-prefix CK5-32 210 #ifdef CK5 211 212 // CK5-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4] 213 // Map types: OMP_MAP_BYCOPY = 128 214 // CK5-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 128] 215 216 // CK5-LABEL: implicit_maps_nested_integer_and_enum 217 void implicit_maps_nested_integer_and_enum (int a){ 218 enum Bla { 219 SomeEnum = 0x09 220 }; 221 222 // Using an enum should not change the mapping information. 223 int i = a; 224 225 // CK5-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}}) 226 // CK5-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 227 // CK5-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 228 // CK5-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 229 // CK5-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 230 // CK5-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]], 231 // CK5-DAG: store i8* [[VALP:%.+]], i8** [[P1]], 232 // CK5-DAG: [[VALBP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8* 233 // CK5-DAG: [[VALP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8* 234 // CK5-DAG: [[VAL]] = load i[[sz]], i[[sz]]* [[ADDR:%.+]], 235 // CK5-64-DAG: [[CADDR:%.+]] = bitcast i[[sz]]* [[ADDR]] to i32* 236 // CK5-64-DAG: store i32 {{.+}}, i32* [[CADDR]], 237 238 // CK5: call void [[KERNEL:@.+]](i[[sz]] [[VAL]]) 239 #pragma omp target 240 { 241 ++i; 242 i += SomeEnum; 243 } 244 } 245 246 // CK5: define internal void [[KERNEL]](i[[sz]] [[ARG:%.+]]) 247 // CK5: [[ADDR:%.+]] = alloca i[[sz]], 248 // CK5: store i[[sz]] [[ARG]], i[[sz]]* [[ADDR]], 249 // CK5-64: [[CADDR:%.+]] = bitcast i64* [[ADDR]] to i32* 250 // CK5-64: {{.+}} = load i32, i32* [[CADDR]], 251 // CK5-32: {{.+}} = load i32, i32* [[ADDR]], 252 253 #endif 254 ///==========================================================================/// 255 // RUN: %clang_cc1 -DCK6 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK6 --check-prefix CK6-64 256 // RUN: %clang_cc1 -DCK6 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 257 // 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 CK6 --check-prefix CK6-64 258 // RUN: %clang_cc1 -DCK6 -verify -fopenmp -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK6 --check-prefix CK6-32 259 // RUN: %clang_cc1 -DCK6 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 260 // 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 CK6 --check-prefix CK6-32 261 #ifdef CK6 262 // CK6-DAG: [[GBL:@Gi]] = global i32 0 263 // CK6-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4] 264 // Map types: OMP_MAP_BYCOPY = 128 265 // CK6-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 128] 266 267 // CK6-LABEL: implicit_maps_host_global 268 int Gi; 269 void implicit_maps_host_global (int a){ 270 // CK6-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}}) 271 // CK6-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 272 // CK6-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 273 // CK6-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 274 // CK6-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 275 // CK6-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]], 276 // CK6-DAG: store i8* [[VALP:%.+]], i8** [[P1]], 277 // CK6-DAG: [[VALBP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8* 278 // CK6-DAG: [[VALP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8* 279 // CK6-64-DAG: [[VAL]] = load i[[sz]], i[[sz]]* [[ADDR:%.+]], 280 // CK6-64-DAG: [[CADDR:%.+]] = bitcast i[[sz]]* [[ADDR]] to i32* 281 // CK6-64-DAG: store i32 [[GBLVAL:%.+]], i32* [[CADDR]], 282 // CK6-64-DAG: [[GBLVAL]] = load i32, i32* [[GBL]], 283 // CK6-32-DAG: [[VAL]] = load i[[sz]], i[[sz]]* [[GBLVAL:%.+]], 284 285 // CK6: call void [[KERNEL:@.+]](i[[sz]] [[VAL]]) 286 #pragma omp target 287 { 288 ++Gi; 289 } 290 } 291 292 // CK6: define internal void [[KERNEL]](i[[sz]] [[ARG:%.+]]) 293 // CK6: [[ADDR:%.+]] = alloca i[[sz]], 294 // CK6: store i[[sz]] [[ARG]], i[[sz]]* [[ADDR]], 295 // CK6-64: [[CADDR:%.+]] = bitcast i64* [[ADDR]] to i32* 296 // CK6-64: {{.+}} = load i32, i32* [[CADDR]], 297 // CK6-32: {{.+}} = load i32, i32* [[ADDR]], 298 299 #endif 300 ///==========================================================================/// 301 // RUN: %clang_cc1 -DCK7 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK7 --check-prefix CK7-64 302 // RUN: %clang_cc1 -DCK7 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 303 // 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 CK7 --check-prefix CK7-64 304 // RUN: %clang_cc1 -DCK7 -verify -fopenmp -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK7 --check-prefix CK7-32 305 // RUN: %clang_cc1 -DCK7 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 306 // 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 CK7 --check-prefix CK7-32 307 #ifdef CK7 308 309 // For a 32-bit targets, the value doesn't fit the size of the pointer, 310 // therefore it is passed by reference with a map 'to' specification. 311 312 // CK7-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 8] 313 // Map types: OMP_MAP_BYCOPY = 128 314 // CK7-64-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 128] 315 // Map types: OMP_MAP_TO = 1 316 // CK7-32-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 1] 317 318 // CK7-LABEL: implicit_maps_double 319 void implicit_maps_double (int a){ 320 double d = (double)a; 321 322 // CK7-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}}) 323 // CK7-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 324 // CK7-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 325 // CK7-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 326 // CK7-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 327 328 // CK7-64-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]], 329 // CK7-64-DAG: store i8* [[VALP:%.+]], i8** [[P1]], 330 // CK7-64-DAG: [[VALBP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8* 331 // CK7-64-DAG: [[VALP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8* 332 // CK7-64-DAG: [[VAL]] = load i[[sz]], i[[sz]]* [[ADDR:%.+]], 333 // CK7-64-64-DAG: [[CADDR:%.+]] = bitcast i[[sz]]* [[ADDR]] to double* 334 // CK7-64-64-DAG: store double {{.+}}, double* [[CADDR]], 335 336 // CK7-32-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]], 337 // CK7-32-DAG: store i8* [[VALP:%.+]], i8** [[P1]], 338 // CK7-32-DAG: [[VALBP]] = bitcast double* [[DECL:%.+]] to i8* 339 // CK7-32-DAG: [[VALP]] = bitcast double* [[DECL]] to i8* 340 341 // CK7-64: call void [[KERNEL:@.+]](i[[sz]] [[VAL]]) 342 // CK7-32: call void [[KERNEL:@.+]](double* [[DECL]]) 343 #pragma omp target 344 { 345 d += 1.0; 346 } 347 } 348 349 // CK7-64: define internal void [[KERNEL]](i[[sz]] [[ARG:%.+]]) 350 // CK7-64: [[ADDR:%.+]] = alloca i[[sz]], 351 // CK7-64: store i[[sz]] [[ARG]], i[[sz]]* [[ADDR]], 352 // CK7-64: [[CADDR:%.+]] = bitcast i64* [[ADDR]] to double* 353 // CK7-64: {{.+}} = load double, double* [[CADDR]], 354 355 // CK7-32: define internal void [[KERNEL]](double* {{.+}}[[ARG:%.+]]) 356 // CK7-32: [[ADDR:%.+]] = alloca double*, 357 // CK7-32: store double* [[ARG]], double** [[ADDR]], 358 // CK7-32: [[REF:%.+]] = load double*, double** [[ADDR]], 359 // CK7-32: {{.+}} = load double, double* [[REF]], 360 361 #endif 362 ///==========================================================================/// 363 // RUN: %clang_cc1 -DCK8 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK8 364 // RUN: %clang_cc1 -DCK8 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 365 // 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 CK8 366 // RUN: %clang_cc1 -DCK8 -verify -fopenmp -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK8 367 // RUN: %clang_cc1 -DCK8 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 368 // 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 CK8 369 #ifdef CK8 370 371 // CK8-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4] 372 // Map types: OMP_MAP_BYCOPY = 128 373 // CK8-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 128] 374 375 // CK8-LABEL: implicit_maps_float 376 void implicit_maps_float (int a){ 377 float f = (float)a; 378 379 // CK8-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}}) 380 // CK8-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 381 // CK8-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 382 // CK8-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 383 // CK8-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 384 // CK8-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]], 385 // CK8-DAG: store i8* [[VALP:%.+]], i8** [[P1]], 386 // CK8-DAG: [[VALBP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8* 387 // CK8-DAG: [[VALP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8* 388 // CK8-DAG: [[VAL]] = load i[[sz]], i[[sz]]* [[ADDR:%.+]], 389 // CK8-DAG: [[CADDR:%.+]] = bitcast i[[sz]]* [[ADDR]] to float* 390 // CK8-DAG: store float {{.+}}, float* [[CADDR]], 391 392 // CK8: call void [[KERNEL:@.+]](i[[sz]] [[VAL]]) 393 #pragma omp target 394 { 395 f += 1.0; 396 } 397 } 398 399 // CK8: define internal void [[KERNEL]](i[[sz]] [[ARG:%.+]]) 400 // CK8: [[ADDR:%.+]] = alloca i[[sz]], 401 // CK8: store i[[sz]] [[ARG]], i[[sz]]* [[ADDR]], 402 // CK8: [[CADDR:%.+]] = bitcast i[[sz]]* [[ADDR]] to float* 403 // CK8: {{.+}} = load float, float* [[CADDR]], 404 405 #endif 406 ///==========================================================================/// 407 // RUN: %clang_cc1 -DCK9 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK9 408 // RUN: %clang_cc1 -DCK9 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 409 // 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 CK9 410 // RUN: %clang_cc1 -DCK9 -verify -fopenmp -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK9 411 // RUN: %clang_cc1 -DCK9 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 412 // 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 CK9 413 #ifdef CK9 414 415 // CK9-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 16] 416 // Map types: OMP_MAP_TO + OMP_MAP_FROM = 2 + 1 417 // CK9-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 3] 418 419 // CK9-LABEL: implicit_maps_array 420 void implicit_maps_array (int a){ 421 double darr[2] = {(double)a, (double)a}; 422 423 // CK9-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}}) 424 // CK9-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 425 // CK9-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 426 // CK9-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 427 // CK9-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 428 // CK9-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]], 429 // CK9-DAG: store i8* [[VALP:%.+]], i8** [[P1]], 430 // CK9-DAG: [[VALBP]] = bitcast [2 x double]* [[DECL:%.+]] to i8* 431 // CK9-DAG: [[VALP]] = bitcast [2 x double]* [[DECL]] to i8* 432 433 // CK9: call void [[KERNEL:@.+]]([2 x double]* [[DECL]]) 434 #pragma omp target 435 { 436 darr[0] += 1.0; 437 darr[1] += 1.0; 438 } 439 } 440 441 // CK9: define internal void [[KERNEL]]([2 x double]* {{.+}}[[ARG:%.+]]) 442 // CK9: [[ADDR:%.+]] = alloca [2 x double]*, 443 // CK9: store [2 x double]* [[ARG]], [2 x double]** [[ADDR]], 444 // CK9: [[REF:%.+]] = load [2 x double]*, [2 x double]** [[ADDR]], 445 // CK9: {{.+}} = getelementptr inbounds [2 x double], [2 x double]* [[REF]], i[[sz]] 0, i[[sz]] 0 446 #endif 447 ///==========================================================================/// 448 // RUN: %clang_cc1 -DCK10 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK10 449 // RUN: %clang_cc1 -DCK10 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 450 // 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 CK10 451 // RUN: %clang_cc1 -DCK10 -verify -fopenmp -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK10 452 // RUN: %clang_cc1 -DCK10 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 453 // 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 CK10 454 #ifdef CK10 455 456 // CK10-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} {{8|4}}] 457 // Map types: OMP_MAP_BYCOPY | OMP_MAP_PTR = 128 + 32 458 // CK10-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 160] 459 460 // CK10-LABEL: implicit_maps_pointer 461 void implicit_maps_pointer (){ 462 double *ddyn; 463 464 // CK10-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}}) 465 // CK10-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 466 // CK10-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 467 // CK10-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 468 // CK10-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 469 // CK10-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]], 470 // CK10-DAG: store i8* [[VALP:%.+]], i8** [[P1]], 471 // CK10-DAG: [[VALBP]] = bitcast double* [[PTR:%.+]] to i8* 472 // CK10-DAG: [[VALP]] = bitcast double* [[PTR]] to i8* 473 474 // CK10: call void [[KERNEL:@.+]](double* [[PTR]]) 475 #pragma omp target 476 { 477 ddyn[0] += 1.0; 478 ddyn[1] += 1.0; 479 } 480 } 481 482 // CK10: define internal void [[KERNEL]](double* {{.*}}[[ARG:%.+]]) 483 // CK10: [[ADDR:%.+]] = alloca double*, 484 // CK10: store double* [[ARG]], double** [[ADDR]], 485 // CK10: [[REF:%.+]] = load double*, double** [[ADDR]], 486 // CK10: {{.+}} = getelementptr inbounds double, double* [[REF]], i[[sz]] 0 487 488 #endif 489 ///==========================================================================/// 490 // RUN: %clang_cc1 -DCK11 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK11 491 // RUN: %clang_cc1 -DCK11 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 492 // 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 CK11 493 // RUN: %clang_cc1 -DCK11 -verify -fopenmp -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK11 494 // RUN: %clang_cc1 -DCK11 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 495 // 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 CK11 496 #ifdef CK11 497 498 // CK11-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 16] 499 // Map types: OMP_MAP_TO = 1 500 // CK11-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 1] 501 502 // CK11-LABEL: implicit_maps_double_complex 503 void implicit_maps_double_complex (int a){ 504 double _Complex dc = (double)a; 505 506 // CK11-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}}) 507 // CK11-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 508 // CK11-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 509 // CK11-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 510 // CK11-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 511 // CK11-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]], 512 // CK11-DAG: store i8* [[VALP:%.+]], i8** [[P1]], 513 // CK11-DAG: [[VALBP]] = bitcast { double, double }* [[PTR:%.+]] to i8* 514 // CK11-DAG: [[VALP]] = bitcast { double, double }* [[PTR]] to i8* 515 516 // CK11: call void [[KERNEL:@.+]]({ double, double }* [[PTR]]) 517 #pragma omp target 518 { 519 dc *= dc; 520 } 521 } 522 523 // CK11: define internal void [[KERNEL]]({ double, double }* {{.*}}[[ARG:%.+]]) 524 // CK11: [[ADDR:%.+]] = alloca { double, double }*, 525 // CK11: store { double, double }* [[ARG]], { double, double }** [[ADDR]], 526 // CK11: [[REF:%.+]] = load { double, double }*, { double, double }** [[ADDR]], 527 // CK11: {{.+}} = getelementptr inbounds { double, double }, { double, double }* [[REF]], i32 0, i32 0 528 #endif 529 ///==========================================================================/// 530 // RUN: %clang_cc1 -DCK12 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK12 --check-prefix CK12-64 531 // RUN: %clang_cc1 -DCK12 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 532 // 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 CK12 --check-prefix CK12-64 533 // RUN: %clang_cc1 -DCK12 -verify -fopenmp -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK12 --check-prefix CK12-32 534 // RUN: %clang_cc1 -DCK12 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 535 // 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 CK12 --check-prefix CK12-32 536 #ifdef CK12 537 538 // For a 32-bit targets, the value doesn't fit the size of the pointer, 539 // therefore it is passed by reference with a map 'to' specification. 540 541 // CK12-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 8] 542 // Map types: OMP_MAP_BYCOPY = 128 543 // CK12-64-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 128] 544 // Map types: OMP_MAP_TO = 1 545 // CK12-32-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 1] 546 547 // CK12-LABEL: implicit_maps_float_complex 548 void implicit_maps_float_complex (int a){ 549 float _Complex fc = (float)a; 550 551 // CK12-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}}) 552 // CK12-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 553 // CK12-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 554 // CK12-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 555 // CK12-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 556 557 // CK12-64-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]], 558 // CK12-64-DAG: store i8* [[VALP:%.+]], i8** [[P1]], 559 // CK12-64-DAG: [[VALBP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8* 560 // CK12-64-DAG: [[VALP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8* 561 // CK12-64-DAG: [[VAL]] = load i[[sz]], i[[sz]]* [[ADDR:%.+]], 562 // CK12-64-DAG: [[CADDR:%.+]] = bitcast i[[sz]]* [[ADDR]] to { float, float }* 563 // CK12-64-DAG: store { float, float } {{.+}}, { float, float }* [[CADDR]], 564 565 // CK12-32-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]], 566 // CK12-32-DAG: store i8* [[VALP:%.+]], i8** [[P1]], 567 // CK12-32-DAG: [[VALBP]] = bitcast { float, float }* [[DECL:%.+]] to i8* 568 // CK12-32-DAG: [[VALP]] = bitcast { float, float }* [[DECL]] to i8* 569 570 // CK12-64: call void [[KERNEL:@.+]](i[[sz]] [[VAL]]) 571 // CK12-32: call void [[KERNEL:@.+]]({ float, float }* [[DECL]]) 572 #pragma omp target 573 { 574 fc *= fc; 575 } 576 } 577 578 // CK12-64: define internal void [[KERNEL]](i[[sz]] [[ARG:%.+]]) 579 // CK12-64: [[ADDR:%.+]] = alloca i[[sz]], 580 // CK12-64: store i[[sz]] [[ARG]], i[[sz]]* [[ADDR]], 581 // CK12-64: [[CADDR:%.+]] = bitcast i[[sz]]* [[ADDR]] to { float, float }* 582 // CK12-64: {{.+}} = getelementptr inbounds { float, float }, { float, float }* [[CADDR]], i32 0, i32 0 583 584 // CK12-32: define internal void [[KERNEL]]({ float, float }* {{.+}}[[ARG:%.+]]) 585 // CK12-32: [[ADDR:%.+]] = alloca { float, float }*, 586 // CK12-32: store { float, float }* [[ARG]], { float, float }** [[ADDR]], 587 // CK12-32: [[REF:%.+]] = load { float, float }*, { float, float }** [[ADDR]], 588 // CK12-32: {{.+}} = getelementptr inbounds { float, float }, { float, float }* [[REF]], i32 0, i32 0 589 #endif 590 ///==========================================================================/// 591 // RUN: %clang_cc1 -DCK13 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK13 592 // RUN: %clang_cc1 -DCK13 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 593 // 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 CK13 594 // RUN: %clang_cc1 -DCK13 -verify -fopenmp -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK13 595 // RUN: %clang_cc1 -DCK13 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 596 // 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 CK13 597 #ifdef CK13 598 599 // We don't have a constant map size for VLAs. 600 // Map types: 601 // - OMP_MAP_BYCOPY = 128 (vla size) 602 // - OMP_MAP_BYCOPY = 128 (vla size) 603 // - OMP_MAP_TO + OMP_MAP_FROM = 2 + 1 604 // CK13-DAG: [[TYPES:@.+]] = {{.+}}constant [3 x i32] [i32 128, i32 128, i32 3] 605 606 // CK13-LABEL: implicit_maps_variable_length_array 607 void implicit_maps_variable_length_array (int a){ 608 double vla[2][a]; 609 610 // CK13-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 3, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], i[[sz:64|32]]* [[SGEP:%[^,]+]], {{.+}}[[TYPES]]{{.+}}) 611 // CK13-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 612 // CK13-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 613 // CK13-DAG: [[SGEP]] = getelementptr inbounds {{.+}}[[SS:%[^,]+]], i32 0, i32 0 614 615 // CK13-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 616 // CK13-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 617 // CK13-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[SS]], i32 0, i32 0 618 // CK13-DAG: store i8* inttoptr (i[[sz]] 2 to i8*), i8** [[BP0]], 619 // CK13-DAG: store i8* inttoptr (i[[sz]] 2 to i8*), i8** [[P0]], 620 // CK13-DAG: store i[[sz]] {{8|4}}, i[[sz]]* [[S0]], 621 622 // CK13-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 1 623 // CK13-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 1 624 // CK13-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[SS]], i32 0, i32 1 625 // CK13-DAG: store i8* [[VALBP1:%.+]], i8** [[BP1]], 626 // CK13-DAG: store i8* [[VALP1:%.+]], i8** [[P1]], 627 // CK13-DAG: [[VALBP1]] = inttoptr i[[sz]] [[VAL:%.+]] to i8* 628 // CK13-DAG: [[VALP1]] = inttoptr i[[sz]] [[VAL:%.+]] to i8* 629 // CK13-DAG: store i[[sz]] {{8|4}}, i[[sz]]* [[S1]], 630 631 // CK13-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 2 632 // CK13-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 2 633 // CK13-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[SS]], i32 0, i32 2 634 // CK13-DAG: store i8* [[VALBP2:%.+]], i8** [[BP2]], 635 // CK13-DAG: store i8* [[VALP2:%.+]], i8** [[P2]], 636 // CK13-DAG: store i[[sz]] [[VALS2:%.+]], i[[sz]]* [[S2]], 637 // CK13-DAG: [[VALBP2]] = bitcast double* [[DECL:%.+]] to i8* 638 // CK13-DAG: [[VALP2]] = bitcast double* [[DECL]] to i8* 639 // CK13-DAG: [[VALS2]] = mul nuw i[[sz]] %{{.+}}, 8 640 641 // CK13: call void [[KERNEL:@.+]](i[[sz]] {{.+}}, i[[sz]] {{.+}}, double* [[DECL]]) 642 #pragma omp target 643 { 644 vla[1][3] += 1.0; 645 } 646 } 647 648 // CK13: define internal void [[KERNEL]](i[[sz]] [[VLA0:%.+]], i[[sz]] [[VLA1:%.+]], double* {{.+}}[[ARG:%.+]]) 649 // CK13: [[ADDR0:%.+]] = alloca i[[sz]], 650 // CK13: [[ADDR1:%.+]] = alloca i[[sz]], 651 // CK13: [[ADDR2:%.+]] = alloca double*, 652 // CK13: store i[[sz]] [[VLA0]], i[[sz]]* [[ADDR0]], 653 // CK13: store i[[sz]] [[VLA1]], i[[sz]]* [[ADDR1]], 654 // CK13: store double* [[ARG]], double** [[ADDR2]], 655 // CK13: {{.+}} = load i[[sz]], i[[sz]]* [[ADDR0]], 656 // CK13: {{.+}} = load i[[sz]], i[[sz]]* [[ADDR1]], 657 // CK13: [[REF:%.+]] = load double*, double** [[ADDR2]], 658 // CK13: {{.+}} = getelementptr inbounds double, double* [[REF]], i[[sz]] %{{.+}} 659 #endif 660 ///==========================================================================/// 661 // RUN: %clang_cc1 -DCK14 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK14 --check-prefix CK14-64 662 // RUN: %clang_cc1 -DCK14 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 663 // 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 CK14 --check-prefix CK14-64 664 // RUN: %clang_cc1 -DCK14 -verify -fopenmp -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK14 --check-prefix CK14-32 665 // RUN: %clang_cc1 -DCK14 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 666 // 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 CK14 --check-prefix CK14-32 667 #ifdef CK14 668 669 // CK14-DAG: [[ST:%.+]] = type { i32, double } 670 // CK14-DAG: [[SIZES:@.+]] = {{.+}}constant [2 x i[[sz:64|32]]] [i{{64|32}} {{16|12}}, i{{64|32}} 4] 671 // Map types: 672 // - OMP_MAP_TO | OMP_MAP_FROM = 1 + 2 673 // - OMP_MAP_BYCOPY = 128 674 // CK14-DAG: [[TYPES:@.+]] = {{.+}}constant [2 x i32] [i32 3, i32 128] 675 676 class SSS { 677 public: 678 int a; 679 double b; 680 681 void foo(int c) { 682 #pragma omp target 683 { 684 a += c; 685 b += (double)c; 686 } 687 } 688 689 SSS(int a, double b) : a(a), b(b) {} 690 }; 691 692 // CK14-LABEL: implicit_maps_class 693 void implicit_maps_class (int a){ 694 SSS sss(a, (double)a); 695 696 // CK14: define {{.*}}void @{{.+}}foo{{.+}}([[ST]]* {{[^,]+}}, i32 {{[^,]+}}) 697 // CK14-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 2, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}}) 698 // CK14-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 699 // CK14-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 700 701 // CK14-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 702 // CK14-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 703 // CK14-DAG: store i8* [[VALBP0:%.+]], i8** [[BP0]], 704 // CK14-DAG: store i8* [[VALP0:%.+]], i8** [[P0]], 705 // CK14-DAG: [[VALBP0]] = bitcast [[ST]]* [[DECL:%.+]] to i8* 706 // CK14-DAG: [[VALP0]] = bitcast [[ST]]* [[DECL]] to i8* 707 708 // CK14-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 1 709 // CK14-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 1 710 // CK14-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]], 711 // CK14-DAG: store i8* [[VALP:%.+]], i8** [[P1]], 712 // CK14-DAG: [[VALBP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8* 713 // CK14-DAG: [[VALP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8* 714 // CK14-DAG: [[VAL]] = load i[[sz]], i[[sz]]* [[ADDR:%.+]], 715 // CK14-64-DAG: [[CADDR:%.+]] = bitcast i[[sz]]* [[ADDR]] to i32* 716 // CK14-64-DAG: store i32 {{.+}}, i32* [[CADDR]], 717 718 // CK14: call void [[KERNEL:@.+]]([[ST]]* [[DECL]], i[[sz]] {{.+}}) 719 sss.foo(123); 720 } 721 722 // CK14: define internal void [[KERNEL]]([[ST]]* [[THIS:%.+]], i[[sz]] [[ARG:%.+]]) 723 // CK14: [[ADDR0:%.+]] = alloca [[ST]]*, 724 // CK14: [[ADDR1:%.+]] = alloca i[[sz]], 725 // CK14: store [[ST]]* [[THIS]], [[ST]]** [[ADDR0]], 726 // CK14: store i[[sz]] [[ARG]], i[[sz]]* [[ADDR1]], 727 // CK14: [[REF0:%.+]] = load [[ST]]*, [[ST]]** [[ADDR0]], 728 // CK14-64: [[CADDR1:%.+]] = bitcast i[[sz]]* [[ADDR1]] to i32* 729 // CK14-64: {{.+}} = load i32, i32* [[CADDR1]], 730 // CK14-32: {{.+}} = load i32, i32* [[ADDR1]], 731 // CK14: {{.+}} = getelementptr inbounds [[ST]], [[ST]]* [[REF0]], i32 0, i32 0 732 733 #endif 734 ///==========================================================================/// 735 // RUN: %clang_cc1 -DCK15 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK15 --check-prefix CK15-64 736 // RUN: %clang_cc1 -DCK15 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 737 // 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 CK15 --check-prefix CK15-64 738 // RUN: %clang_cc1 -DCK15 -verify -fopenmp -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK15 --check-prefix CK15-32 739 // RUN: %clang_cc1 -DCK15 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 740 // 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 CK15 --check-prefix CK15-32 741 #ifdef CK15 742 743 // CK15: [[ST:%.+]] = type { i32, double, i32* } 744 // CK15: [[SIZES:@.+]] = {{.+}}constant [2 x i[[sz:64|32]]] [i{{64|32}} {{24|16}}, i{{64|32}} 4] 745 // Map types: 746 // - OMP_MAP_TO | OMP_MAP_FROM = 1 + 2 747 // - OMP_MAP_BYCOPY = 128 748 // CK15: [[TYPES:@.+]] = {{.+}}constant [2 x i32] [i32 3, i32 128] 749 750 // CK15: [[SIZES2:@.+]] = {{.+}}constant [2 x i[[sz]]] [i{{64|32}} {{24|16}}, i{{64|32}} 4] 751 // Map types: 752 // - OMP_MAP_TO | OMP_MAP_FROM = 1 + 2 753 // - OMP_MAP_BYCOPY = 128 754 // CK15: [[TYPES2:@.+]] = {{.+}}constant [2 x i32] [i32 3, i32 128] 755 756 template<int x> 757 class SSST { 758 public: 759 int a; 760 double b; 761 int &r; 762 763 void foo(int c) { 764 #pragma omp target 765 { 766 a += c + x; 767 b += (double)(c + x); 768 r += x; 769 } 770 } 771 template<int y> 772 void bar(int c) { 773 #pragma omp target 774 { 775 a += c + x + y; 776 b += (double)(c + x + y); 777 r += x + y; 778 } 779 } 780 781 SSST(int a, double b, int &r) : a(a), b(b), r(r) {} 782 }; 783 784 // CK15-LABEL: implicit_maps_templated_class 785 void implicit_maps_templated_class (int a){ 786 SSST<123> ssst(a, (double)a, a); 787 788 // CK15: define {{.*}}void @{{.+}}foo{{.+}}([[ST]]* {{[^,]+}}, i32 {{[^,]+}}) 789 // CK15-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 2, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}}) 790 // CK15-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 791 // CK15-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 792 793 // CK15-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 794 // CK15-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 795 // CK15-DAG: store i8* [[VALBP0:%.+]], i8** [[BP0]], 796 // CK15-DAG: store i8* [[VALP0:%.+]], i8** [[P0]], 797 // CK15-DAG: [[VALBP0]] = bitcast [[ST]]* [[DECL:%.+]] to i8* 798 // CK15-DAG: [[VALP0]] = bitcast [[ST]]* [[DECL]] to i8* 799 800 // CK15-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 1 801 // CK15-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 1 802 // CK15-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]], 803 // CK15-DAG: store i8* [[VALP:%.+]], i8** [[P1]], 804 // CK15-DAG: [[VALBP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8* 805 // CK15-DAG: [[VALP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8* 806 // CK15-DAG: [[VAL]] = load i[[sz]], i[[sz]]* [[ADDR:%.+]], 807 // CK15-64-DAG: [[CADDR:%.+]] = bitcast i[[sz]]* [[ADDR]] to i32* 808 // CK15-64-DAG: store i32 {{.+}}, i32* [[CADDR]], 809 810 // CK15: call void [[KERNEL:@.+]]([[ST]]* [[DECL]], i[[sz]] {{.+}}) 811 ssst.foo(456); 812 813 // CK15: define {{.*}}void @{{.+}}bar{{.+}}([[ST]]* {{[^,]+}}, i32 {{[^,]+}}) 814 // CK15-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 2, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES2]]{{.+}}, {{.+}}[[TYPES2]]{{.+}}) 815 // CK15-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 816 // CK15-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 817 818 // CK15-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 819 // CK15-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 820 // CK15-DAG: store i8* [[VALBP0:%.+]], i8** [[BP0]], 821 // CK15-DAG: store i8* [[VALP0:%.+]], i8** [[P0]], 822 // CK15-DAG: [[VALBP0]] = bitcast [[ST]]* [[DECL:%.+]] to i8* 823 // CK15-DAG: [[VALP0]] = bitcast [[ST]]* [[DECL]] to i8* 824 825 // CK15-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 1 826 // CK15-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 1 827 // CK15-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]], 828 // CK15-DAG: store i8* [[VALP:%.+]], i8** [[P1]], 829 // CK15-DAG: [[VALBP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8* 830 // CK15-DAG: [[VALP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8* 831 // CK15-DAG: [[VAL]] = load i[[sz]], i[[sz]]* [[ADDR:%.+]], 832 // CK15-64-DAG: [[CADDR:%.+]] = bitcast i[[sz]]* [[ADDR]] to i32* 833 // CK15-64-DAG: store i32 {{.+}}, i32* [[CADDR]], 834 835 // CK15: call void [[KERNEL2:@.+]]([[ST]]* [[DECL]], i[[sz]] {{.+}}) 836 ssst.bar<210>(789); 837 } 838 839 // CK15: define internal void [[KERNEL]]([[ST]]* [[THIS:%.+]], i[[sz]] [[ARG:%.+]]) 840 // CK15: [[ADDR0:%.+]] = alloca [[ST]]*, 841 // CK15: [[ADDR1:%.+]] = alloca i[[sz]], 842 // CK15: store [[ST]]* [[THIS]], [[ST]]** [[ADDR0]], 843 // CK15: store i[[sz]] [[ARG]], i[[sz]]* [[ADDR1]], 844 // CK15: [[REF0:%.+]] = load [[ST]]*, [[ST]]** [[ADDR0]], 845 // CK15-64: [[CADDR1:%.+]] = bitcast i[[sz]]* [[ADDR1]] to i32* 846 // CK15-64: {{.+}} = load i32, i32* [[CADDR1]], 847 // CK15-32: {{.+}} = load i32, i32* [[ADDR1]], 848 // CK15: {{.+}} = getelementptr inbounds [[ST]], [[ST]]* [[REF0]], i32 0, i32 0 849 850 // CK15: define internal void [[KERNEL2]]([[ST]]* [[THIS:%.+]], i[[sz]] [[ARG:%.+]]) 851 // CK15: [[ADDR0:%.+]] = alloca [[ST]]*, 852 // CK15: [[ADDR1:%.+]] = alloca i[[sz]], 853 // CK15: store [[ST]]* [[THIS]], [[ST]]** [[ADDR0]], 854 // CK15: store i[[sz]] [[ARG]], i[[sz]]* [[ADDR1]], 855 // CK15: [[REF0:%.+]] = load [[ST]]*, [[ST]]** [[ADDR0]], 856 // CK15-64: [[CADDR1:%.+]] = bitcast i[[sz]]* [[ADDR1]] to i32* 857 // CK15-64: {{.+}} = load i32, i32* [[CADDR1]], 858 // CK15-32: {{.+}} = load i32, i32* [[ADDR1]], 859 // CK15: {{.+}} = getelementptr inbounds [[ST]], [[ST]]* [[REF0]], i32 0, i32 0 860 861 #endif 862 ///==========================================================================/// 863 // RUN: %clang_cc1 -DCK16 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK16 --check-prefix CK16-64 864 // RUN: %clang_cc1 -DCK16 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 865 // 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 CK16 --check-prefix CK16-64 866 // RUN: %clang_cc1 -DCK16 -verify -fopenmp -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK16 --check-prefix CK16-32 867 // RUN: %clang_cc1 -DCK16 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 868 // 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 CK16 --check-prefix CK16-32 869 #ifdef CK16 870 871 // CK16-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4] 872 // Map types: 873 // - OMP_MAP_BYCOPY = 128 874 // CK16-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 128] 875 876 template<int y> 877 int foo(int d) { 878 int res = d; 879 #pragma omp target 880 { 881 res += y; 882 } 883 return res; 884 } 885 // CK16-LABEL: implicit_maps_templated_function 886 void implicit_maps_templated_function (int a){ 887 int i = a; 888 889 // CK16: define {{.*}}i32 @{{.+}}foo{{.+}}(i32 {{[^,]+}}) 890 // CK16-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}}) 891 // CK16-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 892 // CK16-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 893 894 // CK16-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 895 // CK16-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 896 // CK16-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]], 897 // CK16-DAG: store i8* [[VALP:%.+]], i8** [[P1]], 898 // CK16-DAG: [[VALBP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8* 899 // CK16-DAG: [[VALP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8* 900 // CK16-DAG: [[VAL]] = load i[[sz]], i[[sz]]* [[ADDR:%.+]], 901 // CK16-64-DAG: [[CADDR:%.+]] = bitcast i[[sz]]* [[ADDR]] to i32* 902 // CK16-64-DAG: store i32 {{.+}}, i32* [[CADDR]], 903 904 // CK16: call void [[KERNEL:@.+]](i[[sz]] [[VAL]]) 905 i = foo<543>(i); 906 } 907 // CK16: define internal void [[KERNEL]](i[[sz]] [[ARG:%.+]]) 908 // CK16: [[ADDR:%.+]] = alloca i[[sz]], 909 // CK16: store i[[sz]] [[ARG]], i[[sz]]* [[ADDR]], 910 // CK16-64: [[CADDR:%.+]] = bitcast i64* [[ADDR]] to i32* 911 // CK16-64: {{.+}} = load i32, i32* [[CADDR]], 912 // CK16-32: {{.+}} = load i32, i32* [[ADDR]], 913 914 #endif 915 ///==========================================================================/// 916 // RUN: %clang_cc1 -DCK17 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK17 917 // RUN: %clang_cc1 -DCK17 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 918 // 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 CK17 919 // RUN: %clang_cc1 -DCK17 -verify -fopenmp -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK17 920 // RUN: %clang_cc1 -DCK17 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 921 // 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 CK17 922 #ifdef CK17 923 924 // CK17-DAG: [[ST:%.+]] = type { i32, double } 925 // CK17-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} {{16|12}}] 926 // Map types: OMP_MAP_TO + OMP_MAP_FROM = 2 + 1 927 // CK17-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 3] 928 929 class SSS { 930 public: 931 int a; 932 double b; 933 }; 934 935 // CK17-LABEL: implicit_maps_struct 936 void implicit_maps_struct (int a){ 937 SSS s = {a, (double)a}; 938 939 // CK17-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}}) 940 // CK17-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 941 // CK17-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 942 // CK17-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 943 // CK17-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 944 // CK17-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]], 945 // CK17-DAG: store i8* [[VALP:%.+]], i8** [[P1]], 946 // CK17-DAG: [[VALBP]] = bitcast [[ST]]* [[DECL:%.+]] to i8* 947 // CK17-DAG: [[VALP]] = bitcast [[ST]]* [[DECL]] to i8* 948 949 // CK17: call void [[KERNEL:@.+]]([[ST]]* [[DECL]]) 950 #pragma omp target 951 { 952 s.a += 1; 953 s.b += 1.0; 954 } 955 } 956 957 // CK17: define internal void [[KERNEL]]([[ST]]* {{.+}}[[ARG:%.+]]) 958 // CK17: [[ADDR:%.+]] = alloca [[ST]]*, 959 // CK17: store [[ST]]* [[ARG]], [[ST]]** [[ADDR]], 960 // CK17: [[REF:%.+]] = load [[ST]]*, [[ST]]** [[ADDR]], 961 // CK17: {{.+}} = getelementptr inbounds [[ST]], [[ST]]* [[REF]], i32 0, i32 0 962 #endif 963 ///==========================================================================/// 964 // RUN: %clang_cc1 -DCK18 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK18 --check-prefix CK18-64 965 // RUN: %clang_cc1 -DCK18 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 966 // 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 CK18 --check-prefix CK18-64 967 // RUN: %clang_cc1 -DCK18 -verify -fopenmp -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK18 --check-prefix CK18-32 968 // RUN: %clang_cc1 -DCK18 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 969 // 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 CK18 --check-prefix CK18-32 970 #ifdef CK18 971 972 // CK18-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4] 973 // Map types: 974 // - OMP_MAP_BYCOPY = 128 975 // CK18-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 128] 976 977 template<typename T> 978 int foo(T d) { 979 #pragma omp target 980 { 981 d += (T)1; 982 } 983 return d; 984 } 985 // CK18-LABEL: implicit_maps_template_type_capture 986 void implicit_maps_template_type_capture (int a){ 987 int i = a; 988 989 // CK18: define {{.*}}i32 @{{.+}}foo{{.+}}(i32 {{[^,]+}}) 990 // CK18-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}}) 991 // CK18-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 992 // CK18-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 993 994 // CK18-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 995 // CK18-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 996 // CK18-DAG: store i8* [[VALBP:%.+]], i8** [[BP1]], 997 // CK18-DAG: store i8* [[VALP:%.+]], i8** [[P1]], 998 // CK18-DAG: [[VALBP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8* 999 // CK18-DAG: [[VALP]] = inttoptr i[[sz]] [[VAL:%.+]] to i8* 1000 // CK18-DAG: [[VAL]] = load i[[sz]], i[[sz]]* [[ADDR:%.+]], 1001 // CK18-64-DAG: [[CADDR:%.+]] = bitcast i[[sz]]* [[ADDR]] to i32* 1002 // CK18-64-DAG: store i32 {{.+}}, i32* [[CADDR]], 1003 1004 // CK18: call void [[KERNEL:@.+]](i[[sz]] [[VAL]]) 1005 i = foo(i); 1006 } 1007 // CK18: define internal void [[KERNEL]](i[[sz]] [[ARG:%.+]]) 1008 // CK18: [[ADDR:%.+]] = alloca i[[sz]], 1009 // CK18: store i[[sz]] [[ARG]], i[[sz]]* [[ADDR]], 1010 // CK18-64: [[CADDR:%.+]] = bitcast i64* [[ADDR]] to i32* 1011 // CK18-64: {{.+}} = load i32, i32* [[CADDR]], 1012 // CK18-32: {{.+}} = load i32, i32* [[ADDR]], 1013 1014 #endif 1015 #endif 1016
