forked from OSchip/llvm-project
[OpenMP] Fix `omp target update` for array extension
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. Reviewed By: grokos, jdoerfert Differential Revision: https://reviews.llvm.org/D85246
This commit is contained in:
Joel E. Denny
parent
03bb545b68
commit
5ab43989c3
|
|
@ -236,14 +236,16 @@ void *DeviceTy::getOrAllocTgtPtr(void *HstPtrBegin, void *HstPtrBase,
|
|||
// Return the target pointer begin (where the data will be moved).
|
||||
// Decrement the reference counter if called from targetDataEnd.
|
||||
void *DeviceTy::getTgtPtrBegin(void *HstPtrBegin, int64_t Size, bool &IsLast,
|
||||
bool UpdateRefCount, bool &IsHostPtr) {
|
||||
bool UpdateRefCount, bool &IsHostPtr,
|
||||
bool MustContain) {
|
||||
void *rc = NULL;
|
||||
IsHostPtr = false;
|
||||
IsLast = false;
|
||||
DataMapMtx.lock();
|
||||
LookupResult lr = lookupMapping(HstPtrBegin, Size);
|
||||
|
||||
if (lr.Flags.IsContained || lr.Flags.ExtendsBefore || lr.Flags.ExtendsAfter) {
|
||||
if (lr.Flags.IsContained ||
|
||||
(!MustContain && (lr.Flags.ExtendsBefore || lr.Flags.ExtendsAfter))) {
|
||||
auto &HT = *lr.Entry;
|
||||
IsLast = HT.getRefCount() == 1;
|
||||
|
||||
|
|
|
|||
|
|
@ -182,7 +182,8 @@ struct DeviceTy {
|
|||
bool HasPresentModifier);
|
||||
void *getTgtPtrBegin(void *HstPtrBegin, int64_t Size);
|
||||
void *getTgtPtrBegin(void *HstPtrBegin, int64_t Size, bool &IsLast,
|
||||
bool UpdateRefCount, bool &IsHostPtr);
|
||||
bool UpdateRefCount, bool &IsHostPtr,
|
||||
bool MustContain = false);
|
||||
int deallocTgtPtr(void *TgtPtrBegin, int64_t Size, bool ForceDelete,
|
||||
bool HasCloseModifier = false);
|
||||
int associatePtr(void *HstPtrBegin, void *TgtPtrBegin, int64_t Size);
|
||||
|
|
|
|||
|
|
@ -670,8 +670,8 @@ int target_data_update(DeviceTy &Device, int32_t arg_num,
|
|||
void *HstPtrBegin = args[i];
|
||||
int64_t MapSize = arg_sizes[i];
|
||||
bool IsLast, IsHostPtr;
|
||||
void *TgtPtrBegin = Device.getTgtPtrBegin(HstPtrBegin, MapSize, IsLast,
|
||||
false, IsHostPtr);
|
||||
void *TgtPtrBegin = Device.getTgtPtrBegin(
|
||||
HstPtrBegin, MapSize, IsLast, false, IsHostPtr, /*MustContain=*/true);
|
||||
if (!TgtPtrBegin) {
|
||||
DP("hst data:" DPxMOD " not found, becomes a noop\n", DPxPTR(HstPtrBegin));
|
||||
if (arg_types[i] & OMP_TGT_MAPTYPE_PRESENT) {
|
||||
|
|
|
|||
|
|
@ -0,0 +1,140 @@
|
|||
// --------------------------------------------------
|
||||
// Check 'to' and extends before
|
||||
// --------------------------------------------------
|
||||
|
||||
// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu \
|
||||
// RUN: -fopenmp-version=51 -DCLAUSE=to -DEXTENDS=BEFORE
|
||||
// RUN: %libomptarget-run-fail-aarch64-unknown-linux-gnu 2>&1 \
|
||||
// RUN: | %fcheck-aarch64-unknown-linux-gnu
|
||||
|
||||
// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu \
|
||||
// RUN: -fopenmp-version=51 -DCLAUSE=to -DEXTENDS=BEFORE
|
||||
// RUN: %libomptarget-run-fail-powerpc64-ibm-linux-gnu 2>&1 \
|
||||
// RUN: | %fcheck-powerpc64-ibm-linux-gnu
|
||||
|
||||
// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu \
|
||||
// RUN: -fopenmp-version=51 -DCLAUSE=to -DEXTENDS=BEFORE
|
||||
// RUN: %libomptarget-run-fail-powerpc64le-ibm-linux-gnu 2>&1 \
|
||||
// RUN: | %fcheck-powerpc64le-ibm-linux-gnu
|
||||
|
||||
// RUN: %libomptarget-compile-x86_64-pc-linux-gnu \
|
||||
// RUN: -fopenmp-version=51 -DCLAUSE=to -DEXTENDS=BEFORE
|
||||
// RUN: %libomptarget-run-fail-x86_64-pc-linux-gnu 2>&1 \
|
||||
// RUN: | %fcheck-x86_64-pc-linux-gnu
|
||||
|
||||
// --------------------------------------------------
|
||||
// Check 'from' and extends before
|
||||
// --------------------------------------------------
|
||||
|
||||
// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu \
|
||||
// RUN: -fopenmp-version=51 -DCLAUSE=from -DEXTENDS=BEFORE
|
||||
// RUN: %libomptarget-run-fail-aarch64-unknown-linux-gnu 2>&1 \
|
||||
// RUN: | %fcheck-aarch64-unknown-linux-gnu
|
||||
|
||||
// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu \
|
||||
// RUN: -fopenmp-version=51 -DCLAUSE=from -DEXTENDS=BEFORE
|
||||
// RUN: %libomptarget-run-fail-powerpc64-ibm-linux-gnu 2>&1 \
|
||||
// RUN: | %fcheck-powerpc64-ibm-linux-gnu
|
||||
|
||||
// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu \
|
||||
// RUN: -fopenmp-version=51 -DCLAUSE=from -DEXTENDS=BEFORE
|
||||
// RUN: %libomptarget-run-fail-powerpc64le-ibm-linux-gnu 2>&1 \
|
||||
// RUN: | %fcheck-powerpc64le-ibm-linux-gnu
|
||||
|
||||
// RUN: %libomptarget-compile-x86_64-pc-linux-gnu \
|
||||
// RUN: -fopenmp-version=51 -DCLAUSE=from -DEXTENDS=BEFORE
|
||||
// RUN: %libomptarget-run-fail-x86_64-pc-linux-gnu 2>&1 \
|
||||
// RUN: | %fcheck-x86_64-pc-linux-gnu
|
||||
|
||||
// --------------------------------------------------
|
||||
// Check 'to' and extends after
|
||||
// --------------------------------------------------
|
||||
|
||||
// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu \
|
||||
// RUN: -fopenmp-version=51 -DCLAUSE=to -DEXTENDS=AFTER
|
||||
// RUN: %libomptarget-run-fail-aarch64-unknown-linux-gnu 2>&1 \
|
||||
// RUN: | %fcheck-aarch64-unknown-linux-gnu
|
||||
|
||||
// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu \
|
||||
// RUN: -fopenmp-version=51 -DCLAUSE=to -DEXTENDS=AFTER
|
||||
// RUN: %libomptarget-run-fail-powerpc64-ibm-linux-gnu 2>&1 \
|
||||
// RUN: | %fcheck-powerpc64-ibm-linux-gnu
|
||||
|
||||
// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu \
|
||||
// RUN: -fopenmp-version=51 -DCLAUSE=to -DEXTENDS=AFTER
|
||||
// RUN: %libomptarget-run-fail-powerpc64le-ibm-linux-gnu 2>&1 \
|
||||
// RUN: | %fcheck-powerpc64le-ibm-linux-gnu
|
||||
|
||||
// RUN: %libomptarget-compile-x86_64-pc-linux-gnu \
|
||||
// RUN: -fopenmp-version=51 -DCLAUSE=to -DEXTENDS=AFTER
|
||||
// RUN: %libomptarget-run-fail-x86_64-pc-linux-gnu 2>&1 \
|
||||
// RUN: | %fcheck-x86_64-pc-linux-gnu
|
||||
|
||||
// --------------------------------------------------
|
||||
// Check 'from' and extends after
|
||||
// --------------------------------------------------
|
||||
|
||||
// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu \
|
||||
// RUN: -fopenmp-version=51 -DCLAUSE=from -DEXTENDS=AFTER
|
||||
// RUN: %libomptarget-run-fail-aarch64-unknown-linux-gnu 2>&1 \
|
||||
// RUN: | %fcheck-aarch64-unknown-linux-gnu
|
||||
|
||||
// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu \
|
||||
// RUN: -fopenmp-version=51 -DCLAUSE=from -DEXTENDS=AFTER
|
||||
// RUN: %libomptarget-run-fail-powerpc64-ibm-linux-gnu 2>&1 \
|
||||
// RUN: | %fcheck-powerpc64-ibm-linux-gnu
|
||||
|
||||
// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu \
|
||||
// RUN: -fopenmp-version=51 -DCLAUSE=from -DEXTENDS=AFTER
|
||||
// RUN: %libomptarget-run-fail-powerpc64le-ibm-linux-gnu 2>&1 \
|
||||
// RUN: | %fcheck-powerpc64le-ibm-linux-gnu
|
||||
|
||||
// RUN: %libomptarget-compile-x86_64-pc-linux-gnu \
|
||||
// RUN: -fopenmp-version=51 -DCLAUSE=from -DEXTENDS=AFTER
|
||||
// RUN: %libomptarget-run-fail-x86_64-pc-linux-gnu 2>&1 \
|
||||
// RUN: | %fcheck-x86_64-pc-linux-gnu
|
||||
|
||||
// END.
|
||||
|
||||
#include <stdio.h>
|
||||
|
||||
#define BEFORE 0
|
||||
#define AFTER 1
|
||||
|
||||
#if EXTENDS == BEFORE
|
||||
# define SMALL 2:3
|
||||
# define LARGE 0:5
|
||||
#elif EXTENDS == AFTER
|
||||
# define SMALL 0:3
|
||||
# define LARGE 0:5
|
||||
#else
|
||||
# error EXTENDS undefined
|
||||
#endif
|
||||
|
||||
int main() {
|
||||
int arr[5];
|
||||
|
||||
// CHECK: addr=0x[[#%x,HOST_ADDR:]], size=[[#%u,SIZE:]]
|
||||
fprintf(stderr, "addr=%p, size=%ld\n", arr, sizeof arr);
|
||||
|
||||
// CHECK-NOT: Libomptarget
|
||||
#pragma omp target data map(alloc: arr[LARGE])
|
||||
{
|
||||
#pragma omp target update CLAUSE(present: arr[SMALL])
|
||||
}
|
||||
|
||||
// CHECK: arr is present
|
||||
fprintf(stderr, "arr is present\n");
|
||||
|
||||
// CHECK: Libomptarget message: device mapping required by 'present' motion modifier does not exist for host address 0x{{0*}}[[#HOST_ADDR]] ([[#SIZE]] bytes)
|
||||
// CHECK: Libomptarget fatal error 1: failure of target construct while offloading is mandatory
|
||||
#pragma omp target data map(alloc: arr[SMALL])
|
||||
{
|
||||
#pragma omp target update CLAUSE(present: arr[LARGE])
|
||||
}
|
||||
|
||||
// CHECK-NOT: arr is present
|
||||
fprintf(stderr, "arr is present\n");
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
|
@ -0,0 +1,136 @@
|
|||
// --------------------------------------------------
|
||||
// Check 'to' and extends before
|
||||
// --------------------------------------------------
|
||||
|
||||
// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu \
|
||||
// RUN: -DCLAUSE=to -DEXTENDS=BEFORE
|
||||
// RUN: %libomptarget-run-aarch64-unknown-linux-gnu 2>&1 \
|
||||
// RUN: | %fcheck-aarch64-unknown-linux-gnu
|
||||
|
||||
// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu \
|
||||
// RUN: -DCLAUSE=to -DEXTENDS=BEFORE
|
||||
// RUN: %libomptarget-run-powerpc64-ibm-linux-gnu 2>&1 \
|
||||
// RUN: | %fcheck-powerpc64-ibm-linux-gnu
|
||||
|
||||
// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu \
|
||||
// RUN: -DCLAUSE=to -DEXTENDS=BEFORE
|
||||
// RUN: %libomptarget-run-powerpc64le-ibm-linux-gnu 2>&1 \
|
||||
// RUN: | %fcheck-powerpc64le-ibm-linux-gnu
|
||||
|
||||
// RUN: %libomptarget-compile-x86_64-pc-linux-gnu \
|
||||
// RUN: -DCLAUSE=to -DEXTENDS=BEFORE
|
||||
// RUN: %libomptarget-run-x86_64-pc-linux-gnu 2>&1 \
|
||||
// RUN: | %fcheck-x86_64-pc-linux-gnu
|
||||
|
||||
// --------------------------------------------------
|
||||
// Check 'from' and extends before
|
||||
// --------------------------------------------------
|
||||
|
||||
// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu \
|
||||
// RUN: -DCLAUSE=from -DEXTENDS=BEFORE
|
||||
// RUN: %libomptarget-run-aarch64-unknown-linux-gnu 2>&1 \
|
||||
// RUN: | %fcheck-aarch64-unknown-linux-gnu
|
||||
|
||||
// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu \
|
||||
// RUN: -DCLAUSE=from -DEXTENDS=BEFORE
|
||||
// RUN: %libomptarget-run-powerpc64-ibm-linux-gnu 2>&1 \
|
||||
// RUN: | %fcheck-powerpc64-ibm-linux-gnu
|
||||
|
||||
// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu \
|
||||
// RUN: -DCLAUSE=from -DEXTENDS=BEFORE
|
||||
// RUN: %libomptarget-run-powerpc64le-ibm-linux-gnu 2>&1 \
|
||||
// RUN: | %fcheck-powerpc64le-ibm-linux-gnu
|
||||
|
||||
// RUN: %libomptarget-compile-x86_64-pc-linux-gnu \
|
||||
// RUN: -DCLAUSE=from -DEXTENDS=BEFORE
|
||||
// RUN: %libomptarget-run-x86_64-pc-linux-gnu 2>&1 \
|
||||
// RUN: | %fcheck-x86_64-pc-linux-gnu
|
||||
|
||||
// --------------------------------------------------
|
||||
// Check 'to' and extends after
|
||||
// --------------------------------------------------
|
||||
|
||||
// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu \
|
||||
// RUN: -DCLAUSE=to -DEXTENDS=AFTER
|
||||
// RUN: %libomptarget-run-aarch64-unknown-linux-gnu 2>&1 \
|
||||
// RUN: | %fcheck-aarch64-unknown-linux-gnu
|
||||
|
||||
// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu \
|
||||
// RUN: -DCLAUSE=to -DEXTENDS=AFTER
|
||||
// RUN: %libomptarget-run-powerpc64-ibm-linux-gnu 2>&1 \
|
||||
// RUN: | %fcheck-powerpc64-ibm-linux-gnu
|
||||
|
||||
// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu \
|
||||
// RUN: -DCLAUSE=to -DEXTENDS=AFTER
|
||||
// RUN: %libomptarget-run-powerpc64le-ibm-linux-gnu 2>&1 \
|
||||
// RUN: | %fcheck-powerpc64le-ibm-linux-gnu
|
||||
|
||||
// RUN: %libomptarget-compile-x86_64-pc-linux-gnu \
|
||||
// RUN: -DCLAUSE=to -DEXTENDS=AFTER
|
||||
// RUN: %libomptarget-run-x86_64-pc-linux-gnu 2>&1 \
|
||||
// RUN: | %fcheck-x86_64-pc-linux-gnu
|
||||
|
||||
// --------------------------------------------------
|
||||
// Check 'from' and extends after
|
||||
// --------------------------------------------------
|
||||
|
||||
// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu \
|
||||
// RUN: -DCLAUSE=from -DEXTENDS=AFTER
|
||||
// RUN: %libomptarget-run-aarch64-unknown-linux-gnu 2>&1 \
|
||||
// RUN: | %fcheck-aarch64-unknown-linux-gnu
|
||||
|
||||
// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu \
|
||||
// RUN: -DCLAUSE=from -DEXTENDS=AFTER
|
||||
// RUN: %libomptarget-run-powerpc64-ibm-linux-gnu 2>&1 \
|
||||
// RUN: | %fcheck-powerpc64-ibm-linux-gnu
|
||||
|
||||
// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu \
|
||||
// RUN: -DCLAUSE=from -DEXTENDS=AFTER
|
||||
// RUN: %libomptarget-run-powerpc64le-ibm-linux-gnu 2>&1 \
|
||||
// RUN: | %fcheck-powerpc64le-ibm-linux-gnu
|
||||
|
||||
// RUN: %libomptarget-compile-x86_64-pc-linux-gnu \
|
||||
// RUN: -DCLAUSE=from -DEXTENDS=AFTER
|
||||
// RUN: %libomptarget-run-x86_64-pc-linux-gnu 2>&1 \
|
||||
// RUN: | %fcheck-x86_64-pc-linux-gnu
|
||||
|
||||
// END.
|
||||
|
||||
#include <stdio.h>
|
||||
|
||||
#define BEFORE 0
|
||||
#define AFTER 1
|
||||
|
||||
#if EXTENDS == BEFORE
|
||||
# define SMALL 2:3
|
||||
# define LARGE 0:5
|
||||
#elif EXTENDS == AFTER
|
||||
# define SMALL 0:3
|
||||
# define LARGE 0:5
|
||||
#else
|
||||
# error EXTENDS undefined
|
||||
#endif
|
||||
|
||||
int main() {
|
||||
int arr[5];
|
||||
|
||||
// CHECK-NOT: Libomptarget
|
||||
#pragma omp target data map(alloc: arr[LARGE])
|
||||
{
|
||||
#pragma omp target update CLAUSE(arr[SMALL])
|
||||
}
|
||||
|
||||
// CHECK: success
|
||||
fprintf(stderr, "success\n");
|
||||
|
||||
// CHECK-NOT: Libomptarget
|
||||
#pragma omp target data map(alloc: arr[SMALL])
|
||||
{
|
||||
#pragma omp target update CLAUSE(arr[LARGE])
|
||||
}
|
||||
|
||||
// CHECK: success
|
||||
fprintf(stderr, "success\n");
|
||||
|
||||
return 0;
|
||||
}
|
||||
Loading…
Reference in New Issue