[OPENMP] Driver support for OpenMP offloading

[sfantao]

Summary

With a full implementation of OpenMP 3.1. already available upstream, we aim at continuing that work and add support for OpenMP 4.0 as well. One important component introduced by OpenMP 4.0 is offloading which enables the execution of a given structured block to be transferred to a device other than the host.

An implementation for OpenMP offloading infrastructure in clang is proposed in http://goo.gl/L1rnKJ. This document is already a second iteration that includes contributions from several vendors and members of the LLVM community. It was published in http://lists.cs.uiuc.edu/pipermail/llvmdev/2015-April/084304.html for discussion by the community, and so far we didn’t have any major concern about the design.

Unlike other OpenMP components, offloading requires support from the compiler driver given that for the same source file, several (host and target) objects will be generated using potentially different toolchains. At the same time, the compiler needs to have a mechanism to relate variables in the host with the ones generated with target, so communication between toolchains is required. The way this relation is supported by the driver will also have implications in the code generation.

This patch proposes an implementation of the driver support for offloading. The following summarizes the main changes this patch introduces:

a) clang can be invoked with -fopenmp=libiom5 -omptargets=triple1,…,tripleN, where triplei are the target triples the user wants to be able to offload to.

b) driver detects whether the offloading triples are valid or not and if the corresponding toolchain is prepared to offload. This patch only enables offloading for Linux toolchains.

c) Each target compiler phase takes the host IR (result of the host compiler phase) as a second input. This will enable the host generation to specify the variables that should be emitted for the target in the form of metadata and this metadata could be read by the target frontend.

d) Given that the same host IR result info is used by the different toolchains, the driver keeps a cache of results in order to avoid the job that generates a given result to be emitted twice.

e) Offloading leverages the argument translation functionality in order to convert host arguments into target arguments. This is currently used to make sure a shared library is always produced by the target toolchain - a library that can be loaded by the OpenMP runtime library.

f) The target shared libraries are embedded into the host binary by using a linker script produced by the driver and passed to the host linker.

g) The driver passes to the frontend offloading a command that specify if the frontend is producing code for a target. This is required as the code generation for target and host have to be different.

h) A full path to the original source file is passed to the frontend so it can be used to produce unique IDs that are the same for the host and targets.

Thanks!
Samuel

[rjmccall]

Hmm. Using the host IR as an implicit line of communication is an interesting approach. Can you expand on what kind of information needs to flow from host to target there, or at least link to a place in the previous discussion?

[sfantao]

Hi John

Thanks for looking into this patch!

Sure, let me expand on the host-target communication. Just a little bit of context before I do that:

During code generation, the target frontend has to decide whether a given declaration or target region has to be emitted or not. Without any information communicated from the host frontend, this decision become complicated for cases like:

• #pragma omp target regions in static functions or class members;
• static declarations delimitted by #pragma omp declare target regions that end up not being used;
• #pragma omp target in template functions

In order for the target frontend to correctly identify all the declarations that need to be emitted it would have to, somehow, emulate the actions done by the host frontend which would turn the code generation messy in places that do not even relate with OpenMP.

On top of that, in order to have an efficient mapping between host and target entries (global declarations/target regions)
table (this is discussed in the document, in section 5.1, where __tgt_offload_entry is introduced) the compiler would have to emit the corresponding entries in the host and target side in the same order. This is useful for devices whose toolchain maintain the order of the symbols given that the order of the entries in the host and target tables will be the same after linking. So knowing an index would be enough to do the mapping. In order for that to happen, the target frontend would have to know that order, which would be also hard to extract if no information is communicated form the host.

So, the information that needs to be propagated to make what I described above possible is basically i) declaration mangled names and ii) order they were emitted. This information could be communicated in the form of metadata that is emitted by the host frontend when the module is released and loaded by the target frontend when CGOpenMPRuntime is created. This information has however to be coded in slightly different ways for different kinds of declarations. Let me explain this with an example:

//######################################
#pragma omp declare target
struct MyClass{
  ...

  MyClass &add(MyClass &op){...}
  
  MyClass &add(int (&op)[N]){...}
  
  bool eq(MyClass &op){...}

  MyClass() {...}
  
  ~MyClass() {...}
};

MyClass C;
MyClass D;
#pragma omp end declare target 	

void foo(){
  int AA[N];
  MyClass H, T;
  MyClass HC;
  
  ...

  #pragma omp target
  {
    MyClass TC;
    T.add(AA).add(HC);
  }
  
  if (H.eq(T)) {...}
  
  #pragma omp target
  {
    T.add(AA);
  } 
}
//######################################

I was planning the metadata for this example to look more or less like this:

; Named metadata that encloses all the offloading information
!openmp.offloading.info = !{!1, !2, !3, !4, !5, !6, !7, !8, !9, !10}

; Global variables that require a map between host and target:
; Entry 0 -> ID for this type of metadata (0)
; Entry 1 -> Mangled name of the variable
; Entry 2 -> Order it was emitted
!1 = !{i32 0, !"C", i32 0}
!2 = !{i32 0, !"D", i32 2}

; Functions with target regions
; Entry 0 -> ID for this type of metadata (1)
; Entry 1 -> Mangled name of the function that was emitted for the host and encloses target regions
; Entry 2-n -> Order the target regions in the functions (in the same sequence the statements are found) are emitted 
!3 = !{i32 1, !"_Z3foov", i32 4, i32 5}

; Global initializers
; Entry 0 -> ID for this type of metadata (2)
; Entry 1-n -> Order the initializers are emitted in descending order of priority (we will require a target region per set of initializers with the same priority)
!4 = !{i32 2, i32 6}

; Global Dtors
; Entry 0 -> ID for this type of metadata (3)
; Entry 1 -> Mangled name of the variable to be destructed 
; Entry 2 -> Order the destructor was emitted (we will have a target region per variable being destroyed - this can probably be optimized)
!5 = !{i32 3, !"C", i32 1}
!6 = !{i32 3, !"D", i32 3}

; Other functions that should be emitted in the target but do not require to be mapped to the host
; Entry 0 -> ID for this type of metadata (4)
; Entry 1 -> Mangled name of the function that has to be emitted.
!7 = !{i32 4, !"_ZN7MyClass3addERA64_i"}
!8 = !{i32 4, !"_ZN7MyClass3addERS_"}
!9 = !{i32 4, !"_ZN7MyClassC2Ev"}
!10 = !{i32 4, !"_ZN7MyClassD2Ev"}

I realize this is the kind of information I should propose as a patch to the codegen part of offloading, but I think it makes sense to discuss it now as the driver has to enable it.

I also foresee the communication between target and host to be useful for other cases, like the propagation of alias information from host to target. I don’t have have however a proposal for that at this moment.

Hope I haven’t been either too brief or too exhaustive! Let me know if I can clarify anything else for you.

Thanks!
Samuel

[sfantao]

Are there any other comments or questions about this patch?

Many thanks!
Samuel

[sfantao]

I've just noticed Chad is owning the Compiler driver, so I believe he should also be added to the list of reviewer of this patch.

Thanks!
Samuel

[echristo]

Quite a big patch, I'd definitely like to take a look at this as well. It relates to how some of the cuda work is progressing too.

Thanks!

-eric

[sfantao]

In D9888#188244, @echristo wrote:

Quite a big patch, I'd definitely like to take a look at this as well. It relates to how some of the cuda work is progressing too.

Thanks!

-eric

Thanks eric,

Please let me know any comments you may have.

I agree the patch is quite big... I had a hard time trying to find a better partition that would make sense - this requires small but related changes in several places - and mapped to something meaningful in terms of the regression tests. If you see a good way to partition the patch let me know and I would gladly do it.

Thanks again!
Samuel

[Hahnfeld]

I think this has to be updated for the current trunk...

[sfantao]

This diff refactors the original patch and is rebased on top of the latests offloading changes inserted for CUDA.

Here I don't touch the CUDA support. I tried, however, to have the implementation modular enough so that it could eventually be combined with the CUDA implementation. In my view OpenMP offloading is more general in the sense that it does not refer to a given tool chain, instead it uses existing toolchains to generate code for offloading devices. So, I believe that a tool chain (which I did not include in this patch) targeting NVPTX will be able to handle both CUDA and OpenMP offloading models.

Chris, Art, I understand you have worked out the latest CUDA changes so any feedback from you is greatly appreciated!

Here are few more details about this diff:

a) Add tool to bundle and unbundle corresponding host and device files into a single one.

One of the goals of OpenMP offloading is to enable users to offload with little effort, by annotating the code with a few pragmas. I'd also like to save users the trouble of changing their existent applications' build system. So having the compiler always return a single file instead of one for the host and each target even if the user is doing separate compilation is desirable.

This diff includes a tool named clang-offload-bundled (happy to change the name or even include it in the driver if someone thinks it is the best direction to go) that is used on all input files that are not source files to unbundle them, and on top level jobs that are not linking jobs to bundle the results obtained for host and each target.

The format of the bundled files is currently very simple: text formats are concatenated with comments that have a magic string and target identifying triple in between, and binary formats have a header that contains the triple and the offset and size of the code for host and each target.

This tool still has to be improved in the future to deal with archive files so that each individual file in the archive is properly dealt with. We see that archives are very commonly used in current application to combine separate compilation results. So I'm convinced users would enjoy this feature.

b ) The building of the driver actions is unchanged.

I don't create device specific actions. Instead only the bundling/unbundling are inserted as first or last action if the file type requires that.

c) Add offloading kind to ToolChain

Offloading does not require a new toolchain to be created. Existent toolchains are used and the offloading kind is used to drive specific behavior in each toolchain so that valid device code is generated.

This is a major difference from what is currently done for CUDA. But I guess the CUDA implementation easily fits this design and the Nvidia GPU toolchain could be reused for both CUDA and OpenMP offloading.

d) Use Job results cache to easily use host results in device actions and vice-versa.

An array of the results for each job is kept so that the device job can use the result previously generated for the host and used it as input or vice-versa.

In OpenMP the device declarations have be communicated from the host frontend to the device frontend. So this is used to conveniently pass that information. Unlike CUDA, OpenMP doesn't have already outline functions with "device" attributes that the frontend can rely on to make the decision on what to be emitted or not.

The result cache can also be updated to keep the required information for the CUDA implementation to decide host/device binaries combining (injection is the term used in the code). I don't have a concrete proposal for that however, given that is not clear to me what are the plans for CUDA to support separate compilation, I understand that the CUDA binary is inserted directly in host IR (Art, can you shed some light on this?).

e) Use compiler generated linker script to do the device/host code combining and correctly support separate compilation.

Currently the OpenMP support in the toolchains is only implemented for Generic GCC targets and a linker script is used to embed the resulting device images into the host binary ELF sections. Also, the linker script defines the symbols that are emitted during code generation so that the address of the images can be easily retrieved.

f) Minor refactoring of the existing code to enable reusing.

I've outlined some of the exiting code into static function so that it could be reused by the new offloading related hooks.

Any comments/remarks are very welcome!

Thanks!
Samuel

[Hahnfeld]

Currently trying to test, but

1. Offloading to the same target isn't supported (x86_64-unknown-linux-gnu as host and device) - this was working with clang-omp

The produced IR isn't showing any calls to the target library and on linkage it complains:

undefined reference to `.omp_offloading.img_start.x86_64-unknown-linux-gnu'
undefined reference to `.omp_offloading.img_end.x86_64-unknown-linux-gnu'
undefined reference to `.omp_offloading.entries_begin'
undefined reference to `.omp_offloading.entries_end'
undefined reference to `.omp_offloading.entries_begin'
undefined reference to `.omp_offloading.entries_end'

(btw: clang-offload-bundler saves the IR file to $TMP with -S -emit-llvm, this seems to be a bug - I had to use --save-temps)

2. I can't seem to figure out the target triple for NVIDIA GPUs. It should be nvptx[64]-nvidia-cuda which gives me

include/llvm/Option/Option.h:101: const llvm::opt::Option llvm::opt::Option::getAlias() const: Assertion `Info && "Must have a valid info!"' failed.

In clang-omp it was nvptxsm_35-nvidia-cuda but this is now invalid...

[tra]

In D9888#257904, @sfantao wrote:

This diff refactors the original patch and is rebased on top of the latests offloading changes inserted for CUDA.

Here I don't touch the CUDA support. I tried, however, to have the implementation modular enough so that it could eventually be combined with the CUDA implementation. In my view OpenMP offloading is more general in the sense that it does not refer to a given tool chain, instead it uses existing toolchains to generate code for offloading devices. So, I believe that a tool chain (which I did not include in this patch) targeting NVPTX will be able to handle both CUDA and OpenMP offloading models.

What do you mean by "does not refer to a given toolchain"? Do you have the toolchain patch available?

Creating a separate toolchain for CUDA was a crutch that was available to craft appropriate cc1 command line for device-side compilation using existing toolchain. It works, but it's rather rigid arrangement. Creating a NVPTX toolchain which can be parameterized to produce CUDA or OpenMP would be an improvement.

Ideally toolchain tweaking should probably be done outside of the toolchain itself so that it can be used with any combination of {CUDA or OpenMP target tweaks}x{toolchains capable of generating target code}.

b ) The building of the driver actions is unchanged.

I don't create device specific actions. Instead only the bundling/unbundling are inserted as first or last action if the file type requires that.

Could you elaborate on that? The way I read it, the driver sees linear chain of compilation steps plus bundling/unbundling at the beginning/end and that each action would result in multiple compiler invocations, presumably per target.

If that's the case, then it may present a bit of a challenge in case one part of compilation depends on results of another. That's the case for CUDA where results of device-side compilation must be present for host-side compilation so we can generate additional code to initialize it at runtime.

c) Add offloading kind to ToolChain

Offloading does not require a new toolchain to be created. Existent toolchains are used and the offloading kind is used to drive specific behavior in each toolchain so that valid device code is generated.

This is a major difference from what is currently done for CUDA. But I guess the CUDA implementation easily fits this design and the Nvidia GPU toolchain could be reused for both CUDA and OpenMP offloading.

Sounds good. I'd be happy to make necessary make CUDA support use it.

d) Use Job results cache to easily use host results in device actions and vice-versa.

An array of the results for each job is kept so that the device job can use the result previously generated for the host and used it as input or vice-versa.

Nice. That's something that will be handy for CUDA and may help to avoid passing bits of info about other jobs explicitly throughout the driver.

The result cache can also be updated to keep the required information for the CUDA implementation to decide host/device binaries combining (injection is the term used in the code). I don't have a concrete proposal for that however, given that is not clear to me what are the plans for CUDA to support separate compilation, I understand that the CUDA binary is inserted directly in host IR (Art, can you shed some light on this?).

Currently CUDA depends on libcudart which assumes that GPU code and its initialization is done the way nvcc does it. Currently we do include PTX assembly (as in readable text) generated on device side into host-side IR *and* generate some host data structures and init code to register GPU binaries with libcudart. I haven't figured out a way to compile host/device sides of CUDA without a host-side compilation depending on device results.

Long-term we're considering implementing CUDA runtime support based on plain driver interface which would give us more control over where we keep GPU code and how we initialize it. Then we could simplify things and, for example, incorporate GPU code via linker script. Alas for the time being we're stuck with libcudart and sequential device and host compilation phases.

As for separate compilation -- compilation part is doable. It's using the results of such compilation that becomes tricky. CUDA's triple-bracket kernel launch syntax depends on libcudart and will not work, because we would not generate init code. You can still launch kernels manually using raw driver API, but it's quite a bit more convoluted.

--Artem

include/clang/Driver/Driver.h

206

re -> are

208

"If offloading is not supported" perhaps?

lib/Driver/Driver.cpp

1849

"has to be"

[sfantao]

Make the offloading ELF sections consistent with what is in http://reviews.llvm.org/D12614.

Fix bug in AtTopLevel flag, so that the bundling job is considered always top level job.

Fix several typos.

[sfantao]

Art, Jonas,

Thanks for the comments!

In D9888#261434, @Hahnfeld wrote:

Currently trying to test, but

1. Offloading to the same target isn't supported (x86_64-unknown-linux-gnu as host and device) - this was working with clang-omp

The produced IR isn't showing any calls to the target library and on linkage it complains:

undefined reference to `.omp_offloading.img_start.x86_64-unknown-linux-gnu'
undefined reference to `.omp_offloading.img_end.x86_64-unknown-linux-gnu'
undefined reference to `.omp_offloading.entries_begin'
undefined reference to `.omp_offloading.entries_end'
undefined reference to `.omp_offloading.entries_begin'
undefined reference to `.omp_offloading.entries_end'

I assume you were trying this using the diff in http://reviews.llvm.org/D12614. There was an inconsistency in the names of the ELF sections and symbols defined by the linker script in these two patches. This is now fixed.

Note that if you are using the libomptarget library from clang-omp, you need to replace in the code .openmptgt_host_entries by .omp_offloading.entries. I changed the names so that all of them are consistent with what is already in place for other OpenMP directives.

I also changed the files generation so that different files are used even if target and host have the same triple.

Please, let me know if it still does not work for you.

(btw: clang-offload-bundler saves the IR file to $TMP with -S -emit-llvm, this seems to be a bug - I had to use --save-temps)

Yes, the bundling job was not being marked as top level. It is now fixed!

2. I can't seem to figure out the target triple for NVIDIA GPUs. It should be nvptx[64]-nvidia-cuda which gives me

include/llvm/Option/Option.h:101: const llvm::opt::Option llvm::opt::Option::getAlias() const: Assertion `Info && "Must have a valid info!"' failed.

In clang-omp it was nvptxsm_35-nvidia-cuda but this is now invalid...

I didn't implement the triples logic for the nvptx targets yet. I'll port that from clang-omp once we have the basic functionality working upstream.

I'll address Art's comments in a separate message.

Thanks again,
Samuel

include/clang/Driver/Driver.h

206

Fixed!

208

Fixed!

lib/Driver/Driver.cpp

1849

Fixed!

[sfantao]

In D9888#262325, @tra wrote:

In D9888#257904, @sfantao wrote:

This diff refactors the original patch and is rebased on top of the latests offloading changes inserted for CUDA.

Here I don't touch the CUDA support. I tried, however, to have the implementation modular enough so that it could eventually be combined with the CUDA implementation. In my view OpenMP offloading is more general in the sense that it does not refer to a given tool chain, instead it uses existing toolchains to generate code for offloading devices. So, I believe that a tool chain (which I did not include in this patch) targeting NVPTX will be able to handle both CUDA and OpenMP offloading models.

What do you mean by "does not refer to a given toolchain"? Do you have the toolchain patch available?

I mean not having to create a toolchain for a specific offloading model. OpenMP offloading is meant for any target and possibility many different targets simultaneously, so having a toolchain for each combination would be overwhelming.

I don't have a patch for the toolchain out for review yet. I'm planing to port what we have in clang-omp for the NVPTX toolchain once I have the host functionality in place. In there (https://github.com/clang-omp/clang_trunk/tree/master/lib/Driver) the Driver is implemented in a different way, I guess the version I'm proposing here is much cleaner. However, the ToolChains shouldn't be that different. All the tweaking is moved to the Tool itself, and I imagine I can drive that using the ToolChain offloading kind I'm proposing here. In https://github.com/clang-omp/clang_trunk/blob/master/lib/Driver/Tools.cpp I basically pick some arguments to forward to the tool and do some tricks to include libdevice in compilation when required. Do you think something like that could also work for CUDA?

Creating a separate toolchain for CUDA was a crutch that was available to craft appropriate cc1 command line for device-side compilation using existing toolchain. It works, but it's rather rigid arrangement. Creating a NVPTX toolchain which can be parameterized to produce CUDA or OpenMP would be an improvement.

Ideally toolchain tweaking should probably be done outside of the toolchain itself so that it can be used with any combination of {CUDA or OpenMP target tweaks}x{toolchains capable of generating target code}.

I agree. I decided to move all the offloading tweaking to the tools, given that that is what clang tool already does: customizes the arguments based on the ToolChain that is passed to it.

b ) The building of the driver actions is unchanged.

I don't create device specific actions. Instead only the bundling/unbundling are inserted as first or last action if the file type requires that.

Could you elaborate on that? The way I read it, the driver sees linear chain of compilation steps plus bundling/unbundling at the beginning/end and that each action would result in multiple compiler invocations, presumably per target.

If that's the case, then it may present a bit of a challenge in case one part of compilation depends on results of another. That's the case for CUDA where results of device-side compilation must be present for host-side compilation so we can generate additional code to initialize it at runtime.

That's right. I try to tackle the challenge of passing host/device results to device/host jobs by using a cache of results as I had described in d). The goal here is to add the flexibility required to accommodate different offloading models. In OpenMP we use host compile results in device compile jobs, and device link results in host link jobs whereas in CUDA the assemble result is used in compile job. I believe that we can have that cache to include whatever information is required to suit all needs.

c) Add offloading kind to ToolChain

Offloading does not require a new toolchain to be created. Existent toolchains are used and the offloading kind is used to drive specific behavior in each toolchain so that valid device code is generated.

This is a major difference from what is currently done for CUDA. But I guess the CUDA implementation easily fits this design and the Nvidia GPU toolchain could be reused for both CUDA and OpenMP offloading.

