[OpenMP] Initial implementation of OpenMP offloading library - libomptarget.

[grokos]

Summary

This patches proposes an initial implementation of the libomptarget based on what we have working today in https://github.com/clang-omp.

I tried to follow what has been done in openmp/runtime in terms of the cmake build system.

The library sources are placed under libomptarget and it consists of three components:

• device agnostic library (the interface that clang uses)
• device plugins - right now we create a plug-in for CUDA-enabled devices, and a generic 64-bit plugin that works on powerpcBE/LE and x86_64 (mostly for testing purposes)
• device runtime library for CUDA enabled devices.

The interface of the plugins and target agnostic library are documented in http://goo.gl/L1rnKJ.

I included the logic for testing based on llvm-lit and FileCheck. Currently there are only two regression tests that basically check that offloading failed. This is because the target side codegen (http://reviews.llvm.org/D12614) and driver support (http://reviews.llvm.org/D9888) are still under review in clang. Once that functionality is in clang I plan to include more tests to exercise the library.

In this patch, I tried to organize things based on what makes sense to me and also based on previous discussions with other OpenMP project contributors. I am happy to organize things in a different way if suggested to do so - the goal here is also open the discussion on how this library should be contributed to the the project.

Let me now any suggestions/comments you may have,

Thanks!
Samuel

[Hahnfeld]

Great work!
Would it be possible to split this into multiple, smaller reviews (for main library, plugins and device runtime for nvptx)?

Greetings
Jonas

libomptarget/deviceRTLs/nvptx/src/parallel.cu

151–158 ↗

(On Diff #38275)

Commented out - should this be committed?

libomptarget/deviceRTLs/nvptx/src/reduction.cu

71–117 ↗

(On Diff #38275)

Commented or #if 0

157–178 ↗

(On Diff #38275)

Not ready for shipping?

780–795 ↗

(On Diff #38275)

#if 0

1001–1028 ↗

(On Diff #38275)

some more #if 0

1239–1246 ↗

(On Diff #38275)

#if 1?

[sfantao]

Partition the original patch and do some code cleanup as suggested by Jonas.

The plugins patch is posted in http://reviews.llvm.org/D14253.

The device RTLs patch is posted in http://reviews.llvm.org/D14254.

[sfantao]

Rebase and add libomptarget install rule that was missing in the previous diff.

[sfantao]

Ping!

[Hahnfeld]

I'm fine otherwise - but definitely not the one for any final decision!

libomptarget/cmake/Modules/LibomptargetGetDependencies.cmake

62–77

In previous versions you have been using PkgConfig - is there a particular reason that this was removed?

(I'm asking because RHEL 6 is using a rather weird path: /usr/lib64/libffi-3.0.5/include/ffi.h)

libomptarget/src/omptarget.cpp

197–198

auto and range-based loop?

227–228

auto and range-based loop?

257–258

auto and range-based loop?

1022–1029

Should this be committed for development?

[sfantao]

Rewrite a few loops to use auto and iterate on ranges. Add hint entries to cmake for the detection of libffi following issue reported by Jonas Hahnfeld.

[sfantao]

Hi Jonas,

Thanks for reviewing the patch!

libomptarget/cmake/Modules/LibomptargetGetDependencies.cmake

62–77

I was using the same logic I used for libelf and that was working well in the environments I tested this with. I restored the hint based mechanism so your problem may now be solved. Let me know if not.

Thanks for reporting the issue.

libomptarget/src/omptarget.cpp

197–198

Done!

227–228

I am using auto now but have to keep the iterator so that the entry in the map can be removed in the body of the loop.

257–258

Done!

1022–1029

I've removed these two debug calls.

[Hahnfeld]

libomptarget/cmake/Modules/LibomptargetGetDependencies.cmake

64–79

After looking at this quite some time and wondering why it didn't work: It's CMake, they use curly braces. With this small change everything builds correctly out-of-the-box, thanks!

[sfantao]

Use curly braces in variables when looking for libffi.

[sfantao]

Hi Jonas sorry for the mess up with the braces. It should be good now.

Thanks!
Samuel

libomptarget/cmake/Modules/LibomptargetGetDependencies.cmake

64–79

Oh, right, sorry, my bad.... The new diff fixes it.

[sfantao]

Remove considerations about the format (e.g. ELF) of the target binary. The format is only going to be evaluated by the plugins. This makes the target agnostic part of the library to be really target agnostic.

[Hahnfeld]

I just came across an interesting problem...

libomptarget/src/omptarget.cpp

724–727

I think this assumption may be wrong: The standard defines that each device has an initial data environment (section 1.4.2).

Therefore a #pragma omp target update may be a first (and valid) statement which means that it has to be init here as well. Currently the library will in such a case return at this point and ignore the data transfer.

[RaviNarayanaswamy]

OpenMP says : " If the corresponding list item is not present in the device data environment then no assignment occurs to or from the original list item."

Device.IsInit being false can be one's implementation as the list item not present in the device data environment.

Ravi

[Hahnfeld]

In D14031#320528, @RaviNarayanaswamy wrote:

OpenMP says : " If the corresponding list item is not present in the device data environment then no assignment occurs to or from the original list item."

Device.IsInit being false can be one's implementation as the list item not present in the device data environment.

Ravi

Sorry for not directly posting the code in question that I had in mind:

#include <stdio.h>

#pragma omp declare target
int a = 1;
#pragma omp end declare target

int main(int argc, char* argv[]) {
	a = 2;
	#pragma omp target update to(a)

	#pragma omp target
	{
		printf("a = %d\n", a);
	}

	return 0;
}

What would the expected result be? However I just noticed that yesterday's simple init() didn't change it to 2 because the code and global variables are still only mapped in the first omp target...

[sfantao]

Hi Jonas,

I think you are correct, we need to register the device and initialize it
for #target update and #target data directives not only the #target
directive. I'll send a fix for that soon.

Thanks,
Samuel

From: Jonas Hahnfeld <[e-mail]>
To: Samuel F Antao/Watson/IBM@IBMUS, [e-mail],

[e-mail], [e-mail],
[e-mail], [e-mail],
[e-mail]

Cc: [e-mail], Alexandre

Eichenberger/Watson/IBM@IBMUS, [e-mail],
[e-mail], [e-mail], Kevin K
O'Brien/Watson/IBM@IBMUS, Carlo Bertolli/Watson/IBM@IBMUS,
[e-mail], [e-mail]

Date: 01/07/2016 02:53 AM
Subject: Re: [PATCH] D14031: [OpenMP] Initial implementation of OpenMP

offloading library - libomptarget.

Hahnfeld added a comment.

In http://reviews.llvm.org/D14031#320528, @RaviNarayanaswamy wrote:

OpenMP says : " If the corresponding list item is not present in the

device data environment then no assignment occurs to or from the original
list item."

Device.IsInit being false can be one's implementation as the list item

not present in the device data environment.

Ravi

Sorry for not directly posting the code in question that I had in mind:

#include <stdio.h>

#pragma omp declare target
int a = 1;
#pragma omp end declare target

int main(int argc, char* argv[]) {
		 a = 2;
		 #pragma omp target update to(a)

		 #pragma omp target
		 {
		 		 printf("a = %d\n", a);
		 }

		 return 0;
}

What would the expected result be? However I just noticed that yesterday's
simple init() didn't change it to 2 because the code and global
variables are still only mapped in the first omp target...

http://reviews.llvm.org/D14031

[Hahnfeld]

Below some comments on the tests

libomptarget/test/CMakeLists.txt

29–35

Maybe we could later on think about sharing some testing setup between libomp and libomptarget...

65–68

I think this can safely assume the relative location of ../runtime/src

libomptarget/test/offloading/offloading_success.c

10

Maybe make this isHost = 1. clang 3.7.1 seems to completely ignore the target code, so omp_is_initial_device is never called...

libomptarget/test/offloading/offloading_success.cpp

10

likewise

[• cbergstrom]

Please take this comment as more a nitpick instead of having tested it

In general I see a lot of really dirty code - Why so many casts? What makes "long long" universally correct and portable.. are there standard types we can use to represent pointers?

https://docs.oracle.com/cd/E19683-01/806-6543/chapter3-10/index.html
http://stackoverflow.com/questions/25211015/using-x-to-print-the-hex-address-contained-in-a-pointer

uintptr_t
printf("%" PRIxPTR "\n", (uintptr_t) p);

---

Is it really a good idea to have two functions which are almost the same, but one is blocking and the other isn't?

tgt_target_data_begin_nowait vs tgt_target_data_begin

Would it not be "better" or more clean for it to take a parameter?

libomptarget/src/omptarget.cpp

349

Instead of casting to long long and friends - what about making the type uintptr_t ?

[sergos]

libomptarget/src/omptarget.cpp

724–727

not quite relevant. the initial data environment should be initialized at the point of __tgt_register_lib() call that should prepend any data/control transfer to/from target. that's why we can say the target image was not initialized yet.

[Hahnfeld]

libomptarget/src/omptarget.cpp

724–727

I think the plan was to lazily initialize the device once it is really needed.

I support this because there may be more libs registered than actually used. So it would be wasted to initialize all of their devices and transfer data to them.

[sfantao]

libomptarget/src/omptarget.cpp

724–727

I agree with Jonas, we should allow the device to be initialized in target update so that we only copy the data for that device and not other devices that may not be used at all by the application. I'll fix this.

[pawosm01]

This looks interesting. Will it ever be applied upstream?

[hyviquel]

libomptarget/src/omptarget.cpp

930

Why the offloading is continuing if the mapping of data already failed...
target_data_begin could return OFFLOAD_FAIL or OFFLOAD_SUCCESS to allow the offloading to stop earlier

981

you return OFFLOAD_SUCCESS even if something went wrong in target_data_end

[grokos]

libomptarget/cmake/Modules/LibomptargetGetDependencies.cmake

64–79

So I'm marking this issue as "done".

libomptarget/src/omptarget.cpp

349

Done in the new diff.

724–727

Done in new diff.

930

I fixed that behavior in the new diff. If target_data_begin fails, the offloading stops.

981

Fixed.

libomptarget/test/CMakeLists.txt

29–35

Sure, let's keep this on our list for future implementation.

65–68

Actually, it's ../../runtime/src. I've hardcoded it as:

${CMAKE_CURRENT_BINARY_DIR}/../../runtime/src

libomptarget/test/offloading/offloading_success.c

10

With the current version of clang this should work. Let me know if I'm mistaken and I'll push another patch.

libomptarget/test/offloading/offloading_success.cpp

10

Same here...

[grokos]

Addressed comments from previous revision plus:

1. support for OpenMP 4.5 target map types
2. implementation of target API functions (omp_target_*)
3. protection of internal data structures against data races

[Hahnfeld]

This diff looks wrong, maybe it shows the changes compared to the last revision?

[grokos]

My bad, I've upload the correct diff.

[Hahnfeld]

Full review and comments all over

libomptarget/Build_With_CMake.txt

103–116

This should probably go under NVPTX device RTL specific

libomptarget/CMakeLists.txt

97–98

There is LLVM_LIBDIR_SUFFIX which should probably be respected (compare to LIBOMP_LIBDIR_SUFFIX)

107–109

These are not part of this patch and should therefore not be added here

libomptarget/README.txt

61

Is that true? I remember there were some changes to the section names in the fatbinary?

libomptarget/src/omptarget.cpp

123–131

Could this reuse the assignment operator? Or otherwise do this in the initializer list

256–258

Please move this up, directly below struct DeviceTy

284

This is not in the standard, should this really start with OMP_?

285–286

can be collapsed

422–423

can be collapsed

495

Do we want to check device_is_ready here?

509–512

Should we have a list of all allocated addresses above?

514

Do we want to check device_is_ready here?

533–534

can be collapsed

539–540

likewise

911

bool ForceDelete?

1107

if no RTL was found?

1113–1137

Can't this be done immediately after the corresponding RTL was found the same way we register the image into the translation table? Maybe also refactor the code into another static function?

1115

I think this should either be Devices[FoundRTL->Idx + i] or *FoundRTL->Devices[i]?

1178

I think this should either be Devices[FoundRTL->Idx + i] or *FoundRTL->Devices[i]?

1194–1206

I think this is only needed once per desc and not per img, isn't it?

1355–1356

The Device should only be fetched after the device_id has been proved valid

1358–1378

This looks the same as device_is_ready, can this be reused?

1384–1385

can be collapsed

1394–1396

Backwards-compatible on the first check-in? That sounds weird...

2117

break afterwards?

2132

likewise?

libomptarget/src/omptarget.h

87

Maybe rename this to NumDeviceImages?

89–90

Can those be renamed to have Host in the same so that I'm not that much confused? :D

167

Only as data_end, not as data_update?

libomptarget/test/CMakeLists.txt

71–77

1. This will abort CMake which probably isn't a good idea.
2. MSVC is possible for standalone builds?

[sfantao]

George,

Thanks for the patch. Here are few comments following Jonas review.

libomptarget/CMakeLists.txt

39

Following Jonas suggestion, we can select the install suffix here like libomp does. Something like:

if(${LIBOMPTARGET_STANDALONE_BUILD})
  set(LIBOMPTARGET_ENABLE_WERROR FALSE CACHE BOOL
    "Enable -Werror flags to turn warnings into errors for supporting compilers.")
  # CMAKE_BUILD_TYPE was not defined, set default to Release
  if(NOT CMAKE_BUILD_TYPE)
    set(CMAKE_BUILD_TYPE Release)
  endif()
  set(LIBOMPTARGET_LIBDIR_SUFFIX "" CACHE STRING
    "suffix of lib installation directory e.g., 64 => lib64")
else()
  set(LIBOMPTARGET_ENABLE_WERROR ${LLVM_ENABLE_WERROR})
  # If building in tree, we honor the same install suffix LLVM uses.
  set(LIBOMPTARGET_LIBDIR_SUFFIX ${LLVM_LIBDIR_SUFFIX})
endif()

Note that we have LIBOMP_ENABLE_WERROR, that should be LIBOMPTARGET_ENABLE_WERROR.

97–98

Right, following what we selected before we can use lib${LIBOMPTARGET_LIBDIR_SUFFIX} instead of just lib.

libomptarget/README.txt

61

Good catch. The library is compatible with what is https://github.com/clang-ykt, we should probably use that instead. Unlike http://clang-omp.github.io/, it contains the latest Sema/CodeGen reimplementation that has been happening in the trunk.

Also, as Jonas mentioned, some section naming is different and the offload entries descriptors contain more information.

63

This is only true with the change in https://reviews.llvm.org/D28298 we need to make sure that is committed first.

libomptarget/test/CMakeLists.txt

71–77

1. Ok, we should probably replace this by a warning instead of using an error.

2. It should be possible. The reason of the error/warning is that it is untested. We try not to do anything that precludes MSVC builds - we also try to pave the way for that to happen by preparing some MSVC options as currently used by other components of LLVM and libomp. However, the goal of our contribution is to add support for linux machines which are the machines we and our users have access to. This is true for this library but is also true for the compiler support - it will need tuning/features to fully work on Windows. We expect someone with interest in having support for Windows to complete the work by contributing the features and testing infrastructure for that.

[grokos]

Addressed more comments, awaiting new review.

libomptarget/Build_With_CMake.txt

103–116

We only have one instruction file for the whole library. After all, cmake is invoked once from libomptarget's root directory, all -D definitions are passed to the "root" cmake.

I am happy to revise this scheme if you think it makes more sense to split the instructions, although I strongly prefer the current implementation.

libomptarget/CMakeLists.txt

97–98

OK, I incorporated the changes in the new diff. Thanks for the suggestion!

107–109

OK, I'll make them part of the corresponding patches for plugins and RTLs.

libomptarget/README.txt

61

I updated the documentation in the new diff.

libomptarget/src/omptarget.cpp

123–131

It's not efficient to reuse the assignment operator. I've done this in the initializer list instead.

256–258

Done in new diff.

284

Our buildbot already uses this env var name, but I am happy to change it to whatever name you propose. However, I don't see any problem with OMP_, especially if there is any chance that future revisions of the standard introduce such an env var.

422–423

Done.

495

Right, thanks for pointing it out!

509–512

No, omp_target_is_present checks whether some host address has corresponding storage on the device, it's got nothing to do with memory allocated via omp_target_alloc. So, in order to check whether a host address has been mapped we just use the device's HostDataToTargetMap.

514

It's not necessary. omp_target_is_present only accesses the device's HostDataToTargetMap to determine whether an address has been mapped. If the device is not initialized, HostDataToTargetMap has zero contents and omp_target_is_present correctly returns false. The function does not make any call to RTL functions (like data_alloc or data_free), so even if the device has not been initialized there is no problem.

533–534

Done.

539–540

Done.

911

I've changed most instances where long served as a bool, i.e. variables IsNew, IsLast, Forcedelete, IsImplicit etc.

I also realized that long was used for array sizes, whereas the correct datatype should be int64_t (that's the datatype used by clang to communicate with the interface functions). I fixed that too.

1107

Thanks for spotting this.

1113–1137

Done. I refactored the code into a new static function which is called right after registering the image into the translation table.

1115

Correct, fixed in new diff.

1178

Fixed.

1194–1206

Correct.

1355–1356

Fixed.

1358–1378

Correct, I've reused the former function in the new diff.

1384–1385

Done.

1394–1396

Backwards compatible with internal, unreleased versions of clang. Oh well....

2117

Thanks for spotting this, obviously if offload fails at some point there is no need to continue with the loop.

2132

Done.

libomptarget/src/omptarget.h

87

You're right, some variable names caused confusion, I've changed them :)

89–90

Done.

167

Right, it's only data_end.

libomptarget/test/CMakeLists.txt

71–77

OK, I replaced the fatal error with a warning.

[Hahnfeld]

Only a minor comment for the code. And I think you may have missed some changes for the uploaded diff?

libomptarget/Build_With_CMake.txt

103–116

I just meant that these options should be moved a few lines below so that they are listed under the secion NVPTX device RTL specific in this document

libomptarget/CMakeLists.txt

39

LIBOMP_ENABLE_WERROR is not yet fixed

97–98

I don't see the suffix here, maybe you forgot to add some changes to this file?

libomptarget/README.txt

61

I don't see that neither

libomptarget/src/omptarget.cpp

284

Ok, let's keep it like that. But if that ever happens, we can't ensure any backwards compatibility...

514

Alright. I think we should then check for device_num < Devices.size() because operator [] will not perform any bounds check. Just in case the user supplies an invalid parameter?

1376–1377

I think this has not yet been collapsed...

1394–1396

Hmm, when is that planned for removal then?

libomptarget/test/CMakeLists.txt

71–77

I still see libomptarget_error_say here...

[grokos]

Jonas, thanks for the comments once again. It seems I had forgotten to add the CMakeLists files to the git commit. Here is the latest revision of the patch.

libomptarget/Build_With_CMake.txt

103–116

OK, I moved them.

libomptarget/src/omptarget.cpp

514

Good catch! I added a check.

1376–1377

Done.

1394–1396

It will be removed once clang starts using internally the new map types. As far as I know, we haven't started any work on that yet. The async interface is our next priority, upgrading clang to use the new map types will come afterwards.

[grokos]

Took care of the last few comments.

[Hahnfeld]

LGTM, thanks for your patience!

CMakeLists.txt

 cmake_minimum_required(VERSION 2.8 FATAL_ERROR)
 add_subdirectory(runtime)
+add_subdirectory(libomptarget)
 No newline at end of file

libomptarget/Build_With_CMake.txt

+#
+#//===----------------------------------------------------------------------===//
+#//
+#//                     The LLVM Compiler Infrastructure
+#//
+#// This file is dual licensed under the MIT and the University of Illinois Open
+#// Source Licenses. See LICENSE.txt for details.
+#//
+#//===----------------------------------------------------------------------===//
+#
+
+=====================================================================
+How to Build the LLVM* OpenMP* Offloading Runtime Library using CMake
+=====================================================================
+
+==== Version of CMake required: v2.8.0 or above ====
+ 
+============================================
+How to call cmake initially, then repeatedly
+============================================
+- When calling cmake for the first time, all needed compiler options
+  must be specified on the command line.  After this initial call to
+  cmake, the compiler definitions must not be included for further calls
+  to cmake.  Other options can be specified on the command line multiple
+  times including all definitions in the Build options section below.
+- Example of configuring, building, reconfiguring, rebuilding:
+  $ mkdir build
+  $ cd build
+  $ cmake -DCMAKE_C_COMPILER=gcc -DCMAKE_CXX_COMPILER=g++ ..  # Initial configuration
+  $ make
+  ...
+  $ make clean
+  $ cmake -DCMAKE_BUILD_TYPE=Debug ..                         # Second configuration
+  $ make
+  ...
+  $ rm -rf *
+  $ cmake -DCMAKE_C_COMPILER=gcc -DCMAKE_CXX_COMPILER=g++ ..  # Third configuration
+  $ make
+- Notice in the example how the compiler definitions are only specified
+  for an empty build directory, but other Build options are used at any time.
+- The file CMakeCache.txt which is created after the first call to cmake
+  is a configuration file which holds all the values for the Build options.
+  These configuration values can be changed using a text editor to modify 
+  CMakeCache.txt as opposed to using definitions on the command line.
+- To have cmake create a particular type of build generator file simply 
+  inlude the -G <Generator name> option:
+  $ cmake -G "Unix Makefiles" ...
+  You can see a list of generators cmake supports by executing cmake with
+  no arguments and a list will be printed.
+
+=====================
+Instructions to Build
+=====================
+ $ cd libomptarget_top_level/ [ directory with plugins/ , deviceRTLs/ , etc. ]
+ $ mkdir build
+ $ cd build
+
+ [ Unix* Libraries ]
+ $ cmake -DCMAKE_C_COMPILER=<C Compiler> -DCMAKE_CXX_COMPILER=<C++ Compiler> ..
+
+ $ make
+ $ make install
+
+===========
+Tests
+===========
+After the library has been built, there are optional tests that can be 
+performed.  Some will be skipped based upon the platform.
+To run the tests,
+$ make check-libomptarget
+
+=============
+CMake options
+=============
+-DCMAKE_C_COMPILER=<C compiler name>
+Specify the C compiler
+
+-DCMAKE_CXX_COMPILER=<C++ compiler name>
+Specify the C++ compiler
+
+==== First values listed are the default value ====
+-DCMAKE_BUILD_TYPE=Release|Debug|RelWithDebInfo
+Build type can be Release, Debug, or RelWithDebInfo.
+
+-DLIBOMPTARGET_ENABLE_WERROR=true|false
+Should consider warnings as errors.
+
+-DLIBOMPTARGET_LLVM_LIT_EXECUTABLE=""
+Full path to the llvm-lit tool. Required for testing in out-of-tree builds.
+
+-DLIBOMPTARGET_FILECHECK_EXECUTABLE=""
+Full path to the FileCheck tool. Required for testing in out-of-tree builds.
+
+-DLIBOMPTARGET_OPENMP_HEADER_FOLDER=""
+Path of the folder that contains omp.h. This is required for testing 
+out-of-tree builds.
+
+-DLIBOMPTARGET_OPENMP_HOST_RTL_FOLDER=""
+Path of the folder that contains libomp.so. This is required for testing 
+out-of-tree builds.
+
+==== NVPTX device RTL specific ====
+-DLIBOMPTARGET_NVPTX_ALTERNATE_HOST_COMPILER=""
+Host compiler to use with NVCC. This compiler is not going to be used to produce
+any binary. Instead, this is used to overcome the input compiler checks done by
+NVCC. E.g. if using a default host compiler that is not compatible with NVCC,
+this option can be use to pass to NVCC a valid compiler to avoid the error. 
+
+-DLIBOMPTARGET_NVPTX_COMPUTE_CAPABILITY="35"
+Comma-separated list of CUDA compute capabilities that should be supported by
+the NVPTX device RTL. E.g. for compute capabilities 3.0 and 3.5, the option 
+"30,35" should be used.
+
+=======================
+Example usages of CMake
+=======================
+---- Typical usage ----
+cmake -DCMAKE_C_COMPILER=gcc -DCMAKE_CXX_COMPILER=g++ ..
+cmake -DCMAKE_C_COMPILER=clang -DCMAKE_CXX_COMPILER=clang++ ..
+
+---- Request an NVPTX runtime library that supports compute capability 5.0 ----
+cmake -DCMAKE_C_COMPILER=clang -DCMAKE_CXX_COMPILER=clang++ -DLIBOMPTARGET_NVPTX_COMPUTE_CAPABILITY="50"
+
+=========
+Footnotes
+=========
+[*] Other names and brands may be claimed as the property of others.

libomptarget/CMakeLists.txt

+##===----------------------------------------------------------------------===##
+# 
+#                     The LLVM Compiler Infrastructure
+#
+# This file is dual licensed under the MIT and the University of Illinois Open
+# Source Licenses. See LICENSE.txt for details.
+# 
+##===----------------------------------------------------------------------===##
+#
+# Build offloading library libomptarget.so.
+#
+##===----------------------------------------------------------------------===##
+
+# CMAKE libomptarget
+cmake_minimum_required(VERSION 2.8 FATAL_ERROR)
+
+# Add cmake directory to search for custom cmake functions.
+set(CMAKE_MODULE_PATH ${CMAKE_CURRENT_SOURCE_DIR}/cmake/Modules ${CMAKE_MODULE_PATH})
+
+# Standalone build or part of LLVM?
+set(LIBOMPTARGET_STANDALONE_BUILD FALSE)
+if("${CMAKE_SOURCE_DIR}" STREQUAL "${CMAKE_CURRENT_SOURCE_DIR}" OR
+   "${CMAKE_SOURCE_DIR}/libomptarget" STREQUAL "${CMAKE_CURRENT_SOURCE_DIR}")
+  project(libomptarget C CXX)
+  set(LIBOMPTARGET_STANDALONE_BUILD TRUE)
+endif()
+  
+  
+if(${LIBOMPTARGET_STANDALONE_BUILD})
+  set(LIBOMPTARGET_ENABLE_WERROR FALSE CACHE BOOL
+    "Enable -Werror flags to turn warnings into errors for supporting compilers.")
+  # CMAKE_BUILD_TYPE was not defined, set default to Release
+  if(NOT CMAKE_BUILD_TYPE)
+    set(CMAKE_BUILD_TYPE Release)
+  endif()
+else()
+  set(LIBOMP_ENABLE_WERROR ${LLVM_ENABLE_WERROR})
+endif()
+
+# Compiler flag checks.
+include(config-ix)
+
+# Message utilities.
+include(LibomptargetUtils)
+
+# Get dependencies for the different components of the project.
+include(LibomptargetGetDependencies)
+
+# This is a list of all the targets that are supported/tested right now.
+set (LIBOMPTARGET_ALL_TARGETS "${LIBOMPTARGET_ALL_TARGETS} powerpc64le-ibm-linux-gnu")
+set (LIBOMPTARGET_ALL_TARGETS "${LIBOMPTARGET_ALL_TARGETS} powerpc64-ibm-linux-gnu")
+set (LIBOMPTARGET_ALL_TARGETS "${LIBOMPTARGET_ALL_TARGETS} x86_64-pc-linux-gnu")
+set (LIBOMPTARGET_ALL_TARGETS "${LIBOMPTARGET_ALL_TARGETS} nvptx64-nvidia-cuda")
+
+# Once the plugins for the different targets are validated, they will be added to
+# the list of supported targets in the current system.
+set (LIBOMPTARGET_SYSTEM_TARGETS "")
+
+# Set base directories - required for lit to locate the tests.
+set(LIBOMPTARGET_BASE_DIR ${CMAKE_CURRENT_SOURCE_DIR})
+set(LIBOMPTARGET_BINARY_DIR ${CMAKE_CURRENT_BINARY_DIR})
+
+# We need C++11 support.
+if(LIBOMPTARGET_HAVE_STD_CPP11_FLAG)
+
+  # We can only build the offloading library if libelf is available.
+  if(LIBOMPTARGET_DEP_LIBELF_FOUND)
+  
+    libomptarget_say("Building offloading runtime library libomptarget.")
+  
+    # Add libelf include directories.
+    include_directories(${LIBOMPTARGET_DEP_LIBELF_INCLUDE_DIRS})
+  
+    # Enable support for C++11.
+    add_definitions(-std=c++11)
+    
+    if(LIBOMPTARGET_ENABLE_WERROR AND LIBOMPTARGET_HAVE_WERROR_FLAG)
+      add_definitions(-Werror)
+    endif()
+    
+    # If building this library in debug mode, we define a macro to enable
+    # dumping progress messages at runtime.
+    string( TOLOWER "${CMAKE_BUILD_TYPE}" LIBOMPTARGET_CMAKE_BUILD_TYPE)
+    if(LIBOMPTARGET_CMAKE_BUILD_TYPE MATCHES debug)
+      add_definitions(-DOMPTARGET_DEBUG)
+      add_definitions(-g)
+      add_definitions(-O0)
+    endif()
+    
+    set(src_files
+      src/omptarget.cpp
+      src/targets_info.cpp
+    )
+    
+    include_directories(src/)
+    
+    # Build libomptarget library with libelf and libdl dependencies.
+    add_library(omptarget SHARED ${src_files})
+    target_link_libraries(omptarget 
+      ${LIBOMPTARGET_DEP_LIBELF_LIBRARIES} 
+      dl
+      "-Wl,--version-script=${CMAKE_CURRENT_SOURCE_DIR}/exports")
+    
+    # Install libomptarget under the lib destination folder.
+    install(TARGETS omptarget LIBRARY DESTINATION "lib")
+    
+    # Retrieve the path to the resulting library so that it can be used for 
+    # testing.
+    get_target_property(LIBOMPTARGET_LIBRARY_DIR omptarget LIBRARY_OUTPUT_DIRECTORY)
+    if(NOT LIBOMPTARGET_LIBRARY_DIR)
+      set(LIBOMPTARGET_LIBRARY_DIR ${CMAKE_CURRENT_BINARY_DIR})
+    endif()
+    
+    # Add tests.
+    add_subdirectory(test)
+    
+  else(LIBOMPTARGET_DEP_LIBELF_FOUND)
+    libomptarget_say("Not building offloading runtime library libomptarget: libelf dependency wasn't found.")
+  endif(LIBOMPTARGET_DEP_LIBELF_FOUND)
+  
+else(LIBOMPTARGET_HAVE_STD_CPP11_FLAG)
+  libomptarget_say("Not building offloading runtime library libomptarget: host compiler must have c++11 support.")
+endif(LIBOMPTARGET_HAVE_STD_CPP11_FLAG)

libomptarget/README.txt

+
+    README for the LLVM* OpenMP* Offloading Runtime Library (libomptarget)
+    ======================================================================
+
+How to Build the LLVM* OpenMP* Offloading Runtime Library (libomptarget)
+========================================================================
+In-tree build:
+
+$ cd where-you-want-to-live
+Check out openmp (libomptarget lives under ./libomptarget) into llvm/projects
+$ cd where-you-want-to-build
+$ mkdir build && cd build
+$ cmake path/to/llvm -DCMAKE_C_COMPILER=<C compiler> -DCMAKE_CXX_COMPILER=<C++ compiler>
+$ make omptarget
+
+Out-of-tree build:
+
+$ cd where-you-want-to-live
+Check out openmp (libomptarget lives under ./libomptarget)
+$ cd where-you-want-to-live/openmp/libomptarget
+$ mkdir build && cd build
+$ cmake path/to/openmp -DCMAKE_C_COMPILER=<C compiler> -DCMAKE_CXX_COMPILER=<C++ compiler>
+$ make
+
+For details about building, please look at Build_With_CMake.txt
+
+Architectures Supported
+=======================
+The current library has been only tested in Linux operating system and the 
+following host architectures: 
+* Intel(R) 64 architecture
+* IBM(R) Power architecture (big endian)
+* IBM(R) Power architecture (little endian)
+
+The currently supported offloading device architectures are:
+* Intel(R) 64 architecture (generic 64-bit plugin - mostly for testing purposes)
+* IBM(R) Power architecture (big endian) (generic 64-bit plugin - mostly for testing purposes)
+* IBM(R) Power architecture (little endian) (generic 64-bit plugin - mostly for testing purposes)
+* CUDA(R) enabled 64-bit NVIDIA(R) GPU architectures
+
+Supported RTL Build Configurations
+==================================
+Supported Architectures: Intel(R) 64, IBM(R) Power 7 and Power 8
+
+              ---------------------------
+              |   gcc      |   clang    |
+--------------|------------|------------|
+| Linux* OS   |  Yes(1)    |  Yes(2)    |
+-----------------------------------------
+
+(1) gcc version 4.8.2 or later is supported.
+(2) clang version 3.7 or later is supported.
+
+
+Front-end Compilers that work with this RTL
+===========================================
+
+The following compilers are known to do compatible code generation for
+this RTL: 
+  - clang (from the OpenMP development branch at http://clang-omp.github.io/ )
+  - clang (development branch at http://clang.llvm.org - several features still 
+    under development)
+
+-----------------------------------------------------------------------
+
+Notices
+=======
+This library and related compiler support is still under development, so the 
+employed interface is likely to change in the future.
+
+*Other names and brands may be claimed as the property of others.
+

libomptarget/cmake/Modules/LibomptargetGetDependencies.cmake

+#
+#//===----------------------------------------------------------------------===//
+#//
+#//                     The LLVM Compiler Infrastructure
+#//
+#// This file is dual licensed under the MIT and the University of Illinois Open
+#// Source Licenses. See LICENSE.txt for details.
+#//
+#//===----------------------------------------------------------------------===//
+#
+
+# Try to detect in the system several dependencies required by the different
+# components of libomptarget. These are the dependencies we have:
+#
+# libelf : required by some targets to handle the ELF files at runtime.
+# libffi : required to launch target kernels given function and argument 
+#          pointers.
+# CUDA : required to control offloading to NVIDIA GPUs.
+
+include (FindPackageHandleStandardArgs)
+
+################################################################################
+# Looking for libelf...
+################################################################################
+
+find_path (
+  LIBOMPTARGET_DEP_LIBELF_INCLUDE_DIR
+  NAMES
+    libelf.h
+  PATHS
+    /usr/include
+    /usr/local/include
+    /opt/local/include
+    /sw/include
+    ENV CPATH
+  PATH_SUFFIXES
+    libelf)
+
+find_library (
+  LIBOMPTARGET_DEP_LIBELF_LIBRARIES
+  NAMES
+    elf
+  PATHS
+    /usr/lib
+    /usr/local/lib
+    /opt/local/lib
+    /sw/lib
+    ENV LIBRARY_PATH
+    ENV LD_LIBRARY_PATH)
+    
+set(LIBOMPTARGET_DEP_LIBELF_INCLUDE_DIRS ${LIBOMPTARGET_DEP_LIBELF_INCLUDE_DIR})
+find_package_handle_standard_args(
+  LIBOMPTARGET_DEP_LIBELF 
+  DEFAULT_MSG
+  LIBOMPTARGET_DEP_LIBELF_LIBRARIES
+  LIBOMPTARGET_DEP_LIBELF_INCLUDE_DIRS)
+
+mark_as_advanced(
+  LIBOMPTARGET_DEP_LIBELF_INCLUDE_DIRS 
+  LIBOMPTARGET_DEP_LIBELF_LIBRARIES)
+  
+################################################################################
+# Looking for libffi...
+################################################################################
+
+find_path (
+  LIBOMPTARGET_DEP_LIBFFI_INCLUDE_DIR
+  NAMES
+    ffi.h
+  PATHS
+    /usr/include
+    /usr/local/include
+    /opt/local/include
+    /sw/include
+    ENV CPATH
+  PATH_SUFFIXES
+    libffi)
+
+find_library (
+  LIBOMPTARGET_DEP_LIBFFI_LIBRARIES
+  NAMES
+    ffi
+  PATHS
+    /usr/lib
+    /usr/local/lib
+    /opt/local/lib
+    /sw/lib
+    ENV LIBRARY_PATH
+    ENV LD_LIBRARY_PATH)
+
+set(LIBOMPTARGET_DEP_LIBFFI_INCLUDE_DIRS ${LIBOMPTARGET_DEP_LIBFFI_INCLUDE_DIR})
+find_package_handle_standard_args(
+  LIBOMPTARGET_DEP_LIBFFI 
+  DEFAULT_MSG
+  LIBOMPTARGET_DEP_LIBFFI_LIBRARIES
+  LIBOMPTARGET_DEP_LIBFFI_INCLUDE_DIRS)
+
+mark_as_advanced(
+  LIBOMPTARGET_DEP_LIBFFI_INCLUDE_DIRS 
+  LIBOMPTARGET_DEP_LIBFFI_LIBRARIES)
+  
+################################################################################
+# Looking for CUDA...
+################################################################################
+find_package(CUDA QUIET)
+
+set(LIBOMPTARGET_DEP_CUDA_FOUND ${CUDA_FOUND})
+set(LIBOMPTARGET_DEP_CUDA_LIBRARIES ${CUDA_LIBRARIES})
+set(LIBOMPTARGET_DEP_CUDA_INCLUDE_DIRS ${CUDA_INCLUDE_DIRS})
+
+mark_as_advanced(
+  LIBOMPTARGET_DEP_CUDA_FOUND 
+  LIBOMPTARGET_DEP_CUDA_INCLUDE_DIRS
+  LIBOMPTARGET_DEP_CUDA_LIBRARIES)

libomptarget/cmake/Modules/LibomptargetUtils.cmake

+#
+#//===----------------------------------------------------------------------===//
+#//
+#//                     The LLVM Compiler Infrastructure
+#//
+#// This file is dual licensed under the MIT and the University of Illinois Open
+#// Source Licenses. See LICENSE.txt for details.
+#//
+#//===----------------------------------------------------------------------===//
+#
+
+# void libomptarget_say(string message_to_user);
+# - prints out message_to_user
+macro(libomptarget_say message_to_user)
+  message(STATUS "LIBOMPTARGET: ${message_to_user}")
+endmacro()
+
+# void libomptarget_warning_say(string message_to_user);
+# - prints out message_to_user with a warning
+macro(libomptarget_warning_say message_to_user)
+  message(WARNING "LIBOMPTARGET: ${message_to_user}")
+endmacro()
+
+# void libomptarget_error_say(string message_to_user);
+# - prints out message_to_user with an error and exits cmake
+macro(libomptarget_error_say message_to_user)
+  message(FATAL_ERROR "LIBOMPTARGET: ${message_to_user}")
+endmacro()

libomptarget/cmake/Modules/config-ix.cmake

+#
+#//===----------------------------------------------------------------------===//
+#//
+#//                     The LLVM Compiler Infrastructure
+#//
+#// This file is dual licensed under the MIT and the University of Illinois Open
+#// Source Licenses. See LICENSE.txt for details.
+#//
+#//===----------------------------------------------------------------------===//
+#
+
+include(CheckCCompilerFlag)
+include(CheckCXXCompilerFlag)
+
+# Checking C, CXX
+check_cxx_compiler_flag(-std=c++11 LIBOMPTARGET_HAVE_STD_CPP11_FLAG)
+check_c_compiler_flag(-Werror LIBOMPTARGET_HAVE_WERROR_FLAG)

libomptarget/exports

+VERS1.0 {
+  global:
+    __tgt_register_lib;
+    __tgt_unregister_lib;
+    __tgt_target_data_begin;
+    __tgt_target_data_end;
+    __tgt_target_data_update;
+    __tgt_target;
+    __tgt_target_teams;
+    __tgt_target_data_begin_nowait;
+    __tgt_target_data_end_nowait;
+    __tgt_target_data_update_nowait;
+    __tgt_target_nowait;
+    __tgt_target_teams_nowait;
+  local:
+    *;
+};
+

libomptarget/src/omptarget.h

+//===-------- omptarget.h - Target independent OpenMP target RTL -- C++ -*-===//
+//
+//                     The LLVM Compiler Infrastructure
+//
+// This file is dual licensed under the MIT and the University of Illinois Open
+// Source Licenses. See LICENSE.txt for details.
+//
+//===----------------------------------------------------------------------===//
+//
+// Interface to be used by Clang during the codegen of a
+// target region.
+//
+//===----------------------------------------------------------------------===//
+
+#include <stdint.h>
+
+#ifndef _OMPTARGET_H_
+#define _OMPTARGET_H_
+
+#define OFFLOAD_SUCCESS (0)
+#define OFFLOAD_FAIL (~0)
+
+#define OFFLOAD_DEVICE_DEFAULT -1
+#define OFFLOAD_DEVICE_CONSTRUCTOR -2
+#define OFFLOAD_DEVICE_DESTRUCTOR -3
+
+/// Data attributes for each data reference used in an OpenMP target region.
+enum tgt_map_type {
+  tgt_map_alloc = 0x00, // allocate memory in the device for this reference
+  tgt_map_to =
+      0x01, // copy the data to the device but do not update the host memory
+  tgt_map_from =
+      0x02, // copy the data to the host but do not update the device memory
+  tgt_map_always = 0x04,
+  tgt_map_release = 0x08,
+  tgt_map_delete = 0x18,
+  tgt_map_pointer = 0x20,
+  tgt_map_extra = 0x40
+};
+
+/// 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 {
+  void *addr;   // Pointer to the offload entry info (function or global)
+  char *name;   // Name of the function or global
+  int64_t size; // Size of the entry info (0 if it a function)
+};
+
+/// This struct is a record of the device image information
+struct __tgt_device_image {
+  void *ImageStart;                  // Pointer to the target code start
+  void *ImageEnd;                    // Pointer to the target code end
+  __tgt_offload_entry *EntriesBegin; // Begin of the table with all the entries
+  __tgt_offload_entry
+      *EntriesEnd; // End of the table with all the entries (non inclusive)
+};
+
+/// This struct is a record of all the host code that may be offloaded to a
+/// target.
+struct __tgt_bin_desc {
+  int32_t NumDevices;                // Number of devices supported
+  __tgt_device_image *DeviceImages;  // Arrays of device images (one per device)
+  __tgt_offload_entry *EntriesBegin; // Begin of the table with all the entries
+  __tgt_offload_entry
+      *EntriesEnd; // End of the table with all the entries (non inclusive)
+};
+
+/// This struct contains the offload entries identified by the target runtime
+struct __tgt_target_table {
+  __tgt_offload_entry *EntriesBegin; // Begin of the table with all the entries
+  __tgt_offload_entry
+      *EntriesEnd; // End of the table with all the entries (non inclusive)
+};
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+void omp_set_default_device(int device_num);
+int omp_get_default_device(void);
+int omp_get_num_devices(void);
+int omp_is_initial_device(void);
+
+/// adds a target shared library to the target execution image
+void __tgt_register_lib(__tgt_bin_desc *desc);
+
+/// removes a target shared library to the target execution image
+void __tgt_unregister_lib(__tgt_bin_desc *desc);
+
+// creates host to the target data mapping, store it in the
+// libtarget.so internal structure (an entry in a stack of data maps) and
+// passes the data to the device;
+void __tgt_target_data_begin(int32_t device_id, int32_t arg_num,
+                             void **args_base, void **args, int64_t *arg_sizes,
+                             int32_t *arg_types);
+void __tgt_target_data_begin_nowait(int32_t device_id, int32_t arg_num,
+                                    void **args_base, void **args,
+                                    int64_t *arg_sizes, int32_t *arg_types,
+                                    int32_t depNum, void *depList,
+                                    int32_t noAliasDepNum,
+                                    void *noAliasDepList);
+
+// passes data from the target, release target memory and destroys the
+// host-target mapping (top entry from the stack of data maps) created by
+// the last __tgt_target_data_begin
+void __tgt_target_data_end(int32_t device_id, int32_t arg_num, void **args_base,
+                           void **args, int64_t *arg_sizes, int32_t *arg_types);
+void __tgt_target_data_end_nowait(int32_t device_id, int32_t arg_num,
+                                  void **args_base, void **args,
+                                  int64_t *arg_sizes, int32_t *arg_types,
+                                  int32_t depNum, void *depList,
+                                  int32_t noAliasDepNum, void *noAliasDepList);
+
+/// passes data to/from the target
+void __tgt_target_data_update(int32_t device_id, int32_t arg_num,
+                              void **args_base, void **args, int64_t *arg_sizes,
+                              int32_t *arg_types);
+void __tgt_target_data_update_nowait(int32_t device_id, int32_t arg_num,
+                                     void **args_base, void **args,
+                                     int64_t *arg_sizes, int32_t *arg_types,
+                                     int32_t depNum, void *depList,
+                                     int32_t noAliasDepNum,
+                                     void *noAliasDepList);
+
+// performs the same actions as data_begin in case arg_num is non-zero
+// and initiates run of offloaded region on target platform; if arg_num
+// is non-zero after the region execution is done it also performs the
+// same action as data_update and data_end aboveThe following types are
+// used; this function return 0 if it was able to transfer the execution
+// to a target and an int different from zero otherwise
+int __tgt_target(int32_t device_id, void *host_ptr, int32_t arg_num,
+                 void **args_base, void **args, int64_t *arg_sizes,
+                 int32_t *arg_types);
+int __tgt_target_nowait(int32_t device_id, void *host_ptr, int32_t arg_num,
+                        void **args_base, void **args, int64_t *arg_sizes,
+                        int32_t *arg_types, int32_t depNum, void *depList,
+                        int32_t noAliasDepNum, void *noAliasDepList);
+
+int __tgt_target_teams(int32_t device_id, void *host_ptr, int32_t arg_num,
+                       void **args_base, void **args, int64_t *arg_sizes,
+                       int32_t *arg_types, int32_t num_teams,
+                       int32_t thread_limit);
+int __tgt_target_teams_nowait(int32_t device_id, void *host_ptr,
+                              int32_t arg_num, void **args_base, void **args,
+                              int64_t *arg_sizes, int32_t *arg_types,
+                              int32_t num_teams, int32_t thread_limit,
+                              int32_t depNum, void *depList,
+                              int32_t noAliasDepNum, void *noAliasDepList);
+
+#ifdef __cplusplus
+}
+#endif
+
+#ifdef OMPTARGET_DEBUG
+#define DEBUGP(prefix, ...)                                                    \
+  {                                                                            \
+    fprintf(stderr, "%s --> ", prefix);                                        \
+    fprintf(stderr, __VA_ARGS__);                                              \
+  }
+#else
+#define DEBUGP(prefix, ...)                                                    \
+  {}
+#endif
+
+#ifdef __cplusplus
+#define EXTERN extern "C"
+#else
+#define EXTERN extern
+#endif
+
+#endif // _OMPTARGET_H_

libomptarget/src/omptarget.cpp

+//===------ omptarget.cpp - Target independent OpenMP target RTL -- C++ -*-===//
+//
+//                     The LLVM Compiler Infrastructure
+//
+// This file is dual licensed under the MIT and the University of Illinois Open
+// Source Licenses. See LICENSE.txt for details.
+//
+//===----------------------------------------------------------------------===//
+//
+// Implementation of the interface to be used by Clang during the codegen of a
+// target region.
+//
+//===----------------------------------------------------------------------===//
+
+#include <assert.h>
+#include <dlfcn.h>
+#include <gelf.h>
+#include <libelf.h>
+#include <list>
+#include <map>
+#include <set>
+#include <sstream>
+#include <stdio.h>
+#include <stdlib.h>
+#include <string>
+#include <vector>
+#include <queue>
+#include <mutex>
+
+/* Header file global to this project */
+#include "omptarget.h"
+#include "targets_info.h"
+
+#define DP(...) DEBUGP("Libomptarget", __VA_ARGS__)
+
+extern targets_info_table targets_info;
+
+struct RTLInfoTy;
+
+/// Map between host data and target data
+struct HostDataToTargetTy {
+  long HstPtrBase; // host info
+  long HstPtrBegin;
+  long HstPtrEnd; // non-inclusive
+
+  long TgtPtrBegin; // target info
+  long TgtPtrEnd;   // non-inclusive (aee: not needed???)
+
+  long RefCount;
+
+  HostDataToTargetTy()
+      : HstPtrBase(0), HstPtrBegin(0), HstPtrEnd(0), TgtPtrBegin(0),
+        TgtPtrEnd(0), RefCount(0) {}
+  HostDataToTargetTy(long BP, long B, long E, long TB, long TE)
+      : HstPtrBase(BP), HstPtrBegin(B), HstPtrEnd(E), TgtPtrBegin(TB),
+        TgtPtrEnd(TE), RefCount(1) {}
+};
+
+typedef std::list<HostDataToTargetTy> HostDataToTargetListTy;
+
+struct DeviceTy {
+  int32_t DeviceID;
+  RTLInfoTy *RTL;
+  int32_t RTLDeviceID;
+
+  bool IsInit;
+  HostDataToTargetListTy HostDataToTargetMap;
+  std::list<void *> PendingConstrDestrHostPtrList;
+
+  DeviceTy(RTLInfoTy *RTL)
+      : DeviceID(-1), RTL(RTL), RTLDeviceID(-1), IsInit(false),
+        HostDataToTargetMap(), PendingConstrDestrHostPtrList() {}
+
+  void *getOrAllocTgtPtr(void *HstPtrBegin, void *HstPtrBase, long Size,
+                         long &IsNew, long UpdateRefCount = true);
+  void *getTgtPtrBegin(void *HstPtrBegin, long Size);
+  void *getTgtPtrBegin(void *HstPtrBegin, long Size, long &IsLast,
+                       long UpdateRefCount = true);
+  void deallocTgtPtr(void *TgtPtrBegin, long Size, long ForceDelete);
+
+  // calls to RTL
+  int32_t init();
+  __tgt_target_table *load_binary(void *Img);
+
+  int32_t data_submit(void *TgtPtrBegin, void *HstPtrBegin, int64_t Size);
+  int32_t data_retrieve(void *HstPtrBegin, void *TgtPtrBegin, int64_t Size);
+
+  int32_t run_region(void *TgtEntryPtr, void **TgtVarsPtr, int32_t TgtVarsSize);
+  int32_t run_team_region(void *TgtEntryPtr, void **TgtVarsPtr,
+                          int32_t TgtVarsSize, int32_t NumTeams,
+                          int32_t ThreadLimit);
+};
+
+struct RTLInfoTy {
+  typedef int32_t(device_type_ty)(int32_t);
+  typedef int32_t(number_of_devices_ty)();
+  typedef int32_t(init_device_ty)(int32_t);
+  typedef __tgt_target_table *(load_binary_ty)(int32_t, void *);
+  typedef void *(data_alloc_ty)(int32_t, int64_t);
+  typedef int32_t(data_submit_ty)(int32_t, void *, void *, int64_t);
+  typedef int32_t(data_retrieve_ty)(int32_t, void *, void *, int64_t);
+  typedef int32_t(data_delete_ty)(int32_t, void *);
+  typedef int32_t(run_region_ty)(int32_t, void *, void **, int32_t);
+  typedef int32_t(run_team_region_ty)(int32_t, void *, void **, int32_t,
+                                      int32_t, int32_t);
+
+  int32_t Idx;                     // RTL index, index is the number of devices
+                                   // of other RTLs that were registered before
+  int32_t NumberOfDevices;         // Number of devices this RTL deal with
+  std::vector<DeviceTy *> Devices; // one per device (NumberOfDevices in total)
+
+  void *LibraryHandler;
+  // Functions implemented in the RTL
+  device_type_ty *device_type;
+  number_of_devices_ty *number_of_devices;
+  init_device_ty *init_device;
+  load_binary_ty *load_binary;
+  data_alloc_ty *data_alloc;
+  data_submit_ty *data_submit;
+  data_retrieve_ty *data_retrieve;
+  data_delete_ty *data_delete;
+  run_region_ty *run_region;
+  run_team_region_ty *run_team_region;
+};
+
+/// RTLs identified in the system
+typedef std::list<RTLInfoTy> RTLsTy;
+static RTLsTy RTLs;
+
+/// Map between device type (elf id) and RTL
+typedef std::map<uint16_t, RTLInfoTy *> DTypeToRTLMapTy;
+static DTypeToRTLMapTy DTypeToRTLMap;
+
+/// Map between Device ID (i.e. openmp device id) and its DeviceTy
+typedef std::vector<DeviceTy> DevicesTy;
+static DevicesTy Devices;
+
+/// Map between the host entry begin and the translation table. Each
+/// registered library get one TranslationTable. Use the map from
+/// __tgt_offload_entry so that we may quickly determine if we are
+/// trying to re=register a existing lib, or we really have a new one
+struct TranslationTable {
+  __tgt_target_table HostTable;
+
+  // Image assigned to a given device
+  std::vector<__tgt_device_image *> TargetsImages; // One image per device ID
+
+  // Table of entry points or NULL if it was not already computed
+  std::vector<__tgt_target_table *> TargetsTable; // One table per device ID
+};
+typedef std::map<__tgt_offload_entry *, TranslationTable>
+    HostEntriesBeginToTransTableTy;
+static HostEntriesBeginToTransTableTy HostEntriesBeginToTransTable;
+
+/// Map between the host ptr and a table index
+struct TableMap {
+  TranslationTable *Table; // table associated with host ptr
+  uint32_t Index; // index in which the host ptr translated entry is found
+  TableMap() : Table(0), Index(0) {}
+  TableMap(TranslationTable *table, uint32_t index)
+      : Table(table), Index(index) {}
+};
+typedef std::map<void *, TableMap> HostPtrToTableMapTy;
+static HostPtrToTableMapTy HostPtrToTableMap;
+
+////////////////////////////////////////////////////////////////////////////////
+// getter and setter
+//
+// aee: non compliant. This need to be integrated in KMPC; can keep it
+// here for the moment
+
+static int DefaultDevice = 0;
+
+// aee non-compliant
+void omp_set_default_device(int device_num) { DefaultDevice = device_num; }
+
+// aee non-compliant
+int omp_get_default_device(void) { return DefaultDevice; }
+
+int omp_get_num_devices(void) { return Devices.size(); }
+
+int omp_is_initial_device(void) { return true; }
+
+////////////////////////////////////////////////////////////////////////////////
+// functionality for device
+
+// return the target pointer begin (where the data will be moved)
+void *DeviceTy::getTgtPtrBegin(void *HstPtrBegin, long Size) {
+  long IsLast;
+  return getTgtPtrBegin(HstPtrBegin, Size, IsLast, false);
+}
+void *DeviceTy::getTgtPtrBegin(void *HstPtrBegin, long Size, long &IsLast,
+                               long UpdateRefCount) {
+  long hp = (long)HstPtrBegin;
+  IsLast = false;
+
+  for (HostDataToTargetListTy::iterator ii = HostDataToTargetMap.begin(),
+                                        ie = HostDataToTargetMap.end();
+       ii != ie; ++ii) {
+    HostDataToTargetTy &HT = *ii;
+    if (hp >= HT.HstPtrBegin && hp < HT.HstPtrEnd) {
+      if ((hp + Size) > HT.HstPtrEnd) {
+        DP("WARNING: Array contain pointer but does not contain the complete "
+           "section\n");
+      }
+
+      IsLast = !(HT.RefCount > 1);
+
+      if (HT.RefCount > 1 && UpdateRefCount)
+        --HT.RefCount;
+
+      long tp = HT.TgtPtrBegin + (hp - HT.HstPtrBegin);
+      return (void *)tp;
+    }
+  }
+
+  return NULL;
+}
+
+// return the target pointer begin (where the data will be moved)
+void *DeviceTy::getOrAllocTgtPtr(void *HstPtrBegin, void *HstPtrBase, long Size,
+                                 long &IsNew, long UpdateRefCount) {
+  long hp = (long)HstPtrBegin;
+  IsNew = false;
+
+  // Check if the pointer is contained
+  for (HostDataToTargetListTy::iterator ii = HostDataToTargetMap.begin(),
+                                        ie = HostDataToTargetMap.end();
+       ii != ie; ++ii) {
+    HostDataToTargetTy &HT = *ii;
+
+    // It is contained?
+    if (hp >= HT.HstPtrBegin && hp < HT.HstPtrEnd) {
+      if ((hp + Size) > HT.HstPtrEnd) {
+        DP("WARNING: Array contain pointer but does not contain the complete "
+           "section\n");
+      }
+      if (UpdateRefCount)
+        ++HT.RefCount;
+      long tp = HT.TgtPtrBegin + (hp - HT.HstPtrBegin);
+      return (void *)tp;
+    }
+  }
+
+  // It is not contained we should create a new entry for it
+  IsNew = true;
+  long tp = (long)RTL->data_alloc(RTLDeviceID, Size);
+  HostDataToTargetMap.push_front(
+      HostDataToTargetTy((long)HstPtrBase, hp, hp + Size, tp, tp + Size));
+  return (void *)tp;
+}
+
+void DeviceTy::deallocTgtPtr(void *HstPtrBegin, long Size, long ForceDelete) {
+  long hp = (long)HstPtrBegin;
+
+  // Check if the pointer is contained in any sub-nodes
+  for (HostDataToTargetListTy::iterator ii = HostDataToTargetMap.begin(),
+                                        ie = HostDataToTargetMap.end();
+       ii != ie; ++ii) {
+    HostDataToTargetTy &HT = *ii;
+
+    // It is contained?
+    if (hp >= HT.HstPtrBegin && hp < HT.HstPtrEnd) {
+      if ((hp + Size) > HT.HstPtrEnd) {
+        DP("WARNING: Array contain pointer but does not contain the complete "
+           "section\n");
+      }
+      if (ForceDelete)
+        HT.RefCount = 1;
+      if (--HT.RefCount <= 0) {
+        assert(HT.RefCount == 0 && "did not expect a negative ref count");
+        DP("Deleting tgt data 0x%016llx of size %lld\n",
+           (long long)HT.TgtPtrBegin, (long long)Size);
+        RTL->data_delete(RTLDeviceID, (void *)HT.TgtPtrBegin);
+        HostDataToTargetMap.erase(ii);
+      }
+      return;
+    }
+  }
+  DP("Section to delete (hst addr 0x%llx) does not exist in the allocated "
+     "memory\n",
+     (unsigned long long)hp);
+}
+
+// init device
+int32_t DeviceTy::init() {
+  int32_t rc = RTL->init_device(RTLDeviceID);
+  if (rc == OFFLOAD_SUCCESS) {
+    IsInit = true;
+  }
+  return rc;
+}
+
+// load binary to device
+__tgt_target_table *DeviceTy::load_binary(void *Img) {
+  return RTL->load_binary(RTLDeviceID, Img);
+}
+
+// submit data to device
+int32_t DeviceTy::data_submit(void *TgtPtrBegin, void *HstPtrBegin,
+                              int64_t Size) {
+  return RTL->data_submit(RTLDeviceID, TgtPtrBegin, HstPtrBegin, Size);
+}
+
+// retrieve data from device
+int32_t DeviceTy::data_retrieve(void *HstPtrBegin, void *TgtPtrBegin,
+                                int64_t Size) {
+  return RTL->data_retrieve(RTLDeviceID, HstPtrBegin, TgtPtrBegin, Size);
+}
+
+// run region on device
+int32_t DeviceTy::run_region(void *TgtEntryPtr, void **TgtVarsPtr,
+                             int32_t TgtVarsSize) {
+  return RTL->run_region(RTLDeviceID, TgtEntryPtr, TgtVarsPtr, TgtVarsSize);
+}
+
+// run team region on device
+int32_t DeviceTy::run_team_region(void *TgtEntryPtr, void **TgtVarsPtr,
+                                  int32_t TgtVarsSize, int32_t NumTeams,
+                                  int32_t ThreadLimit) {
+  return RTL->run_team_region(RTLDeviceID, TgtEntryPtr, TgtVarsPtr, TgtVarsSize,
+                              NumTeams, ThreadLimit);
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// functionality for registering libs
+
+static void RegisterImageIntoTranslationTable(TranslationTable &TT,
+                                              RTLInfoTy &RTL,
+                                              __tgt_device_image *image) {
+
+  // sane size, as when we increase the one, we also increase the other
+  assert(TT.TargetsTable.size() == TT.TargetsImages.size() &&
+         "We should have as many images as we have tables!");
+
+  // Resize the Targets Table and Images to accommodate the new targets if
+  // required
+  unsigned TargetsTableMinimumSize = RTL.Idx + RTL.NumberOfDevices;
+
+  if (TT.TargetsTable.size() < TargetsTableMinimumSize) {
+    TT.TargetsImages.resize(TargetsTableMinimumSize, 0);
+    TT.TargetsTable.resize(TargetsTableMinimumSize, 0);
+  }
+
+  // Register the image in all devices for this target type
+  for (int32_t i = 0; i < RTL.NumberOfDevices; ++i) {
+    // If we are changing the image we are also invalidating the target table
+    if (TT.TargetsImages[RTL.Idx + i] != image) {
+      TT.TargetsImages[RTL.Idx + i] = image;
+      TT.TargetsTable[RTL.Idx + i] = 0; // lazy initialization of target table
+    }
+  }
+}
+
+////////////////////////////////////////////////////////////////////////////////
+/// adds a target shared library to the target execution image
+EXTERN void __tgt_register_lib(__tgt_bin_desc *desc) {
+
+  // Is the library version incompatible with the header file?
+  if (elf_version(EV_CURRENT) == EV_NONE) {
+    DP("Incompatible ELF library!\n");
+    return;
+  }
+
+  // Can't read the host entries??
+  if (!desc) {
+    DP("Invalid descr!\n");
+    return;
+  }
+  if (!desc->EntriesBegin) {
+    DP("Invalid Host entries!\n");
+    return;
+  }
+
+  DP("Registering entries starting at %016lx...\n",
+     (Elf64_Addr)desc->EntriesBegin->addr);
+
+  HostEntriesBeginToTransTableTy::iterator TransTableIt =
+      HostEntriesBeginToTransTable.find(desc->EntriesBegin);
+
+  if (TransTableIt != HostEntriesBeginToTransTable.end()) {
+    DP("Already registered!\n");
+    return;
+  }
+
+  // Initialize translation table for this
+  TranslationTable &TransTable =
+      HostEntriesBeginToTransTable[desc->EntriesBegin];
+  TransTable.HostTable.EntriesBegin = desc->EntriesBegin;
+  TransTable.HostTable.EntriesEnd = desc->EntriesEnd;
+
+  // Scan all the device images
+  for (int32_t i = 0; i < desc->NumDevices; ++i) {
+
+    __tgt_device_image &img = desc->DeviceImages[i];
+    char *img_begin = (char *)img.ImageStart;
+    char *img_end = (char *)img.ImageEnd;
+    size_t img_size = img_end - img_begin;
+
+    // Obtain elf handler
+    Elf *e = elf_memory(img_begin, img_size);
+    if (!e) {
+      DP("Unable to get ELF handle: %s!\n", elf_errmsg(-1));
+      continue;
+    }
+
+    // Check if ELF is the right kind
+    if (elf_kind(e) != ELF_K_ELF) {
+      DP("Unexpected ELF type!\n");
+      continue;
+    }
+    Elf64_Ehdr *eh64 = elf64_getehdr(e);
+    Elf32_Ehdr *eh32 = elf32_getehdr(e);
+    uint16_t MachineID;
+    if (eh64 && !eh32)
+      MachineID = eh64->e_machine;
+    else if (eh32 && !eh64)
+      MachineID = eh32->e_machine;
+    else {
+      DP("Ambiguous ELF header!\n");
+      continue;
+    }
+
+    DTypeToRTLMapTy::iterator RTLHandler = DTypeToRTLMap.find(MachineID);
+
+    // We already have an handler for this device's RTL?
+    // If so insert the Image <-> RTL Map entry and continue
+    if (RTLHandler != DTypeToRTLMap.end()) {
+
+      // We were unable to find a runtime for this device before...
+      if (RTLHandler->second == 0)
+        continue;
+
+      RegisterImageIntoTranslationTable(TransTable, *RTLHandler->second, &img);
+      continue;
+    }
+
+    // Locate the library name and create an handler
+    void *dynlib_handle = 0;
+
+    DP("Looking for RTL libraries: with machine id %d\n", MachineID);
+    for (int32_t j = 0; j < targets_info.Number_of_Entries; ++j) {
+      if (targets_info.Entries[j].Machine_Elf_ID == MachineID) {
+        DP("Found library %s!\n", targets_info.Entries[j].Machine_RTL_Lib);
+
+        dynlib_handle =
+            dlopen(targets_info.Entries[j].Machine_RTL_Lib, RTLD_NOW);
+
+        if (!dynlib_handle)
+          DP("Error opening library for target %d: %s!\n", MachineID,
+             dlerror());
+
+        break;
+      } else
+        DP("Machine ID %d != %d!\n", targets_info.Entries[j].Machine_Elf_ID,
+           MachineID);
+    }
+
+    RTLInfoTy *&RTL = DTypeToRTLMap[MachineID];
+    RTL = 0;
+
+    if (!dynlib_handle) {
+      DP("No valid library for target %d, %s!\n", MachineID, dlerror());
+      continue;
+    }
+
+    // Create the handler and try to retrieve the functions
+    RTLInfoTy R;
+
+    R.LibraryHandler = dynlib_handle;
+    if (!(R.device_type = (RTLInfoTy::device_type_ty *)dlsym(
+              dynlib_handle, "__tgt_rtl_device_type")))
+      continue;
+    if (!(R.number_of_devices = (RTLInfoTy::number_of_devices_ty *)dlsym(
+              dynlib_handle, "__tgt_rtl_number_of_devices")))
+      continue;
+    if (!(R.init_device = (RTLInfoTy::init_device_ty *)dlsym(
+              dynlib_handle, "__tgt_rtl_init_device")))
+      continue;
+    if (!(R.load_binary = (RTLInfoTy::load_binary_ty *)dlsym(
+              dynlib_handle, "__tgt_rtl_load_binary")))
+      continue;
+    if (!(R.data_alloc = (RTLInfoTy::data_alloc_ty *)dlsym(
+              dynlib_handle, "__tgt_rtl_data_alloc")))
+      continue;
+    if (!(R.data_submit = (RTLInfoTy::data_submit_ty *)dlsym(
+              dynlib_handle, "__tgt_rtl_data_submit")))
+      continue;
+    if (!(R.data_retrieve = (RTLInfoTy::data_retrieve_ty *)dlsym(
+              dynlib_handle, "__tgt_rtl_data_retrieve")))
+      continue;
+    if (!(R.data_delete = (RTLInfoTy::data_delete_ty *)dlsym(
+              dynlib_handle, "__tgt_rtl_data_delete")))
+      continue;
+    if (!(R.run_region = (RTLInfoTy::run_region_ty *)dlsym(
+              dynlib_handle, "__tgt_rtl_run_target_region")))
+      continue;
+    if (!(R.run_team_region = (RTLInfoTy::run_team_region_ty *)dlsym(
+              dynlib_handle, "__tgt_rtl_run_target_team_region")))
+      continue;
+
+    // No devices are supported by this RTL?
+    if (!(R.NumberOfDevices = R.number_of_devices())) {
+      DP("No device support this RTL\n");
+      continue;
+    }
+
+    R.Idx = (RTLs.empty()) ? 0 : RTLs.back().Idx + RTLs.back().NumberOfDevices;
+
+    DP("Registered RTL supporting %d devices at index %d!\n", R.NumberOfDevices,
+       R.Idx);
+
+    // The RTL is valid!
+    RTLs.push_back(R);
+    RTL = &RTLs.back();
+
+    // Add devices
+    DeviceTy device(RTL);
+
+    size_t start = Devices.size();
+    Devices.resize(start + RTL->NumberOfDevices, device);
+    for (int32_t device_id = 0; device_id < RTL->NumberOfDevices; device_id++) {
+      // global device ID
+      Devices[start + device_id].DeviceID = start + device_id;
+      // RTL local device ID
+      Devices[start + device_id].RTLDeviceID = device_id;
+
+      // Save pointer to device in RTL in case we want to unregister the RTL
+      RTL->Devices.push_back(&Devices[start + device_id]);
+    }
+
+    DP("Registering image %016lx with RTL!\n", (Elf64_Addr)img_begin);
+    RegisterImageIntoTranslationTable(TransTable, *RTL, &img);
+  }
+
+  DP("Done register entries!\n");
+}
+
+////////////////////////////////////////////////////////////////////////////////
+/// unloads a target shared library
+EXTERN void __tgt_unregister_lib(__tgt_bin_desc *desc) {
+  DP("Unloading target library!\n");
+  return;
+}
+
+/// Internal function to do the mapping and transfer the data to the device
+static void target_data_begin(DeviceTy &Device, int32_t arg_num,
+                              void **args_base, void **args, int64_t *arg_sizes,
+                              int32_t *arg_types) {
+  // process each input
+  for (int32_t i = 0; i < arg_num; ++i) {
+    void *HstPtrBegin = args[i];
+    void *HstPtrBase = args_base[i];
+    void *Pointer_TgtPtrBegin;
+    long IsNew, Pointer_IsNew;
+    if (arg_types[i] & tgt_map_pointer) {
+      DP("has a pointer entry: \n");
+      // base is address of poiner
+      Pointer_TgtPtrBegin = Device.getOrAllocTgtPtr(
+          HstPtrBase, HstPtrBase, sizeof(void *), Pointer_IsNew);
+      DP("There are %ld bytes allocated at target address %016lx\n",
+         (long)sizeof(void *), (long)Pointer_TgtPtrBegin);
+      assert(Pointer_TgtPtrBegin &&
+             "Data allocation by RTL returned invalid ptr");
+      // modify current entry
+      HstPtrBase = *(void **)HstPtrBase;
+    }
+
+    void *TgtPtrBegin =
+        Device.getOrAllocTgtPtr(HstPtrBegin, HstPtrBase, arg_sizes[i], IsNew);
+    DP("There are %ld bytes allocated at target address %016lx - is new %ld\n",
+       (long)arg_sizes[i], (long)TgtPtrBegin, IsNew);
+    assert(TgtPtrBegin && "Data allocation by RTL returned invalid ptr");
+
+    if ((arg_types[i] & tgt_map_to) &&
+        (IsNew || (arg_types[i] & tgt_map_always))) {
+      DP("Moving %ld bytes (hst:%016lx) -> (tgt:%016lx)\n", (long)arg_sizes[i],
+         (long)HstPtrBegin, (long)TgtPtrBegin);
+      Device.data_submit(TgtPtrBegin, HstPtrBegin, arg_sizes[i]);
+    }
+
+    if (arg_types[i] & tgt_map_pointer /* && Pointer_IsNew */) {
+      DP("Update pointer (%016lx) -> [%016lx]\n", (long)Pointer_TgtPtrBegin,
+         (long)TgtPtrBegin);
+      uint64_t Delta = (uint64_t)HstPtrBegin - (uint64_t)HstPtrBase;
+      void *TgrPtrBase_Value = (void *)((uint64_t)TgtPtrBegin - Delta);
+      Device.data_submit(Pointer_TgtPtrBegin, &TgrPtrBase_Value,
+                         sizeof(void *));
+    }
+  }
+}
+
+EXTERN void __tgt_target_data_begin_nowait(int32_t device_id, int32_t arg_num,
+                                           void **args_base, void **args,
+                                           int64_t *arg_sizes,
+                                           int32_t *arg_types, int32_t depNum,
+                                           void *depList, int32_t noAliasDepNum,
+                                           void *noAliasDepList) {
+  __tgt_target_data_begin(device_id, arg_num, args_base, args, arg_sizes,
+                          arg_types);
+}
+
+/// creates host to the target data mapping, store it in the
+/// libtarget.so internal structure (an entry in a stack of data maps)
+/// and passes the data to the device;
+EXTERN void __tgt_target_data_begin(int32_t device_id, int32_t arg_num,
+                                    void **args_base, void **args,
+                                    int64_t *arg_sizes, int32_t *arg_types) {
+  DP("Entering data begin region for device %d with %d mappings\n", device_id,
+     arg_num);
+
+  // No devices available?
+  if (device_id == OFFLOAD_DEVICE_DEFAULT) {
+    device_id = omp_get_default_device();
+    DP("Use default device id %d\n", device_id);
+  }
+  if (Devices.size() <= (size_t)device_id) {
+    DP("Device ID  %d does not have a matching RTL.\n", device_id);
+    return;
+  }
+
+  // Get device info
+  DeviceTy &Device = Devices[device_id];
+  // Init the device if not done before
+  if (!Device.IsInit) {
+    if (Device.init() != OFFLOAD_SUCCESS) {
+      DP("failed to init device %d\n", device_id);
+      return;
+    }
+  }
+
+  target_data_begin(Device, arg_num, args_base, args, arg_sizes, arg_types);
+}
+
+/// Internal function to undo the mapping and retrieve the data from the device
+static void target_data_end(DeviceTy &Device, int32_t arg_num, void **args_base,
+                            void **args, int64_t *arg_sizes,
+                            int32_t *arg_types) {
+  // process each input
+  for (int32_t i = 0; i < arg_num; ++i) {
+    void *HstPtrBegin = args[i];
+    void *HstPtrBase = args_base[i];
+    long IsLast;
+    long ForceDelete = arg_types[i] & tgt_map_delete;
+    if (arg_types[i] & tgt_map_pointer) {
+      // base is pointer begin
+      Device.getTgtPtrBegin(HstPtrBase, sizeof(void *), IsLast);
+      if (IsLast || ForceDelete) {
+        Device.deallocTgtPtr(HstPtrBase, sizeof(void *), ForceDelete);
+      }
+    }
+    void *TgtPtrBegin =
+        Device.getTgtPtrBegin(HstPtrBegin, arg_sizes[i], IsLast);
+
+    DP("There are %ld bytes allocated at target address %016lx - is last %ld\n",
+       (long)arg_sizes[i], (long)TgtPtrBegin, IsLast);
+
+    long Always = arg_types[i] & tgt_map_always;
+    if ((arg_types[i] & tgt_map_from) && (IsLast || ForceDelete || Always)) {
+      DP("Moving %ld bytes (tgt:%016lx) -> (hst:%016lx)\n", (long)arg_sizes[i],
+         (long)TgtPtrBegin, (long)HstPtrBegin);
+      Device.data_retrieve(HstPtrBegin, TgtPtrBegin, arg_sizes[i]);
+    }
+    if (IsLast || ForceDelete) {
+      Device.deallocTgtPtr(HstPtrBegin, arg_sizes[i], ForceDelete);
+    }
+  }
+}
+
+/// passes data from the target, release target memory and destroys
+/// the host-target mapping (top entry from the stack of data maps)
+/// created by the last __tgt_target_data_begin
+EXTERN void __tgt_target_data_end(int32_t device_id, int32_t arg_num,
+                                  void **args_base, void **args,
+                                  int64_t *arg_sizes, int32_t *arg_types) {
+  DP("Entering data end region with %d mappings\n", arg_num);
+
+  // No devices available?
+  if (device_id == OFFLOAD_DEVICE_DEFAULT) {
+    device_id = omp_get_default_device();
+  }
+  if (Devices.size() <= (size_t)device_id) {
+    DP("Device ID  %d does not have a matching RTL.\n", device_id);
+    return;
+  }
+
+  DeviceTy &Device = Devices[device_id];
+  if (!Device.IsInit) {
+    DP("uninit device: ignore");
+    return;
+  }
+
+  target_data_end(Device, arg_num, args_base, args, arg_sizes, arg_types);
+}
+
+EXTERN void __tgt_target_data_end_nowait(int32_t device_id, int32_t arg_num,
+                                         void **args_base, void **args,
+                                         int64_t *arg_sizes, int32_t *arg_types,
+                                         int32_t depNum, void *depList,
+                                         int32_t noAliasDepNum,
+                                         void *noAliasDepList) {
+  __tgt_target_data_end(device_id, arg_num, args_base, args, arg_sizes,
+                        arg_types);
+}
+
+/// passes data to/from the target
+EXTERN void __tgt_target_data_update(int32_t device_id, int32_t arg_num,
+                                     void **args_base, void **args,
+                                     int64_t *arg_sizes, int32_t *arg_types) {
+  DP("Entering data update with %d mappings\n", arg_num);
+
+  // No devices available?
+  if (device_id == OFFLOAD_DEVICE_DEFAULT) {
+    device_id = omp_get_default_device();
+  }
+  if (Devices.size() <= (size_t)device_id) {
+    DP("Device ID  %d does not have a matching RTL.\n", device_id);
+    return;
+  }
+
+  // Get device info
+  DeviceTy &Device = Devices[device_id];
+  if (!Device.IsInit) {
+    DP("uninit device: ignore");
+    return;
+  }
+
+  // process each input
+  for (int32_t i = 0; i < arg_num; ++i) {
+    void *HstPtrBegin = args[i];
+    // void *HstPtrBase = args_base[i];
+    long IsLast;
+    void *TgtPtrBegin =
+        Device.getTgtPtrBegin(HstPtrBegin, arg_sizes[i], IsLast, false);
+    if (arg_types[i] & tgt_map_from) {
+      DP("Moving %ld bytes (tgt:%016lx) -> (hst:%016lx)\n", (long)arg_sizes[i],
+         (long)TgtPtrBegin, (long)HstPtrBegin);
+      Device.data_retrieve(HstPtrBegin, TgtPtrBegin, arg_sizes[i]);
+    }
+    if (arg_types[i] & tgt_map_to) {
+      DP("Moving %ld bytes (hst:%016lx) -> (tgt:%016lx)\n", (long)arg_sizes[i],
+         (long)HstPtrBegin, (long)TgtPtrBegin);
+      Device.data_submit(TgtPtrBegin, HstPtrBegin, arg_sizes[i]);
+    }
+  }
+}
+
+EXTERN void __tgt_target_data_update_nowait(
+    int32_t device_id, int32_t arg_num, void **args_base, void **args,
+    int64_t *arg_sizes, int32_t *arg_types, int32_t depNum, void *depList,
+    int32_t noAliasDepNum, void *noAliasDepList) {
+  __tgt_target_data_update(device_id, arg_num, args_base, args, arg_sizes,
+                           arg_types);
+}
+
+/// performs the same actions as data_begin in case arg_num is
+/// non-zero and initiates run of offloaded region on target platform;
+/// if arg_num is non-zero after the region execution is done it also
+/// performs the same action as data_update and data_end aboveThe
+/// following types are used; this function return 0 if it was able to
+/// transfer the execution to a target and an int different from zero
+/// otherwise
+static int target(int32_t device_id, void *host_ptr, int32_t arg_num,
+                  void **args_base, void **args, int64_t *arg_sizes,
+                  int32_t *arg_types, int32_t team_num, int32_t thread_limit,
+                  int IsTeamConstruct, int IsConstrDestrRecursiveCall) {
+  DP("Entering target region with entry point %016lx and device Id %d\n",
+     (Elf64_Addr)host_ptr, device_id);
+
+  if (device_id == OFFLOAD_DEVICE_DEFAULT) {
+    device_id = omp_get_default_device();
+  }
+  // got a new constructor/destructor?
+  if (device_id == OFFLOAD_DEVICE_CONSTRUCTOR ||
+      device_id == OFFLOAD_DEVICE_DESTRUCTOR) {
+    DP("Got a constructor/destructor\n");
+    for (unsigned D = 0; D < Devices.size(); D++) {
+      DeviceTy &Device = Devices[D];
+      DP("device %d: enqueue constr/destr\n", D);
+      Device.PendingConstrDestrHostPtrList.push_back(host_ptr);
+    }
+    DP("Done with constructor/destructor\n");
+    return OFFLOAD_SUCCESS;
+  }
+
+  // No devices available?
+  if (!(device_id >= 0 && (size_t)device_id < Devices.size())) {
+    DP("Device ID %d does not have a matching RTL.\n", device_id);
+    return OFFLOAD_FAIL;
+  }
+
+  // Get device info
+  DeviceTy &Device = Devices[device_id];
+  DP("Is the device %d (local is %d) initialized? %d\n", device_id,
+     Device.RTLDeviceID, (int)Device.IsInit);
+
+  // Init the device if not done before
+  if (!Device.IsInit) {
+    assert(!IsConstrDestrRecursiveCall && "constr & destr should not init RT");
+    if (Device.init() != OFFLOAD_SUCCESS) {
+      DP("failed to init device %d\n", device_id);
+      return OFFLOAD_FAIL;
+    }
+  }
+
+  if (!IsConstrDestrRecursiveCall &&
+      !Device.PendingConstrDestrHostPtrList.empty()) {
+    DP("has pending constr/destr... call now\n");
+    for (std::list<void *>::iterator
+             ii = Device.PendingConstrDestrHostPtrList.begin(),
+             ie = Device.PendingConstrDestrHostPtrList.end();
+         ii != ie; ++ii) {
+      void *ConstrDestrHostPtr = *ii;
+      int rc = target(device_id, ConstrDestrHostPtr, 0, NULL, NULL, NULL, NULL,
+                      1, 1, true /*team*/, true /*recursive*/);
+      if (rc != OFFLOAD_SUCCESS) {
+        DP("failed to run constr/destr... enqueue it\n");
+        return OFFLOAD_FAIL;
+      }
+    }
+    DP("done with pending constr/destr\n");
+    Device.PendingConstrDestrHostPtrList.clear();
+  }
+
+  // Find the table information in the map or look it up in the translation
+  // tables
+  TableMap *TM = 0;
+  HostPtrToTableMapTy::iterator TableMapIt = HostPtrToTableMap.find(host_ptr);
+  if (TableMapIt == HostPtrToTableMap.end()) {
+    // We don't have a map. So search all the registered libraries
+    for (HostEntriesBeginToTransTableTy::iterator
+             ii = HostEntriesBeginToTransTable.begin(),
+             ie = HostEntriesBeginToTransTable.end();
+         !TM && ii != ie; ++ii) {
+      // get the translation table (which contains all the good info)
+      TranslationTable *TransTable = &ii->second;
+      // iterate over all the host table entries to see if we can locate the
+      // host_ptr
+      __tgt_offload_entry *begin = TransTable->HostTable.EntriesBegin;
+      __tgt_offload_entry *end = TransTable->HostTable.EntriesEnd;
+      __tgt_offload_entry *cur = begin;
+      for (uint32_t i = 0; cur < end; ++cur, ++i) {
+        if (cur->addr != host_ptr)
+          continue;
+        // we got a match, now fill the HostPtrToTableMap so that we
+        // may avoid this search next time.
+        TM = &HostPtrToTableMap[host_ptr];
+        TM->Table = TransTable;
+        TM->Index = i;
+        break;
+      }
+    }
+  } else {
+    TM = &TableMapIt->second;
+  }
+  // No map for this host pointer found!
+  if (!TM) {
+    DP("Host ptr %016lx does not have a matching target pointer.\n",
+       (Elf64_Addr)host_ptr);
+    return OFFLOAD_FAIL;
+  }
+
+  // get target table
+  assert(TM->Table->TargetsTable.size() > (size_t)device_id &&
+         "Not expecting a device ID outside the tables bounds!");
+  __tgt_target_table *TargetTable = TM->Table->TargetsTable[device_id];
+  // if first call, need to move the data
+  if (!TargetTable) {
+    // 1) get image
+    assert(TM->Table->TargetsImages.size() > (size_t)device_id &&
+           "Not expecting a device ID outside the tables bounds!");
+    __tgt_device_image *img = TM->Table->TargetsImages[device_id];
+    if (!img) {
+      DP("No image loaded for device id %d.\n", device_id);
+      return OFFLOAD_FAIL;
+    }
+    // 2) load image into the target table
+    TargetTable = TM->Table->TargetsTable[device_id] = Device.load_binary(img);
+    // Unable to get table for this image: invalidate image and fail
+    if (!TargetTable) {
+      DP("Unable to generate entries table for device id %d.\n", device_id);
+      TM->Table->TargetsImages[device_id] = 0;
+      return OFFLOAD_FAIL;
+    }
+
+    // Verify if the two tables sizes match
+    size_t hsize =
+        TM->Table->HostTable.EntriesEnd - TM->Table->HostTable.EntriesBegin;
+    size_t tsize = TargetTable->EntriesEnd - TargetTable->EntriesBegin;
+
+    // Invalid image for this host entries!
+    if (hsize != tsize) {
+      DP("Host and Target tables mismatch for device id %d [%lx != %lx].\n",
+         device_id, hsize, tsize);
+      TM->Table->TargetsImages[device_id] = 0;
+      TM->Table->TargetsTable[device_id] = 0;
+      return OFFLOAD_FAIL;
+    }
+    assert(TM->Index < hsize &&
+           "Not expecting index greater than the table size");
+
+    // process global data that needs to be mapped
+    __tgt_target_table *HostTable = &TM->Table->HostTable;
+    for (__tgt_offload_entry *CurrDeviceEntry = TargetTable->EntriesBegin,
+                             *CurrHostEntry = HostTable->EntriesBegin,
+                             *EntryDeviceEnd = TargetTable->EntriesEnd;
+         CurrDeviceEntry != EntryDeviceEnd;
+         CurrDeviceEntry++, CurrHostEntry++) {
+      if (CurrDeviceEntry->size != 0) {
+        // has data
+        assert(CurrDeviceEntry->size == CurrHostEntry->size &&
+               "data size mismatch");
+        assert(Device.getTgtPtrBegin(CurrHostEntry->addr,
+                                     CurrHostEntry->size) == NULL &&
+               "data in declared target should not be already mapped");
+        // add entry to map
+        DP("add mapping from host 0x%llx to 0x%llx with size %lld\n\n",
+           (unsigned long long)CurrHostEntry->addr,
+           (unsigned long long)CurrDeviceEntry->addr,
+           (unsigned long long)CurrDeviceEntry->size);
+        Device.HostDataToTargetMap.push_front(HostDataToTargetTy(
+            (long)CurrHostEntry->addr, (long)CurrHostEntry->addr,
+            (long)CurrHostEntry->addr + CurrHostEntry->size,
+            (long)CurrDeviceEntry->addr,
+            (long)CurrDeviceEntry->addr + CurrDeviceEntry->size));
+      }
+    }
+  }
+
+  // Move data to device
+  target_data_begin(Device, arg_num, args_base, args, arg_sizes, arg_types);
+
+  std::vector<void *> tgt_args;
+
+  for (int32_t i = 0; i < arg_num; ++i) {
+
+    if (arg_types[i] & tgt_map_extra)
+      continue;
+
+    void *HstPtrBegin = args[i];
+    void *HstPtrBase = args_base[i];
+    void *TgtPtrBase;
+    long IsLast; // unused
+    if (arg_types[i] & tgt_map_pointer) {
+      DP("Obtaining target argument from host pointer %016lx to object %016lx "
+         "\n",
+         (long)HstPtrBase, (long)HstPtrBegin);
+      void *TgtPtrBegin =
+          Device.getTgtPtrBegin(HstPtrBase, sizeof(void *), IsLast, false);
+      TgtPtrBase = TgtPtrBegin; // no offset for ptrs
+    } else {
+      DP("Obtaining target argument from host pointer %016lx\n",
+         (long)HstPtrBegin);
+      void *TgtPtrBegin =
+          Device.getTgtPtrBegin(HstPtrBegin, arg_sizes[i], IsLast, false);
+      assert(TgtPtrBegin && "NULL argument for hst ptr");
+      uint64_t PtrDelta = (uint64_t)HstPtrBegin - (uint64_t)HstPtrBase;
+      TgtPtrBase = (void *)((uint64_t)TgtPtrBegin - PtrDelta);
+    }
+    tgt_args.push_back(TgtPtrBase);
+  }
+  // Push omp handle
+  tgt_args.push_back((void *)0);
+
+  // Launch device execution
+  int rc;
+  DP("Launching target execution with pointer %016lx (index=%d).\n",
+     (Elf64_Addr)TargetTable->EntriesBegin[TM->Index].addr, TM->Index);
+  if (IsTeamConstruct) {
+    rc = Device.run_team_region(TargetTable->EntriesBegin[TM->Index].addr,
+                                &tgt_args[0], tgt_args.size(), team_num,
+                                thread_limit);
+  } else {
+    rc = Device.run_region(TargetTable->EntriesBegin[TM->Index].addr,
+                           &tgt_args[0], tgt_args.size());
+  }
+
+  if (rc)
+    return OFFLOAD_FAIL;
+
+  // Move data from device
+  target_data_end(Device, arg_num, args_base, args, arg_sizes, arg_types);
+  return OFFLOAD_SUCCESS;
+}
+
+EXTERN int __tgt_target(int32_t device_id, void *host_ptr, int32_t arg_num,
+                        void **args_base, void **args, int64_t *arg_sizes,
+                        int32_t *arg_types) {
+  return target(device_id, host_ptr, arg_num, args_base, args, arg_sizes,
+                arg_types, 0, 0, false /*team*/, false /*recursive*/);
+}
+
+EXTERN int __tgt_target_nowait(int32_t device_id, void *host_ptr,
+                               int32_t arg_num, void **args_base, void **args,
+                               int64_t *arg_sizes, int32_t *arg_types,
+                               int32_t depNum, void *depList,
+                               int32_t noAliasDepNum, void *noAliasDepList) {
+  return __tgt_target(device_id, host_ptr, arg_num, args_base, args, arg_sizes,
+                      arg_types);
+}
+
+EXTERN int __tgt_target_teams(int32_t device_id, void *host_ptr,
+                              int32_t arg_num, void **args_base, void **args,
+                              int64_t *arg_sizes, int32_t *arg_types,
+                              int32_t team_num, int32_t thread_limit) {
+  return target(device_id, host_ptr, arg_num, args_base, args, arg_sizes,
+                arg_types, team_num, thread_limit, true /*team*/,
+                false /*recursive*/);
+}
+
+EXTERN int __tgt_target_teams_nowait(int32_t device_id, void *host_ptr,
+                                     int32_t arg_num, void **args_base,
+                                     void **args, int64_t *arg_sizes,
+                                     int32_t *arg_types, int32_t team_num,
+                                     int32_t thread_limit, int32_t depNum,
+                                     void *depList, int32_t noAliasDepNum,
+                                     void *noAliasDepList) {
+  return __tgt_target_teams(device_id, host_ptr, arg_num, args_base, args,
+                            arg_sizes, arg_types, team_num, thread_limit);
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// temporary for debugging (matching the ones in omptarget-nvptx
+
+EXTERN void __kmpc_kernel_print(char *title) { DP(" %s\n", title); }
+
+EXTERN void __kmpc_kernel_print_int8(char *title, int64_t data) {
+  DP(" %s val=%lld\n", title, (long long)data);
+}

libomptarget/src/targets_info.h

+//===-------- targets_info.h - Information about Target RTLs ------ C++ -*-===//
+//
+//                     The LLVM Compiler Infrastructure
+//
+// This file is dual licensed under the MIT and the University of Illinois Open
+// Source Licenses. See LICENSE.txt for details.
+//
+//===----------------------------------------------------------------------===//
+//
+// Map between ELF machine IDs and the RTL library that supports it
+//
+//===----------------------------------------------------------------------===//
+
+#include <stdint.h>
+
+#ifndef _TARGETS_INFO_H_
+#define _TARGETS_INFO_H_
+
+struct targets_info_table_entry {
+  uint16_t Machine_Elf_ID;
+  const char *Machine_RTL_Lib;
+};
+struct targets_info_table {
+  int32_t Number_of_Entries;
+  targets_info_table_entry *Entries;
+};
+
+#endif

libomptarget/src/targets_info.cpp

+//===-------- targets_info.cpp - Information about Target RTLs ---- C++ -*-===//
+//
+//                     The LLVM Compiler Infrastructure
+//
+// This file is dual licensed under the MIT and the University of Illinois Open
+// Source Licenses. See LICENSE.txt for details.
+//
+//===----------------------------------------------------------------------===//
+//
+// Map between ELF machine IDs and the RTL library that supports it
+//
+//===----------------------------------------------------------------------===//
+
+#include "targets_info.h"
+
+static targets_info_table_entry targets_info_entries[] = {
+    {21 /* EM_PPC64 */, "libomptarget.rtl.ppc64.so"},
+    {62 /* EM_X86_64*/, "libomptarget.rtl.x86_64.so"},
+    {190 /* EM_CUDA */, "libomptarget.rtl.cuda.so"}};
+
+targets_info_table targets_info = {sizeof(targets_info_entries) /
+                                       sizeof(targets_info_table_entry),
+                                   &targets_info_entries[0]};

libomptarget/test/CMakeLists.txt

+# CMakeLists.txt file for unit testing OpenMP Library
+include(FindPythonInterp)
+include(CheckTypeSize)
+if(NOT PYTHONINTERP_FOUND)
+  libomptarget_warning_say("Could not find Python.")
+  libomptarget_warning_say("The check-libomptarget target will not be available!")
+  return()
+endif()
+
+set(LIBOMPTARGET_TEST_CFLAGS "" CACHE STRING
+  "Extra compiler flags to send to the test compiler")
+
+if(${LIBOMPTARGET_STANDALONE_BUILD})
+  # Make sure we can use the console pool for recent cmake and ninja > 1.5
+  if(CMAKE_VERSION VERSION_LESS 3.1.20141117)
+    set(cmake_3_2_USES_TERMINAL)
+  else()
+    set(cmake_3_2_USES_TERMINAL USES_TERMINAL)
+  endif()
+  set(LIBOMPTARGET_TEST_C_COMPILER ${CMAKE_C_COMPILER} CACHE STRING
+    "C compiler to use for testing OpenMP offloading library")
+  set(LIBOMPTARGET_TEST_CXX_COMPILER ${CMAKE_CXX_COMPILER} CACHE STRING
+    "C++ compiler to use for testing OpenMP offloading library")
+  set(LIBOMPTARGET_TEST_OPENMP_FLAG -fopenmp CACHE STRING
+    "OpenMP compiler flag to use for testing OpenMP offloading library")
+  set(LIBOMPTARGET_LLVM_LIT_EXECUTABLE "" CACHE STRING
+    "Path to llvm-lit")
+  find_program(LIT_EXECUTABLE NAMES llvm-lit ${LIBOMPTARGET_LLVM_LIT_EXECUTABLE})
+  if(NOT LIT_EXECUTABLE)
+    libomptarget_say("Cannot find llvm-lit.")
+    libomptarget_say("Please put llvm-lit in your PATH or set LIBOMPTARGET_LLVM_LIT_EXECUTABLE to its full path")
+    libomptarget_warning_say("The check-libomptarget target will not be available!")
+    return()
+  endif()
+  
+  set(LIBOMPTARGET_FILECHECK_EXECUTABLE "" CACHE STRING
+    "Path to FileCheck")
+  find_program(LIBOMPTARGET_FILECHECK NAMES FileCheck ${LIBOMPTARGET_FILECHECK_EXECUTABLE})
+  if(NOT LIBOMPTARGET_FILECHECK)
+    libomptarget_say("Cannot find FileCheck.")
+    libomptarget_say("Please put FileCheck in your PATH or set LIBOMPTARGET_FILECHECK_EXECUTABLE to its full path")
+    libomptarget_warning_say("The check-libomptarget target will not be available!")
+    return()
+  endif()
+  
+  # Set lit arguments
+  # The -j 1 lets the actual tests run with the entire machine.
+  # We have one test thread that spawns the tests serially.  This allows
+  # Each test to use the entire machine.
+  set(LIBOMPTARGET_LIT_ARGS_DEFAULT "-sv --show-unsupported --show-xfail -j 1")
+  if(MSVC OR XCODE)
+    set(LIBOMPTARGET_LIT_ARGS_DEFAULT "${LIBOMPTARGET_LIT_ARGS_DEFAULT} --no-progress-bar")
+  endif()
+  set(LIBOMPTARGET_LIT_ARGS "${LIBOMPTARGET_LIT_ARGS_DEFAULT}" CACHE STRING
+    "Default options for lit")
+  separate_arguments(LIBOMPTARGET_LIT_ARGS)
+  add_custom_target(check-libomptarget
+    COMMAND ${PYTHON_EXECUTABLE} ${LIT_EXECUTABLE} ${LIBOMPTARGET_LIT_ARGS} ${CMAKE_CURRENT_BINARY_DIR}
+    DEPENDS omptarget
+    COMMENT "Running libomptarget tests"
+    ${cmake_3_2_USES_TERMINAL}
+  )
+  
+  set(LIBOMPTARGET_OPENMP_HEADER_FOLDER "" CACHE STRING
+    "Path to folder containing omp.h")
+  set(LIBOMPTARGET_OPENMP_HOST_RTL_FOLDER "" CACHE STRING
+    "Path to folder containing libomp.h")
+else()
+  # LLVM source tree build, test just-built clang
+  if(NOT MSVC)
+    set(LIBOMPTARGET_TEST_C_COMPILER ${LLVM_RUNTIME_OUTPUT_INTDIR}/clang)
+    set(LIBOMPTARGET_TEST_CXX_COMPILER ${LLVM_RUNTIME_OUTPUT_INTDIR}/clang++)
+    set(LIBOMPTARGET_FILECHECK ${LLVM_RUNTIME_OUTPUT_INTDIR}/FileCheck)
+  else()
+    libomptarget_error_say("Not prepared to run tests on Windows systems.")
+  endif()
+  set(LIBOMPTARGET_TEST_OPENMP_FLAG -fopenmp=libomp)
+  # Use add_lit_testsuite() from LLVM CMake. This also depends on OpenMP
+  # implementation because it uses omp.h.
+  add_lit_testsuite(check-libomptarget
+    "Running libomptarget tests"
+    ${CMAKE_CURRENT_BINARY_DIR}
+    ARGS "-j 1"
+    DEPENDS omptarget omp
+  )
+  
+  set(LIBOMPTARGET_OPENMP_HEADER_FOLDER "${LIBOMPTARGET_BINARY_DIR}/../runtime/src")
+endif()
+
+# Configure the lit.site.cfg.in file
+set(AUTO_GEN_COMMENT "## Autogenerated by libomptarget configuration.\n# Do not edit!")
+configure_file(lit.site.cfg.in lit.site.cfg @ONLY)
+

libomptarget/test/lit.cfg

+# -*- Python -*- vim: set ft=python ts=4 sw=4 expandtab tw=79:
+# Configuration file for the 'lit' test runner.
+
+import os
+import lit.formats
+
+# Tell pylint that we know config and lit_config exist somewhere.
+if 'PYLINT_IMPORT' in os.environ:
+    config = object()
+    lit_config = object()
+
+def append_dynamic_library_path(name, value, sep):
+    if name in config.environment:
+        config.environment[name] = value + sep + config.environment[name]
+    else:
+        config.environment[name] = value
+
+# name: The name of this test suite.
+config.name = 'libomptarget'
+
+# suffixes: A list of file extensions to treat as test files.
+config.suffixes = ['.c', '.cpp', '.cc']
+
+# test_source_root: The root path where tests are located.
+config.test_source_root = os.path.dirname(__file__)
+
+# test_exec_root: The root object directory where output is placed
+config.test_exec_root = config.libomptarget_obj_root
+
+# test format
+config.test_format = lit.formats.ShTest()
+
+# compiler flags
+config.test_cflags = config.test_openmp_flag + \
+    " -I " + config.test_source_root + \
+    " -I " + config.omp_header_directory + \
+    " -L " + config.library_dir;
+    
+if config.omp_host_rtl_directory:
+    config.test_cflags = config.test_cflags + " -L " + \
+        config.omp_host_rtl_directory  
+    
+config.test_cflags = config.test_cflags + " " + config.test_extra_cflags
+
+# Setup environment to find dynamic library at runtime
+if config.operating_system == 'Windows':
+    append_dynamic_library_path('PATH', config.library_dir, ";")
+    append_dynamic_library_path('PATH', config.omp_host_rtl_directory, ";")
+elif config.operating_system == 'Darwin':
+    append_dynamic_library_path('DYLD_LIBRARY_PATH', config.library_dir, ":")
+    append_dynamic_library_path('DYLD_LIBRARY_PATH', \
+        config.omp_host_rtl_directory, ";")
+    config.test_cflags += " -Wl,-rpath," + config.library_dir
+    config.test_cflags += " -Wl,-rpath," + config.omp_host_rtl_directory
+else: # Unices
+    append_dynamic_library_path('LD_LIBRARY_PATH', config.library_dir, ":")
+    append_dynamic_library_path('LD_LIBRARY_PATH', \
+        config.omp_host_rtl_directory, ":")
+    
+# substitutions
+# - for targets that exist in the system create the actual command.
+# - for valid targets that do not exist in the system, return false, so that the
+#   same test can be used for different targets.
+
+# Scan all the valid targets.
+for libomptarget_target in config.libomptarget_all_targets:
+    # Is this target in the current system? If so create a compile, run and test
+    # command. Otherwise create command that return false.
+    if libomptarget_target in config.libomptarget_system_targets:
+        config.substitutions.append(("%libomptarget-compilexx-run-and-check-" + \
+            libomptarget_target, \
+            "%libomptarget-compilexx-and-run-" + libomptarget_target + \
+            " | " + config.libomptarget_filecheck + " %s"))
+        config.substitutions.append(("%libomptarget-compile-run-and-check-" + \
+            libomptarget_target, \
+            "%libomptarget-compile-and-run-" + libomptarget_target + \
+            " | " + config.libomptarget_filecheck + " %s"))      
+        config.substitutions.append(("%libomptarget-compilexx-and-run-" + \
+            libomptarget_target, \
+            "%clangxx-" + libomptarget_target + " %s -o %t-" + \
+            libomptarget_target + " && %t-" + libomptarget_target))    
+        config.substitutions.append(("%libomptarget-compile-and-run-" + \
+            libomptarget_target, \
+            "%clang-" + libomptarget_target + " %s -o %t-" + \
+            libomptarget_target + " && %t-" + libomptarget_target))
+        config.substitutions.append(("%clangxx-" + libomptarget_target, \
+            "%clangxx %cflags -omptargets=" + libomptarget_target))    
+        config.substitutions.append(("%clang-" + libomptarget_target, \
+            "%clang %cflags -omptargets=" + libomptarget_target))
+        config.substitutions.append(("%fcheck-" + libomptarget_target, \
+            config.libomptarget_filecheck + " %s"))
+    else:
+        config.substitutions.append(("%libomptarget-compile-run-and-check-" + \
+            libomptarget_target, \
+            "echo ignored-command"))
+        config.substitutions.append(("%libomptarget-compilexx-run-and-check-" + \
+            libomptarget_target, \
+            "echo ignored-command"))
+        config.substitutions.append(("%libomptarget-compile-and-run-" + \
+            libomptarget_target, \
+            "echo ignored-command"))
+        config.substitutions.append(("%libomptarget-compilexx-and-run-" + \
+            libomptarget_target, \
+            "echo ignored-command"))
+        config.substitutions.append(("%clang-" + libomptarget_target, \
+            "echo ignored-command"))
+        config.substitutions.append(("%clangxx-" + libomptarget_target, \
+            "echo ignored-command"))
+        config.substitutions.append(("%fcheck-" + libomptarget_target, \
+            "echo ignored-command"))
+
+config.substitutions.append(("%clangxx", config.test_cxx_compiler))
+config.substitutions.append(("%clang", config.test_c_compiler))
+config.substitutions.append(("%openmp_flag", config.test_openmp_flag))
+config.substitutions.append(("%cflags", config.test_cflags))
+

libomptarget/test/lit.site.cfg.in

+@AUTO_GEN_COMMENT@
+
+config.test_c_compiler = "@LIBOMPTARGET_TEST_C_COMPILER@"
+config.test_cxx_compiler = "@LIBOMPTARGET_TEST_CXX_COMPILER@"
+config.test_openmp_flag = "@LIBOMPTARGET_TEST_OPENMP_FLAG@"
+# For the moment we still need to pass libomptarget explicitly. Once the driver
+# patch, lands, this is not required anymore.
+config.test_extra_cflags = "-lomptarget @LIBOMPTARGET_TEST_CFLAGS@"
+config.libomptarget_obj_root = "@CMAKE_CURRENT_BINARY_DIR@"
+config.library_dir = "@LIBOMPTARGET_LIBRARY_DIR@"
+config.omp_header_directory = "@LIBOMPTARGET_OPENMP_HEADER_FOLDER@"
+config.omp_host_rtl_directory = "@LIBOMPTARGET_OPENMP_HOST_RTL_FOLDER@"
+config.operating_system = "@CMAKE_SYSTEM_NAME@"
+config.libomptarget_all_targets = "@LIBOMPTARGET_ALL_TARGETS@".split()
+config.libomptarget_system_targets = "@LIBOMPTARGET_SYSTEM_TARGETS@".split()
+config.libomptarget_filecheck = "@LIBOMPTARGET_FILECHECK@"
+
+# Let the main config do the real work.
+lit_config.load_config(config, "@LIBOMPTARGET_BASE_DIR@/test/lit.cfg")
+

libomptarget/test/offloading/offloading_success.c

+// RUN: %libomptarget-compile-run-and-check-powerpc64-ibm-linux-gnu
+// RUN: %libomptarget-compile-run-and-check-powerpc64le-ibm-linux-gnu
+// RUN: %libomptarget-compile-run-and-check-x86_64-pc-linux-gnu
+
+#include <stdio.h>
+#include <omp.h>
+
+int main(void) {
+  int isHost = 0;
+
+#pragma omp target
+  { isHost = omp_is_initial_device(); }
+
+  // The compiler doesn't have support to launch the target region on the
+  // device.
+  // CHECK: Target region executed on the host
+  printf("Target region executed on the %s\n", isHost ? "host" : "device");
+
+  return !isHost;
+}

libomptarget/test/offloading/offloading_success.cpp

+// RUN: %libomptarget-compilexx-run-and-check-powerpc64-ibm-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-powerpc64le-ibm-linux-gnu
+// RUN: %libomptarget-compilexx-run-and-check-x86_64-pc-linux-gnu
+
+#include <stdio.h>
+#include <omp.h>
+
+int main(void) {
+  int isHost = 0;
+
+#pragma omp target
+  { isHost = omp_is_initial_device(); }
+
+  // The compiler doesn't have support to launch the target region on the
+  // device.
+  // CHECK: Target region executed on the host
+  printf("Target region executed on the %s\n", isHost ? "host" : "device");
+
+  return !isHost;
+}