1 // RUN: %clang_cc1 -DCK1 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s 2 // RUN: %clang_cc1 -DCK1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 3 // RUN: %clang_cc1 -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 4 5 // RUN: %clang_cc1 -DCK1 -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 6 // RUN: %clang_cc1 -DCK1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 7 // RUN: %clang_cc1 -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 8 // SIMD-ONLY0-NOT: {{__kmpc|__tgt}} 9 // expected-no-diagnostics 10 11 #ifndef HEADER 12 #define HEADER 13 14 // CHECK-DAG: [[SIZES1:@.+]] = private unnamed_addr constant [6 x i64] [i64 4, i64 16, i64 4, i64 4, i64 0, i64 4] 15 // 64 = 0x40 = OMP_MAP_RETURN_PARAM 16 // CHECK-DAG: [[MAPTYPES1:@.+]] = private unnamed_addr constant [6 x i64] [i64 67, i64 115, i64 51, i64 67, i64 67, i64 67] 17 // CHECK-DAG: [[SIZES2:@.+]] = private unnamed_addr constant [6 x i64] [i64 0, i64 4, i64 16, i64 4, i64 4, i64 0] 18 // 0 = OMP_MAP_NONE 19 // 281474976710720 = 0x1000000000040 = OMP_MAP_MEMBER_OF | OMP_MAP_RETURN_PARAM 20 // CHECK-DAG: [[MAPTYPES2:@.+]] = private unnamed_addr constant [6 x i64] [i64 0, i64 281474976710723, i64 281474976710739, i64 281474976710739, i64 281474976710675, i64 281474976710723] 21 struct S { 22 int a = 0; 23 int *ptr = &a; 24 int &ref = a; 25 int arr[4]; 26 S() {} 27 void foo() { 28 #pragma omp target data map(tofrom: a, ptr [3:4], ref, ptr[0], arr[:a]) use_device_addr(a, ptr [3:4], ref, ptr[0], arr[:a]) 29 ++a, ++*ptr, ++ref, ++arr[0]; 30 } 31 }; 32 33 int main() { 34 float a = 0; 35 float *ptr = &a; 36 float &ref = a; 37 float arr[4]; 38 float vla[(int)a]; 39 S s; 40 s.foo(); 41 #pragma omp target data map(tofrom: a, ptr [3:4], ref, ptr[0], arr[:(int)a], vla[0]) use_device_addr(a, ptr [3:4], ref, ptr[0], arr[:(int)a], vla[0]) 42 ++a, ++*ptr, ++ref, ++arr[0], ++vla[0]; 43 return a; 44 } 45 46 // CHECK-LABEL: @main() 47 // CHECK: [[A_ADDR:%.+]] = alloca float, 48 // CHECK: [[PTR_ADDR:%.+]] = alloca ptr, 49 // CHECK: [[REF_ADDR:%.+]] = alloca ptr, 50 // CHECK: [[ARR_ADDR:%.+]] = alloca [4 x float], 51 // CHECK: [[BPTRS:%.+]] = alloca [6 x ptr], 52 // CHECK: [[PTRS:%.+]] = alloca [6 x ptr], 53 // CHECK: [[MAP_PTRS:%.+]] = alloca [6 x ptr], 54 // CHECK: [[SIZES:%.+]] = alloca [6 x i64], 55 // CHECK: [[VLA_ADDR:%.+]] = alloca float, i64 %{{.+}}, 56 // CHECK: [[PTR:%.+]] = load ptr, ptr [[PTR_ADDR]], 57 // CHECK-NEXT: [[ARR_IDX:%.+]] = getelementptr inbounds nuw float, ptr [[PTR]], i64 3 58 // CHECK: [[P5:%.+]] = load ptr, ptr [[PTR_ADDR]], align 8 59 // CHECK-NEXT: [[ARR_IDX1:%.+]] = getelementptr inbounds float, ptr [[P5]], i64 0 60 // CHECK: [[P7:%.+]] = load ptr, ptr [[REF_ADDR]], 61 // CHECK-NEXT: [[REF:%.+]] = load ptr, ptr [[REF_ADDR]], 62 // CHECK-NEXT: [[ARR_IDX2:%.+]] = getelementptr inbounds nuw [4 x float], ptr [[ARR_ADDR]], i64 0, i64 0 63 // CHECK: [[P10:%.+]] = mul nuw i64 {{.+}}, 4 64 // CHECK-NEXT: [[ARR_IDX5:%.+]] = getelementptr inbounds float, ptr [[VLA_ADDR]], i64 0 65 // CHECK-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[SIZES]], ptr align 8 [[SIZES1]], i64 48, i1 false) 66 // CHECK: [[BPTR0:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 0 67 // CHECK: store ptr [[A_ADDR]], ptr [[BPTR0]], 68 // CHECK: [[PTR0:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 0 69 // CHECK: store ptr [[A_ADDR]], ptr [[PTR0]], 70 // CHECK: [[BPTR1:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 1 71 // CHECK: store ptr [[PTR_ADDR]], ptr [[BPTR1]], 72 // CHECK: [[PTR1:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 1 73 // CHECK: store ptr [[ARR_IDX]], ptr [[PTR1]], 74 // CHECK: [[BPTR2:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 2 75 // CHECK: store ptr [[PTR_ADDR]], ptr [[BPTR2]], 76 // CHECK: [[PTR2:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 2 77 // CHECK: store ptr [[ARR_IDX1]], ptr [[PTR2]], 78 // CHECK: [[BPTR3:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 3 79 // CHECK: store ptr [[P7]], ptr [[BPTR3]], 80 // CHECK: [[PTR3:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 3 81 // CHECK: store ptr [[REF]], ptr [[PTR3]], 82 // CHECK: [[BPTR4:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 4 83 // CHECK: store ptr [[ARR_ADDR]], ptr [[BPTR4]], align 84 // CHECK: [[PTR4:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 4 85 // CHECK: store ptr [[ARR_IDX2]], ptr [[PTR4]], align 8 86 // CHECK: [[SIZE_PTR:%.+]] = getelementptr inbounds [6 x i64], ptr [[SIZES]], i32 0, i32 4 87 // CHECK: store i64 [[P10:%.+]], ptr [[SIZE_PTR]], align 8 88 // CHECK: [[MAP_PTR:%.+]] = getelementptr inbounds [6 x ptr], ptr [[MAP_PTRS]], i64 0, i64 4 89 // CHECK: store ptr null, ptr [[MAP_PTR]], align 8 90 // CHECK: [[BPTR5:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 5 91 // CHECK: store ptr [[VLA_ADDR]], ptr [[BPTR5]], 92 // CHECK: [[PTR5:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 5 93 // CHECK: store ptr [[ARR_IDX5]], ptr [[PTR5]], 94 95 // CHECK: [[BPTR:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 0 96 // CHECK: [[PTR:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 0 97 // CHECK: [[SIZE:%.+]] = getelementptr inbounds [6 x i64], ptr [[SIZES]], i32 0, i32 0 98 // CHECK: call void @__tgt_target_data_begin_mapper(ptr @{{.+}}, i64 -1, i32 6, ptr [[BPTR]], ptr [[PTR]], ptr [[SIZE]], ptr [[MAPTYPES1]], ptr null, ptr null) 99 // CHECK: [[A_REF:%.+]] = load ptr, ptr [[BPTR0]], 100 // CHECK: [[REF_REF:%.+]] = load ptr, ptr [[BPTR3]], 101 // CHECK: store ptr [[REF_REF]], ptr [[TMP_REF_ADDR:%.+]], 102 // CHECK: [[ARR_REF:%.+]] = load ptr, ptr [[BPTR4]], 103 // CHECK: [[VLA_REF:%.+]] = load ptr, ptr [[BPTR5]], 104 // CHECK: [[A:%.+]] = load float, ptr [[A_REF]], 105 // CHECK: [[INC:%.+]] = fadd float [[A]], 1.000000e+00 106 // CHECK: store float [[INC]], ptr [[A_REF]], 107 // CHECK: [[PTR:%.+]] = load ptr, ptr [[BPTR1]], 108 // CHECK: [[VAL:%.+]] = load float, ptr [[PTR]], 109 // CHECK: [[INC:%.+]] = fadd float [[VAL]], 1.000000e+00 110 // CHECK: store float [[INC]], ptr [[PTR]], 111 // CHECK: [[REF_ADDR:%.+]] = load ptr, ptr [[TMP_REF_ADDR]], 112 // CHECK: [[REF:%.+]] = load float, ptr [[REF_ADDR]], 113 // CHECK: [[INC:%.+]] = fadd float [[REF]], 1.000000e+00 114 // CHECK: store float [[INC]], ptr [[REF_ADDR]], 115 // CHECK: [[ARR0_ADDR:%.+]] = getelementptr inbounds [4 x float], ptr [[ARR_REF]], i64 0, i64 0 116 // CHECK: [[ARR0:%.+]] = load float, ptr [[ARR0_ADDR]], 117 // CHECK: [[INC:%.+]] = fadd float [[ARR0]], 1.000000e+00 118 // CHECK: store float [[INC]], ptr [[ARR0_ADDR]], 119 // CHECK: [[VLA0_ADDR:%.+]] = getelementptr inbounds float, ptr [[VLA_REF]], i64 0 120 // CHECK: [[VLA0:%.+]] = load float, ptr [[VLA0_ADDR]], 121 // CHECK: [[INC:%.+]] = fadd float [[VLA0]], 1.000000e+00 122 // CHECK: store float [[INC]], ptr [[VLA0_ADDR]], 123 // CHECK: [[BPTR:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 0 124 // CHECK: [[PTR:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 0 125 // CHECK: [[SIZE:%.+]] = getelementptr inbounds [6 x i64], ptr [[SIZES]], i32 0, i32 0 126 // CHECK: call void @__tgt_target_data_end_mapper(ptr @{{.+}}, i64 -1, i32 6, ptr [[BPTR]], ptr [[PTR]], ptr [[SIZE]], ptr [[MAPTYPES1]], ptr null, ptr null) 127 128 // CHECK: foo 129 // CHECK: [[BPTRS:%.+]] = alloca [6 x ptr], 130 // CHECK: [[PTRS:%.+]] = alloca [6 x ptr], 131 // CHECK: [[MAP_PTRS:%.+]] = alloca [6 x ptr], 132 // CHECK: [[SIZES:%.+]] = alloca [6 x i64], 133 // CHECK: [[A_ADDR:%.+]] = getelementptr inbounds nuw %struct.S, ptr [[THIS:%.+]], i32 0, i32 0 134 // CHECK: [[PTR_ADDR:%.+]] = getelementptr inbounds nuw %struct.S, ptr [[THIS]], i32 0, i32 1 135 // CHECK: [[ARR_IDX:%.+]] = getelementptr inbounds nuw i32, ptr %{{.+}}, i64 3 136 // CHECK: [[REF_REF:%.+]] = getelementptr inbounds nuw %struct.S, ptr [[THIS]], i32 0, i32 2 137 // CHECK: [[REF_PTR:%.+]] = load ptr, ptr [[REF_REF]], 138 // CHECK-NEXT: [[P3:%.+]] = getelementptr inbounds nuw %struct.S, ptr [[THIS]], i32 0, i32 1 139 // CHECK: [[ARR_IDX5:%.+]] = getelementptr inbounds i32, ptr {{.+}}, i64 0 140 // CHECK: [[ARR_ADDR:%.+]] = getelementptr inbounds nuw %struct.S, ptr [[THIS]], i32 0, i32 3 141 142 // CHECK: [[ARR_IDX6:%.+]] = getelementptr inbounds nuw [4 x i32], ptr [[ARR_ADDR]], i64 0, i64 0 143 // CHECK: [[A_ADDR2:%.+]] = getelementptr inbounds nuw %struct.S, ptr [[THIS]], i32 0, i32 0 144 // CHECK: [[P4:%.+]] = mul nuw i64 [[CONV:%.+]], 4 145 // CHECK: [[A_ADDR3:%.+]] = getelementptr inbounds nuw %struct.S, ptr [[THIS]], i32 0, i32 0 146 // CHECK: [[L5:%.+]] = load i32, ptr [[A_ADDR3]] 147 // CHECK: [[L6:%.+]] = sext i32 [[L5]] to i64 148 // CHECK: [[LB_ADD_LEN:%lb_add_len]] = add nsw i64 -1, [[L6]] 149 // CHECK: [[ARR_ADDR9:%.+]] = getelementptr inbounds nuw %struct.S, ptr [[THIS]], i32 0, i32 3 150 // CHECK: [[ARR_IDX10:%arrayidx.+]] = getelementptr inbounds nuw [4 x i32], ptr [[ARR_ADDR9]], i64 0, i64 %lb_add_len 151 // CHECK: [[ARR_END:%.+]] = getelementptr i32, ptr [[ARR_IDX10]], i32 1 152 // CHECK: [[E:%.+]] = ptrtoint ptr [[ARR_END]] to i64 153 // CHECK: [[B:%.+]] = ptrtoint ptr [[A_ADDR]] to i64 154 // CHECK: [[DIFF:%.+]] = sub i64 [[E]], [[B]] 155 // CHECK: [[SZ:%.+]] = sdiv exact i64 [[DIFF]], ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64) 156 // CHECK: [[BPTR0:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 0 157 // CHECK: store ptr [[THIS]], ptr [[BPTR0]], 158 // CHECK: [[PTR0:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 0 159 // CHECK: store ptr [[A_ADDR]], ptr [[PTR0]], 160 // CHECK: [[SIZE0:%.+]] = getelementptr inbounds [6 x i64], ptr [[SIZES]], i32 0, i32 0 161 // CHECK: store i64 [[SZ]], ptr [[SIZE0]], 162 // CHECK: [[BPTR1:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 1 163 // CHECK: store ptr [[THIS]], ptr [[BPTR1]] 164 // CHECK: [[PTR1:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 1 165 // CHECK: store ptr [[A_ADDR]], ptr [[PTR1]], 166 // CHECK: [[BPTR2:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 2 167 // CHECK: store ptr [[PTR_ADDR]], ptr [[BPTR2]], 168 // CHECK: [[PTR2:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 2 169 // CHECK: store ptr [[ARR_IDX]], ptr [[PTR2]], 170 // CHECK: [[BPTR3:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 3 171 // CHECK: store ptr [[THIS]], ptr [[BPTR3]] 172 // CHECK: [[PTR3:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 3 173 // CHECK: store ptr [[REF_PTR]], ptr [[PTR3]], 174 // CHECK: [[BPTR4:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 4 175 // CHECK: store ptr [[P3]], ptr [[BPTR4]], 176 // CHECK: [[PTR4:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 4 177 // CHECK: store ptr [[ARR_IDX5]], ptr [[PTR4]] 178 179 // CHECK: [[BPTR5:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 5 180 // CHECK: store ptr [[THIS]], ptr [[BPTR5]], align 8 181 // CHECK: [[PTR5:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 5 182 // CHECK: store ptr [[ARR_IDX6]], ptr [[PTR5]], align 8 183 // CHECK: [[SIZE1:%.+]] = getelementptr inbounds [6 x i64], ptr [[SIZES]], i32 0, i32 5 184 // CHECK: store i64 [[P4]], ptr [[SIZE1]], align 8 185 // CHECK: [[BPTR:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 0 186 // CHECK: [[PTR:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 0 187 // CHECK: [[SIZE:%.+]] = getelementptr inbounds [6 x i64], ptr [[SIZES]], i32 0, i32 0 188 // CHECK: call void @__tgt_target_data_begin_mapper(ptr @{{.+}}, i64 -1, i32 6, ptr [[BPTR]], ptr [[PTR]], ptr [[SIZE]], ptr [[MAPTYPES2]], ptr null, ptr null) 189 // CHECK: [[A_ADDR:%.+]] = load ptr, ptr [[BPTR1]], 190 // CHECK: store ptr [[A_ADDR]], ptr [[A_REF:%.+]], 191 // CHECK: [[PTR_ADDR:%.+]] = load ptr, ptr [[BPTR2]], 192 // CHECK: store ptr [[PTR_ADDR]], ptr [[PTR_REF:%.+]], 193 // CHECK: [[REF_PTR:%.+]] = load ptr, ptr [[BPTR3]], 194 // CHECK: store ptr [[REF_PTR]], ptr [[REF_REF:%.+]], 195 // CHECK: [[PTR_ADDR:%.+]] = load ptr, ptr [[BPTR2]], 196 // CHECK: store ptr [[PTR_ADDR]], ptr [[PTR_REF2:%.+]], 197 // CHECK: [[ARR_ADDR:%.+]] = load ptr, ptr [[BPTR5]], 198 // CHECK: store ptr [[ARR_ADDR]], ptr [[ARR_REF:%.+]], 199 // CHECK: [[A_ADDR:%.+]] = load ptr, ptr [[A_REF]], 200 // CHECK: [[A:%.+]] = load i32, ptr [[A_ADDR]], 201 // CHECK: [[INC:%.+]] = add nsw i32 [[A]], 1 202 // CHECK: store i32 [[INC]], ptr [[A_ADDR]], 203 // CHECK: [[PTR_PTR:%.+]] = load ptr, ptr [[PTR_REF2]], 204 // CHECK: [[PTR:%.+]] = load ptr, ptr [[PTR_PTR]], 205 // CHECK: [[VAL:%.+]] = load i32, ptr [[PTR]], 206 // CHECK: [[INC:%.+]] = add nsw i32 [[VAL]], 1 207 // CHECK: store i32 [[INC]], ptr [[PTR]], 208 // CHECK: [[REF_PTR:%.+]] = load ptr, ptr [[REF_REF]], 209 // CHECK: [[VAL:%.+]] = load i32, ptr [[REF_PTR]], 210 // CHECK: [[INC:%.+]] = add nsw i32 [[VAL]], 1 211 // CHECK: store i32 [[INC]], ptr [[REF_PTR]], 212 // CHECK: [[ARR_ADDR:%.+]] = load ptr, ptr [[ARR_REF]], 213 // CHECK: [[ARR0_ADDR:%.+]] = getelementptr inbounds [4 x i32], ptr [[ARR_ADDR]], i64 0, i64 0 214 // CHECK: [[VAL:%.+]] = load i32, ptr [[ARR0_ADDR]], 215 // CHECK: [[INC:%.+]] = add nsw i32 [[VAL]], 1 216 // CHECK: store i32 [[INC]], ptr [[ARR0_ADDR]], 217 // CHECK: [[BPTR:%.+]] = getelementptr inbounds [6 x ptr], ptr [[BPTRS]], i32 0, i32 0 218 // CHECK: [[PTR:%.+]] = getelementptr inbounds [6 x ptr], ptr [[PTRS]], i32 0, i32 0 219 // CHECK: [[SIZE:%.+]] = getelementptr inbounds [6 x i64], ptr [[SIZES]], i32 0, i32 0 220 // CHECK: call void @__tgt_target_data_end_mapper(ptr @{{.+}}, i64 -1, i32 6, ptr [[BPTR]], ptr [[PTR]], ptr [[SIZE]], ptr [[MAPTYPES2]], ptr null, ptr null) 221 222 #endif 223