[OpenMP] Target directive host codegen

[sfantao]

Summary

This patch implements the outlining for offloading functions for code annotated with the OpenMP target directive. It uses a temporary naming of the outlined functions that will have to be updated later on once target side codegen and registration of offloading libraries is implemented - the naming needs to be made unique in the produced library.

Unlike other captured regions, target offloading cannot use directly the Capture declaration, as each captured field has to be passed explicitly to the runtime library and associated with potentially different mapping types (to/from/alloc...). Therefore, some tweaking in the function prologue codegen is required.

The current implementation still do not support capturing of global variables. That requires a modification in Sema that I will propose in a separate patch.

Thanks!
Samuel

[ABataev]

lib/CodeGen/CGOpenMPRuntime.h

188–212

Seems to me these enums are not used right now, must be removed

lib/CodeGen/CGStmt.cpp

2108–2109

Why do you need all this stuff here? Why you can't use variables captured in CapturedDecl?

[sfantao]

Remove unused enums.

[sfantao]

Alexey,

Thanks for the review! I am addressing your remarks as inline comments.

Thanks again,
Samuel

lib/CodeGen/CGOpenMPRuntime.h

188–212

Done!

lib/CodeGen/CGStmt.cpp

2108–2109

There is only one argument that comes with the CaptureDecl and that is the anonymous struct that captures all the references.

For the target outlined functions we have to pass each capture separately because not only are they arguments but also data mappings. Therefore we don't have to generate the anonymous struct at all and use the captures directly, signaling the proper map types. This follows what was discussed in the offloading infrastructure proposal - the compiler will generate arrays with each argument/mapping and forward that to the target specific plugin. The target function is therefore expected to have the same signature of the host function that is being generated by this patch, so that the order in the array is consistent with the signature.

The code in this patch that is generating the host version of the target region can therefore be completely reused to generate the device version.

[ABataev]

lib/CodeGen/CGOpenMPRuntime.h

188–203

Seems to me these enums are used in CGOpenMPRuntime.cpp, so we should not declare them in CGOpenMPRuntime interface.

lib/CodeGen/CGStmt.cpp

2108–2109

Could we add required functionality to CapturedStmt to make it pass mappings also, if required? Currently we have to expose some things that better to stay hidden like VLAMap, modify StartFunction(), when we could add required parameters in Sema analysis and get all required stuff for free, without any additional changes in codegen

[rjmccall]

lib/CodeGen/CGOpenMPRuntime.cpp

3070

You already know the element types for all of these. The code will be much more readable if you just use those types directly.

3100

Please move the creation of this global out of this initializer list.

lib/CodeGen/CGStmt.cpp

2108–2109

Okay, so from what I understand, __tgt_target is just a placeholder call that some later pass will rewrite to invoke a target-specific plugin that behaves like the host function. Fine, but:

• Doesn't it still need to get passed the host function?
• Why does it return an i32 instead of just an i1, if you're going to rewrite it anyway?
• It would be much clearer to name it something actually intrinsic-like, like "omp.target.call".

As for the signature: as we've discussed before, because the captured statement invocation function isn't actually exposed anywhere, its physical signature is really completely an implementation detail of the captured statement emitter. So I have no problem with this kind of change, but you should really design it to be flexible for different captured-statement emitters (by, say, taking a lambda or something similar) rather than hardcoding the specific needs of this one.

lib/CodeGen/CodeGenFunction.cpp

770

Please find some way of doing this that does not rely on changing common infrastructure like StartFunction.

1512

What do you need this for that you can't use getVLASize for?

[rjmccall]

lib/CodeGen/CGOpenMPRuntime.cpp

894–895

Spurious "to" at the start.

2954

getTypeStoreSize()

2963

You should ask the ASTContext to compute this size instead of making assumptions about the layout size of the IR type.

Also, what are the semantics supposed to be for mapping to and from? Do referents need to be trivially copyable? What if there are pointers or references? What happens to virtual bases?

2973

Same thing, please ask the ASTContext.

Also, you might need to care about variably-sized types here.

3042

This is creating a bunch of dynamic allocas instead of just temporary values. Please call CreateMemTemp with an array type instead.

3047

The sizes aren't constant if you've captured a VLA. But this comment is actually just wrong, because this isn't building something for the sizes at all; it's building something for the flags.

That said, I think you ought to be able to do the same thing with the sizes when you don't have a VLA.

3050

Please comment boolean arguments like this:

/*constant*/ true

And please mark this variable unnamed_addr.

3077

It makes more sense to emit the store next to the GEP. This will also let you more easily skip all the code associated with an array when you can emit it as a constant.

[ABataev]

lib/CodeGen/CGStmt.cpp

2108–2109

John, __tgt_target is not a placeholder, this is a real name of offloading runtime function. Runtime library for offloading is not a part of LLVM yet. You can find it here https://github.com/clang-omp/libomptarget. Runtime functions are declared here https://github.com/clang-omp/libomptarget/blob/master/src/omptarget.h

[rjmccall]

lib/CodeGen/CGStmt.cpp

2108–2109

Oh, I see. It's the host-pointer argument that's meant to uniquely determine the operation to perform on the target, and there's going to be some separate registration step that will associate that with a target-specific implementation?

Are you really adding this otherwise-unnecessary use of libffi just so that you can pass the captures as (indirect) arguments for no particular reason? You're free to make that decision, but I think it's a poor one. If you're assuming that all the migrating captures are POD — and I don't see how that's avoidable — then you'd be better off putting all the captured variables into a single allocation and presenting the target with the base address of that.

[sfantao]

This patch tries to avoid as much as possible changing the common infrastructure, by adapting the CapturedDecl creation in SEMA and by adding support to a second type of capture - ImplicitParamDecl (on top of the existent FieldDecl).

Also adds codegen for device clause as the Parsing and SEMA support was added in the meantime.

The regression test was not updated yet. I wanted to make sure the direction this is taking is approved before diving into that.

Thanks!
Samuel

[sfantao]

Alexey, John,

Thanks for the review! I've tried to address your concerns in the last diff. Please, check the inlined comments to find answers for the remarks of the previous diff.

Thanks again!
Samuel

lib/CodeGen/CGOpenMPRuntime.cpp

894–895

Fixed!

2954

Now using getTypeSizeInChars from the ASTcontext as suggested bellow.

2963

Using using getTypeSizeInChars from the ASTcontext.

This patch only deals with trivially copiable types. By default, a variable that is captured in the target region is mapped "by value" using a to-from policy. In order to do something different than the default, the user has to use a map clause (I'll send out a patch for it once we have the Parsing/SEMA in place). As per the current spec, the map clause allows a user to map the pointee of a pointer as well as only mapping a section of an array or pointee. In the next version of the OpenMP spec we will have the ability to map aggregate fields and more support for "deep" copy. We expect to be able to handle all the cases with the proper flags in OpenMPOffloadMappingFlags.

About virtual bases: OpenMP 4.0 forbids virtual members in mappable variable. Nevertheless, it is possible this constraint be lifted in future versions.

2973

Now using ASTContext.

I was planing to deal with the VLA types once I add support for the map clause, but I agree it makes more sense to do it now. Thanks for the suggestion.

3042

Done!

3047

Sorry for the error in the comment. It is fixed now.

Also, I added code to deal with constant sizes for when we don't have VLAs.

3050

Added comments for the boolean.

Setting unnamed_addr now, thanks!

3070

Ok, now using the types explicitly.

3077

Moved the stores next to the GEP. Also skipping the sizes for when no VLAs are used.

3100

Done!

lib/CodeGen/CGStmt.cpp

2108–2109

Correct, there is going to be a separate registration step where all the host pointers are going to be captured in a descriptor and passed to the OpenMP runtime library (__tgt_register_lib()). I'll contribute that in a separate patch.

libffi is not required unless one wants to do offloading to some host thread. The host version of the target region is not launched from the runtime library. If the runtime decides to execute the target region on the host it should return an error code in tgt_target and it will just work. What I think is key here is:

• To have the runtime be more efficient while scanning the mapped pointers/sizes/map types, and using arrays is a good solution for that.

• For offloading APIs I have experience with, e.g. CUDA, having an array of pointers is exactly what is expected. By passing the parameters separately will let the implementation decide what is best for the target, usually use some dedicated memory or registers for parameters as defined in the calling/launching convention it has in place (in the limit it could implement a single allocation and load the parameters from there as you are suggesting).

Let me know if you need me to clarify any detail about the interface.

lib/CodeGen/CodeGenFunction.cpp

770

In the new patch CaptureDecl was changed and CodeGenFunction::PrepareOMPTargetDirectiveBodyEmission was created so that this code is not required anymore.

1512

What I need is the size of each dimension and what getVLASize returns is the size of the type. I can probably find ways of doing this without this new member function, but that will be definitely more messy. Should I do that? Any suggestion?

[ABataev]

lib/CodeGen/CGExpr.cpp

1969

Samuel, why you don't want to capture all used variables in CapturedDecl instead of creating ImplicitParamDecl for each captured variable? Also, you will solve possible trouble with VLAs automatically using CapturedDecl.

lib/CodeGen/CGOpenMPRuntime.cpp

2953

Use CGM.getContext().getTypeSizeInChars() instead of CGM.getContext().getTypeSize() / 8.

2954

Maybe llvm::ConstantInt::get(CGF.SizeTy, SizeVal)?

2963

Use CGM.getContext().getTypeSizeInChars() instead of CGM.getContext().getTypeSize() / 8.

2980

The same

lib/CodeGen/CGStmtOpenMP.cpp

2144–2145

MakeNaturalAlignAddrLValue(ThisRef, ri->getType())?

2147

Extra semicolon

[sfantao]

Adress reviewers concerns.

Also fix issue with target regions with no arguments and in the VLA size computation I found in the meantime.

[sfantao]

Alexey,

Thanks for the review! Find my comments inlined.

Thanks again!
Samuel

lib/CodeGen/CGExpr.cpp

1969

Alexey, I'm not sure I understand what you mean here. Unlike the other captured regions, the target region outlined function does not take a context that captures all the variables in fields of a record as argument. Instead, it takes all the captured references as arguments. This will enable the device runtime library to decide what is best to pass the arguments to the device (see my response to John's question in my previous diff).

It happens that all the machinery in the common infrastructure that creates the outlined functions (CodeGenFunction::StartFunction and GenerateCapturedStmtFunction) is prepared to get the context record from the CapuredDecl. Therefore, in order to not disrupt the common infrastructure, in Sema::ActOnOpenMPRegionEnd I am creating a new CapturedDecl that contains implicit parameters. I gather the information to build the new CapturedDecl from the CapturedDecl that is created with the context argument and the RecordDecl fields so that I don't need to touch the capturing code in Sema.

Having CapturedDecl with implicit parameters will drive StartFunction to create the outlined region with the right signature without having to change anything in there. I only had to guard the initialization of VLAs and 'this' in GenerateCapturedStmtFunction to not do anything that expects the context argument. However, during the emission of the VLAs that happens in StartFunction, the emission of these implicit parameters is attempted based on references that are marked as refersToEnclosingVariableOrCapture- this is the reason for the change in EmitDeclRefLValue.

Given that the references in the outlined function statements are still expecting the VLAs expression used in the caller of the outlined function, PrepareOMPTargetDirectiveBodyEmission will make sure that the mapped values to those expressions will be the same as the ones that use the new expression based on implicit parameters.

Let me know if you need me to clarify anything.
Thanks!

lib/CodeGen/CGOpenMPRuntime.cpp

2953

Done!

2954

I agree, it makes more sense to use size_t. Thanks for the suggestion!

2963

done!

2980

Done!

lib/CodeGen/CGStmtOpenMP.cpp

2144–2145

Now using MakeNaturalAlignAddrLValue.

2147

Fixed.

[ABataev]

lib/CodeGen/CGExpr.cpp

1969–1971

Instead I would do the same thing I did for tasks.
Generate captured function as is. But also create the second function with the profile required for target codegen. This function must gather all its parameters into a single record and then call auto generated captured function. This captured function must be marked as AlwaysInline.
In this case you don't need to make some additional changes in Sema for particular processing of target directives, you will just need to generate simple function in codegen.
This may result in a little bit slower performance, but we may improve it later, when we have time to improve codegen for outlined functions for CapturedDecls.
I don't like the idea of reinventing of features, that are invented already, like capturing of VLAs, exposing some private functions (like getVLASizeMap()) etc.
so the code would be like this:

void .omp_outlined.(%captures *ptr) always_inline {
 <captured code>;
}
void .target_omp_outlined.(int* a, float* b, ...) {
  %captures rec;
  rec.a_ref = a;
  rec.b_ref = b;
  ...
  .omp_outlined.(&rec);
}

lib/CodeGen/CGOpenMPRuntime.cpp

3065

Maybe CGF.getContext().getSizeType()?

3163

CGF.Builder.CreateIsNotNull()?

test/OpenMP/target_codegen.cpp

2–4

It would be good to see the tests for 32 bit target.

[sfantao]

Implement proxy function for target directive.

Move the creation of the target region parameters from CGOpenMPRuntime::emitTargetCall to CodeGenFunction::EmitTargetDirective because we need to access the VLA Maps of the target enclosing function.

Update regression test and add run directive for 32-bit target.

[sfantao]

Thanks for review. The new diff now uses a proxy function. See other comments inlined.

Thanks again!
Samuel

lib/CodeGen/CGExpr.cpp

1969–1971

Ok, I am now using the proxy function.

lib/CodeGen/CGOpenMPRuntime.cpp

3065

Done!

3163

Done!

test/OpenMP/target_codegen.cpp

2–4

Done!

[ABataev]

lib/CodeGen/CGOpenMPRuntime.cpp

2909

I don't think you need this argument. You're emitting a new outlined function here and don't need info about your current function.

2928–2933

You'd better to emit internal function separately in a new static function. Then you don't need to create TargetAuxCGF and TargetMainCGF. You should use just CGF everywhere. One CodeGenFunction instance per function.

2992–2994

Variable names should start with an upper case letter (e.g. Leader or Boats).

3092–3129

Try instead:

SizesArrayInit = llvm::ConstantArray::get(llvm::ArrayType::get(CGM.SizeTy, Sizes.size()), Sizes);

3183–3186

llvm::ConstantPointerNull::get(<type>);

lib/CodeGen/CGOpenMPRuntime.h

190–204

These enums must not be exposed by CGOpenMPRuntime until they are used in arguments of runtime member functions.

766

I don't like the idea of using 'unsigned' as a map type. You should create some particular OpenMPMapClauseKind (just like OpenMPDefaultClauseKind, OpenMPDependClauseKind, OpenMPProcBindClauseKind etc.) and use it where required.

lib/CodeGen/CGStmtOpenMP.cpp

2154–2218

I think this one can be moved to runtime codegen

[ABataev]

