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
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
