Page MenuHomePhabricator

[OpenMP] Fix `omp target update` for array extension
ClosedPublic

Authored by jdenny on Aug 4 2020, 2:06 PM.

Details

Summary

OpenMP TR8 sec. 2.15.6 "target update Construct", p. 183, L3-4 states:

If the corresponding list item is not present in the device data
environment and there is no present modifier in the clause, then no
assignment occurs to or from the original list item.

L10-11 states:

If a present modifier appears in the clause and the corresponding
list item is not present in the device data environment then an
error occurs and the program termintates.

(OpenMP 5.0 also has the first passage but without mention of the
present modifier of course.)

In both passages, I assume "is not present" includes the case of
partially but not entirely present. However, without this patch, the
target update directive misbehaves in this case both with and without
the present modifier. For example:

#pragma omp target enter data map(to:arr[0:3])
#pragma omp target update to(arr[0:5]) // might fail on data transfer
#pragma omp target update to(present:arr[0:5]) // might fail on data transfer

The problem is that DeviceTy::getTgtPtrBegin does not return a null
pointer in that case, so target_data_update sees the data as fully
present, and the data transfer then might fail depending on the target
device. However, without the present modifier, there should never be
a failure. Moreover, with the present modifier, there should always
be a failure, and the diagnostic should mention the present modifier.

This patch fixes DeviceTy::getTgtPtrBegin to return null when
target_data_update is the caller. I'm wondering if it should do the
same for more callers.

Diff Detail

Event Timeline

jdenny requested review of this revision.Aug 4 2020, 2:06 PM
jdenny created this revision.
jdenny set the repository for this revision to rG LLVM Github Monorepo.
Herald added a project: Restricted Project. · View Herald TranscriptAug 4 2020, 2:07 PM
grokos accepted this revision.Aug 4 2020, 3:06 PM

I suppose the same must be applied to all cases where present is used. Without a modification like this, present just confirms that the starting address of the object is already mapped, although a size may be specified in the map clause.

#pragma omp target map(present: p[0:10]) // present will succeed even if only p[0] is mapped and p[1]-p[9] are not

So the question here is: should present apply to the size as well as the begin address? I would say yes.

This revision is now accepted and ready to land.Aug 4 2020, 3:06 PM
jdenny added a comment.Aug 4 2020, 3:49 PM

Thanks for the quick review.

I suppose the same must be applied to all cases where present is used.

Agreed. Fortunately, when present is instead a map type modifier, DeviceTy::getOrAllocTgtPtr is called instead of DeviceTy::getTgtPtrBegin and, as far as I know, it already returns a null pointer for this case. However, I see now that we get the general failure that happens regardless of the present modifier, and we don't see the present-specific diagnostic. I suppose that ought to change.

Without a modification like this, present just confirms that the starting address of the object is already mapped, although a size may be specified in the map clause.

It doesn't even confirm the start address exactly. That is, in the extends-before case, the starting address might be before the array section that was already mapped. Without this patch, the present motion modifier permits that case too.

#pragma omp target map(present: p[0:10]) // present will succeed even if only p[0] is mapped and p[1]-p[9] are not

So the question here is: should present apply to the size as well as the begin address? I would say yes.

I agree, and it seems to already except for the diagnostic.

Maybe change the commit title to OpenMP to confuse less people.

jdenny retitled this revision from [OpenACC] Fix `omp target update` for array extension to [OpenMP] Fix `omp target update` for array extension.Aug 4 2020, 4:56 PM
jdenny added a comment.Aug 4 2020, 4:57 PM

Maybe change the commit title to OpenMP to confuse less people.

Thanks for catching that. Done.

jdenny added a comment.Aug 4 2020, 5:40 PM

This patch fixes DeviceTy::getTgtPtrBegin to return null when
target_data_update is the caller. I'm wondering if it should do the
same for more callers.

Here's an example of what I'm thinking here:

#pragma omp target data map(tofrom:arr[0:20])
{
  #pragma omp target exit data map(delete:arr[0:20])
  #pragma omp target enter data map(alloc:arr[10:2])
}

The transfer at the end of the target data region currently fails. According to my read of the spec, the transfer shouldn't even be attempted because arr[0:20] isn't (fully) present. If DeviceTy::getTgtPtrBegin were to return null for that case too, then this problem would go away.

This revision was automatically updated to reflect the committed changes.