Sounds good. I'd be happy to make necessary make CUDA support use it.

Great! Thanks.

d) Use Job results cache to easily use host results in device actions and vice-versa.

An array of the results for each job is kept so that the device job can use the result previously generated for the host and used it as input or vice-versa.

Nice. That's something that will be handy for CUDA and may help to avoid passing bits of info about other jobs explicitly throughout the driver.

The result cache can also be updated to keep the required information for the CUDA implementation to decide host/device binaries combining (injection is the term used in the code). I don't have a concrete proposal for that however, given that is not clear to me what are the plans for CUDA to support separate compilation, I understand that the CUDA binary is inserted directly in host IR (Art, can you shed some light on this?).

Currently CUDA depends on libcudart which assumes that GPU code and its initialization is done the way nvcc does it. Currently we do include PTX assembly (as in readable text) generated on device side into host-side IR *and* generate some host data structures and init code to register GPU binaries with libcudart. I haven't figured out a way to compile host/device sides of CUDA without a host-side compilation depending on device results.

Long-term we're considering implementing CUDA runtime support based on plain driver interface which would give us more control over where we keep GPU code and how we initialize it. Then we could simplify things and, for example, incorporate GPU code via linker script. Alas for the time being we're stuck with libcudart and sequential device and host compilation phases.

As for separate compilation -- compilation part is doable. It's using the results of such compilation that becomes tricky. CUDA's triple-bracket kernel launch syntax depends on libcudart and will not work, because we would not generate init code. You can still launch kernels manually using raw driver API, but it's quite a bit more convoluted.

Ok, I see. I am not aware of what exactly libcudart does, but I can elaborate on what the OpenMP offloading implementation we have in place does:

We have a descriptor that is registered with the runtime library (we generate a function for that called before any global initializers are executed ), this descriptor has (among other things) fields that are initialized with the symbols defined by the linker script (so that the runtime library can immediately get the CUDA module) and also the names of the kernels (in OpenMP with don't have user-defined names for these kernels, so we generate some mangling to make sure they are unique). While launching the kernel, the runtime gets a pointer from which he can easily retrieve the name, and the CUDA driver API is used to get the CUDA function to be launched. We have been successfully generating a CUDA module that works well with separate compilation using ptxas and nvlink.

Part of my work is also port the runtime library in clang-omp to the LLLVM OpenMP project. I see CUDA as a simplified version of what OpenMP does, given that the user controls the data mappings explicitly, so I am sure we can find some synergies in the runtime library too and you may be able to use something that we already have in there.

Thanks!
Samuel

--Artem

[Hahnfeld]

In D9888#262389, @sfantao wrote:

[...]

I assume you were trying this using the diff in http://reviews.llvm.org/D12614. There was an inconsistency in the names of the ELF sections and symbols defined by the linker script in these two patches. This is now fixed.

Note that if you are using the libomptarget library from clang-omp, you need to replace in the code .openmptgt_host_entries by .omp_offloading.entries. I changed the names so that all of them are consistent with what is already in place for other OpenMP directives.

I also changed the files generation so that different files are used even if target and host have the same triple.

Please, let me know if it still does not work for you.

Thanks for your help, a small test program now seems to work!

[...]

I didn't implement the triples logic for the nvptx targets yet. I'll port that from clang-omp once we have the basic functionality working upstream.

Ok, I'll wait then. Thanks for your work and finally upstreaming this!
Jonas

[sfantao]

Are there any more comments/suggestions about this patch?

Thanks!
Samuel

[sfantao]

Move clang-offload-bundler to to a separate review: http://reviews.llvm.org/D13909.

This patch depends on http://reviews.llvm.org/D13909.

[sfantao]

Rebase.

[sfantao]

Rebase.

[sfantao]

This comment was removed by sfantao.

[Hahnfeld]

Will this somewhen receive a final review and get merged?

[Hahnfeld]

@rsmith could you possibly take a look at this one? It has been around for roughly 8 months now and hasn't received much feedback

[rsmith]

@echristo, you asked for time to review this; if you still want to, please can you do so?
@tra, it looks like you're happy with this design (and with moving the CUDA offloading support in this direction), please let us know if not!

include/clang/Driver/Options.td

1495–1496

This is an unfortunate flag name; -oblah already means something. Is this name chosen for compatibility with some other system, or could we change it to, say, -fopenmp-targets=?

lib/Driver/Tools.cpp

299

s -> is

[sfantao]

Hi Richard,

Thanks for your review. I partitioned some of the stuff I am proposing here in smaller patches:

http://reviews.llvm.org/D18170
http://reviews.llvm.org/D18171
http://reviews.llvm.org/D18172

These patches already try to incorporate the feedback I got in http://lists.llvm.org/pipermail/cfe-dev/2016-February/047547.html related with the generation of actions.

Thanks again,
Samuel

include/clang/Driver/Options.td

1495–1496

You are right, we are now using -fomptargets in codegen exactly because of that. I can change it to -fopenmp-targets= we don't have any compatibility issues at this point.

lib/Driver/Tools.cpp

299

I'll fix it.

[mkuron]

The three smaller patches into which you divided this one appear to be missing some things. For example, AddOpenMPLinkerScript in lib/Driver/Tools.cpp from this patch appears to still be necessary to get the desired functionality, but it is not present in any of the three.

[sfantao]

Hi Michael,

In D9888#380225, @mkuron wrote:

The three smaller patches into which you divided this one appear to be missing some things. For example, AddOpenMPLinkerScript in lib/Driver/Tools.cpp from this patch appears to still be necessary to get the desired functionality, but it is not present in any of the three.

Those three patches do not add any OpenMP specific code yet, so they do not cover the whole implementation I have here. I am doing things in a slightly different way in the new patches given the feedback I had in the mailing list and I am waiting to review to see if the approach I have in there is acceptable. If so, I'll continue with the OpenMP related patches afterwards.

Thanks,
Samuel

[echristo]

First I'd like to note that the code quality here is really high, most of my comments are higher level design decisions going with the driver and the implementation here rather than that.

One meta comment: offload appears to be something that could be used for CUDA and OpenMP (and OpenACC etc) as a term. I think we should either merge these concepts or pick a different name :)

Thanks for all of your work and patience here! The rest of the comments are inline.

-eric

include/clang/Driver/Driver.h

209–212

Example?

215–216

Any reason?

385–393

This function is starting to get a little silly. Perhaps we should look into refactoring such that this doesn't need to be "the one function that rules them all". Perhaps a different ownership model for the things that are arguments here?

lib/Driver/Compilation.cpp

66–67

Hmm?

lib/Driver/Driver.cpp

222–223

This can probably be done separately? Can you split this out and make it generally useful?

1773–1774

Might be time to make some specialized versions of this function. This may take it from "ridiculously confusing" to "code no one should ever look at" :)

lib/Driver/Tools.cpp

5374

Should we get the offload bundler in first so that the interface is there and testable? (Honest question, no particular opinion here). Though the command lines there will affect how this code is written.

test/OpenMP/target_driver.c

42–48

Do we really think the phases should be a DAG check?

55

How do you pass options to individual omptargets? e.g. -mvsx or -mavx2?

[sfantao]

Hi Eric,

Thanks for the review!

As you are probably a aware, I started partitioning this patch following your initial concern related with the size of this patch and the feedback I got from http://lists.llvm.org/pipermail/cfe-dev/2016-February/047547.html. I am keeping this patch as it shows the big picture of what I am trying to accomplish, so if you prefer to add other higher level suggesting here that's perfectly fine. Let me know if there is a more proper way to link patches.

So, I am incorporating your suggestions here in the partioned patches as specified in the inline comments. The partitioned patches are http://reviews.llvm.org/D18170, http://reviews.llvm.org/D18171 and http://reviews.llvm.org/D18172.

One meta comment: offload appears to be something that could be used for CUDA and OpenMP (and OpenACC etc) as a term. I think we should either merge these concepts or pick a different name :)

Yes, I agree. I am now using offloading. I only refer to the programming model name if the code relates to something specific of that programming model.

Thanks again,
Samuel

include/clang/Driver/Driver.h

209–212

I got rid of this extra toolchain cache and I am organizing it in a multimap by offload kind as Art suggested in http://reviews.llvm.org/D18170. That avoids the multiple containers for the offloading toolchains (this one and the ordered one).

215–216

Currently in OpenMP any directive that relates with offloading supports a device() clause that basically specifies which device to use for that region or data transfer. E.g.

void foo() {
 ...
}

void bar(int i) {
  #pragma omp target device(i)
   foo();
}

... here foo is going to be executed on the device i. The problem is that the device is an integer - it does not tell which device type it is - so it is up to the implementation to decide how i is interpreted. So, if we have a system with two GPUs and two DSP devices. We may bind 0-1 to the GPUs and 2-3 to the DSPs.

My goal with preserving the order of the toolchains was to allow codegen to leverage that information and make a better decision on how to bind devices to integers. Maybe, if the user requests the GPU toolchain first he may be interested in prioritizing its use, so the first IDs would map to GPUs. Making a long story short, this is only about preserving information so that codegen can use it.

In any case, this is going to change in the future as the OpenMP language committee is working on having a device identifier to use instead of an integer. So, if you prefer remove the ordered out of the name, I am not opposed to that.

385–393

This has changed a little in recent CUDA work, in the version http://reviews.llvm.org/D18171 is based on, Result is returned instead of being passed by reference, and we have a `string/action-result map. I'll have to add to that string the offloading kind eventually, but in the partitioned patches I didn't touch that yet.

Do you suggest having that cache owned by the driver instead of passing it along?

lib/Driver/Compilation.cpp

66–67

This relates in some extend to your other question: how do we pass device-specific options.

So, right now we are relying on the host options to derive device-specific options. This hook was meant to make the tuning of the host options so that things that do not make sense on the device are filtered. Also, the device resulting image is usually a shared library so it that can be easily loaded, this hook is also used to specify the options that result in a shared library, even if the host options don't ask for a host shared library.

Can you think of a better way to abstract this?

lib/Driver/Driver.cpp

222–223

Given the feedback I got in http://lists.llvm.org/pipermail/cfe-dev/2016-February/047547.html, I end up moving most the functionality that I have in jobs creation to the creation of actions. Having a action graph that shows the offloading specifics was desired feature. As a result, what gets more complex is the dump of the actions.

In http://reviews.llvm.org/D18171 I have an example on how that dump looks like. That patch also proposes a unified offloading action that should be reused by the different offloading programming models. Does this address your concern?

1773–1774

I agree. This function is really messy... :S

In http://reviews.llvm.org/D18171 I am proposing collapseOffloadingAction that drives the collapsing of offload actions and abstracts some of the complexity in selectToolForJob. Do you think that goes in the right direction, or you think I should do something else?

lib/Driver/Tools.cpp

5374

Yes, sure, I proposed an implementation of the bundler, using a generic format in http://reviews.llvm.org/D13909. Let me know any comments you have about that specific component.

I still need to add testing specific to http://reviews.llvm.org/D13909, which I didn't yet because I didn't know where it was supposed to live - maybe in the Driver? Do you have an opinion about that?

Also, in http://lists.llvm.org/pipermail/cfe-dev/2016-February/047547.html, the generic opinion was that the bundler should use the host object format to bundle whenever possible. So, I also have to add a default behavior for the binary bundler when the input is an object file. For the other input types, I don't think there were any strong opinions. Do you happen to have one?

In any case, I was planing to add the object file specific bundling in a separate patch, which seems to me a natural way to partition the bundler functionality. Does that sound like a good plan?

test/OpenMP/target_driver.c

42–48

Using a DAG seemed to me a robust way to test that. I'd have to double check, but several map containers are used for the inputs and actions, so the order may depend on the implementation of the container. I was just trying to use a safe way to test.

Do you prefer to change this to the exact sequence I am getting?

55

Well, currently I don't. In http://lists.llvm.org/pipermail/cfe-dev/2016-February/047547.html I was proposing something to tackle that, but the opinion was that it was somewhat secondary and the driver design should be settled first.

What I as proposing was some sort of group option associated with the device triple. The idea was to avoid proliferation of device specific options and reuse what we already have, just organize it groups so that i could be forwarded to the right tool chain. The goal was to make things like this possible:

clang -mcpu=pwr8 -target-offload=nvptx64-nvidia-cuda -fopenmp -mcpu=sm_35 -target-offload=nvptx64-nvidia-cuda -fcuda -mcpu=sm_32 a.c

... where mcpu is used to specify the cpu/gpu for the different tool chains and programing models. This would also be useful to specify include and library paths that only make sense to the device.

Do you have any opinion about that?

[Hahnfeld]

I think these changes have been contributed to trunk in multiple commits so this can be closed?

[sfantao]

Hi Jonas,

In D9888#581809, @Hahnfeld wrote:

I think these changes have been contributed to trunk in multiple commits so this can be closed?

You're right, this can be closed now.

Thanks!
Samuel

include/clang/Basic/DiagnosticDriverKinds.td

def err_drv_mg_requires_m_or_mm : Error<
   "option '-MG' requires '-M' or '-MM'">;
 def err_drv_unknown_objc_runtime : Error<
   "unknown or ill-formed Objective-C runtime '%0'">;
 def err_drv_emit_llvm_link : Error<
    "-emit-llvm cannot be used when linking">;
 def err_drv_optimization_remark_pattern : Error<
   "%0 in '%1'">;
 def err_drv_no_neon_modifier : Error<"[no]neon is not accepted as modifier, please use [no]simd instead">;
+def err_drv_invalid_omp_target : Error<
+  "OpenMP target is invalid: '%0'">;
+def err_drv_omp_target_toolchain_not_available : Error<
+  "Toolchain for target '%0' is not supporting OpenMP offloading.">;
+def warn_drv_target_file_found : Warning<
+  "OpenMP target file '%0' is being implicitly used in the '%1' toolchain.">, 
+  InGroup<OpenMPOffloading>;
 
 def warn_O4_is_O3 : Warning<"-O4 is equivalent to -O3">, InGroup<Deprecated>;
 def warn_drv_optimization_value : Warning<"optimization level '%0' is not supported; using '%1%2' instead">,
   InGroup<InvalidCommandLineArgument>;
 def warn_ignored_gcc_optimization : Warning<"optimization flag '%0' is not supported">,
   InGroup<IgnoredOptimizationArgument>;
 def warn_c_kext : Warning<
   "ignoring -fapple-kext which is valid for C++ and Objective-C++ only">;

include/clang/Basic/DiagnosticGroups.td

 def ASM : DiagGroup<"asm", [
     ASMOperandWidths
   ]>;
 
 // OpenMP warnings.
 def SourceUsesOpenMP : DiagGroup<"source-uses-openmp">;
 def OpenMPClauses : DiagGroup<"openmp-clauses">;
 def OpenMPLoopForm : DiagGroup<"openmp-loop-form">;
+def OpenMPOffloading : DiagGroup<"openmp-offloading">;
 
 // Backend warnings.
 def BackendInlineAsm : DiagGroup<"inline-asm">;
 def BackendFrameLargerThanEQ : DiagGroup<"frame-larger-than=">;
 def BackendPlugin : DiagGroup<"backend-plugin">;
 def RemarkBackendPlugin : DiagGroup<"remark-backend-plugin">;
 def BackendOptimizationRemark : DiagGroup<"pass">;
 def BackendOptimizationRemarkMissed : DiagGroup<"pass-missed">;

include/clang/Driver/Action.h

 public:
   typedef ActionList::size_type size_type;
   typedef ActionList::iterator iterator;
   typedef ActionList::const_iterator const_iterator;
 
   enum ActionClass {
     InputClass = 0,
     BindArchClass,
+    BindTargetClass,
     PreprocessJobClass,
     PrecompileJobClass,
     AnalyzeJobClass,
     MigrateJobClass,
     CompileJobClass,
     BackendJobClass,
     AssembleJobClass,
     LinkJobClass,
private:
 
   /// The output type of this action.
   types::ID Type;
 
   ActionList Inputs;
 
   unsigned OwnsInputs : 1;
 
+  /// Is this action referring to the main host or an OpenMP offloading device
+  const char *OffloadingDevice;
+
 protected:
   Action(ActionClass _Kind, types::ID _Type)
-    : Kind(_Kind), Type(_Type), OwnsInputs(true)  {}
+      : Kind(_Kind), Type(_Type), OwnsInputs(true), OffloadingDevice(0) {}
   Action(ActionClass _Kind, std::unique_ptr<Action> Input, types::ID _Type)
-      : Kind(_Kind), Type(_Type), Inputs(1, Input.release()), OwnsInputs(true) {
-  }
+      : Kind(_Kind), Type(_Type), Inputs(1, Input.release()), OwnsInputs(true),
+        OffloadingDevice(0) {}
   Action(ActionClass _Kind, std::unique_ptr<Action> Input)
       : Kind(_Kind), Type(Input->getType()), Inputs(1, Input.release()),
-        OwnsInputs(true) {}
+        OwnsInputs(true), OffloadingDevice(0) {}
   Action(ActionClass _Kind, const ActionList &_Inputs, types::ID _Type)
-    : Kind(_Kind), Type(_Type), Inputs(_Inputs), OwnsInputs(true) {}
+      : Kind(_Kind), Type(_Type), Inputs(_Inputs), OwnsInputs(true),
+        OffloadingDevice(0) {}
+
 public:
   virtual ~Action();
 
   const char *getClassName() const { return Action::getClassName(getKind()); }
 
-  bool getOwnsInputs() { return OwnsInputs; }
+  bool getOwnsInputs() const { return OwnsInputs; }
   void setOwnsInputs(bool Value) { OwnsInputs = Value; }
 
+  const char *getOffloadingDevice() const { return OffloadingDevice; }
+  void setOffloadingDevice(const char *Value) { OffloadingDevice = Value; }
+
   ActionClass getKind() const { return Kind; }
   types::ID getType() const { return Type; }
 
   ActionList &getInputs() { return Inputs; }
   const ActionList &getInputs() const { return Inputs; }
 
   size_type size() const { return Inputs.size(); }
 
public:
 
   const char *getArchName() const { return ArchName; }
 
   static bool classof(const Action *A) {
     return A->getKind() == BindArchClass;
   }
 };
 
