[PATCH] D79972: [OpenMP5.0] map item can be non-contiguous for target update

[Chi Chun Chen via Phabricator via cfe-commits]

cchen marked 2 inline comments as done.
cchen added inline comments.


================
Comment at: clang/lib/CodeGen/CGOpenMPRuntime.cpp:7834
+            llvm::APInt Size = CAT->getSize();
+            SizeV = llvm::ConstantInt::get(CGF.SizeTy, Size);
+          } else if (VAT) {
----------------
ABataev wrote:
> cchen wrote:
> > ABataev wrote:
> > > cchen wrote:
> > > > ABataev wrote:
> > > > > Create directly as of `CGF.Int64Ty` type.
> > > > Doing this I'll get assertion error in this exact line if on a 32-bits target.
> > > Hmm, why, can you investigate?
> > My comment was not accurate, I've updated it. What I want to convey is that we can only have `CAT, VAT, or pointer` here, since analysis in Sema has a restriction for it. (SemaOpenMP line 16623)
> It does not relate to the comments thread but I got it. Anyway, try to investigate why the compiler crashes if you try to cr4eate a constant ща ]СПАюШте64Ен] directly.
I'll investigate it, thanks.


================
Comment at: clang/lib/CodeGen/CGOpenMPRuntime.cpp:7872-7873
+          } else {
+            Offset = CGF.Builder.CreateIntCast(CGF.EmitScalarExpr(OffsetExpr),
+                                               CGF.Int64Ty,
+                                               /*isSigned=*/false);
----------------
ABataev wrote:
> cchen wrote:
> > ABataev wrote:
> > > cchen wrote:
> > > > ABataev wrote:
> > > > > Do you really to pass real offsets here? Can we use pointers instead?
> > > > Do you mean I should set the type of Offset to Expr*?
> > > Currently, you're passing offsets to the runtime. Can we pass pointers instead? I mean, for `a[b]` you pass `b` to the runtime, can we pass `&a[b]` instead?
> > Yes, I'm fine either passing index or passing address, though I'm curious why you're recommending passing address.
> It is going to simplify the codegen. Currently, to get the offset, you need to dig through all the elements of the array section. If, instead, you use the pointers, you would not need to do this and you can rely on something like `CGF.EmitArraySectionLValue()`. At least, I hope so.
After discussed with my colleagues, I think passing relative offset makes more sense.

For a 1-dim array, storing the offset as a pointer could work, but it seems strange to me to store as a pointer when there are 2+ dimensions with multiple disjoint chunks of memory because the pointer can only point to the offset for the first chunk. That is, a pointer would refer to an absolute location in a single chunk, whereas the offset is relative to the start of any chunk.

For example:

int a[4][4];
#pragma omp target update to(a[1:2][1:2])

This is two disjoint chunks of memory:

XXXX
XOOX
XOOX
XXXX

The offset for the outer dimension could be store as a pointer, since there is only one instance of that dimension:

Dim1: Offset=&a[1]

But, the inner dimension is "instantiated" twice, once for each element in the outer dimension. So, there are really two absolute pointers, depending on which instance (element in the outer dimension) you're talking about:

Dim2: Offset=&a[1][1]
Dim2: Offset=&a[2][1]

We could set the policy that the absolute offset would always be expressed as the offset in the first instance, but then wouldn't we need to refer to that location when computing the offset for all of the other instances? That seems unintuitive to me, and potentially complicates the implementation. The relative offset makes a lot more senes to me - for a starting point, what relative offset is needed for each dimension. The starting point for the outermost dimension does require the base address, but all inner dimensions have a variable starting pointer based on which element in the outer dimensions you're currently looking at.


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D79972