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

Alexey Bataev via Phabricator via Openmp-commits openmp-commits at lists.llvm.org
Tue Feb 16 05:37:11 PST 2021


ABataev marked 2 inline comments as done.
ABataev added a comment.

In D86119#2562187 <https://reviews.llvm.org/D86119#2562187>, @grokos wrote:

> In D86119#2561163 <https://reviews.llvm.org/D86119#2561163>, @abhinavgaba wrote:
>
>> 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>
>>   }
>
> @ABataev This is a nice complex example, I think it's worth including it in the runtime tests (under libomptarget).
>
> @abhinavgaba Thanks for providing it!

Ok, will add it as a part of the patch



================
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.");
----------------
abhinavgaba wrote:
> Tabs should probably be spaces. Same for a few other places in the changeset.
These are not tabs. Looks like this is how Phabricators shows some format changes.


================
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,
----------------
abhinavgaba wrote:
> Commented-out code intentionally left in?
Yeah, forgot to remove it, thanks.


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D86119



More information about the Openmp-commits mailing list