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