[Openmp-commits] [openmp] r361294 - [OpenMP][libomptarget] Enable requires flags for target libraries.
Gheorghe-Teodor Bercea via Openmp-commits
openmp-commits at lists.llvm.org
Tue May 21 12:35:02 PDT 2019
Author: gbercea
Date: Tue May 21 12:35:02 2019
New Revision: 361294
URL: http://llvm.org/viewvc/llvm-project?rev=361294&view=rev
Log:
[OpenMP][libomptarget] Enable requires flags for target libraries.
Summary:
Target link variables are currently implemented by creating a copy of the variables on the device side and unified memory never gets exploited.
When the prgram uses the:
```
#pragma omp requires unified_shared_memory
```
directive in conjunction with a declare target link, the linked variable is no longer allocated on the device and the host version is used instead.
This behavior is overridden by performing an explicit mapping.
A Clang side patch is required.
Reviewers: ABataev, AlexEichenberger, grokos, Hahnfeld
Reviewed By: AlexEichenberger, grokos, Hahnfeld
Subscribers: Hahnfeld, jfb, guansong, jdoerfert, openmp-commits
Tags: #openmp
Differential Revision: https://reviews.llvm.org/D60223
Added:
openmp/trunk/libomptarget/test/offloading/requires.c
Modified:
openmp/trunk/libomptarget/include/omptarget.h
openmp/trunk/libomptarget/src/device.cpp
openmp/trunk/libomptarget/src/device.h
openmp/trunk/libomptarget/src/exports
openmp/trunk/libomptarget/src/interface.cpp
openmp/trunk/libomptarget/src/rtl.cpp
openmp/trunk/libomptarget/src/rtl.h
Modified: openmp/trunk/libomptarget/include/omptarget.h
URL: http://llvm.org/viewvc/llvm-project/openmp/trunk/libomptarget/include/omptarget.h?rev=361294&r1=361293&r2=361294&view=diff
==============================================================================
--- openmp/trunk/libomptarget/include/omptarget.h (original)
+++ openmp/trunk/libomptarget/include/omptarget.h Tue May 21 12:35:02 2019
@@ -60,6 +60,21 @@ enum OpenMPOffloadingDeclareTargetFlags
OMP_DECLARE_TARGET_DTOR = 0x04
};
+enum OpenMPOffloadingRequiresDirFlags {
+ /// flag undefined.
+ OMP_REQ_UNDEFINED = 0x000,
+ /// no requires directive present.
+ OMP_REQ_NONE = 0x001,
+ /// reverse_offload clause.
+ OMP_REQ_REVERSE_OFFLOAD = 0x002,
+ /// unified_address clause.
+ OMP_REQ_UNIFIED_ADDRESS = 0x004,
+ /// unified_shared_memory clause.
+ OMP_REQ_UNIFIED_SHARED_MEMORY = 0x008,
+ /// dynamic_allocators clause.
+ OMP_REQ_DYNAMIC_ALLOCATORS = 0x010
+};
+
/// This struct is a record of an entry point or global. For a function
/// entry point the size is expected to be zero
struct __tgt_offload_entry {
@@ -113,6 +128,9 @@ int omp_target_associate_ptr(void *host_
size_t device_offset, int device_num);
int omp_target_disassociate_ptr(void *host_ptr, int device_num);
+/// add the clauses of the requires directives in a given file
+void __tgt_register_requires(int64_t flags);
+
/// adds a target shared library to the target execution image
void __tgt_register_lib(__tgt_bin_desc *desc);
Modified: openmp/trunk/libomptarget/src/device.cpp
URL: http://llvm.org/viewvc/llvm-project/openmp/trunk/libomptarget/src/device.cpp?rev=361294&r1=361293&r2=361294&view=diff
==============================================================================
--- openmp/trunk/libomptarget/src/device.cpp (original)
+++ openmp/trunk/libomptarget/src/device.cpp Tue May 21 12:35:02 2019
@@ -152,7 +152,7 @@ LookupResult DeviceTy::lookupMapping(voi
// Used by target_data_begin
// Return the target pointer begin (where the data will be moved).
-// Allocate memory if this is the first occurrence if this mapping.
+// Allocate memory if this is the first occurrence of this mapping.
// Increment the reference counter.
// If NULL is returned, then either data allocation failed or the user tried
// to do an illegal mapping.
Modified: openmp/trunk/libomptarget/src/device.h
URL: http://llvm.org/viewvc/llvm-project/openmp/trunk/libomptarget/src/device.h?rev=361294&r1=361293&r2=361294&view=diff
==============================================================================
--- openmp/trunk/libomptarget/src/device.h (original)
+++ openmp/trunk/libomptarget/src/device.h Tue May 21 12:35:02 2019
@@ -98,11 +98,13 @@ struct DeviceTy {
uint64_t loopTripCnt;
+ int64_t RTLRequiresFlags;
+
DeviceTy(RTLInfoTy *RTL)
: DeviceID(-1), RTL(RTL), RTLDeviceID(-1), IsInit(false), InitFlag(),
HasPendingGlobals(false), HostDataToTargetMap(),
PendingCtorsDtors(), ShadowPtrMap(), DataMapMtx(), PendingGlobalsMtx(),
- ShadowMtx(), loopTripCnt(0) {}
+ ShadowMtx(), loopTripCnt(0), RTLRequiresFlags(0) {}
// The existence of mutexes makes DeviceTy non-copyable. We need to
// provide a copy constructor and an assignment operator explicitly.
@@ -112,7 +114,8 @@ struct DeviceTy {
HostDataToTargetMap(d.HostDataToTargetMap),
PendingCtorsDtors(d.PendingCtorsDtors), ShadowPtrMap(d.ShadowPtrMap),
DataMapMtx(), PendingGlobalsMtx(),
- ShadowMtx(), loopTripCnt(d.loopTripCnt) {}
+ ShadowMtx(), loopTripCnt(d.loopTripCnt),
+ RTLRequiresFlags(d.RTLRequiresFlags) {}
DeviceTy& operator=(const DeviceTy &d) {
DeviceID = d.DeviceID;
@@ -124,6 +127,7 @@ struct DeviceTy {
PendingCtorsDtors = d.PendingCtorsDtors;
ShadowPtrMap = d.ShadowPtrMap;
loopTripCnt = d.loopTripCnt;
+ RTLRequiresFlags = d.RTLRequiresFlags;
return *this;
}
Modified: openmp/trunk/libomptarget/src/exports
URL: http://llvm.org/viewvc/llvm-project/openmp/trunk/libomptarget/src/exports?rev=361294&r1=361293&r2=361294&view=diff
==============================================================================
--- openmp/trunk/libomptarget/src/exports (original)
+++ openmp/trunk/libomptarget/src/exports Tue May 21 12:35:02 2019
@@ -1,5 +1,6 @@
VERS1.0 {
global:
+ __tgt_register_requires;
__tgt_register_lib;
__tgt_unregister_lib;
__tgt_target_data_begin;
Modified: openmp/trunk/libomptarget/src/interface.cpp
URL: http://llvm.org/viewvc/llvm-project/openmp/trunk/libomptarget/src/interface.cpp?rev=361294&r1=361293&r2=361294&view=diff
==============================================================================
--- openmp/trunk/libomptarget/src/interface.cpp (original)
+++ openmp/trunk/libomptarget/src/interface.cpp Tue May 21 12:35:02 2019
@@ -57,7 +57,7 @@ static void HandleTargetOutcome(bool suc
}
break;
case tgt_default:
- FATAL_MESSAGE0(1, "default offloading policy must switched to "
+ FATAL_MESSAGE0(1, "default offloading policy must switched to "
"mandatory or disabled");
break;
case tgt_mandatory:
@@ -69,6 +69,12 @@ static void HandleTargetOutcome(bool suc
}
////////////////////////////////////////////////////////////////////////////////
+/// adds requires flags
+EXTERN void __tgt_register_requires(int64_t flags) {
+ RTLs.RegisterRequires(flags);
+}
+
+////////////////////////////////////////////////////////////////////////////////
/// adds a target shared library to the target execution image
EXTERN void __tgt_register_lib(__tgt_bin_desc *desc) {
RTLs.RegisterLib(desc);
Modified: openmp/trunk/libomptarget/src/rtl.cpp
URL: http://llvm.org/viewvc/llvm-project/openmp/trunk/libomptarget/src/rtl.cpp?rev=361294&r1=361293&r2=361294&view=diff
==============================================================================
--- openmp/trunk/libomptarget/src/rtl.cpp (original)
+++ openmp/trunk/libomptarget/src/rtl.cpp Tue May 21 12:35:02 2019
@@ -186,6 +186,46 @@ static void RegisterGlobalCtorsDtorsForI
}
}
+void RTLsTy::RegisterRequires(int64_t flags) {
+ // TODO: add more elaborate check.
+ // Minimal check: only set requires flags if previous value
+ // is undefined. This ensures that only the first call to this
+ // function will set the requires flags. All subsequent calls
+ // will be checked for compatibility.
+ assert(flags != OMP_REQ_UNDEFINED &&
+ "illegal undefined flag for requires directive!");
+ if (RequiresFlags == OMP_REQ_UNDEFINED) {
+ RequiresFlags = flags;
+ return;
+ }
+
+ // If multiple compilation units are present enforce
+ // consistency across all of them for require clauses:
+ // - reverse_offload
+ // - unified_address
+ // - unified_shared_memory
+ if ((RequiresFlags & OMP_REQ_REVERSE_OFFLOAD) !=
+ (flags & OMP_REQ_REVERSE_OFFLOAD)) {
+ FATAL_MESSAGE0(1,
+ "'#pragma omp requires reverse_offload' not used consistently!");
+ }
+ if ((RequiresFlags & OMP_REQ_UNIFIED_ADDRESS) !=
+ (flags & OMP_REQ_UNIFIED_ADDRESS)) {
+ FATAL_MESSAGE0(1,
+ "'#pragma omp requires unified_address' not used consistently!");
+ }
+ if ((RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY) !=
+ (flags & OMP_REQ_UNIFIED_SHARED_MEMORY)) {
+ FATAL_MESSAGE0(1,
+ "'#pragma omp requires unified_shared_memory' not used consistently!");
+ }
+
+ // TODO: insert any other missing checks
+
+ DP("New requires flags %ld compatible with existing %ld!\n",
+ flags, RequiresFlags);
+}
+
void RTLsTy::RegisterLib(__tgt_bin_desc *desc) {
// Attempt to load all plugins available in the system.
std::call_once(initFlag, &RTLsTy::LoadRTLs, this);
@@ -222,6 +262,8 @@ void RTLsTy::RegisterLib(__tgt_bin_desc
Devices[start + device_id].DeviceID = start + device_id;
// RTL local device ID
Devices[start + device_id].RTLDeviceID = device_id;
+ // RTL requires flags
+ Devices[start + device_id].RTLRequiresFlags = RequiresFlags;
}
// Initialize the index of this RTL and save it in the used RTLs.
Modified: openmp/trunk/libomptarget/src/rtl.h
URL: http://llvm.org/viewvc/llvm-project/openmp/trunk/libomptarget/src/rtl.h?rev=361294&r1=361293&r2=361294&view=diff
==============================================================================
--- openmp/trunk/libomptarget/src/rtl.h (original)
+++ openmp/trunk/libomptarget/src/rtl.h Tue May 21 12:35:02 2019
@@ -118,8 +118,13 @@ public:
// binaries.
std::vector<RTLInfoTy *> UsedRTLs;
+ int64_t RequiresFlags;
+
explicit RTLsTy() {}
+ // Register the clauses of the requires directive.
+ void RegisterRequires(int64_t flags);
+
// Register a shared library with all (compatible) RTLs.
void RegisterLib(__tgt_bin_desc *desc);
Added: openmp/trunk/libomptarget/test/offloading/requires.c
URL: http://llvm.org/viewvc/llvm-project/openmp/trunk/libomptarget/test/offloading/requires.c?rev=361294&view=auto
==============================================================================
--- openmp/trunk/libomptarget/test/offloading/requires.c (added)
+++ openmp/trunk/libomptarget/test/offloading/requires.c Tue May 21 12:35:02 2019
@@ -0,0 +1,46 @@
+// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-aarch64-unknown-linux-gnu 2>&1 | %fcheck-aarch64-unknown-linux-gnu -allow-empty -check-prefix=DEBUG
+// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-powerpc64-ibm-linux-gnu 2>&1 | %fcheck-powerpc64-ibm-linux-gnu -allow-empty -check-prefix=DEBUG
+// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-powerpc64le-ibm-linux-gnu 2>&1 | %fcheck-powerpc64le-ibm-linux-gnu -allow-empty -check-prefix=DEBUG
+// RUN: %libomptarget-compile-x86_64-pc-linux-gnu && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-x86_64-pc-linux-gnu 2>&1 | %fcheck-x86_64-pc-linux-gnu -allow-empty -check-prefix=DEBUG
+// REQUIRES: libomptarget-debug
+
+/*
+ Test for the 'requires' clause check.
+ When a target region is used, the requires flags are set in the
+ runtime for the entire compilation unit. If the flags are set again,
+ (for whatever reason) the set must be consistent with previously
+ set values.
+*/
+#include <stdio.h>
+#include <omp.h>
+
+// ---------------------------------------------------------------------------
+// Various definitions copied from OpenMP RTL
+
+extern void __tgt_register_requires(int64_t);
+
+// End of definitions copied from OpenMP RTL.
+// ---------------------------------------------------------------------------
+
+void run_reg_requires() {
+ // Before the target region is registered, the requires registers the status
+ // of the requires clauses. Since there are no requires clauses in this file
+ // the flags state can only be OMP_REQ_NONE i.e. 1.
+
+ // This is the 2nd time this function is called so it should print the debug
+ // info belonging to the check.
+ __tgt_register_requires(1);
+ __tgt_register_requires(1);
+ // DEBUG: New requires flags 1 compatible with existing 1!
+}
+
+// ---------------------------------------------------------------------------
+int main() {
+ run_reg_requires();
+
+// This also runs reg requires for the first time.
+#pragma omp target
+ {}
+
+ return 0;
+}
\ No newline at end of file
More information about the Openmp-commits
mailing list