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