+class BindTargetAction : public Action {
+  virtual void anchor();
+  /// The offloading target to bind, or 0 if the default (host) architecture
+  /// should be bound.
+  const char *TargetName;
+
+public:
+  BindTargetAction(std::unique_ptr<Action> Input, const char *_TargetName);
+
+  const char *getTargetName() const { return TargetName; }
+
+  static bool classof(const Action *A) {
+    return A->getKind() == BindTargetClass;
+  }
+};
+
 class JobAction : public Action {
   virtual void anchor();
 protected:
   JobAction(ActionClass Kind, std::unique_ptr<Action> Input, types::ID Type);
   JobAction(ActionClass Kind, const ActionList &Inputs, types::ID Type);
 
 public:
   static bool classof(const Action *A) {
static bool classof(const Action *A) {
     return A->getKind() == MigrateJobClass;
   }
 };
 
 class CompileJobAction : public JobAction {
   void anchor() override;
 public:
   CompileJobAction(std::unique_ptr<Action> Input, types::ID OutputType);
+  // If offloading is being used, the compiler phase will take the host
+  // IR file as input in addition to the preprocessed source file. In this
+  // case, a list of inputs is required instead of a single action.
+  CompileJobAction(ActionList &Inputs, types::ID OutputType);
 
   static bool classof(const Action *A) {
     return A->getKind() == CompileJobClass;
   }
 };
 
 class BackendJobAction : public JobAction {
   void anchor() override;

include/clang/Driver/Compilation.h

class Compilation {
 
   /// The list of actions.
   ActionList Actions;
 
   /// The root list of jobs.
   JobList Jobs;
 
   /// Cache of translated arguments for a particular tool chain and bound
-  /// architecture.
-  llvm::DenseMap<std::pair<const ToolChain *, const char *>,
+  /// architecture. We use the second element of the pair to record whether
+  /// the toolchain belongs to an offloading target, as the same toolchain
+  /// may be used for host and targets, each using different arguments.
+  typedef std::pair<const ToolChain *, int> ToolChainWithTargetInfo;
+  llvm::DenseMap<std::pair<ToolChainWithTargetInfo, const char *>,
                  llvm::opt::DerivedArgList *> TCArgs;
 
   /// Temporary files which should be removed on exit.
   llvm::opt::ArgStringList TempFiles;
 
   /// Result files which should be removed on failure.
   ArgStringMap ResultFiles;
 
public:
 
   /// Returns the sysroot path.
   StringRef getSysRoot() const;
 
   /// getArgsForToolChain - Return the derived argument list for the
   /// tool chain \p TC (or the default tool chain, if TC is not specified).
   ///
   /// \param BoundArch - The bound architecture name, or 0.
+  /// \param isOpenMPTarget - True if this tool chain refer to an OpenMP target
+  /// \param isSuccess - set to true if the arguments were successfully obtained
   const llvm::opt::DerivedArgList &getArgsForToolChain(const ToolChain *TC,
-                                                       const char *BoundArch);
+                                                       const char *BoundArch,
+                                                       bool isOpenMPTarget,
+                                                       bool &isSuccess);
 
   /// addTempFile - Add a file to remove on exit, and returns its
   /// argument.
   const char *addTempFile(const char *Name) {
     TempFiles.push_back(Name);
     return Name;
   }
 

include/clang/Driver/Driver.h

 #ifndef LLVM_CLANG_DRIVER_DRIVER_H
 #define LLVM_CLANG_DRIVER_DRIVER_H
 
 #include "clang/Basic/Diagnostic.h"
 #include "clang/Basic/LLVM.h"
 #include "clang/Driver/Phases.h"
 #include "clang/Driver/Types.h"
 #include "clang/Driver/Util.h"
+#include "llvm/ADT/DenseMap.h"
 #include "llvm/ADT/StringMap.h"
 #include "llvm/ADT/StringRef.h"
 #include "llvm/ADT/Triple.h"
 #include "llvm/Support/Path.h" // FIXME: Kill when CompilationInfo
 #include <memory>
                               // lands.
 #include <list>
 #include <set>
private:
   std::list<std::string> ResultFiles;
 
   /// \brief Cache of all the ToolChains in use by the driver.
   ///
   /// This maps from the string representation of a triple to a ToolChain
   /// created targeting that triple. The driver owns all the ToolChain objects
   /// stored in it, and will clean them up when torn down.
   mutable llvm::StringMap<ToolChain *> ToolChains;
+  mutable llvm::StringMap<ToolChain *> ToolChainsOffloading;
+
+  /// \brief ToolChain used by the host
+  const ToolChain *HostToolChain;
+
+  /// \brief Array of the triples of offloading targets in the order they were
+  /// requested
+  /// by the user.
+  SmallVector<const char *, 4> OffloadingTriples;
+
+  /// \brief Cache that records the outputs obtained for a given action in a
+  /// given toolchain to avoid computing that multiple times
+  typedef llvm::DenseMap<const ToolChain *, InputInfo>
+      OutputsForActionPerToolChainTy;
+  typedef llvm::DenseMap<const Action *, OutputsForActionPerToolChainTy>
+      OutputsForActionCacheTy;
+  mutable OutputsForActionCacheTy OutputsForActionCache;
 
 private:
   /// TranslateInputArgs - Create a new derived argument list from the input
   /// arguments, after applying the standard argument translations.
   llvm::opt::DerivedArgList *
   TranslateInputArgs(const llvm::opt::InputArgList &Args) const;
 
   // getFinalPhase - Determine which compilation mode we are in and record 
   // which option we used to determine the final phase.
   phases::ID getFinalPhase(const llvm::opt::DerivedArgList &DAL,
                            llvm::opt::Arg **FinalPhaseArg = nullptr) const;
 
   // Before executing jobs, sets up response files for commands that need them.
   void setUpResponseFiles(Compilation &C, Job &J);
 
   void generatePrefixedToolNames(const char *Tool, const ToolChain &TC,
                                  SmallVectorImpl<std::string> &Names) const;
 
 public:
   Driver(StringRef _ClangExecutable,
          StringRef _DefaultTargetTriple,
          DiagnosticsEngine &_Diags);
public:
   /// treated before building actions or binding tools.
   ///
   /// \return Whether any compilation should be built for this
   /// invocation.
   bool HandleImmediateArgs(const Compilation &C);
 
   /// ConstructAction - Construct the appropriate action to do for
   /// \p Phase on the \p Input, taking in to account arguments
-  /// like -fsyntax-only or --analyze.
+  /// like -fsyntax-only or --analyze as well as any dependency
+  /// related with offloading (host-target dependencies).
   std::unique_ptr<Action>
   ConstructPhaseAction(const ToolChain &TC, const llvm::opt::ArgList &Args,
-                       phases::ID Phase, std::unique_ptr<Action> Input) const;
+                       phases::ID Phase, std::unique_ptr<Action> Input,
+                       std::unique_ptr<Action> OffloadingDepInput) const;
 
   /// BuildJobsForAction - Construct the jobs to perform for the
   /// action \p A.
   void BuildJobsForAction(Compilation &C,
                           const Action *A,
                           const ToolChain *TC,
                           const char *BoundArch,
                           bool AtTopLevel,
                           bool MultipleArchs,
                           const char *LinkingOutput,
                           InputInfo &Result) const;
 
   /// Returns the default name for linked images (e.g., "a.out").
   const char *getDefaultImageName() const;
 
   /// GetNamedOutputPath - Return the name to use for the output of
   /// the action \p JA. The result is appended to the compilation's
   /// list of temporary or result files, as appropriate.
   ///
   /// \param C - The compilation.
   /// \param JA - The action of interest.
   /// \param BaseInput - The original input file that this action was
   /// triggered by.
   /// \param BoundArch - The bound architecture. 
   /// \param AtTopLevel - Whether this is a "top-level" action.
   /// \param MultipleArchs - Whether multiple -arch options were supplied.
-  const char *GetNamedOutputPath(Compilation &C,
-                                 const JobAction &JA,
-                                 const char *BaseInput,
-                                 const char *BoundArch,
-                                 bool AtTopLevel,
-                                 bool MultipleArchs) const;
+  /// \param HasOffloadingTargetSuffixApended - Whether the base input was
+  /// already appended with an offloading target suffix.
+  const char *GetNamedOutputPath(Compilation &C, const JobAction &JA,
+                                 const char *BaseInput, const char *BoundArch,
+                                 bool AtTopLevel, bool MultipleArchs,
+                                 bool HasOffloadingTargetSuffixApended) const;
 
   /// GetTemporaryPath - Return the pathname of a temporary file to use 
   /// as part of compilation; the file will have the given prefix and suffix.
   ///
   /// GCC goes to extra lengths here to be a bit more robust.
   std::string GetTemporaryPath(StringRef Prefix, const char *Suffix) const;
 
   /// ShouldUseClangCompiler - Should the clang compiler be used to
   /// handle this action.
   bool ShouldUseClangCompiler(const JobAction &JA) const;
 
   bool IsUsingLTO(const ToolChain &TC, const llvm::opt::ArgList &Args) const;
 
 private:
   /// \brief Retrieves a ToolChain for a particular target triple.
   ///
   /// Will cache ToolChains for the life of the driver object, and create them
-  /// on-demand.
+  /// on-demand. If \p TripleString is provided, the triple is obtained
+  /// from it.
   const ToolChain &getToolChain(const llvm::opt::ArgList &Args,
-                                StringRef DarwinArchName = "") const;
+                                StringRef DarwinArchName = "",
+                                const char *TripleString = nullptr) const;
 
   /// @}
 
   /// \brief Get bitmasks for which option flags to include and exclude based on
   /// the driver mode.
   std::pair<unsigned, unsigned> getIncludeExcludeOptionFlagMasks() const;
 
 public:

include/clang/Driver/Options.td

def Wnonportable_cfstrings : Joined<["-"], "Wnonportable-cfstrings">, Group<W_Group>,
   Flags<[CC1Option]>;
 def Wp_COMMA : CommaJoined<["-"], "Wp,">,
   HelpText<"Pass the comma separated arguments in <arg> to the preprocessor">,
   MetaVarName<"<arg>">;
 def Wwrite_strings : Flag<["-"], "Wwrite-strings">, Group<W_Group>, Flags<[CC1Option]>;
 def Wno_write_strings : Flag<["-"], "Wno-write-strings">, Group<W_Group>, Flags<[CC1Option]>;
 def W_Joined : Joined<["-"], "W">, Group<W_Group>, Flags<[CC1Option, CoreOption]>,
   MetaVarName<"<warning>">, HelpText<"Enable the specified warning">;
+def Womp_implicit_target_files : Flag<["-"], "Womp-implicit-target-files">, Group<W_Group>, 
+  Flags<[CC1Option, NoArgumentUnused]>;
 def Xanalyzer : Separate<["-"], "Xanalyzer">,
   HelpText<"Pass <arg> to the static analyzer">, MetaVarName<"<arg>">;
 def Xarch__ : JoinedAndSeparate<["-"], "Xarch_">, Flags<[DriverOption]>;
 def Xassembler : Separate<["-"], "Xassembler">,
   HelpText<"Pass <arg> to the assembler">, MetaVarName<"<arg>">;
 def Xclang : Separate<["-"], "Xclang">,
   HelpText<"Pass <arg> to the clang compiler">, MetaVarName<"<arg>">,
   Flags<[DriverOption, CoreOption]>;

 def nostdinc : Flag<["-"], "nostdinc">;
 def nostdlibinc : Flag<["-"], "nostdlibinc">;
 def nostdincxx : Flag<["-"], "nostdinc++">, Flags<[CC1Option]>,
   HelpText<"Disable standard #include directories for the C++ standard library">;
 def nostdlib : Flag<["-"], "nostdlib">;
 def object : Flag<["-"], "object">;
 def o : JoinedOrSeparate<["-"], "o">, Flags<[DriverOption, RenderAsInput, CC1Option, CC1AsOption]>,
   HelpText<"Write output to <file>">, MetaVarName<"<file>">;
+def omptargets_EQ : CommaJoined<["-"], "omptargets=">, Flags<[DriverOption, CC1Option]>;
 def pagezero__size : JoinedOrSeparate<["-"], "pagezero_size">;
 def pass_exit_codes : Flag<["-", "--"], "pass-exit-codes">, Flags<[Unsupported]>;
 def pedantic_errors : Flag<["-", "--"], "pedantic-errors">, Group<pedantic_Group>, Flags<[CC1Option]>;
 def pedantic : Flag<["-", "--"], "pedantic">, Group<pedantic_Group>, Flags<[CC1Option]>;
 def pg : Flag<["-"], "pg">, HelpText<"Enable mcount instrumentation">, Flags<[CC1Option]>;
 def pipe : Flag<["-", "--"], "pipe">,
   HelpText<"Use pipes between commands, when possible">;
 def prebind__all__twolevel__modules : Flag<["-"], "prebind_all_twolevel_modules">;
 def prebind : Flag<["-"], "prebind">;

include/clang/Driver/ToolChain.h

private:
   mutable std::unique_ptr<Tool> Link;
   Tool *getClang() const;
   Tool *getAssemble() const;
   Tool *getLink() const;
   Tool *getClangAs() const;
 
   mutable std::unique_ptr<SanitizerArgs> SanitizerArguments;
 
