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