[PATCH] D86119: [OPENMP50]Allow overlapping mapping in target constrcuts.

[Abhinav Gaba via Phabricator via cfe-commits]

abhinavgaba added a comment.

Thanks for the changes, Alexey! I tried the patch locally, and it looks stable. It handled several tests I tried, including the following case involving array section on a pointer to pointer base, and nested mappers with `PTR_AND_OBJ` maps successfully:

  #include <stdio.h>
  
  typedef struct { int a; double *b; } C;
  #pragma omp declare mapper(id1: C s) map(to:s.a) map(from:s.b[0:2])
  
  typedef struct { int e; C f; int h; short *g; } D;
  #pragma omp declare mapper(default: D r) map(from:r.e) map(mapper(id1), tofrom:r.f) map(tofrom: r.g[0:r.h])
  
  int main() {
    constexpr int N = 10;
    D s;
    s.e = 111;
    s.f.a = 222;
    double x[2]; x[1] = 20;
    short y[N]; y[1] = 30;
    s.f.b = &x[0];
    s.g = &y[0];
    s.h = N;
  
    D* sp = &s;
    D** spp = &sp;
  
    printf("%d %d %lf %p %d %p\n", spp[0][0].e, spp[0][0].f.a, spp[0][0].f.b[1], spp[0][0].f.b, spp[0][0].g[1], spp[0][0].g);
    // Expected: 111 222 20.0 <host_addr1> 30 <host_addr2>
  
    #pragma omp target map(tofrom:spp[0][0])
    {
      printf("%d %d %lf %p %d %p\n", spp[0][0].e, spp[0][0].f.a, spp[0][0].f.b[1], spp[0][0].f.b, spp[0][0].g[1], spp[0][0].g);
      // Expected: <not 111> 222 <not 20.0> <dev_addr1> 30 <dev_addr2>
      spp[0][0].e = 333;
      spp[0][0].f.a = 444;
      spp[0][0].f.b[1] = 40;
      spp[0][0].g[1] = 50;
    }
    printf("%d %d %lf %p %d %p\n", spp[0][0].e, spp[0][0].f.a, spp[0][0].f.b[1], spp[0][0].f.b, spp[0][0].g[1], spp[0][0].g);
    // Expected: 333 222 40.0 <host_addr1> 50 <host_addr2>
  }



================
Comment at: clang/lib/CodeGen/CGOpenMPRuntime.cpp:8483-8484
+      for (const auto &M : Data.second) {
+        for (const MapInfo &L : M) {
+          assert(!L.Components.empty() &&
+                 "Not expecting declaration with no component lists.");
----------------
Tabs should probably be spaces. Same for a few other places in the changeset.


================
Comment at: clang/lib/CodeGen/CGOpenMPRuntime.cpp:9726
     llvm::Value *OriMapType = MapperCGF.Builder.getInt64(Info.Types[I]);
-    llvm::Value *Member = MapperCGF.Builder.CreateAnd(
-        OriMapType,
-        MapperCGF.Builder.getInt64(MappableExprsHandler::OMP_MAP_MEMBER_OF));
-    llvm::BasicBlock *MemberCombineBB =
-        MapperCGF.createBasicBlock("omp.member.combine");
-    llvm::BasicBlock *TypeBB = MapperCGF.createBasicBlock("omp.type");
-    llvm::Value *IsMember = MapperCGF.Builder.CreateIsNull(Member);
-    MapperCGF.Builder.CreateCondBr(IsMember, TypeBB, MemberCombineBB);
-    // Add the number of pre-existing components to the MEMBER_OF field if it
-    // is valid.
-    MapperCGF.EmitBlock(MemberCombineBB);
-    llvm::Value *CombinedMember =
-        MapperCGF.Builder.CreateNUWAdd(OriMapType, ShiftedPreviousSize);
+    // llvm::Value *Member = MapperCGF.Builder.CreateAnd(
+    //     OriMapType,
----------------
Commented-out code intentionally left in?


================
Comment at: clang/lib/CodeGen/CGOpenMPRuntime.cpp:9568
 ///   // Allocate space for an array section first.
-///   if (size > 1 && !maptype.IsDelete)
+///   if ((size > 1 || base != begin) && !maptype.IsDeleten)
 ///     __tgt_push_mapper_component(rt_mapper_handle, base, begin,
----------------
IsDelete


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D86119