[clang] c964f30 - [OPENMP]Use the real pointer value as base, not indexed value.

[Alexey Bataev via cfe-commits]

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;
 }