+  /// This is set to true when the toolchain is created if it refers to an
+  /// offloading target toolchain
+  unsigned IsOffloadingTargetToolchain : 1;
+
 protected:
   MultilibSet Multilibs;
 
   ToolChain(const Driver &D, const llvm::Triple &T,
-            const llvm::opt::ArgList &Args);
+            const llvm::opt::ArgList &Args,
+            bool IsOffloadingTargetToolchain = false);
 
   virtual Tool *buildAssembler() const;
   virtual Tool *buildLinker() const;
   virtual Tool *getTool(Action::ActionClass AC) const;
 
   /// \name Utilities for implementing subclasses.
   ///@{
   static void addSystemInclude(const llvm::opt::ArgList &DriverArgs,

 
 public:
   virtual ~ToolChain();
 
   // Accessors
 
   const Driver &getDriver() const;
   const llvm::Triple &getTriple() const { return Triple; }
+  bool isOffloadingTargetToolchain() const {
+    return IsOffloadingTargetToolchain;
+  }
 
   llvm::Triple::ArchType getArch() const { return Triple.getArch(); }
   StringRef getArchName() const { return Triple.getArchName(); }
   StringRef getPlatform() const { return Triple.getVendorName(); }
   StringRef getOS() const { return Triple.getOSName(); }
 
   /// \brief Provide the default architecture name (as expected by -arch) for
   /// this toolchain. Note t
public:
 
   // Tool access.
 
   /// TranslateArgs - Create a new derived argument list for any argument
   /// translations this ToolChain may wish to perform, or 0 if no tool chain
   /// specific translations are needed.
   ///
   /// \param BoundArch - The bound architecture name, or 0.
+  /// \param isOffloadingTarget - True if this toolchain is an offloading
+  /// target.
+  /// \param isSuccess - set to True if the arguments were successfully
+  /// translated.
   virtual llvm::opt::DerivedArgList *
-  TranslateArgs(const llvm::opt::DerivedArgList &Args,
-                const char *BoundArch) const {
-    return nullptr;
-  }
+  TranslateArgs(const llvm::opt::DerivedArgList &Args, const char *BoundArch,
+                bool isOffloadingTarget, bool &isSuccess) const;
 
   /// Choose a tool to use to handle the action \p JA.
   Tool *SelectTool(const JobAction &JA) const;
 
   // Helper methods
 
   std::string GetFilePath(const char *Name) const;
   std::string GetProgramPath(const char *Name) const;

include/clang/Driver/Types.def

 TYPE("ast",                      AST,          INVALID,         "ast",   "u")
 TYPE("pcm",                      ModuleFile,   INVALID,         "pcm",   "u")
 TYPE("plist",                    Plist,        INVALID,         "plist", "")
 TYPE("rewritten-objc",           RewrittenObjC,INVALID,         "cpp",   "")
 TYPE("rewritten-legacy-objc",    RewrittenLegacyObjC,INVALID,   "cpp",   "")
 TYPE("remap",                    Remap,        INVALID,         "remap", "")
 TYPE("precompiled-header",       PCH,          INVALID,         "gch",   "A")
 TYPE("object",                   Object,       INVALID,         "o",     "")
+TYPE("shared-object",            SObject,      INVALID,         "so",    "")
 TYPE("treelang",                 Treelang,     INVALID,         nullptr, "u")
 TYPE("image",                    Image,        INVALID,         "out",   "")
 TYPE("dSYM",                     dSYM,         INVALID,         "dSYM",  "A")
 TYPE("dependencies",             Dependencies, INVALID,         "d",     "")
 TYPE("none",                     Nothing,      INVALID,         nullptr, "u")

lib/Driver/Action.cpp

for (iterator it = begin(), ie = end(); it != ie; ++it)
       delete *it;
   }
 }
 
 const char *Action::getClassName(ActionClass AC) {
   switch (AC) {
   case InputClass: return "input";
   case BindArchClass: return "bind-arch";
+  case BindTargetClass:
+    return "bind-target";
   case PreprocessJobClass: return "preprocessor";
   case PrecompileJobClass: return "precompiler";
   case AnalyzeJobClass: return "analyzer";
   case MigrateJobClass: return "migrator";
   case CompileJobClass: return "compiler";
   case BackendJobClass: return "backend";
   case AssembleJobClass: return "assembler";
   case LinkJobClass: return "linker";

 }
 
 void BindArchAction::anchor() {}
 
 BindArchAction::BindArchAction(std::unique_ptr<Action> Input,
                                const char *_ArchName)
     : Action(BindArchClass, std::move(Input)), ArchName(_ArchName) {}
 
+void BindTargetAction::anchor() {}
+
+BindTargetAction::BindTargetAction(std::unique_ptr<Action> Input,
+                                   const char *_TargetName)
+    : Action(BindTargetClass, std::move(Input)), TargetName(_TargetName) {
+  setOffloadingDevice(_TargetName);
+}
+
 void JobAction::anchor() {}
 
 JobAction::JobAction(ActionClass Kind, std::unique_ptr<Action> Input,
                      types::ID Type)
     : Action(Kind, std::move(Input), Type) {}
 
 JobAction::JobAction(ActionClass Kind, const ActionList &Inputs, types::ID Type)
   : Action(Kind, Inputs, Type) {
MigrateJobAction::MigrateJobAction(std::unique_ptr<Action> Input,
     : JobAction(MigrateJobClass, std::move(Input), OutputType) {}
 
 void CompileJobAction::anchor() {}
 
 CompileJobAction::CompileJobAction(std::unique_ptr<Action> Input,
                                    types::ID OutputType)
     : JobAction(CompileJobClass, std::move(Input), OutputType) {}
 
+CompileJobAction::CompileJobAction(ActionList &Inputs, types::ID OutputType)
+    : JobAction(CompileJobClass, Inputs, OutputType) {}
+
 void BackendJobAction::anchor() {}
 
 BackendJobAction::BackendJobAction(std::unique_ptr<Action> Input,
                                    types::ID OutputType)
     : JobAction(BackendJobClass, std::move(Input), OutputType) {}
 
 void AssembleJobAction::anchor() {}
 

lib/Driver/Compilation.cpp

: TheDriver(D), DefaultToolChain(_DefaultToolChain), Args(_Args),
       TranslatedArgs(_TranslatedArgs), Redirects(nullptr),
       ForDiagnostics(false) {}
 
 Compilation::~Compilation() {
   delete TranslatedArgs;
   delete Args;
 
   // Free any derived arg lists.
-  for (llvm::DenseMap<std::pair<const ToolChain*, const char*>,
-                      DerivedArgList*>::iterator it = TCArgs.begin(),
-         ie = TCArgs.end(); it != ie; ++it)
+  for (llvm::DenseMap<std::pair<ToolChainWithTargetInfo, const char *>,
+                      DerivedArgList *>::iterator it = TCArgs.begin(),
+                                                  ie = TCArgs.end();
+       it != ie; ++it)
     if (it->second != TranslatedArgs)
       delete it->second;
 
   // Free the actions, if built.
   for (ActionList::iterator it = Actions.begin(), ie = Actions.end();
        it != ie; ++it)
     delete *it;
 
   // Free redirections of stdout/stderr.
   if (Redirects) {
     delete Redirects[1];
     delete Redirects[2];
     delete [] Redirects;
   }
 }
 
 const DerivedArgList &Compilation::getArgsForToolChain(const ToolChain *TC,
-                                                       const char *BoundArch) {
+                                                       const char *BoundArch,
+                                                       bool isOffloadingTarget,
+                                                       bool &isSuccess) {
   if (!TC)
     TC = &DefaultToolChain;
 
-  DerivedArgList *&Entry = TCArgs[std::make_pair(TC, BoundArch)];
-  if (!Entry) {
-    Entry = TC->TranslateArgs(*TranslatedArgs, BoundArch);
-    if (!Entry)
-      Entry = TranslatedArgs;
-  }
+  // Try to locate the args in the toolchain map
+  llvm::DenseMap<std::pair<ToolChainWithTargetInfo, const char *>,
+                 llvm::opt::DerivedArgList *>::iterator it =
+      TCArgs.find(
+          std::make_pair(std::make_pair(TC, isOffloadingTarget), BoundArch));
+
+  // We have already tried to translate these args
+  if (it != TCArgs.end()) {
+    // if null, we already failed to translate
+    isSuccess = (it->second);
+    return *(it->second);
+  }
+
+  // This is the first attempt to translate these arguments
+  DerivedArgList *&Entry =
+      TCArgs[std::make_pair(std::make_pair(TC, isOffloadingTarget), BoundArch)];
+
+  Entry = TC->TranslateArgs(*TranslatedArgs, BoundArch, isOffloadingTarget,
+                            isSuccess);
+  // If we didn't get derived arguments but the translation is signaled as
+  // successful, we should used the untranslated arguments
+  if (!Entry && isSuccess)
 
   return *Entry;
 }
 
 bool Compilation::CleanupFile(const char *File, bool IssueErrors) const {
   // FIXME: Why are we trying to remove files that we have not created? For
   // example we should only try to remove a temporary assembly file if
   // "clang -cc1" succeed in writing it. Was this a workaround for when

lib/Driver/Driver.cpp

Driver::Driver(StringRef ClangExecutable, StringRef DefaultTargetTriple,
     : Opts(createDriverOptTable()), Diags(Diags), Mode(GCCMode),
       SaveTemps(SaveTempsNone), ClangExecutable(ClangExecutable),
       SysRoot(DEFAULT_SYSROOT), UseStdLib(true),
       DefaultTargetTriple(DefaultTargetTriple),
       DriverTitle("clang LLVM compiler"), CCPrintOptionsFilename(nullptr),
       CCPrintHeadersFilename(nullptr), CCLogDiagnosticsFilename(nullptr),
       CCCPrintBindings(false), CCPrintHeaders(false), CCLogDiagnostics(false),
       CCGenDiagnostics(false), CCCGenericGCCName(""), CheckInputsExist(true),
-      CCCUsePCH(true), SuppressMissingInputWarning(false) {
+      CCCUsePCH(true), SuppressMissingInputWarning(false),
+      HostToolChain(nullptr) {
 
   Name = llvm::sys::path::filename(ClangExecutable);
   Dir  = llvm::sys::path::parent_path(ClangExecutable);
 
   // Compute the path to the resource directory.
   StringRef ClangResourceDir(CLANG_RESOURCE_DIR);
   SmallString<128> P(Dir);
   if (ClangResourceDir != "") {
     llvm::sys::path::append(P, ClangResourceDir);
   } else {
     StringRef ClangLibdirSuffix(CLANG_LIBDIR_SUFFIX);
     llvm::sys::path::append(P, "..", Twine("lib") + ClangLibdirSuffix, "clang",
                             CLANG_VERSION_STRING);
   }
   ResourceDir = P.str();
 }
 
 Driver::~Driver() {
   delete Opts;
 
   llvm::DeleteContainerSeconds(ToolChains);
+  llvm::DeleteContainerSeconds(ToolChainsOffloading);
 }
 
 void Driver::ParseDriverMode(ArrayRef<const char *> Args) {
   const std::string OptName =
     getOpts().getOption(options::OPT_driver_mode).getPrefixedName();
 
   for (size_t I = 0, E = Args.size(); I != E; ++I) {
     // Ingore nullptrs, they are response file's EOL markers
InputArgList *Driver::ParseArgStrings(ArrayRef<const char *> ArgStrings) {
 
   // Check for unsupported options.
   for (const Arg *A : *Args) {
     if (A->getOption().hasFlag(options::Unsupported)) {
       Diag(clang::diag::err_drv_unsupported_opt) << A->getAsString(*Args);
       continue;
     }
 
-    // Warn about -mcpu= without an argument.
-    if (A->getOption().matches(options::OPT_mcpu_EQ) &&
-        A->containsValue("")) {
+    // Warn about -mcpu= and -omptargets= without an argument.
+    if ((A->getOption().matches(options::OPT_mcpu_EQ) &&
+         A->containsValue("")) ||
+        (A->getOption().matches(options::OPT_omptargets_EQ) &&
+         !A->getNumValues())) {
       Diag(clang::diag::warn_drv_empty_joined_argument) <<
         A->getAsString(*Args);
     }
   }
 
   for (arg_iterator it = Args->filtered_begin(options::OPT_UNKNOWN),
          ie = Args->filtered_end(); it != ie; ++it) {
     Diags.Report(diag::err_drv_unknown_argument) << (*it) ->getAsString(*Args);
DerivedArgList *Driver::TranslateInputArgs(const InputArgList &Args) const {
   for (Arg *A : Args) {
     // Unfortunately, we have to parse some forwarding options (-Xassembler,
     // -Xlinker, -Xpreprocessor) because we either integrate their functionality
     // (assembler and preprocessor), or bypass a previous driver ('collect2').
 
     // Rewrite linker options, to replace --no-demangle with a custom internal
     // option.
     if ((A->getOption().matches(options::OPT_Wl_COMMA) ||
          A->getOption().matches(options::OPT_Xlinker)) &&
         A->containsValue("--no-demangle")) {
       // Add the rewritten no-demangle argument.
       DAL->AddFlagArg(A, Opts->getOption(options::OPT_Z_Xlinker__no_demangle));
 
       // Add the remaining values as Xlinker arguments.
       for (unsigned i = 0, e = A->getNumValues(); i != e; ++i)
         if (StringRef(A->getValue(i)) != "--no-demangle")
           DAL->AddSeparateArg(A, Opts->getOption(options::OPT_Xlinker),
                               A->getValue(i));
SaveTemps = llvm::StringSwitch<SaveTempsMode>(A->getValue())
                     .Case("obj", SaveTempsObj)
                     .Default(SaveTempsCwd);
   }
 
   // Perform the default argument translations.
   DerivedArgList *TranslatedArgs = TranslateInputArgs(*Args);
 
   // Owned by the host.
-  const ToolChain &TC = getToolChain(*Args);
+  HostToolChain = &getToolChain(*Args);
+
+  // Get the triples for the offloading targets if any
+  OffloadingTriples.clear();
+
+  if (const Arg *A = Args->getLastArg(options::OPT_fopenmp_EQ)) {
+    if (StringRef(A->getValue()) == "libiomp5") {
+      // check if there is any openmp target we care generating code to
+      Arg *Tgts = Args->getLastArg(options::OPT_omptargets_EQ);
+
+      // If omptargets was specified use only the required targets
+      if (Tgts && Tgts->getNumValues()) {
+        for (unsigned v = 0; v < Tgts->getNumValues(); ++v) {
+          std::string error;
+          const char *val = Tgts->getValue(v);
+
+          llvm::Triple TT(val);
+
+          // If the specified target is invalid, emit error
+          if (TT.getArch() == llvm::Triple::UnknownArch)
+            Diag(clang::diag::err_drv_invalid_omp_target) << val;
+          else {
+            OffloadingTriples.push_back(val);
+          }
+        }
+      }
+    }
+  }
 
   // The compilation takes ownership of Args.
-  Compilation *C = new Compilation(*this, TC, Args, TranslatedArgs);
+  Compilation *C = new Compilation(*this, *HostToolChain, Args, TranslatedArgs);
 
   if (!HandleImmediateArgs(*C))
     return C;
 
   // Construct the list of inputs.
   InputList Inputs;
   BuildInputs(C->getDefaultToolChain(), *TranslatedArgs, Inputs);
 
   // Construct the list of abstract actions to perform for this compilation. On
   // MachO targets this uses the driver-driver and universal actions.
-  if (TC.getTriple().isOSBinFormatMachO())
+  if (HostToolChain->getTriple().isOSBinFormatMachO())
     BuildUniversalActions(C->getDefaultToolChain(), C->getArgs(),
                           Inputs, C->getActions());
   else
     BuildActions(C->getDefaultToolChain(), C->getArgs(), Inputs,
                  C->getActions());
 
   if (CCCPrintActions) {
     PrintActions(*C);
void Driver::BuildActions(const ToolChain &TC, DerivedArgList &Args,
   if (Arg *A = Args.getLastArg(options::OPT__SLASH_o)) {
     if (A->getValue()[0] == '\0') {
       // It has to have a value.
       Diag(clang::diag::err_drv_missing_argument) << A->getSpelling() << 1;
       Args.eraseArg(options::OPT__SLASH_o);
     }
   }
 
-  // Construct the actions to perform.
-  ActionList LinkerInputs;
+  // We need an array of actions to trace the actions for the main target
+  // and each available omp target detected before
+  unsigned OffloadingTargetsNum = OffloadingTriples.size();
+  std::unique_ptr<Action> *ActionsForTarget =
+      new std::unique_ptr<Action>[1 + OffloadingTargetsNum];
+  ActionList *LinkerInputsForTarget = new ActionList[1 + OffloadingTargetsNum];
 
+  // Construct the actions to perform.
   llvm::SmallVector<phases::ID, phases::MaxNumberOfPhases> PL;
   for (unsigned i = 0, e = Inputs.size(); i != e; ++i) {
     types::ID InputType = Inputs[i].first;
     const Arg *InputArg = Inputs[i].second;
 
     PL.clear();
     types::getCompilationPhases(InputType, PL);
 
if (InitialPhase > FinalPhase) {
           << InputArg->getAsString(Args)
           << getPhaseName(InitialPhase)
           << !!FinalPhaseArg
           << (FinalPhaseArg ? FinalPhaseArg->getOption().getName() : "");
       continue;
     }
 
     // Build the pipeline for this file.
-    std::unique_ptr<Action> Current(new InputAction(*InputArg, InputType));
+
+    // Initialize with the current input
+    ActionsForTarget[0].reset(new InputAction(*InputArg, InputType));
+
+    // The number of targets we should take into account.
+    // Before the preprocessor phase there is only one and from
+    // the OffloadingToolChainNum phase on we will have N+1 targets
+    // where N is all the available omp targets
+    unsigned PreviousTotalTargets = 1;
+    unsigned TotalTargets = 1;
+
     for (SmallVectorImpl<phases::ID>::iterator
            i = PL.begin(), e = PL.end(); i != e; ++i) {
       phases::ID Phase = *i;
 
       // We are done if this step is past what the user requested.
       if (Phase > FinalPhase)
         break;
 
+      bool NothingElseToDo = false;
+
+      // keep track of the number of targets in the previous phase
+      PreviousTotalTargets = TotalTargets;
+
+      // if it is the preprocess phase or after, the total targets also contain
+      // the omp targets
+      if (Phase >= phases::Preprocess)
+        TotalTargets = 1 + OffloadingTargetsNum;
+
+      // if this is after a preprocessing phase, and only one input is available
+      // for all targets, we need to locate the target corresponding inputs.
+      // We have to make sure we only do this for input files and not other
+      // options.
+      if (Phase > phases::Preprocess && PreviousTotalTargets < TotalTargets &&
+          InputArg->getOption().getKind() == llvm::opt::Option::InputClass) {
+
+        // The unique available input has to be an input action
+        InputAction *IA = cast<InputAction>(ActionsForTarget[0].get());
+
+        const Arg &HostArg = IA->getInputArg();
+        StringRef HostFileName(HostArg.getValue(0));
+
+        // Lets find out if we have an input for each target
+        bool HaveDifferentInputsForEachTarget = true;
+        llvm::SmallVector<SmallString<128>, 4> TargetFileNames(
+            OffloadingTargetsNum);
+
+        for (unsigned tgt = 0; tgt < OffloadingTargetsNum; ++tgt) {
+          // Name of the file
+          TargetFileNames[tgt] += HostFileName;
+
+          // Target suffix
+          TargetFileNames[tgt] += ".tgt-";
+          TargetFileNames[tgt] += OffloadingTriples[tgt];
+
+          // Check if the file exists. If not, don't bother trying to find other
+          // files as will not be able to use them.
+          if (!llvm::sys::fs::exists(TargetFileNames[tgt].c_str())) {
+            HaveDifferentInputsForEachTarget = false;
+            break;
+          }
+        }
+
+        for (unsigned tgt = 0; tgt < OffloadingTargetsNum; ++tgt) {
+
+          // If we do not have an input file for each target, we use null to
+          // indicate it and potential
+          if (!HaveDifferentInputsForEachTarget) {
+            ActionsForTarget[1 + tgt].reset(nullptr);
+            continue;
+          }
+
+          // Notify the user we found a target file that we will be using
+          if (Args.hasArg(options::OPT_Womp_implicit_target_files))
+            Diag(clang::diag::warn_drv_target_file_found)
+                << TargetFileNames[tgt].c_str() << OffloadingTriples[tgt];
+
+          Arg *TargetArg = Args.MakePositionalArg(
+              &HostArg, HostArg.getOption(),
+              Args.MakeArgString(TargetFileNames[tgt].c_str()));
+
+          std::unique_ptr<Action> TgtIA;
+          TgtIA.reset(new InputAction(*TargetArg, IA->getType()));
+          TgtIA.get()->setOffloadingDevice(OffloadingTriples[tgt]);
+
+          BindTargetAction *BindedTgtIA =
+              new BindTargetAction(std::move(TgtIA), OffloadingTriples[tgt]);
+
+          ActionsForTarget[1 + tgt].reset(BindedTgtIA);
+        }
+
+        // If we have an input for each target, this is like supporting multiple
+        // targets in a previous phase. If we are ahead of the compiling phase
+        // we cannot use the host files anymore, so we need to have target
+        // files already. If the input is set to null this will result in the
+        // target actions not to be generated at all in the code below.
+        if (HaveDifferentInputsForEachTarget || Phase > phases::Compile)
+          PreviousTotalTargets = TotalTargets;
+      }
+
+      // If we have multiple OpenMP targets we use the host compile output
+      // as a target input in order to host information can be used during
+      // target codegen
+      Action *LastHostInput = nullptr;
+      Action *LastHostOutput = nullptr;
+
+      for (unsigned tgt = 0; tgt < TotalTargets; ++tgt) {
+        std::unique_ptr<Action> CurrentInput;
+
+        // Do we need to reuse the host target file for all other targets?
+        if (PreviousTotalTargets < TotalTargets) {
+          // Each target picks its input file from the host
+          if (tgt == 0)
+            CurrentInput.reset(LastHostInput = ActionsForTarget[0].release());
+          else
+            CurrentInput.reset(LastHostInput);
+
+          assert(CurrentInput.get() && "Expecting an input to be defined");
+        } else {
+          // Each target uses its own target input file
+          CurrentInput.reset(ActionsForTarget[tgt].release());
+
+          assert((tgt || CurrentInput.get()) && "Host must have an input!");
+
+          // If no action for this target is defined we just move to the next
+          // target
+          if (!CurrentInput.get())
+            continue;
+        }
+
         // Queue linker inputs.
         if (Phase == phases::Link) {
           assert((i + 1) == e && "linking must be final compilation step.");
-        LinkerInputs.push_back(Current.release());
-        break;
+          LinkerInputsForTarget[tgt].push_back(CurrentInput.release());
+          NothingElseToDo = true;
+          continue;
         }
 
         // Some types skip the assembler phase (e.g., llvm-bc), but we can't
         // encode this in the steps because the intermediate type depends on
         // arguments. Just special case here.
-      if (Phase == phases::Assemble && Current->getType() != types::TY_PP_Asm)
+        if (Phase == phases::Assemble &&
+            CurrentInput->getType() != types::TY_PP_Asm) {
+          // Just store the current input to be used directly in the next phase
+          ActionsForTarget[tgt].reset(CurrentInput.release());
           continue;
+        }
+
+        // Build the phase action. We pass the last llvm IR file that was
+        // produced for the host for an compile phase for the OpenMP targets
+        std::unique_ptr<Action> DepAction;
+        if (Phase == phases::Compile && tgt > 0 &&
+            LastHostOutput->getType() == types::TY_LLVM_BC) {
+          DepAction.reset(LastHostOutput);
+          // Bind action to the host toolchain - nullptr
+          DepAction.reset(new BindTargetAction(std::move(DepAction), nullptr));
+          DepAction.get()->setOwnsInputs(false);
+        }
+        ActionsForTarget[tgt].reset(
+            ConstructPhaseAction(TC, Args, Phase, std::move(CurrentInput),
+                                 std::move(DepAction))
+                .release());
+
+        // Save the last host output has it may be required for some OpenMP
+        // host-target dependency
+        if (tgt == 0)
+          LastHostOutput = ActionsForTarget[0].get();
+
+        if (ActionsForTarget[tgt]->getType() == types::TY_Nothing)
+          NothingElseToDo = true;
+
+        // If we are coming from a single to multiple target phase set the input
+        // ownership flag to false.
+        // Also, if this is a compile phase for the OpenMP host and we are
+        // supporting multiple targets, we also need to set this to false
+        ActionsForTarget[tgt]->setOwnsInputs(
+            !((Phase == phases::Compile && !tgt && TotalTargets > 1 &&
+               ActionsForTarget[0].get()->getType() == types::TY_LLVM_BC) ||
+              (PreviousTotalTargets < TotalTargets)));
+
+        // If we are processing a target action, we need to bind it to the
+        // target it refers to
+        if (tgt > 0) {
+          ActionsForTarget[tgt]->setOffloadingDevice(
+              OffloadingTriples[tgt - 1]);
+          ActionsForTarget[tgt].reset(new BindTargetAction(
+              std::move(ActionsForTarget[tgt]), OffloadingTriples[tgt - 1]));
+        }
+      }
 
-      // Otherwise construct the appropriate action.
-      Current = ConstructPhaseAction(TC, Args, Phase, std::move(Current));
-      if (Current->getType() == types::TY_Nothing)
+      if (NothingElseToDo)
         break;
     }
 
     // If we ended with something, add to the output list.
-    if (Current)
-      Actions.push_back(Current.release());
+    for (unsigned tgt = 0; tgt < TotalTargets; ++tgt) {
+      if (ActionsForTarget[tgt])
+        Actions.push_back(ActionsForTarget[tgt].release());
+    }
+  }
+
+  // release the Actions-For-Target array
+  delete[] ActionsForTarget;
+
+  // Create link action for each target if any
+  for (unsigned tgt = 0; tgt < OffloadingTargetsNum; ++tgt) {
+
+    if (LinkerInputsForTarget[tgt + 1].empty())
+      continue;
+
+    std::unique_ptr<Action> TgtLinkAction;
+    // Link target action: produces a shared library
+    TgtLinkAction.reset(
+        new LinkJobAction(LinkerInputsForTarget[tgt + 1], types::TY_SObject));
+    TgtLinkAction.get()->setOffloadingDevice(OffloadingTriples[tgt]);
+
+    // if the target link phase takes an input that is not binded to it, it
+    // means it does not own it, as it may be used by other targets and host too
+    for (const auto &II : LinkerInputsForTarget[tgt + 1]) {
+      if (!isa<BindTargetAction>(II))
+        TgtLinkAction.get()->setOwnsInputs(false);
+    }
+
+    // Bind action to target
+    TgtLinkAction.reset(
+        new BindTargetAction(std::move(TgtLinkAction), OffloadingTriples[tgt]));
+    // Include the resulting object as part of the host linking
+    LinkerInputsForTarget[0].push_back(TgtLinkAction.release());
   }
 
   // Add a link action if necessary.
-  if (!LinkerInputs.empty())
-    Actions.push_back(new LinkJobAction(LinkerInputs, types::TY_Image));
+  if (!LinkerInputsForTarget[0].empty())
+    Actions.push_back(
+        new LinkJobAction(LinkerInputsForTarget[0], types::TY_Image));
+
+  // Release the linking input arrays
+  delete[] LinkerInputsForTarget;
 
   // If we are linking, claim any options which are obviously only used for
   // compilation.
   if (FinalPhase == phases::Link && PL.size() == 1) {
     Args.ClaimAllArgs(options::OPT_CompileOnly_Group);
     Args.ClaimAllArgs(options::OPT_cl_compile_Group);
   }
 
   // Claim ignored clang-cl options.
   Args.ClaimAllArgs(options::OPT_cl_ignored_Group);
 }
 
 std::unique_ptr<Action>
 Driver::ConstructPhaseAction(const ToolChain &TC, const ArgList &Args,
-                             phases::ID Phase,
-                             std::unique_ptr<Action> Input) const {
+                             phases::ID Phase, std::unique_ptr<Action> Input,
+                             std::unique_ptr<Action> OffloadingDepInput) const {
   llvm::PrettyStackTraceString CrashInfo("Constructing phase actions");
   // Build the appropriate action.
   switch (Phase) {
   case phases::Link: llvm_unreachable("link action invalid here.");
   case phases::Preprocess: {
     types::ID OutputTy;
     // -{M, MM} alter the output type.
     if (Args.hasArg(options::OPT_M, options::OPT_MM)) {
if (Args.hasArg(options::OPT_emit_ast))
       return llvm::make_unique<CompileJobAction>(std::move(Input),
                                                  types::TY_AST);
     if (Args.hasArg(options::OPT_module_file_info))
       return llvm::make_unique<CompileJobAction>(std::move(Input),
                                                  types::TY_ModuleFile);
     if (Args.hasArg(options::OPT_verify_pch))
       return llvm::make_unique<VerifyPCHJobAction>(std::move(Input),
                                                    types::TY_Nothing);
-    return llvm::make_unique<CompileJobAction>(std::move(Input),
-                                               types::TY_LLVM_BC);
+    // In a compile phase, an offloading dependency file may be provided. If so,
+    // we add it to the input list.
+    ActionList AL;
+    AL.push_back(Input.release());
+    if (OffloadingDepInput.get())
+      AL.push_back(OffloadingDepInput.release());
+    return llvm::make_unique<CompileJobAction>(AL, types::TY_LLVM_BC);
   }
   case phases::Backend: {
     if (IsUsingLTO(TC, Args)) {
       types::ID Output =
         Args.hasArg(options::OPT_S) ? types::TY_LTO_IR : types::TY_LTO_BC;
       return llvm::make_unique<BackendJobAction>(std::move(Input), Output);
     }
     if (Args.hasArg(options::OPT_emit_llvm)) {
if (TC.getSanitizerArgs().needsLTO())
     return true;
 
   if (Args.hasFlag(options::OPT_flto, options::OPT_fno_lto, false))
     return true;
 
   return false;
 }
 
+//// cache that records the outputs obtained for a given action in a given
+//// toolchain to avoid computing that multiple times
+// typedef llvm::DenseMap<const ToolChain *, InputInfo>
+//    OutputsForActionPerToolChainTy;
+// typedef llvm::DenseMap<const Action *, OutputsForActionPerToolChainTy>
+//    OutputsForActionCacheTy;
+// static OutputsForActionCacheTy OutputsForActionCache;
+
 void Driver::BuildJobs(Compilation &C) const {
   llvm::PrettyStackTraceString CrashInfo("Building compilation jobs");
 
   Arg *FinalOutput = C.getArgs().getLastArg(options::OPT_o);
 
   // It is an error to provide a -o option if we are making multiple output
-  // files.
+  // files, except if the action producing the output is an offloading
+  // one.
   if (FinalOutput) {
     unsigned NumOutputs = 0;
     for (const Action *A : C.getActions())
-      if (A->getType() != types::TY_Nothing)
+      if (A->getType() != types::TY_Nothing && !A->getOffloadingDevice())
         ++NumOutputs;
 
     if (NumOutputs > 1) {
       Diag(clang::diag::err_drv_output_argument_with_multiple_files);
       FinalOutput = nullptr;
     }
   }
 
   // Collect the list of architectures.
   llvm::StringSet<> ArchNames;
   if (C.getDefaultToolChain().getTriple().isOSBinFormatMachO())
     for (const Arg *A : C.getArgs())
       if (A->getOption().matches(options::OPT_arch))
         ArchNames.insert(A->getValue());
 
+  // Make sure the outputs cache is empty prior starting computing jobs and
+  // inputs.
+  OutputsForActionCache.clear();
+
   for (Action *A : C.getActions()) {
     // If we are linking an image for multiple archs then the linker wants
     // -arch_multiple and -final_output <final image name>. Unfortunately, this
     // doesn't fit in cleanly because we have to pass this information down.
     //
     // FIXME: This is a hack; find a cleaner way to integrate this into the
     // process.
     const char *LinkingOutput = nullptr;
if (!A->isClaimed()) {
       }
 
       Diag(clang::diag::warn_drv_unused_argument)
         << A->getAsString(C.getArgs());
     }
   }
 }
 
+static Action *GetPreviousSingleAction(const ActionList *&Inputs) {
+  // Determine the single action that generated the current input, if any
+  Action *PrevSingleAction = nullptr;
+  if (Inputs->size() == 1) {
+    PrevSingleAction = dyn_cast<Action>(*Inputs->begin());
+    // Bypass the binding of a target
+    if (PrevSingleAction && isa<BindTargetAction>(PrevSingleAction))
+      PrevSingleAction = *PrevSingleAction->getInputs().begin();
+  }
+  return PrevSingleAction;
+}
+
 static const Tool *SelectToolForJob(Compilation &C, bool SaveTemps,
+                                    bool isLegalToMergeCompilerAndBackend,
                                     const ToolChain *TC, const JobAction *JA,
                                     const ActionList *&Inputs) {
   const Tool *ToolForJob = nullptr;
+  Action *PrevSingleAction = GetPreviousSingleAction(Inputs);
 
   // See if we should look for a compiler with an integrated assembler. We match
   // bottom up, so what we are actually looking for is an assembler job with a
   // compiler input.
 
-  if (TC->useIntegratedAs() &&
-      !SaveTemps &&
+  if (TC->useIntegratedAs() && !SaveTemps &&
       !C.getArgs().hasArg(options::OPT_via_file_asm) &&
       !C.getArgs().hasArg(options::OPT__SLASH_FA) &&
       !C.getArgs().hasArg(options::OPT__SLASH_Fa) &&
-      isa<AssembleJobAction>(JA) &&
-      Inputs->size() == 1 && isa<BackendJobAction>(*Inputs->begin())) {
+      isa<AssembleJobAction>(JA) && PrevSingleAction &&
+      isa<BackendJobAction>(PrevSingleAction)) {
     // A BackendJob is always preceded by a CompileJob, and without
     // -save-temps they will always get combined together, so instead of
-    // checking the backend tool, check if the tool for the CompileJob
-    // has an integrated assembler.
-    const ActionList *BackendInputs = &(*Inputs)[0]->getInputs();
-    JobAction *CompileJA = cast<CompileJobAction>(*BackendInputs->begin());
+    // checking the backend tool, check if the tool for the CompileJobt
+    // has an integrated assembler. The isLegalToMergeCompileAndBackend flag
+    // has to be set so that Backend and Compile phases can integrated into a
+    // single one though.
+    if (isLegalToMergeCompilerAndBackend) {
+      const ActionList *BackendInputs = &PrevSingleAction->getInputs();
+      Action *PrevPrevSingleAction = GetPreviousSingleAction(BackendInputs);
+
+      JobAction *CompileJA = cast<CompileJobAction>(PrevPrevSingleAction);
       const Tool *Compiler = TC->SelectTool(*CompileJA);
       if (!Compiler)
         return nullptr;
       if (Compiler->hasIntegratedAssembler()) {
-      Inputs = &(*BackendInputs)[0]->getInputs();
+        Inputs = &PrevPrevSingleAction->getInputs();
+        PrevSingleAction = GetPreviousSingleAction(Inputs);
         ToolForJob = Compiler;
       }
+    } else {
+      JobAction *BackendJA = cast<BackendJobAction>(PrevSingleAction);
+      const Tool *Backend = TC->SelectTool(*BackendJA);
+      if (!Backend)
+        return nullptr;
+      if (Backend->hasIntegratedAssembler()) {
+        Inputs = &PrevSingleAction->getInputs();
+        return Backend;
+      }
+    }
   }
 
   // A backend job should always be combined with the preceding compile job
   // unless OPT_save_temps is enabled and the compiler is capable of emitting
   // LLVM IR as an intermediate output.
-  if (isa<BackendJobAction>(JA)) {
+  if (isa<BackendJobAction>(JA) && isLegalToMergeCompilerAndBackend) {
     // Check if the compiler supports emitting LLVM IR.
-    assert(Inputs->size() == 1);
-    JobAction *CompileJA = cast<CompileJobAction>(*Inputs->begin());
+    assert(PrevSingleAction);
+    JobAction *CompileJA = cast<CompileJobAction>(PrevSingleAction);
     const Tool *Compiler = TC->SelectTool(*CompileJA);
     if (!Compiler)
       return nullptr;
     if (!Compiler->canEmitIR() || !SaveTemps) {
-      Inputs = &(*Inputs)[0]->getInputs();
+      Inputs = &PrevSingleAction->getInputs();
+      PrevSingleAction = GetPreviousSingleAction(Inputs);
       ToolForJob = Compiler;
     }
   }
 
   // Otherwise use the tool for the current job.
   if (!ToolForJob)
     ToolForJob = TC->SelectTool(*JA);
 
   // See if we should use an integrated preprocessor. We do so when we have
   // exactly one input, since this is the only use case we care about
   // (irrelevant since we don't support combine yet).
-  if (Inputs->size() == 1 && isa<PreprocessJobAction>(*Inputs->begin()) &&
+  if (PrevSingleAction && isa<PreprocessJobAction>(PrevSingleAction) &&
       !C.getArgs().hasArg(options::OPT_no_integrated_cpp) &&
-      !C.getArgs().hasArg(options::OPT_traditional_cpp) &&
-      !SaveTemps &&
+      !C.getArgs().hasArg(options::OPT_traditional_cpp) && !SaveTemps &&
       !C.getArgs().hasArg(options::OPT_rewrite_objc) &&
       ToolForJob->hasIntegratedCPP())
-    Inputs = &(*Inputs)[0]->getInputs();
+    Inputs = &PrevSingleAction->getInputs();
 
   return ToolForJob;
 }
 
 void Driver::BuildJobsForAction(Compilation &C,
                                 const Action *A,
                                 const ToolChain *TC,
                                 const char *BoundArch,
                                 bool AtTopLevel,
                                 bool MultipleArchs,
                                 const char *LinkingOutput,
                                 InputInfo &Result) const {
   llvm::PrettyStackTraceString CrashInfo("Building compilation jobs");
 
+  // Find out whether we already created a job for this exact action and
+  // toolchain.
+  {
+    auto PrevOutputsPerToolchainIt =
+        OutputsForActionCache.find(const_cast<Action *>(A));
+    // Any toolchains used to generate outputs for this actions?
+    if (PrevOutputsPerToolchainIt != OutputsForActionCache.end()) {
+      auto &PrevOutputsPerToolchain = PrevOutputsPerToolchainIt->getSecond();
+      auto PrevOutputsIt =
+          PrevOutputsPerToolchain.find(const_cast<ToolChain *>(TC));
+      // Is there any output for the requested toolchain?
+      if (PrevOutputsIt != PrevOutputsPerToolchain.end()) {
+        Result = PrevOutputsIt->getSecond();
+        return;
+      }
+    }
+  }
+
   if (const InputAction *IA = dyn_cast<InputAction>(A)) {
     // FIXME: It would be nice to not claim this here; maybe the old scheme of
     // just using Args was better?
     const Arg &Input = IA->getInputArg();
     Input.claim();
+
+    // If the input action has an offloading device associated, it means that
+    // the file is already a target file and therefore has the target suffix
+    // already appended to it.
     if (Input.getOption().matches(options::OPT_INPUT)) {
       const char *Name = Input.getValue();
-      Result = InputInfo(Name, A->getType(), Name);
+      Result = InputInfo(Name, A, Name, !IA->getOffloadingDevice());
     } else
-      Result = InputInfo(&Input, A->getType(), "");
+      Result = InputInfo(&Input, A, "", !IA->getOffloadingDevice());
+
     return;
   }
 
   if (const BindArchAction *BAA = dyn_cast<BindArchAction>(A)) {
     const ToolChain *TC;
     const char *ArchName = BAA->getArchName();
 
     if (ArchName)
       TC = &getToolChain(C.getArgs(), ArchName);
     else
       TC = &C.getDefaultToolChain();
 
     BuildJobsForAction(C, *BAA->begin(), TC, BAA->getArchName(),
                        AtTopLevel, MultipleArchs, LinkingOutput, Result);
     return;
   }
 
+  if (const BindTargetAction *BTA = dyn_cast<BindTargetAction>(A)) {
+    const ToolChain *TC;
+    const char *OffloadingTargetName = BTA->getTargetName();
+
+    if (OffloadingTargetName) {
+      TC = &getToolChain(C.getArgs(), "", OffloadingTargetName);
+      // If the toolchain does not have isOffloadingTargetToolchain set it means
+      // there was some problem creating the toolchain
+      if (!TC->isOffloadingTargetToolchain()) {
+        Diag(clang::diag::err_drv_omp_target_toolchain_not_available)
+            << OffloadingTargetName;
+        return;
+      }
+    } else {
+      TC = HostToolChain;
+      OffloadingTargetName = nullptr;
+    }
+
+    BuildJobsForAction(C, *BTA->begin(), TC, OffloadingTargetName, AtTopLevel,
+                       false, LinkingOutput, Result);
+    return;
+  }
+
   const ActionList *Inputs = &A->getInputs();
 
   const JobAction *JA = cast<JobAction>(A);
-  const Tool *T = SelectToolForJob(C, isSaveTempsEnabled(), TC, JA, Inputs);
+  const Tool *T = SelectToolForJob(
+      C, isSaveTempsEnabled(),
+      /*isLegalToMergeCompilerAndBackend=*/OffloadingTriples.empty(), TC, JA,
+      Inputs);
   if (!T)
     return;
 
   // Only use pipes when there is exactly one input.
   InputInfoList InputInfos;
   for (const Action *Input : *Inputs) {
     // Treat dsymutil and verify sub-jobs as being at the top-level too, they
     // shouldn't get temporary output names.
     // FIXME: Clean this up.
     bool SubJobAtTopLevel = false;
     if (AtTopLevel && (isa<DsymutilJobAction>(A) || isa<VerifyJobAction>(A)))
       SubJobAtTopLevel = true;
 
     InputInfo II;
     BuildJobsForAction(C, Input, TC, BoundArch, SubJobAtTopLevel, MultipleArchs,
                        LinkingOutput, II);
     InputInfos.push_back(II);
   }
 
   // Always use the first input as the base input.
   const char *BaseInput = InputInfos[0].getBaseInput();
 
+  // We need to propagate the information on whether a target suffix has to be
+  // added to the name of the files.
+  bool HasTargetSuffixApended =
+      InputInfos[0].hasOffloadingTargetSuffixApended();
+
   // ... except dsymutil actions, which use their actual input as the base
   // input.
   if (JA->getType() == types::TY_dSYM)
     BaseInput = InputInfos[0].getFilename();
 
   // Determine the place to write output to, if any.
   if (JA->getType() == types::TY_Nothing)
-    Result = InputInfo(A->getType(), BaseInput);
+    Result = InputInfo(A, BaseInput, HasTargetSuffixApended);
   else
-    Result = InputInfo(GetNamedOutputPath(C, *JA, BaseInput, BoundArch,
-                                          AtTopLevel, MultipleArchs),
-                       A->getType(), BaseInput);
+    Result =
+        InputInfo(GetNamedOutputPath(C, *JA, BaseInput, BoundArch, AtTopLevel,
+                                     MultipleArchs, HasTargetSuffixApended),
+                  A, BaseInput, HasTargetSuffixApended);
+
+  OutputsForActionCache[const_cast<Action *>(A)][const_cast<ToolChain *>(TC)] =
+      Result;
 
   if (CCCPrintBindings && !CCGenDiagnostics) {
     llvm::errs() << "# \"" << T->getToolChain().getTripleString() << '"'
                  << " - \"" << T->getName() << "\", inputs: [";
     for (unsigned i = 0, e = InputInfos.size(); i != e; ++i) {
       llvm::errs() << InputInfos[i].getAsString();
       if (i + 1 != e)
         llvm::errs() << ", ";
     }
     llvm::errs() << "], output: " << Result.getAsString() << "\n";
   } else {
-    T->ConstructJob(C, *JA, Result, InputInfos,
-                    C.getArgsForToolChain(TC, BoundArch), LinkingOutput);
+
+    bool isSuccess;
+
+    // Get the derived args for the target
+    const DerivedArgList &DArgs = C.getArgsForToolChain(
+        TC, BoundArch, JA->getOffloadingDevice(), isSuccess);
+
+    // If we were unable to successfully get the arguments we cannot create the
+    // job. This will happen if using an offloading target toolchain that is not
+    // prepared to translate host arguments into target's.
+    if (!isSuccess)
+      return;
+
+    T->ConstructJob(C, *JA, Result, InputInfos, DArgs, LinkingOutput);
   }
 }
 
 const char *Driver::getDefaultImageName() const {
   llvm::Triple Target(llvm::Triple::normalize(DefaultTargetTriple));
   return Target.isOSWindows() ? "a.exe" : "a.out";
 }
 
if (!llvm::sys::path::has_extension(ArgValue)) {
     }
 
     llvm::sys::path::replace_extension(Filename, Extension);
   }
 
   return Args.MakeArgString(Filename.c_str());
 }
 
-const char *Driver::GetNamedOutputPath(Compilation &C,
-                                       const JobAction &JA,
-                                       const char *BaseInput,
-                                       const char *BoundArch,
-                                       bool AtTopLevel,
-                                       bool MultipleArchs) const {
+const char *
+Driver::GetNamedOutputPath(Compilation &C, const JobAction &JA,
+                           const char *BaseInput, const char *BoundArch,
+                           bool AtTopLevel, bool MultipleArchs,
+                           bool HasOffloadingTargetSuffixApended) const {
   llvm::PrettyStackTraceString CrashInfo("Computing output path");
   // Output to a user requested destination?
   if (AtTopLevel && !isa<DsymutilJobAction>(JA) &&
       !isa<VerifyJobAction>(JA)) {
-    if (Arg *FinalOutput = C.getArgs().getLastArg(options::OPT_o))
-      return C.addResultFile(FinalOutput->getValue(), &JA);
+    if (Arg *FinalOutput = C.getArgs().getLastArg(options::OPT_o)) {
+      SmallString<128> Suffixed(FinalOutput->getValue());
+
+      // Append offloading target suffix to output file
+      if (JA.getOffloadingDevice()) {
+        Suffixed += ".tgt-";
+        Suffixed.append(BoundArch);
+      }
+
+      return C.addResultFile(C.getArgs().MakeArgString(Suffixed.c_str()), &JA);
+    }
   }
 
   // For /P, preprocess to file named after BaseInput.
   if (C.getArgs().hasArg(options::OPT__SLASH_P)) {
     assert(AtTopLevel && isa<PreprocessJobAction>(JA));
     StringRef BaseName = llvm::sys::path::filename(BaseInput);
     StringRef NameArg;
     if (Arg *A = C.getArgs().getLastArg(options::OPT__SLASH_Fi,
if (IsCLMode()) {
       NamedOutput = C.getArgs().MakeArgString(Output.c_str());
     } else
       NamedOutput = getDefaultImageName();
   } else {
     const char *Suffix = types::getTypeTempSuffix(JA.getType(), IsCLMode());
     assert(Suffix && "All types used for output should have a suffix.");
 
     std::string::size_type End = std::string::npos;
-    if (!types::appendSuffixForType(JA.getType()))
+    if (!types::appendSuffixForType(JA.getType())) {
       End = BaseName.rfind('.');
+      // If this is an offloading target job that already contains the target
+      // suffix we need to skip both type and target suffix
+      if (JA.getOffloadingDevice() && !HasOffloadingTargetSuffixApended)
+        End = BaseName.substr(0, End).rfind('.');
+    }
     SmallString<128> Suffixed(BaseName.substr(0, End));
     if (MultipleArchs && BoundArch) {
       Suffixed += "-";
       Suffixed.append(BoundArch);
     }
     // When using both -save-temps and -emit-llvm, use a ".tmp.bc" suffix for
     // the unoptimized bitcode so that it does not get overwritten by the ".bc"
     // optimized bitcode output.
     if (!AtTopLevel && C.getArgs().hasArg(options::OPT_emit_llvm) &&
         JA.getType() == types::TY_LLVM_BC)
       Suffixed += ".tmp";
     Suffixed += '.';
     Suffixed += Suffix;
+
+    // Append offloading target suffix to output file
+    if (JA.getOffloadingDevice()) {
+      Suffixed += ".tgt-";
+      Suffixed.append(BoundArch);
+    }
+
     NamedOutput = C.getArgs().MakeArgString(Suffixed.c_str());
   }
 
   // Prepend object file path if -save-temps=obj
   if (!AtTopLevel && isSaveTempsObj() && C.getArgs().hasArg(options::OPT_o) &&
       JA.getType() != types::TY_PCH) {
     Arg *FinalOutput = C.getArgs().getLastArg(options::OPT_o);
     SmallString<128> TempPath(FinalOutput->getValue());
if (Arg *A = Args.getLastArg(options::OPT_m64, options::OPT_mx32,
 
     if (AT != llvm::Triple::UnknownArch && AT != Target.getArch())
       Target.setArch(AT);
   }
 
   return Target;
 }
 
-const ToolChain &Driver::getToolChain(const ArgList &Args,
-                                      StringRef DarwinArchName) const {
-  llvm::Triple Target = computeTargetTriple(DefaultTargetTriple, Args,
-                                            DarwinArchName);
+const ToolChain &
+Driver::getToolChain(const ArgList &Args, StringRef DarwinArchName,
+                     const char *OffloadingTripleString) const {
+
+  llvm::Triple Target;
+  ToolChain *TC;
+  bool IsOffloadingTargetToolchain = OffloadingTripleString != nullptr;
+
+  // if a specific triple is passed, that means it was already parsed
+  // before while creating the actions for offloading targets, therefore we
+  // should
+  // be able to get the architecture directly from it.
+
+  if (IsOffloadingTargetToolchain) {
+    Target = llvm::Triple(OffloadingTripleString);
+    assert(Target.getArch() != llvm::Triple::UnknownArch &&
+           "Target unknown - impossible to infer toolchain!");
+
+    TC = ToolChainsOffloading[Target.str()];
+  } else {
+    Target = computeTargetTriple(DefaultTargetTriple, Args, DarwinArchName);
+    TC = ToolChains[Target.str()];
+  }
 
-  ToolChain *&TC = ToolChains[Target.str()];
   if (!TC) {
     switch (Target.getOS()) {
     case llvm::Triple::CloudABI:
       TC = new toolchains::CloudABI(*this, Target, Args);
       break;
     case llvm::Triple::Darwin:
     case llvm::Triple::MacOSX:
     case llvm::Triple::IOS:
case llvm::Triple::FreeBSD:
       break;
     case llvm::Triple::Minix:
       TC = new toolchains::Minix(*this, Target, Args);
       break;
     case llvm::Triple::Linux:
       if (Target.getArch() == llvm::Triple::hexagon)
         TC = new toolchains::Hexagon_TC(*this, Target, Args);
       else
-        TC = new toolchains::Linux(*this, Target, Args);
+        TC = new toolchains::Linux(*this, Target, Args,
+                                   IsOffloadingTargetToolchain);
       break;
     case llvm::Triple::NaCl:
       TC = new toolchains::NaCl_TC(*this, Target, Args);
       break;
     case llvm::Triple::Solaris:
       TC = new toolchains::Solaris(*this, Target, Args);
       break;
     case llvm::Triple::Win32:
default:
         TC = new toolchains::Hexagon_TC(*this, Target, Args);
         break;
       }
       if (Target.getArch() == llvm::Triple::xcore) {
         TC = new toolchains::XCore(*this, Target, Args);
         break;
       }
       if (Target.isOSBinFormatELF()) {
-        TC = new toolchains::Generic_ELF(*this, Target, Args);
+        TC = new toolchains::Generic_ELF(*this, Target, Args,
+                                         IsOffloadingTargetToolchain);
         break;
       }
       if (Target.isOSBinFormatMachO()) {
         TC = new toolchains::MachO(*this, Target, Args);
         break;
       }
-      TC = new toolchains::Generic_GCC(*this, Target, Args);
+      TC = new toolchains::Generic_GCC(*this, Target, Args,
+                                       IsOffloadingTargetToolchain);
       break;
     }
   }
+
+  // Store the toolchain in the map in case it is requested later
+  if (IsOffloadingTargetToolchain)
+    ToolChainsOffloading[Target.str()] = TC;
+  else
+    ToolChains[Target.str()] = TC;
+
+  // If this is not an OffloadingToolchain the information in TC needs to
+  // be consistent with the flag IsOffloadingTargetToolchain
+  assert((IsOffloadingTargetToolchain ||
+          IsOffloadingTargetToolchain == TC->isOffloadingTargetToolchain()) &&
+         "Unable to initialize toolchain for offloading!!!");
+
   return *TC;
 }
 
 bool Driver::ShouldUseClangCompiler(const JobAction &JA) const {
   // Check if user requested no clang, or clang doesn't understand this type (we
-  // only handle single inputs for now).
-  if (JA.size() != 1 ||
+  // only handle single inputs for now except for ofloading target compile
+  // phases).
+  if (!(JA.size() == 1 || (JA.size() == 2 && JA.getOffloadingDevice() &&
+                           isa<CompileJobAction>(JA))) ||
       !types::isAcceptedByClang((*JA.begin())->getType()))
     return false;
 
   // Otherwise make sure this is an action clang understands.
   if (!isa<PreprocessJobAction>(JA) && !isa<PrecompileJobAction>(JA) &&
       !isa<CompileJobAction>(JA) && !isa<BackendJobAction>(JA))
     return false;
 

lib/Driver/InputInfo.h

 //===--- InputInfo.h - Input Source & Type Information ----------*- C++ -*-===//
 //
 //                     The LLVM Compiler Infrastructure
 //
 // This file is distributed under the University of Illinois Open Source
 // License. See LICENSE.TXT for details.
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef LLVM_CLANG_LIB_DRIVER_INPUTINFO_H
 #define LLVM_CLANG_LIB_DRIVER_INPUTINFO_H
 
+#include "clang/Driver/Action.h"
 #include "clang/Driver/Types.h"
 #include "llvm/Option/Arg.h"
 #include <cassert>
 #include <string>
 
 namespace clang {
 namespace driver {
 
enum Class {
     Pipe
   };
 
   union {
     const char *Filename;
     const llvm::opt::Arg *InputArg;
   } Data;
   Class Kind;
-  types::ID Type;
   const char *BaseInput;
 
+  // Action that originates this info
+  const Action *OrigAction;
+
+  // True if the filename associated with this info has the offloading target
+  // suffix appended
+  bool HasOffloadingTargetSuffixApended;
+
 public:
   InputInfo() {}
-  InputInfo(types::ID _Type, const char *_BaseInput)
-    : Kind(Nothing), Type(_Type), BaseInput(_BaseInput) {
-  }
-  InputInfo(const char *_Filename, types::ID _Type, const char *_BaseInput)
-    : Kind(Filename), Type(_Type), BaseInput(_BaseInput) {
+  InputInfo(const Action *_OrigAction, const char *_BaseInput,
+            bool _HasOffloadingTargetSuffixApended)
+      : Kind(Nothing), BaseInput(_BaseInput), OrigAction(_OrigAction),
+        HasOffloadingTargetSuffixApended(_HasOffloadingTargetSuffixApended) {}
+  InputInfo(const char *_Filename, const Action *_OrigAction,
+            const char *_BaseInput, bool _HasOffloadingTargetSuffixApended)
+      : Kind(Filename), BaseInput(_BaseInput), OrigAction(_OrigAction),
+        HasOffloadingTargetSuffixApended(_HasOffloadingTargetSuffixApended) {
     Data.Filename = _Filename;
   }
-  InputInfo(const llvm::opt::Arg *_InputArg, types::ID _Type,
-            const char *_BaseInput)
-      : Kind(InputArg), Type(_Type), BaseInput(_BaseInput) {
+  InputInfo(const llvm::opt::Arg *_InputArg, const Action *_OrigAction,
+            const char *_BaseInput, bool _HasOffloadingTargetSuffixApended)
+      : Kind(InputArg), BaseInput(_BaseInput), OrigAction(_OrigAction),
+        HasOffloadingTargetSuffixApended(_HasOffloadingTargetSuffixApended) {
     Data.InputArg = _InputArg;
   }
 
   bool isNothing() const { return Kind == Nothing; }
   bool isFilename() const { return Kind == Filename; }
   bool isInputArg() const { return Kind == InputArg; }
-  types::ID getType() const { return Type; }
+  types::ID getType() const {
+    return OrigAction ? OrigAction->getType() : types::TY_Nothing;
+  }
   const char *getBaseInput() const { return BaseInput; }
 
   const char *getFilename() const {
     assert(isFilename() && "Invalid accessor.");
     return Data.Filename;
   }
   const llvm::opt::Arg &getInputArg() const {
     assert(isInputArg() && "Invalid accessor.");
     return *Data.InputArg;
   }
 
   /// getAsString - Return a string name for this input, for
   /// debugging.
   std::string getAsString() const {
     if (isFilename())
       return std::string("\"") + getFilename() + '"';
     else if (isInputArg())
       return "(input arg)";
     else
       return "(nothing)";
   }
+
+  /// getOriginalAction - Return the action that produced this input info
+  const Action *getOriginalAction() const { return OrigAction; }
+
+  /// hasOffloadingTargetSuffixApended - Return true if the target suffix is
+  /// appended to
+  /// to this input filename
+  bool hasOffloadingTargetSuffixApended() {
+    return HasOffloadingTargetSuffixApended;
+  }
 };
 
 } // end namespace driver
 } // end namespace clang
 
 #endif

lib/Driver/ToolChain.cpp

if (Exceptions &&
       (Exceptions->getOption().matches(options::OPT_fexceptions) ||
        Exceptions->getOption().matches(options::OPT_fcxx_exceptions)))
     return ToolChain::RM_EnabledImplicitly;
 
   return ToolChain::RM_DisabledImplicitly;
 }
 
 ToolChain::ToolChain(const Driver &D, const llvm::Triple &T,
-                     const ArgList &Args)
+                     const ArgList &Args, bool IsOffloadingTargetToolchain)
     : D(D), Triple(T), Args(Args), CachedRTTIArg(GetRTTIArgument(Args)),
-      CachedRTTIMode(CalculateRTTIMode(Args, Triple, CachedRTTIArg)) {
+      CachedRTTIMode(CalculateRTTIMode(Args, Triple, CachedRTTIArg)),
+      IsOffloadingTargetToolchain(IsOffloadingTargetToolchain) {
   if (Arg *A = Args.getLastArg(options::OPT_mthread_model))
     if (!isThreadModelSupported(A->getValue()))
       D.Diag(diag::err_drv_invalid_thread_model_for_target)
           << A->getValue() << A->getAsString(Args);
 }
 
 ToolChain::~ToolChain() {
 }
Tool *ToolChain::getTool(Action::ActionClass AC) const {
   case Action::AssembleJobClass:
     return getAssemble();
 
   case Action::LinkJobClass:
     return getLink();
 
   case Action::InputClass:
   case Action::BindArchClass:
+  case Action::BindTargetClass:
   case Action::LipoJobClass:
   case Action::DsymutilJobClass:
   case Action::VerifyDebugInfoJobClass:
     llvm_unreachable("Invalid tool kind.");
 
   case Action::CompileJobClass:
   case Action::PrecompileJobClass:
   case Action::PreprocessJobClass:
   case Action::AnalyzeJobClass:
   case Action::MigrateJobClass:
   case Action::VerifyPCHJobClass:
   case Action::BackendJobClass:
     return getClang();
   }
 
   llvm_unreachable("Invalid tool kind.");
 }
 
+llvm::opt::DerivedArgList *
+ToolChain::TranslateArgs(const llvm::opt::DerivedArgList &Args,
+                         const char *BoundArch, bool isOffloadingTarget,
+                         bool &isSuccess) const {
+  // In the event this toolchain refers to an offloading target, a translation
+  // from host to target arguments should be implemented. This can be used.
+  // among other things, to force the resulting image to be a shared library
+  // that can be loaded by the offloading runtime libraries.
+  if (isOffloadingTarget) {
+    isSuccess = false;
+    llvm_unreachable(
+        "Toolchain not prepared to translate offloading arguments.");
+    return nullptr;
+  }
+
+  isSuccess = true;
+  return nullptr;
+}
+
 Tool *ToolChain::SelectTool(const JobAction &JA) const {
   if (getDriver().ShouldUseClangCompiler(JA))
     return getClang();
   Action::ActionClass AC = JA.getKind();
   if (AC == Action::AssembleJobClass && useIntegratedAs())
     return getClangAs();
   return getTool(AC);
 }

lib/Driver/ToolChains.h

class GCCInstallationDetector {
     std::set<std::string> CandidateGCCInstallPaths;
 
     /// The set of multilibs that the detected installation supports.
     MultilibSet Multilibs;
 
   public:
     GCCInstallationDetector() : IsValid(false) {}
     void init(const Driver &D, const llvm::Triple &TargetTriple,
-                            const llvm::opt::ArgList &Args);
+              const llvm::opt::ArgList &Args, bool isOffloadingTarget);
 
     /// \brief Check whether we detected a valid GCC install.
     bool isValid() const { return IsValid; }
 
     /// \brief Get the GCC triple for the detected install.
     const llvm::Triple &getTriple() const { return GCCTriple; }
 
     /// \brief Get the detected GCC installation path.
CollectLibDirsAndTriples(const llvm::Triple &TargetTriple,
                              SmallVectorImpl<StringRef> &TripleAliases,
                              SmallVectorImpl<StringRef> &BiarchLibDirs,
                              SmallVectorImpl<StringRef> &BiarchTripleAliases);
 
     void ScanLibDirForGCCTriple(const llvm::Triple &TargetArch,
                                 const llvm::opt::ArgList &Args,
                                 const std::string &LibDir,
                                 StringRef CandidateTriple,
+                                bool isOffloadingTarget,
                                 bool NeedsBiarchSuffix = false);
   };
 
   GCCInstallationDetector GCCInstallation;
 
 public:
   Generic_GCC(const Driver &D, const llvm::Triple &Triple,
-              const llvm::opt::ArgList &Args);
+              const llvm::opt::ArgList &Args,
+              bool IsOffloadingTargetToolchain = false);
   ~Generic_GCC() override;
 
   void printVerboseInfo(raw_ostream &OS) const override;
 
   bool IsUnwindTablesDefault() const override;
   bool isPICDefault() const override;
   bool isPIEDefault() const override;
   bool isPICDefaultForced() const override;
   bool IsIntegratedAssemblerDefault() const override;
 
+  virtual llvm::opt::DerivedArgList *
+  TranslateArgs(const llvm::opt::DerivedArgList &Args, const char *BoundArch,
+                bool isOffloadingTarget, bool &isSuccess) const override;
+
 protected:
   Tool *getTool(Action::ActionClass AC) const override;
   Tool *buildAssembler() const override;
   Tool *buildLinker() const override;
 
   /// \name ToolChain Implementation Helper Functions
   /// @{
 
public:
   std::string ComputeEffectiveClangTriple(const llvm::opt::ArgList &Args,
                                           types::ID InputType) const override;
 
   types::ID LookupTypeForExtension(const char *Ext) const override;
 
   bool HasNativeLLVMSupport() const override;
 
   llvm::opt::DerivedArgList *
-  TranslateArgs(const llvm::opt::DerivedArgList &Args,
-                const char *BoundArch) const override;
+  TranslateArgs(const llvm::opt::DerivedArgList &Args, const char *BoundArch,
+                bool isOffloadingTarget, bool &isSuccess) const override;
 
   bool IsBlocksDefault() const override {
     // Always allow blocks on Apple; users interested in versioning are
     // expected to use /usr/include/Block.h.
     return true;
   }
   bool IsIntegratedAssemblerDefault() const override {
     // Default integrated assembler to on for Apple's MachO targets.
public:
   /// {
 
   // Darwin tools support multiple architecture (e.g., i386 and x86_64) and
   // most development is done against SDKs, so compiling for a different
   // architecture should not get any special treatment.
   bool isCrossCompiling() const override { return false; }
 
   llvm::opt::DerivedArgList *
-  TranslateArgs(const llvm::opt::DerivedArgList &Args,
-                const char *BoundArch) const override;
+  TranslateArgs(const llvm::opt::DerivedArgList &Args, const char *BoundArch,
+                bool isOffloadingTarget, bool &isSuccess) const override;
 
   ObjCRuntime getDefaultObjCRuntime(bool isNonFragile) const override;
   bool hasBlocksRuntime() const override;
 
   bool UseObjCMixedDispatch() const override {
     // This is only used with the non-fragile ABI and non-legacy dispatch.
 
     // Mixed dispatch is used everywhere except OS X before 10.6.
void AddLinkSanitizerLibArgs(const llvm::opt::ArgList &Args,
                                llvm::opt::ArgStringList &CmdArgs,
                                StringRef Sanitizer) const;
 };
 
 class LLVM_LIBRARY_VISIBILITY Generic_ELF : public Generic_GCC {
   virtual void anchor();
 public:
   Generic_ELF(const Driver &D, const llvm::Triple &Triple,
-              const llvm::opt::ArgList &Args)
-      : Generic_GCC(D, Triple, Args) {}
+              const llvm::opt::ArgList &Args,
+              bool IsOffloadingTargetToolchain = false)
+      : Generic_GCC(D, Triple, Args, IsOffloadingTargetToolchain) {}
 
   void addClangTargetOptions(const llvm::opt::ArgList &DriverArgs,
                              llvm::opt::ArgStringList &CC1Args) const override;
 };
 
 class LLVM_LIBRARY_VISIBILITY CloudABI : public Generic_ELF {
 public:
   CloudABI(const Driver &D, const llvm::Triple &Triple,

 protected:
   Tool *buildAssembler() const override;
   Tool *buildLinker() const override;
 };
 
 class LLVM_LIBRARY_VISIBILITY Linux : public Generic_ELF {
 public:
   Linux(const Driver &D, const llvm::Triple &Triple,
-        const llvm::opt::ArgList &Args);
+        const llvm::opt::ArgList &Args,
+        bool IsOffloadingTargetToolchain = false);
 
   bool HasNativeLLVMSupport() const override;
 
   void
   AddClangSystemIncludeArgs(const llvm::opt::ArgList &DriverArgs,
                             llvm::opt::ArgStringList &CC1Args) const override;
   void
   AddClangCXXStdlibIncludeArgs(const llvm::opt::ArgList &DriverArgs,

lib/Driver/ToolChains.cpp

void DarwinClang::AddCCKextLibArgs(const ArgList &Args,
 
   // For now, allow missing resource libraries to support developers who may
   // not have compiler-rt checked out or integrated into their build.
   if (llvm::sys::fs::exists(P))
     CmdArgs.push_back(Args.MakeArgString(P));
 }
 
 DerivedArgList *MachO::TranslateArgs(const DerivedArgList &Args,
-                                     const char *BoundArch) const {
+                                     const char *BoundArch,
+                                     bool isOffloadingTarget,
+                                     bool &isSuccess) const {
+
+  if (isOffloadingTarget) {
+    // This translation is not dealing with offloading targets yet.
+    isSuccess = false;
+    llvm_unreachable(
+        "Toolchain not prepared to translate offloading arguments.");
+    return nullptr;
+  }
+
+  isSuccess = true;
+
   DerivedArgList *DAL = new DerivedArgList(Args.getBaseArgs());
   const OptTable &Opts = getDriver().getOpts();
 
   // FIXME: We really want to get out of the tool chain level argument
   // translation business, as it makes the driver functionality much
   // more opaque. For now, we follow gcc closely solely for the
   // purpose of easily achieving feature parity & testability. Once we
   // have something that works, we should reevaluate each translation
CompilerRT +=
       tools::arm::getARMFloatABI(getDriver(), Args, getTriple()) == "hard"
           ? "hard"
           : "soft";
   CompilerRT += Args.hasArg(options::OPT_fPIC) ? "_pic.a" : "_static.a";
 
   AddLinkRuntimeLib(Args, CmdArgs, CompilerRT, false, true);
 }
 
-
 DerivedArgList *Darwin::TranslateArgs(const DerivedArgList &Args,
-                                      const char *BoundArch) const {
+                                      const char *BoundArch,
+                                      bool isOffloadingTarget,
+                                      bool &isSuccess) const {
+
+  if (isOffloadingTarget) {
+    // This translation is not dealing with offloading targets yet.
+    isSuccess = false;
+    llvm_unreachable(
+        "Toolchain not prepared to translate offloading arguments.");
+    return nullptr;
+  }
+
   // First get the generic Apple args, before moving onto Darwin-specific ones.
-  DerivedArgList *DAL = MachO::TranslateArgs(Args, BoundArch);
+  DerivedArgList *DAL =
+      MachO::TranslateArgs(Args, BoundArch, isOffloadingTarget, isSuccess);
   const OptTable &Opts = getDriver().getOpts();
 
   // If no architecture is bound, none of the translations here are relevant.
   if (!BoundArch)
     return DAL;
 
   // Add an explicit version min argument for the deployment target. We do this
   // after argument translation because -Xarch_ arguments may add a version min

 ///
 /// This performs all of the autodetection and sets up the various paths.
 /// Once constructed, a GCCInstallationDetector is essentially immutable.
 ///
 /// FIXME: We shouldn't need an explicit TargetTriple parameter here, and
 /// should instead pull the target out of the driver. This is currently
 /// necessary because the driver doesn't store the final version of the target
 /// triple.
-void
-Generic_GCC::GCCInstallationDetector::init(
-    const Driver &D, const llvm::Triple &TargetTriple, const ArgList &Args) {
+void Generic_GCC::GCCInstallationDetector::init(
+    const Driver &D, const llvm::Triple &TargetTriple, const ArgList &Args,
+    bool isOffloadingTarget) {
   llvm::Triple BiarchVariantTriple =
       TargetTriple.isArch32Bit() ? TargetTriple.get64BitArchVariant()
                                  : TargetTriple.get32BitArchVariant();
   // The library directories which may contain GCC installations.
   SmallVector<StringRef, 4> CandidateLibDirs, CandidateBiarchLibDirs;
   // The compatible GCC triples for this particular architecture.
   SmallVector<StringRef, 16> CandidateTripleAliases;
   SmallVector<StringRef, 16> CandidateBiarchTripleAliases;
for (unsigned i = 0, ie = Prefixes.size(); i < ie; ++i) {
     if (!llvm::sys::fs::exists(Prefixes[i]))
       continue;
     for (unsigned j = 0, je = CandidateLibDirs.size(); j < je; ++j) {
       const std::string LibDir = Prefixes[i] + CandidateLibDirs[j].str();
       if (!llvm::sys::fs::exists(LibDir))
         continue;
       for (unsigned k = 0, ke = CandidateTripleAliases.size(); k < ke; ++k)
         ScanLibDirForGCCTriple(TargetTriple, Args, LibDir,
-                               CandidateTripleAliases[k]);
+                               CandidateTripleAliases[k], isOffloadingTarget);
     }
     for (unsigned j = 0, je = CandidateBiarchLibDirs.size(); j < je; ++j) {
       const std::string LibDir = Prefixes[i] + CandidateBiarchLibDirs[j].str();
       if (!llvm::sys::fs::exists(LibDir))
         continue;
       for (unsigned k = 0, ke = CandidateBiarchTripleAliases.size(); k < ke;
            ++k)
         ScanLibDirForGCCTriple(TargetTriple, Args, LibDir,
                                CandidateBiarchTripleAliases[k],
-                               /*NeedsBiarchSuffix=*/ true);
+                               isOffloadingTarget,
+                               /*NeedsBiarchSuffix=*/true);
     }
   }
 }
 
 void Generic_GCC::GCCInstallationDetector::print(raw_ostream &OS) const {
   for (const auto &InstallPath : CandidateGCCInstallPaths)
     OS << "Found candidate GCC installation: " << InstallPath << "\n";
 
if (Result.SelectedMultilib == Alt64 ||
     Result.BiarchSibling = Default;
 
   return true;
 }
 
 void Generic_GCC::GCCInstallationDetector::ScanLibDirForGCCTriple(
     const llvm::Triple &TargetTriple, const ArgList &Args,
     const std::string &LibDir, StringRef CandidateTriple,
-    bool NeedsBiarchSuffix) {
+    bool isOffloadingTarget, bool NeedsBiarchSuffix) {
   llvm::Triple::ArchType TargetArch = TargetTriple.getArch();
   // There are various different suffixes involving the triple we
   // check for. We also record what is necessary to walk from each back
   // up to the lib directory.
   const std::string LibSuffixes[] = {
     "/gcc/" + CandidateTriple.str(),
     // Debian puts cross-compilers in gcc-cross
     "/gcc-cross/" + CandidateTriple.str(),
for (llvm::sys::fs::directory_iterator LI(LibDir + LibSuffix, EC), LE;
       // Linux.
       GCCInstallPath = LibDir + LibSuffixes[i] + "/" + VersionText.str();
       GCCParentLibPath = GCCInstallPath + InstallSuffixes[i];
       IsValid = true;
     }
   }
 }
 
-Generic_GCC::Generic_GCC(const Driver &D, const llvm::Triple& Triple,
-                         const ArgList &Args)
-  : ToolChain(D, Triple, Args), GCCInstallation() {
+Generic_GCC::Generic_GCC(const Driver &D, const llvm::Triple &Triple,
+                         const ArgList &Args, bool IsOffloadingTargetToolchain)
+    : ToolChain(D, Triple, Args, IsOffloadingTargetToolchain),
+      GCCInstallation() {
   getProgramPaths().push_back(getDriver().getInstalledDir());
   if (getDriver().getInstalledDir() != getDriver().Dir)
     getProgramPaths().push_back(getDriver().Dir);
 }
 
 Generic_GCC::~Generic_GCC() {
 }
 

 bool Generic_GCC::isPIEDefault() const {
   return false;
 }
 
 bool Generic_GCC::isPICDefaultForced() const {
   return false;
 }
 
+llvm::opt::DerivedArgList *
+Generic_GCC::TranslateArgs(const llvm::opt::DerivedArgList &Args,
+                           const char *BoundArch, bool isOffloadingTarget,
+                           bool &isSuccess) const {
+
+  isSuccess = true;
+
+  // If not a target tool chain we can use the arguments directly without
+  // translation
+  if (!isOffloadingTarget)
+    return 0;
+
+  DerivedArgList *DAL = new DerivedArgList(Args.getBaseArgs());
+  const OptTable &Opts = getDriver().getOpts();
+
+  // Make sure we always generate a shared library for an offloading target
+  // regardless the commands the user passed to the host.
+  DAL->AddFlagArg(0, Opts.getOption(options::OPT_shared));
+  DAL->AddFlagArg(0, Opts.getOption(options::OPT_fPIC));
+
+  // Filter all the arguments we don't care passing to the offloading toolchain
+  // as they can mess up with the creation of a shared library.
+  for (ArgList::const_iterator it = Args.begin(), ie = Args.end(); it != ie;
+       ++it) {
+    Arg *A = *it;
+
+    switch ((options::ID)A->getOption().getID()) {
+    default:
+      DAL->append(A);
+      break;
+    case options::OPT_shared:
+    case options::OPT_static:
+    case options::OPT_fPIC:
+    case options::OPT_fno_PIC:
+    case options::OPT_fpic:
+    case options::OPT_fno_pic:
+    case options::OPT_fPIE:
+    case options::OPT_fno_PIE:
+    case options::OPT_fpie:
+    case options::OPT_fno_pie:
+      break;
+    }
+  }
+
+  return DAL;
+}
+
 bool Generic_GCC::IsIntegratedAssemblerDefault() const {
   return getTriple().getArch() == llvm::Triple::x86 ||
          getTriple().getArch() == llvm::Triple::x86_64 ||
          getTriple().getArch() == llvm::Triple::aarch64 ||
          getTriple().getArch() == llvm::Triple::aarch64_be ||
          getTriple().getArch() == llvm::Triple::arm ||
          getTriple().getArch() == llvm::Triple::armeb ||
          getTriple().getArch() == llvm::Triple::thumb ||
static StringRef getOSLibDir(const llvm::Triple &Triple, const ArgList &Args) {
 
   if (Triple.getArch() == llvm::Triple::x86_64 &&
       Triple.getEnvironment() == llvm::Triple::GNUX32)
     return "libx32";
 
   return Triple.isArch32Bit() ? "lib" : "lib64";
 }
 
-Linux::Linux(const Driver &D, const llvm::Triple &Triple, const ArgList &Args)
-  : Generic_ELF(D, Triple, Args) {
-  GCCInstallation.init(D, Triple, Args);
+Linux::Linux(const Driver &D, const llvm::Triple &Triple, const ArgList &Args,
+             bool IsOffloadingTargetToolchain)
+    : Generic_ELF(D, Triple, Args, IsOffloadingTargetToolchain) {
+  GCCInstallation.init(D, Triple, Args, IsOffloadingTargetToolchain);
   Multilibs = GCCInstallation.getMultilibs();
   llvm::Triple::ArchType Arch = Triple.getArch();
   std::string SysRoot = computeSysRoot();
 
   // Cross-compiling binutils and GCC installations (vanilla and openSUSE at
   // least) put various tools in a triple-prefixed directory off of the parent
   // of the GCC installation. We use the GCC triple here to ensure that we end
   // up with tools that support the same amount of cross compiling as the

lib/Driver/Tools.cpp

if (CombinedArg) {
       CmdArgs.push_back(Args.MakeArgString(std::string(ArgName) + Dirs));
     } else {
       CmdArgs.push_back(ArgName);
       CmdArgs.push_back(Args.MakeArgString(Dirs));
     }
   }
 }
 
-static void AddLinkerInputs(const ToolChain &TC,
-                            const InputInfoList &Inputs, const ArgList &Args,
-                            ArgStringList &CmdArgs) {
+static void AddLinkerInputs(const ToolChain &TC, const InputInfoList &Inputs,
+                            const ArgList &Args, ArgStringList &CmdArgs,
+                            bool isOffloadingTargetLinkage = false) {
   const Driver &D = TC.getDriver();
 
   // Add extra linker input arguments which are not treated as inputs
   // (constructed via -Xarch_).
   Args.AddAllArgValues(CmdArgs, options::OPT_Zlinker_input);
 
   for (const auto &II : Inputs) {
     if (!TC.HasNativeLLVMSupport()) {
       // Don't try to pass LLVM inputs unless we have native support.
       if (II.getType() == types::TY_LLVM_IR ||
           II.getType() == types::TY_LTO_IR ||
           II.getType() == types::TY_LLVM_BC ||
           II.getType() == types::TY_LTO_BC)
         D.Diag(diag::err_drv_no_linker_llvm_support)
           << TC.getTripleString();
     }
 
-    // Add filenames immediately.
+    // Add filenames immediately except if this is a host linker phase and the
+    // the object files (inputs) are produced by a target assembler. Those are
+    // handled by a linker script.
     if (II.isFilename()) {
+
+      if (!isOffloadingTargetLinkage &&
+          II.getOriginalAction()->getOffloadingDevice())
+        continue;
+
       CmdArgs.push_back(II.getFilename());
       continue;
     }
 
     // Otherwise, this is a linker input argument.
     const Arg &A = II.getInputArg();
 
     // Handle reserved library options.
static void AddLinkerInputs(const ToolChain &TC, const InputInfoList &Inputs,
   }
 
   // LIBRARY_PATH - included following the user specified library paths.
   //                and only supported on native toolchains.
   if (!TC.isCrossCompiling())
     addDirectoryList(Args, CmdArgs, "-L", "LIBRARY_PATH");
 }
 
+static void AddOpenMPLinkerScript(const ToolChain &TC, Compilation &C,
+                                  const JobAction &JA, const InputInfo &Output,
+                                  const InputInfoList &Inputs,
+                                  const ArgList &Args, ArgStringList &CmdArgs) {
+
+  const Arg *OpenMPArg = Args.getLastArg(options::OPT_fopenmp_EQ);
+
+  if (!OpenMPArg)
+    return;
+  if (!(StringRef(OpenMPArg->getValue()) == "libiomp5"))
+    return;
+
+  // This is the linkage for the target
+  if (JA.getOffloadingDevice())
+    return;
+
+  // Add OpenMP target arguments by employing
+
+  // FIXME: check if the toolchain supports a linker script
+
+  // Gather the pairs (target triple)-(file name)
+
+  std::vector<std::pair<const char *, const char *>> Targets;
+
+  for (InputInfoList::const_iterator it = Inputs.begin(), ie = Inputs.end();
+       it != ie; ++it) {
+    const InputInfo &II = *it;
+
+    if (const char *tname = II.getOriginalAction()->getOffloadingDevice()) {
+      Targets.push_back(
+          std::pair<const char *, const char *>(tname, II.getFilename()));
+    }
+  }
+
+  if (Targets.empty())
+    return;
+
+  // Create temporary linker script
+  StringRef Name = llvm::sys::path::filename(Output.getFilename());
+  std::pair<StringRef, StringRef> Split = Name.rsplit('.');
+  std::string TmpName = C.getDriver().GetTemporaryPath(Split.first, "lk");
+  const char *LKS = C.addTempFile(C.getArgs().MakeArgString(TmpName.c_str()));
+
+  // Open script file in order to write contents
+  std::error_code EC;
+  llvm::raw_fd_ostream lksf(LKS, EC, llvm::sys::fs::F_None);
+
+  if (EC) {
+    // FIXME: maybe use a file open failure message here
+    C.getDriver().Diag(clang::diag::err_unable_to_make_temp) << EC.message();
+    return;
+  }
+
+  // Add commands to embed target binaries. We ensure that each section and
+  // image s 16-byte aligned. This is not mandatory, but increases the
+  // likelihood of data to be aligned with a cache block in several main host
+  // machines.
+  lksf << "TARGET(binary)\n";
+  for (unsigned i = 0; i < Targets.size(); ++i)
+    lksf << "INPUT(" << Targets[i].second << ")\n";
+
+  lksf << "SECTIONS\n";
+  lksf << "{\n";
+  lksf << "  .openmptgt :\n";
+  lksf << "  ALIGN(0x10)\n";
+  lksf << "  {\n";
+
+  for (unsigned i = 0; i < Targets.size(); ++i) {
+    std::string tgt_name(Targets[i].first);
+    std::replace(tgt_name.begin(), tgt_name.end(), '-', '_');
+    lksf << "    . = ALIGN(0x10);\n";
+    lksf << "    PROVIDE_HIDDEN(__omptgt__img_start_" << tgt_name << " = .);\n";
+    lksf << "    " << Targets[i].second << "\n";
+    lksf << "    PROVIDE_HIDDEN(__omptgt__img_end_" << tgt_name << " = .);\n";
+  }
+
+  lksf << "  }\n";
+  // Add commands to define host entries begin and end
+  lksf << "  .openmptgt_host_entries :\n";
+  lksf << "  ALIGN(0x10)\n";
+  lksf << "  SUBALIGN(0x01)\n";
+  lksf << "  {\n";
+  lksf << "    PROVIDE_HIDDEN(__omptgt__host_entries_begin = .);\n";
+  lksf << "    *(.openmptgt_host_entries)\n";
+  lksf << "    PROVIDE_HIDDEN(__omptgt__host_entries_end = .);\n";
+  lksf << "  }\n";
+  lksf << "}\n";
+  lksf << "INSERT BEFORE .data\n";
+
+  lksf.close();
+
+  CmdArgs.push_back("-T");
+  CmdArgs.push_back(LKS);
+}
+
 /// \brief Determine whether Objective-C automated reference counting is
 /// enabled.
 static bool isObjCAutoRefCount(const ArgList &Args) {
   return Args.hasFlag(options::OPT_fobjc_arc, options::OPT_fno_objc_arc, false);
 }
 
 /// \brief Determine whether we are linking the ObjC runtime.
 static bool isObjCRuntimeLinked(const ArgList &Args) {
void Clang::ConstructJob(Compilation &C, const JobAction &JA,
   const Driver &D = getToolChain().getDriver();
   ArgStringList CmdArgs;
 
   bool IsWindowsGNU = getToolChain().getTriple().isWindowsGNUEnvironment();
   bool IsWindowsCygnus =
       getToolChain().getTriple().isWindowsCygwinEnvironment();
   bool IsWindowsMSVC = getToolChain().getTriple().isWindowsMSVCEnvironment();
 
-  assert(Inputs.size() == 1 && "Unable to handle multiple inputs.");
+  assert(((Inputs.size() == 1) ||
+          (Inputs.size() == 2 && JA.getOffloadingDevice() &&
+           isa<CompileJobAction>(JA))) &&
+         "Expecting 1 or 2 inputs (offloading host file)");
   const InputInfo &Input = Inputs[0];
 
   // Invoke ourselves in -cc1 mode.
   //
   // FIXME: Implement custom jobs for internal actions.
   CmdArgs.push_back("-cc1");
 
+  // Add the OpenMP arguments, including the ones related to offloading.
+  if (const Arg *A = Args.getLastArg(options::OPT_fopenmp_EQ)) {
+    if (StringRef(A->getValue()) == "libiomp5") {
+      CmdArgs.push_back("-fopenmp=libiomp5");
+
+      // pass the targets we are generating code to
+      if (Arg *Tgts = Args.getLastArg(options::OPT_omptargets_EQ)) {
+
+        ArrayRef<const char *> Vals = Tgts->getValues();
+
+        if (!Vals.empty()) {
+          std::string S("-omptargets=");
+          S += Vals[0];
+          for (unsigned i = 1; i < Vals.size(); ++i) {
+            S += ',';
+            S += Vals[i];
+          }
+          CmdArgs.push_back(Args.MakeArgString(S));
+        }
+      }
+
+      // Inform the frontend we are generating code for a target.
+      if (JA.getOffloadingDevice()) {
+        CmdArgs.push_back("-omp-target-mode");
+        // For compile jobs, the host resulting IR is also passed, so that the
+        // required elements that require to be emitted for the target are read
+        // from metadata presented in there.
+        if (isa<CompileJobAction>(JA)) {
+          CmdArgs.push_back("-omp-host-output-file-path");
+          CmdArgs.push_back(Args.MakeArgString(Inputs[1].getFilename()));
+        }
+      }
+
+      // The frontend components needs to know the path of the original source
+      // file given that the target functions use that to generate a unique
+      // names.
+      CmdArgs.push_back("-omp-main-file-path");
+      CmdArgs.push_back(Args.MakeArgString(Input.getBaseInput()));
+    }
+  }
+
   // Add the "effective" target triple.
   CmdArgs.push_back("-triple");
   std::string TripleStr = getToolChain().ComputeEffectiveClangTriple(Args);
   CmdArgs.push_back(Args.MakeArgString(TripleStr));
 
   const llvm::Triple TT(TripleStr);
   if (TT.isOSWindows() && (TT.getArch() == llvm::Triple::arm ||
                            TT.getArch() == llvm::Triple::thumb)) {
if (Output.getType() == types::TY_Dependencies) {
     // Handled with other dependency code.
   } else if (Output.isFilename()) {
     CmdArgs.push_back("-o");
     CmdArgs.push_back(Output.getFilename());
   } else {
     assert(Output.isNothing() && "Invalid output.");
   }
 
-  for (const auto &II : Inputs) {
-    addDashXForInput(Args, II, CmdArgs);
+  addDashXForInput(Args, Input, CmdArgs);
 
-    if (II.isFilename())
-      CmdArgs.push_back(II.getFilename());
-    else
-      II.getInputArg().renderAsInput(Args, CmdArgs);
-  }
+  if (Input.isFilename())
+    CmdArgs.push_back(Input.getFilename());
+    Input.getInputArg().renderAsInput(Args, CmdArgs);
 
   Args.AddAllArgs(CmdArgs, options::OPT_undef);
 
   const char *Exec = getToolChain().getDriver().getClangProgramPath();
 
   // Optionally embed the -cc1 level arguments into the debug info, for build
   // analysis.
   if (getToolChain().UseDwarfDebugFlags()) {

 
 void GnuTool::anchor() {}
 
 void gcc::Common::ConstructJob(Compilation &C, const JobAction &JA,
                                const InputInfo &Output,
                                const InputInfoList &Inputs,
                                const ArgList &Args,
                                const char *LinkingOutput) const {
   const Driver &D = getToolChain().getDriver();
   ArgStringList CmdArgs;
 
+  bool isLinkJob = JA.getKind() == Action::LinkJobClass;
+
   for (const auto &A : Args) {
     if (forwardToGCC(A->getOption())) {
       // Don't forward any -g arguments to assembly steps.
       if (isa<AssembleJobAction>(JA) &&
           A->getOption().matches(options::OPT_g_Group))
         continue;
 
       // Don't forward any -W arguments to assembly and link steps.
else if (II.getType() == types::TY_ModuleFile)
       D.Diag(diag::err_drv_no_module_support)
         << getToolChain().getTripleString();
 
     if (types::canTypeBeUserSpecified(II.getType())) {
       CmdArgs.push_back("-x");
       CmdArgs.push_back(types::getTypeName(II.getType()));
     }
 
-    if (II.isFilename())
+    if (II.isFilename()) {
+      bool isTargetLinkage = isLinkJob && JA.getOffloadingDevice();
+      if (!isTargetLinkage && II.getOriginalAction()->getOffloadingDevice())
+        continue;
+
       CmdArgs.push_back(II.getFilename());
-    else {
+    } else {
       const Arg &A = II.getInputArg();
 
       // Reverse translate some rewritten options.
       if (A.getOption().matches(options::OPT_Z_reserved_lib_stdcxx)) {
         CmdArgs.push_back("-lstdc++");
         continue;
       }
 
       // Don't render as input, we need gcc to do the translations.
       A.render(Args, CmdArgs);
     }
   }
 
   const std::string customGCCName = D.getCCCGenericGCCName();
   const char *GCCName;
   if (!customGCCName.empty())
     GCCName = customGCCName.c_str();
   else if (D.CCCIsCXX()) {
     GCCName = "g++";
   } else
     GCCName = "gcc";
 
+  if (isLinkJob)
+    AddOpenMPLinkerScript(getToolChain(), C, JA, Output, Inputs, Args, CmdArgs);
+
   const char *Exec =
     Args.MakeArgString(getToolChain().GetProgramPath(GCCName));
   C.addCommand(llvm::make_unique<Command>(JA, *this, Exec, CmdArgs));
 }
 
 void gcc::Preprocess::RenderExtraToolArgs(const JobAction &JA,
                                           ArgStringList &CmdArgs) const {
   CmdArgs.push_back("-E");
case LibGOMP:
     break;
   case LibIOMP5:
     CmdArgs.push_back("-liomp5");
     break;
   case LibUnknown:
     break;
   }
 
-  AddLinkerInputs(getToolChain(), Inputs, Args, CmdArgs);
+  AddLinkerInputs(getToolChain(), Inputs, Args, CmdArgs,
+                  JA.getOffloadingDevice());
   // Build the input file for -filelist (list of linker input files) in case we
   // need it later
   for (const auto &II : Inputs) {
     if (!II.isFilename()) {
       // This is a linker input argument.
       // We cannot mix input arguments and file names in a -filelist input, thus
       // we prematurely stop our list (remaining files shall be passed as
       // arguments).
void darwin::Link::ConstructJob(Compilation &C, const JobAction &JA,
   if (!Args.hasArg(options::OPT_nostdlib) &&
       !Args.hasArg(options::OPT_nostartfiles)) {
     // endfile_spec is empty.
   }
 
   Args.AddAllArgs(CmdArgs, options::OPT_T_Group);
   Args.AddAllArgs(CmdArgs, options::OPT_F);
 
+  AddOpenMPLinkerScript(getToolChain(), C, JA, Output, Inputs, Args, CmdArgs);
+
   // -iframework should be forwarded as -F.
   for (auto it = Args.filtered_begin(options::OPT_iframework),
          ie = Args.filtered_end(); it != ie; ++it)
     CmdArgs.push_back(Args.MakeArgString(std::string("-F") +
                                          (*it)->getValue()));
 
   if (!Args.hasArg(options::OPT_nostdlib) &&
       !Args.hasArg(options::OPT_nodefaultlibs)) {
void gnutools::Link::ConstructJob(Compilation &C, const JobAction &JA,
 
   if (D.IsUsingLTO(getToolChain(), Args))
     AddGoldPlugin(ToolChain, Args, CmdArgs);
 
   if (Args.hasArg(options::OPT_Z_Xlinker__no_demangle))
     CmdArgs.push_back("--no-demangle");
 
   bool NeedsSanitizerDeps = addSanitizerRuntimes(ToolChain, Args, CmdArgs);
-  AddLinkerInputs(ToolChain, Inputs, Args, CmdArgs);
+  AddLinkerInputs(ToolChain, Inputs, Args, CmdArgs, JA.getOffloadingDevice());
   // The profile runtime also needs access to system libraries.
   addProfileRT(getToolChain(), Args, CmdArgs);
 
   if (D.CCCIsCXX() &&
       !Args.hasArg(options::OPT_nostdlib) &&
       !Args.hasArg(options::OPT_nodefaultlibs)) {
     bool OnlyLibstdcxxStatic = Args.hasArg(options::OPT_static_libstdcxx) &&
       !Args.hasArg(options::OPT_static);
if (!Args.hasArg(options::OPT_nodefaultlibs)) {
         CmdArgs.push_back("-lgomp");
 
         // FIXME: Exclude this for platforms with libgomp that don't require
         // librt. Most modern Linux platforms require it, but some may not.
         CmdArgs.push_back("-lrt");
         break;
       case LibIOMP5:
         CmdArgs.push_back("-liomp5");
+        if (Args.hasArg(options::OPT_omptargets_EQ) &&
+            !JA.getOffloadingDevice())
+          CmdArgs.push_back("-lomptarget");
         break;
       case LibUnknown:
         break;
       }
       AddRunTimeLibs(ToolChain, D, CmdArgs, Args);
 
       if ((Args.hasArg(options::OPT_pthread) ||
            Args.hasArg(options::OPT_pthreads) || UsedOpenMPLib != LibUnknown) &&
if (!Args.hasArg(options::OPT_nostartfiles)) {
         crtend = isAndroid ? "crtend_android.o" : "crtend.o";
 
       CmdArgs.push_back(Args.MakeArgString(ToolChain.GetFilePath(crtend)));
       if (!isAndroid)
         CmdArgs.push_back(Args.MakeArgString(ToolChain.GetFilePath("crtn.o")));
     }
   }
 
+  AddOpenMPLinkerScript(getToolChain(), C, JA, Output, Inputs, Args, CmdArgs);
+
   C.addCommand(
       llvm::make_unique<Command>(JA, *this, ToolChain.Linker.c_str(), CmdArgs));
 }
 
 
 // NaCl ARM assembly (inline or standalone) can be written with a set of macros
 // for the various SFI requirements like register masking. The assembly tool
 // inserts the file containing the macros as an input into all the assembly
 // jobs.
 void nacltools::AssembleARM::ConstructJob(Compilation &C, const JobAction &JA,
                                           const InputInfo &Output,
                                           const InputInfoList &Inputs,
                                           const ArgList &Args,
                                           const char *LinkingOutput) const {
   const toolchains::NaCl_TC& ToolChain =
     static_cast<const toolchains::NaCl_TC&>(getToolChain());
-  InputInfo NaClMacros(ToolChain.GetNaClArmMacrosPath(), types::TY_PP_Asm,
-                       "nacl-arm-macros.s");
+  InputInfo NaClMacros(ToolChain.GetNaClArmMacrosPath(), &JA,
+                       "nacl-arm-macros.s", false);
   InputInfoList NewInputs;
   NewInputs.push_back(NaClMacros);
   NewInputs.append(Inputs.begin(), Inputs.end());
   gnutools::Assemble::ConstructJob(C, JA, Output, NewInputs, Args,
                                    LinkingOutput);
 }
 
 

test/OpenMP/target_driver.c

+///
+/// Perform several driver tests for OpenMP offloading
+///
+
+/// ###########################################################################
+
+/// Check whether an invalid OpenMP target is specified:
+// RUN:   %clang -### -fopenmp=libiomp5 -omptargets=aaa-bbb-ccc-ddd %s 2>&1 \
+// RUN:   | FileCheck -check-prefix=CHK-INVALID-TARGET %s
+// CHK-INVALID-TARGET: error: OpenMP target is invalid: 'aaa-bbb-ccc-ddd'
+
+/// ###########################################################################
+
+/// Check error for empty -omptargets
+// RUN:   %clang -### -fopenmp=libiomp5 -omptargets=  %s 2>&1 \
+// RUN:   | FileCheck -check-prefix=CHK-EMPTY-OMPTARGETS %s
+// CHK-EMPTY-OMPTARGETS: warning: joined argument expects additional value: '-omptargets='
+
+/// ###########################################################################
+
+/// Check whether we are using a target whose toolchain was not prepared to
+/// to support offloading - e.g. x86_64-apple-darwin:
+// RUN:   %clang -### -fopenmp=libiomp5 -omptargets=x86_64-apple-darwin %s 2>&1 \
+// RUN:   | FileCheck -check-prefix=CHK-NO-SUPPORT %s
+// CHK-NO-SUPPORT: error: Toolchain for target 'x86_64-apple-darwin' is not supporting OpenMP offloading.
+
+/// ###########################################################################
+
+/// Check the phases graph when using a single target, different from the host.
+/// Each target phase must be binded to a target and linked into a shared
+/// library. The host compiler phase result is used in the compiler phase of the
+/// the target
+// RUN:   %clang -ccc-print-phases -fopenmp=libiomp5 -target powerpc64-ibm-linux-gnu -omptargets=x86_64-pc-linux-gnu %s 2>&1 \
+// RUN:   | FileCheck -check-prefix=CHK-PHASES %s
+
+// Host linking
+// CHK-PHASES-DAG: {{.*}}: linker, {[[A0:[0-9]+]], [[BL1:[0-9]+]]}, image
+
+// Target 1 library generation
+// CHK-PHASES-DAG: [[BL1]]: bind-target, {[[L1:[0-9]+]]}, shared-object
+// CHK-PHASES-DAG: [[L1]]: linker, {[[BA1:[0-9]+]]}, shared-object
+// CHK-PHASES-DAG: [[BA1]]: bind-target, {[[A1:[0-9]+]]}, object
+// CHK-PHASES-DAG: [[A1]]: assembler, {[[BB1:[0-9]+]]}, object
+// CHK-PHASES-DAG: [[BB1]]: bind-target, {[[B1:[0-9]+]]}, assembler
+// CHK-PHASES-DAG: [[B1]]: backend, {[[BC1:[0-9]+]]}, assembler
+// CHK-PHASES-DAG: [[BC1]]: bind-target, {[[C1:[0-9]+]]}, ir
+// CHK-PHASES-DAG: [[C1]]: compiler, {[[BP1:[0-9]+]], [[BC01:[0-9]+]]}, ir
+// CHK-PHASES-DAG: [[BC01]]: bind-target, {[[C0:[0-9]+]]}, ir
+// CHK-PHASES-DAG: [[BP1]]: bind-target, {[[P1:[0-9]+]]}, cpp-output
+// CHK-PHASES-DAG: [[P1]]: preprocessor, {[[I:[0-9]+]]}, cpp-output
+
+// Host objects generation:
+// CHK-PHASES-DAG: [[A0]]: assembler, {[[B0:[0-9]+]]}, object
+// CHK-PHASES-DAG: [[B0]]: backend, {[[C0]]}, assembler
+// CHK-PHASES-DAG: [[C0]]: compiler, {[[P0:[0-9]+]]}, ir
+// CHK-PHASES-DAG: [[P0]]: preprocessor, {[[I]]}, cpp-output
+
+// Single input file:
+// CHK-PHASES-DAG: [[I]]: input, {{.*}}, c
+
+/// ###########################################################################
+
+/// Check the phases graph when using two targets, and one of them is the same
+/// as the host. We also add a library to make sure they are not treated has
+/// inputs.
+/// Each target phase must be binded to a target and linked into a shared
+/// library. The host compiler phase result is used in the compiler phase of the
+/// the target
+// RUN:   %clang -ccc-print-phases -lm -fopenmp=libiomp5 -target powerpc64-ibm-linux-gnu -omptargets=x86_64-pc-linux-gnu,powerpc64-ibm-linux-gnu %s 2>&1 \
+// RUN:   | FileCheck -check-prefix=CHK-PHASES2 %s
+
+// Host linking
+// CHK-PHASES2-DAG: {{.*}}: linker, {[[L:[0-9]+]], [[A0:[0-9]+]], [[BL1:[0-9]+]], [[BL2:[0-9]+]]}, image
+
+// Target 2 library generation
+// CHK-PHASES2-DAG: [[BL2]]: bind-target, {[[L2:[0-9]+]]}, shared-object
+// CHK-PHASES2-DAG: [[L2]]: linker, {[[L]], [[BA2:[0-9]+]]}, shared-object
+// CHK-PHASES2-DAG: [[BA2]]: bind-target, {[[A2:[0-9]+]]}, object
+// CHK-PHASES2-DAG: [[A2]]: assembler, {[[BB2:[0-9]+]]}, object
+// CHK-PHASES2-DAG: [[BB2]]: bind-target, {[[B2:[0-9]+]]}, assembler
+// CHK-PHASES2-DAG: [[B2]]: backend, {[[BC2:[0-9]+]]}, assembler
+// CHK-PHASES2-DAG: [[BC2]]: bind-target, {[[C2:[0-9]+]]}, ir
+// CHK-PHASES2-DAG: [[C2]]: compiler, {[[BP2:[0-9]+]], [[BC02:[0-9]+]]}, ir
+// CHK-PHASES2-DAG: [[BC02]]: bind-target, {[[C0:[0-9]+]]}, ir
+// CHK-PHASES2-DAG: [[BP2]]: bind-target, {[[P2:[0-9]+]]}, cpp-output
+// CHK-PHASES2-DAG: [[P2]]: preprocessor, {[[I:[0-9]+]]}, cpp-output
+
+// Target 1 library generation
+// CHK-PHASES2-DAG: [[BL1]]: bind-target, {[[L1:[0-9]+]]}, shared-object
+// CHK-PHASES2-DAG: [[L1]]: linker, {[[L]], [[BA1:[0-9]+]]}, shared-object
+// CHK-PHASES2-DAG: [[BA1]]: bind-target, {[[A1:[0-9]+]]}, object
+// CHK-PHASES2-DAG: [[A1]]: assembler, {[[BB1:[0-9]+]]}, object
+// CHK-PHASES2-DAG: [[BB1]]: bind-target, {[[B1:[0-9]+]]}, assembler
+// CHK-PHASES2-DAG: [[B1]]: backend, {[[BC1:[0-9]+]]}, assembler
+// CHK-PHASES2-DAG: [[BC1]]: bind-target, {[[C1:[0-9]+]]}, ir
+// CHK-PHASES2-DAG: [[C1]]: compiler, {[[BP1:[0-9]+]], [[BC01:[0-9]+]]}, ir
+// CHK-PHASES2-DAG: [[BC01]]: bind-target, {[[C0]]}, ir
+// CHK-PHASES2-DAG: [[BP1]]: bind-target, {[[P1:[0-9]+]]}, cpp-output
+// CHK-PHASES2-DAG: [[P1]]: preprocessor, {[[I:[0-9]+]]}, cpp-output
+
+// Host objects generation:
+// CHK-PHASES2-DAG: [[A0]]: assembler, {[[B0:[0-9]+]]}, object
+// CHK-PHASES2-DAG: [[B0]]: backend, {[[C0]]}, assembler
+// CHK-PHASES2-DAG: [[C0]]: compiler, {[[P0:[0-9]+]]}, ir
+// CHK-PHASES2-DAG: [[P0]]: preprocessor, {[[I]]}, cpp-output
+
+// Single input file:
+// CHK-PHASES2-DAG: [[I]]: input, {{.*}}, c
+// CHK-PHASES2-DAG: [[L]]: input, "m", object
+
+/// ###########################################################################
+
+/// Check of the commands passed to each tool when using valid OpenMP targets.
+/// Here we also check that offloading does not break the use of integrated
+/// assembler. It does however preclude the use of integrated preprocessor as
+/// host IR is shared by all the compile phases. There several offloading
+/// specific commands:
+/// -omp-target-mode: will tell the frontend that it will generate code for a
+/// target.
+/// -omp-main-file-path: the original source file that relates with that
+/// frontend run, will be used to generate unique variable names (IDs) that are
+/// the same for all targets.
+/// -omp-host-output-file-path: specifies the host IR file that can be loaded by
+/// the target code generation to gather information about which declaration
+/// really need to be emitted.
+///
+// RUN:   %clang -### -fopenmp=libiomp5 -target powerpc64le-linux -omptargets=powerpc64le-ibm-linux-gnu,x86_64-pc-linux-gnu %s 2>&1 \
+// RUN:   | FileCheck -check-prefix=CHK-COMMANDS %s
+//
+
+// Final linking - host (ppc64le)
+// CHK-COMMANDS-DAG: ld" {{.*}}"-m" "elf64lppc" {{.*}}"-o" "a.out" {{.*}}"[[HSTOBJ:.+]].o" "-liomp5" "-lomptarget" {{.*}}"-T" "[[LKSCRIPT:.+]].lk"
+
+// Target 2 commands (x86_64)
+// CHK-COMMANDS-DAG: ld" {{.*}}"-m" "elf_x86_64" {{.*}}"-shared" {{.*}}"-o" "[[T2LIB:.+]].so" {{.*}}"[[T2OBJ:.+]].o" {{.*}}"-liomp5"
+// CHK-COMMANDS-DAG: clang{{.*}}" "-cc1" {{.*}}"-fopenmp=libiomp5" "-omptargets=powerpc64le-ibm-linux-gnu,x86_64-pc-linux-gnu" "-omp-target-mode" "-omp-main-file-path" "[[SRC:.+]].c" {{.*}}"-triple" "x86_64-pc-linux-gnu" {{.*}}"-emit-obj" {{.*}}"-o" "[[T2OBJ]].o" "-x" "ir" "[[T2BC:.+]].bc"
+// CHK-COMMANDS-DAG: clang{{.*}}" "-cc1" {{.*}}"-fopenmp=libiomp5" "-omptargets=powerpc64le-ibm-linux-gnu,x86_64-pc-linux-gnu" "-omp-target-mode" "-omp-host-output-file-path" "[[HSTBC:.+]].bc" "-omp-main-file-path" "[[SRC]].c" {{.*}}"-triple" "x86_64-pc-linux-gnu" {{.*}}"-emit-llvm-bc" {{.*}}"-o" "[[T2BC]].bc" "-x" "cpp-output" "[[T2PP:.+]].i"
+// CHK-COMMANDS-DAG: clang{{.*}}" "-cc1" {{.*}}"-fopenmp=libiomp5" "-omptargets=powerpc64le-ibm-linux-gnu,x86_64-pc-linux-gnu" "-omp-target-mode" "-omp-main-file-path" "[[SRC]].c" {{.*}}"-triple" "x86_64-pc-linux-gnu" {{.*}}"-E" {{.*}}"-o" "[[T2PP]].i" "-x" "c" "[[SRC]].c"
+
+// Target 1 commands (ppc64le)
+// CHK-COMMANDS-DAG: ld" {{.*}}"-m" "elf64lppc" {{.*}}"-shared" {{.*}}"-o" "[[T1LIB:.+]].so" {{.*}}"[[T1OBJ:.+]].o" {{.*}}"-liomp5"
+// CHK-COMMANDS-DAG: clang{{.*}}" "-cc1" {{.*}}"-fopenmp=libiomp5" "-omptargets=powerpc64le-ibm-linux-gnu,x86_64-pc-linux-gnu" "-omp-target-mode" "-omp-main-file-path" "[[SRC]].c" {{.*}}"-triple" "powerpc64le-ibm-linux-gnu" {{.*}}"-emit-obj" {{.*}}"-o" "[[T1OBJ]].o" "-x" "ir" "[[T1BC:.+]].bc"
+// CHK-COMMANDS-DAG: clang{{.*}}" "-cc1" {{.*}}"-fopenmp=libiomp5" "-omptargets=powerpc64le-ibm-linux-gnu,x86_64-pc-linux-gnu" "-omp-target-mode" "-omp-host-output-file-path" "[[HSTBC]].bc" "-omp-main-file-path" "[[SRC]].c" {{.*}}"-triple" "powerpc64le-ibm-linux-gnu" {{.*}}"-emit-llvm-bc" {{.*}}"-o" "[[T1BC]].bc" "-x" "cpp-output" "[[T1PP:.+]].i"
+// CHK-COMMANDS-DAG: clang{{.*}}" "-cc1" {{.*}}"-fopenmp=libiomp5" "-omptargets=powerpc64le-ibm-linux-gnu,x86_64-pc-linux-gnu" "-omp-target-mode" "-omp-main-file-path" "[[SRC]].c" {{.*}}"-triple" "powerpc64le-ibm-linux-gnu" {{.*}}"-E" {{.*}}"-o" "[[T1PP]].i" "-x" "c" "[[SRC]].c"
+
+// Host object generation
+// CHK-COMMANDS-DAG: clang{{.*}}" "-cc1" {{.*}}"-fopenmp=libiomp5" "-omptargets=powerpc64le-ibm-linux-gnu,x86_64-pc-linux-gnu" "-omp-main-file-path" "[[SRC]].c" {{.*}}"-triple" "powerpc64le--linux" {{.*}}"-emit-obj" {{.*}}"-o" "[[HSTOBJ]].o" "-x" "ir" "[[HSTBC]].bc"
+// CHK-COMMANDS-DAG: clang{{.*}}" "-cc1" {{.*}}"-fopenmp=libiomp5" "-omptargets=powerpc64le-ibm-linux-gnu,x86_64-pc-linux-gnu" "-omp-main-file-path" "[[SRC]].c" {{.*}}"-triple" "powerpc64le--linux" {{.*}}"-emit-llvm-bc" {{.*}}"-o" "[[HSTBC]].bc" "-x" "c" "[[SRC]].c"
+
+/// ###########################################################################
+
+/// Check the automatic detection of target files. The driver will automatically
+/// detect if a target file is in the same path as the host file and include
+/// that in the compilation. The user can choose to have the compiler generating
+/// a warning if such file is included.
+/// Create dummy host and target files.
+// RUN:   echo ' ' > %t.i
+// RUN:   echo ' ' > %t.i.tgt-x86_64-pc-linux-gnu
+// RUN:   echo ' ' > %t.bc
+// RUN:   echo ' ' > %t.bc.tgt-x86_64-pc-linux-gnu
+// RUN:   echo ' ' > %t.s
+// RUN:   echo ' ' > %t.s.tgt-x86_64-pc-linux-gnu
+// RUN:   echo ' ' > %t.o
+// RUN:   echo ' ' > %t.o.tgt-x86_64-pc-linux-gnu
+
+// RUN:   %clang -### -fopenmp=libiomp5 -target powerpc64-linux -omptargets=x86_64-pc-linux-gnu %t.i -c -emit-llvm -Womp-implicit-target-files 2>&1 \
+// RUN:   | FileCheck -check-prefix=CHK-TARGET-WARN-IMPLICIT %s
+// RUN:   %clang -### -fopenmp=libiomp5 -target powerpc64-linux -omptargets=x86_64-pc-linux-gnu %t.i -c -emit-llvm 2>&1 \
+// RUN:   | FileCheck -check-prefix=CHK-TARGET-NOTWARN-IMPLICIT %s
+// RUN:   %clang -### -fopenmp=libiomp5 -target powerpc64-linux -omptargets=x86_64-pc-linux-gnu %t.bc -S -Womp-implicit-target-files 2>&1 \
+// RUN:   | FileCheck -check-prefix=CHK-TARGET-WARN-IMPLICIT %s
+// RUN:   %clang -### -fopenmp=libiomp5 -target powerpc64-linux -omptargets=x86_64-pc-linux-gnu %t.bc -S 2>&1 \
+// RUN:   | FileCheck -check-prefix=CHK-TARGET-NOTWARN-IMPLICIT %s
+// RUN:   %clang -### -fopenmp=libiomp5 -target powerpc64-linux -omptargets=x86_64-pc-linux-gnu %t.s -c -Womp-implicit-target-files 2>&1 \
+// RUN:   | FileCheck -check-prefix=CHK-TARGET-WARN-IMPLICIT %s
+// RUN:   %clang -### -fopenmp=libiomp5 -target powerpc64-linux -omptargets=x86_64-pc-linux-gnu %t.s -c 2>&1 \
+// RUN:   | FileCheck -check-prefix=CHK-TARGET-NOTWARN-IMPLICIT %s
+// RUN:   %clang -### -fopenmp=libiomp5 -target powerpc64-linux -omptargets=x86_64-pc-linux-gnu %t.o -Womp-implicit-target-files 2>&1 \
+// RUN:   | FileCheck -check-prefix=CHK-TARGET-WARN-IMPLICIT %s
+// RUN:   %clang -### -fopenmp=libiomp5 -target powerpc64-linux -omptargets=x86_64-pc-linux-gnu %t.o 2>&1 \
+// RUN:   | FileCheck -check-prefix=CHK-TARGET-NOTWARN-IMPLICIT %s
+
+// CHK-TARGET-WARN-IMPLICIT: warning: OpenMP target file '{{.*}}.tgt-x86_64-pc-linux-gnu' is being implicitly used in the 'x86_64-pc-linux-gnu' toolchain.
+// CHK-TARGET-NOTWARN-IMPLICIT-NOT: warning: OpenMP target file '{{.*}}.tgt-x86_64-pc-linux-gnu' is being implicitly used in the 'x86_64-pc-linux-gnu' toolchain.
+
+/// ###########################################################################
+
+/// Check separate compilation feature - the ability of the driver to assign
+/// host and target files to different phases. Only the host files are passed
+/// to the driver. The driver will detect all the target files
+// RUN:   echo ' ' > %t.1.s
+// RUN:   echo ' ' > %t.1.s.tgt-x86_64-pc-linux-gnu
+// RUN:   echo ' ' > %t.2.o
+// RUN:   %clang -### -fopenmp=libiomp5 -target powerpc64-linux -omptargets=x86_64-pc-linux-gnu %t.1.s %t.2.o 2>&1 \
+// RUN:   | FileCheck -check-prefix=CHK-SEP-COMPILATION %s
+
+// Final linking
+// CHK-SEP-COMPILATION-DAG: ld" {{.*}}"-m" "elf64ppc" {{.*}}"[[HOSTOBJ:.+]].o" "[[HOSTOBJ2:.+]].o" "-liomp5" "-lomptarget" {{.*}}"-T" "[[LKS:.+]].lk"
+
+// Target image generation
+// CHK-SEP-COMPILATION-DAG: ld" {{.*}}"-m" "elf_x86_64" {{.*}}"-shared" {{.*}}"-o" "[[TGTSO:.+]].so" {{.*}}"[[TGTOBJ:.+]].o" {{.*}}"-liomp5"
+// CHK-SEP-COMPILATION-DAG: clang{{.*}}" "-cc1as" {{.*}}"-triple" "x86_64-pc-linux-gnu" {{.*}}"-o" "[[TGTOBJ]].o" "[[TGTASM:.+]].s.tgt-x86_64-pc-linux-gnu"
+
+// Host object generation
+// CHK-SEP-COMPILATION-DAG: clang{{.*}}" "-cc1as" {{.*}}"-triple" "powerpc64--linux" {{.*}}"-o" "[[HOSTOBJ]].o" "[[HOSTASM:.+]].s"