[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

2212–2240 ↗

(On Diff #30172)

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

2215–2245 ↗

(On Diff #30264)

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

2215–2245 ↗

(On Diff #30264)

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

3041

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

3071

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

lib/CodeGen/CGStmt.cpp

2215–2245 ↗

(On Diff #30264)

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

769 ↗

(On Diff #30264)

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

1550 ↗

(On Diff #30264)

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

[rjmccall]

lib/CodeGen/CGOpenMPRuntime.cpp

865–866

Spurious "to" at the start.

2925

getTypeStoreSize()

2934

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?

2944

Same thing, please ask the ASTContext.

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

3013

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

3018

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.

3021

Please comment boolean arguments like this:

/*constant*/ true

And please mark this variable unnamed_addr.

3048

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

2215–2245 ↗

(On Diff #30264)

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

2215–2245 ↗

(On Diff #30264)

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

865–866

Fixed!

2925

Now using getTypeSizeInChars from the ASTcontext as suggested bellow.

2934

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.

2944

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.

3013

Done!

3018

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.

3021

Added comments for the boolean.

Setting unnamed_addr now, thanks!

3041

Ok, now using the types explicitly.

3048

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

3071

Done!

lib/CodeGen/CGStmt.cpp

2215–2245 ↗

(On Diff #30264)

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

769 ↗

(On Diff #30264)

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

1550 ↗

(On Diff #30264)

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 ↗

(On Diff #32211)

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

2924

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

2925

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

2934

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

2951

The same

lib/CodeGen/CGStmtOpenMP.cpp

2140–2141

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

2143

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 ↗

(On Diff #32211)

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

2924

Done!

2925

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

2934

done!

2951

Done!

lib/CodeGen/CGStmtOpenMP.cpp

2140–2141

Now using MakeNaturalAlignAddrLValue.

2143

Fixed.

[ABataev]

lib/CodeGen/CGExpr.cpp

1969–1970 ↗

(On Diff #32322)

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

3036

Maybe CGF.getContext().getSizeType()?

3134

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–1970 ↗

(On Diff #32322)

Ok, I am now using the proxy function.

lib/CodeGen/CGOpenMPRuntime.cpp

3036

Done!

3134

Done!

test/OpenMP/target_codegen.cpp

2–4

Done!

[ABataev]

lib/CodeGen/CGOpenMPRuntime.cpp

2880

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.

2899–2904

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.

2963–2965

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

3063–3100

Try instead:

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

3154–3157

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.

751

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

2150–2214

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

2880

Done!

2899–2904

Done!

2963–2965

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

3063–3100

Done!

3154–3157

Done!

lib/CodeGen/CGStmtOpenMP.cpp

2150–2214

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.

746

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.

746

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!

746

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

2991–3005

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

2991–3005

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

2135–2136

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

2141–2142

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

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, size_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
   };
+
   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
   ///                                  member  */
   ///    kmp_int32 reserved_2;   /**<  not really used in Fortran any more;
   ///                                  see above */
   ///#if USE_ITT_BUILD
   ///                            /*  but currently used for storing
   ///                                region-specific ITT */
   ///                            /*  contextual information. */
   ///#endif /* USE_ITT_BUILD */
   ///    kmp_int32 reserved_3;   /**< source[4] in Fortran, do not use for
   ///                                 C++  */
   ///    char const *psource;    /**< String describing the source location.
   ///                            The string is composed of semi-colon separated
   //                             fields which describe the source file,
   ///                            the function and a pair of line numbers that
   ///                            delimit the construct.
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(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,
+                              ArrayRef<llvm::Value *> VLASizesInit);
 };
 
 } // 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:
   }
 
 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(CS, TargetRegion, CodeGen, OMPD_target) {}
+
+  /// \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 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, size_t *arg_sizes, int32_t
+    // *arg_types);
+    llvm::Type *TypeParams[] = {CGM.Int32Ty,
+                                CGM.VoidPtrTy,
+                                CGM.Int32Ty,
+                                CGM.VoidPtrPtrTy,
+                                CGM.VoidPtrPtrTy,
+                                CGM.SizeTy->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");
if (CancelDest.isValid()) {
       emitBarrierCall(CGF, Loc, OMPD_unknown, /*CheckForCancel=*/false);
       // exit from construct;
       CGF.EmitBranchThroughCleanup(CancelDest);
       CGF.EmitBlock(ContBB, /*IsFinished=*/true);
     }
   }
 }
 
+/// \brief Emit a proxy function for target regions, which accepts each capture
+/// as an argument and maps each argument to a new context to pass to the
+/// main outlined function.
+/// \code
+/// void .omp_offloading_entry.(ty1 capture1, ..., tyn capture n) {
+///   Context ctx;
+///   // fillup the context with all the arguments.
+///   .omp_offloading.(ctx);
+///   return;
+/// }
+/// \endcode
+static llvm::Value *emitProxyTargetFunction(CodeGenModule &CGM,
+                                            const CapturedStmt &CS,
+                                            SourceLocation Loc,
+                                            llvm::Value *TargetFunction) {
+  auto &C = CGM.getContext();
+
+  // Collect the arguments of the main function.
+  FunctionArgList Args;
+  auto *RDecl = CS.getCapturedRecordDecl();
+  auto RI = RDecl->field_begin();
+  for (CapturedStmt::const_capture_iterator CI = CS.capture_begin(),
+                                            CE = CS.capture_end();
+       CI != CE; ++CI, ++RI) {
+    StringRef Name;
+    QualType Ty;
+    if (CI->capturesVariableArrayType()) {
+      Ty = C.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 a variable-length array type, we use a
+      // pointer type to the element type of the VLA. The resulting IR type is
+      // the same, will avoid the trouble of emiting the VLAs, and we won't
+      // be dereferencing the the VLA in this function anyway.
+
+      // Get the element type of the variable-length array type.
+      while (auto *VAT = dyn_cast<VariableArrayType>(Ty.getTypePtr()))
+        Ty = VAT->getElementType();
+      Ty = C.getPointerType(Ty);
+      Name = VD->getName();
+    }
+
+    IdentifierInfo *ParamName = &C.Idents.get(Name);
+    ImplicitParamDecl *Param =
+        ImplicitParamDecl::Create(C, /*DC=*/nullptr, Loc, ParamName, Ty);
+    Args.push_back(Param);
+  }
+
+  // Create the main function.
+  FunctionType::ExtInfo Info;
+  auto &FnInfo =
+      CGM.getTypes().arrangeFreeFunctionDeclaration(C.VoidTy, Args, Info,
+                                                    /*isVariadic=*/false);
+  auto *FnTy = CGM.getTypes().GetFunctionType(FnInfo);
+  auto *Fn = llvm::Function::Create(FnTy, llvm::GlobalValue::InternalLinkage,
+                                    ".omp_offloading_entry.", &CGM.getModule());
+  CGM.SetLLVMFunctionAttributes(/*D=*/nullptr, FnInfo, Fn);
+  CodeGenFunction CGF(CGM);
+  CGF.disableDebugInfo();
+  CGF.StartFunction(GlobalDecl(), C.VoidTy, Fn, FnInfo, Args);
+
+  // Create the context.
+  auto ContextType = C.getTagDeclType(RDecl);
+  auto *Context = CGF.CreateMemTemp(ContextType, ".omp_offloading.context");
+
+  LValue ContextLV = CGF.MakeNaturalAlignAddrLValue(Context, ContextType);
+
+  auto AI = Args.begin();
+  for (RecordDecl::field_iterator RI = RDecl->field_begin(),
+                                  RE = RDecl->field_end();
+       RI != RE; ++RI, ++AI) {
+
+    // Get the address of the field.
+    auto Dst = CGF.EmitLValueForFieldInitialization(ContextLV, *RI);
+    auto ArgRef = CGF.MakeNaturalAlignAddrLValue(CGF.GetAddrOfLocalVar(*AI),
+                                                 (*AI)->getType());
+    auto Arg = CGF.EmitLoadOfLValue(ArgRef, Loc);
+    if (RI->hasCapturedVLAType()) {
+      auto *Val = CGF.EmitLoadOfScalar(
+          Arg.getScalarVal(), /*Volatile=*/false,
+          C.getTypeAlignInChars(RI->getType()).getQuantity(), RI->getType(),
+          Loc);
+      CGF.EmitStoreOfScalar(Val, Dst);
+      continue;
+    }
+    CGF.EmitStoreOfScalar(Arg.getScalarVal(), Dst);
+  }
+  CGF.EmitCallOrInvoke(TargetFunction, Context);
+  CGF.FinishFunction();
+  return Fn;
+}
+
+llvm::Value *
+CGOpenMPRuntime::emitTargetOutlinedFunction(const OMPExecutableDirective &D,
+                                            const RegionCodeGenTy &CodeGen) {
+  const CapturedStmt &CS = *cast<CapturedStmt>(D.getAssociatedStmt());
+
+  CodeGenFunction CGF(CGM, true);
+  CGOpenMPTargetRegionInfo CGInfo(CS, CodeGen);
+  CodeGenFunction::CGCapturedStmtRAII CapInfoRAII(CGF, &CGInfo);
+  auto *Fn = CGF.GenerateCapturedStmtFunction(CS);
+  Fn->addFnAttr(llvm::Attribute::AlwaysInline);
+
+  return emitProxyTargetFunction(CGM, CS, D.getLocStart(), Fn);
+}
+
+void CGOpenMPRuntime::emitTargetCall(CodeGenFunction &CGF,
+                                     const OMPExecutableDirective &D,
+                                     llvm::Value *OutlinedFn,
+                                     const Expr *IfCond, const Expr *Device,
+                                     ArrayRef<llvm::Value *> VLASizesInit) {
+
+  /// \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,
+  };
+
+  // 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 = VLASizesInit.size();
+
+  const CapturedStmt &CS = *cast<CapturedStmt>(D.getAssociatedStmt());
+  auto RI = CS.getCapturedRecordDecl()->field_begin();
+  auto II = CS.capture_init_begin();
+  auto VI = VLASizesInit.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()) {
+      LValue LV = CGF.MakeNaturalAlignAddrLValue(
+          CGF.CreateMemTemp(RI->getType(), "__vla_size"), RI->getType());
+      CGF.EmitStoreThroughLValue(RValue::get(*VI), LV);
+      ++VI;
+      BasePointer = Pointer = LV.getAddress();
+      uint64_t SizeVal =
+          CGM.getContext().getTypeSizeInChars(RI->getType()).getQuantity();
+      Size = llvm::ConstantInt::get(CGM.SizeTy, SizeVal);
+
+      hasVLACaptures = true;
+      // VLA sizes don't need to be copied back from the device.
+      MapType = OMP_MAP_TO;
+    } else if (CI->capturesThis()) {
+      BasePointer = Pointer = CGF.LoadCXXThis();
+      const PointerType *PtrTy = cast<PointerType>(RI->getType().getTypePtr());
+      uint64_t SizeVal = CGM.getContext()
+                             .getTypeSizeInChars(PtrTy->getPointeeType())
+                             .getQuantity();
+      Size = llvm::ConstantInt::get(CGM.SizeTy, SizeVal);
+
+      // Default map type.
+      MapType = OMP_MAP_TO | 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())) {
+        auto VATInfo = CGF.getVLASize(VAT);
+        Size = llvm::ConstantInt::get(
+            CGM.SizeTy,
+            CGM.getContext().getTypeSizeInChars(VATInfo.second).getQuantity());
+        Size = CGF.Builder.CreateNUWMul(Size, VATInfo.first);
+      } else {
+        uint64_t ElementTypeSize =
+            CGM.getContext().getTypeSizeInChars(ElementType).getQuantity();
+        Size = llvm::ConstantInt::get(CGM.SizeTy, ElementTypeSize);
+      }
+
+      // Default map type.
+      MapType = OMP_MAP_TO | 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().getSizeType(), 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<llvm::Constant *, 16> ConstSizes;
+      for (auto S : Sizes)
+        ConstSizes.push_back(cast<llvm::Constant>(S));
+
+      auto *SizesArrayInit = llvm::ConstantArray::get(
+          llvm::ArrayType::get(CGM.SizeTy, ConstSizes.size()), 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.SizeTy, PointerNumVal), SizesArray, 0, i);
+        CGF.Builder.CreateStore(
+            CGF.Builder.CreateIntCast(Sizes[i], CGM.SizeTy, /*isSigned=*/true),
+            S);
+      }
+    }
+
+    BasePointersArray = CGF.Builder.CreateConstInBoundsGEP2_32(
+        llvm::ArrayType::get(CGM.VoidPtrTy, PointerNumVal), BasePointersArray,
+        0, 0);
+    PointersArray = CGF.Builder.CreateConstInBoundsGEP2_32(
+        llvm::ArrayType::get(CGM.VoidPtrTy, PointerNumVal), PointersArray, 0,
+        0);
+    SizesArray = CGF.Builder.CreateConstInBoundsGEP2_32(
+        llvm::ArrayType::get(CGM.SizeTy, PointerNumVal), SizesArray, 0, 0);
+    MapTypesArray = CGF.Builder.CreateConstInBoundsGEP2_32(
+        llvm::ArrayType::get(CGM.Int32Ty, PointerNumVal), MapTypesArray, 0, 0);
+
+  } else {
+    BasePointersArray = llvm::ConstantPointerNull::get(CGM.VoidPtrPtrTy);
+    PointersArray = llvm::ConstantPointerNull::get(CGM.VoidPtrPtrTy);
+    SizesArray = llvm::ConstantPointerNull::get(CGM.SizeTy->getPointerTo());
+    MapTypesArray = llvm::ConstantPointerNull::get(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,
+                                   BasePointersArray, PointersArray, SizesArray,
+                                   MapTypesArray};
+  auto Return = CGF.EmitRuntimeCall(createRuntimeFunction(OMPRTL__tgt_target),
+                                    OffloadingArgs);
+  auto Error = CGF.Builder.CreateIsNotNull(Return);
+  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/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::EmitOMPTargetDirective(const OMPTargetDirective &S) {
+  LexicalScope Scope(*this, S.getSourceRange());
+  const CapturedStmt &CS = *cast<CapturedStmt>(S.getAssociatedStmt());
+
+  // Emit target region as a standalone region.
+  auto &&CodeGen = [&CS](CodeGenFunction &CGF) {
+    CGF.EmitStmt(CS.getCapturedStmt());
+  };
+
+  // Obtain the target region outlined function.
+  llvm::Value *Fn =
+      CGM.getOpenMPRuntime().emitTargetOutlinedFunction(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();
+  }
+
+  SmallVector<llvm::Value *, 8> VLASizesInit;
+  for (auto F : CS.getCapturedRecordDecl()->fields())
+    if (F->hasCapturedVLAType()) {
+      auto *V = VLASizeMap[F->getCapturedVLAType()->getSizeExpr()];
+      assert(V && "VLA size must exist!");
+      VLASizesInit.push_back(V);
+    }
+
+  CGM.getOpenMPRuntime().emitTargetCall(*this, S, Fn, IfCond, Device,
+                                        VLASizesInit);
 }
 
 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.
 void CodeGenFunction::EmitOMPTargetDataDirective(
     const OMPTargetDataDirective &S) {
 
   // emit the code inside the construct for now
   auto CS = cast<CapturedStmt>(S.getAssociatedStmt());
   CGM.getOpenMPRuntime().emitInlinedDirective(
       *this, OMPD_target_data,
       [&CS](CodeGenFunction &CGF) { CGF.EmitStmt(CS->getCapturedStmt()); });
 }

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
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
+// expected-no-diagnostics
+#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, only 6 will have mapped arguments, and only 4 have all-constant map
+// sizes.
+
+// CHECK-DAG: [[SIZET2:@.+]] = private unnamed_addr constant [1 x i{{32|64}}] [i[[SZ:32|64]] 2]
+// CHECK-DAG: [[MAPT2:@.+]] = private unnamed_addr constant [1 x i32] [i32 3]
+// CHECK-DAG: [[SIZET3:@.+]] = private unnamed_addr constant [2 x i[[SZ]]] [i[[SZ]] 4, i[[SZ]] 2]
+// CHECK-DAG: [[MAPT3:@.+]] = private unnamed_addr constant [2 x i32] [i32 3, i32 3]
+// CHECK-DAG: [[MAPT4:@.+]] = private unnamed_addr constant [9 x i32] [i32 3, i32 3, i32 1, i32 3, i32 3, i32 1, i32 1, i32 3, i32 3]
+// CHECK-DAG: [[SIZET5:@.+]] = private unnamed_addr constant [3 x i[[SZ]]] [i[[SZ]] 4, i[[SZ]] 2, i[[SZ]] 40]
+// CHECK-DAG: [[MAPT5:@.+]] = private unnamed_addr constant [3 x i32] [i32 3, i32 3, i32 3]
+// CHECK-DAG: [[SIZET6:@.+]] = private unnamed_addr constant [4 x i[[SZ]]] [i[[SZ]] 4, i[[SZ]] 2, i[[SZ]] 1, i[[SZ]] 40]
+// CHECK-DAG: [[MAPT6:@.+]] = private unnamed_addr constant [4 x i32] [i32 3, i32 3, i32 3, i32 3]
+// CHECK-DAG: [[MAPT7:@.+]] = private unnamed_addr 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 long, char> d;
+
+  // CHECK:       br label %[[TRY:[^,]+]]
+  // CHECK:       [[TRY]]
+  // CHECK:       [[RET:%.+]] = call i32 @__tgt_target(i32 -1, i8* @{{[^,]+}}, i32 0, i8** null, i8** null, i[[SZ]]* 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:%[^,]+]], i[[SZ]]* getelementptr inbounds ([1 x i[[SZ]]], [1 x i[[SZ]]]* [[SIZET2]], i32 0, i32 0), i32* getelementptr inbounds ([1 x i32], [1 x i32]* [[MAPT2]], i32 0, i32 0))
+  // CHECK-DAG:   [[BP]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[BPR:%[^,]+]], i32 0, i32 0
+  // CHECK-DAG:   [[P]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[PR:%[^,]+]], i32 0, i32 0
+  // CHECK-DAG:   [[BPADDR0:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[BPR]], i32 0, i32 [[IDX0:[0-9]+]]
+  // CHECK-DAG:   [[PADDR0:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[PR]], i32 0, i32 [[IDX0]]
+  // CHECK-DAG:   store i8* [[BP0:%[^,]+]], i8** [[BPADDR0]]
+  // CHECK-DAG:   store i8* [[P0:%[^,]+]], i8** [[PADDR0]]
+  // CHECK-DAG:   [[BP0]] = bitcast i16* %{{.+}} to i8*
+  // CHECK-DAG:   [[P0]] = 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 [[HVT2:@.+]](i16* {{[^,]+}})
+  // CHECK-NEXT:  br label %[[END]]
+  // CHECK:       [[END]]
+  #pragma omp target if(1)
+  {
+    aa += 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** [[BPR:%[^,]+]], i8** [[PR:%[^,]+]], i[[SZ]]* getelementptr inbounds ([2 x i[[SZ]]], [2 x i[[SZ]]]* [[SIZET3]], i32 0, i32 0), i32* getelementptr inbounds ([2 x i32], [2 x i32]* [[MAPT3]], i32 0, i32 0))
+  // CHECK-DAG:   [[BPR]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BP:%[^,]+]], i32 0, i32 0
+  // CHECK-DAG:   [[PR]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[P:%[^,]+]], i32 0, i32 0
+
+  // CHECK-DAG:   [[BPADDR0:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BP]], i32 0, i32 0
+  // CHECK-DAG:   [[PADDR0:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[P]], i32 0, i32 0
+  // 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:   [[BPADDR1:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BP]], i32 0, i32 1
+  // CHECK-DAG:   [[PADDR1:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[P]], i32 0, i32 1
+  // 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:       store i[[SZ]] [[BNELEMSIZE:%.+]], i[[SZ]]* [[VLA0:%[^,]+]]
+  // CHECK:       [[BNSIZE:%.+]] = mul nuw i[[SZ]] 4, [[BNELEMSIZE:%.+]]
+
+  // CHECK:       store i[[SZ]] 5, i[[SZ]]* [[VLA1:%[^,]+]]
+  // CHECK:       store i[[SZ]] [[CNELEMSIZE1:%.+]], i[[SZ]]* [[VLA2:%[^,]+]]
+  // CHECK:       [[CNELEMSIZE2:%.+]] = mul nuw i[[SZ]] 5, [[CNELEMSIZE1]]
+  // CHECK:       [[CNSIZE:%.+]] = mul nuw i[[SZ]] 8, [[CNELEMSIZE2]]
+
+  // 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** [[BPR:%[^,]+]], i8** [[PR:%[^,]+]], i[[SZ]]* [[SR:%[^,]+]], i32* getelementptr inbounds ([9 x i32], [9 x i32]* [[MAPT4]], i32 0, i32 0))
+  // CHECK-DAG:   [[BPR]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[BP:%[^,]+]], i32 0, i32 0
+  // CHECK-DAG:   [[PR]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[P:%[^,]+]], i32 0, i32 0
+  // CHECK-DAG:   [[SR]] = getelementptr inbounds [9 x i[[SZ]]], [9 x i[[SZ]]]* [[S:%[^,]+]], i32 0, i32 0
+
+  // CHECK-DAG:   [[SADDR0:%.+]] = getelementptr inbounds [9 x i[[SZ]]], [9 x i[[SZ]]]* [[S]], i32 0, i32 [[IDX0:[0-9]+]]
+  // CHECK-DAG:   [[BPADDR0:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[BP]], i32 0, i32 [[IDX0]]
+  // CHECK-DAG:   [[PADDR0:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[P]], i32 0, i32 [[IDX0]]
+  // CHECK-DAG:   [[SADDR1:%.+]] = getelementptr inbounds [9 x i[[SZ]]], [9 x i[[SZ]]]* [[S]], i32 0, i32 [[IDX1:[0-9]+]]
+  // CHECK-DAG:   [[BPADDR1:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[BP]], i32 0, i32 [[IDX1]]
+  // CHECK-DAG:   [[PADDR1:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[P]], i32 0, i32 [[IDX1]]
+  // CHECK-DAG:   [[SADDR2:%.+]] = getelementptr inbounds [9 x i[[SZ]]], [9 x i[[SZ]]]* [[S]], i32 0, i32 [[IDX2:[0-9]+]]
+  // CHECK-DAG:   [[BPADDR2:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[BP]], i32 0, i32 [[IDX2]]
+  // CHECK-DAG:   [[PADDR2:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[P]], i32 0, i32 [[IDX2]]
+  // CHECK-DAG:   [[SADDR3:%.+]] = getelementptr inbounds [9 x i[[SZ]]], [9 x i[[SZ]]]* [[S]], i32 0, i32 [[IDX3:[0-9]+]]
+  // CHECK-DAG:   [[BPADDR3:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[BP]], i32 0, i32 [[IDX3]]
+  // CHECK-DAG:   [[PADDR3:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[P]], i32 0, i32 [[IDX3]]
+  // CHECK-DAG:   [[SADDR4:%.+]] = getelementptr inbounds [9 x i[[SZ]]], [9 x i[[SZ]]]* [[S]], i32 0, i32 [[IDX4:[0-9]+]]
+  // CHECK-DAG:   [[BPADDR4:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[BP]], i32 0, i32 [[IDX4]]
+  // CHECK-DAG:   [[PADDR4:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[P]], i32 0, i32 [[IDX4]]
+  // CHECK-DAG:   [[SADDR5:%.+]] = getelementptr inbounds [9 x i[[SZ]]], [9 x i[[SZ]]]* [[S]], i32 0, i32 [[IDX5:[0-9]+]]
+  // CHECK-DAG:   [[BPADDR5:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[BP]], i32 0, i32 [[IDX5]]
+  // CHECK-DAG:   [[PADDR5:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[P]], i32 0, i32 [[IDX5]]
+  // CHECK-DAG:   [[SADDR6:%.+]] = getelementptr inbounds [9 x i[[SZ]]], [9 x i[[SZ]]]* [[S]], i32 0, i32 [[IDX6:[0-9]+]]
+  // CHECK-DAG:   [[BPADDR6:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[BP]], i32 0, i32 [[IDX6]]
+  // CHECK-DAG:   [[PADDR6:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[P]], i32 0, i32 [[IDX6]]
+  // CHECK-DAG:   [[SADDR7:%.+]] = getelementptr inbounds [9 x i[[SZ]]], [9 x i[[SZ]]]* [[S]], i32 0, i32 [[IDX7:[0-9]+]]
+  // CHECK-DAG:   [[BPADDR7:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[BP]], i32 0, i32 [[IDX7]]
+  // CHECK-DAG:   [[PADDR7:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[P]], i32 0, i32 [[IDX7]]
+  // CHECK-DAG:   [[SADDR8:%.+]] = getelementptr inbounds [9 x i[[SZ]]], [9 x i[[SZ]]]* [[S]], i32 0, i32 [[IDX8:[0-9]+]]
+  // CHECK-DAG:   [[BPADDR8:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[BP]], i32 0, i32 [[IDX8]]
+  // CHECK-DAG:   [[PADDR8:%.+]] = getelementptr inbounds [9 x i8*], [9 x i8*]* [[P]], i32 0, 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 i[[SZ]]* [[VLA0]] to i8*
+  // CHECK-DAG:   [[P0:%[^,]+]] = bitcast i[[SZ]]* [[VLA0]] to i8*
+  // CHECK-DAG:   store i8* [[BP0]], i8** {{%[^,]+}}
+  // CHECK-DAG:   store i8* [[P0]], i8** {{%[^,]+}}
+  // CHECK-DAG:   store i[[SZ]] {{4|8}}, i[[SZ]]* {{%[^,]+}}
+
+  // CHECK-DAG:   [[BP1:%[^,]+]] = bitcast i[[SZ]]* [[VLA1]] to i8*
+  // CHECK-DAG:   [[P1:%[^,]+]] = bitcast i[[SZ]]* [[VLA1]] to i8*
+  // CHECK-DAG:   store i8* [[BP1]], i8** {{%[^,]+}}
+  // CHECK-DAG:   store i8* [[P1]], i8** {{%[^,]+}}
+  // CHECK-DAG:   store i[[SZ]] {{4|8}}, i[[SZ]]* {{%[^,]+}}
+
+  // CHECK-DAG:   [[BP2:%[^,]+]] = bitcast i[[SZ]]* [[VLA2]] to i8*
+  // CHECK-DAG:   [[P2:%[^,]+]] = bitcast i[[SZ]]* [[VLA2]] to i8*
+  // CHECK-DAG:   store i8* [[BP2]], i8** {{%[^,]+}}
+  // CHECK-DAG:   store i8* [[P2]], i8** {{%[^,]+}}
+  // CHECK-DAG:   store i[[SZ]] {{4|8}}, i[[SZ]]* {{%[^,]+}}
+
+  // 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 i[[SZ]] 4, i[[SZ]]* {{%[^,]+}}
+
+  // 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 i[[SZ]] 40, i[[SZ]]* {{%[^,]+}}
+
+  // 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 i[[SZ]] [[BNSIZE]], i[[SZ]]* {{%[^,]+}}
+
+  // 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 i[[SZ]] 400, i[[SZ]]* {{%[^,]+}}
+
+  // 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 i[[SZ]] [[CNSIZE]], i[[SZ]]* {{%[^,]+}}
+
+  // 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 i[[SZ]] {{12|16}}, i[[SZ]]* {{%[^,]+}}
+
+  // 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]]
+// Create context and reference.
+// CHECK-DAG:   [[CAPREF:%.+]] = alloca [[CAPTY:%[^,].+]]*,
+// CHECK-DAG:   [[CAP:%.+]] = alloca [[CAPTY]],
+// CHECK-DAG:   store [[CAPTY]]* [[CAP]], [[CAPTY]]** [[CAPREF]]
+// Create local storage for each capture.
+// CHECK-DAG:   [[LOCAL_A:%.+]] = alloca i32*
+// CHECK-DAG:   store i32* [[ARG_A:%.+]], i32** [[LOCAL_A]]
+// Store captures in the context.
+// CHECK-DAG:   [[FIELD_A:%.+]] = getelementptr inbounds [[CAPTY]], [[CAPTY]]* [[CAP]], i32 0, i32 0
+// CHECK-DAG:   [[REF_A:%.+]] = load i32*, i32** [[LOCAL_A]],
+// CHECK-DAG:   store i32* [[REF_A]], i32** [[FIELD_A]],
+// Get capture from context.
+// CHECK:       [[CAP2:%.+]] = load [[CAPTY]]*, [[CAPTY]]** [[CAPREF]]
+// CHECK-DAG:   [[FIELD2_A:%.+]] = getelementptr inbounds [[CAPTY]], [[CAPTY]]* [[CAP2]], i32 0, i32 0
+// CHECK-DAG:   [[REF2_A:%.+]] = load i32*, i32** [[FIELD2_A]]
+// Use captures.
+// CHECK-DAG:   load i32, i32* [[REF2_A]]
+
+// CHECK:       define internal void [[HVT2]]
+// CHECK-DAG:   [[CAPREF:%.+]] = alloca [[CAPTY:%[^,].+]]*,
+// CHECK-DAG:   [[CAP:%.+]] = alloca [[CAPTY]],
+// CHECK-DAG:   store [[CAPTY]]* [[CAP]], [[CAPTY]]** [[CAPREF]]
+// Create local storage for each capture.
+// CHECK-DAG:   [[LOCAL_AA:%.+]] = alloca i16*
+// CHECK-DAG:   store i16* [[ARG_AA:%.+]], i16** [[LOCAL_AA]]
+// Store captures in the context.
+// CHECK-DAG:   [[FIELD_AA:%.+]] = getelementptr inbounds [[CAPTY]], [[CAPTY]]* [[CAP]], i32 0, i32 0
+// CHECK-DAG:   [[REF_AA:%.+]] = load i16*, i16** [[LOCAL_AA]],
+// CHECK-DAG:   store i16* [[REF_AA]], i16** [[FIELD_AA]],
+// Get capture from context.
+// CHECK:       [[CAP2:%.+]] = load [[CAPTY]]*, [[CAPTY]]** [[CAPREF]]
+// CHECK-DAG:   [[FIELD2_AA:%.+]] = getelementptr inbounds [[CAPTY]], [[CAPTY]]* [[CAP2]], i32 0, i32 0
+// CHECK-DAG:   [[REF2_AA:%.+]] = load i16*, i16** [[FIELD2_AA]]
+// Use captures.
+// CHECK-DAG:   load i16, i16* [[REF2_AA]]
+
+// CHECK:       define internal void [[HVT3]]
+// Create context and reference.
+// CHECK-DAG:   [[CAPREF:%.+]] = alloca [[CAPTY:%[^,].+]]*,
+// CHECK-DAG:   [[CAP:%.+]] = alloca [[CAPTY]],
+// CHECK-DAG:   store [[CAPTY]]* [[CAP]], [[CAPTY]]** [[CAPREF]]
+// Create local storage for each capture.
+// CHECK-DAG:   [[LOCAL_A:%.+]] = alloca i32*
+// CHECK-DAG:   [[LOCAL_AA:%.+]] = alloca i16*
+// CHECK-DAG:   store i32* [[ARG_A:%.+]], i32** [[LOCAL_A]]
+// CHECK-DAG:   store i16* [[ARG_AA:%.+]], i16** [[LOCAL_AA]]
+// Store captures in the context.
+// CHECK-DAG:   [[REF_A:%.+]] = load i32*, i32** [[LOCAL_A]],
+// CHECK-DAG:   [[REF_AA:%.+]] = load i16*, i16** [[LOCAL_AA]],
+// CHECK-DAG:   store i32* [[REF_A]], i32**  [[FIELD_A:%.+]],
+// CHECK-DAG:   store i16* [[REF_AA]], i16** [[FIELD_AA:%.+]],
+// CHECK-DAG:   [[FIELD_A]] = getelementptr inbounds [[CAPTY]], [[CAPTY]]* [[CAP]], i32 0, i32 {{[0-9]}}
+// CHECK-DAG:   [[FIELD_AA]] = getelementptr inbounds [[CAPTY]], [[CAPTY]]* [[CAP]], i32 0, i32 {{[0-9]}}
+// Get capture from context.
+// CHECK:       [[CAP2:%.+]] = load [[CAPTY]]*, [[CAPTY]]** [[CAPREF]]
+// CHECK-DAG:   [[REF2_A:%.+]] = load i32*, i32** [[FIELD2_A:%.+]],
+// CHECK-DAG:   [[REF2_AA:%.+]] = load i16*, i16** [[FIELD2_AA:%.+]],
+// CHECK-DAG:   [[FIELD2_A]] = getelementptr inbounds [[CAPTY]], [[CAPTY]]* [[CAP2]], i32 0, i32 {{[0-9]}}
+// CHECK-DAG:   [[FIELD2_AA]] = getelementptr inbounds [[CAPTY]], [[CAPTY]]* [[CAP2]], i32 0, i32 {{[0-9]}}
+// Use captures.
+// CHECK-DAG:   load i32, i32* [[REF2_A]]
+// CHECK-DAG:   load i16, i16* [[REF2_AA]]
+
+// CHECK:       define internal void [[HVT4]]
+// Create context and reference.
+// CHECK:       [[CAPREF:%.+]] = alloca [[CAPTY:%[^,].+]]*,
+// CHECK:       [[LOCAL_A:%.+]] = alloca i32*
+// CHECK-DAG:   [[CAP:%.+]] = alloca [[CAPTY]],
+// CHECK-DAG:   store [[CAPTY]]* [[CAP]], [[CAPTY]]** [[CAPREF]]
+// Create local storage for each capture.
+// CHECK-DAG:   [[LOCAL_B:%.+]] = alloca [10 x float]*
+// CHECK-DAG:   [[LOCAL_VLA1:%.+]] = alloca i[[SZ]]*
+// CHECK-DAG:   [[LOCAL_BN:%.+]] = alloca float*
+// CHECK-DAG:   [[LOCAL_C:%.+]] = alloca [5 x [10 x double]]*
+// CHECK-DAG:   [[LOCAL_VLA2:%.+]] = alloca i[[SZ]]*
+// CHECK-DAG:   [[LOCAL_VLA3:%.+]] = alloca i[[SZ]]*
+// CHECK-DAG:   [[LOCAL_CN:%.+]] = alloca double*
+// CHECK-DAG:   [[LOCAL_D:%.+]] = alloca [[TT]]*
+// CHECK-DAG:   store i32* [[ARG_A:%.+]], i32** [[LOCAL_A]]
+// CHECK-DAG:   store [10 x float]* [[ARG_B:%.+]], [10 x float]** [[LOCAL_B]]
+// CHECK-DAG:   store i[[SZ]]* [[ARG_VLA1:%.+]], i[[SZ]]** [[LOCAL_VLA1]]
+// CHECK-DAG:   store float* [[ARG_BN:%.+]], float** [[LOCAL_BN]]
+// CHECK-DAG:   store [5 x [10 x double]]* [[ARG_C:%.+]], [5 x [10 x double]]** [[LOCAL_C]]
+// CHECK-DAG:   store i[[SZ]]* [[ARG_VLA2:%.+]], i[[SZ]]** [[LOCAL_VLA2]]
+// CHECK-DAG:   store i[[SZ]]* [[ARG_VLA3:%.+]], i[[SZ]]** [[LOCAL_VLA3]]
+// CHECK-DAG:   store double* [[ARG_CN:%.+]], double** [[LOCAL_CN]]
+// CHECK-DAG:   store [[TT]]* [[ARG_D:%.+]], [[TT]]** [[LOCAL_D]]
+// Store captures in the context.
+// CHECK-DAG:   [[REF_A:%.+]] = load i32*, i32** [[LOCAL_A]],
+// CHECK-DAG:   [[REF_B:%.+]] = load [10 x float]*, [10 x float]** [[LOCAL_B]],
+// CHECK-DAG:   [[REF_VLA1:%.+]] = load i[[SZ]]*, i[[SZ]]** [[LOCAL_VLA1]],
+// CHECK-DAG:   [[VAL_VLA1:%.+]] = load i[[SZ]], i[[SZ]]* [[REF_VLA1]],
+// CHECK-DAG:   [[REF_BN:%.+]] = load float*, float** [[LOCAL_BN]],
+// CHECK-DAG:   [[REF_C:%.+]] = load [5 x [10 x double]]*, [5 x [10 x double]]** [[LOCAL_C]],
+// CHECK-DAG:   [[REF_VLA2:%.+]] = load i[[SZ]]*, i[[SZ]]** [[LOCAL_VLA2]],
+// CHECK-DAG:   [[VAL_VLA2:%.+]] = load i[[SZ]], i[[SZ]]* [[REF_VLA2]],
+// CHECK-DAG:   [[REF_VLA3:%.+]] = load i[[SZ]]*, i[[SZ]]** [[LOCAL_VLA3]],
+// CHECK-DAG:   [[VAL_VLA3:%.+]] = load i[[SZ]], i[[SZ]]* [[REF_VLA3]],
+// CHECK-DAG:   [[REF_CN:%.+]] = load double*, double** [[LOCAL_CN]],
+// CHECK-DAG:   [[REF_D:%.+]] = load [[TT]]*, [[TT]]** [[LOCAL_D]],
+// CHECK-DAG:   store i32* [[REF_A]], i32**  [[FIELD_A:%.+]],
+// CHECK-DAG:   store [10 x float]* [[REF_B]], [10 x float]** [[FIELD_B:%.+]],
+// CHECK-DAG:   store i[[SZ]] [[VAL_VLA1]], i[[SZ]]* [[FIELD_VLA1:%.+]],
+// CHECK-DAG:   store float* [[REF_BN]], float**  [[FIELD_BN:%.+]],
+// CHECK-DAG:   store [5 x [10 x double]]* [[REF_C]], [5 x [10 x double]]** [[FIELD_C:%.+]],
+// CHECK-DAG:   store i[[SZ]] [[VAL_VLA2]], i[[SZ]]*  [[FIELD_VLA2:%.+]],
+// CHECK-DAG:   store i[[SZ]] [[VAL_VLA3]], i[[SZ]]*  [[FIELD_VLA3:%.+]],
+// CHECK-DAG:   store double* [[REF_CN]], double**  [[FIELD_CN:%.+]],
+// CHECK-DAG:   store [[TT]]* [[REF_D]], [[TT]]**  [[FIELD_D:%.+]],
+// CHECK-DAG:   [[FIELD_A]] = getelementptr inbounds [[CAPTY]], [[CAPTY]]* [[CAP]], i32 0, i32 {{[0-9]}}
+// CHECK-DAG:   [[FIELD_B]] = getelementptr inbounds [[CAPTY]], [[CAPTY]]* [[CAP]], i32 0, i32 {{[0-9]}}
+// CHECK-DAG:   [[FIELD_VLA1]] = getelementptr inbounds [[CAPTY]], [[CAPTY]]* [[CAP]], i32 0, i32 {{[0-9]}}
+// CHECK-DAG:   [[FIELD_BN]] = getelementptr inbounds [[CAPTY]], [[CAPTY]]* [[CAP]], i32 0, i32 {{[0-9]}}
+// CHECK-DAG:   [[FIELD_C]] = getelementptr inbounds [[CAPTY]], [[CAPTY]]* [[CAP]], i32 0, i32 {{[0-9]}}
+// CHECK-DAG:   [[FIELD_VLA2]] = getelementptr inbounds [[CAPTY]], [[CAPTY]]* [[CAP]], i32 0, i32 {{[0-9]}}
+// CHECK-DAG:   [[FIELD_VLA3]] = getelementptr inbounds [[CAPTY]], [[CAPTY]]* [[CAP]], i32 0, i32 {{[0-9]}}
+// CHECK-DAG:   [[FIELD_CN]] = getelementptr inbounds [[CAPTY]], [[CAPTY]]* [[CAP]], i32 0, i32 {{[0-9]}}
+// CHECK-DAG:   [[FIELD_D]] = getelementptr inbounds [[CAPTY]], [[CAPTY]]* [[CAP]], i32 0, i32 {{[0-9]}}
+// Get capture from context.
+// CHECK:       [[CAP2:%.+]] = load [[CAPTY]]*, [[CAPTY]]** [[CAPREF]]
+// CHECK-DAG:   [[FIELD2_A:%.+]] = getelementptr inbounds [[CAPTY]], [[CAPTY]]* [[CAP2]], i32 0, i32 0
+// CHECK-DAG:   [[FIELD2_B:%.+]] = getelementptr inbounds [[CAPTY]], [[CAPTY]]* [[CAP2]], i32 0, i32 1
+// CHECK-DAG:   [[FIELD2_VLA1:%.+]] = getelementptr inbounds [[CAPTY]], [[CAPTY]]* [[CAP2]], i32 0, i32 2
+// CHECK-DAG:   [[FIELD2_BN:%.+]] = getelementptr inbounds [[CAPTY]], [[CAPTY]]* [[CAP2]], i32 0, i32 3
+// CHECK-DAG:   [[FIELD2_C:%.+]] = getelementptr inbounds [[CAPTY]], [[CAPTY]]* [[CAP2]], i32 0, i32 4
+// CHECK-DAG:   [[FIELD2_VLA2:%.+]] = getelementptr inbounds [[CAPTY]], [[CAPTY]]* [[CAP2]], i32 0, i32 5
+// CHECK-DAG:   [[FIELD2_VLA3:%.+]] = getelementptr inbounds [[CAPTY]], [[CAPTY]]* [[CAP2]], i32 0, i32 6
+// CHECK-DAG:   [[FIELD2_CN:%.+]] = getelementptr inbounds [[CAPTY]], [[CAPTY]]* [[CAP2]], i32 0, i32 7
+// CHECK-DAG:   [[FIELD2_D:%.+]] = getelementptr inbounds [[CAPTY]], [[CAPTY]]* [[CAP2]], i32 0, i32 8
+// CHECK-DAG:   [[REF2_A:%.+]] = load i32*, i32** [[FIELD2_A]],
+// CHECK-DAG:   [[REF2_B:%.+]] = load [10 x float]*, [10 x float]** [[FIELD2_B]],
+// CHECK-DAG:   [[REF2_VLA1:%.+]] = load i[[SZ]], i[[SZ]]* [[FIELD2_VLA1]],
+// CHECK-DAG:   [[REF2_BN:%.+]] = load float*, float** [[FIELD2_BN]],
+// CHECK-DAG:   [[REF2_C:%.+]] = load [5 x [10 x double]]*, [5 x [10 x double]]** [[FIELD2_C]],
+// CHECK-DAG:   [[REF2_VLA2:%.+]] = load i[[SZ]], i[[SZ]]* [[FIELD2_VLA2]],
+// CHECK-DAG:   [[REF2_VLA3:%.+]] = load i[[SZ]], i[[SZ]]* [[FIELD2_VLA3]],
+// CHECK-DAG:   [[REF2_CN:%.+]] = load double*, double** [[FIELD2_CN]],
+// CHECK-DAG:   [[REF2_D:%.+]] = load [[TT]]*, [[TT]]** [[FIELD2_D]],
+// Use captures.
+// CHECK-DAG:   load i32, i32* [[REF2_A]]
+// CHECK-DAG:   getelementptr inbounds [10 x float], [10 x float]* [[REF2_B]], i32 0, i[[SZ]] 2
+// CHECK-DAG:   getelementptr inbounds float, float* [[REF2_BN]], i[[SZ]] 3
+// CHECK-DAG:   getelementptr inbounds [5 x [10 x double]], [5 x [10 x double]]* [[REF2_C]], i32 0, i[[SZ]] 1
+// CHECK-DAG:   getelementptr inbounds double, double* [[REF2_CN]], i[[SZ]] %{{.+}}
+// CHECK-DAG:   getelementptr inbounds [[TT]], [[TT]]* [[REF2_D]], i32 0, i32 0
+
+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:       store i[[SZ]] 2, i[[SZ]]* [[VLA0:%[^,]+]]
+// CHECK:       store i[[SZ]] [[CELEMSIZE1:%.+]], i[[SZ]]* [[VLA1:%[^,]+]]
+// CHECK:       [[CELEMSIZE2:%.+]] = mul nuw i[[SZ]] 2, [[CELEMSIZE1]]
+// CHECK:       [[CSIZE:%.+]] = mul nuw i[[SZ]] 2, [[CELEMSIZE2]]
+
+// 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** [[BPR:%[^,]+]], i8** [[PR:%[^,]+]], i[[SZ]]* [[SR:%[^,]+]], i32* getelementptr inbounds ([5 x i32], [5 x i32]* [[MAPT7]], i32 0, i32 0))
+// CHECK-DAG:   [[BPR]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BP:%.+]], i32 0, i32 0
+// CHECK-DAG:   [[PR]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[P:%.+]], i32 0, i32 0
+// CHECK-DAG:   [[SR]] = getelementptr inbounds [5 x i[[SZ]]], [5 x i[[SZ]]]* [[S:%.+]], i32 0, i32 0
+// CHECK-DAG:   [[SADDR0:%.+]] = getelementptr inbounds [5 x i[[SZ]]], [5 x i[[SZ]]]* [[S]], i32 [[IDX0:[0-9]+]]
+// CHECK-DAG:   [[BPADDR0:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BP]], i32 [[IDX0]]
+// CHECK-DAG:   [[PADDR0:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[P]], i32 [[IDX0]]
+// CHECK-DAG:   [[SADDR1:%.+]] = getelementptr inbounds [5 x i[[SZ]]], [5 x i[[SZ]]]* [[S]], i32 [[IDX1:[0-9]+]]
+// CHECK-DAG:   [[BPADDR1:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BP]], i32 [[IDX1]]
+// CHECK-DAG:   [[PADDR1:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[P]], i32 [[IDX1]]
+// CHECK-DAG:   [[SADDR2:%.+]] = getelementptr inbounds [5 x i[[SZ]]], [5 x i[[SZ]]]* [[S]], i32 [[IDX2:[0-9]+]]
+// CHECK-DAG:   [[BPADDR2:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BP]], i32 [[IDX2]]
+// CHECK-DAG:   [[PADDR2:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[P]], i32 [[IDX2]]
+// CHECK-DAG:   [[SADDR3:%.+]] = getelementptr inbounds [5 x i[[SZ]]], [5 x i[[SZ]]]* [[S]], i32 [[IDX3:[0-9]+]]
+// CHECK-DAG:   [[BPADDR3:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BP]], i32 [[IDX3]]
+// CHECK-DAG:   [[PADDR3:%.+]] = getelementptr inbounds [5 x i8*], [5 x 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 i[[SZ]]* [[VLA0]] to i8*
+// CHECK-DAG:   [[P0:%[^,]+]] = bitcast i[[SZ]]* [[VLA0]] to i8*
+// CHECK-DAG:   store i8* [[BP0]], i8** {{%[^,]+}}
+// CHECK-DAG:   store i8* [[P0]], i8** {{%[^,]+}}
+// CHECK-DAG:   store i[[SZ]] {{4|8}}, i[[SZ]]* {{%[^,]+}}
+
+// CHECK-DAG:   [[BP1:%[^,]+]] = bitcast i[[SZ]]* [[VLA1]] to i8*
+// CHECK-DAG:   [[P1:%[^,]+]] = bitcast i[[SZ]]* [[VLA1]] to i8*
+// CHECK-DAG:   store i8* [[BP1]], i8** {{%[^,]+}}
+// CHECK-DAG:   store i8* [[P1]], i8** {{%[^,]+}}
+// CHECK-DAG:   store i[[SZ]] {{4|8}}, i[[SZ]]* {{%[^,]+}}
+
+// 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 i[[SZ]] 4, i[[SZ]]* {{%[^,]+}}
+
+// 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 i[[SZ]] 8, i[[SZ]]* {{%[^,]+}}
+
+// 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 i[[SZ]] [[CSIZE]], i[[SZ]]* {{%[^,]+}}
+
+// 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** [[BPR:%[^,]+]], i8** [[PR:%[^,]+]], i[[SZ]]* getelementptr inbounds ([4 x i[[SZ]]], [4 x i[[SZ]]]* [[SIZET6]], i32 0, i32 0), i32* getelementptr inbounds ([4 x i32], [4 x i32]* [[MAPT6]], i32 0, i32 0))
+// CHECK-DAG:   [[BPR]] = getelementptr inbounds [4 x i8*], [4 x i8*]* [[BP:%.+]], i32 0, i32 0
+// CHECK-DAG:   [[PR]] = getelementptr inbounds [4 x i8*], [4 x i8*]* [[P:%.+]], i32 0, i32 0
+
+// CHECK-DAG:   [[BPADDR0:%.+]] = getelementptr inbounds [4 x i8*], [4 x i8*]* [[BP]], i32 0, i32 0
+// CHECK-DAG:   [[PADDR0:%.+]] = getelementptr inbounds [4 x i8*], [4 x i8*]* [[P]], i32 0, i32 0
+// 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:   [[BPADDR1:%.+]] = getelementptr inbounds [4 x i8*], [4 x i8*]* [[BP]], i32 0, i32 1
+// CHECK-DAG:   [[PADDR1:%.+]] = getelementptr inbounds [4 x i8*], [4 x i8*]* [[P]], i32 0, i32 1
+// 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:   [[BPADDR2:%.+]] = getelementptr inbounds [4 x i8*], [4 x i8*]* [[BP]], i32 0, i32 2
+// CHECK-DAG:   [[PADDR2:%.+]] = getelementptr inbounds [4 x i8*], [4 x i8*]* [[P]], i32 0, i32 2
+// CHECK-DAG:   store i8* [[BP2:%[^,]+]], i8** [[BPADDR2]]
+// CHECK-DAG:   store i8* [[P2:%[^,]+]], i8** [[PADDR2]]
+
+// CHECK-DAG:   [[BPADDR3:%.+]] = getelementptr inbounds [4 x i8*], [4 x i8*]* [[BP]], i32 0, i32 3
+// CHECK-DAG:   [[PADDR3:%.+]] = getelementptr inbounds [4 x i8*], [4 x i8*]* [[P]], i32 0, i32 3
+// 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** [[BPR:%[^,]+]], i8** [[PR:%[^,]+]], i[[SZ]]* getelementptr inbounds ([3 x i[[SZ]]], [3 x i[[SZ]]]* [[SIZET5]], i32 0, i32 0), i32* getelementptr inbounds ([3 x i32], [3 x i32]* [[MAPT5]], i32 0, i32 0))
+// CHECK-DAG:   [[BPR]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BP:%.+]], i32 0, i32 0
+// CHECK-DAG:   [[PR]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[P:%.+]], i32 0, i32 0
+
+// CHECK-DAG:   [[BPADDR0:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BP]], i32 0, i32 0
+// CHECK-DAG:   [[PADDR0:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[P]], i32 0, i32 0
+// 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:   [[BPADDR1:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BP]], i32 0, i32 1
+// CHECK-DAG:   [[PADDR1:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[P]], i32 0, i32 1
+// 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:   [[BPADDR2:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BP]], i32 0, i32 2
+// CHECK-DAG:   [[PADDR2:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[P]], i32 0, i32 2
+// 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]]
+// Create context and reference.
+// CHECK:       [[CAPREF:%.+]] = alloca [[CAPTY:%[^,].+]]*,
+
+// Create local storage for each capture.
+// CHECK:       [[LOCAL_THIS:%.+]] = alloca [[S1]]*
+// CHECK:       [[LOCAL_B:%.+]] = alloca i32*
+// CHECK:       [[LOCAL_VLA1:%.+]] = alloca i[[SZ]]*
+// CHECK:       [[LOCAL_VLA2:%.+]] = alloca i[[SZ]]*
+// CHECK:       [[LOCAL_C:%.+]] = alloca i16*
+// CHECK:       [[CAP:%.+]] = alloca [[CAPTY]],
+// CHECK-DAG:   store [[CAPTY]]* [[CAP]], [[CAPTY]]** [[CAPREF]]
+// CHECK-DAG:   store [[S1]]* [[ARG_THIS:%.+]], [[S1]]** [[LOCAL_THIS]]
+// CHECK-DAG:   store i32* [[ARG_B:%.+]], i32** [[LOCAL_B]]
+// CHECK-DAG:   store i[[SZ]]* [[ARG_VLA1:%.+]], i[[SZ]]** [[LOCAL_VLA1]]
+// CHECK-DAG:   store i[[SZ]]* [[ARG_VLA2:%.+]], i[[SZ]]** [[LOCAL_VLA2]]
+// CHECK-DAG:   store i16* [[ARG_C:%.+]], i16** [[LOCAL_C]]
+// Store captures in the context.
+// CHECK-DAG:   [[REF_THIS:%.+]] = load [[S1]]*, [[S1]]** [[LOCAL_THIS]],
+// CHECK-DAG:   [[REF_B:%.+]] = load i32*, i32** [[LOCAL_B]],
+// CHECK-DAG:   [[REF_VLA1:%.+]] = load i[[SZ]]*, i[[SZ]]** [[LOCAL_VLA1]],
+// CHECK-DAG:   [[VAL_VLA1:%.+]] = load i[[SZ]], i[[SZ]]* [[REF_VLA1]],
+// CHECK-DAG:   [[REF_VLA2:%.+]] = load i[[SZ]]*, i[[SZ]]** [[LOCAL_VLA2]],
+// CHECK-DAG:   [[VAL_VLA2:%.+]] = load i[[SZ]], i[[SZ]]* [[REF_VLA2]],
+// CHECK-DAG:   [[REF_C:%.+]] = load i16*, i16** [[LOCAL_C]],
+// CHECK-DAG:   store [[S1]]* [[REF_THIS]], [[S1]]**  [[FIELD_THIS:%.+]],
+// CHECK-DAG:   store i32* [[REF_B]], i32** [[FIELD_B:%.+]],
+// CHECK-DAG:   store i[[SZ]] [[VAL_VLA1]], i[[SZ]]*  [[FIELD_VLA1:%.+]],
+// CHECK-DAG:   store i[[SZ]] [[VAL_VLA2]], i[[SZ]]*  [[FIELD_VLA2:%.+]],
+// CHECK-DAG:   store i16* [[REF_C]], i16**  [[FIELD_C:%.+]],
+// CHECK-DAG:   [[FIELD_THIS]] = getelementptr inbounds [[CAPTY]], [[CAPTY]]* [[CAP]], i32 0, i32 {{[0-9]}}
+// CHECK-DAG:   [[FIELD_B]] = getelementptr inbounds [[CAPTY]], [[CAPTY]]* [[CAP]], i32 0, i32 {{[0-9]}}
+// CHECK-DAG:   [[FIELD_VLA1]] = getelementptr inbounds [[CAPTY]], [[CAPTY]]* [[CAP]], i32 0, i32 {{[0-9]}}
+// CHECK-DAG:   [[FIELD_VLA2]] = getelementptr inbounds [[CAPTY]], [[CAPTY]]* [[CAP]], i32 0, i32 {{[0-9]}}
+// CHECK-DAG:   [[FIELD_C]] = getelementptr inbounds [[CAPTY]], [[CAPTY]]* [[CAP]], i32 0, i32 {{[0-9]}}
+// Get capture from context.
+// CHECK:       [[CAP2:%.+]] = load [[CAPTY]]*, [[CAPTY]]** [[CAPREF]]
+// CHECK-DAG:   [[FIELD2_THIS:%.+]] = getelementptr inbounds [[CAPTY]], [[CAPTY]]* [[CAP2]], i32 0, i32 0
+// CHECK-DAG:   [[FIELD2_B:%.+]] = getelementptr inbounds [[CAPTY]], [[CAPTY]]* [[CAP2]], i32 0, i32 1
+// CHECK-DAG:   [[FIELD2_VLA1:%.+]] = getelementptr inbounds [[CAPTY]], [[CAPTY]]* [[CAP2]], i32 0, i32 2
+// CHECK-DAG:   [[FIELD2_VLA2:%.+]] = getelementptr inbounds [[CAPTY]], [[CAPTY]]* [[CAP2]], i32 0, i32 3
+// CHECK-DAG:   [[FIELD2_C:%.+]] = getelementptr inbounds [[CAPTY]], [[CAPTY]]* [[CAP2]], i32 0, i32 4
+// CHECK-DAG:   [[REF2_THIS:%.+]] = load [[S1]]*, [[S1]]** [[FIELD2_THIS]],
+// CHECK-DAG:   [[REF2_B:%.+]] = load i32*, i32** [[FIELD2_B]],
+// CHECK-DAG:   [[REF2_VLA1:%.+]] = load i[[SZ]], i[[SZ]]* [[FIELD2_VLA1]],
+// CHECK-DAG:   [[REF2_VLA2:%.+]] = load i[[SZ]], i[[SZ]]* [[FIELD2_VLA2]],
+// CHECK-DAG:   [[REF2_C:%.+]] = load i16*, i16** [[FIELD2_C]],
+// Use captures.
+// CHECK-DAG:   getelementptr inbounds [[S1]], [[S1]]* [[REF2_THIS]], i32 0, i32 0
+// CHECK-DAG:   load i32, i32* [[REF2_B]]
+// CHECK-DAG:   getelementptr inbounds i16, i16* [[REF2_C]], i[[SZ]] %{{.+}}
+
+
+// CHECK:       define internal void [[HVT6]]
+// Create context and reference.
+// CHECK-DAG:   [[CAPREF:%.+]] = alloca [[CAPTY:%[^,].+]]*,
+// CHECK-DAG:   [[CAP:%.+]] = alloca [[CAPTY]],
+// CHECK-DAG:   store [[CAPTY]]* [[CAP]], [[CAPTY]]** [[CAPREF]]
+// Create local storage for each capture.
+// CHECK-DAG:   [[LOCAL_A:%.+]] = alloca i32*
+// CHECK-DAG:   [[LOCAL_AA:%.+]] = alloca i16*
+// CHECK-DAG:   [[LOCAL_AAA:%.+]] = alloca i8*
+// CHECK-DAG:   [[LOCAL_B:%.+]] = alloca [10 x i32]*
+// CHECK-DAG:   store i32* [[ARG_A:%.+]], i32** [[LOCAL_A]]
+// CHECK-DAG:   store i16* [[ARG_AA:%.+]], i16** [[LOCAL_AA]]
+// CHECK-DAG:   store i8* [[ARG_AAA:%.+]], i8** [[LOCAL_AAA]]
+// CHECK-DAG:   store [10 x i32]* [[ARG_B:%.+]], [10 x i32]** [[LOCAL_B]]
+// Store captures in the context.
+// CHECK-DAG:   [[REF_A:%.+]] = load i32*, i32** [[LOCAL_A]],
+// CHECK-DAG:   [[REF_AA:%.+]] = load i16*, i16** [[LOCAL_AA]],
+// CHECK-DAG:   [[REF_AAA:%.+]] = load i8*, i8** [[LOCAL_AAA]],
+// CHECK-DAG:   [[REF_B:%.+]] = load [10 x i32]*, [10 x i32]** [[LOCAL_B]],
+// CHECK-DAG:   store i32* [[REF_A]], i32**  [[FIELD_A:%.+]],
+// CHECK-DAG:   store i16* [[REF_AA]], i16** [[FIELD_AA:%.+]],
+// CHECK-DAG:   store i8* [[REF_AAA]], i8** [[FIELD_AAA:%.+]],
+// CHECK-DAG:   store [10 x i32]* [[REF_B]], [10 x i32]** [[FIELD_B:%.+]],
+// CHECK-DAG:   [[FIELD_A]] = getelementptr inbounds [[CAPTY]], [[CAPTY]]* [[CAP]], i32 0, i32 {{[0-9]}}
+// CHECK-DAG:   [[FIELD_AA]] = getelementptr inbounds [[CAPTY]], [[CAPTY]]* [[CAP]], i32 0, i32 {{[0-9]}}
+// CHECK-DAG:   [[FIELD_AAA]] = getelementptr inbounds [[CAPTY]], [[CAPTY]]* [[CAP]], i32 0, i32 {{[0-9]}}
+// CHECK-DAG:   [[FIELD_B]] = getelementptr inbounds [[CAPTY]], [[CAPTY]]* [[CAP]], i32 0, i32 {{[0-9]}}
+// Get capture from context.
+// CHECK:       [[CAP2:%.+]] = load [[CAPTY]]*, [[CAPTY]]** [[CAPREF]]
+// CHECK-DAG:   [[REF2_A:%.+]] = load i32*, i32** [[FIELD2_A:%.+]],
+// CHECK-DAG:   [[REF2_AA:%.+]] = load i16*, i16** [[FIELD2_AA:%.+]],
+// CHECK-DAG:   [[REF2_AAA:%.+]] = load i8*, i8** [[FIELD2_AAA:%.+]],
+// CHECK-DAG:   [[REF2_B:%.+]] = load [10 x i32]*, [10 x i32]** [[FIELD2_B:%.+]],
+// CHECK-DAG:   [[FIELD2_A]] = getelementptr inbounds [[CAPTY]], [[CAPTY]]* [[CAP2]], i32 0, i32 {{[0-9]}}
+// CHECK-DAG:   [[FIELD2_AA]] = getelementptr inbounds [[CAPTY]], [[CAPTY]]* [[CAP2]], i32 0, i32 {{[0-9]}}
+// CHECK-DAG:   [[FIELD2_AAA]] = getelementptr inbounds [[CAPTY]], [[CAPTY]]* [[CAP2]], i32 0, i32 {{[0-9]}}
+// CHECK-DAG:   [[FIELD2_B]] = getelementptr inbounds [[CAPTY]], [[CAPTY]]* [[CAP2]], i32 0, i32 {{[0-9]}}
+// Use captures.
+// CHECK-DAG:   load i32, i32* [[REF2_A]]
+// CHECK-DAG:   load i16, i16* [[REF2_AA]]
+// CHECK-DAG:   load i8, i8* [[REF2_AAA]]
+// CHECK-DAG:   getelementptr inbounds [10 x i32], [10 x i32]* [[REF2_B]], i32 0, i[[SZ]] 2
+
+// CHECK:       define internal void [[HVT5]]
+// Create context and reference.
+// CHECK-DAG:   [[CAPREF:%.+]] = alloca [[CAPTY:%[^,].+]]*,
+// CHECK-DAG:   [[CAP:%.+]] = alloca [[CAPTY]],
+// CHECK-DAG:   store [[CAPTY]]* [[CAP]], [[CAPTY]]** [[CAPREF]]
+// Create local storage for each capture.
+// CHECK-DAG:   [[LOCAL_A:%.+]] = alloca i32*
+// CHECK-DAG:   [[LOCAL_AA:%.+]] = alloca i16*
+// CHECK-DAG:   [[LOCAL_B:%.+]] = alloca [10 x i32]*
+// CHECK-DAG:   store i32* [[ARG_A:%.+]], i32** [[LOCAL_A]]
+// CHECK-DAG:   store i16* [[ARG_AA:%.+]], i16** [[LOCAL_AA]]
+// CHECK-DAG:   store [10 x i32]* [[ARG_B:%.+]], [10 x i32]** [[LOCAL_B]]
+// Store captures in the context.
+// CHECK-DAG:   [[REF_A:%.+]] = load i32*, i32** [[LOCAL_A]],
+// CHECK-DAG:   [[REF_AA:%.+]] = load i16*, i16** [[LOCAL_AA]],
+// CHECK-DAG:   [[REF_B:%.+]] = load [10 x i32]*, [10 x i32]** [[LOCAL_B]],
+// CHECK-DAG:   store i32* [[REF_A]], i32**  [[FIELD_A:%.+]],
+// CHECK-DAG:   store i16* [[REF_AA]], i16** [[FIELD_AA:%.+]],
+// CHECK-DAG:   store [10 x i32]* [[REF_B]], [10 x i32]** [[FIELD_B:%.+]],
+// CHECK-DAG:   [[FIELD_A]] = getelementptr inbounds [[CAPTY]], [[CAPTY]]* [[CAP]], i32 0, i32 {{[0-9]}}
+// CHECK-DAG:   [[FIELD_AA]] = getelementptr inbounds [[CAPTY]], [[CAPTY]]* [[CAP]], i32 0, i32 {{[0-9]}}
+// CHECK-DAG:   [[FIELD_B]] = getelementptr inbounds [[CAPTY]], [[CAPTY]]* [[CAP]], i32 0, i32 {{[0-9]}}
+// Get capture from context.
+// CHECK:       [[CAP2:%.+]] = load [[CAPTY]]*, [[CAPTY]]** [[CAPREF]]
+// CHECK-DAG:   [[REF2_A:%.+]] = load i32*, i32** [[FIELD2_A:%.+]],
+// CHECK-DAG:   [[REF2_AA:%.+]] = load i16*, i16** [[FIELD2_AA:%.+]],
+// CHECK-DAG:   [[REF2_B:%.+]] = load [10 x i32]*, [10 x i32]** [[FIELD2_B:%.+]],
+// CHECK-DAG:   [[FIELD2_A]] = getelementptr inbounds [[CAPTY]], [[CAPTY]]* [[CAP2]], i32 0, i32 {{[0-9]}}
+// CHECK-DAG:   [[FIELD2_AA]] = getelementptr inbounds [[CAPTY]], [[CAPTY]]* [[CAP2]], i32 0, i32 {{[0-9]}}
+// CHECK-DAG:   [[FIELD2_B]] = getelementptr inbounds [[CAPTY]], [[CAPTY]]* [[CAP2]], i32 0, i32 {{[0-9]}}
+// Use captures.
+// CHECK-DAG:   load i32, i32* [[REF2_A]]
+// CHECK-DAG:   load i16, i16* [[REF2_AA]]
+// CHECK-DAG:   getelementptr inbounds [10 x i32], [10 x i32]* [[REF2_B]], i32 0, i[[SZ]] 2
+#endif