Author: Alexey Bataev Date: 2020-11-20T11:34:14-08:00 New Revision: c964f308141578f24932c68a03af5fae7f876011
URL: https://github.com/llvm/llvm-project/commit/c964f308141578f24932c68a03af5fae7f876011 DIFF: https://github.com/llvm/llvm-project/commit/c964f308141578f24932c68a03af5fae7f876011.diff LOG: [OPENMP]Use the real pointer value as base, not indexed value. After fix for PR48174 the base pointer for pointer-based array-sections/array-subscripts will be emitted as `&ptr[idx]`, but actually it should be just `ptr`, i.e. the address stored in the ponter to point correctly to the beginning of the array. Currently it may lead to a crash in the runtime. Differential Revision: https://reviews.llvm.org/D91805 Added: Modified: clang/lib/CodeGen/CGOpenMPRuntime.cpp clang/test/OpenMP/target_data_map_pointer_array_subscript_codegen.cpp Removed: ################################################################################ diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp index 84733369a4a1..c15f6350b95e 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -7904,8 +7904,11 @@ class MappableExprsHandler { IsCaptureFirstInfo = false; FirstPointerInComplexData = false; } else if (FirstPointerInComplexData) { - BP = CGF.EmitOMPSharedLValue(I->getAssociatedExpression()) - .getAddress(CGF); + QualType Ty = Components.rbegin() + ->getAssociatedDeclaration() + ->getType() + .getNonReferenceType(); + BP = CGF.EmitLoadOfPointer(BP, Ty->castAs<PointerType>()); FirstPointerInComplexData = false; } } diff --git a/clang/test/OpenMP/target_data_map_pointer_array_subscript_codegen.cpp b/clang/test/OpenMP/target_data_map_pointer_array_subscript_codegen.cpp index 14960191af03..77d416a685aa 100644 --- a/clang/test/OpenMP/target_data_map_pointer_array_subscript_codegen.cpp +++ b/clang/test/OpenMP/target_data_map_pointer_array_subscript_codegen.cpp @@ -38,10 +38,17 @@ MyObject *objects; // CHECK-DAG: [[MAPS1:@.+]] = private unnamed_addr constant [2 x i64] [i64 32, i64 281474976710673] // CHECK: @main int main(void) { +// CHECK: [[BPTR0:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* %{{.+}}, i32 0, i32 0 +// CHECK: [[BPTRC0:%.+]] = bitcast i8** [[BPTR0]] to %struct.MyObject*** +// CHECK: store %struct.MyObject** @objects, %struct.MyObject*** [[BPTRC0]], // CHECK: call void @__tgt_target_data_begin_mapper(%struct.ident_t* @{{.+}}, i64 -1, i32 1, i8** %{{.+}}, i8** %{{.+}}, i64* getelementptr inbounds ([1 x i64], [1 x i64]* [[SIZES0]], i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* [[MAPS0]], i32 0, i32 0), i8** null, i8** null) -#pragma omp target enter data map(to : objects [0:1]) +#pragma omp target enter data map(to : objects [1:1]) +// CHECK: [[OBJ:%.+]] = load %struct.MyObject*, %struct.MyObject** @objects, +// CHECK: [[BPTR0:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 0 +// CHECK: [[BPTRC0:%.+]] = bitcast i8** [[BPTR0]] to %struct.MyObject** +// CHECK: store %struct.MyObject* [[OBJ]], %struct.MyObject** [[BPTRC0]], // CHECK: call void @__tgt_target_data_begin_mapper(%struct.ident_t* @{{.+}}, i64 -1, i32 2, i8** %{{.+}}, i8** %{{.+}}, i64* %{{.+}}, i64* getelementptr inbounds ([2 x i64], [2 x i64]* [[MAPS1]], i32 0, i32 0), i8** null, i8** null) -#pragma omp target enter data map(to : objects[0].arr [0:1]) +#pragma omp target enter data map(to : objects[1].arr [0:1]) return 0; } _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
