From 41b1aefecb9447620dd182b0352abed0df05665c Mon Sep 17 00:00:00 2001 From: "Joel E. Denny" Date: Wed, 5 Aug 2020 16:47:29 -0400 Subject: [PATCH] [OpenMP] Fix `present` diagnostic for array extension For example, without this patch, the following fails as expected with or without the `present` modifier, but the `present` modifier doesn't produce its usual diagnostic: ``` #pragma omp target data map(alloc: arr[0:2]) { #pragma omp target map(present, tofrom: arr[0:100]) // not fully present ; } ``` Reviewed By: grokos, vzakhari Differential Revision: https://reviews.llvm.org/D85320 --- openmp/libomptarget/src/device.cpp | 10 +- .../test/mapping/present/target_array_extension.c | 112 +++++++++++++++++++++ .../mapping/present/target_data_array_extension.c | 112 +++++++++++++++++++++ 3 files changed, 233 insertions(+), 1 deletion(-) create mode 100644 openmp/libomptarget/test/mapping/present/target_array_extension.c create mode 100644 openmp/libomptarget/test/mapping/present/target_data_array_extension.c diff --git a/openmp/libomptarget/src/device.cpp b/openmp/libomptarget/src/device.cpp index 7049463..5a01257 100644 --- a/openmp/libomptarget/src/device.cpp +++ b/openmp/libomptarget/src/device.cpp @@ -191,7 +191,15 @@ void *DeviceTy::getOrAllocTgtPtr(void *HstPtrBegin, void *HstPtrBase, rc = (void *)tp; } else if ((lr.Flags.ExtendsBefore || lr.Flags.ExtendsAfter) && !IsImplicit) { // Explicit extension of mapped data - not allowed. - DP("Explicit extension of mapping is not allowed.\n"); + MESSAGE("explicit extension not allowed: host address specified is " DPxMOD + " (%" PRId64 " bytes), but device allocation maps to host at " + DPxMOD " (%" PRId64 " bytes)", + DPxPTR(HstPtrBegin), Size, DPxPTR(lr.Entry->HstPtrBegin), + lr.Entry->HstPtrEnd - lr.Entry->HstPtrBegin); + if (HasPresentModifier) + MESSAGE("device mapping required by 'present' map type modifier does not " + "exist for host address " DPxMOD " (%" PRId64 " bytes)", + DPxPTR(HstPtrBegin), Size); } else if (RTLs->RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY && !HasCloseModifier) { // If unified shared memory is active, implicitly mapped variables that are diff --git a/openmp/libomptarget/test/mapping/present/target_array_extension.c b/openmp/libomptarget/test/mapping/present/target_array_extension.c new file mode 100644 index 0000000..870be39 --- /dev/null +++ b/openmp/libomptarget/test/mapping/present/target_array_extension.c @@ -0,0 +1,112 @@ +// -------------------------------------------------- +// Check extends before +// -------------------------------------------------- + +// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu \ +// RUN: -fopenmp-version=51 -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 -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 -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 -DEXTENDS=BEFORE +// RUN: %libomptarget-run-fail-x86_64-pc-linux-gnu 2>&1 \ +// RUN: | %fcheck-x86_64-pc-linux-gnu + +// -------------------------------------------------- +// Check extends after +// -------------------------------------------------- + +// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu \ +// RUN: -fopenmp-version=51 -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 -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 -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 -DEXTENDS=AFTER +// RUN: %libomptarget-run-fail-x86_64-pc-linux-gnu 2>&1 \ +// RUN: | %fcheck-x86_64-pc-linux-gnu + +// END. + +#include + +#define BEFORE 0 +#define AFTER 1 + +#define SIZE 100 + +#if EXTENDS == BEFORE +# define SMALL_BEG (SIZE-2) +# define SMALL_END SIZE +# define LARGE_BEG 0 +# define LARGE_END SIZE +#elif EXTENDS == AFTER +# define SMALL_BEG 0 +# define SMALL_END 2 +# define LARGE_BEG 0 +# define LARGE_END SIZE +#else +# error EXTENDS undefined +#endif + +#define SMALL_SIZE (SMALL_END-SMALL_BEG) +#define LARGE_SIZE (LARGE_END-LARGE_BEG) + +#define SMALL SMALL_BEG:SMALL_SIZE +#define LARGE LARGE_BEG:LARGE_SIZE + +int main() { + int arr[SIZE]; + + // CHECK: addr=0x[[#%x,SMALL_ADDR:]], size=[[#%u,SMALL_BYTES:]] + fprintf(stderr, "addr=%p, size=%ld\n", &arr[SMALL_BEG], + SMALL_SIZE * sizeof arr[0]); + + // CHECK: addr=0x[[#%x,LARGE_ADDR:]], size=[[#%u,LARGE_BYTES:]] + fprintf(stderr, "addr=%p, size=%ld\n", &arr[LARGE_BEG], + LARGE_SIZE * sizeof arr[0]); + + // CHECK-NOT: Libomptarget +#pragma omp target data map(alloc: arr[LARGE]) + { +#pragma omp target map(present, tofrom: arr[SMALL]) + ; + } + + // CHECK: arr is present + fprintf(stderr, "arr is present\n"); + + // CHECK: Libomptarget message: explicit extension not allowed: host address specified is 0x{{0*}}[[#LARGE_ADDR]] ([[#LARGE_BYTES]] bytes), but device allocation maps to host at 0x{{0*}}[[#SMALL_ADDR]] ([[#SMALL_BYTES]] bytes) + // CHECK: Libomptarget message: device mapping required by 'present' map type modifier does not exist for host address 0x{{0*}}[[#LARGE_ADDR]] ([[#LARGE_BYTES]] 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 map(present, tofrom: arr[LARGE]) + ; + } + + // CHECK-NOT: arr is present + fprintf(stderr, "arr is present\n"); + + return 0; +} diff --git a/openmp/libomptarget/test/mapping/present/target_data_array_extension.c b/openmp/libomptarget/test/mapping/present/target_data_array_extension.c new file mode 100644 index 0000000..3aef777 --- /dev/null +++ b/openmp/libomptarget/test/mapping/present/target_data_array_extension.c @@ -0,0 +1,112 @@ +// -------------------------------------------------- +// Check extends before +// -------------------------------------------------- + +// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu \ +// RUN: -fopenmp-version=51 -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 -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 -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 -DEXTENDS=BEFORE +// RUN: %libomptarget-run-fail-x86_64-pc-linux-gnu 2>&1 \ +// RUN: | %fcheck-x86_64-pc-linux-gnu + +// -------------------------------------------------- +// Check extends after +// -------------------------------------------------- + +// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu \ +// RUN: -fopenmp-version=51 -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 -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 -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 -DEXTENDS=AFTER +// RUN: %libomptarget-run-fail-x86_64-pc-linux-gnu 2>&1 \ +// RUN: | %fcheck-x86_64-pc-linux-gnu + +// END. + +#include + +#define BEFORE 0 +#define AFTER 1 + +#define SIZE 100 + +#if EXTENDS == BEFORE +# define SMALL_BEG (SIZE-2) +# define SMALL_END SIZE +# define LARGE_BEG 0 +# define LARGE_END SIZE +#elif EXTENDS == AFTER +# define SMALL_BEG 0 +# define SMALL_END 2 +# define LARGE_BEG 0 +# define LARGE_END SIZE +#else +# error EXTENDS undefined +#endif + +#define SMALL_SIZE (SMALL_END-SMALL_BEG) +#define LARGE_SIZE (LARGE_END-LARGE_BEG) + +#define SMALL SMALL_BEG:SMALL_SIZE +#define LARGE LARGE_BEG:LARGE_SIZE + +int main() { + int arr[SIZE]; + + // CHECK: addr=0x[[#%x,SMALL_ADDR:]], size=[[#%u,SMALL_BYTES:]] + fprintf(stderr, "addr=%p, size=%ld\n", &arr[SMALL_BEG], + SMALL_SIZE * sizeof arr[0]); + + // CHECK: addr=0x[[#%x,LARGE_ADDR:]], size=[[#%u,LARGE_BYTES:]] + fprintf(stderr, "addr=%p, size=%ld\n", &arr[LARGE_BEG], + LARGE_SIZE * sizeof arr[0]); + + // CHECK-NOT: Libomptarget +#pragma omp target data map(alloc: arr[LARGE]) + { +#pragma omp target data map(present, tofrom: arr[SMALL]) + ; + } + + // CHECK: arr is present + fprintf(stderr, "arr is present\n"); + + // CHECK: Libomptarget message: explicit extension not allowed: host address specified is 0x{{0*}}[[#LARGE_ADDR]] ([[#LARGE_BYTES]] bytes), but device allocation maps to host at 0x{{0*}}[[#SMALL_ADDR]] ([[#SMALL_BYTES]] bytes) + // CHECK: Libomptarget message: device mapping required by 'present' map type modifier does not exist for host address 0x{{0*}}[[#LARGE_ADDR]] ([[#LARGE_BYTES]] 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 data map(present, tofrom: arr[LARGE]) + ; + } + + // CHECK-NOT: arr is present + fprintf(stderr, "arr is present\n"); + + return 0; +} -- 2.7.4