[PATCH] D86119: [OPENMP50]Allow overlapping mapping in target constrcuts.
Alexey Bataev via Phabricator via cfe-commits
cfe-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 cfe-commits
mailing list