1 // expected-no-diagnostics 2 #ifndef HEADER 3 #define HEADER 4 5 #ifdef CK1 6 7 ///==========================================================================/// 8 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK1 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK1 9 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 10 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK1 11 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK1 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK1 12 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 13 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK1 14 15 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK1 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY10 %s 16 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 17 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY10 %s 18 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK1 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY10 %s 19 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK1 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 20 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY10 %s 21 // SIMD-ONLY10-NOT: {{__kmpc|__tgt}} 22 23 // CK1-LABEL: @.__omp_offloading_{{.*}}implicit_maps_double_complex{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0 24 25 // CK1-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16] 26 // Map types: OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 544 27 // CK1-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 544] 28 29 // CK1-LABEL: implicit_maps_double_complex{{.*}}( 30 void implicit_maps_double_complex (int a){ 31 double _Complex dc = (double)a; 32 33 // CK1-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]]) 34 // CK1-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 35 // CK1-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]] 36 // CK1-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 37 // CK1-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]] 38 // CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 39 // CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 40 // CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 41 // CK1-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 42 // CK1-DAG: store ptr [[PTR:%[^,]+]], ptr [[BP1]] 43 // CK1-DAG: store ptr [[PTR]], ptr [[P1]] 44 45 // CK1: call void [[KERNEL:@.+]](ptr [[PTR]]) 46 #pragma omp target defaultmap(alloc \ 47 : scalar) 48 { 49 dc *= dc; 50 } 51 } 52 53 // CK1: define internal void [[KERNEL]](ptr {{.*}}[[ARG:%.+]]) 54 // CK1: [[ADDR:%.+]] = alloca ptr, 55 // CK1: store ptr [[ARG]], ptr [[ADDR]], 56 // CK1: [[REF:%.+]] = load ptr, ptr [[ADDR]], 57 // CK1: {{.+}} = getelementptr inbounds nuw { double, double }, ptr [[REF]], i32 0, i32 0 58 #endif 59 ///==========================================================================/// 60 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK2 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK2 61 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK2 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 62 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK2 63 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK2 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK2 64 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK2 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 65 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK2 66 67 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK2 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY10 %s 68 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK2 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 69 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY10 %s 70 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK2 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY10 %s 71 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK2 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 72 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY10 %s 73 // SIMD-ONLY10-NOT: {{__kmpc|__tgt}} 74 #ifdef CK2 75 76 // CK2-LABEL: @.__omp_offloading_{{.*}}implicit_maps_double_complex{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0 77 78 // CK2-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16] 79 // Map types: OMP_MAP_TO | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 545 80 // CK2-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 545] 81 82 // CK2-LABEL: implicit_maps_double_complex{{.*}}( 83 void implicit_maps_double_complex (int a){ 84 double _Complex dc = (double)a; 85 86 // CK2-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]]) 87 // CK2-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 88 // CK2-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]] 89 // CK2-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 90 // CK2-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]] 91 // CK2-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 92 // CK2-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 93 // CK2-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 94 // CK2-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 95 // CK2-DAG: store ptr [[PTR:%[^,]+]], ptr [[BP1]] 96 // CK2-DAG: store ptr [[PTR]], ptr [[P1]] 97 98 // CK2: call void [[KERNEL:@.+]](ptr [[PTR]]) 99 #pragma omp target defaultmap(to \ 100 : scalar) 101 { 102 dc *= dc; 103 } 104 } 105 106 // CK2: define internal void [[KERNEL]](ptr {{.*}}[[ARG:%.+]]) 107 // CK2: [[ADDR:%.+]] = alloca ptr, 108 // CK2: store ptr [[ARG]], ptr [[ADDR]], 109 // CK2: [[REF:%.+]] = load ptr, ptr [[ADDR]], 110 // CK2: {{.+}} = getelementptr inbounds nuw { double, double }, ptr [[REF]], i32 0, i32 0 111 #endif 112 ///==========================================================================/// 113 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK3 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK3 114 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK3 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 115 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK3 116 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK3 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK3 117 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK3 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 118 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK3 119 120 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK3 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY10 %s 121 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK3 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 122 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY10 %s 123 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK3 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY10 %s 124 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK3 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 125 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY10 %s 126 // SIMD-ONLY10-NOT: {{__kmpc|__tgt}} 127 #ifdef CK3 128 129 // CK3-LABEL: @.__omp_offloading_{{.*}}implicit_maps_double_complex{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0 130 131 // CK3-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16] 132 // Map types: OMP_MAP_FROM | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 546 133 // CK3-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 546] 134 135 // CK3-LABEL: implicit_maps_double_complex{{.*}}( 136 void implicit_maps_double_complex (int a){ 137 double _Complex dc = (double)a; 138 139 // CK3-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]]) 140 // CK3-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 141 // CK3-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]] 142 // CK3-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 143 // CK3-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]] 144 // CK3-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 145 // CK3-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 146 // CK3-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 147 // CK3-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 148 // CK3-DAG: store ptr [[PTR:%[^,]+]], ptr [[BP1]] 149 // CK3-DAG: store ptr [[PTR]], ptr [[P1]] 150 151 // CK3: call void [[KERNEL:@.+]](ptr [[PTR]]) 152 #pragma omp target defaultmap(from \ 153 : scalar) 154 { 155 dc *= dc; 156 } 157 } 158 159 // CK3: define internal void [[KERNEL]](ptr {{.*}}[[ARG:%.+]]) 160 // CK3: [[ADDR:%.+]] = alloca ptr, 161 // CK3: store ptr [[ARG]], ptr [[ADDR]], 162 // CK3: [[REF:%.+]] = load ptr, ptr [[ADDR]], 163 // CK3: {{.+}} = getelementptr inbounds nuw { double, double }, ptr [[REF]], i32 0, i32 0 164 #endif 165 ///==========================================================================/// 166 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK4 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK4 --check-prefix CK4-64 167 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK4 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 168 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK4 --check-prefix CK4-64 169 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK4 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK4 --check-prefix CK4-32 170 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK4 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 171 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK4 --check-prefix CK4-32 172 173 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK4 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY6 %s 174 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK4 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 175 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY6 %s 176 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK4 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY6 %s 177 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK4 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 178 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY6 %s 179 // SIMD-ONLY6-NOT: {{__kmpc|__tgt}} 180 #ifdef CK4 181 182 // For a 32-bit targets, the value doesn't fit the size of the pointer, 183 // therefore it is passed by reference with a map 'to' specification. 184 185 // CK4-LABEL: @.__omp_offloading_{{.*}}implicit_maps_double{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0 186 187 // CK4-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8] 188 // Map types: OMP_MAP_PRIVATE_VAL | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 800 189 // CK4-64-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 800] 190 // Map types: OMP_MAP_TO | OMP_MAP_PRIVATE | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 673 191 // CK4-32-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 673] 192 193 // CK4-LABEL: implicit_maps_double{{.*}}( 194 void implicit_maps_double (int a){ 195 double d = (double)a; 196 197 // CK4-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]]) 198 // CK4-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 199 // CK4-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]] 200 // CK4-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 201 // CK4-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]] 202 // CK4-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 203 // CK4-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 204 // CK4-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 205 // CK4-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 206 207 // CK4-64-DAG: store i[[sz:64|32]] [[VAL:%[^,]+]], ptr [[BP1]] 208 // CK4-64-DAG: store i[[sz]] [[VAL]], ptr [[P1]] 209 // CK4-64-DAG: [[VAL]] = load i[[sz]], ptr [[ADDR:%.+]], 210 // CK4-64-64-DAG: store double {{.+}}, ptr [[ADDR]], 211 212 // CK4-32-DAG: store ptr [[DECL:%[^,]+]], ptr [[BP1]] 213 // CK4-32-DAG: store ptr [[DECL]], ptr [[P1]] 214 215 // CK4-64: call void [[KERNEL:@.+]](i[[sz]] [[VAL]]) 216 // CK4-32: call void [[KERNEL:@.+]](ptr [[DECL]]) 217 #pragma omp target defaultmap(firstprivate \ 218 : scalar) 219 { 220 d += 1.0; 221 } 222 } 223 224 // CK4-64: define internal void [[KERNEL]](i[[sz]] [[ARG:%.+]]) 225 // CK4-64: [[ADDR:%.+]] = alloca i[[sz]], 226 // CK4-64: store i[[sz]] [[ARG]], ptr [[ADDR]], 227 // CK4-64: {{.+}} = load double, ptr [[ADDR]], 228 229 // CK4-32: define internal void [[KERNEL]](ptr {{.+}}[[ARG:%.+]]) 230 // CK4-32: [[ADDR:%.+]] = alloca ptr, 231 // CK4-32: store ptr [[ARG]], ptr [[ADDR]], 232 // CK4-32: [[REF:%.+]] = load ptr, ptr [[ADDR]], 233 // CK4-32: {{.+}} = load double, ptr [[REF]], 234 235 #endif 236 ///==========================================================================/// 237 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK5 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK5 238 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK5 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 239 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK5 240 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK5 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK5 241 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK5 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 242 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK5 243 244 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK5 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 245 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK5 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 246 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 247 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK5 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 248 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK5 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 249 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 250 // SIMD-ONLY8-NOT: {{__kmpc|__tgt}} 251 #ifdef CK5 252 253 // CK5-LABEL: @.__omp_offloading_{{.*}}implicit_maps_array{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0 254 255 // CK5-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16] 256 // Map types: OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 544 257 // CK5-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 544] 258 259 // CK5-LABEL: implicit_maps_array{{.*}}( 260 void implicit_maps_array (int a){ 261 double darr[2] = {(double)a, (double)a}; 262 263 // CK5-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]]) 264 // CK5-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 265 // CK5-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]] 266 // CK5-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 267 // CK5-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]] 268 // CK5-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 269 // CK5-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 270 // CK5-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 271 // CK5-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 272 // CK5-DAG: store ptr [[DECL:%[^,]+]], ptr [[BP1]] 273 // CK5-DAG: store ptr [[DECL]], ptr [[P1]] 274 275 // CK5: call void [[KERNEL:@.+]](ptr [[DECL]]) 276 #pragma omp target defaultmap(alloc \ 277 : aggregate) 278 { 279 darr[0] += 1.0; 280 darr[1] += 1.0; 281 } 282 } 283 284 // CK5: define internal void [[KERNEL]](ptr {{.+}}[[ARG:%.+]]) 285 // CK5: [[ADDR:%.+]] = alloca ptr, 286 // CK5: store ptr [[ARG]], ptr [[ADDR]], 287 // CK5: [[REF:%.+]] = load ptr, ptr [[ADDR]], 288 // CK5: {{.+}} = getelementptr inbounds [2 x double], ptr [[REF]], i{{64|32}} 0, i{{64|32}} 0 289 #endif 290 ///==========================================================================/// 291 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK6 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK6 292 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK6 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 293 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK6 294 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK6 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK6 295 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK6 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 296 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK6 297 298 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK6 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 299 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK6 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 300 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 301 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK6 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 302 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK6 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 303 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 304 // SIMD-ONLY8-NOT: {{__kmpc|__tgt}} 305 #ifdef CK6 306 307 // CK6-LABEL: @.__omp_offloading_{{.*}}implicit_maps_array{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0 308 309 // CK6-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16] 310 // Map types: OMP_MAP_TO | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 545 311 // CK6-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 545] 312 313 // CK6-LABEL: implicit_maps_array{{.*}}( 314 void implicit_maps_array (int a){ 315 double darr[2] = {(double)a, (double)a}; 316 317 // CK6-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]]) 318 // CK6-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 319 // CK6-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]] 320 // CK6-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 321 // CK6-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]] 322 // CK6-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 323 // CK6-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 324 // CK6-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 325 // CK6-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 326 // CK6-DAG: store ptr [[DECL:%[^,]+]], ptr [[BP1]] 327 // CK6-DAG: store ptr [[DECL]], ptr [[P1]] 328 329 // CK6: call void [[KERNEL:@.+]](ptr [[DECL]]) 330 #pragma omp target defaultmap(to) 331 { 332 darr[0] += 1.0; 333 darr[1] += 1.0; 334 } 335 } 336 337 // CK6: define internal void [[KERNEL]](ptr {{.+}}[[ARG:%.+]]) 338 // CK6: [[ADDR:%.+]] = alloca ptr, 339 // CK6: store ptr [[ARG]], ptr [[ADDR]], 340 // CK6: [[REF:%.+]] = load ptr, ptr [[ADDR]], 341 // CK6: {{.+}} = getelementptr inbounds [2 x double], ptr [[REF]], i{{64|32}} 0, i{{64|32}} 0 342 #endif 343 ///==========================================================================/// 344 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK7 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK7 345 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK7 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 346 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK7 347 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK7 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK7 348 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK7 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 349 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK7 350 351 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK7 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 352 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK7 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 353 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 354 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK7 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 355 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK7 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 356 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 357 // SIMD-ONLY8-NOT: {{__kmpc|__tgt}} 358 #ifdef CK7 359 360 // CK7-LABEL: @.__omp_offloading_{{.*}}implicit_maps_array{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0 361 362 // CK7-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16] 363 // Map types: OMP_MAP_FROM | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 546 364 // CK7-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 546] 365 366 // CK7-LABEL: implicit_maps_array{{.*}}( 367 void implicit_maps_array (int a){ 368 double darr[2] = {(double)a, (double)a}; 369 370 // CK7-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]]) 371 // CK7-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 372 // CK7-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]] 373 // CK7-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 374 // CK7-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]] 375 // CK7-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 376 // CK7-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 377 // CK7-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 378 // CK7-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 379 // CK7-DAG: store ptr [[DECL:%[^,]+]], ptr [[BP1]] 380 // CK7-DAG: store ptr [[DECL]], ptr [[P1]] 381 382 // CK7: call void [[KERNEL:@.+]](ptr [[DECL]]) 383 #pragma omp target defaultmap(from \ 384 : aggregate) 385 { 386 darr[0] += 1.0; 387 darr[1] += 1.0; 388 } 389 } 390 391 // CK7: define internal void [[KERNEL]](ptr {{.+}}[[ARG:%.+]]) 392 // CK7: [[ADDR:%.+]] = alloca ptr, 393 // CK7: store ptr [[ARG]], ptr [[ADDR]], 394 // CK7: [[REF:%.+]] = load ptr, ptr [[ADDR]], 395 // CK7: {{.+}} = getelementptr inbounds [2 x double], ptr [[REF]], i{{64|32}} 0, i{{64|32}} 0 396 #endif 397 ///==========================================================================/// 398 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK8 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK8 399 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK8 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 400 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK8 401 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK8 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK8 402 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK8 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 403 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK8 404 405 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK8 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 406 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK8 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 407 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 408 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK8 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 409 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK8 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 410 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 411 // SIMD-ONLY8-NOT: {{__kmpc|__tgt}} 412 #ifdef CK8 413 414 // CK8-LABEL: @.__omp_offloading_{{.*}}implicit_maps_array{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0 415 416 // CK8-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16] 417 // Map types: OMP_MAP_TO | OMP_MAP_FROM | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 547 418 // CK8-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 547] 419 420 // CK8-LABEL: implicit_maps_array{{.*}}( 421 void implicit_maps_array (int a){ 422 double darr[2] = {(double)a, (double)a}; 423 424 // CK8-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]]) 425 // CK8-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 426 // CK8-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]] 427 // CK8-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 428 // CK8-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]] 429 // CK8-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 430 // CK8-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 431 // CK8-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 432 // CK8-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 433 // CK8-DAG: store ptr [[DECL:%[^,]+]], ptr [[BP1]] 434 // CK8-DAG: store ptr [[DECL]], ptr [[P1]] 435 436 // CK8: call void [[KERNEL:@.+]](ptr [[DECL]]) 437 #pragma omp target defaultmap(tofrom) 438 { 439 darr[0] += 1.0; 440 darr[1] += 1.0; 441 } 442 } 443 444 // CK8: define internal void [[KERNEL]](ptr {{.+}}[[ARG:%.+]]) 445 // CK8: [[ADDR:%.+]] = alloca ptr, 446 // CK8: store ptr [[ARG]], ptr [[ADDR]], 447 // CK8: [[REF:%.+]] = load ptr, ptr [[ADDR]], 448 // CK8: {{.+}} = getelementptr inbounds [2 x double], ptr [[REF]], i{{64|32}} 0, i{{64|32}} 0 449 #endif 450 451 ///==========================================================================/// 452 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK9 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK9 --check-prefix CK9-64 453 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK9 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 454 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK9 --check-prefix CK9-64 455 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK9 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK9 --check-prefix CK9-32 456 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK9 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 457 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK9 --check-prefix CK9-32 458 459 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK9 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY26 %s 460 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK9 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 461 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY26 %s 462 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK9 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY26 %s 463 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK9 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 464 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY26 %s 465 // SIMD-ONLY26-NOT: {{__kmpc|__tgt}} 466 467 #ifdef CK9 468 469 470 // CK9-LABEL: @.__omp_offloading_{{.*}}zero_size_section_and_private_maps{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0 471 // CK9: [[SIZE09:@.+]] = private {{.*}}constant [1 x i64] [i64 40] 472 // CK9: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i64] [i64 673] 473 474 475 // CK9-LABEL: zero_size_section_and_private_maps{{.*}}( 476 void zero_size_section_and_private_maps (int ii){ 477 int pvtArr[10]; 478 479 // Region 09 480 // CK9-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]]) 481 // CK9-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 482 // CK9-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]] 483 // CK9-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 484 // CK9-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]] 485 // CK9-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 486 // CK9-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 487 488 // CK9-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 489 // CK9-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 490 // CK9-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]] 491 // CK9-DAG: store ptr [[VAR0]], ptr [[P0]] 492 493 // CK9: call void [[CALL09:@.+]](ptr {{[^,]+}}) 494 #pragma omp target defaultmap(firstprivate \ 495 : aggregate) 496 { 497 pvtArr[5]++; 498 } 499 500 } 501 502 503 #endif 504 ///==========================================================================/// 505 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK10 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK10 506 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK10 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 507 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK10 508 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK10 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK10 509 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK10 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 510 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK10 511 512 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK10 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 513 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK10 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 514 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 515 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK10 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 516 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK10 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 517 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 518 // SIMD-ONLY8-NOT: {{__kmpc|__tgt}} 519 #ifdef CK10 520 521 522 // CK10-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0 523 524 // CK10: [[SIZE:@.+]] = private {{.*}}constant [1 x i64] [i64 {{8|4}}] 525 // Map types: OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 544 526 // CK10: [[MTYPE:@.+]] = private {{.*}}constant [1 x i64] [i64 544] 527 528 // CK10-LABEL: explicit_maps_single{{.*}}( 529 void explicit_maps_single (){ 530 int *pa; 531 532 // CK10-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]]) 533 // CK10-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 534 // CK10-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]] 535 // CK10-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 536 // CK10-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]] 537 // CK10-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 538 // CK10-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 539 540 // CK10-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 541 // CK10-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 542 // CK10-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]] 543 // CK10-DAG: store ptr [[VAR0]], ptr [[P0]] 544 545 // CK10: call void [[CALL:@.+]](ptr {{[^,]+}}) 546 #pragma omp target defaultmap(alloc \ 547 : pointer) 548 { 549 pa[50]++; 550 } 551 } 552 553 // CK10: define {{.+}}[[CALL]] 554 #endif 555 ///==========================================================================/// 556 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK11 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK11 557 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK11 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 558 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK11 559 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK11 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK11 560 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK11 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 561 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK11 562 563 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK11 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 564 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK11 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 565 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 566 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK11 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 567 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK11 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 568 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 569 // SIMD-ONLY8-NOT: {{__kmpc|__tgt}} 570 #ifdef CK11 571 572 573 // CK11-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0 574 575 // CK11: [[SIZE09:@.+]] = private {{.*}}constant [1 x i64] [i64 {{8|4}}] 576 // Map types: OMP_MAP_TO | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 545 577 // CK11: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i64] [i64 545] 578 579 // CK11-LABEL: explicit_maps_single{{.*}}( 580 void explicit_maps_single (){ 581 int *pa; 582 583 // CK11-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]]) 584 // CK11-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 585 // CK11-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]] 586 // CK11-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 587 // CK11-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]] 588 // CK11-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 589 // CK11-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 590 591 // CK11-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 592 // CK11-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 593 // CK11-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]] 594 // CK11-DAG: store ptr [[VAR0]], ptr [[P0]] 595 596 // CK11: call void [[CALL09:@.+]](ptr {{[^,]+}}) 597 #pragma omp target defaultmap(to \ 598 : pointer) 599 { 600 pa[50]++; 601 } 602 } 603 604 // CK11: define {{.+}}[[CALL09]] 605 #endif 606 ///==========================================================================/// 607 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK12 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK12 608 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK12 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 609 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK12 610 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK12 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK12 611 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK12 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 612 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK12 613 614 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK12 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 615 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK12 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 616 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 617 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK12 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 618 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK12 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 619 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 620 // SIMD-ONLY8-NOT: {{__kmpc|__tgt}} 621 #ifdef CK12 622 623 624 // CK12-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0 625 626 // CK12: [[SIZE09:@.+]] = private {{.*}}constant [1 x i64] [i64 {{8|4}}] 627 // Map types: OMP_MAP_FROM | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 546 628 // CK12: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i64] [i64 546] 629 630 // CK12-LABEL: explicit_maps_single{{.*}}( 631 void explicit_maps_single (){ 632 int *pa; 633 634 // CK12-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]]) 635 // CK12-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 636 // CK12-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]] 637 // CK12-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 638 // CK12-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]] 639 // CK12-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 640 // CK12-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 641 642 // CK12-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 643 // CK12-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 644 // CK12-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]] 645 // CK12-DAG: store ptr [[VAR0]], ptr [[P0]] 646 647 // CK12: call void [[CALL09:@.+]](ptr {{[^,]+}}) 648 #pragma omp target defaultmap(from \ 649 : pointer) 650 { 651 pa[50]++; 652 } 653 } 654 655 // CK12: define {{.+}}[[CALL09]] 656 #endif 657 ///==========================================================================/// 658 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK13 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK13 659 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK13 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 660 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK13 661 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK13 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK13 662 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK13 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 663 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK13 664 665 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK13 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 666 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK13 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 667 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 668 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK13 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 669 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK13 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 670 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 671 // SIMD-ONLY8-NOT: {{__kmpc|__tgt}} 672 #ifdef CK13 673 674 675 // CK13-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0 676 677 // CK13: [[SIZE09:@.+]] = private {{.*}}constant [1 x i64] [i64 {{8|4}}] 678 // Map types: OMP_MAP_TO | OMP_MAP_FROM | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 547 679 // CK13: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i64] [i64 547] 680 681 // CK13-LABEL: explicit_maps_single{{.*}}( 682 void explicit_maps_single (){ 683 int *pa; 684 685 // CK13-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]]) 686 // CK13-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 687 // CK13-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]] 688 // CK13-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 689 // CK13-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]] 690 // CK13-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 691 // CK13-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 692 693 // CK13-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 694 // CK13-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 695 // CK13-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]] 696 // CK13-DAG: store ptr [[VAR0]], ptr [[P0]] 697 698 // CK13: call void [[CALL09:@.+]](ptr {{[^,]+}}) 699 #pragma omp target defaultmap(tofrom \ 700 : pointer) 701 { 702 pa[50]++; 703 } 704 } 705 706 // CK13: define {{.+}}[[CALL09]] 707 #endif 708 ///==========================================================================/// 709 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK14 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK14 710 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK14 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 711 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK14 712 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK14 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK14 713 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK14 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 714 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK14 715 716 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK14 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 717 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK14 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 718 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 719 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK14 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 720 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK14 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 721 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 722 // SIMD-ONLY8-NOT: {{__kmpc|__tgt}} 723 #ifdef CK14 724 725 726 // CK14-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0 727 728 // CK14: [[SIZE09:@.+]] = private {{.*}}constant [1 x i64] zeroinitializer 729 // Map types: OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 544 730 // CK14: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i64] [i64 544] 731 732 // CK14-LABEL: explicit_maps_single{{.*}}( 733 void explicit_maps_single (){ 734 int *pa; 735 736 // CK14-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]]) 737 // CK14-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 738 // CK14-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]] 739 // CK14-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 740 // CK14-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]] 741 // CK14-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 742 // CK14-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 743 744 // CK14-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 745 // CK14-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 746 // CK14-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]] 747 // CK14-DAG: store ptr [[VAR0]], ptr [[P0]] 748 749 // CK14: call void [[CALL09:@.+]](ptr {{[^,]+}}) 750 #pragma omp target defaultmap(firstprivate \ 751 : pointer) 752 { 753 pa[50]++; 754 } 755 } 756 757 // CK14: define {{.+}}[[CALL09]] 758 #endif 759 ///==========================================================================/// 760 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK15 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK15 761 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK15 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 762 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK15 763 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK15 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK15 764 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK15 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 765 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK15 766 767 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK15 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY12 %s 768 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK15 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 769 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY12 %s 770 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK15 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY12 %s 771 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK15 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 772 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY12 %s 773 // SIMD-ONLY12-NOT: {{__kmpc|__tgt}} 774 #ifdef CK15 775 776 // CK15-LABEL: @.__omp_offloading_{{.*}}implicit_maps_variable_length_array{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0 777 778 // CK15-DAG: [[SIZES:@.+]] = {{.+}}constant [3 x i64] [i64 {{4|8}}, i64 {{4|8}}, i64 0] 779 // Map types: 780 // - OMP_MAP_LITERAL + OMP_MAP_TARGET_PARAM + OMP_MAP_IMPLICIT = 800 (vla size) 781 // - OMP_MAP_LITERAL + OMP_MAP_TARGET_PARAM + OMP_MAP_IMPLICIT = 800 (vla size) 782 // - OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 544 783 // CK15-DAG: [[TYPES:@.+]] = {{.+}}constant [3 x i64] [i64 800, i64 800, i64 544] 784 785 // CK15-LABEL: implicit_maps_variable_length_array{{.*}}( 786 void implicit_maps_variable_length_array (int a){ 787 double vla[2][a]; 788 789 // CK15-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]]) 790 // CK15-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 791 // CK15-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]] 792 // CK15-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 793 // CK15-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]] 794 // CK15-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4 795 // CK15-DAG: store ptr [[SGEP:%.+]], ptr [[SARG]] 796 // CK15-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 797 // CK15-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 798 // CK15-DAG: [[SGEP]] = getelementptr inbounds {{.+}}[[SS:%[^,]+]], i32 0, i32 0 799 800 // CK15-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 801 // CK15-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 802 // CK15-DAG: store i[[sz:64|32]] 2, ptr [[BP0]] 803 // CK15-DAG: store i[[sz]] 2, ptr [[P0]] 804 805 // CK15-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 1 806 // CK15-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 1 807 // CK15-DAG: store i[[sz]] [[VAL:%.+]], ptr [[BP1]] 808 // CK15-DAG: store i[[sz]] [[VAL]], ptr [[P1]] 809 810 // CK15-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 2 811 // CK15-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 2 812 // CK15-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[SS]], i32 0, i32 2 813 // CK15-DAG: store ptr [[DECL:%.+]], ptr [[BP2]] 814 // CK15-DAG: store ptr [[DECL]], ptr [[P2]] 815 // CK15-DAG: store i64 [[VALS2:%.+]], ptr [[S2]], 816 // CK15-DAG: [[VALS2]] = {{mul nuw i64 %.+, 8|sext i32 %.+ to i64}} 817 818 // CK15: call void [[KERNEL:@.+]](i[[sz]] {{.+}}, i[[sz]] {{.+}}, ptr [[DECL]]) 819 #pragma omp target defaultmap(alloc \ 820 : aggregate) 821 { 822 vla[1][3] += 1.0; 823 } 824 } 825 826 // CK15: define internal void [[KERNEL]](i[[sz]] [[VLA0:%.+]], i[[sz]] [[VLA1:%.+]], ptr {{.*}}[[ARG:%.+]]) 827 // CK15: [[ADDR0:%.+]] = alloca i[[sz]], 828 // CK15: [[ADDR1:%.+]] = alloca i[[sz]], 829 // CK15: [[ADDR2:%.+]] = alloca ptr, 830 // CK15: store i[[sz]] [[VLA0]], ptr [[ADDR0]], 831 // CK15: store i[[sz]] [[VLA1]], ptr [[ADDR1]], 832 // CK15: store ptr [[ARG]], ptr [[ADDR2]], 833 // CK15: {{.+}} = load i[[sz]], ptr [[ADDR0]], 834 // CK15: {{.+}} = load i[[sz]], ptr [[ADDR1]], 835 // CK15: [[REF:%.+]] = load ptr, ptr [[ADDR2]], 836 // CK15: {{.+}} = getelementptr inbounds double, ptr [[REF]], i[[sz]] %{{.+}} 837 #endif 838 ///==========================================================================/// 839 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK16 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK16 840 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK16 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 841 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK16 842 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK16 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK16 843 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK16 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 844 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK16 845 846 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK16 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY16 %s 847 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK16 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 848 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY16 %s 849 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK16 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY16 %s 850 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK16 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 851 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY16 %s 852 // SIMD-ONLY16-NOT: {{__kmpc|__tgt}} 853 #ifdef CK16 854 855 // CK16-DAG: [[ST:%.+]] = type { i32, double } 856 // CK16-LABEL: @.__omp_offloading_{{.*}}implicit_maps_struct{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0 857 // CK16-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 {{16|12}}] 858 // Map types: OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 544 859 // CK16-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 544] 860 861 class SSS { 862 public: 863 int a; 864 double b; 865 }; 866 867 // CK16-LABEL: implicit_maps_struct{{.*}}( 868 void implicit_maps_struct (int a){ 869 SSS s = {a, (double)a}; 870 871 // CK16-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]]) 872 // CK16-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 873 // CK16-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]] 874 // CK16-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 875 // CK16-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]] 876 // CK16-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 877 // CK16-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 878 // CK16-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 879 // CK16-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 880 // CK16-DAG: store ptr [[DECL:%.+]], ptr [[BP1]] 881 // CK16-DAG: store ptr [[DECL]], ptr [[P1]] 882 883 // CK16: call void [[KERNEL:@.+]](ptr [[DECL]]) 884 #pragma omp target defaultmap(alloc \ 885 : aggregate) 886 { 887 s.a += 1; 888 s.b += 1.0; 889 } 890 } 891 892 // CK16: define internal void [[KERNEL]](ptr {{.+}}[[ARG:%.+]]) 893 // CK16: [[ADDR:%.+]] = alloca ptr, 894 // CK16: store ptr [[ARG]], ptr [[ADDR]], 895 // CK16: [[REF:%.+]] = load ptr, ptr [[ADDR]], 896 // CK16: {{.+}} = getelementptr inbounds nuw [[ST]], ptr [[REF]], i32 0, i32 0 897 #endif 898 ///==========================================================================/// 899 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK17 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK17 900 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK17 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 901 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK17 902 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK17 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK17 903 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK17 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 904 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK17 905 906 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK17 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY16 %s 907 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK17 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 908 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY16 %s 909 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK17 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY16 %s 910 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK17 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 911 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY16 %s 912 // SIMD-ONLY16-NOT: {{__kmpc|__tgt}} 913 #ifdef CK17 914 915 // CK17-DAG: [[ST:%.+]] = type { i32, double } 916 // CK17-LABEL: @.__omp_offloading_{{.*}}implicit_maps_struct{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0 917 // CK17-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 {{16|12}}] 918 // Map types: OMP_MAP_TO | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 545 919 // CK17-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 545] 920 921 class SSS { 922 public: 923 int a; 924 double b; 925 }; 926 927 // CK17-LABEL: implicit_maps_struct{{.*}}( 928 void implicit_maps_struct (int a){ 929 SSS s = {a, (double)a}; 930 931 // CK17-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]]) 932 // CK17-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 933 // CK17-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]] 934 // CK17-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 935 // CK17-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]] 936 // CK17-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 937 // CK17-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 938 // CK17-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 939 // CK17-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 940 // CK17-DAG: store ptr [[DECL:%.+]], ptr [[BP1]] 941 // CK17-DAG: store ptr [[DECL]], ptr [[P1]] 942 943 // CK17: call void [[KERNEL:@.+]](ptr [[DECL]]) 944 #pragma omp target defaultmap(to \ 945 : aggregate) 946 { 947 s.a += 1; 948 s.b += 1.0; 949 } 950 } 951 952 // CK17: define internal void [[KERNEL]](ptr {{.+}}[[ARG:%.+]]) 953 // CK17: [[ADDR:%.+]] = alloca ptr, 954 // CK17: store ptr [[ARG]], ptr [[ADDR]], 955 // CK17: [[REF:%.+]] = load ptr, ptr [[ADDR]], 956 // CK17: {{.+}} = getelementptr inbounds nuw [[ST]], ptr [[REF]], i32 0, i32 0 957 #endif 958 ///==========================================================================/// 959 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK18 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK18 960 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK18 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 961 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK18 962 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK18 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK18 963 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK18 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 964 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK18 965 966 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK18 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY16 %s 967 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK18 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 968 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY16 %s 969 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK18 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY16 %s 970 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK18 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 971 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY16 %s 972 // SIMD-ONLY16-NOT: {{__kmpc|__tgt}} 973 #ifdef CK18 974 975 // CK18-DAG: [[ST:%.+]] = type { i32, double } 976 // CK18-LABEL: @.__omp_offloading_{{.*}}implicit_maps_struct{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0 977 // CK18-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 {{16|12}}] 978 // Map types: OMP_MAP_FROM | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 546 979 // CK18-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 546] 980 981 class SSS { 982 public: 983 int a; 984 double b; 985 }; 986 987 // CK18-LABEL: implicit_maps_struct{{.*}}( 988 void implicit_maps_struct (int a){ 989 SSS s = {a, (double)a}; 990 991 // CK18-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]]) 992 // CK18-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 993 // CK18-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]] 994 // CK18-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 995 // CK18-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]] 996 // CK18-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 997 // CK18-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 998 // CK18-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 999 // CK18-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 1000 // CK18-DAG: store ptr [[DECL:%.+]], ptr [[BP1]] 1001 // CK18-DAG: store ptr [[DECL]], ptr [[P1]] 1002 1003 // CK18: call void [[KERNEL:@.+]](ptr [[DECL]]) 1004 #pragma omp target defaultmap(from \ 1005 : aggregate) 1006 { 1007 s.a += 1; 1008 s.b += 1.0; 1009 } 1010 } 1011 1012 // CK18: define internal void [[KERNEL]](ptr {{.+}}[[ARG:%.+]]) 1013 // CK18: [[ADDR:%.+]] = alloca ptr, 1014 // CK18: store ptr [[ARG]], ptr [[ADDR]], 1015 // CK18: [[REF:%.+]] = load ptr, ptr [[ADDR]], 1016 // CK18: {{.+}} = getelementptr inbounds nuw [[ST]], ptr [[REF]], i32 0, i32 0 1017 #endif 1018 ///==========================================================================/// 1019 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK19 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK19 1020 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK19 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 1021 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK19 1022 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK19 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK19 1023 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK19 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 1024 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK19 1025 1026 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK19 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY16 %s 1027 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK19 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 1028 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY16 %s 1029 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK19 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY16 %s 1030 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK19 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 1031 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY16 %s 1032 // SIMD-ONLY16-NOT: {{__kmpc|__tgt}} 1033 #ifdef CK19 1034 1035 // CK19-DAG: [[ST:%.+]] = type { i32, double } 1036 // CK19-LABEL: @.__omp_offloading_{{.*}}implicit_maps_struct{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0 1037 // CK19-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 {{16|12}}] 1038 // Map types: OMP_MAP_TO | OMP_MAP_FROM | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 547 1039 // CK19-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 547] 1040 1041 class SSS { 1042 public: 1043 int a; 1044 double b; 1045 }; 1046 1047 // CK19-LABEL: implicit_maps_struct{{.*}}( 1048 void implicit_maps_struct (int a){ 1049 SSS s = {a, (double)a}; 1050 1051 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]]) 1052 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 1053 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]] 1054 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 1055 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]] 1056 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 1057 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 1058 // CK19-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 1059 // CK19-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 1060 // CK19-DAG: store ptr [[DECL:%.+]], ptr [[BP1]] 1061 // CK19-DAG: store ptr [[DECL]], ptr [[P1]] 1062 1063 // CK19: call void [[KERNEL:@.+]](ptr [[DECL]]) 1064 #pragma omp target defaultmap(tofrom \ 1065 : aggregate) 1066 { 1067 s.a += 1; 1068 s.b += 1.0; 1069 } 1070 } 1071 1072 // CK19: define internal void [[KERNEL]](ptr {{.+}}[[ARG:%.+]]) 1073 // CK19: [[ADDR:%.+]] = alloca ptr, 1074 // CK19: store ptr [[ARG]], ptr [[ADDR]], 1075 // CK19: [[REF:%.+]] = load ptr, ptr [[ADDR]], 1076 // CK19: {{.+}} = getelementptr inbounds nuw [[ST]], ptr [[REF]], i32 0, i32 0 1077 #endif 1078 ///==========================================================================/// 1079 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK20 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK20 --check-prefix CK20-64 1080 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK20 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 1081 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK20 --check-prefix CK20-64 1082 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK20 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK20 --check-prefix CK20-32 1083 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK20 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 1084 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK20 --check-prefix CK20-32 1085 1086 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK20 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY6 %s 1087 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK20 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 1088 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY6 %s 1089 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK20 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY6 %s 1090 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK20 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 1091 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY6 %s 1092 // SIMD-ONLY6-NOT: {{__kmpc|__tgt}} 1093 #ifdef CK20 1094 1095 // For a 32-bit targets, the value doesn't fit the size of the pointer, 1096 // therefore it is passed by reference with a map 'to' specification. 1097 1098 // CK20-LABEL: @.__omp_offloading_{{.*}}implicit_maps_double{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0 1099 1100 // CK20-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8] 1101 // Map types: OMP_MAP_LITERAL | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 800 1102 // CK20-64-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 800] 1103 // Map types: OMP_MAP_TO | OMP_MAP_PRIVATE | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 673 1104 // CK20-32-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 673] 1105 1106 // CK20-LABEL: implicit_maps_double{{.*}}( 1107 void implicit_maps_double (int a){ 1108 double d = (double)a; 1109 1110 // CK20-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]]) 1111 // CK20-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 1112 // CK20-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]] 1113 // CK20-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 1114 // CK20-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]] 1115 // CK20-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 1116 // CK20-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 1117 // CK20-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 1118 // CK20-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 1119 1120 // CK20-64-DAG: store i[[sz:64|32]] [[VAL:%[^,]+]], ptr [[BP1]] 1121 // CK20-64-DAG: store i[[sz]] [[VAL]], ptr [[P1]] 1122 // CK20-64-DAG: [[VAL]] = load i[[sz]], ptr [[ADDR:%.+]], 1123 // CK20-64-64-DAG: store double {{.+}}, ptr [[ADDR]], 1124 1125 // CK20-32-DAG: store ptr [[DECL:%[^,]+]], ptr [[BP1]] 1126 // CK20-32-DAG: store ptr [[DECL]], ptr [[P1]] 1127 1128 // CK20-64: call void [[KERNEL:@.+]](i[[sz]] [[VAL]]) 1129 // CK20-32: call void [[KERNEL:@.+]](ptr [[DECL]]) 1130 #pragma omp target defaultmap(default \ 1131 : scalar) 1132 { 1133 d += 1.0; 1134 } 1135 } 1136 1137 // CK20-64: define internal void [[KERNEL]](i[[sz]] [[ARG:%.+]]) 1138 // CK20-64: [[ADDR:%.+]] = alloca i[[sz]], 1139 // CK20-64: store i[[sz]] [[ARG]], ptr [[ADDR]], 1140 // CK20-64: {{.+}} = load double, ptr [[ADDR]], 1141 1142 // CK20-32: define internal void [[KERNEL]](ptr {{.+}}[[ARG:%.+]]) 1143 // CK20-32: [[ADDR:%.+]] = alloca ptr, 1144 // CK20-32: store ptr [[ARG]], ptr [[ADDR]], 1145 // CK20-32: [[REF:%.+]] = load ptr, ptr [[ADDR]], 1146 // CK20-32: {{.+}} = load double, ptr [[REF]], 1147 1148 #endif 1149 ///==========================================================================/// 1150 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK21 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK21 1151 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK21 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 1152 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK21 1153 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK21 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK21 1154 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK21 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 1155 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK21 1156 1157 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK21 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY16 %s 1158 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK21 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 1159 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY16 %s 1160 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK21 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY16 %s 1161 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK21 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 1162 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY16 %s 1163 // SIMD-ONLY16-NOT: {{__kmpc|__tgt}} 1164 #ifdef CK21 1165 1166 // CK21-DAG: [[ST:%.+]] = type { i32, double } 1167 // CK21-LABEL: @.__omp_offloading_{{.*}}implicit_maps_struct{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0 1168 // CK21-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 {{16|12}}] 1169 // Map types: OMP_MAP_TO + OMP_MAP_FROM + OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 547 1170 // CK21-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 547] 1171 1172 class SSS { 1173 public: 1174 int a; 1175 double b; 1176 }; 1177 1178 // CK21-LABEL: implicit_maps_struct{{.*}}( 1179 void implicit_maps_struct (int a){ 1180 SSS s = {a, (double)a}; 1181 1182 // CK21-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]]) 1183 // CK21-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 1184 // CK21-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]] 1185 // CK21-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 1186 // CK21-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]] 1187 // CK21-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 1188 // CK21-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 1189 // CK21-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 1190 // CK21-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 1191 // CK21-DAG: store ptr [[DECL:%.+]], ptr [[BP1]] 1192 // CK21-DAG: store ptr [[DECL]], ptr [[P1]] 1193 1194 // CK21: call void [[KERNEL:@.+]](ptr [[DECL]]) 1195 #pragma omp target defaultmap(default \ 1196 : aggregate) 1197 { 1198 s.a += 1; 1199 s.b += 1.0; 1200 } 1201 } 1202 1203 // CK21: define internal void [[KERNEL]](ptr {{.+}}[[ARG:%.+]]) 1204 // CK21: [[ADDR:%.+]] = alloca ptr, 1205 // CK21: store ptr [[ARG]], ptr [[ADDR]], 1206 // CK21: [[REF:%.+]] = load ptr, ptr [[ADDR]], 1207 // CK21: {{.+}} = getelementptr inbounds nuw [[ST]], ptr [[REF]], i32 0, i32 0 1208 #endif 1209 ///==========================================================================/// 1210 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK22 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK22 1211 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK22 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 1212 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK22 1213 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK22 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK22 1214 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK22 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 1215 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK22 1216 1217 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK22 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY9 %s 1218 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK22 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 1219 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY9 %s 1220 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK22 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY9 %s 1221 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK22 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 1222 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY9 %s 1223 // SIMD-ONLY9-NOT: {{__kmpc|__tgt}} 1224 #ifdef CK22 1225 1226 // CK22-LABEL: @.__omp_offloading_{{.*}}implicit_maps_pointer{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0 1227 1228 // CK22-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] zeroinitializer 1229 // Map types: OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 544 1230 // CK22-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 544] 1231 1232 // CK22-LABEL: implicit_maps_pointer{{.*}}( 1233 void implicit_maps_pointer (){ 1234 double *ddyn; 1235 1236 // CK22-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]]) 1237 // CK22-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 1238 // CK22-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]] 1239 // CK22-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 1240 // CK22-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]] 1241 // CK22-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 1242 // CK22-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 1243 // CK22-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 1244 // CK22-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 1245 // CK22-DAG: store ptr [[PTR:%[^,]+]], ptr [[BP1]] 1246 // CK22-DAG: store ptr [[PTR]], ptr [[P1]] 1247 1248 // CK22: call void [[KERNEL:@.+]](ptr [[PTR]]) 1249 #pragma omp target defaultmap(default \ 1250 : pointer) 1251 { 1252 ddyn[0] += 1.0; 1253 ddyn[1] += 1.0; 1254 } 1255 } 1256 1257 // CK22: define internal void [[KERNEL]](ptr {{.*}}[[ARG:%.+]]) 1258 // CK22: [[ADDR:%.+]] = alloca ptr, 1259 // CK22: store ptr [[ARG]], ptr [[ADDR]], 1260 // CK22: [[REF:%.+]] = load ptr, ptr [[ADDR]], 1261 // CK22: {{.+}} = getelementptr inbounds double, ptr [[REF]], i{{64|32}} 0 1262 1263 #endif 1264 ///==========================================================================/// 1265 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK23 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK23 --check-prefix CK23-64 1266 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK23 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 1267 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck %s --check-prefix CK23 --check-prefix CK23-64 1268 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK23 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK23 --check-prefix CK23-32 1269 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK23 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 1270 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck %s --check-prefix CK23 --check-prefix CK23-32 1271 1272 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK23 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s 1273 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK23 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 1274 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s 1275 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK23 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s 1276 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK23 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 1277 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s 1278 // SIMD-ONLY0-NOT: {{__kmpc|__tgt}} 1279 #ifdef CK23 1280 1281 double *g; 1282 1283 // CK23: @g ={{.*}} global ptr 1284 // CK23: [[SIZES00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} {{8|4}}] 1285 // CK23: [[TYPES00:@.+]] = {{.+}}constant [1 x i64] [i64 288] 1286 1287 // CK23: [[SIZES01:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] {{8|4}}] 1288 // CK23: [[TYPES01:@.+]] = {{.+}}constant [1 x i64] [i64 288] 1289 1290 // CK23: [[SIZES02:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] {{8|4}}] 1291 // CK23: [[TYPES02:@.+]] = {{.+}}constant [1 x i64] [i64 288] 1292 1293 // CK23: [[SIZES03:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] {{8|4}}] 1294 // CK23: [[TYPES03:@.+]] = {{.+}}constant [1 x i64] [i64 288] 1295 1296 // CK23: [[SIZES04:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] {{8|4}}] 1297 // CK23: [[TYPES04:@.+]] = {{.+}}constant [1 x i64] [i64 288] 1298 1299 // CK23: [[SIZES05:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] {{8|4}}] 1300 // CK23: [[TYPES05:@.+]] = {{.+}}constant [1 x i64] [i64 288] 1301 1302 // CK23: [[SIZES06:@.+]] = {{.+}}constant [2 x i[[sz]]] [i[[sz]] {{8|4}}, i[[sz]] {{8|4}}] 1303 // CK23: [[TYPES06:@.+]] = {{.+}}constant [2 x i64] [i64 288, i64 288] 1304 1305 // CK23-LABEL: @_Z3foo{{.*}}( 1306 template<typename T> 1307 void foo(float *&lr, T *&tr) { 1308 float *l; 1309 T *t; 1310 1311 // CK23-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]]) 1312 // CK23-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 1313 // CK23-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]] 1314 // CK23-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 1315 // CK23-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]] 1316 // CK23-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 1317 // CK23-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 1318 // CK23-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 1319 // CK23-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 1320 // CK23-DAG: store ptr [[VAL:%.+]], ptr [[BP1]] 1321 // CK23-DAG: store ptr [[VAL]], ptr [[P1]] 1322 // CK23-DAG: [[VAL]] = load ptr, ptr [[ADDR:@g]], 1323 1324 // CK23: call void [[KERNEL:@.+]](ptr [[VAL]]) 1325 #pragma omp target is_device_ptr(g) defaultmap(none \ 1326 : pointer) 1327 { 1328 ++g; 1329 } 1330 1331 // CK23-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]]) 1332 // CK23-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 1333 // CK23-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]] 1334 // CK23-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 1335 // CK23-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]] 1336 // CK23-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 1337 // CK23-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 1338 // CK23-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 1339 // CK23-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 1340 // CK23-DAG: store ptr [[VAL:%.+]], ptr [[BP1]] 1341 // CK23-DAG: store ptr [[VAL]], ptr [[P1]] 1342 // CK23-DAG: [[VAL]] = load ptr, ptr [[ADDR:%.+]], 1343 1344 // CK23: call void [[KERNEL:@.+]](ptr [[VAL]]) 1345 #pragma omp target is_device_ptr(l) defaultmap(none \ 1346 : pointer) 1347 { 1348 ++l; 1349 } 1350 1351 // CK23-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]]) 1352 // CK23-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 1353 // CK23-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]] 1354 // CK23-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 1355 // CK23-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]] 1356 // CK23-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 1357 // CK23-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 1358 // CK23-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 1359 // CK23-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 1360 // CK23-DAG: store ptr [[VAL:%.+]], ptr [[BP1]] 1361 // CK23-DAG: store ptr [[VAL]], ptr [[P1]] 1362 // CK23-DAG: [[VAL]] = load ptr, ptr [[ADDR:%.+]], 1363 1364 // CK23: call void [[KERNEL:@.+]](ptr [[VAL]]) 1365 #pragma omp target is_device_ptr(t) defaultmap(none \ 1366 : pointer) 1367 { 1368 ++t; 1369 } 1370 1371 // CK23-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]]) 1372 // CK23-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 1373 // CK23-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]] 1374 // CK23-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 1375 // CK23-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]] 1376 // CK23-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 1377 // CK23-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 1378 // CK23-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 1379 // CK23-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 1380 // CK23-DAG: store ptr [[VAL:%.+]], ptr [[BP1]] 1381 // CK23-DAG: store ptr [[VAL]], ptr [[P1]] 1382 // CK23-DAG: [[VAL]] = load ptr, ptr [[ADDR:%.+]], 1383 // CK23-DAG: [[ADDR]] = load ptr, ptr [[ADDR2:%.+]], 1384 1385 // CK23: call void [[KERNEL:@.+]](ptr [[VAL]]) 1386 #pragma omp target is_device_ptr(lr) defaultmap(none \ 1387 : pointer) 1388 { 1389 ++lr; 1390 } 1391 1392 // CK23-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]]) 1393 // CK23-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 1394 // CK23-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]] 1395 // CK23-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 1396 // CK23-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]] 1397 // CK23-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 1398 // CK23-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 1399 // CK23-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 1400 // CK23-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 1401 // CK23-DAG: store ptr [[VAL:%.+]], ptr [[BP1]] 1402 // CK23-DAG: store ptr [[VAL]], ptr [[P1]] 1403 // CK23-DAG: [[VAL]] = load ptr, ptr [[ADDR:%.+]], 1404 // CK23-DAG: [[ADDR]] = load ptr, ptr [[ADDR2:%.+]], 1405 1406 // CK23: call void [[KERNEL:@.+]](ptr [[VAL]]) 1407 #pragma omp target is_device_ptr(tr) defaultmap(none \ 1408 : pointer) 1409 { 1410 ++tr; 1411 } 1412 1413 // CK23-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]]) 1414 // CK23-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 1415 // CK23-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]] 1416 // CK23-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 1417 // CK23-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]] 1418 // CK23-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 1419 // CK23-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 1420 // CK23-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 1421 // CK23-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 1422 // CK23-DAG: store ptr [[VAL:%.+]], ptr [[BP1]] 1423 // CK23-DAG: store ptr [[VAL]], ptr [[P1]] 1424 // CK23-DAG: [[VAL]] = load ptr, ptr [[ADDR:%.+]], 1425 // CK23-DAG: [[ADDR]] = load ptr, ptr [[ADDR2:%.+]], 1426 1427 // CK23: call void [[KERNEL:@.+]](ptr [[VAL]]) 1428 #pragma omp target is_device_ptr(tr, lr) defaultmap(none \ 1429 : pointer) 1430 { 1431 ++tr; 1432 } 1433 1434 // CK23-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]]) 1435 // CK23-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 1436 // CK23-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]] 1437 // CK23-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 1438 // CK23-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]] 1439 // CK23-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 1440 // CK23-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 1441 // CK23-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 1442 // CK23-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 1443 // CK23-DAG: store ptr [[VAL:%.+]], ptr [[BP1]] 1444 // CK23-DAG: store ptr [[VAL]], ptr [[P1]] 1445 // CK23-DAG: [[VAL]] = load ptr, ptr [[ADDR:%.+]], 1446 // CK23-DAG: [[ADDR]] = load ptr, ptr [[ADDR2:%.+]], 1447 1448 // CK23-DAG: [[_BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 1 1449 // CK23-DAG: [[_P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 1 1450 // CK23-DAG: store ptr [[_VAL:%.+]], ptr [[_BP1]] 1451 // CK23-DAG: store ptr [[_VAL]], ptr [[_P1]] 1452 // CK23-DAG: [[_VAL]] = load ptr, ptr [[_ADDR:%.+]], 1453 // CK23-DAG: [[_ADDR]] = load ptr, ptr [[_ADDR2:%.+]], 1454 1455 // CK23: call void [[KERNEL:@.+]](ptr [[VAL]], ptr [[_VAL]]) 1456 #pragma omp target is_device_ptr(tr, lr) defaultmap(none \ 1457 : pointer) 1458 { 1459 ++tr,++lr; 1460 } 1461 } 1462 1463 void bar(float *&a, int *&b) { 1464 foo<int>(a,b); 1465 } 1466 1467 #endif 1468 ///==========================================================================/// 1469 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK24 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK24 --check-prefix CK24-64 1470 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK24 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 1471 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK24 --check-prefix CK24-64 1472 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK24 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK24 --check-prefix CK24-32 1473 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK24 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 1474 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK24 --check-prefix CK24-32 1475 1476 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK24 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s 1477 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK24 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 1478 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s 1479 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK24 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s 1480 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK24 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 1481 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s 1482 // SIMD-ONLY18-NOT: {{__kmpc|__tgt}} 1483 #ifdef CK24 1484 1485 // CK24-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0 1486 // CK24: [[SIZE00:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4] 1487 // CK24: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i64] [i64 1059] 1488 1489 // CK24-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0 1490 // CK24: [[SIZE01:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4] 1491 // CK24: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i64] [i64 1063] 1492 1493 // CK24-LABEL: explicit_maps_single{{.*}}( 1494 void explicit_maps_single (int ii){ 1495 // Map of a scalar. 1496 int a = ii; 1497 1498 // Close. 1499 // Region 00 1500 // CK24-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]]) 1501 // CK24-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 1502 // CK24-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]] 1503 // CK24-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 1504 // CK24-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]] 1505 // CK24-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 1506 // CK24-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 1507 1508 // CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 1509 // CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 1510 // CK24-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]] 1511 // CK24-DAG: store ptr [[VAR0]], ptr [[P0]] 1512 1513 // CK24: call void [[CALL00:@.+]](ptr {{[^,]+}}) 1514 #pragma omp target map(close, tofrom \ 1515 : a) defaultmap(none \ 1516 : scalar) 1517 { 1518 a++; 1519 } 1520 1521 // Always Close. 1522 // Region 01 1523 // CK24-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]]) 1524 // CK24-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 1525 // CK24-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]] 1526 // CK24-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 1527 // CK24-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]] 1528 // CK24-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 1529 // CK24-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 1530 1531 // CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 1532 // CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 1533 // CK24-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]] 1534 // CK24-DAG: store ptr [[VAR0]], ptr [[P0]] 1535 1536 // CK24: call void [[CALL01:@.+]](ptr {{[^,]+}}) 1537 #pragma omp target map(always close tofrom \ 1538 : a) defaultmap(none \ 1539 : scalar) 1540 { 1541 a++; 1542 } 1543 } 1544 // CK24: define {{.+}}[[CALL00]] 1545 // CK24: define {{.+}}[[CALL01]] 1546 1547 #endif 1548 ///==========================================================================/// 1549 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK25 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK25 --check-prefix CK25-64 1550 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK25 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 1551 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK25 --check-prefix CK25-64 1552 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK25 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK25 --check-prefix CK25-32 1553 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK25 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 1554 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK25 --check-prefix CK25-32 1555 1556 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK25 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s 1557 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK25 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 1558 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s 1559 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK25 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s 1560 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK25 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 1561 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s 1562 // SIMD-ONLY18-NOT: {{__kmpc|__tgt}} 1563 #ifdef CK25 1564 1565 extern int x; 1566 #pragma omp declare target to(x) 1567 1568 // CK25-LABEL: @.__omp_offloading_{{.*}}declare_target_to{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0 1569 1570 void declare_target_to() 1571 { 1572 // CK25: [[C:%.+]] = load i32, ptr @x 1573 // CK25: [[INC:%.+]] = add nsw i32 [[C]], 1 1574 // CK25: store i32 [[INC]], ptr @x 1575 // CK25: ret void 1576 #pragma omp target defaultmap(none : scalar) 1577 { 1578 x++; 1579 } 1580 } 1581 1582 #endif 1583 ///==========================================================================/// 1584 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK26 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK26 --check-prefix CK26-64 1585 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK26 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 1586 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK26 --check-prefix CK26-64 1587 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK26 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK26 --check-prefix CK26-32 1588 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK26 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 1589 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK26 --check-prefix CK26-32 1590 1591 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK26 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s 1592 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK26 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 1593 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s 1594 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK26 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s 1595 // RUN: %clang_cc1 -no-enable-noundef-analysis -DCK26 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 1596 // RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s 1597 // SIMD-ONLY18-NOT: {{__kmpc|__tgt}} 1598 #ifdef CK26 1599 1600 // CK26-DAG: [[SIZES:@.+]] = {{.+}}constant [3 x i64] [i64 4, i64 4096, i64 {{.+}}] 1601 // Map types: OMP_MAP_TO | OMP_MAP_FROM | OMP_MAP_PTR_AND_OBJ | OMP_MAP_IMPLICIT = 531 1602 // CK26-DAG: [[TYPES:@.+]] = {{.+}}constant [3 x i64] [i64 531, i64 531, i64 531] 1603 1604 float Vector[1024]; 1605 #pragma omp declare target link(Vector) 1606 1607 extern int a; 1608 #pragma omp declare target link(a) 1609 1610 double *ptr; 1611 #pragma omp declare target link(ptr) 1612 1613 void declare_target_link() 1614 { 1615 #pragma omp target defaultmap(none:scalar) defaultmap(none:aggregate) defaultmap(none:pointer) 1616 { 1617 1618 // CK26-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]]) 1619 // CK26-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 1620 // CK26-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]] 1621 // CK26-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 1622 // CK26-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]] 1623 // CK26-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 1624 // CK26-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 1625 // CK26-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 2 1626 // CK26-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 2 1627 // CK26-DAG: store ptr @ptr_decl_tgt_ref_ptr, ptr [[BP1]] 1628 // CK26-DAG: store ptr @ptr, ptr [[P1]] 1629 1630 Vector[a]++; 1631 ptr++; 1632 } 1633 } 1634 1635 #endif 1636 #endif 1637