1 // expected-no-diagnostics 2 #ifndef HEADER 3 #define HEADER 4 5 ///==========================================================================/// 6 // RUN: %clang_cc1 -no-opaque-pointers -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 -no-opaque-pointers -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 -no-opaque-pointers -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 -no-opaque-pointers -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 -no-opaque-pointers -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 -no-opaque-pointers -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 13 // RUN: %clang_cc1 -no-opaque-pointers -DCK1 -verify -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s 14 // RUN: %clang_cc1 -no-opaque-pointers -DCK1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 15 // RUN: %clang_cc1 -no-opaque-pointers -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s 16 // RUN: %clang_cc1 -no-opaque-pointers -DCK1 -verify -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s 17 // RUN: %clang_cc1 -no-opaque-pointers -DCK1 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 18 // RUN: %clang_cc1 -no-opaque-pointers -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s 19 // SIMD-ONLY0-NOT: {{__kmpc|__tgt}} 20 #ifdef CK1 21 22 double *g; 23 24 // CK1: @g ={{.*}} global double* 25 // CK1: [[SIZES00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} {{8|4}}] 26 // CK1: [[TYPES00:@.+]] = {{.+}}constant [1 x i64] [i64 288] 27 28 // CK1: [[SIZES01:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] {{8|4}}] 29 // CK1: [[TYPES01:@.+]] = {{.+}}constant [1 x i64] [i64 288] 30 31 // CK1: [[SIZES02:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] {{8|4}}] 32 // CK1: [[TYPES02:@.+]] = {{.+}}constant [1 x i64] [i64 288] 33 34 // CK1: [[SIZES03:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] {{8|4}}] 35 // CK1: [[TYPES03:@.+]] = {{.+}}constant [1 x i64] [i64 288] 36 37 // CK1: [[SIZES04:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] {{8|4}}] 38 // CK1: [[TYPES04:@.+]] = {{.+}}constant [1 x i64] [i64 288] 39 40 // CK1: [[SIZES05:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] {{8|4}}] 41 // CK1: [[TYPES05:@.+]] = {{.+}}constant [1 x i64] [i64 288] 42 43 // CK1: [[SIZES06:@.+]] = {{.+}}constant [2 x i[[sz]]] [i[[sz]] {{8|4}}, i[[sz]] {{8|4}}] 44 // CK1: [[TYPES06:@.+]] = {{.+}}constant [2 x i64] [i64 288, i64 288] 45 46 // CK1-LABEL: @_Z3foo{{.*}}( 47 template<typename T> 48 void foo(float *&lr, T *&tr) { 49 float *l; 50 T *t; 51 52 // CK1-DAG: [[RET:%.+]] = call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 53 // CK1-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 54 // CK1-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 55 // CK1-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 56 // CK1-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 57 // CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 58 // CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 59 // CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 60 // CK1-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 61 // CK1-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to double** 62 // CK1-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to double** 63 // CK1-DAG: store double* [[VAL:%.+]], double** [[CBP1]] 64 // CK1-DAG: store double* [[VAL]], double** [[CP1]] 65 // CK1-DAG: [[VAL]] = load double*, double** [[ADDR:@g]], 66 67 // CK1: call void [[KERNEL:@.+]](double* [[VAL]]) 68 #pragma omp target is_device_ptr(g) 69 { 70 ++g; 71 } 72 73 // CK1-DAG: [[RET:%.+]] = call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 74 // CK1-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 75 // CK1-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 76 // CK1-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 77 // CK1-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 78 // CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 79 // CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 80 // CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 81 // CK1-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 82 // CK1-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to float** 83 // CK1-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to float** 84 // CK1-DAG: store float* [[VAL:%.+]], float** [[CBP1]] 85 // CK1-DAG: store float* [[VAL]], float** [[CP1]] 86 // CK1-DAG: [[VAL]] = load float*, float** [[ADDR:%.+]], 87 88 // CK1: call void [[KERNEL:@.+]](float* [[VAL]]) 89 #pragma omp target is_device_ptr(l) 90 { 91 ++l; 92 } 93 94 // CK1-DAG: [[RET:%.+]] = call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 95 // CK1-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 96 // CK1-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 97 // CK1-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 98 // CK1-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 99 // CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 100 // CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 101 // CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 102 // CK1-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 103 // CK1-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to i32** 104 // CK1-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32** 105 // CK1-DAG: store i32* [[VAL:%.+]], i32** [[CBP1]] 106 // CK1-DAG: store i32* [[VAL]], i32** [[CP1]] 107 // CK1-DAG: [[VAL]] = load i32*, i32** [[ADDR:%.+]], 108 109 // CK1: call void [[KERNEL:@.+]](i32* [[VAL]]) 110 #pragma omp target is_device_ptr(t) 111 { 112 ++t; 113 } 114 115 // CK1-DAG: [[RET:%.+]] = call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 116 // CK1-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 117 // CK1-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 118 // CK1-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 119 // CK1-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 120 // CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 121 // CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 122 // CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 123 // CK1-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 124 // CK1-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to float** 125 // CK1-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to float** 126 // CK1-DAG: store float* [[VAL:%.+]], float** [[CBP1]] 127 // CK1-DAG: store float* [[VAL]], float** [[CP1]] 128 // CK1-DAG: [[VAL]] = load float*, float** [[ADDR:%.+]], 129 // CK1-DAG: [[ADDR]] = load float**, float*** [[ADDR2:%.+]], 130 131 // CK1: call void [[KERNEL:@.+]](float* [[VAL]]) 132 #pragma omp target is_device_ptr(lr) 133 { 134 ++lr; 135 } 136 137 // CK1-DAG: [[RET:%.+]] = call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 138 // CK1-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 139 // CK1-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 140 // CK1-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 141 // CK1-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 142 // CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 143 // CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 144 // CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 145 // CK1-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 146 // CK1-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to i32** 147 // CK1-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32** 148 // CK1-DAG: store i32* [[VAL:%.+]], i32** [[CBP1]] 149 // CK1-DAG: store i32* [[VAL]], i32** [[CP1]] 150 // CK1-DAG: [[VAL]] = load i32*, i32** [[ADDR:%.+]], 151 // CK1-DAG: [[ADDR]] = load i32**, i32*** [[ADDR2:%.+]], 152 153 // CK1: call void [[KERNEL:@.+]](i32* [[VAL]]) 154 #pragma omp target is_device_ptr(tr) 155 { 156 ++tr; 157 } 158 159 // CK1-DAG: [[RET:%.+]] = call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 160 // CK1-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 161 // CK1-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 162 // CK1-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 163 // CK1-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 164 // CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 165 // CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 166 // CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 167 // CK1-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 168 // CK1-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to i32** 169 // CK1-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32** 170 // CK1-DAG: store i32* [[VAL:%.+]], i32** [[CBP1]] 171 // CK1-DAG: store i32* [[VAL]], i32** [[CP1]] 172 // CK1-DAG: [[VAL]] = load i32*, i32** [[ADDR:%.+]], 173 // CK1-DAG: [[ADDR]] = load i32**, i32*** [[ADDR2:%.+]], 174 175 // CK1: call void [[KERNEL:@.+]](i32* [[VAL]]) 176 #pragma omp target is_device_ptr(tr, lr) 177 { 178 ++tr; 179 } 180 181 // CK1-DAG: [[RET:%.+]] = call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 182 // CK1-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 183 // CK1-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 184 // CK1-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 185 // CK1-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 186 // CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 187 // CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 188 // CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 189 // CK1-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 190 // CK1-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to i32** 191 // CK1-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32** 192 // CK1-DAG: store i32* [[VAL:%.+]], i32** [[CBP1]] 193 // CK1-DAG: store i32* [[VAL]], i32** [[CP1]] 194 // CK1-DAG: [[VAL]] = load i32*, i32** [[ADDR:%.+]], 195 // CK1-DAG: [[ADDR]] = load i32**, i32*** [[ADDR2:%.+]], 196 197 // CK1-DAG: [[_BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 1 198 // CK1-DAG: [[_P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 1 199 // CK1-DAG: [[_CBP1:%.+]] = bitcast i8** [[_BP1]] to float** 200 // CK1-DAG: [[_CP1:%.+]] = bitcast i8** [[_P1]] to float** 201 // CK1-DAG: store float* [[_VAL:%.+]], float** [[_CBP1]] 202 // CK1-DAG: store float* [[_VAL]], float** [[_CP1]] 203 // CK1-DAG: [[_VAL]] = load float*, float** [[_ADDR:%.+]], 204 // CK1-DAG: [[_ADDR]] = load float**, float*** [[_ADDR2:%.+]], 205 206 // CK1: call void [[KERNEL:@.+]](i32* [[VAL]], float* [[_VAL]]) 207 #pragma omp target is_device_ptr(tr, lr) 208 { 209 ++tr,++lr; 210 } 211 } 212 213 void bar(float *&a, int *&b) { 214 foo<int>(a,b); 215 } 216 217 #endif 218 ///==========================================================================/// 219 // RUN: %clang_cc1 -no-opaque-pointers -DCK2 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK2 --check-prefix CK2-64 220 // RUN: %clang_cc1 -no-opaque-pointers -DCK2 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 221 // RUN: %clang_cc1 -no-opaque-pointers -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK2 --check-prefix CK2-64 222 // RUN: %clang_cc1 -no-opaque-pointers -DCK2 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK2 --check-prefix CK2-32 223 // RUN: %clang_cc1 -no-opaque-pointers -DCK2 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 224 // RUN: %clang_cc1 -no-opaque-pointers -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK2 --check-prefix CK2-32 225 226 // RUN: %clang_cc1 -no-opaque-pointers -DCK2 -verify -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY1 %s 227 // RUN: %clang_cc1 -no-opaque-pointers -DCK2 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 228 // RUN: %clang_cc1 -no-opaque-pointers -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY1 %s 229 // RUN: %clang_cc1 -no-opaque-pointers -DCK2 -verify -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY1 %s 230 // RUN: %clang_cc1 -no-opaque-pointers -DCK2 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 231 // RUN: %clang_cc1 -no-opaque-pointers -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY1 %s 232 // SIMD-ONLY1-NOT: {{__kmpc|__tgt}} 233 #ifdef CK2 234 235 // CK2: [[ST:%.+]] = type { double*, double** } 236 237 template <typename T> 238 struct ST { 239 T *a; 240 double *&b; 241 ST(double *&b) : a(0), b(b) {} 242 243 // CK2-LABEL: @{{.*}}foo{{.*}} 244 void foo(double *&arg) { 245 int *la = 0; 246 247 // CK2-DAG: [[RET:%.+]] = call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 248 // CK2-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 249 // CK2-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 250 // CK2-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 251 // CK2-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 252 // CK2-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 253 // CK2-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 254 255 // CK2-DAG: [[A:%.*]] = getelementptr inbounds [[STRUCT_ST:%.*]], %struct.ST* [[THIS1:%.+]], i32 0, i32 0 256 // CK2-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 257 // CK2-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 258 // CK2-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]** 259 // CK2-DAG: store [[ST]]* [[THIS1]], [[ST]]** [[CBP0]] 260 // CK2-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to double** 261 // CK2-DAG: store double** [[A]], double*** [[CP0]] 262 #pragma omp target is_device_ptr(a) 263 { 264 a++; 265 } 266 267 // CK2-DAG: [[RET:%.+]] = call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 268 // CK2-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 269 // CK2-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 270 // CK2-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 271 // CK2-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 272 // CK2-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4 273 // CK2-DAG: store i64* [[SIZE:%.+]], i64** [[SARG]] 274 // CK2-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 275 // CK2-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 276 277 // CK2-DAG: [[S:%[^,]+]] = sdiv exact i64 [[SZ:%.+]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64) 278 // CK2-DAG: [[SIZE:%[^,]+]] = getelementptr inbounds [2 x i64], [2 x i64]* %.offload_sizes, i32 0, i32 0 279 // CK2-DAG: store i64 [[S]], i64* [[SIZE]] 280 // CK2-DAG: [[B:%.*]] = getelementptr inbounds [[STRUCT_ST]], %struct.ST* [[THIS1]], i32 0, i32 1 281 // CK2-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 282 // CK2-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 283 // CK2-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]** 284 // CK2-DAG: store %struct.ST* [[THIS1]], %struct.ST** [[CBP0]] 285 // CK2-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to double*** 286 // CK2-DAG: store double*** [[B]], double**** [[CP0]] 287 #pragma omp target is_device_ptr(b) 288 { 289 b++; 290 } 291 292 // CK2-DAG: [[RET:%.+]] = call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 293 // CK2-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 294 // CK2-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 295 // CK2-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 296 // CK2-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 297 // CK2-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4 298 // CK2-DAG: store i64* [[SIZE:%.+]], i64** [[SARG]] 299 // CK2-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 300 // CK2-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 301 302 // CK2-DAG: [[A8:%.*]] = getelementptr inbounds [[STRUCT_ST]], %struct.ST* [[THIS1]], i32 0, i32 0 303 // CK2-DAG: [[B9:%.*]] = getelementptr inbounds [[STRUCT_ST]], %struct.ST* [[THIS1]], i32 0, i32 1 304 // CK2-DAG: [[S:%[^,]+]] = sdiv exact i64 [[SZ:%.+]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64) 305 // CK2-DAG: store i64 [[S]], i64* [[SIZE:%.+]] 306 307 // CK2-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 308 // CK2-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 309 // CK2-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]** 310 // CK2-DAG: store %struct.ST* [[THIS1]], %struct.ST** [[CBP0]] 311 // CH2-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to to double*** 312 // CK2-DAG: store double** [[A8]], double*** [[TMP64:%.+]] 313 #pragma omp target is_device_ptr(a, b) 314 { 315 a++; 316 b++; 317 } 318 } 319 }; 320 321 void bar(double *arg){ 322 ST<double> A(arg); 323 A.foo(arg); 324 ++arg; 325 } 326 #endif 327 ///==========================================================================/// 328 // RUN: %clang_cc1 -no-opaque-pointers -DCK3 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK3 --check-prefix CK3-64 329 // RUN: %clang_cc1 -no-opaque-pointers -DCK3 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 330 // RUN: %clang_cc1 -no-opaque-pointers -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 CK3 --check-prefix CK3-64 331 // RUN: %clang_cc1 -no-opaque-pointers -DCK3 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK3 --check-prefix CK3-32 332 // RUN: %clang_cc1 -no-opaque-pointers -DCK3 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 333 // RUN: %clang_cc1 -no-opaque-pointers -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 CK3 --check-prefix CK3-32 334 335 // RUN: %clang_cc1 -no-opaque-pointers -DCK3 -verify -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY1 %s 336 // RUN: %clang_cc1 -no-opaque-pointers -DCK3 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 337 // RUN: %clang_cc1 -no-opaque-pointers -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY1 %s 338 // RUN: %clang_cc1 -no-opaque-pointers -DCK3 -verify -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY1 %s 339 // RUN: %clang_cc1 -no-opaque-pointers -DCK3 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 340 // RUN: %clang_cc1 -no-opaque-pointers -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY1 %s 341 // SIMD-ONLY1-NOT: {{__kmpc|__tgt}} 342 #ifdef CK3 343 344 // CK3-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[SZ:64|32]]] [i{{64|32}} {{8|4}}] 345 // OMP_MAP_TARGET_PARAM = 0x20 | OMP_MAP_TO = 0x1 = 0x21 346 // CK3-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x21]]] 347 void bar() { 348 __attribute__((aligned(64))) double *ptr; 349 // CK3-DAG: [[RET:%.+]] = call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 350 // CK3-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 351 // CK3-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 352 // CK3-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 353 // CK3-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 354 // CK3-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 355 // CK3-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 356 // CK3-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 357 // CK3-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 358 // CK3-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to double*** 359 // CK3-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to double*** 360 // CK3-DAG: store double** [[PTR:%.+]], double*** [[CBP1]] 361 // CK3-DAG: store double** [[PTR]], double*** [[CP1]] 362 363 // CK3: call void [[KERNEL:@.+]](double** [[PTR]]) 364 #pragma omp target is_device_ptr(ptr) 365 *ptr = 0; 366 } 367 #endif 368 #endif 369