[PATCH] D91805: [OPENMP]Use the real pointer value as base, not indexed value.

[Alexey Bataev via Phabricator via cfe-commits]

This revision was landed with ongoing or failed builds.
This revision was automatically updated to reflect the committed changes.
Closed by commit rGc964f3081415: [OPENMP]Use the real pointer value as base, not indexed value. (authored by ABataev).

Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D91805/new/

https://reviews.llvm.org/D91805

Files:
  clang/lib/CodeGen/CGOpenMPRuntime.cpp
  clang/test/OpenMP/target_data_map_pointer_array_subscript_codegen.cpp


Index: clang/test/OpenMP/target_data_map_pointer_array_subscript_codegen.cpp
===================================================================
--- clang/test/OpenMP/target_data_map_pointer_array_subscript_codegen.cpp
+++ clang/test/OpenMP/target_data_map_pointer_array_subscript_codegen.cpp
@@ -38,10 +38,17 @@
 // 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;
 }
Index: clang/lib/CodeGen/CGOpenMPRuntime.cpp
===================================================================
--- clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -7904,8 +7904,11 @@
         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;
       }
     }


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D91805.306752.patch
Type: text/x-patch
Size: 2570 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20201120/e2014d1c/attachment.bin>