Another one thing I forget to mention. Current implementation of CGOpenMPRuntime is libomp-specific. You're trying to add functionality that is libtarget-specific. Maybe it is a good idea to separate support for libomp and libtarget runtime libraries?

[sfantao]

Address reviewer concerns.

[sfantao]

In D11361#229744, @ABataev wrote:

Another one thing I forget to mention. Current implementation of CGOpenMPRuntime is libomp-specific. You're trying to add functionality that is libtarget-specific. Maybe it is a good idea to separate support for libomp and libtarget runtime libraries?

Not sure what do you mean by separation. Different files? Different codegen class? My perspective is that the two things should be together given that they both address the same specification, and I see that interaction is required between the two components. E.g. teams codegen will have to interact with the target codegen (communicate number of teams/threads ) and the teams codegen will require the libomp interface in its implementation. We could always separate the two things in the future if we see that is a better way to organize the code.

lib/CodeGen/CGOpenMPRuntime.cpp

2909

Done!

2928–2933

Done!

2992–2994

Ok, thought iterators were an exception to that rule. Fixed now!

3092–3129

Done!

3183–3186

Done!

lib/CodeGen/CGStmtOpenMP.cpp

2154–2218

I moved this to codegen, but I still had to forward some VLA information that I can only get from CGF, so that the VLAMaps are exposed. I need that information so that I can initialize the target region parameters.

[sfantao]

Two more inlined comments that I forgot to integrate in my previous response.

Thanks!
Samuel

lib/CodeGen/CGOpenMPRuntime.h

190–204

Got it, not exposed anymore.

761

Unlike the other enums, more than one map types need to be combined. E.g., to/from are two different enums. Once the map clause and 4.1 get to be support, we will have more combinations. I see two options here: add enums for all combinations or use a typedef each time an ineger refer to map types, so the code is more readable. Let me know your thoughts.

[ABataev]

Samuel, Yes, I thought about different files and different classes. Runtime for offloading codegen is not a part of libomp and it would be good to have separate runtime handler class for target codegen also. We need to think about it in future.

lib/CodeGen/CGOpenMPRuntime.h

188–203

Move them to .cpp file.

761

Yes, I think we need to add separate enums for different combination in Basic/OpenMPKinds.{def, h} for AST. In runtime support library we can represent these combinations as a bit-or of single mapping types.

[sfantao]

Move map type and device id enums from CGOpenMPRuntime.h to CGOpenMPRuntime.cpp.

[sfantao]

Thanks for the review!

In D11361#232045, @ABataev wrote:

Samuel, Yes, I thought about different files and different classes. Runtime for offloading codegen is not a part of libomp and it would be good to have separate runtime handler class for target codegen also. We need to think about it in future.

Sure, we can definitely do something like that in the future if we detect a clear place to separate. I guess that when we do that, the design of the offloading-unrelated lib has to change a little. There are things like the capture statement info that may make sense to share between the two implementation.

lib/CodeGen/CGOpenMPRuntime.h

188–203

Done!

761

Ok, I see, it makes sense to do that from the AST for the map clause SEMA. I'm keeping the bit-or in the runtime library call codegen as you say.

[ABataev]

lib/CodeGen/CGOpenMPRuntime.cpp

3020–3034

Move them to CGOpenMPRuntime::emitTargetCall(), they can be made local

test/OpenMP/target_codegen.cpp

9

Some of your tests has triple i386, they don't need PowerPC target

[sfantao]

Address last review comments.

lib/CodeGen/CGOpenMPRuntime.cpp

3020–3034

Ok, done!

test/OpenMP/target_codegen.cpp

9

True, I'm not using any target specific property here. So I guess it is safe to remove the requirement. Not using // REQUIRES anymore. Thanks.

[ABataev]

Seems good to me, but it would be good if John McCall could look at the patch.

[Hahnfeld]

Needs two small changes to work with current trunk

lib/CodeGen/CGStmtOpenMP.cpp

2139–2140

This now has to be S.getSingleClause<OMPIfClause>() and we can therefore omit the cast...

2145–2146

Likewise S.getSingleClause<OMPDeviceClause>() without the need for an extra cast

[rjmccall]

Sorry for putting off the final review on this; I was heads-down trying to get the alignment patch done. It's looking good; obviously you'll need to update it to work with Addresses properly, but hopefully that won't be too much of a problem.

When you do, maybe you should start a new review; I think there's some way to do that in Phabricator that ties it to the old one. Phabricator seems to not be very happy with the extent to which the code has changed, and the old comments now just make it harder to review the current patch.

[sfantao]

Closing revision. It has been replaced by http://reviews.llvm.org/D12871 has suggested by John.

Thanks!
Samuel

include/clang/AST/Decl.h

private:
   /// \brief The number of parameters to the outlined function.
   unsigned NumParams;
   /// \brief The position of context parameter in list of parameters.
   unsigned ContextParam;
   /// \brief The body of the outlined function.
   llvm::PointerIntPair<Stmt *, 1, bool> BodyAndNothrow;
 
   explicit CapturedDecl(DeclContext *DC, unsigned NumParams)
       : Decl(Captured, DC, SourceLocation()), DeclContext(Captured),
-      NumParams(NumParams), ContextParam(0), BodyAndNothrow(nullptr, false) { }
+        NumParams(NumParams), ContextParam(-1u),
+        BodyAndNothrow(nullptr, false) {}
 
   ImplicitParamDecl *const *getParams() const {
     return getTrailingObjects<ImplicitParamDecl *>();
   }
 
   ImplicitParamDecl **getParams() {
     return getTrailingObjects<ImplicitParamDecl *>();
   }
ImplicitParamDecl *getParam(unsigned i) const {
     assert(i < NumParams);
     return getParams()[i];
   }
   void setParam(unsigned i, ImplicitParamDecl *P) {
     assert(i < NumParams);
     getParams()[i] = P;
   }
 
+  bool hasContextParam() const { return ContextParam < NumParams; }
+
   /// \brief Retrieve the parameter containing captured variables.
   ImplicitParamDecl *getContextParam() const {
-    assert(ContextParam < NumParams);
+    assert(hasContextParam());
     return getParam(ContextParam);
   }
   void setContextParam(unsigned i, ImplicitParamDecl *P) {
     assert(i < NumParams);
     ContextParam = i;
     setParam(i, P);
   }
   unsigned getContextParamPosition() const { return ContextParam; }

include/clang/AST/Stmt.h

public:
   };
 
 private:
   /// \brief The number of variable captured, including 'this'.
   unsigned NumCaptures;
 
   /// \brief The pointer part is the implicit the outlined function and the 
   /// int part is the captured region kind, 'CR_Default' etc.
-  llvm::PointerIntPair<CapturedDecl *, 1, CapturedRegionKind> CapDeclAndKind;
+  llvm::PointerIntPair<CapturedDecl *, 2, CapturedRegionKind> CapDeclAndKind;
 
   /// \brief The record for captured variables, a RecordDecl or CXXRecordDecl.
   RecordDecl *TheRecordDecl;
 
   /// \brief Construct a captured statement.
   CapturedStmt(Stmt *S, CapturedRegionKind Kind, ArrayRef<Capture> Captures,
                ArrayRef<Expr *> CaptureInits, CapturedDecl *CD, RecordDecl *RD);
 

include/clang/Basic/CapturedStmt.h

 #ifndef LLVM_CLANG_BASIC_CAPTUREDSTMT_H
 #define LLVM_CLANG_BASIC_CAPTUREDSTMT_H
 
 namespace clang {
 
 /// \brief The different kinds of captured statement.
 enum CapturedRegionKind {
   CR_Default,
-  CR_OpenMP
+  CR_OpenMP,
+  CR_OpenMP_Target
 };
 
 } // end namespace clang
 
 #endif // LLVM_CLANG_BASIC_CAPTUREDSTMT_H

include/clang/Sema/ScopeInfo.h

public:
 
   /// \brief A descriptive name for the kind of captured region this is.
   StringRef getRegionName() const {
     switch (CapRegionKind) {
     case CR_Default:
       return "default captured statement";
     case CR_OpenMP:
       return "OpenMP region";
+    case CR_OpenMP_Target:
+      return "OpenMP target region";
     }
     llvm_unreachable("Invalid captured region kind!");
   }
 
   static bool classof(const FunctionScopeInfo *FSI) {
     return FSI->Kind == SK_CapturedRegion;
   }
 };

lib/CodeGen/CGExpr.cpp

 
 static LValue EmitCapturedFieldLValue(CodeGenFunction &CGF, const FieldDecl *FD,
                                       llvm::Value *ThisValue) {
   QualType TagType = CGF.getContext().getTagDeclType(FD->getParent());
   LValue LV = CGF.MakeNaturalAlignAddrLValue(ThisValue, TagType);
   return CGF.EmitLValueForField(LV, FD);
 }
 
+static LValue EmitCapturedImplicitParamLValue(CodeGenFunction &CGF,
+                                              const ImplicitParamDecl *PD) {
+  // If the captured declaration is an implicit parameter, it should already
+  // exist in the local declaration map.
+  LValue LV = LValue::MakeAddr(CGF.GetAddrOfLocalVar(PD), PD->getType(),
+                               CharUnits(), CGF.getContext());
+  RValue RV = CGF.EmitLoadOfLValue(LV, PD->getLocStart());
+  return LValue::MakeAddr(RV.getScalarVal(), PD->getType(), CharUnits(),
+                          CGF.getContext());
+}
+
+static LValue EmitCapturedValue(CodeGenFunction &CGF, const Decl *D,
+                                llvm::Value *ThisValue) {
+  switch (D->getKind()) {
+  default:
+    llvm_unreachable("Unexpected declaration kind for capture!");
+  case Decl::ImplicitParam:
+    return EmitCapturedImplicitParamLValue(CGF, cast<ImplicitParamDecl>(D));
+  case Decl::Field:
+    return EmitCapturedFieldLValue(CGF, cast<FieldDecl>(D), ThisValue);
+  }
+}
+
 /// Named Registers are named metadata pointing to the register name
 /// which will be read from/written to as an argument to the intrinsic
 /// @llvm.read/write_register.
 /// So far, only the name is being passed down, but other options such as
 /// register type, allocation type or even optimization options could be
 /// passed down via the metadata node.
 static LValue EmitGlobalNamedRegister(const VarDecl *VD,
                                       CodeGenModule &CGM,
if (const auto *VD = dyn_cast<VarDecl>(ND)) {
     // Check for captured variables.
     if (E->refersToEnclosingVariableOrCapture()) {
       if (auto *FD = LambdaCaptureFields.lookup(VD))
         return EmitCapturedFieldLValue(*this, FD, CXXABIThisValue);
       else if (CapturedStmtInfo) {
         if (auto *V = LocalDeclMap.lookup(VD))
           return MakeAddrLValue(V, T, Alignment);
         else
-          return EmitCapturedFieldLValue(*this, CapturedStmtInfo->lookup(VD),
-                                         CapturedStmtInfo->getContextValue());
+          return EmitCapturedValue(*this, CapturedStmtInfo->lookup(VD),
       }
       assert(isa<BlockDecl>(CurCodeDecl));
       return MakeAddrLValue(GetAddrOfBlockDecl(VD, VD->hasAttr<BlocksAttr>()),
                             T, Alignment);
     }
   }
 
   // FIXME: We should be able to assert this for FunctionDecls as well!
   // FIXME: We should be able to assert this for all DeclRefExprs, not just

lib/CodeGen/CGOpenMPRuntime.h

enum OpenMPRTLFunction {
     // ndeps_noalias, kmp_depend_info_t *noalias_dep_list);
     OMPRTL__kmpc_omp_wait_deps,
     // Call to kmp_int32 __kmpc_cancellationpoint(ident_t *loc, kmp_int32
     // global_tid, kmp_int32 cncl_kind);
     OMPRTL__kmpc_cancellationpoint,
     // Call to kmp_int32 __kmpc_cancel(ident_t *loc, kmp_int32 global_tid,
     // kmp_int32 cncl_kind);
     OMPRTL__kmpc_cancel,
+
+    //
+    // Offloading related calls
+    //
+    // Call to int32_t __tgt_target(int32_t device_id, void *host_ptr, int32_t
+    // arg_num, void** args_base, void **args, int64_t *arg_sizes, int32_t
+    // *arg_types);
+    OMPRTL__tgt_target,
   };
 
   /// \brief Values for bit flags used in the ident_t to describe the fields.
   /// All enumeric elements are named and described in accordance with the code
   /// from http://llvm.org/svn/llvm-project/openmp/trunk/runtime/src/kmp.h
   enum OpenMPLocationFlags {
     /// \brief Use trampoline for internal microtask.
     OMP_IDENT_IMD = 0x01,
     /// \brief Use c-style ident structure.
     OMP_IDENT_KMPC = 0x02,
     /// \brief Atomic reduction option for kmpc_reduce.
     OMP_ATOMIC_REDUCE = 0x10,
     /// \brief Explicit 'barrier' directive.
     OMP_IDENT_BARRIER_EXPL = 0x20,
     /// \brief Implicit barrier in code.
     OMP_IDENT_BARRIER_IMPL = 0x40,
     /// \brief Implicit barrier in 'for' directive.
     OMP_IDENT_BARRIER_IMPL_FOR = 0x40,
     /// \brief Implicit barrier in 'sections' directive.
     OMP_IDENT_BARRIER_IMPL_SECTIONS = 0xC0,
     /// \brief Implicit barrier in 'single' directive.
     OMP_IDENT_BARRIER_IMPL_SINGLE = 0x140
   };
+
+  /// \brief Values for bit flags used to specify the mapping type for
+  /// offloading.
+  enum OpenMPOffloadMappingFlags {
+    /// \brief Allocate memory on the device and move data from host to device.
+    OMP_MAP_TO = 0x01,
+    /// \brief Allocate memory on the device and move data from device to host.
+    OMP_MAP_FROM = 0x02,
+  };
+
+  enum OpenMPOffloadingReservedDeviceIDs {
+    /// \brief Device ID if the device was not defined, runtime should get it
+    /// from environment variables in the spec.
+    OMP_DEVICEID_UNDEF = -1,
+  };
+
   CodeGenModule &CGM;
   /// \brief Default const ident_t object used for initialization of all other
   /// ident_t objects.
   llvm::Constant *DefaultOpenMPPSource;
   /// \brief Map of flags and corresponding default locations.
   typedef llvm::DenseMap<unsigned, llvm::Value *> OpenMPDefaultLocMapTy;
   OpenMPDefaultLocMapTy OpenMPDefaultLocMap;
   llvm::Value *getOrCreateDefaultLocation(OpenMPLocationFlags Flags);
   /// \brief Describes ident structure that describes a source location.
   /// All descriptions are taken from
   /// http://llvm.org/svn/llvm-project/openmp/trunk/runtime/src/kmp.h
   /// Original structure:
   /// typedef struct ident {
   ///    kmp_int32 reserved_1;   /**<  might be used in Fortran;
   ///                                  see above  */
   ///    kmp_int32 flags;        /**<  also f.flags; KMP_IDENT_xxx flags;
   ///                                  KMP_IDENT_KMPC identifies this union
virtual void emitCancellationPointCall(CodeGenFunction &CGF,
                                          SourceLocation Loc,
                                          OpenMPDirectiveKind CancelRegion);
 
   /// \brief Emit code for 'cancel' construct.
   /// \param CancelRegion Region kind for which the cancel must be emitted.
   ///
   virtual void emitCancelCall(CodeGenFunction &CGF, SourceLocation Loc,
                               OpenMPDirectiveKind CancelRegion);
+
+  /// \brief Emit outilined function for 'target' directive.
+  /// \param D Directive to emit.
+  /// \param CodeGen Code generation sequence for the \a D directive.
+  virtual llvm::Value *
+  emitTargetOutlinedFunction(CodeGenFunction &CGF,
+                             const OMPExecutableDirective &D,
+                             const RegionCodeGenTy &CodeGen);
+
+  /// \brief Emit the target offloading code associated with \a D. The emitted
+  /// code attempts offloading the execution to the device, an the event of
+  /// a failure it executes the host version outlined in \a OutlinedFn.
+  /// \param D Directive to emit.
+  /// \param OutlinedFn Host version of the code to be offloaded.
+  /// \param IfCond Expression evaluated in if clause associated with the target
+  /// directive, or null if no if clause is used.
+  /// \param Device Expression evaluated in device clause associated with the
+  /// target directive, or null if no device clause is used.
+  virtual void emitTargetCall(CodeGenFunction &CGF,
+                              const OMPExecutableDirective &D,
+                              llvm::Value *OutlinedFn, const Expr *IfCond,
+                              const Expr *Device);
 };
 
 } // namespace CodeGen
 } // namespace clang
 
 #endif

