1 // RUN: %clang_cc1 -no-opaque-pointers -DCK1 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s 2 // RUN: %clang_cc1 -no-opaque-pointers -DCK1 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 3 // RUN: %clang_cc1 -no-opaque-pointers -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s 4 5 // RUN: %clang_cc1 -no-opaque-pointers -DCK1 -verify -fopenmp-simd -fopenmp-version=50 -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 -no-opaque-pointers -DCK1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 7 // RUN: %clang_cc1 -no-opaque-pointers -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s 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 [5 x i64] zeroinitializer 15 // 64 = 0x40 = OMP_MAP_RETURN_PARAM 16 // CHECK-DAG: [[MAPTYPES1:@.+]] = private unnamed_addr constant [5 x i64] [i64 64, i64 64, i64 64, i64 64, i64 64] 17 // CHECK-DAG: [[SIZES2:@.+]] = private unnamed_addr constant [5 x i64] zeroinitializer 18 // 0 = OMP_MAP_NONE 19 // 281474976710720 = 0x1000000000040 = OMP_MAP_MEMBER_OF | OMP_MAP_RETURN_PARAM 20 // CHECK-DAG: [[MAPTYPES2:@.+]] = private unnamed_addr constant [5 x i64] [i64 0, i64 281474976710720, i64 281474976710720, i64 281474976710720, i64 281474976710720] 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 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 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 float*, 49 // CHECK: [[REF_ADDR:%.+]] = alloca float*, 50 // CHECK: [[ARR_ADDR:%.+]] = alloca [4 x float], 51 // CHECK: [[BPTRS:%.+]] = alloca [5 x i8*], 52 // CHECK: [[PTRS:%.+]] = alloca [5 x i8*], 53 // CHECK: [[VLA_ADDR:%.+]] = alloca float, i64 %{{.+}}, 54 // CHECK: [[PTR:%.+]] = load float*, float** [[PTR_ADDR]], 55 // CHECK: [[REF:%.+]] = load float*, float** [[REF_ADDR]], 56 // CHECK: [[ARR:%.+]] = getelementptr inbounds [4 x float], [4 x float]* [[ARR_ADDR]], i64 0, i64 0 57 // CHECK: [[BPTR0:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BPTRS]], i32 0, i32 0 58 // CHECK: [[BPTR0_A_ADDR:%.+]] = bitcast i8** [[BPTR0]] to float** 59 // CHECK: store float* [[A_ADDR]], float** [[BPTR0_A_ADDR]], 60 // CHECK: [[PTR0:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[PTRS]], i32 0, i32 0 61 // CHECK: [[PTR0_A_ADDR:%.+]] = bitcast i8** [[PTR0]] to float** 62 // CHECK: store float* [[A_ADDR]], float** [[PTR0_A_ADDR]], 63 // CHECK: [[BPTR1:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BPTRS]], i32 0, i32 1 64 // CHECK: [[BPTR1_PTR_ADDR:%.+]] = bitcast i8** [[BPTR1]] to float** 65 // CHECK: store float* [[PTR]], float** [[BPTR1_PTR_ADDR]], 66 // CHECK: [[PTR1:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[PTRS]], i32 0, i32 1 67 // CHECK: [[PTR1_PTR_ADDR:%.+]] = bitcast i8** [[PTR1]] to float** 68 // CHECK: store float* [[PTR]], float** [[PTR1_PTR_ADDR]], 69 // CHECK: [[BPTR2:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BPTRS]], i32 0, i32 2 70 // CHECK: [[BPTR2_REF_ADDR:%.+]] = bitcast i8** [[BPTR2]] to float** 71 // CHECK: store float* [[REF]], float** [[BPTR2_REF_ADDR]], 72 // CHECK: [[PTR2:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[PTRS]], i32 0, i32 2 73 // CHECK: [[PTR2_REF_ADDR:%.+]] = bitcast i8** [[PTR2]] to float** 74 // CHECK: store float* [[REF]], float** [[PTR2_REF_ADDR]], 75 // CHECK: [[BPTR3:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BPTRS]], i32 0, i32 3 76 // CHECK: [[BPTR3_ARR_ADDR:%.+]] = bitcast i8** [[BPTR3]] to float** 77 // CHECK: store float* [[ARR]], float** [[BPTR3_ARR_ADDR]], 78 // CHECK: [[PTR3:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[PTRS]], i32 0, i32 3 79 // CHECK: [[PTR3_ARR_ADDR:%.+]] = bitcast i8** [[PTR3]] to float** 80 // CHECK: store float* [[ARR]], float** [[PTR3_ARR_ADDR]], 81 // CHECK: [[BPTR4:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BPTRS]], i32 0, i32 4 82 // CHECK: [[BPTR4_VLA_ADDR:%.+]] = bitcast i8** [[BPTR4]] to float** 83 // CHECK: store float* [[VLA_ADDR]], float** [[BPTR4_VLA_ADDR]], 84 // CHECK: [[PTR4:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[PTRS]], i32 0, i32 4 85 // CHECK: [[PTR4_VLA_ADDR:%.+]] = bitcast i8** [[PTR4]] to float** 86 // CHECK: store float* [[VLA_ADDR]], float** [[PTR4_VLA_ADDR]], 87 // CHECK: [[BPTR:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BPTRS]], i32 0, i32 0 88 // CHECK: [[PTR:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[PTRS]], i32 0, i32 0 89 // CHECK: call void @__tgt_target_data_begin_mapper(%struct.ident_t* @{{.+}}, i64 -1, i32 5, i8** [[BPTR]], i8** [[PTR]], i64* getelementptr inbounds ([5 x i64], [5 x i64]* [[SIZES1]], i32 0, i32 0), i64* getelementptr inbounds ([5 x i64], [5 x i64]* [[MAPTYPES1]], i32 0, i32 0), i8** null, i8** null) 90 // CHECK: [[A_REF:%.+]] = load float*, float** [[BPTR0_A_ADDR]], 91 // CHECK: [[REF_REF:%.+]] = load float*, float** [[BPTR2_REF_ADDR]], 92 // CHECK: store float* [[REF_REF]], float** [[TMP_REF_ADDR:%.+]], 93 // CHECK: [[BPTR3_ARR_ADDR_CAST:%.+]] = bitcast float** [[BPTR3_ARR_ADDR]] to [4 x float]** 94 // CHECK: [[ARR_REF:%.+]] = load [4 x float]*, [4 x float]** [[BPTR3_ARR_ADDR_CAST]], 95 // CHECK: [[VLA_REF:%.+]] = load float*, float** [[BPTR4_VLA_ADDR]], 96 // CHECK: [[A:%.+]] = load float, float* [[A_REF]], 97 // CHECK: [[INC:%.+]] = fadd float [[A]], 1.000000e+00 98 // CHECK: store float [[INC]], float* [[A_REF]], 99 // CHECK: [[PTR_ADDR:%.+]] = load float*, float** [[BPTR1_PTR_ADDR]], 100 // CHECK: [[VAL:%.+]] = load float, float* [[PTR_ADDR]], 101 // CHECK: [[INC:%.+]] = fadd float [[VAL]], 1.000000e+00 102 // CHECK: store float [[INC]], float* [[PTR_ADDR]], 103 // CHECK: [[REF_ADDR:%.+]] = load float*, float** [[TMP_REF_ADDR]], 104 // CHECK: [[REF:%.+]] = load float, float* [[REF_ADDR]], 105 // CHECK: [[INC:%.+]] = fadd float [[REF]], 1.000000e+00 106 // CHECK: store float [[INC]], float* [[REF_ADDR]], 107 // CHECK: [[ARR0_ADDR:%.+]] = getelementptr inbounds [4 x float], [4 x float]* [[ARR_REF]], i64 0, i64 0 108 // CHECK: [[ARR0:%.+]] = load float, float* [[ARR0_ADDR]], 109 // CHECK: [[INC:%.+]] = fadd float [[ARR0]], 1.000000e+00 110 // CHECK: store float [[INC]], float* [[ARR0_ADDR]], 111 // CHECK: [[VLA0_ADDR:%.+]] = getelementptr inbounds float, float* [[VLA_REF]], i64 0 112 // CHECK: [[VLA0:%.+]] = load float, float* [[VLA0_ADDR]], 113 // CHECK: [[INC:%.+]] = fadd float [[VLA0]], 1.000000e+00 114 // CHECK: store float [[INC]], float* [[VLA0_ADDR]], 115 // CHECK: [[BPTR:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BPTRS]], i32 0, i32 0 116 // CHECK: [[PTR:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[PTRS]], i32 0, i32 0 117 // CHECK: call void @__tgt_target_data_end_mapper(%struct.ident_t* @{{.+}}, i64 -1, i32 5, i8** [[BPTR]], i8** [[PTR]], i64* getelementptr inbounds ([5 x i64], [5 x i64]* [[SIZES1]], i32 0, i32 0), i64* getelementptr inbounds ([5 x i64], [5 x i64]* [[MAPTYPES1]], i32 0, i32 0), i8** null, i8** null) 118 119 // CHECK: foo 120 // %this.addr = alloca %struct.S*, align 8 121 // CHECK: [[BPTRS:%.+]] = alloca [5 x i8*], 122 // CHECK: [[PTRS:%.+]] = alloca [5 x i8*], 123 // CHECK: [[SIZES:%.+]] = alloca [5 x i64], 124 // %tmp = alloca i32*, align 8 125 // %tmp6 = alloca i32**, align 8 126 // %tmp7 = alloca i32*, align 8 127 // %tmp8 = alloca i32**, align 8 128 // %tmp9 = alloca [4 x i32]*, align 8 129 // store %struct.S* %this, %struct.S** %this.addr, align 8 130 // %this1 = load %struct.S*, %struct.S** %this.addr, align 8 131 // CHECK: [[A_ADDR:%.+]] = getelementptr inbounds %struct.S, %struct.S* [[THIS:%.+]], i32 0, i32 0 132 // %ptr = getelementptr inbounds %struct.S, %struct.S* %this1, i32 0, i32 1 133 // %ref = getelementptr inbounds %struct.S, %struct.S* %this1, i32 0, i32 2 134 // %0 = load i32*, i32** %ref, align 8 135 // CHECK: [[ARR_ADDR:%.+]] = getelementptr inbounds %struct.S, %struct.S* [[THIS]], i32 0, i32 3 136 // CHECK: [[A_ADDR2:%.+]] = getelementptr inbounds %struct.S, %struct.S* [[THIS]], i32 0, i32 0 137 // CHECK: [[PTR_ADDR:%.+]] = getelementptr inbounds %struct.S, %struct.S* [[THIS]], i32 0, i32 1 138 // CHECK: [[REF_REF:%.+]] = getelementptr inbounds %struct.S, %struct.S* [[THIS]], i32 0, i32 2 139 // CHECK: [[REF_PTR:%.+]] = load i32*, i32** [[REF_REF]], 140 // CHECK: [[ARR_ADDR2:%.+]] = getelementptr inbounds %struct.S, %struct.S* [[THIS]], i32 0, i32 3 141 // CHECK: [[ARR_END:%.+]] = getelementptr [4 x i32], [4 x i32]* [[ARR_ADDR]], i32 1 142 // CHECK: [[BEGIN:%.+]] = bitcast i32* [[A_ADDR]] to i8* 143 // CHECK: [[END:%.+]] = bitcast [4 x i32]* [[ARR_END]] to i8* 144 // CHECK: [[E:%.+]] = ptrtoint i8* [[END]] to i64 145 // CHECK: [[B:%.+]] = ptrtoint i8* [[BEGIN]] to i64 146 // CHECK: [[DIFF:%.+]] = sub i64 [[E]], [[B]] 147 // CHECK: [[SZ:%.+]] = sdiv exact i64 [[DIFF]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64) 148 // CHECK: [[BPTR0:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BPTRS]], i32 0, i32 0 149 // CHECK: [[BPTR0_S:%.+]] = bitcast i8** [[BPTR0]] to %struct.S** 150 // CHECK: store %struct.S* [[THIS]], %struct.S** [[BPTR0_S]], 151 // CHECK: [[PTR0:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[PTRS]], i32 0, i32 0 152 // CHECK: [[PTR0_BEGIN:%.+]] = bitcast i8** [[PTR0]] to i32** 153 // CHECK: store i32* [[A_ADDR]], i32** [[PTR0_BEGIN]], 154 // CHECK: [[SIZE0:%.+]] = getelementptr inbounds [5 x i64], [5 x i64]* [[SIZES]], i32 0, i32 0 155 // CHECK: store i64 [[SZ]], i64* [[SIZE0]], 156 // CHECK: [[BPTR1:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BPTRS]], i32 0, i32 1 157 // CHECK: [[BPTR1_A_ADDR:%.+]] = bitcast i8** [[BPTR1]] to i32** 158 // CHECK: store i32* [[A_ADDR2]], i32** [[BPTR1_A_ADDR]], 159 // CHECK: [[PTR1:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[PTRS]], i32 0, i32 1 160 // CHECK: [[PTR1_A_ADDR:%.+]] = bitcast i8** [[PTR1]] to i32** 161 // CHECK: store i32* [[A_ADDR2]], i32** [[PTR1_A_ADDR]], 162 // CHECK: [[BPTR2:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BPTRS]], i32 0, i32 2 163 // CHECK: [[BPTR2_PTR_ADDR:%.+]] = bitcast i8** [[BPTR2]] to i32*** 164 // CHECK: store i32** [[PTR_ADDR]], i32*** [[BPTR2_PTR_ADDR]], 165 // CHECK: [[PTR2:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[PTRS]], i32 0, i32 2 166 // CHECK: [[PTR2_PTR_ADDR:%.+]] = bitcast i8** [[PTR2]] to i32*** 167 // CHECK: store i32** [[PTR_ADDR]], i32*** [[PTR2_PTR_ADDR]], 168 // CHECK: [[BPTR3:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BPTRS]], i32 0, i32 3 169 // CHECK: [[BPTR3_REF_PTR:%.+]] = bitcast i8** [[BPTR3]] to i32** 170 // CHECK: store i32* [[REF_PTR]], i32** [[BPTR3_REF_PTR]], 171 // CHECK: [[PTR3:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[PTRS]], i32 0, i32 3 172 // CHECK: [[PTR3_REF_PTR:%.+]] = bitcast i8** [[PTR3]] to i32** 173 // CHECK: store i32* [[REF_PTR]], i32** [[PTR3_REF_PTR]], 174 // CHECK: [[BPTR4:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BPTRS]], i32 0, i32 4 175 // CHECK: [[BPTR4_ARR_ADDR:%.+]] = bitcast i8** [[BPTR4]] to [4 x i32]** 176 // CHECK: store [4 x i32]* [[ARR_ADDR2]], [4 x i32]** [[BPTR4_ARR_ADDR]], 177 // CHECK: [[PTR4:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[PTRS]], i32 0, i32 4 178 // CHECK: [[PTR4_ARR_ADDR:%.+]] = bitcast i8** [[PTR4]] to [4 x i32]** 179 // CHECK: store [4 x i32]* [[ARR_ADDR2]], [4 x i32]** [[PTR4_ARR_ADDR]], 180 // CHECK: [[BPTR:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BPTRS]], i32 0, i32 0 181 // CHECK: [[PTR:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[PTRS]], i32 0, i32 0 182 // CHECK: [[SIZE:%.+]] = getelementptr inbounds [5 x i64], [5 x i64]* [[SIZES]], i32 0, i32 0 183 // CHECK: call void @__tgt_target_data_begin_mapper(%struct.ident_t* @{{.+}}, i64 -1, i32 5, i8** [[BPTR]], i8** [[PTR]], i64* [[SIZE]], i64* getelementptr inbounds ([5 x i64], [5 x i64]* [[MAPTYPES2]], i32 0, i32 0), i8** null, i8** null) 184 // CHECK: [[A_ADDR:%.+]] = load i32*, i32** [[BPTR1_A_ADDR]], 185 // CHECK: store i32* [[A_ADDR]], i32** [[A_REF:%.+]], 186 // CHECK: [[PTR_ADDR:%.+]] = load i32**, i32*** [[BPTR2_PTR_ADDR]], 187 // CHECK: store i32** [[PTR_ADDR]], i32*** [[PTR_REF:%.+]], 188 // CHECK: [[REF_PTR:%.+]] = load i32*, i32** [[BPTR3_REF_PTR]], 189 // CHECK: store i32* [[REF_PTR]], i32** [[REF_REF:%.+]], 190 // CHECK: [[PTR_ADDR:%.+]] = load i32**, i32*** [[BPTR2_PTR_ADDR]], 191 // CHECK: store i32** [[PTR_ADDR]], i32*** [[PTR_REF2:%.+]], 192 // CHECK: [[ARR_ADDR:%.+]] = load [4 x i32]*, [4 x i32]** [[BPTR4_ARR_ADDR]], 193 // CHECK: store [4 x i32]* [[ARR_ADDR]], [4 x i32]** [[ARR_REF:%.+]], 194 // CHECK: [[A_ADDR:%.+]] = load i32*, i32** [[A_REF]], 195 // CHECK: [[A:%.+]] = load i32, i32* [[A_ADDR]], 196 // CHECK: [[INC:%.+]] = add nsw i32 [[A]], 1 197 // CHECK: store i32 [[INC]], i32* [[A_ADDR]], 198 // CHECK: [[PTR_PTR:%.+]] = load i32**, i32*** [[PTR_REF2]], 199 // CHECK: [[PTR:%.+]] = load i32*, i32** [[PTR_PTR]], 200 // CHECK: [[VAL:%.+]] = load i32, i32* [[PTR]], 201 // CHECK: [[INC:%.+]] = add nsw i32 [[VAL]], 1 202 // CHECK: store i32 [[INC]], i32* [[PTR]], 203 // CHECK: [[REF_PTR:%.+]] = load i32*, i32** [[REF_REF]], 204 // CHECK: [[VAL:%.+]] = load i32, i32* [[REF_PTR]], 205 // CHECK: [[INC:%.+]] = add nsw i32 [[VAL]], 1 206 // CHECK: store i32 [[INC]], i32* [[REF_PTR]], 207 // CHECK: [[ARR_ADDR:%.+]] = load [4 x i32]*, [4 x i32]** [[ARR_REF]], 208 // CHECK: [[ARR0_ADDR:%.+]] = getelementptr inbounds [4 x i32], [4 x i32]* [[ARR_ADDR]], i64 0, i64 0 209 // CHECK: [[VAL:%.+]] = load i32, i32* [[ARR0_ADDR]], 210 // CHECK: [[INC:%.+]] = add nsw i32 [[VAL]], 1 211 // CHECK: store i32 [[INC]], i32* [[ARR0_ADDR]], 212 // CHECK: [[BPTR:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BPTRS]], i32 0, i32 0 213 // CHECK: [[PTR:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[PTRS]], i32 0, i32 0 214 // CHECK: [[SIZE:%.+]] = getelementptr inbounds [5 x i64], [5 x i64]* [[SIZES]], i32 0, i32 0 215 // CHECK: call void @__tgt_target_data_end_mapper(%struct.ident_t* @{{.+}}, i64 -1, i32 5, i8** [[BPTR]], i8** [[PTR]], i64* [[SIZE]], i64* getelementptr inbounds ([5 x i64], [5 x i64]* [[MAPTYPES2]], i32 0, i32 0), i8** null, i8** null) 216 217 #endif 218