lib/CodeGen/CGOpenMPRuntime.cpp

enum CGOpenMPRegionKind {
     /// \brief Region with outlined function for standalone 'parallel'
     /// directive.
     ParallelOutlinedRegion,
     /// \brief Region with outlined function for standalone 'task' directive.
     TaskOutlinedRegion,
     /// \brief Region for constructs that do not require function outlining,
     /// like 'for', 'sections', 'atomic' etc. directives.
     InlinedRegion,
+    /// \brief Region with outlined function for standalone 'target' directive.
+    TargetRegion,
   };
 
   CGOpenMPRegionInfo(const CapturedStmt &CS,
                      const CGOpenMPRegionKind RegionKind,
                      const RegionCodeGenTy &CodeGen, OpenMPDirectiveKind Kind)
       : CGCapturedStmtInfo(CS, CR_OpenMP), RegionKind(RegionKind),
         CodeGen(CodeGen), Kind(Kind) {}
 
public:
   virtual void setContextValue(llvm::Value *V) override {
     if (OuterRegionInfo) {
       OuterRegionInfo->setContextValue(V);
       return;
     }
     llvm_unreachable("No context value for inlined OpenMP region");
   }
   /// \brief Lookup the captured field decl for a variable.
-  const FieldDecl *lookup(const VarDecl *VD) const override {
+  const Decl *lookup(const VarDecl *VD) const override {
     if (OuterRegionInfo)
       return OuterRegionInfo->lookup(VD);
     // If there is no outer outlined region,no need to lookup in a list of
     // captured variables, we can use the original one.
     return nullptr;
   }
-  FieldDecl *getThisFieldDecl() const override {
+  Decl *getThisDecl() const override {
     if (OuterRegionInfo)
-      return OuterRegionInfo->getThisFieldDecl();
+      return OuterRegionInfo->getThisDecl();
     return nullptr;
   }
   /// \brief Get a variable or parameter for storing global thread id
   /// inside OpenMP construct.
   const VarDecl *getThreadIDVariable() const override {
     if (OuterRegionInfo)
       return OuterRegionInfo->getThreadIDVariable();
     return nullptr;
public:
   }
 
 private:
   /// \brief CodeGen info about outer OpenMP region.
   CodeGenFunction::CGCapturedStmtInfo *OldCSI;
   CGOpenMPRegionInfo *OuterRegionInfo;
 };
 
+/// \brief API for captured statement code generation in OpenMP target
+/// constructs. For this captures, implicit parameters are used instead of the
+/// captured fields.
+class CGOpenMPTargetRegionInfo : public CGOpenMPRegionInfo {
+public:
+  CGOpenMPTargetRegionInfo(const CapturedStmt &CS,
+                           const RegionCodeGenTy &CodeGen)
+      : CGOpenMPRegionInfo(TargetRegion, CodeGen, OMPD_target),
+        CXXThisImplicitParamDecl(nullptr) {
+
+    CapturedDecl::param_iterator Param = CS.getCapturedDecl()->param_begin();
+    for (CapturedStmt::const_capture_iterator I = CS.capture_begin(),
+                                              E = CS.capture_end();
+         I != E; ++I, ++Param) {
+      if (I->capturesThis())
+        CXXThisImplicitParamDecl = *Param;
+      else if (I->capturesVariable())
+        CaptureImplicitParams[I->getCapturedVar()] = *Param;
+    }
+  }
+
+  /// \brief This is unused for target regions because each starts executing
+  /// with a single thread.
+  const VarDecl *getThreadIDVariable() const override { return nullptr; }
+
+  /// \brief Get the name of the capture helper.
+  StringRef getHelperName() const override { return ".omp_offloading."; }
+
+  static bool classof(const CGCapturedStmtInfo *Info) {
+    return CGOpenMPRegionInfo::classof(Info) &&
+           cast<CGOpenMPRegionInfo>(Info)->getRegionKind() == TargetRegion;
+  }
+
+  /// \brief Lookup the captured implicit parameter declaration for a variable.
+  virtual const Decl *lookup(const VarDecl *VD) const override {
+    return CaptureImplicitParams.lookup(VD);
+  }
+
+  virtual Decl *getThisDecl() const override {
+    return CXXThisImplicitParamDecl;
+  }
+
+private:
+  /// \brief Keep the map between VarDecl and FieldDecl.
+  llvm::SmallDenseMap<const VarDecl *, ImplicitParamDecl *>
+      CaptureImplicitParams;
+
+  /// \brief Captured 'this' type.
+  ImplicitParamDecl *CXXThisImplicitParamDecl;
+};
+
 /// \brief RAII for emitting code of OpenMP constructs.
 class InlinedOpenMPRegionRAII {
   CodeGenFunction &CGF;
 
 public:
   /// \brief Constructs region for combined constructs.
   /// \param CodeGen Code generation sequence for combined directives. Includes
   /// a list of functions used for code generation of implicitly inlined
case OMPRTL__kmpc_cancel: {
     // Build kmp_int32 __kmpc_cancel(ident_t *loc, kmp_int32 global_tid,
     // kmp_int32 cncl_kind)
     llvm::Type *TypeParams[] = {getIdentTyPointerTy(), CGM.Int32Ty, CGM.IntTy};
     llvm::FunctionType *FnTy =
         llvm::FunctionType::get(CGM.Int32Ty, TypeParams, /*isVarArg*/ false);
     RTLFn = CGM.CreateRuntimeFunction(FnTy, "__kmpc_cancel");
     break;
   }
+  case OMPRTL__tgt_target: {
+    // Build int32_t __tgt_target(int32_t device_id, void *host_ptr, int32_t
+    // arg_num, void** args_base, void **args, int64_t *arg_sizes, int32_t
+    // *arg_types);
+    llvm::Type *TypeParams[] = {CGM.Int32Ty,
+                                CGM.VoidPtrTy,
+                                CGM.Int32Ty,
+                                CGM.VoidPtrPtrTy,
+                                CGM.VoidPtrPtrTy,
+                                CGM.Int64Ty->getPointerTo(),
+                                CGM.Int32Ty->getPointerTo()};
+    llvm::FunctionType *FnTy =
+        llvm::FunctionType::get(CGM.Int32Ty, TypeParams, /*isVarArg*/ false);
+    RTLFn = CGM.CreateRuntimeFunction(FnTy, "__tgt_target");
+    break;
+  }
   }
   return RTLFn;
 }
 
 llvm::Constant *CGOpenMPRuntime::createForStaticInitFunction(unsigned IVSize,
                                                              bool IVSigned) {
   assert((IVSize == 32 || IVSize == 64) &&
          "IV size is not compatible with the omp runtime");
CodeGenFunction::CGCapturedStmtInfo CapturesInfo(
         cast<CapturedStmt>(*D.getAssociatedStmt()));
     for (auto &&Pair : Privates) {
       auto *VD = Pair.second.PrivateCopy;
       auto *Init = VD->getAnyInitializer();
       LValue PrivateLValue = CGF.EmitLValueForField(PrivatesBase, *FI);
       if (Init) {
         if (auto *Elem = Pair.second.PrivateElemInit) {
           auto *OriginalVD = Pair.second.Original;
-          auto *SharedField = CapturesInfo.lookup(OriginalVD);
+          auto *SharedField = cast<FieldDecl>(CapturesInfo.lookup(OriginalVD));
           auto SharedRefLValue =
               CGF.EmitLValueForField(SharedsBase, SharedField);
           QualType Type = OriginalVD->getType();
           if (Type->isArrayType()) {
             // Initialize firstprivate array.
             if (!isa<CXXConstructExpr>(Init) ||
                 CGF.isTrivialInitializer(Init)) {
               // Perform simple memcpy.
if (CancelDest.isValid()) {
       emitBarrierCall(CGF, Loc, OMPD_unknown, /*CheckForCancel=*/false);
       // exit from construct;
       CGF.EmitBranchThroughCleanup(CancelDest);
       CGF.EmitBlock(ContBB, /*IsFinished=*/true);
     }
   }
 }
 
+llvm::Value *
+CGOpenMPRuntime::emitTargetOutlinedFunction(CodeGenFunction &CGF,
+                                            const OMPExecutableDirective &D,
+                                            const RegionCodeGenTy &CodeGen) {
+
+  const CapturedStmt *CS = cast<CapturedStmt>(D.getAssociatedStmt());
+
+  // Generate the outlined target offloading function.
+  CodeGenFunction TargetCGF(CGM, true);
+  CGOpenMPTargetRegionInfo CGInfo(*CS, CodeGen);
+  CodeGenFunction::CGCapturedStmtRAII CapInfoRAII(TargetCGF, &CGInfo);
+  return TargetCGF.GenerateCapturedStmtFunction(*CS);
+}
+
+void CGOpenMPRuntime::emitTargetCall(CodeGenFunction &CGF,
+                                     const OMPExecutableDirective &D,
+                                     llvm::Value *OutlinedFn,
+                                     const Expr *IfCond, const Expr *Device) {
+
+  // Fill up the arrays with the all the captured variables.
+  SmallVector<llvm::Value *, 16> BasePointers;
+  SmallVector<llvm::Value *, 16> Pointers;
+  SmallVector<llvm::Value *, 16> Sizes;
+  SmallVector<unsigned, 16> MapTypes;
+
+  bool hasVLACaptures = false;
+  const CapturedStmt &CS = *cast<CapturedStmt>(D.getAssociatedStmt());
+  auto ri = CS.getCapturedRecordDecl()->field_begin();
+  auto ii = CS.capture_init_begin();
+  for (CapturedStmt::const_capture_iterator ci = CS.capture_begin(),
+                                            ce = CS.capture_end();
+       ci != ce; ++ci, ++ri, ++ii) {
+    StringRef Name;
+    QualType Ty;
+    llvm::Value *BasePointer;
+    llvm::Value *Pointer;
+    llvm::Value *Size;
+    unsigned MapType;
+
+    if (ci->capturesVariableArrayType()) {
+      llvm::Value *V = CGF.getVLASizeMap(ri->getCapturedVLAType());
+      LValue LV = CGF.MakeNaturalAlignAddrLValue(
+          CGF.CreateMemTemp(ri->getType(), "__vla_size"), ri->getType());
+      CGF.EmitStoreThroughLValue(RValue::get(V), LV);
+      BasePointer = Pointer = LV.getAddress();
+      uint64_t SizeVal = CGM.getContext().getTypeSize(ri->getType()) / 8;
+      Size = CGF.Builder.getInt64(SizeVal);
+
+      hasVLACaptures = true;
+      // VLA sizes don't need to be copied back from the device.
+      MapType = CGOpenMPRuntime::OMP_MAP_TO;
+    } else if (ci->capturesThis()) {
+      BasePointer = Pointer = CGF.LoadCXXThis();
+      const PointerType *PtrTy = cast<PointerType>(ri->getType().getTypePtr());
+      uint64_t SizeVal =
+          CGM.getContext().getTypeSize(PtrTy->getPointeeType()) / 8;
+      Size = CGF.Builder.getInt64(SizeVal);
+
+      // Default map type.
+      MapType = CGOpenMPRuntime::OMP_MAP_TO | CGOpenMPRuntime::OMP_MAP_FROM;
+    } else {
+      BasePointer = Pointer =
+          CGF.EmitLValue(cast<DeclRefExpr>(*ii)).getAddress();
+
+      const ReferenceType *PtrTy =
+          cast<ReferenceType>(ri->getType().getTypePtr());
+      QualType ElementType = PtrTy->getPointeeType();
+
+      if (auto *VAT = dyn_cast<VariableArrayType>(ElementType.getTypePtr()))
+        Size = CGF.getVLASize(VAT).first;
+      else {
+        uint64_t ElementTypeSize =
+            CGM.getContext().getTypeSize(ElementType) / 8;
+        Size = CGF.Builder.getInt64(ElementTypeSize);
+      }
+
+      // Default map type.
+      MapType = CGOpenMPRuntime::OMP_MAP_TO | CGOpenMPRuntime::OMP_MAP_FROM;
+    }
+
+    BasePointers.push_back(BasePointer);
+    Pointers.push_back(Pointer);
+    Sizes.push_back(Size);
+    MapTypes.push_back(MapType);
+  }
+
+  if (IfCond) {
+    // Check if the if clause conditional always evaluates to true or false.
+    // If it evaluates to false, we only need to emit the host version of the
+    // target region. If it evaluates to true, we can proceed with the codegen
+    // as if no if clause was provided.
+    bool CondConstant;
+    if (CGF.ConstantFoldsToSimpleInteger(IfCond, CondConstant)) {
+      if (CondConstant) {
+        IfCond = nullptr;
+      } else {
+        CGF.Builder.CreateCall(OutlinedFn, BasePointers);
+        return;
+      }
+    }
+  }
+
+  // Generate the code to launch the target region. The pattern is the
+  // following:
+  //
+  //   ...
+  //   br IfCond (if any), omp_offload, omp_offload_fail
+  //
+  // omp_offload.try:
+  //   ; create arrays for offloading
+  //   error = __tgt_target(...)
+  //   br error, omp_offload_fail, omp_offload_end
+  //
+  // omp_offload.fail:
+  //   host_version(...)
+  //
+  // omp_offload.end:
+  //   ...
+  //
+
+  auto OffloadTryBlock = CGF.createBasicBlock("omp_offload.try");
+  auto OffloadFailBlock = CGF.createBasicBlock("omp_offload.fail");
+  auto ContBlock = CGF.createBasicBlock("omp_offload.end");
+
+  if (IfCond)
+    CGF.EmitBranchOnBoolExpr(IfCond, OffloadTryBlock, OffloadFailBlock,
+                             /*TrueCount=*/0);
+
+  CGF.EmitBlock(OffloadTryBlock);
+
+  unsigned PointerNumVal = BasePointers.size();
+  llvm::Value *PointerNum = CGF.Builder.getInt32(PointerNumVal);
+  llvm::Value *BasePointersArray;
+  llvm::Value *PointersArray;
+  llvm::Value *SizesArray;
+  llvm::Value *MapTypesArray;
+
+  if (PointerNumVal) {
+    llvm::APInt PointerNumAP(32, PointerNumVal, /*isSigned=*/true);
+    QualType PointerArrayType = CGF.getContext().getConstantArrayType(
+        CGF.getContext().VoidPtrTy, PointerNumAP, ArrayType::Normal,
+        /*IndexTypeQuals=*/0);
+
+    BasePointersArray =
+        CGF.CreateMemTemp(PointerArrayType, ".offload_baseptrs");
+    PointersArray = CGF.CreateMemTemp(PointerArrayType, ".offload_ptrs");
+
+    // If we don't have any VLA types, we can use a constant array for the map
+    // sizes, otherwise we need to fill up the arrays as we do for the pointers.
+    if (hasVLACaptures) {
+      QualType SizeArrayType = CGF.getContext().getConstantArrayType(
+          CGF.getContext().getIntTypeForBitwidth(64, /*Signed=*/true),
+          PointerNumAP, ArrayType::Normal, /*IndexTypeQuals=*/0);
+      SizesArray = CGF.CreateMemTemp(SizeArrayType, ".offload_sizes");
+    } else {
+      // We expect all the sizes to be constant, so we collect them to create
+      // a constant array.
+      SmallVector<uint64_t, 16> ConstSizes;
+      for (auto *V : Sizes)
+        ConstSizes.push_back(cast<llvm::ConstantInt>(V)->getZExtValue());
+
+      llvm::Constant *SizesArrayInit =
+          llvm::ConstantDataArray::get(CGF.Builder.getContext(), ConstSizes);
+      auto *SizesArrayGbl = new llvm::GlobalVariable(
+          CGM.getModule(), SizesArrayInit->getType(),
+          /*isConstant=*/true, llvm::GlobalValue::PrivateLinkage,
+          SizesArrayInit, ".offload_sizes");
+      SizesArrayGbl->setUnnamedAddr(true);
+      SizesArray = SizesArrayGbl;
+    }
+
+    // The map types are always constant so we don't need to generate code to
+    // fill arrays. Instead, we create an array constant.
+    llvm::Constant *MapTypesArrayInit =
+        llvm::ConstantDataArray::get(CGF.Builder.getContext(), MapTypes);
+    auto *MapTypesArrayGbl = new llvm::GlobalVariable(
+        CGM.getModule(), MapTypesArrayInit->getType(),
+        /*isConstant=*/true, llvm::GlobalValue::PrivateLinkage,
+        MapTypesArrayInit, ".offload_maptypes");
+    MapTypesArrayGbl->setUnnamedAddr(true);
+    MapTypesArray = MapTypesArrayGbl;
+
+    for (unsigned i = 0; i < PointerNumVal; ++i) {
+      llvm::Value *BP = CGF.Builder.CreateConstInBoundsGEP2_32(
+          llvm::ArrayType::get(CGM.VoidPtrTy, PointerNumVal), BasePointersArray,
+          0, i);
+      CGF.Builder.CreateStore(
+          CGF.Builder.CreateBitCast(BasePointers[i], CGM.VoidPtrTy), BP);
+
+      llvm::Value *P = CGF.Builder.CreateConstInBoundsGEP2_32(
+          llvm::ArrayType::get(CGM.VoidPtrTy, PointerNumVal), PointersArray, 0,
+          i);
+      CGF.Builder.CreateStore(
+          CGF.Builder.CreateBitCast(Pointers[i], CGM.VoidPtrTy), P);
+
+      if (hasVLACaptures) {
+        llvm::Value *S = CGF.Builder.CreateConstInBoundsGEP2_32(
+            llvm::ArrayType::get(CGM.Int64Ty, PointerNumVal), SizesArray, 0, i);
+        CGF.Builder.CreateStore(
+            CGF.Builder.CreateIntCast(Sizes[i], CGM.Int64Ty, /*isSigned=*/true),
+            S);
+      }
+    }
+  } else {
+    BasePointersArray = llvm::Constant::getNullValue(CGM.VoidPtrPtrTy);
+    PointersArray = llvm::Constant::getNullValue(CGM.VoidPtrPtrTy);
+    SizesArray = llvm::Constant::getNullValue(CGM.Int64Ty->getPointerTo());
+    MapTypesArray = llvm::Constant::getNullValue(CGM.Int32Ty->getPointerTo());
+  }
+
+  // On top of the arrays that were filled up, the target offloading call takes
+  // as arguments the device id as well as the host pointer. The host pointer
+  // is used by the runtime library to identify the current target region, so
+  // it only has to be unique and not necessarily point to anything. It could be
+  // the pointer to the outlined function that implements the target region, but
+  // we aren't using that so that the compiler doesn't need to keep that, and
+  // could therefore inline the host function if proven worthwhile during
+  // optimization.
+
+  llvm::Value *HostPtr = new llvm::GlobalVariable(
+      CGM.getModule(), CGM.Int8Ty, /*isConstant=*/true,
+      llvm::GlobalValue::PrivateLinkage,
+      llvm::Constant::getNullValue(CGM.Int8Ty), ".offload_hstptr");
+
+  // Emit device ID if any.
+  llvm::Value *DeviceID;
+  if (Device)
+    DeviceID = CGF.Builder.CreateIntCast(CGF.EmitScalarExpr(Device),
+                                         CGM.Int32Ty, /*isSigned=*/true);
+  else
+    DeviceID = CGF.Builder.getInt32(OMP_DEVICEID_UNDEF);
+
+  llvm::Value *OffloadingArgs[] = {
+      DeviceID, HostPtr, PointerNum,
+      CGF.Builder.CreateConstGEP2_32(
+          llvm::ArrayType::get(CGM.VoidPtrTy, PointerNumVal), BasePointersArray,
+          0, 0),
+      CGF.Builder.CreateConstGEP2_32(
+          llvm::ArrayType::get(CGM.VoidPtrTy, PointerNumVal), PointersArray, 0,
+          0),
+      CGF.Builder.CreateConstGEP2_32(
+          llvm::ArrayType::get(CGM.Int64Ty, PointerNumVal), SizesArray, 0, 0),
+      CGF.Builder.CreateConstGEP2_32(
+          llvm::ArrayType::get(CGM.Int32Ty, PointerNumVal), MapTypesArray, 0,
+          0)};
+  auto Return = CGF.EmitRuntimeCall(createRuntimeFunction(OMPRTL__tgt_target),
+                                    OffloadingArgs);
+  auto Error = CGF.Builder.CreateICmpNE(Return, CGF.Builder.getInt32(0));
+  CGF.Builder.CreateCondBr(Error, OffloadFailBlock, ContBlock);
+
+  CGF.EmitBlock(OffloadFailBlock);
+  CGF.Builder.CreateCall(OutlinedFn, BasePointers);
+  CGF.EmitBranch(ContBlock);
+
+  CGF.EmitBlock(ContBlock, /*IsFinished=*/true);
+  return;
+}

lib/CodeGen/CGStmt.cpp

CodeGenFunction::GenerateCapturedStmtFunction(const CapturedStmt &S) {
   const CapturedDecl *CD = S.getCapturedDecl();
   const RecordDecl *RD = S.getCapturedRecordDecl();
   SourceLocation Loc = S.getLocStart();
   assert(CD->hasBody() && "missing CapturedDecl body");
 
   // Build the argument list.
   ASTContext &Ctx = CGM.getContext();
   FunctionArgList Args;
   Args.append(CD->param_begin(), CD->param_end());
 
   // Create the function declaration.
   FunctionType::ExtInfo ExtInfo;
   const CGFunctionInfo &FuncInfo =
       CGM.getTypes().arrangeFreeFunctionDeclaration(Ctx.VoidTy, Args, ExtInfo,
                                                     /*IsVariadic=*/false);
   llvm::FunctionType *FuncLLVMTy = CGM.getTypes().GetFunctionType(FuncInfo);
 
   llvm::Function *F =
     llvm::Function::Create(FuncLLVMTy, llvm::GlobalValue::InternalLinkage,
                            CapturedStmtInfo->getHelperName(), &CGM.getModule());
   CGM.SetInternalFunctionAttributes(CD, F, FuncInfo);
   if (CD->isNothrow())
     F->addFnAttr(llvm::Attribute::NoUnwind);
 
   // Generate the function.
   StartFunction(CD, Ctx.VoidTy, F, FuncInfo, Args,
                 CD->getLocation(),
                 CD->getBody()->getLocStart());
-  // Set the context parameter in CapturedStmtInfo.
+
+  // Initialize variable length arrays and 'this' using the context argument,
+  // if any. Otherwise the function implicit parameters will be used.
+  if (CD->hasContextParam()) {
     llvm::Value *DeclPtr = LocalDeclMap[CD->getContextParam()];
     assert(DeclPtr && "missing context parameter for CapturedStmt");
     CapturedStmtInfo->setContextValue(Builder.CreateLoad(DeclPtr));
 
     // Initialize variable-length arrays.
-  LValue Base = MakeNaturalAlignAddrLValue(CapturedStmtInfo->getContextValue(),
-                                           Ctx.getTagDeclType(RD));
+    LValue Base = MakeNaturalAlignAddrLValue(
+        CapturedStmtInfo->getContextValue(), Ctx.getTagDeclType(RD));
     for (auto *FD : RD->fields()) {
       if (FD->hasCapturedVLAType()) {
-      auto *ExprArg = EmitLoadOfLValue(EmitLValueForField(Base, FD),
-                                       S.getLocStart()).getScalarVal();
+        auto *ExprArg =
+            EmitLoadOfLValue(EmitLValueForField(Base, FD), S.getLocStart())
+                .getScalarVal();
         auto VAT = FD->getCapturedVLAType();
         VLASizeMap[VAT->getSizeExpr()] = ExprArg;
       }
     }
 
     // If 'this' is captured, load it into CXXThisValue.
     if (CapturedStmtInfo->isCXXThisExprCaptured()) {
-    FieldDecl *FD = CapturedStmtInfo->getThisFieldDecl();
+      FieldDecl *FD = cast<FieldDecl>(CapturedStmtInfo->getThisDecl());
       LValue ThisLValue = EmitLValueForField(Base, FD);
       CXXThisValue = EmitLoadOfLValue(ThisLValue, Loc).getScalarVal();
     }
+  }
 
   PGO.assignRegionCounters(CD, F);
   CapturedStmtInfo->EmitBody(*this, CD->getBody());
   FinishFunction(CD->getBodyRBrace());
 
   return F;
 }

lib/CodeGen/CGStmtOpenMP.cpp

void CodeGenFunction::EmitOMPAtomicDirective(const OMPAtomicDirective &S) {
   auto &&CodeGen = [&S, Kind, IsSeqCst](CodeGenFunction &CGF) {
     EmitOMPAtomicExpr(CGF, Kind, IsSeqCst, S.isPostfixUpdate(), S.getX(),
                       S.getV(), S.getExpr(), S.getUpdateExpr(),
                       S.isXLHSInRHSPart(), S.getLocStart());
   };
   CGM.getOpenMPRuntime().emitInlinedDirective(*this, OMPD_atomic, CodeGen);
 }
 
-void CodeGenFunction::EmitOMPTargetDirective(const OMPTargetDirective &) {
-  llvm_unreachable("CodeGen for 'omp target' is not supported yet.");
+void CodeGenFunction::PrepareOMPTargetDirectiveBodyEmission(
+    const OMPTargetDirective &S) {
+  const CapturedStmt &CS = *cast<CapturedStmt>(S.getAssociatedStmt());
+  const RecordDecl *RD = CS.getCapturedRecordDecl();
+  const CapturedDecl *CD = CS.getCapturedDecl();
+
+  SmallVector<const FieldDecl *, 3> VLASizeFields;
+  auto ri = RD->field_begin();
+  auto pi = CD->param_begin();
+  for (CapturedStmt::const_capture_iterator ci = CS.capture_begin(),
+                                            ce = CS.capture_end();
+       ci != ce; ++ci, ++ri, ++pi) {
+
+    if (ci->capturesVariableArrayType()) {
+      VLASizeFields.push_back(*ri);
+      continue;
+    }
+
+    if (ci->capturesThis()) {
+      auto *ThisRef = LocalDeclMap[*pi];
+      LValue Addr = LValue::MakeAddr(ThisRef, ri->getType(), CharUnits(),
+                                     CGM.getContext());
+      CXXThisValue = EmitLoadOfLValue(Addr, CS.getLocStart()).getScalarVal();
+      ;
+      continue;
+    }
+
+    // Find the expressions that give the VLA sizes and update the VLASizeMap.
+    if (!VLASizeFields.empty()) {
+      assert(ci->capturesVariable() &&
+             "Expecting declaration that has VLA type!");
+
+      const PointerType *PT = cast<PointerType>((*pi)->getType().getTypePtr());
+      QualType CurTy = PT->getPointeeType();
+
+      for (auto *FD : VLASizeFields) {
+        const VariableArrayType *VAT =
+            cast<VariableArrayType>(CurTy.getTypePtr());
+        auto *V = VLASizeMap[VAT->getSizeExpr()];
+        assert(V && "VLA Size value must exist!");
+        VLASizeMap[FD->getCapturedVLAType()->getSizeExpr()] = V;
+        CurTy = VAT->getElementType();
+      }
+
+      VLASizeFields.clear();
+    }
+  }
+}
+
+void CodeGenFunction::EmitOMPTargetDirective(const OMPTargetDirective &S) {
+  LexicalScope Scope(*this, S.getSourceRange());
+
+  // Emit target region as a standalone region.
+  auto &&CodeGen = [&S](CodeGenFunction &CGF) {
+    CGF.PrepareOMPTargetDirectiveBodyEmission(S);
+    CGF.EmitStmt(cast<CapturedStmt>(S.getAssociatedStmt())->getCapturedStmt());
+  };
+
+  // Obtain the target region outlined function.
+  llvm::Value *Fn =
+      CGM.getOpenMPRuntime().emitTargetOutlinedFunction(*this, S, CodeGen);
+
+  // Check if we have any if clause associated with the directive.
+  const Expr *IfCond = nullptr;
+  if (auto C = S.getSingleClause(OMPC_if)) {
+    IfCond = cast<OMPIfClause>(C)->getCondition();
+  }
+
+  // Check if we have any device clause associated with the directive.
+  const Expr *Device = nullptr;
+  if (auto C = S.getSingleClause(OMPC_device)) {
+    Device = cast<OMPDeviceClause>(C)->getDevice();
+  }
+
+  CGM.getOpenMPRuntime().emitTargetCall(*this, S, Fn, IfCond, Device);
 }
 
 void CodeGenFunction::EmitOMPTeamsDirective(const OMPTeamsDirective &) {
   llvm_unreachable("CodeGen for 'omp teams' is not supported yet.");
 }
 
 void CodeGenFunction::EmitOMPCancellationPointDirective(
     const OMPCancellationPointDirective &S) {
   CGM.getOpenMPRuntime().emitCancellationPointCall(*this, S.getLocStart(),
                                                    S.getCancelRegion());
 }
 
 void CodeGenFunction::EmitOMPCancelDirective(const OMPCancelDirective &S) {
   CGM.getOpenMPRuntime().emitCancelCall(*this, S.getLocStart(),
                                         S.getCancelRegion());
 }
 
 CodeGenFunction::JumpDest
 CodeGenFunction::getOMPCancelDestination(OpenMPDirectiveKind Kind) {
   if (Kind == OMPD_parallel || Kind == OMPD_task)
     return ReturnBlock;
   else if (Kind == OMPD_for || Kind == OMPD_section || Kind == OMPD_sections)
     return BreakContinueStack.empty() ? JumpDest()
                                       : BreakContinueStack.back().BreakBlock;
   return JumpDest();
 }
 
 // Generate the instructions for '#pragma omp target data' directive.

lib/CodeGen/CodeGenFunction.h

public:
 
     CapturedRegionKind getKind() const { return Kind; }
 
     virtual void setContextValue(llvm::Value *V) { ThisValue = V; }
     // \brief Retrieve the value of the context parameter.
     virtual llvm::Value *getContextValue() const { return ThisValue; }
 
     /// \brief Lookup the captured field decl for a variable.
-    virtual const FieldDecl *lookup(const VarDecl *VD) const {
+    virtual const Decl *lookup(const VarDecl *VD) const {
       return CaptureFields.lookup(VD);
     }
 
-    bool isCXXThisExprCaptured() const { return getThisFieldDecl() != nullptr; }
-    virtual FieldDecl *getThisFieldDecl() const { return CXXThisFieldDecl; }
+    bool isCXXThisExprCaptured() const { return getThisDecl() != nullptr; }
+    virtual Decl *getThisDecl() const { return CXXThisFieldDecl; }
 
     static bool classof(const CGCapturedStmtInfo *) {
       return true;
     }
 
     /// \brief Emit the captured statement body.
     virtual void EmitBody(CodeGenFunction &CGF, const Stmt *S) {
       CGF.incrementProfileCounter(S);
#endif
 
   /// getVLASize - Returns an LLVM value that corresponds to the size,
   /// in non-variably-sized elements, of a variable length array type,
   /// plus that largest non-variably-sized element type.  Assumes that
   /// the type has already been emitted with EmitVariablyModifiedType.
   std::pair<llvm::Value*,QualType> getVLASize(const VariableArrayType *vla);
   std::pair<llvm::Value*,QualType> getVLASize(QualType vla);
 
+  /// getVLASizeMap - Returns an LLVM value that corresponds to the size of the
+  /// VLA size of the type \a type. Assumes that the type has already been
+  /// emitted with EmitVariablyModifiedType.
+  llvm::Value *getVLASizeMap(const VariableArrayType *vla);
+
   /// LoadCXXThis - Load the value of 'this'. This function is only valid while
   /// generating code for an C++ member function.
   llvm::Value *LoadCXXThis() {
     assert(CXXThisValue && "no 'this' value for this function");
     return CXXThisValue;
   }
 
   /// LoadCXXVTT - Load the VTT parameter to base constructors/destructors have
#endif
   void EmitOMPTaskDirective(const OMPTaskDirective &S);
   void EmitOMPTaskyieldDirective(const OMPTaskyieldDirective &S);
   void EmitOMPBarrierDirective(const OMPBarrierDirective &S);
   void EmitOMPTaskwaitDirective(const OMPTaskwaitDirective &S);
   void EmitOMPTaskgroupDirective(const OMPTaskgroupDirective &S);
   void EmitOMPFlushDirective(const OMPFlushDirective &S);
   void EmitOMPOrderedDirective(const OMPOrderedDirective &S);
   void EmitOMPAtomicDirective(const OMPAtomicDirective &S);
+  typedef SmallVector<llvm::Value *, 4> OMPTargetDirectiveVLASizes;
+  void PrepareOMPTargetDirectiveBodyEmission(const OMPTargetDirective &S);
   void EmitOMPTargetDirective(const OMPTargetDirective &S);
   void EmitOMPTargetDataDirective(const OMPTargetDataDirective &S);
   void EmitOMPTeamsDirective(const OMPTeamsDirective &S);
   void
   EmitOMPCancellationPointDirective(const OMPCancellationPointDirective &S);
   void EmitOMPCancelDirective(const OMPCancelDirective &S);
 
   /// \brief Emit inner loop of the worksharing/simd construct.

lib/CodeGen/CodeGenFunction.cpp

if (MD->getParent()->isLambda() &&
       // fast register allocator would be happier...
       CXXThisValue = CXXABIThisValue;
     }
   }
 
   // If any of the arguments have a variably modified type, make sure to
   // emit the type size.
   for (FunctionArgList::const_iterator i = Args.begin(), e = Args.end();
        i != e; ++i) {
     const VarDecl *VD = *i;
 
     // Dig out the type as written from ParmVarDecls; it's unclear whether
     // the standard (C99 6.9.1p10) requires this, but we're following the
     // precedent set by gcc.
     QualType Ty;
     if (const ParmVarDecl *PVD = dyn_cast<ParmVarDecl>(VD))
       Ty = PVD->getOriginalType();
if (!numElements) {
       // FIXME: Teach -fsanitize=undefined to trap this.
       numElements = Builder.CreateNUWMul(numElements, vlaSize);
     }
   } while ((type = getContext().getAsVariableArrayType(elementType)));
 
   return std::pair<llvm::Value*,QualType>(numElements, elementType);
 }
 
+llvm::Value *CodeGenFunction::getVLASizeMap(const VariableArrayType *type) {
+  llvm::Value *vlaSize = VLASizeMap[type->getSizeExpr()];
+  assert(vlaSize && "No vla size matching the requested expression!");
+  return vlaSize;
+}
+
 void CodeGenFunction::EmitVariablyModifiedType(QualType type) {
   assert(type->isVariablyModifiedType() &&
          "Must pass variably modified type to EmitVLASizes!");
 
   EnsureInsertPoint();
 
   // We're going to walk down into the type and look for VLA
   // expressions.

lib/Sema/SemaOpenMP.cpp

void Sema::ActOnOpenMPRegionStart(OpenMPDirectiveKind DKind, Scope *CurScope) {
   case OMPD_atomic: {
     Sema::CapturedParamNameType Params[] = {
         std::make_pair(StringRef(), QualType()) // __context with shared vars
     };
     ActOnCapturedRegionStart(DSAStack->getConstructLoc(), CurScope, CR_OpenMP,
                              Params);
     break;
   }
-  case OMPD_target_data:
-  case OMPD_target: {
+  case OMPD_target_data: {
     Sema::CapturedParamNameType Params[] = {
         std::make_pair(StringRef(), QualType()) // __context with shared vars
     };
     ActOnCapturedRegionStart(DSAStack->getConstructLoc(), CurScope, CR_OpenMP,
                              Params);
     break;
   }
+  case OMPD_target: {
+    Sema::CapturedParamNameType Params[] = {
+        std::make_pair(StringRef(), QualType()) // __context with shared vars
+    };
+    ActOnCapturedRegionStart(DSAStack->getConstructLoc(), CurScope,
+                             CR_OpenMP_Target, Params);
+    break;
+  }
   case OMPD_teams: {
     QualType KmpInt32Ty = Context.getIntTypeForBitwidth(32, 1);
     QualType KmpInt32PtrTy = Context.getPointerType(KmpInt32Ty);
     Sema::CapturedParamNameType Params[] = {
         std::make_pair(".global_tid.", KmpInt32PtrTy),
         std::make_pair(".bound_tid.", KmpInt32PtrTy),
         std::make_pair(StringRef(), QualType()) // __context with shared vars
     };
if (isOpenMPPrivate(Clause->getClauseKind()) ||
       // Required for proper codegen of combined directives.
       // TODO: add processing for other clauses.
       if (auto *E = cast_or_null<Expr>(
               cast<OMPScheduleClause>(Clause)->getHelperChunkSize())) {
           MarkDeclarationsReferencedInExpr(E);
         }
     }
   }
-  return ActOnCapturedRegionEnd(S.get());
+
+  CapturedStmt *Res = cast<CapturedStmt>(ActOnCapturedRegionEnd(S.get()).get());
+
+  // If this is an offloading captured region, we need change the captured
+  // declaration so that it uses as parameters each of the captured
+  // declarations.
+  if (Res->getCapturedRegionKind() == CR_OpenMP_Target) {
+    CapturedDecl *OldCD = Res->getCapturedDecl();
+    DeclContext *DC = OldCD->getDeclContext();
+
+    unsigned NumParams = Res->capture_size();
+    const RecordDecl *RD = Res->getCapturedRecordDecl();
+    CapturedDecl *NewCD = CapturedDecl::Create(Context, DC, NumParams);
+    DeclContext *NewDC = CapturedDecl::castToDeclContext(NewCD);
+    SourceLocation Loc = Res->getLocStart();
+
+    SmallVector<const FieldDecl *, 3> VLASizeFields;
+    unsigned i = 0;
+    auto ri = RD->field_begin();
+    for (CapturedStmt::const_capture_iterator ci = Res->capture_begin(),
+                                              ce = Res->capture_end();
+         ci != ce; ++ci, ++ri, ++i) {
+      StringRef Name;
+      QualType Ty;
+      if (ci->capturesVariableArrayType()) {
+        VLASizeFields.push_back(*ri);
+        Ty = Context.getPointerType(ri->getType());
+        Name = "__vla_size";
+      } else if (ci->capturesThis()) {
+        Ty = ri->getType();
+        Name = "__this";
+      } else {
+        const VarDecl *VD = ci->getCapturedVar();
+
+        Ty = VD->getType();
+
+        // If this declaration has an variable-length array type, we need to
+        // adapt the size expression to use the implicit parameter that
+        // captures it. We also need to make sure we adapt all the sizes if
+        // the type has several dimensions.
+        if (!VLASizeFields.empty()) {
+          assert(Ty.getTypePtr() ==
+                     VLASizeFields.front()->getCapturedVLAType() &&
+                 "Invalid VLA information!");
+
+          Ty = VLASizeFields.back()->getCapturedVLAType()->getElementType();
+          for (unsigned t = 1; !VLASizeFields.empty();
+               VLASizeFields.pop_back(), ++t) {
+            auto *OriginalVAT = VLASizeFields.back()->getCapturedVLAType();
+            auto SizeType = VLASizeFields.back()->getType();
+            auto *SizeDecl = NewCD->getParam(i - t);
+            auto *NewSizeRefExpr = buildDeclRefExpr(
+                *this, SizeDecl, SizeDecl->getType(), Loc, true);
+            auto *NewSizeValExpr =
+                new (Context) UnaryOperator(NewSizeRefExpr, UO_Deref, SizeType,
+                                            VK_RValue, OK_Ordinary, Loc);
+            Ty = Context.getVariableArrayType(
+                Ty, NewSizeValExpr, OriginalVAT->getSizeModifier(),
+                OriginalVAT->getIndexTypeCVRQualifiers(),
+                OriginalVAT->getBracketsRange());
+          }
+        }
+        Ty = Context.getPointerType(Ty);
+        Name = VD->getName();
+      }
+
+      IdentifierInfo *ParamName = &Context.Idents.get(Name);
+      ImplicitParamDecl *Param =
+          ImplicitParamDecl::Create(Context, NewDC, Loc, ParamName, Ty);
+      NewCD->setParam(i, Param);
+    }
+
+    NewCD->setBody(Res->getCapturedStmt());
+    Res->setCapturedDecl(NewCD);
+    DC->addDecl(NewCD);
+    DC->removeDecl(OldCD);
+  }
+
+  return Res;
 }
 
 static bool CheckNestingOfRegions(Sema &SemaRef, DSAStackTy *Stack,
                                   OpenMPDirectiveKind CurrentRegion,
                                   const DeclarationNameInfo &CurrentName,
                                   OpenMPDirectiveKind CancelRegion,
                                   SourceLocation StartLoc) {
   // Allowed nesting of constructs

test/OpenMP/target_codegen.cpp

+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
+// expected-no-diagnostics
+// REQUIRES: powerpc-registered-target
+#ifndef HEADER
+#define HEADER
+
+// CHECK-DAG: [[TT:%.+]] = type { i64, i8 }
+// CHECK-DAG: [[S1:%.+]] = type { double }
+
+// We have 8 target regions, but only 7 that actually will generate offloading
+// code, and only 6 will have mapped arguments.
+
+// CHECK-DAG: [[MAPT2:@.+]] = private constant [1 x i32] [i32 3]
+// CHECK-DAG: [[MAPT3:@.+]] = private constant [2 x i32] [i32 3, i32 3]
+// CHECK-DAG: [[MAPT4:@.+]] = private constant [9 x i32] [i32 3, i32 3, i32 1, i32 3, i32 3, i32 1, i32 1, i32 3, i32 3]
+// CHECK-DAG: [[MAPT5:@.+]] = private constant [3 x i32] [i32 3, i32 3, i32 3]
+// CHECK-DAG: [[MAPT6:@.+]] = private constant [4 x i32] [i32 3, i32 3, i32 3, i32 3]
+// CHECK-DAG: [[MAPT7:@.+]] = private constant [5 x i32] [i32 3, i32 3, i32 1, i32 1, i32 3]
+// CHECK-DAG: @{{.*}} = private constant i8 0
+// CHECK-DAG: @{{.*}} = private constant i8 0
+// CHECK-DAG: @{{.*}} = private constant i8 0
+// CHECK-DAG: @{{.*}} = private constant i8 0
+// CHECK-DAG: @{{.*}} = private constant i8 0
+// CHECK-DAG: @{{.*}} = private constant i8 0
+// CHECK-DAG: @{{.*}} = private constant i8 0
+
+template<typename tx, typename ty>
+struct TT{
+  tx X;
+  ty Y;
+};
+
+// CHECK: define {{.*}}[[FOO:@.+]](
+int foo(int n) {
+  int a = 0;
+  short aa = 0;
+  float b[10];
+  float bn[n];
+  double c[5][10];
+  double cn[5][n];
+  TT<long, char> d;
+
+  // CHECK:       br label %[[TRY:[^,]+]]
+  // CHECK:       [[TRY]]
+  // CHECK:       [[RET:%.+]] = call i32 @__tgt_target(i32 -1, i8* @{{[^,]+}}, i32 0, i8** null, i8** null, i64* null, i32* null)
+  // CHECK-NEXT:  [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
+  // CHECK-NEXT:  br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
+  // CHECK:       [[FAIL]]
+  // CHECK:       call void [[HVT0:@.+]]()
+  // CHECK-NEXT:  br label %[[END]]
+  // CHECK:       [[END]]
+  #pragma omp target
+  {
+  }
+
+  // CHECK:       call void [[HVT1:@.+]](i32* {{[^,]+}})
+  #pragma omp target if(0)
+  {
+    a += 1;
+  }
+
+  // CHECK:       br label %[[TRY:[^,]+]]
+  // CHECK:       [[TRY]]
+  // CHECK-DAG:   [[RET:%.+]] = call i32 @__tgt_target(i32 -1, i8* @{{[^,]+}}, i32 1, i8** [[BP:%[^,]+]], i8** [[P:%[^,]+]], i64* [[S:%[^,]+]], i32* getelementptr inbounds ([1 x i32], [1 x i32]* [[MAPT2]], i32 0, i32 0))
+
+  // CHECK-DAG:   store i64 4, i64* [[SADDR0:%.+]]
+  // CHECK-DAG:   [[SADDR0]] = getelementptr inbounds i64, i64* [[S]], i32 [[IDX0:[0-9]+]]
+  // CHECK-DAG:   [[BPADDR0:%.+]] = getelementptr inbounds i8*, i8** [[BP]], i32 [[IDX0]]
+  // CHECK-DAG:   [[PADDR0:%.+]] = getelementptr inbounds i8*, i8** [[P]], i32 [[IDX0]]
+  // CHECK-DAG:   store i8* [[BP0:%[^,]+]], i8** [[BPADDR0]]
+  // CHECK-DAG:   store i8* [[P0:%[^,]+]], i8** [[PADDR0]]
+  // CHECK-DAG:   [[BP0]] = bitcast i32* %{{.+}} to i8*
+  // CHECK-DAG:   [[P0]] = bitcast i32* %{{.+}} to i8*
+
+  // CHECK:       [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
+  // CHECK-NEXT:  br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
+  // CHECK:       [[FAIL]]
+  // CHECK:       call void [[HVT2:@.+]](i32* {{[^,]+}})
+  // CHECK-NEXT:  br label %[[END]]
+  // CHECK:       [[END]]
+  #pragma omp target if(1)
+  {
+    a += 1;
+  }
+
+  // CHECK:       [[IF:%.+]] = icmp sgt i32 {{[^,]+}}, 10
+  // CHECK:       br i1 [[IF]], label %[[TRY:[^,]+]], label %[[FAIL:[^,]+]]
+  // CHECK:       [[TRY]]
+  // CHECK-DAG:   [[RET:%.+]] = call i32 @__tgt_target(i32 -1, i8* @{{[^,]+}}, i32 2, i8** [[BP:%[^,]+]], i8** [[P:%[^,]+]], i64* [[S:%[^,]+]], i32* getelementptr inbounds ([2 x i32], [2 x i32]* [[MAPT3]], i32 0, i32 0))
+
+  // CHECK-DAG:   store i64 4, i64* [[SADDR0:%.+]]
+  // CHECK-DAG:   [[SADDR0]] = getelementptr inbounds i64, i64* [[S]], i32 [[IDX0:[0-9]+]]
+  // CHECK-DAG:   [[BPADDR0:%.+]] = getelementptr inbounds i8*, i8** [[BP]], i32 [[IDX0]]
+  // CHECK-DAG:   [[PADDR0:%.+]] = getelementptr inbounds i8*, i8** [[P]], i32 [[IDX0]]
+  // CHECK-DAG:   store i8* [[BP0:%[^,]+]], i8** [[BPADDR0]]
+  // CHECK-DAG:   store i8* [[P0:%[^,]+]], i8** [[PADDR0]]
+  // CHECK-DAG:   [[BP0]] = bitcast i32* %{{.+}} to i8*
+  // CHECK-DAG:   [[P0]] = bitcast i32* %{{.+}} to i8*
+
+  // CHECK-DAG:   store i64 2, i64* [[SADDR1:%.+]]
+  // CHECK-DAG:   [[SADDR1]] = getelementptr inbounds i64, i64* [[S]], i32 [[IDX1:[0-9]+]]
+  // CHECK-DAG:   [[BPADDR1:%.+]] = getelementptr inbounds i8*, i8** [[BP]], i32 [[IDX1]]
+  // CHECK-DAG:   [[PADDR1:%.+]] = getelementptr inbounds i8*, i8** [[P]], i32 [[IDX1]]
+  // CHECK-DAG:   store i8* [[BP1:%[^,]+]], i8** [[BPADDR1]]
+  // CHECK-DAG:   store i8* [[P1:%[^,]+]], i8** [[PADDR1]]
+  // CHECK-DAG:   [[BP1]] = bitcast i16* %{{.+}} to i8*
+  // CHECK-DAG:   [[P1]] = bitcast i16* %{{.+}} to i8*
+
+  // CHECK:       [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
+  // CHECK-NEXT:  br i1 [[ERROR]], label %[[FAIL]], label %[[END:[^,]+]]
+  // CHECK:       [[FAIL]]
+  // CHECK:       call void [[HVT3:@.+]]({{[^,]+}}, {{[^,]+}})
+  // CHECK-NEXT:  br label %[[END]]
+  // CHECK:       [[END]]
+  #pragma omp target if(n>10)
+  {
+    a += 1;
+    aa += 1;
+  }
+
+  // We capture 3 VLA sizes in this target region
+  // CHECK-DAG:   store i64 %{{[^,]+}}, i64* [[VLA0:%[^,]+]]
+  // CHECK-DAG:   store i64 %{{[^,]+}}, i64* [[VLA1:%[^,]+]]
+  // CHECK-DAG:   store i64 %{{[^,]+}}, i64* [[VLA2:%[^,]+]]
+  // CHECK:       [[IF:%.+]] = icmp sgt i32 {{[^,]+}}, 20
+  // CHECK:       br i1 [[IF]], label %[[TRY:[^,]+]], label %[[FAIL:[^,]+]]
+  // CHECK:       [[TRY]]
+  // CHECK-DAG:   [[RET:%.+]] = call i32 @__tgt_target(i32 -1, i8* @{{[^,]+}}, i32 9, i8** [[BP:%[^,]+]], i8** [[P:%[^,]+]], i64* [[S:%[^,]+]], i32* getelementptr inbounds ([9 x i32], [9 x i32]* [[MAPT4]], i32 0, i32 0))
+
+  // CHECK-DAG:   [[SADDR0:%.+]] = getelementptr inbounds i64, i64* [[S]], i32 [[IDX0:[0-9]+]]
+  // CHECK-DAG:   [[BPADDR0:%.+]] = getelementptr inbounds i8*, i8** [[BP]], i32 [[IDX0]]
+  // CHECK-DAG:   [[PADDR0:%.+]] = getelementptr inbounds i8*, i8** [[P]], i32 [[IDX0]]
+  // CHECK-DAG:   [[SADDR1:%.+]] = getelementptr inbounds i64, i64* [[S]], i32 [[IDX1:[0-9]+]]
+  // CHECK-DAG:   [[BPADDR1:%.+]] = getelementptr inbounds i8*, i8** [[BP]], i32 [[IDX1]]
+  // CHECK-DAG:   [[PADDR1:%.+]] = getelementptr inbounds i8*, i8** [[P]], i32 [[IDX1]]
+  // CHECK-DAG:   [[SADDR2:%.+]] = getelementptr inbounds i64, i64* [[S]], i32 [[IDX2:[0-9]+]]
+  // CHECK-DAG:   [[BPADDR2:%.+]] = getelementptr inbounds i8*, i8** [[BP]], i32 [[IDX2]]
+  // CHECK-DAG:   [[PADDR2:%.+]] = getelementptr inbounds i8*, i8** [[P]], i32 [[IDX2]]
+  // CHECK-DAG:   [[SADDR3:%.+]] = getelementptr inbounds i64, i64* [[S]], i32 [[IDX3:[0-9]+]]
+  // CHECK-DAG:   [[BPADDR3:%.+]] = getelementptr inbounds i8*, i8** [[BP]], i32 [[IDX3]]
+  // CHECK-DAG:   [[PADDR3:%.+]] = getelementptr inbounds i8*, i8** [[P]], i32 [[IDX3]]
+  // CHECK-DAG:   [[SADDR4:%.+]] = getelementptr inbounds i64, i64* [[S]], i32 [[IDX4:[0-9]+]]
+  // CHECK-DAG:   [[BPADDR4:%.+]] = getelementptr inbounds i8*, i8** [[BP]], i32 [[IDX4]]
+  // CHECK-DAG:   [[PADDR4:%.+]] = getelementptr inbounds i8*, i8** [[P]], i32 [[IDX4]]
+  // CHECK-DAG:   [[SADDR5:%.+]] = getelementptr inbounds i64, i64* [[S]], i32 [[IDX5:[0-9]+]]
+  // CHECK-DAG:   [[BPADDR5:%.+]] = getelementptr inbounds i8*, i8** [[BP]], i32 [[IDX5]]
+  // CHECK-DAG:   [[PADDR5:%.+]] = getelementptr inbounds i8*, i8** [[P]], i32 [[IDX5]]
+  // CHECK-DAG:   [[SADDR6:%.+]] = getelementptr inbounds i64, i64* [[S]], i32 [[IDX6:[0-9]+]]
+  // CHECK-DAG:   [[BPADDR6:%.+]] = getelementptr inbounds i8*, i8** [[BP]], i32 [[IDX6]]
+  // CHECK-DAG:   [[PADDR6:%.+]] = getelementptr inbounds i8*, i8** [[P]], i32 [[IDX6]]
+  // CHECK-DAG:   [[SADDR7:%.+]] = getelementptr inbounds i64, i64* [[S]], i32 [[IDX7:[0-9]+]]
+  // CHECK-DAG:   [[BPADDR7:%.+]] = getelementptr inbounds i8*, i8** [[BP]], i32 [[IDX7]]
+  // CHECK-DAG:   [[PADDR7:%.+]] = getelementptr inbounds i8*, i8** [[P]], i32 [[IDX7]]
+  // CHECK-DAG:   [[SADDR8:%.+]] = getelementptr inbounds i64, i64* [[S]], i32 [[IDX8:[0-9]+]]
+  // CHECK-DAG:   [[BPADDR8:%.+]] = getelementptr inbounds i8*, i8** [[BP]], i32 [[IDX8]]
+  // CHECK-DAG:   [[PADDR8:%.+]] = getelementptr inbounds i8*, i8** [[P]], i32 [[IDX8]]
+
+  // The names below are not necessarily consistent with the names used for the
+  // addresses above as some are repeated.
+  // CHECK-DAG:   [[BP0:%[^,]+]] = bitcast i64* [[VLA0]] to i8*
+  // CHECK-DAG:   [[P0:%[^,]+]] = bitcast i64* [[VLA0]] to i8*
+  // CHECK-DAG:   store i8* [[BP0]], i8** {{%[^,]+}}
+  // CHECK-DAG:   store i8* [[P0]], i8** {{%[^,]+}}
+  // CHECK-DAG:   store i64 8, i64* {{%[^,]+}}
+
+  // CHECK-DAG:   [[BP1:%[^,]+]] = bitcast i64* [[VLA1]] to i8*
+  // CHECK-DAG:   [[P1:%[^,]+]] = bitcast i64* [[VLA1]] to i8*
+  // CHECK-DAG:   store i8* [[BP1]], i8** {{%[^,]+}}
+  // CHECK-DAG:   store i8* [[P1]], i8** {{%[^,]+}}
+  // CHECK-DAG:   store i64 8, i64* {{%[^,]+}}
+
+  // CHECK-DAG:   [[BP2:%[^,]+]] = bitcast i64* [[VLA2]] to i8*
+  // CHECK-DAG:   [[P2:%[^,]+]] = bitcast i64* [[VLA2]] to i8*
+  // CHECK-DAG:   store i8* [[BP2]], i8** {{%[^,]+}}
+  // CHECK-DAG:   store i8* [[P2]], i8** {{%[^,]+}}
+  // CHECK-DAG:   store i64 8, i64* {{%[^,]+}}
+
+  // CHECK-DAG:   [[BP3:%[^,]+]] = bitcast i32* %{{.+}} to i8*
+  // CHECK-DAG:   [[P3:%[^,]+]] = bitcast i32* %{{.+}} to i8*
+  // CHECK-DAG:   store i8* [[BP3]], i8** {{%[^,]+}}
+  // CHECK-DAG:   store i8* [[P3]], i8** {{%[^,]+}}
+  // CHECK-DAG:   store i64 4, i64* {{%[^,]+}}
+
+  // CHECK-DAG:   [[BP4:%[^,]+]] = bitcast [10 x float]* %{{.+}} to i8*
+  // CHECK-DAG:   [[P4:%[^,]+]] = bitcast [10 x float]* %{{.+}} to i8*
+  // CHECK-DAG:   store i8* [[BP4]], i8** {{%[^,]+}}
+  // CHECK-DAG:   store i8* [[P4]], i8** {{%[^,]+}}
+  // CHECK-DAG:   store i64 40, i64* {{%[^,]+}}
+
+  // CHECK-DAG:   [[BP5:%[^,]+]] = bitcast float* %{{.+}} to i8*
+  // CHECK-DAG:   [[P5:%[^,]+]] = bitcast float* %{{.+}} to i8*
+  // CHECK-DAG:   store i8* [[BP5]], i8** {{%[^,]+}}
+  // CHECK-DAG:   store i8* [[P5]], i8** {{%[^,]+}}
+  // CHECK-DAG:   store i64 4, i64* {{%[^,]+}}
+
+  // CHECK-DAG:   [[BP6:%[^,]+]] = bitcast [5 x [10 x double]]* %{{.+}} to i8*
+  // CHECK-DAG:   [[P6:%[^,]+]] = bitcast [5 x [10 x double]]* %{{.+}} to i8*
+  // CHECK-DAG:   store i8* [[BP6]], i8** {{%[^,]+}}
+  // CHECK-DAG:   store i8* [[P6]], i8** {{%[^,]+}}
+  // CHECK-DAG:   store i64 400, i64* {{%[^,]+}}
+
+  // CHECK-DAG:   [[BP7:%[^,]+]] = bitcast double* %{{.+}} to i8*
+  // CHECK-DAG:   [[P7:%[^,]+]] = bitcast double* %{{.+}} to i8*
+  // CHECK-DAG:   store i8* [[BP7]], i8** {{%[^,]+}}
+  // CHECK-DAG:   store i8* [[P7]], i8** {{%[^,]+}}
+  // CHECK-DAG:   store i64 8, i64* {{%[^,]+}}
+
+  // CHECK-DAG:   [[BP8:%[^,]+]] = bitcast [[TT]]* %{{.+}} to i8*
+  // CHECK-DAG:   [[P8:%[^,]+]] = bitcast [[TT]]* %{{.+}} to i8*
+  // CHECK-DAG:   store i8* [[BP8]], i8** {{%[^,]+}}
+  // CHECK-DAG:   store i8* [[P8]], i8** {{%[^,]+}}
+  // CHECK-DAG:   store i64 16, i64* {{%[^,]+}}
+
+  // CHECK:       [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
+  // CHECK-NEXT:  br i1 [[ERROR]], label %[[FAIL]], label %[[END:[^,]+]]
+  // CHECK:       [[FAIL]]
+  // CHECK:       call void [[HVT4:@.+]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}})
+  // CHECK-NEXT:  br label %[[END]]
+  // CHECK:       [[END]]
+  #pragma omp target if(n>20)
+  {
+    a += 1;
+    b[2] += 1.0;
+    bn[3] += 1.0;
+    c[1][2] += 1.0;
+    cn[1][3] += 1.0;
+    d.X += 1;
+    d.Y += 1;
+  }
+
+  return a;
+}
+
+// Check that the offloading functions are emitted and that the arguments are
+// correct and loaded correctly for the target regions in foo().
+
+// CHECK:       define internal void [[HVT0]]
+
+// CHECK:       define internal void [[HVT1]]
+// CHECK-DAG:   [[LOCALX_A:%.+]] = alloca i32*
+// CHECK-DAG:   store i32* [[ARG_A:%.+]], i32** [[LOCALX_A]]
+// CHECK-DAG:   [[USE_A:%.+]] = load i32*, i32** [[LOCALX_A:%.+]]
+// CHECK-DAG:   load i32, i32* [[USE_A]]
+
+// CHECK:       define internal void [[HVT2]]
+// CHECK-DAG:   [[LOCALX_A:%.+]] = alloca i32*
+// CHECK-DAG:   store i32* [[ARG_A:%.+]], i32** [[LOCALX_A]]
+// CHECK-DAG:   [[USE_A:%.+]] = load i32*, i32** [[LOCALX_A:%.+]]
+// CHECK-DAG:   load i32, i32* [[USE_A]]
+
+// CHECK:       define internal void [[HVT3]]
+// CHECK-DAG:   [[LOCALX_A:%.+]] = alloca i32*
+// CHECK-DAG:   [[LOCALX_AA:%.+]] = alloca i16*
+// CHECK-DAG:   store i32* [[ARG_A:%.+]], i32** [[LOCALX_A]]
+// CHECK-DAG:   store i16* [[ARG_AA:%.+]], i16** [[LOCALX_AA]]
+// CHECK-DAG:   [[USE_A:%.+]] = load i32*, i32** [[LOCALX_A:%.+]]
+// CHECK-DAG:   [[USE_AA:%.+]] = load i16*, i16** [[LOCALX_AA:%.+]]
+// CHECK-DAG:   load i32, i32* [[USE_A]]
+// CHECK-DAG:   load i16, i16* [[USE_AA]]
+
+// CHECK:       define internal void [[HVT4]]
+// CHECK-DAG:   [[LOCALX_A:%.+]] = alloca i32*
+// CHECK-DAG:   [[LOCALX_B:%.+]] = alloca [10 x float]*
+// CHECK-DAG:   [[LOCALX_BN:%.+]] = alloca float*
+// CHECK-DAG:   [[LOCALX_C:%.+]] = alloca [5 x [10 x double]]*
+// CHECK-DAG:   [[LOCALX_CN:%.+]] = alloca double*
+// CHECK-DAG:   [[LOCALX_D:%.+]] = alloca [[TT]]*
+// CHECK-DAG:   [[LOCALX_VLA1:%.+]] = alloca i64*
+// CHECK-DAG:   [[LOCALX_VLA2:%.+]] = alloca i64*
+// CHECK-DAG:   [[LOCALX_VLA3:%.+]] = alloca i64*
+// CHECK-DAG:   store i32* [[ARG_A:%.+]], i32** [[LOCALX_A]]
+// CHECK-DAG:   store [10 x float]* [[ARG_B:%.+]], [10 x float]** [[LOCALX_B]]
+// CHECK-DAG:   store float* [[ARG_BN:%.+]], float** [[LOCALX_BN]]
+// CHECK-DAG:   store [5 x [10 x double]]* [[ARG_C:%.+]], [5 x [10 x double]]** [[LOCALX_C]]
+// CHECK-DAG:   store double* [[ARG_CN:%.+]], double** [[LOCALX_CN]]
+// CHECK-DAG:   store [[TT]]* [[ARG_D:%.+]], [[TT]]** [[LOCALX_D]]
+// CHECK-DAG:   store i64* [[ARG_VLA1:%.+]], i64** [[LOCALX_VLA1]]
+// CHECK-DAG:   store i64* [[ARG_VLA2:%.+]], i64** [[LOCALX_VLA2]]
+// CHECK-DAG:   store i64* [[ARG_VLA3:%.+]], i64** [[LOCALX_VLA3]]
+// CHECK-DAG:   [[USE_A:%.+]] = load i32*, i32** [[LOCALX_A:%.+]]
+// CHECK-DAG:   [[USE_B:%.+]] = load [10 x float]*, [10 x float]** [[LOCALX_B:%.+]]
+// CHECK-DAG:   [[USE_BN:%.+]] = load float*, float** [[LOCALX_BN:%.+]]
+// CHECK-DAG:   [[USE_C:%.+]] = load [5 x [10 x double]]*, [5 x [10 x double]]** [[LOCALX_C:%.+]]
+// CHECK-DAG:   [[USE_CN:%.+]] = load double*, double** [[LOCALX_CN:%.+]]
+// CHECK-DAG:   [[USE_D:%.+]] = load [[TT]]*, [[TT]]** [[LOCALX_D:%.+]]
+// CHECK-DAG:   [[USE_VLA1:%.+]] = load i64*, i64** [[LOCALX_VLA1:%.+]]
+// CHECK-DAG:   [[USE_VLA2:%.+]] = load i64*, i64** [[LOCALX_VLA2:%.+]]
+// CHECK-DAG:   [[USE_VLA3:%.+]] = load i64*, i64** [[LOCALX_VLA3:%.+]]
+// CHECK-DAG:   load i32, i32* [[USE_A]]
+// CHECK-DAG:   getelementptr inbounds [10 x float], [10 x float]* [[USE_B]], i{{.*}} 0, i{{.*}} 2
+// CHECK-DAG:   getelementptr inbounds float, float* [[USE_BN]], i{{.*}} 3
+// CHECK-DAG:   getelementptr inbounds [5 x [10 x double]], [5 x [10 x double]]* [[USE_C]], i{{.*}} 0, i{{.*}} 1
+// CHECK-DAG:   [[VLAMUL:%.+]] = mul {{.*}}i64 1, %{{.*}}
+// CHECK-DAG:   getelementptr inbounds double, double* [[USE_CN]], i{{.*}} [[VLAMUL]]
+// CHECK-DAG:   load i64, i64* [[USE_VLA1]]
+// CHECK-DAG:   load i64, i64* [[USE_VLA2]]
+// CHECK-DAG:   load i64, i64* [[USE_VLA3]]
+
+template<typename tx>
+tx ftemplate(int n) {
+  tx a = 0;
+  short aa = 0;
+  tx b[10];
+
+  #pragma omp target if(n>40)
+  {
+    a += 1;
+    aa += 1;
+    b[2] += 1;
+  }
+
+  return a;
+}
+
+static
+int fstatic(int n) {
+  int a = 0;
+  short aa = 0;
+  char aaa = 0;
+  int b[10];
+
+  #pragma omp target if(n>50)
+  {
+    a += 1;
+    aa += 1;
+    aaa += 1;
+    b[2] += 1;
+  }
+
+  return a;
+}
+
+struct S1 {
+  double a;
+
+  int r1(int n){
+    int b = n+1;
+    short int c[2][n];
+
+    #pragma omp target if(n>60)
+    {
+      this->a = (double)b + 1.5;
+      c[1][1] = ++a;
+    }
+
+    return c[1][1] + (int)b;
+  }
+};
+
+// CHECK: define {{.*}}@{{.*}}bar{{.*}}
+int bar(int n){
+  int a = 0;
+
+  // CHECK: call {{.*}}i32 [[FOO]](i32 {{.*}})
+  a += foo(n);
+
+  S1 S;
+  // CHECK: call {{.*}}i32 [[FS1:@.+]]([[S1]]* {{.*}}, i32 {{.*}})
+  a += S.r1(n);
+
+  // CHECK: call {{.*}}i32 [[FSTATIC:@.+]](i32 {{.*}})
+  a += fstatic(n);
+
+  // CHECK: call {{.*}}i32 [[FTEMPLATE:@.+]](i32 {{.*}})
+  a += ftemplate<int>(n);
+
+  return a;
+}
+
+//
+// CHECK: define {{.*}}[[FS1]]
+//
+// We capture 2 VLA sizes in this target region
+// CHECK-DAG:   store i64 %{{[^,]+}}, i64* [[VLA0:%[^,]+]]
+// CHECK-DAG:   store i64 %{{[^,]+}}, i64* [[VLA1:%[^,]+]]
+// CHECK:       [[IF:%.+]] = icmp sgt i32 {{[^,]+}}, 60
+// CHECK:       br i1 [[IF]], label %[[TRY:[^,]+]], label %[[FAIL:[^,]+]]
+// CHECK:       [[TRY]]
+// CHECK-DAG:   [[RET:%.+]] = call i32 @__tgt_target(i32 -1, i8* @{{[^,]+}}, i32 5, i8** [[BP:%[^,]+]], i8** [[P:%[^,]+]], i64* [[S:%[^,]+]], i32* getelementptr inbounds ([5 x i32], [5 x i32]* [[MAPT7]], i32 0, i32 0))
+
+// CHECK-DAG:   [[SADDR0:%.+]] = getelementptr inbounds i64, i64* [[S]], i32 [[IDX0:[0-9]+]]
+// CHECK-DAG:   [[BPADDR0:%.+]] = getelementptr inbounds i8*, i8** [[BP]], i32 [[IDX0]]
+// CHECK-DAG:   [[PADDR0:%.+]] = getelementptr inbounds i8*, i8** [[P]], i32 [[IDX0]]
+// CHECK-DAG:   [[SADDR1:%.+]] = getelementptr inbounds i64, i64* [[S]], i32 [[IDX1:[0-9]+]]
+// CHECK-DAG:   [[BPADDR1:%.+]] = getelementptr inbounds i8*, i8** [[BP]], i32 [[IDX1]]
+// CHECK-DAG:   [[PADDR1:%.+]] = getelementptr inbounds i8*, i8** [[P]], i32 [[IDX1]]
+// CHECK-DAG:   [[SADDR2:%.+]] = getelementptr inbounds i64, i64* [[S]], i32 [[IDX2:[0-9]+]]
+// CHECK-DAG:   [[BPADDR2:%.+]] = getelementptr inbounds i8*, i8** [[BP]], i32 [[IDX2]]
+// CHECK-DAG:   [[PADDR2:%.+]] = getelementptr inbounds i8*, i8** [[P]], i32 [[IDX2]]
+// CHECK-DAG:   [[SADDR3:%.+]] = getelementptr inbounds i64, i64* [[S]], i32 [[IDX3:[0-9]+]]
+// CHECK-DAG:   [[BPADDR3:%.+]] = getelementptr inbounds i8*, i8** [[BP]], i32 [[IDX3]]
+// CHECK-DAG:   [[PADDR3:%.+]] = getelementptr inbounds i8*, i8** [[P]], i32 [[IDX3]]
+
+// The names below are not necessarily consistent with the names used for the
+// addresses above as some are repeated.
+// CHECK-DAG:   [[BP0:%[^,]+]] = bitcast i64* [[VLA0]] to i8*
+// CHECK-DAG:   [[P0:%[^,]+]] = bitcast i64* [[VLA0]] to i8*
+// CHECK-DAG:   store i8* [[BP0]], i8** {{%[^,]+}}
+// CHECK-DAG:   store i8* [[P0]], i8** {{%[^,]+}}
+// CHECK-DAG:   store i64 8, i64* {{%[^,]+}}
+
+// CHECK-DAG:   [[BP1:%[^,]+]] = bitcast i64* [[VLA1]] to i8*
+// CHECK-DAG:   [[P1:%[^,]+]] = bitcast i64* [[VLA1]] to i8*
+// CHECK-DAG:   store i8* [[BP1]], i8** {{%[^,]+}}
+// CHECK-DAG:   store i8* [[P1]], i8** {{%[^,]+}}
+// CHECK-DAG:   store i64 8, i64* {{%[^,]+}}
+
+// CHECK-DAG:   [[BP2:%[^,]+]] = bitcast i32* %{{.+}} to i8*
+// CHECK-DAG:   [[P2:%[^,]+]] = bitcast i32* %{{.+}} to i8*
+// CHECK-DAG:   store i8* [[BP2]], i8** {{%[^,]+}}
+// CHECK-DAG:   store i8* [[P2]], i8** {{%[^,]+}}
+// CHECK-DAG:   store i64 4, i64* {{%[^,]+}}
+
+// CHECK-DAG:   [[BP3:%[^,]+]] = bitcast [[S1]]* %{{.+}} to i8*
+// CHECK-DAG:   [[P3:%[^,]+]] = bitcast [[S1]]* %{{.+}} to i8*
+// CHECK-DAG:   store i8* [[BP3]], i8** {{%[^,]+}}
+// CHECK-DAG:   store i8* [[P3]], i8** {{%[^,]+}}
+// CHECK-DAG:   store i64 8, i64* {{%[^,]+}}
+
+// CHECK-DAG:   [[BP4:%[^,]+]] = bitcast i16* %{{.+}} to i8*
+// CHECK-DAG:   [[P4:%[^,]+]] = bitcast i16* %{{.+}} to i8*
+// CHECK-DAG:   store i8* [[BP4]], i8** {{%[^,]+}}
+// CHECK-DAG:   store i8* [[P4]], i8** {{%[^,]+}}
+// CHECK-DAG:   store i64 2, i64* {{%[^,]+}}
+
+// CHECK:       [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
+// CHECK-NEXT:  br i1 [[ERROR]], label %[[FAIL]], label %[[END:[^,]+]]
+// CHECK:       [[FAIL]]
+// CHECK:       call void [[HVT7:@.+]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}})
+// CHECK-NEXT:  br label %[[END]]
+// CHECK:       [[END]]
+
+//
+// CHECK: define {{.*}}[[FSTATIC]]
+//
+// CHECK:       [[IF:%.+]] = icmp sgt i32 {{[^,]+}}, 50
+// CHECK:       br i1 [[IF]], label %[[TRY:[^,]+]], label %[[FAIL:[^,]+]]
+// CHECK:       [[TRY]]
+// CHECK-DAG:   [[RET:%.+]] = call i32 @__tgt_target(i32 -1, i8* @{{[^,]+}}, i32 4, i8** [[BP:%[^,]+]], i8** [[P:%[^,]+]], i64* [[S:%[^,]+]], i32* getelementptr inbounds ([4 x i32], [4 x i32]* [[MAPT6]], i32 0, i32 0))
+
+// CHECK-DAG:   store i64 4, i64* [[SADDR0:%.+]]
+// CHECK-DAG:   [[SADDR0]] = getelementptr inbounds i64, i64* [[S]], i32 [[IDX0:[0-9]+]]
+// CHECK-DAG:   [[BPADDR0:%.+]] = getelementptr inbounds i8*, i8** [[BP]], i32 [[IDX0]]
+// CHECK-DAG:   [[PADDR0:%.+]] = getelementptr inbounds i8*, i8** [[P]], i32 [[IDX0]]
+// CHECK-DAG:   store i8* [[BP0:%[^,]+]], i8** [[BPADDR0]]
+// CHECK-DAG:   store i8* [[P0:%[^,]+]], i8** [[PADDR0]]
+// CHECK-DAG:   [[BP0]] = bitcast i32* %{{.+}} to i8*
+// CHECK-DAG:   [[P0]] = bitcast i32* %{{.+}} to i8*
+
+// CHECK-DAG:   store i64 2, i64* [[SADDR1:%.+]]
+// CHECK-DAG:   [[SADDR1]] = getelementptr inbounds i64, i64* [[S]], i32 [[IDX1:[0-9]+]]
+// CHECK-DAG:   [[BPADDR1:%.+]] = getelementptr inbounds i8*, i8** [[BP]], i32 [[IDX1]]
+// CHECK-DAG:   [[PADDR1:%.+]] = getelementptr inbounds i8*, i8** [[P]], i32 [[IDX1]]
+// CHECK-DAG:   store i8* [[BP1:%[^,]+]], i8** [[BPADDR1]]
+// CHECK-DAG:   store i8* [[P1:%[^,]+]], i8** [[PADDR1]]
+// CHECK-DAG:   [[BP1]] = bitcast i16* %{{.+}} to i8*
+// CHECK-DAG:   [[P1]] = bitcast i16* %{{.+}} to i8*
+
+// CHECK-DAG:   store i64 1, i64* [[SADDR2:%.+]]
+// CHECK-DAG:   [[SADDR2]] = getelementptr inbounds i64, i64* [[S]], i32 [[IDX2:[0-9]+]]
+// CHECK-DAG:   [[BPADDR2:%.+]] = getelementptr inbounds i8*, i8** [[BP]], i32 [[IDX2]]
+// CHECK-DAG:   [[PADDR2:%.+]] = getelementptr inbounds i8*, i8** [[P]], i32 [[IDX2]]
+// CHECK-DAG:   store i8* [[BP2:%[^,]+]], i8** [[BPADDR2]]
+// CHECK-DAG:   store i8* [[P2:%[^,]+]], i8** [[PADDR2]]
+
+// CHECK-DAG:   store i64 40, i64* [[SADDR3:%.+]]
+// CHECK-DAG:   [[SADDR3]] = getelementptr inbounds i64, i64* [[S]], i32 [[IDX3:[0-9]+]]
+// CHECK-DAG:   [[BPADDR3:%.+]] = getelementptr inbounds i8*, i8** [[BP]], i32 [[IDX3]]
+// CHECK-DAG:   [[PADDR3:%.+]] = getelementptr inbounds i8*, i8** [[P]], i32 [[IDX3]]
+// CHECK-DAG:   store i8* [[BP3:%[^,]+]], i8** [[BPADDR3]]
+// CHECK-DAG:   store i8* [[P3:%[^,]+]], i8** [[PADDR3]]
+// CHECK-DAG:   [[BP3]] = bitcast [10 x i32]* %{{.+}} to i8*
+// CHECK-DAG:   [[P3]] = bitcast [10 x i32]* %{{.+}} to i8*
+
+// CHECK:       [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
+// CHECK-NEXT:  br i1 [[ERROR]], label %[[FAIL]], label %[[END:[^,]+]]
+// CHECK:       [[FAIL]]
+// CHECK:       call void [[HVT6:@.+]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}})
+// CHECK-NEXT:  br label %[[END]]
+// CHECK:       [[END]]
+
+//
+// CHECK: define {{.*}}[[FTEMPLATE]]
+//
+// CHECK:       [[IF:%.+]] = icmp sgt i32 {{[^,]+}}, 40
+// CHECK:       br i1 [[IF]], label %[[TRY:[^,]+]], label %[[FAIL:[^,]+]]
+// CHECK:       [[TRY]]
+// CHECK-DAG:   [[RET:%.+]] = call i32 @__tgt_target(i32 -1, i8* @{{[^,]+}}, i32 3, i8** [[BP:%[^,]+]], i8** [[P:%[^,]+]], i64* [[S:%[^,]+]], i32* getelementptr inbounds ([3 x i32], [3 x i32]* [[MAPT5]], i32 0, i32 0))
+
+// CHECK-DAG:   store i64 4, i64* [[SADDR0:%.+]]
+// CHECK-DAG:   [[SADDR0]] = getelementptr inbounds i64, i64* [[S]], i32 [[IDX0:[0-9]+]]
+// CHECK-DAG:   [[BPADDR0:%.+]] = getelementptr inbounds i8*, i8** [[BP]], i32 [[IDX0]]
+// CHECK-DAG:   [[PADDR0:%.+]] = getelementptr inbounds i8*, i8** [[P]], i32 [[IDX0]]
+// CHECK-DAG:   store i8* [[BP0:%[^,]+]], i8** [[BPADDR0]]
+// CHECK-DAG:   store i8* [[P0:%[^,]+]], i8** [[PADDR0]]
+// CHECK-DAG:   [[BP0]] = bitcast i32* %{{.+}} to i8*
+// CHECK-DAG:   [[P0]] = bitcast i32* %{{.+}} to i8*
+
+// CHECK-DAG:   store i64 2, i64* [[SADDR1:%.+]]
+// CHECK-DAG:   [[SADDR1]] = getelementptr inbounds i64, i64* [[S]], i32 [[IDX1:[0-9]+]]
+// CHECK-DAG:   [[BPADDR1:%.+]] = getelementptr inbounds i8*, i8** [[BP]], i32 [[IDX1]]
+// CHECK-DAG:   [[PADDR1:%.+]] = getelementptr inbounds i8*, i8** [[P]], i32 [[IDX1]]
+// CHECK-DAG:   store i8* [[BP1:%[^,]+]], i8** [[BPADDR1]]
+// CHECK-DAG:   store i8* [[P1:%[^,]+]], i8** [[PADDR1]]
+// CHECK-DAG:   [[BP1]] = bitcast i16* %{{.+}} to i8*
+// CHECK-DAG:   [[P1]] = bitcast i16* %{{.+}} to i8*
+
+// CHECK-DAG:   store i64 40, i64* [[SADDR2:%.+]]
+// CHECK-DAG:   [[SADDR2]] = getelementptr inbounds i64, i64* [[S]], i32 [[IDX2:[0-9]+]]
+// CHECK-DAG:   [[BPADDR2:%.+]] = getelementptr inbounds i8*, i8** [[BP]], i32 [[IDX2]]
+// CHECK-DAG:   [[PADDR2:%.+]] = getelementptr inbounds i8*, i8** [[P]], i32 [[IDX2]]
+// CHECK-DAG:   store i8* [[BP2:%[^,]+]], i8** [[BPADDR2]]
+// CHECK-DAG:   store i8* [[P2:%[^,]+]], i8** [[PADDR2]]
+// CHECK-DAG:   [[BP2]] = bitcast [10 x i32]* %{{.+}} to i8*
+// CHECK-DAG:   [[P2]] = bitcast [10 x i32]* %{{.+}} to i8*
+
+// CHECK:       [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
+// CHECK-NEXT:  br i1 [[ERROR]], label %[[FAIL]], label %[[END:[^,]+]]
+// CHECK:       [[FAIL]]
+// CHECK:       call void [[HVT5:@.+]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}})
+// CHECK-NEXT:  br label %[[END]]
+// CHECK:       [[END]]
+
+// Check that the offloading functions are emitted and that the arguments are
+// correct and loaded correctly for the target regions of the callees of bar().
+
+// CHECK:       define internal void [[HVT7]]
+// CHECK-DAG:   [[LOCALX_THIS:%.+]] = alloca [[S1]]*
+// CHECK-DAG:   [[LOCALX_B:%.+]] = alloca i32*
+// CHECK-DAG:   [[LOCALX_C:%.+]] = alloca i16*
+// CHECK-DAG:   [[LOCALX_VLA1:%.+]] = alloca i64*
+// CHECK-DAG:   [[LOCALX_VLA2:%.+]] = alloca i64*
+// CHECK-DAG:   store [[S1]]* [[ARG_THIS:%.+]], [[S1]]** [[LOCALX_THIS]]
+// CHECK-DAG:   store i32* [[ARG_B:%.+]], i32** [[LOCALX_B]]
+// CHECK-DAG:   store i16* [[ARG_C:%.+]], i16** [[LOCALX_C]]
+// CHECK-DAG:   store i64* [[ARG_VLA1:%.+]], i64** [[LOCALX_VLA1]]
+// CHECK-DAG:   store i64* [[ARG_VLA2:%.+]], i64** [[LOCALX_VLA2]]
+// CHECK-DAG:   [[USE_THIS:%.+]] = load [[S1]]*, [[S1]]** [[LOCALX_THIS]]
+// CHECK-DAG:   [[USE_B:%.+]] = load i32*, i32** [[LOCALX_B]]
+// CHECK-DAG:   [[USE_C:%.+]] = load i16*, i16** [[LOCALX_C]]
+// CHECK-DAG:   [[USE_VLA1:%.+]] = load i64*, i64** [[LOCALX_VLA1]]
+// CHECK-DAG:   [[USE_VLA2:%.+]] = load i64*, i64** [[LOCALX_VLA2]]
+// CHECK-DAG:   getelementptr inbounds [[S1]], [[S1]]* [[USE_THIS]], i{{.*}} 0, i{{.*}} 0
+// CHECK-DAG:   load i32, i32* [[USE_B]]
+// CHECK-DAG:   [[VLAMUL:%.+]] = mul {{.*}}i64 1, %{{.*}}
+// CHECK-DAG:   getelementptr inbounds i16, i16* [[USE_C]], i{{.*}} [[VLAMUL]]
+// CHECK-DAG:   load i64, i64* [[USE_VLA1]]
+// CHECK-DAG:   load i64, i64* [[USE_VLA2]]
+
+// CHECK:       define internal void [[HVT6]]
+// CHECK-DAG:   [[LOCALX_A:%.+]] = alloca i32*
+// CHECK-DAG:   [[LOCALX_AA:%.+]] = alloca i16*
+// CHECK-DAG:   [[LOCALX_AAA:%.+]] = alloca i8*
+// CHECK-DAG:   [[LOCALX_B:%.+]] = alloca [10 x i32]*
+// CHECK-DAG:   store i32* [[ARG_A:%.+]], i32** [[LOCALX_A]]
+// CHECK-DAG:   store i16* [[ARG_AA:%.+]], i16** [[LOCALX_AA]]
+// CHECK-DAG:   store i8* [[ARG_AAA:%.+]], i8** [[LOCALX_AAA]]
+// CHECK-DAG:   store [10 x i32]* [[ARG_B:%.+]], [10 x i32]** [[LOCALX_B]]
+// CHECK-DAG:   [[USE_A:%.+]] = load i32*, i32** [[LOCALX_A]]
+// CHECK-DAG:   [[USE_AA:%.+]] = load i16*, i16** [[LOCALX_AA]]
+// CHECK-DAG:   [[USE_AAA:%.+]] = load i8*, i8** [[LOCALX_AAA]]
+// CHECK-DAG:   [[USE_B:%.+]] = load [10 x i32]*, [10 x i32]** [[LOCALX_B]]
+// CHECK-DAG:   load i32, i32* [[USE_A]]
+// CHECK-DAG:   load i16, i16* [[USE_AA]]
+// CHECK-DAG:   load i8, i8* [[USE_AAA]]
+// CHECK-DAG:   getelementptr inbounds [10 x i32], [10 x i32]* [[USE_B]], i{{.*}} 0, i{{.*}} 2
+
+// CHECK:       define internal void [[HVT5]]
+// CHECK-DAG:   [[LOCALX_A:%.+]] = alloca i32*
+// CHECK-DAG:   [[LOCALX_AA:%.+]] = alloca i16*
+// CHECK-DAG:   [[LOCALX_B:%.+]] = alloca [10 x i32]*
+// CHECK-DAG:   store i32* [[ARG_A:%.+]], i32** [[LOCALX_A]]
+// CHECK-DAG:   store i16* [[ARG_AA:%.+]], i16** [[LOCALX_AA]]
+// CHECK-DAG:   store [10 x i32]* [[ARG_B:%.+]], [10 x i32]** [[LOCALX_B]]
+// CHECK-DAG:   [[USE_A:%.+]] = load i32*, i32** [[LOCALX_A]]
+// CHECK-DAG:   [[USE_AA:%.+]] = load i16*, i16** [[LOCALX_AA]]
+// CHECK-DAG:   [[USE_B:%.+]] = load [10 x i32]*, [10 x i32]** [[LOCALX_B]]
+// CHECK-DAG:   load i32, i32* [[USE_A]]
+// CHECK-DAG:   load i16, i16* [[USE_AA]]
+// CHECK-DAG:   getelementptr inbounds [10 x i32], [10 x i32]* [[USE_B]], i{{.*}} 0, i{{.*}} 2
+#endif