[SYCL] Implement OpenCL kernel function generation

[zahiraam]

Summary

All SYCL memory objects shared between host and device (buffers/images, these
objects map to OpenCL buffers and images) must be accessed through special
accessor classes. The "device" side implementation of these classes contain
pointers to the device memory. As there is no way in OpenCL to pass
structures with pointers inside as kernel arguments, all memory objects
shared between host and device must be passed to the kernel as raw
pointers. SYCL also has a special mechanism for passing kernel arguments
from host to the device. In OpenCL kernel arguments are set by calling
clSetKernelArg function for each kernel argument, meanwhile in SYCL all the
kernel arguments are fields of "SYCL kernel function" which can be defined
as a lambda function or a named function object and passed as an argument
to SYCL function for invoking kernels (such as parallel_for or single_task).

To facilitate the mapping of SYCL kernel data members to OpenCL kernel
arguments and overcome OpenCL limitations we added the generation of an
OpenCL kernel function inside the compiler. An OpenCL kernel function
contains the body of the SYCL kernel function, receives OpenCL-like
parameters and additionally does some manipulation to initialize SYCL
kernel data members with these parameters. In some pseudo code the OpenCL
kernel function can look like this:

// SYCL kernel is defined in SYCL headers:
template <typename KernelName, typename KernelType/*, ...*/>
__attribute__((sycl_kernel)) void sycl_kernel_function(KernelType KernelFuncObj) {
  // ...
  KernelFuncObj();
}

// Generated OpenCL kernel function
__kernel KernelName(global int* a) {
  KernelType KernelFuncObj; // Actually kernel function object declaration
  // doesn't have a name in AST.
  // Let the kernel function object have one captured field - accessor A.
  // We need to init it with global pointer from arguments:
  KernelFuncObj.A.__init(a);
  // Body of the SYCL kernel from SYCL headers:
  {
    KernelFuncObj();
  }
}

OpenCL kernel function is generated by the compiler inside the Sema
using AST nodes.

[Fznamznon]

Updated tests using address space attributes added by D71005

[keryell]

clang/lib/Sema/SemaSYCL.cpp

44

Feel free to add more comments in this area up to line 103.

417

After more thought, perhaps the solution proposed by Victor Lomüller during the last SYCL upstreaming meeting about marking accessor classes with some attribute instead of detecting concrete type names is a better generic approach.
I am more convinced now by his argument allowing more experimenting with alternative runtimes, since it looks like it is the only place we detect type names. For example the kernels are marked with an attribute in the runtime instead of concretely detecting the parallel_for() functions and so on.

clang/test/CodeGenSYCL/device-functions.cpp

3

Missing description about the purpose of this test

clang/test/SemaSYCL/fake-accessors.cpp

3

Missing description about the purpose of this test.
I am struggling about understanding what this test is for...
OK, after coming back later, I think I got it. I was confused by the fact that in the kernels you are using both true accessors (A, B & C) *and* some objects with names similar to SYCL accessor.
Is it possible to have some tests without true accessors?

clang/test/SemaSYCL/mangle-kernel.cpp

4

Missing description about the purpose of this test

[Fznamznon]

clang/lib/Sema/SemaSYCL.cpp

417

@Naghasan , what do you think about upstreaming the solution which you proposed?

[bader]

Depends on D89909

Apply improvements developed in out-of-tree repository, which mostly are refactoring of SemaSYCL.cpp file.
Rebase on ToT + D89909 (which only impacts lit test checks).

[bader]

Addressed comment from Ronan in regression tests and synced them with the latest status.

[bader]

clang/test/CodeGenSYCL/device-functions.cpp

3

Done.

clang/test/SemaSYCL/fake-accessors.cpp

3

Added a new test case w/o SYCL accessor and description of the test.

clang/test/SemaSYCL/mangle-kernel.cpp

4

Done.

[ABataev]

clang/include/clang/Sema/Sema.h

13129

ArrayRef<Decl *> getSYCLKernels()

clang/lib/Sema/SemaSYCL.cpp

90

final

115

Make functions static?

117

llvm::find_if

186

Type

[erichkeane]

clang/lib/Sema/SemaSYCL.cpp

45

Isn't there a big rewrite going on downstream of these with sycl_special_class? Why are we trying to upstream this before that happens?

[bader]

clang/lib/Sema/SemaSYCL.cpp

45

Isn't there a big rewrite going on downstream of these with sycl_special_class?

Yes.

Why are we trying to upstream this before that happens?

The downstream work was initiated by this comment: https://reviews.llvm.org/D71016#inline-644645.
This patch was uploaded for review here before refactoring work has started.

[aaron.ballman]

FWIW, it seems like precommit CI is failing with problems in x64 windows > Clang.SemaSYCL::accessors-targets.cpp (which is largely hidden by the spam from the CI pipeline, unfortunately). Also, you should address the tidy warnings.

clang/include/clang/Sema/Sema.h

13124

Is 4 accurate? I would assume that 1 is more accurate as most people aren't going to be using SYCL at all.

13129

Also, this needs a const-correct overload (or surface only the const correct version, if possible).

clang/lib/AST/ASTContext.cpp

11069

clang/lib/Sema/SemaSYCL.cpp

45

So should this patch be abandoned until that refactoring work completes, to avoid churn?

100

Please spell out the type as it's not spelled out in the initialization.

[tschuett]

It feels like you are doing codegen(OpenCL kernel) in Sema. Are OpenCL kernels the only approach.

[bader]

In D71016#3063457, @tschuett wrote:

It feels like you are doing codegen(OpenCL kernel) in Sema. Are OpenCL kernels the only approach.

We need to update the description of the patch to clarify that it applies to other GPU programming models as well. When the patch was uploaded SYCL targeted OpenCL kernels only and today the downstream implementation can target CUDA, HIP and Intel oneAPI Level Zero kernels as well.
SYCL kernel is defined as C++ callable type, but typical GPU kernel interface is a C-like function. In addition to that there might be additional restrictions on passing data from the host system (e.g. image types can be passed as a member of C++ class, etc.). The solution here is emit a wrapper function, which initializes and invokes SYCL callable object.

Does it answer your question or you would like to see changes in addition to the description update?

[tschuett]

In D71016#3063706, @bader wrote:

In D71016#3063457, @tschuett wrote:

It feels like you are doing codegen(OpenCL kernel) in Sema. Are OpenCL kernels the only approach.

We need to update the description of the patch to clarify that it applies to other GPU programming models as well. When the patch was uploaded SYCL targeted OpenCL kernels only and today the downstream implementation can target CUDA, HIP and Intel oneAPI Level Zero kernels as well.
SYCL kernel is defined as C++ callable type, but typical GPU kernel interface is a C-like function. In addition to that there might be additional restrictions on passing data from the host system (e.g. image types can be passed as a member of C++ class, etc.). The solution here is emit a wrapper function, which initializes and invokes SYCL callable object.

Does it answer your question or you would like to see changes in addition to the description update?

Would a codegenSYCL directory help you to separate Sema from code generation?

[lebedev.ri]

Doesn't this make AST non-representable of the reality,
shouldn't the lowering happen in codegen, not in sema?

[bader]

In D71016#3063762, @tschuett wrote:

Would a codegenSYCL directory help you to separate Sema from code generation?

Moving wrapper kernel function generation to CodeGen library make sense to me.

Doesn't this make AST non-representable of the reality,
shouldn't the lowering happen in codegen, not in sema?

I'm not sure I understand what does "make AST non-representable of the reality" mean, but it seems to be the same suggestion as @tschuett proposed.

[zahiraam]

Added in the sycl_special_class attribute. This is still work in progress as I still have a few LIT tests failing and didn't address the issue of the separating sema from codegen work.

clang/include/clang/Basic/Attr.td

 
 def SYCLKernel : InheritableAttr {
   let Spellings = [Clang<"sycl_kernel">];
   let Subjects = SubjectList<[FunctionTmpl]>;
   let LangOpts = [SYCL];
   let Documentation = [SYCLKernelDocs];
 }
 
+def SYCLSpecialClass: InheritableAttr {
+  let Spellings = [Clang<"sycl_special_class">];
+  let Subjects = SubjectList<[CXXRecord]>;
+  let LangOpts = [SYCL];
+  let Documentation = [SYCLSpecialClassDocs];
+}
+
+def SYCLDevice : InheritableAttr {
+  let Spellings = [GNU<"sycl_device">];
+  let Subjects = SubjectList<[Function]>;
+  let LangOpts = [SYCL];
+  let Documentation = [SYCLDeviceDocs];
+}
+
 def C11NoReturn : InheritableAttr {
   let Spellings = [Keyword<"_Noreturn">];
   let Subjects = SubjectList<[Function], ErrorDiag>;
   let SemaHandler = 0;
   let Documentation = [C11NoReturnDocs];
 }
 
 def CXX11NoReturn : InheritableAttr {

clang/include/clang/Basic/AttrDocs.td

- The function must have at least one parameter. The first parameter is
   parameters.
 - The function must return void. The compiler reuses the body of marked functions to
   generate the OpenCL kernel body, and the OpenCL kernel must return ``void``.
 
 The SYCL kernel in the previous code sample meets these expectations.
   }];
 }
 
+def SYCLSpecialClassDocs : Documentation {
+  let Category = DocCatStmt;
+  let Content = [{
+The ``__attribute__((sycl_special_class))`` attribute is used in SYCL
+headers to indicate that a class or a struct needs additional handling when
+it is passed from host to device. Please note that this is an attribute that is
+used for internal implementation and not intended to be used by external users.
+It is used for ``accessor``, ``sampler`` , or ``stream`` classes.
+The types that own this attribute are excluded from device-copyable and other
+type-legalization steps.
+
+.. code-block:: c++
+
+  class __attribute__((sycl_special_class)) accessor {
+    private:
+      void __init() {}
+  };
+  }];
+}
+
 def C11NoReturnDocs : Documentation {
   let Category = DocCatFunction;
   let Content = [{
 A function declared as ``_Noreturn`` shall not return to its caller. The
 compiler will generate a diagnostic for a function declared as ``_Noreturn``
 that appears to be capable of returning to its caller. Despite being a type
 specifier, the ``_Noreturn`` attribute cannot be specified on a function
 pointer type.

 
 This is because, in ``bar``, ``foo`` is resolved prior to template
 instantiation, so the value for ``I`` isn't known (thus, both ``enable_if``
 conditions for ``foo`` fail). However, in ``baz``, ``foo`` is resolved during
 template instantiation, so the value for ``T::number`` is known.
   }];
 }
 
+def SYCLDeviceDocs : Documentation {
+  let Category = DocCatFunction;
+  let Heading = "sycl_device";
+  let Content = [{
+This attribute can only be applied to functions and indicates that the
+function must be treated as a device function and must be emitted even if it has
+no direct uses from other device functions. All ``sycl_device`` function callees
+implicitly inherit this attribute.
+  }];
+}
+
 def DiagnoseIfDocs : Documentation {
   let Category = DocCatFunction;
   let Content = [{
 The ``diagnose_if`` attribute can be placed on function declarations to emit
 warnings or errors at compile-time if calls to the attributed function meet
 certain user-defined criteria. For example:
 
 .. code-block:: c

clang/include/clang/Basic/DiagnosticSemaKinds.td

 // errors of expect.with.probability
 def err_probability_not_constant_float : Error<
    "probability argument to __builtin_expect_with_probability must be constant "
    "floating-point expression">;
 def err_probability_out_of_range : Error<
    "probability argument to __builtin_expect_with_probability is outside the "
    "range [0.0, 1.0]">;
 
+// SYCL-specific diagnostics
+def err_sycl_attribute_internal_function
+    : Error<"%0 attribute cannot be applied to a "
+            "static function or function in an anonymous namespace">;
+def err_sycl_restrict : Error<
+  "SYCL kernel cannot "
+"%select{use a non-const global variable"
+  "|use rtti"
+  "|use a non-const static data variable"
+  "|call a virtual function"
+  "|use exceptions"
+  "|call a recursive function"
+  "|call through a function pointer"
+  "|allocate storage"
+  "|use inline assembly"
+  "|call a dllimport function"
+  "|call a variadic function"
+  "|call an undefined function without SYCL_EXTERNAL attribute"
+  "|use a const static or global variable that is neither zero-initialized "
+  "nor constant-initialized"
+  "}0">;
+
 // TCB warnings
 def err_tcb_conflicting_attributes : Error<
   "attributes '%0(\"%2\")' and '%1(\"%2\")' are mutually exclusive">;
 def warn_tcb_enforcement_violation : Warning<
   "calling %0 is a violation of trusted computing base '%1'">,
   InGroup<DiagGroup<"tcb-enforcement">>;
 
 // RISC-V builtin required extension warning
 def err_riscv_builtin_requires_extension : Error<
   "builtin requires '%0' extension support to be enabled">;
 def err_riscv_builtin_invalid_lmul : Error<
   "LMUL argument must be in the range [0,3] or [5,7]">;
 } // end of sema component.

clang/include/clang/Sema/Sema.h

public:
   /// Describes the reason a calling convention specification was ignored, used
   /// for diagnostics.
   enum class CallingConventionIgnoredReason {
     ForThisTarget = 0,
     VariadicFunction,
     ConstructorDestructor,
     BuiltinFunction
   };
+
+  void checkSYCLDeviceVarDecl(VarDecl *Var);
+
   /// Creates a SemaDiagnosticBuilder that emits the diagnostic if the current
   /// context is "used as device code".
   ///
   /// - If CurLexicalContext is a kernel function or it is known that the
   ///   function will be emitted for the device, emits the diagnostics
   ///   immediately.
   /// - If CurLexicalContext is a function and we are compiling
   ///   for the device, but we don't know that this function will be codegen'ed
public:
   ///   it's never codegen'ed, creates a deferred diagnostic to be emitted if
   ///   and when the caller is codegen'ed, and returns true.
   ///
   /// - Otherwise, returns true without emitting any diagnostics.
   ///
   /// Adds Callee to DeviceCallGraph if we don't know if its caller will be
   /// codegen'ed yet.
   bool checkSYCLDeviceFunction(SourceLocation Loc, FunctionDecl *Callee);
+
+private:
+  /// Contains generated OpenCL kernel functions for SYCL.
+  SmallVector<Decl *, 1> SYCLKernels;
+
+public:
+  void addSYCLKernel(Decl *D) { SYCLKernels.push_back(D); }
+  /// Access to SYCL kernels.
+  //const SmallVectorImpl<Decl *> &getSYCLKernels() { return SYCLKernels; }
+  const ArrayRef<Decl *> getSYCLKernels() { return SYCLKernels; }
+
+  /// Constructs an OpenCL kernel using the KernelCaller function and adds it to
+  /// the SYCL device code.
+  void constructOpenCLKernel(FunctionDecl *KernelCallerFunc, MangleContext &MC);
 };
 
 /// RAII object that enters a new expression evaluation context.
 class EnterExpressionEvaluationContext {
   Sema &Actions;
   bool Entered = true;
 
 public:

clang/lib/AST/ASTContext.cpp

bool ASTContext::DeclMustBeEmitted(const Decl *D) {
   // If this is a member of a class template, we do not need to emit it.
   if (D->getDeclContext()->isDependentContext())
     return false;
 
   // Weak references don't produce any output by themselves.
   if (D->hasAttr<WeakRefAttr>())
     return false;
 
+  // If SYCL, only kernels are required.
+  if (LangOpts.SYCLIsDevice && !D->hasAttr<OpenCLKernelAttr>())
+    return false;
+
   // Aliases and used decls are required.
   if (D->hasAttr<AliasAttr>() || D->hasAttr<UsedAttr>())
     return true;
 
   if (const auto *FD = dyn_cast<FunctionDecl>(D)) {
     // Forward declarations aren't required.
     if (!FD->doesThisDeclarationHaveABody())
       return FD->doesDeclarationForceExternallyVisibleDefinition();

clang/lib/CodeGen/CodeGenModule.cpp

if (auto *DRD = dyn_cast<OMPDeclareReductionDecl>(Global)) {
       return;
     } else if (auto *DMD = dyn_cast<OMPDeclareMapperDecl>(Global)) {
       if (MustBeEmitted(Global))
         EmitOMPDeclareMapper(DMD);
       return;
     }
   }
 
+  if (LangOpts.SYCLIsDevice && Global->hasAttr<OpenCLKernelAttr>() &&
+      MustBeEmitted(Global)) {
+    addDeferredDeclToEmit(GD);
+    return;
+  }
+
   // Ignore declarations, they will be emitted on their first use.
   if (const auto *FD = dyn_cast<FunctionDecl>(Global)) {
     // Forward declarations are emitted lazily on first use.
     if (!FD->doesThisDeclarationHaveABody()) {
       if (!FD->doesDeclarationForceExternallyVisibleDefinition())
         return;
 
       StringRef MangledName = getMangledName(GD);
if (VD->isThisDeclarationADefinition() != VarDecl::Definition &&
       // change linkage, make sure that it's emitted.
       if (Context.getInlineVariableDefinitionKind(VD) ==
           ASTContext::InlineVariableDefinitionKind::Strong)
         GetAddrOfGlobalVar(VD);
       return;
     }
   }
 
+  // clang::ParseAST ensures that we emit the SYCL device functions at the end,
+  // so anything that is a device (or indirectly called) will be handled later.
+  if (LangOpts.SYCLIsDevice && MustBeEmitted(Global)) {
+    addDeferredDeclToEmit(GD);
+    return;
+  }
+
   // Defer code generation to first use when possible, e.g. if this is an inline
   // function. If the global must always be emitted, do it eagerly if possible
   // to benefit from cache locality.
   if (MustBeEmitted(Global) && MayBeEmittedEagerly(Global)) {
     // Emit the definition if it can't be deferred.
     EmitGlobalDefinition(GD);
     return;
   }

clang/lib/Parse/ParseAST.cpp

for (bool AtEOF = P.ParseFirstTopLevelDecl(ADecl); !AtEOF;
         return;
     }
   }
 
   // Process any TopLevelDecls generated by #pragma weak.
   for (Decl *D : S.WeakTopLevelDecls())
     Consumer->HandleTopLevelDecl(DeclGroupRef(D));
 
+  if (S.getLangOpts().SYCLIsDevice)
+    for (Decl *D : S.getSYCLKernels())
+      Consumer->HandleTopLevelDecl(DeclGroupRef(D));
+
   Consumer->HandleTranslationUnit(S.getASTContext());
 
   // Finalize the template instantiation observer chain.
   // FIXME: This (and init.) should be done in the Sema class, but because
   // Sema does not have a reliable "Finalize" function (it has a
   // destructor, but it is not guaranteed to be called ("-disable-free")).
   // So, do the initialization above and do the finalization here:
   finalize(S.TemplateInstCallbacks, S);

clang/lib/Sema/SemaDecl.cpp

if (var->getTypeSourceInfo()->getType()->isBlockPointerType() &&
         !var->hasInit()) {
       Diag(var->getLocation(), diag::err_opencl_invalid_block_declaration)
           << 1 /*Init*/;
       var->setInvalidDecl();
       return;
     }
   }
 
+  if (getLangOpts().SYCLIsDevice)
+    checkSYCLDeviceVarDecl(var);
+
   // In Objective-C, don't allow jumps past the implicit initialization of a
   // local retaining variable.
   if (getLangOpts().ObjC &&
       var->hasLocalStorage()) {
     switch (var->getType().getObjCLifetime()) {
     case Qualifiers::OCL_None:
     case Qualifiers::OCL_ExplicitNone:
     case Qualifiers::OCL_Autoreleasing:

clang/lib/Sema/SemaDeclAttr.cpp

if (MinSizeAttr *MinSize = S.mergeMinSizeAttr(D, AL))
     D->addAttr(MinSize);
 }
 
 static void handleOptimizeNoneAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
   if (OptimizeNoneAttr *Optnone = S.mergeOptimizeNoneAttr(D, AL))
     D->addAttr(Optnone);
 }
 
+static void handleSYCLDeviceAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
+  auto *FD = cast<FunctionDecl>(D);
+  if (!FD->isExternallyVisible()) {
+    S.Diag(AL.getLoc(), diag::err_sycl_attribute_internal_function) << AL;
+    return;
+  }
+
+  handleSimpleAttribute<SYCLDeviceAttr>(S, D, AL);
+}
+
 static void handleConstantAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
   const auto *VD = cast<VarDecl>(D);
   if (VD->hasLocalStorage()) {
     S.Diag(AL.getLoc(), diag::err_cuda_nonstatic_constdev);
     return;
   }
   // constexpr variable may already get an implicit constant attr, which should
   // be replaced by the explicit constant attr.
case ParsedAttr::AT_OptimizeNone:
     handleOptimizeNoneAttr(S, D, AL);
     break;
   case ParsedAttr::AT_EnumExtensibility:
     handleEnumExtensibilityAttr(S, D, AL);
     break;
   case ParsedAttr::AT_SYCLKernel:
     handleSYCLKernelAttr(S, D, AL);
     break;
+  case ParsedAttr::AT_SYCLSpecialClass:
+    handleSimpleAttribute<SYCLSpecialClassAttr>(S, D, AL);
+    break;
+  case ParsedAttr::AT_SYCLDevice:
+    handleSYCLDeviceAttr(S, D, AL);
+    break;
   case ParsedAttr::AT_Format:
     handleFormatAttr(S, D, AL);
     break;
   case ParsedAttr::AT_FormatArg:
     handleFormatArgAttr(S, D, AL);
     break;
   case ParsedAttr::AT_Callback:
     handleCallbackAttr(S, D, AL);

clang/lib/Sema/SemaSYCL.cpp

 //===- SemaSYCL.cpp - Semantic Analysis for SYCL constructs ---------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 // This implements Semantic Analysis for SYCL constructs.
 //===----------------------------------------------------------------------===//
 
+#include "TreeTransform.h"
+#include "clang/AST/AST.h"
 #include "clang/AST/Mangle.h"
+#include "clang/AST/QualTypeNames.h"
+#include "clang/AST/RecordLayout.h"
+#include "clang/AST/RecursiveASTVisitor.h"
+#include "clang/AST/TemplateArgumentVisitor.h"
+#include "clang/AST/TypeVisitor.h"
+#include "clang/Sema/Initialization.h"
 #include "clang/Sema/Sema.h"
 #include "clang/Sema/SemaDiagnostic.h"
 
+#include "llvm/Support/raw_ostream.h"
+
+#include <string>
+
 using namespace clang;
+using namespace std::placeholders;
+
+using ParamDesc = std::tuple<QualType, IdentifierInfo *, TypeSourceInfo *>;
+
+const static std::string InitMethodName = "__init";
+const static std::string FinalizeMethodName = "__finalize";
+
+// anonymous namespace so these don't get linkage.
+namespace {
+
+/// Various utilities.
+class Util {
+public:
+  using DeclContextDesc = std::pair<clang::Decl::Kind, StringRef>;
+
+  /// Checks whether given clang type is a full specialization of a SYCL
+  /// special class.
+  static bool isSyclSpecialType(QualType Ty);
+
+  /// Checks whether given clang type is a full specialization of the SYCL
+  /// half class.
+  static bool isSyclHalfType(const QualType Ty);
+
+  /// Checks whether given clang type is a standard SYCL API class with given
+  /// name.
+  /// \param Ty    the clang type being checked
+  /// \param Name  the class name checked against
+  /// \param Tmpl  whether the class is template instantiation or simple record
+  static bool isSyclType(const QualType &Ty, StringRef Name, bool Tmpl = false);
+
+  // Checks declaration context hierarchy.
+  /// \param DC     the context of the item to be checked.
+  /// \param Scopes the declaration scopes leading from the item context to the
+  ///               translation unit (excluding the latter)
+  static bool matchContext(const DeclContext *DC,
+                           ArrayRef<Util::DeclContextDesc> Scopes);
+
+  /// Checks whether given clang type is declared in the given hierarchy of
+  /// declaration contexts.
+  /// \param Ty         the clang type being checked
+  /// \param Scopes     the declaration scopes leading from the type to the
+  ///     translation unit (excluding the latter)
+  static bool matchQualifiedTypeName(const QualType &Ty,
+                                     ArrayRef<Util::DeclContextDesc> Scopes);
+};
+
+class KernelBodyTransform final : public TreeTransform<KernelBodyTransform> {
+public:
+  KernelBodyTransform(std::pair<DeclaratorDecl *, DeclaratorDecl *> &MPair,
+                      Sema &S)
+      : TreeTransform<KernelBodyTransform>(S), MappingPair(MPair), SemaRef(S) {}
+  bool AlwaysRebuild() { return true; }
+
+  ExprResult TransformDeclRefExpr(DeclRefExpr *DRE) {
+    auto *Ref = dyn_cast<DeclaratorDecl>(DRE->getDecl());
+    if (Ref && Ref == MappingPair.first) {
+      DeclaratorDecl *NewDecl = MappingPair.second;
+      return DeclRefExpr::Create(
+          SemaRef.getASTContext(), DRE->getQualifierLoc(),
+          DRE->getTemplateKeywordLoc(), NewDecl, false, DRE->getNameInfo(),
+          NewDecl->getType(), DRE->getValueKind());
+    }
+    return DRE;
+  }
+
+private:
+  std::pair<DeclaratorDecl *, DeclaratorDecl *> MappingPair;
+  Sema &SemaRef;
+};
+
+/// Return method by name
+static CXXMethodDecl *getMethodByName(const CXXRecordDecl *CRD,
+                                      StringRef MethodName) {
+  CXXMethodDecl *Method;
+  auto It = llvm::find_if(CRD->methods(),
+                         [MethodName](const CXXMethodDecl *Method) {
+                           return Method->getNameAsString() == MethodName;
+                         });
+  Method = (It != CRD->methods().end()) ? *It : nullptr;
+  return Method;
+}
+
+const CXXRecordDecl *getKernelObjectType(FunctionDecl *Caller) {
+  assert(Caller->getNumParams() > 0 && "Insufficient kernel parameters");
+
+  QualType KernelParamTy = Caller->getParamDecl(0)->getType();
+  // In SYCL 2020 kernels are now passed by reference.
+  if (KernelParamTy->isReferenceType())
+    return KernelParamTy->getPointeeCXXRecordDecl();
+
+  // SYCL 1.2.1
+  return KernelParamTy->getAsCXXRecordDecl();
+}
+
+/// Creates a kernel parameter descriptor
+/// \param Src  field declaration to construct name from
+/// \param Ty   the desired parameter type
+/// \return     the constructed descriptor
+ParamDesc makeParamDesc(const FieldDecl *Src, QualType Ty) {
+  ASTContext &Ctx = Src->getASTContext();
+  std::string Name = (Twine("_arg_") + Src->getName()).str();
+  return std::make_tuple(Ty, &Ctx.Idents.get(Name),
+                         Ctx.getTrivialTypeSourceInfo(Ty));
+}
+
+ParamDesc makeParamDesc(ASTContext &Ctx, const CXXBaseSpecifier &Src,
+                        QualType Ty) {
+  // TODO: There is no name for the base available, but duplicate names are
+  // seemingly already possible, so we'll give them all the same name for now.
+  // This only happens with the accessor types.
+  std::string Name = "_arg__base";
+  return std::make_tuple(Ty, &Ctx.Idents.get(Name),
+                         Ctx.getTrivialTypeSourceInfo(Ty));
+}
+
+// The first template argument to the kernel caller function is used to identify
+// the kernel itself.
+QualType calculateKernelNameType(ASTContext &Ctx,
+                                 FunctionDecl *KernelCallerFunc) {
+  const TemplateArgumentList *TAL =
+      KernelCallerFunc->getTemplateSpecializationArgs();
+  assert(TAL && "No template argument info");
+  return TAL->get(0).getAsType().getCanonicalType();
+}
+
+// Gets a name for the OpenCL kernel function, calculated from the first
+// template argument of the kernel caller function.
+std::string constructKernelName(Sema &S, FunctionDecl *KernelCallerFunc,
+                                MangleContext &MC) {
+  QualType KernelNameType =
+      calculateKernelNameType(S.getASTContext(), KernelCallerFunc);
+
+  SmallString<256> Result;
+  llvm::raw_svector_ostream Out(Result);
+
+  MC.mangleTypeName(KernelNameType, Out);
+
+  return std::string(Out.str());
+}
+
+template <typename T> struct bind_param { using type = T; };
+
+template <> struct bind_param<CXXBaseSpecifier &> {
+  using Type = const CXXBaseSpecifier &;
+};
+
+template <> struct bind_param<FieldDecl *&> { using type = FieldDecl *; };
+
+template <> struct bind_param<FieldDecl *const &> { using type = FieldDecl *; };
+
+template <typename T> using bind_param_t = typename bind_param<T>::type;
+
+class KernelObjVisitor {
+  Sema &SemaRef;
+
+  template <typename ParentTy, typename... HandlerTys>
+  void VisitUnionImpl(const CXXRecordDecl *Owner, ParentTy &Parent,
+                      const CXXRecordDecl *Wrapper, HandlerTys &... Handlers) {
+    (void)std::initializer_list<int>{
+        (Handlers.enterUnion(Owner, Parent), 0)...};
+    VisitRecordHelper(Wrapper, Wrapper->fields(), Handlers...);
+    (void)std::initializer_list<int>{
+        (Handlers.leaveUnion(Owner, Parent), 0)...};
+  }
+
+  // These enable handler execution only when previous Handlers succeed.
+  template <typename... Tn>
+  bool handleField(FieldDecl *FD, QualType FDTy, Tn &&... tn) {
+    bool result = true;
+    (void)std::initializer_list<int>{(result = result && tn(FD, FDTy), 0)...};
+    return result;
+  }
+  template <typename... Tn>
+  bool handleField(const CXXBaseSpecifier &BD, QualType BDTy, Tn &&... tn) {
+    bool result = true;
+    std::initializer_list<int>{(result = result && tn(BD, BDTy), 0)...};
+    return result;
+  }
+
+// This definition using std::bind is necessary because of a gcc 7.x bug.
+#define KF_FOR_EACH(FUNC, Item, Qt)                                            \
+  handleField(                                                                 \
+      Item, Qt,                                                                \
+      std::bind(static_cast<bool (std::decay_t<decltype(Handlers)>::*)(        \
+                    bind_param_t<decltype(Item)>, QualType)>(                  \
+                    &std::decay_t<decltype(Handlers)>::FUNC),                  \
+                std::ref(Handlers), _1, _2)...)
+
+  // The following simpler definition works with gcc 8.x and later.
+  //#define KF_FOR_EACH(FUNC) \
+//  handleField(Field, FieldTy, ([&](FieldDecl *FD, QualType FDTy) { \
+//                return Handlers.f(FD, FDTy); \
+//              })...)
+
+  // Parent contains the FieldDecl or CXXBaseSpecifier that was used to enter
+  // the Wrapper structure that we're currently visiting. Owner is the parent
+  // type (which doesn't exist in cases where it is a FieldDecl in the
+  // 'root'), and Wrapper is the current struct being unwrapped.
+  template <typename ParentTy, typename... HandlerTys>
+  void visitComplexRecord(const CXXRecordDecl *Owner, ParentTy &Parent,
+                          const CXXRecordDecl *Wrapper, QualType RecordTy,
+                          HandlerTys &... Handlers) {
+    (void)std::initializer_list<int>{
+        (Handlers.enterStruct(Owner, Parent, RecordTy), 0)...};
+    VisitRecordHelper(Wrapper, Wrapper->bases(), Handlers...);
+    VisitRecordHelper(Wrapper, Wrapper->fields(), Handlers...);
+    (void)std::initializer_list<int>{
+        (Handlers.leaveStruct(Owner, Parent, RecordTy), 0)...};
+  }
+
+  template <typename ParentTy, typename... HandlerTys>
+  void visitSimpleRecord(const CXXRecordDecl *Owner, ParentTy &Parent,
+                         const CXXRecordDecl *Wrapper, QualType RecordTy,
+                         HandlerTys &... Handlers) {
+    (void)std::initializer_list<int>{
+        (Handlers.handleNonDecompStruct(Owner, Parent, RecordTy), 0)...};
+  }
+
+  template <typename ParentTy, typename... HandlerTys>
+  void visitRecord(const CXXRecordDecl *Owner, ParentTy &Parent,
+                   const CXXRecordDecl *Wrapper, QualType RecordTy,
+                   HandlerTys &... Handlers);
+
+  template <typename ParentTy, typename... HandlerTys>
+  void VisitUnion(const CXXRecordDecl *Owner, ParentTy &Parent,
+                  const CXXRecordDecl *Wrapper, HandlerTys &... Handlers);
+
+  template <typename... HandlerTys>
+  void VisitRecordHelper(const CXXRecordDecl *Owner,
+                         clang::CXXRecordDecl::base_class_const_range Range,
+                         HandlerTys &... Handlers) {
+    for (const auto &Base : Range) {
+      QualType BaseTy = Base.getType();
+      // Handle accessor class as base
+      if (Util::isSyclSpecialType(BaseTy)) {
+        (void)std::initializer_list<int>{
+            (Handlers.handleSyclSpecialType(Owner, Base, BaseTy), 0)...};
+      } else
+        // For all other bases, visit the record
+        visitRecord(Owner, Base, BaseTy->getAsCXXRecordDecl(), BaseTy,
+                    Handlers...);
+    }
+  }
+
+  template <typename... HandlerTys>
+  void VisitRecordHelper(const CXXRecordDecl *Owner,
+                         RecordDecl::field_range Range,
+                         HandlerTys &... Handlers) {
+    VisitRecordFields(Owner, Handlers...);
+  }
+
+  // FIXME: Can this be refactored/handled some other way?
+  template <typename ParentTy, typename... HandlerTys>
+  void visitStreamRecord(const CXXRecordDecl *Owner, ParentTy &Parent,
+                         CXXRecordDecl *Wrapper, QualType RecordTy,
+                         HandlerTys &... Handlers) {
+    (void)std::initializer_list<int>{
+        (Handlers.enterStream(Owner, Parent, RecordTy), 0)...};
+    for (const auto &Field : Wrapper->fields()) {
+      QualType FieldTy = Field->getType();
+      // Required to initialize accessors inside streams.
+      if (Util::isSyclAccessorType(FieldTy))
+        KF_FOR_EACH(handleSyclAccessorType, Field, FieldTy);
+    }
+    (void)std::initializer_list<int>{
+        (Handlers.leaveStream(Owner, Parent, RecordTy), 0)...};
+  }
+
+  template <typename... HandlerTys>
+  void visitArrayElementImpl(const CXXRecordDecl *Owner, FieldDecl *ArrayField,
+                             QualType ElementTy, uint64_t Index,
+                             HandlerTys &... Handlers) {
+    (void)std::initializer_list<int>{
+        (Handlers.nextElement(ElementTy, Index), 0)...};
+    visitField(Owner, ArrayField, ElementTy, Handlers...);
+  }
+
+  template <typename... HandlerTys>
+  void visitFirstArrayElement(const CXXRecordDecl *Owner, FieldDecl *ArrayField,
+                              QualType ElementTy, HandlerTys &... Handlers) {
+    visitArrayElementImpl(Owner, ArrayField, ElementTy, 0, Handlers...);
+  }
+
+  template <typename... HandlerTys>
+  void visitNthArrayElement(const CXXRecordDecl *Owner, FieldDecl *ArrayField,
+                            QualType ElementTy, uint64_t Index,
+                            HandlerTys &... Handlers);
+
+  template <typename... HandlerTys>
+  void visitSimpleArray(const CXXRecordDecl *Owner, FieldDecl *Field,
+                        QualType ArrayTy, HandlerTys &... Handlers) {
+    (void)std::initializer_list<int>{
+        (Handlers.handleSimpleArrayType(Field, ArrayTy), 0)...};
+  }
+
+  template <typename... HandlerTys>
+  void visitComplexArray(const CXXRecordDecl *Owner, FieldDecl *Field,
+                         QualType ArrayTy, HandlerTys &... Handlers) {
+    // Array workflow is:
+    // handleArrayType
+    // enterArray
+    // nextElement
+    // VisitField (same as before, note that The FieldDecl is the of array
+    // itself, not the element)
+    // ... repeat per element, opt-out for duplicates.
+    // leaveArray
+
+    if (!KF_FOR_EACH(handleArrayType, Field, ArrayTy))
+      return;
+
+    const ConstantArrayType *CAT =
+        SemaRef.getASTContext().getAsConstantArrayType(ArrayTy);
+    assert(CAT && "Should only be called on constant-size array.");
+    QualType ET = CAT->getElementType();
+    uint64_t ElemCount = CAT->getSize().getZExtValue();
+    assert(ElemCount > 0 && "SYCL prohibits 0 sized arrays");
+
+    (void)std::initializer_list<int>{
+        (Handlers.enterArray(Field, ArrayTy, ET), 0)...};
+
+    visitFirstArrayElement(Owner, Field, ET, Handlers...);
+    for (uint64_t Index = 1; Index < ElemCount; ++Index)
+      visitNthArrayElement(Owner, Field, ET, Index, Handlers...);
+
+    (void)std::initializer_list<int>{
+        (Handlers.leaveArray(Field, ArrayTy, ET), 0)...};
+  }
+
+  template <typename... HandlerTys>
+  void visitArray(const CXXRecordDecl *Owner, FieldDecl *Field,
+                  QualType ArrayTy, HandlerTys &... Handlers);
+
+  template <typename... HandlerTys>
+  void visitField(const CXXRecordDecl *Owner, FieldDecl *Field,
+                  QualType FieldTy, HandlerTys &... Handlers) {
+    if (Util::isSyclSpecialType(FieldTy))
+      KF_FOR_EACH(handleSyclSpecialType, Field, FieldTy);
+    else if (Util::isSyclHalfType(FieldTy))
+      KF_FOR_EACH(handleSyclHalfType, Field, FieldTy);
+    else if (FieldTy->isStructureOrClassType()) {
+      if (KF_FOR_EACH(handleStructType, Field, FieldTy)) {
+        CXXRecordDecl *RD = FieldTy->getAsCXXRecordDecl();
+        visitRecord(Owner, Field, RD, FieldTy, Handlers...);
+      }
+    } else if (FieldTy->isUnionType()) {
+      if (KF_FOR_EACH(handleUnionType, Field, FieldTy)) {
+        CXXRecordDecl *RD = FieldTy->getAsCXXRecordDecl();
+        VisitUnion(Owner, Field, RD, Handlers...);
+      }
+    } else if (FieldTy->isReferenceType())
+      KF_FOR_EACH(handleReferenceType, Field, FieldTy);
+    else if (FieldTy->isPointerType())
+      KF_FOR_EACH(handlePointerType, Field, FieldTy);
+    else if (FieldTy->isArrayType())
+      visitArray(Owner, Field, FieldTy, Handlers...);
+    else if (FieldTy->isScalarType() || FieldTy->isVectorType())
+      KF_FOR_EACH(handleScalarType, Field, FieldTy);
+    else
+      KF_FOR_EACH(handleOtherType, Field, FieldTy);
+  }
+
+public:
+  KernelObjVisitor(Sema &S) : SemaRef(S) {}
+
+  template <typename... HandlerTys>
+  void VisitRecordBases(const CXXRecordDecl *KernelFunctor,
+                        HandlerTys &... Handlers) {
+    VisitRecordHelper(KernelFunctor, KernelFunctor->bases(), Handlers...);
+  }
+
+  // A visitor function that dispatches to functions as defined in
+  // SyclKernelFieldHandler for the purposes of kernel generation.
+  template <typename... HandlerTys>
+  void VisitRecordFields(const CXXRecordDecl *Owner, HandlerTys &... Handlers) {
+    for (const auto Field : Owner->fields())
+      visitField(Owner, Field, Field->getType(), Handlers...);
+  }
+#undef KF_FOR_EACH
+};
+// A base type that the SYCL OpenCL Kernel construction task uses to implement
+// individual tasks.
+class SyclKernelFieldHandlerBase {
+public:
+  static constexpr const bool VisitUnionBody = false;
+  static constexpr const bool VisitNthArrayElement = true;
+  // Opt-in based on whether we should visit inside simple containers (structs,
+  // arrays). All of the 'check' types should likely be true, the int-header,
+  // and kernel decl creation types should not.
+  static constexpr const bool VisitInsideSimpleContainers = true;
+  // Mark these virtual so that we can use override in the implementer classes,
+  // despite virtual dispatch never being used.
+
+  // SYCL special class can be a base class or a field decl, so both must be
+  // handled.
+  virtual bool handleSyclSpecialType(const CXXRecordDecl *,
+                                     const CXXBaseSpecifier &, QualType) {
+    return true;
+  }
+  virtual bool handleSyclSpecialType(FieldDecl *, QualType) { return true; }
+
+
+  virtual bool handleSyclSpecConstantType(FieldDecl *, QualType) {
+    return true;
+  }
+
+  virtual bool handleSyclHalfType(const CXXRecordDecl *,
+                                  const CXXBaseSpecifier &, QualType) {
+    return true;
+  }
+  virtual bool handleSyclHalfType(FieldDecl *, QualType) { return true; }
+  virtual bool handleStructType(FieldDecl *, QualType) { return true; }
+  virtual bool handleUnionType(FieldDecl *, QualType) { return true; }
+  virtual bool handleReferenceType(FieldDecl *, QualType) { return true; }
+  virtual bool handlePointerType(FieldDecl *, QualType) { return true; }
+  virtual bool handleArrayType(FieldDecl *, QualType) { return true; }
+  virtual bool handleScalarType(FieldDecl *, QualType) { return true; }
+  // Most handlers shouldn't be handling this, just the field checker.
+  virtual bool handleOtherType(FieldDecl *, QualType) { return true; }
+
+  // Handle a simple struct that doesn't need to be decomposed, only called on
+  // handlers with VisitInsideSimpleContainers as false.  Replaces
+  // handleStructType, enterStruct, leaveStruct, and visiting of sub-elements.
+  virtual bool handleNonDecompStruct(const CXXRecordDecl *, FieldDecl *,
+                                     QualType) {
+    return true;
+  }
+  virtual bool handleNonDecompStruct(const CXXRecordDecl *,
+                                     const CXXBaseSpecifier &, QualType) {
+    return true;
+  }
+
+  // Instead of handleArrayType, enterArray, leaveArray, and nextElement (plus
+  // descending down the elements), this function gets called in the event of an
+  // array containing simple elements (even in the case of an MD array).
+  virtual bool handleSimpleArrayType(FieldDecl *, QualType) { return true; }
+
+  // The following are only used for keeping track of where we are in the base
+  // class/field graph. Int Headers use this to calculate offset, most others
+  // don't have a need for these.
+
+  virtual bool enterStruct(const CXXRecordDecl *, FieldDecl *, QualType) {
+    return true;
+  }
+  virtual bool leaveStruct(const CXXRecordDecl *, FieldDecl *, QualType) {
+    return true;
+  }
+  virtual bool enterStream(const CXXRecordDecl *, FieldDecl *, QualType) {
+    return true;
+  }
+  virtual bool leaveStream(const CXXRecordDecl *, FieldDecl *, QualType) {
+    return true;
+  }
+  virtual bool enterStruct(const CXXRecordDecl *, const CXXBaseSpecifier &,
+                           QualType) {
+    return true;
+  }
+  virtual bool leaveStruct(const CXXRecordDecl *, const CXXBaseSpecifier &,
+                           QualType) {
+    return true;
+  }
+  virtual bool enterUnion(const CXXRecordDecl *, FieldDecl *) { return true; }
+  virtual bool leaveUnion(const CXXRecordDecl *, FieldDecl *) { return true; }
+
+  // The following are used for stepping through array elements.
+  virtual bool enterArray(FieldDecl *, QualType ArrayTy, QualType ElementTy) {
+    return true;
+  }
+  virtual bool leaveArray(FieldDecl *, QualType ArrayTy, QualType ElementTy) {
+    return true;
+  }
+
+  virtual bool nextElement(QualType, uint64_t) { return true; }
+
+  virtual ~SyclKernelFieldHandlerBase() = default;
+};
+
+// A class to act as the direct base for all the SYCL OpenCL Kernel construction
+// tasks that contains a reference to Sema (and potentially any other
+// universally required data).
+class SyclKernelFieldHandler : public SyclKernelFieldHandlerBase {
+protected:
+  Sema &SemaRef;
+  SyclKernelFieldHandler(Sema &S) : SemaRef(S) {}
+};
+
+// A class to represent the 'do nothing' case for filtering purposes.
+class SyclEmptyHandler final : public SyclKernelFieldHandlerBase {};
+SyclEmptyHandler GlobalEmptyHandler;
+
+template <bool Keep, typename H> struct HandlerFilter;
+template <typename H> struct HandlerFilter<true, H> {
+  H &Handler;
+  HandlerFilter(H &Handler) : Handler(Handler) {}
+};
+template <typename H> struct HandlerFilter<false, H> {
+  SyclEmptyHandler &Handler = GlobalEmptyHandler;
+  HandlerFilter(H &Handler) {}
+};
+
+template <bool B, bool... Rest> struct AnyTrue;
+
+template <bool B> struct AnyTrue<B> { static constexpr bool Value = B; };
+
+template <bool B, bool... Rest> struct AnyTrue {
+  static constexpr bool Value = B || AnyTrue<Rest...>::Value;
+};
+
+template <bool B, bool... Rest> struct AllTrue;
+
+template <bool B> struct AllTrue<B> { static constexpr bool Value = B; };
+
+template <bool B, bool... Rest> struct AllTrue {
+  static constexpr bool Value = B && AllTrue<Rest...>::Value;
+};
+
+template <typename ParentTy, typename... Handlers>
+void KernelObjVisitor::VisitUnion(const CXXRecordDecl *Owner, ParentTy &Parent,
+                                  const CXXRecordDecl *Wrapper,
+                                  Handlers &... handlers) {
+  // Don't continue descending if none of the handlers 'care'. This could be 'if
+  // constexpr' starting in C++17.  Until then, we have to count on the
+  // optimizer to realize "if (false)" is a dead branch.
+  if (AnyTrue<Handlers::VisitUnionBody...>::Value)
+    VisitUnionImpl(
+        Owner, Parent, Wrapper,
+        HandlerFilter<Handlers::VisitUnionBody, Handlers>(handlers).Handler...);
+}
+
+template <typename... Handlers>
+void KernelObjVisitor::visitNthArrayElement(const CXXRecordDecl *Owner,
+                                            FieldDecl *ArrayField,
+                                            QualType ElementTy, uint64_t Index,
+                                            Handlers &... handlers) {
+  // Don't continue descending if none of the handlers 'care'. This could be 'if
+  // constexpr' starting in C++17.  Until then, we have to count on the
+  // optimizer to realize "if (false)" is a dead branch.
+  if (AnyTrue<Handlers::VisitNthArrayElement...>::Value)
+    visitArrayElementImpl(
+        Owner, ArrayField, ElementTy, Index,
+        HandlerFilter<Handlers::VisitNthArrayElement, Handlers>(handlers)
+            .Handler...);
+}
+
+template <typename ParentTy, typename... HandlerTys>
+void KernelObjVisitor::visitRecord(const CXXRecordDecl *Owner, ParentTy &Parent,
+                                   const CXXRecordDecl *Wrapper,
+                                   QualType RecordTy,
+                                   HandlerTys &... Handlers) {
+  RecordDecl *RD = RecordTy->getAsRecordDecl();
+  assert(RD && "should not be null.");
+  // "Simple" Containers are those that do NOT need to be decomposed,
+  // "Complex" containers are those that DO. In the case where the container
+  // does NOT need to be decomposed, we can call VisitSimpleRecord on the
+  // handlers that have opted-out of VisitInsideSimpleContainers. The 'if'
+  // makes sure we only do that if at least 1 has opted out.
+  if (!AllTrue<HandlerTys::VisitInsideSimpleContainers...>::Value)
+    visitSimpleRecord(
+        Owner, Parent, Wrapper, RecordTy,
+        HandlerFilter<!HandlerTys::VisitInsideSimpleContainers, HandlerTys>(
+            Handlers)
+            .Handler...);
+
+  // Even though this is a 'simple' container, some handlers (via
+  // VisitInsideSimpleContainers = true) need to treat it as if it needs
+  // decomposing, so we call VisitComplexRecord iif at least one has.
+  if (AnyTrue<HandlerTys::VisitInsideSimpleContainers...>::Value)
+    visitComplexRecord(
+        Owner, Parent, Wrapper, RecordTy,
+        HandlerFilter<HandlerTys::VisitInsideSimpleContainers, HandlerTys>(
+            Handlers)
+            .Handler...);
+}
+
+template <typename... HandlerTys>
+void KernelObjVisitor::visitArray(const CXXRecordDecl *Owner, FieldDecl *Field,
+                                  QualType ArrayTy, HandlerTys &... Handlers) {
+  if (!AllTrue<HandlerTys::VisitInsideSimpleContainers...>::Value)
+    visitSimpleArray(
+        Owner, Field, ArrayTy,
+        HandlerFilter<!HandlerTys::VisitInsideSimpleContainers, HandlerTys>(
+            Handlers)
+            .Handler...);
+
+  if (AnyTrue<HandlerTys::VisitInsideSimpleContainers...>::Value)
+    visitComplexArray(
+        Owner, Field, ArrayTy,
+        HandlerFilter<HandlerTys::VisitInsideSimpleContainers, HandlerTys>(
+            Handlers)
+            .Handler...);
+}
+
+// A type to Create and own the FunctionDecl for the kernel.
+class SyclKernelDeclCreator : public SyclKernelFieldHandler {
+  FunctionDecl *KernelDecl;
+  llvm::SmallVector<ParmVarDecl *, 8> Params;
+  Sema::ContextRAII FuncContext;
+  // Holds the last handled field's first parameter. This doesn't store an
+  // iterator as push_back invalidates iterators.
+  size_t LastParamIndex = 0;
+  // Keeps track of whether we are currently handling fields inside a struct.
+  int StructDepth = 0;
+
+  void addParam(const FieldDecl *FD, QualType FieldTy) {
+    ParamDesc newParamDesc = makeParamDesc(FD, FieldTy);
+    addParam(newParamDesc, FieldTy);
+  }
+
+  void addParam(const CXXBaseSpecifier &BS, QualType FieldTy) {
+    ParamDesc newParamDesc =
+        makeParamDesc(SemaRef.getASTContext(), BS, FieldTy);
+    addParam(newParamDesc, FieldTy);
+  }
+
+  void addParam(ParamDesc newParamDesc, QualType FieldTy) {
+    // Create a new ParmVarDecl based on the new info.
+    ASTContext &Ctx = SemaRef.getASTContext();
+    auto *NewParam = ParmVarDecl::Create(
+        Ctx, KernelDecl, SourceLocation(), SourceLocation(),
+        std::get<1>(newParamDesc), std::get<0>(newParamDesc),
+        std::get<2>(newParamDesc), SC_None, /*DefArg*/ nullptr);
+    NewParam->setScopeInfo(0, Params.size());
+    NewParam->setIsUsed();
+
+    LastParamIndex = Params.size();
+    Params.push_back(NewParam);
+  }
+
+  // All special SYCL objects must have __init method. We extract types for
+  // kernel parameters from __init method parameters. We will use __init method
+  // and kernel parameters which we build here to initialize special objects in
+  // the kernel body.
+  bool handleSpecialType(FieldDecl *FD, QualType FieldTy,
+                         bool isAccessorType = false) {
+    const auto *RecordDecl = FieldTy->getAsCXXRecordDecl();
+    assert(RecordDecl && "The type must be a RecordDecl");
+    CXXMethodDecl *InitMethod = getMethodByName(RecordDecl, InitMethodName);
+    assert(InitMethod && "The type must have the __init method");
+
+    // Don't do -1 here because we count on this to be the first parameter added
+    // (if any).
+    size_t ParamIndex = Params.size();
+    for (const ParmVarDecl *Param : InitMethod->parameters()) {
+      QualType ParamTy = Param->getType();
+      addParam(FD, ParamTy.getCanonicalType());
+    }
+    LastParamIndex = ParamIndex;
+    return true;
+  }
+
+  static void setKernelImplicitAttrs(ASTContext &Context, FunctionDecl *FD,
+                                     StringRef Name) {
+    // Set implicit attributes.
+    FD->addAttr(OpenCLKernelAttr::CreateImplicit(Context));
+    FD->addAttr(AsmLabelAttr::CreateImplicit(Context, Name));
+    FD->addAttr(ArtificialAttr::CreateImplicit(Context));
+  }
+
+  static FunctionDecl *createKernelDecl(ASTContext &Ctx, StringRef Name,
+                                        SourceLocation Loc, bool IsInline) {
+    // Create this with no prototype, and we can fix this up after we've seen
+    // all the params.
+    FunctionProtoType::ExtProtoInfo Info(CC_OpenCLKernel);
+    QualType FuncType = Ctx.getFunctionType(Ctx.VoidTy, {}, Info);
+
+    FunctionDecl *FD = FunctionDecl::Create(
+        Ctx, Ctx.getTranslationUnitDecl(), Loc, Loc, &Ctx.Idents.get(Name),
+        FuncType, Ctx.getTrivialTypeSourceInfo(Ctx.VoidTy), SC_None);
+    FD->setImplicitlyInline(IsInline);
+    setKernelImplicitAttrs(Ctx, FD, Name);
+
+    // Add kernel to translation unit to see it in AST-dump.
+    Ctx.getTranslationUnitDecl()->addDecl(FD);
+    return FD;
+  }
+
+public:
+  static constexpr const bool VisitInsideSimpleContainers = false;
+  SyclKernelDeclCreator(Sema &S, StringRef Name, SourceLocation Loc,
+                        bool IsInline)
+      : SyclKernelFieldHandler(S),
+        KernelDecl(createKernelDecl(S.getASTContext(), Name, Loc, IsInline)),
+        FuncContext(SemaRef, KernelDecl) {}
+
+  ~SyclKernelDeclCreator() {
+    ASTContext &Ctx = SemaRef.getASTContext();
+    FunctionProtoType::ExtProtoInfo Info(CC_OpenCLKernel);
+
+    SmallVector<QualType, 8> ArgTys;
+    std::transform(std::begin(Params), std::end(Params),
+                   std::back_inserter(ArgTys),
+                   [](const ParmVarDecl *PVD) { return PVD->getType(); });
+
+    QualType FuncType = Ctx.getFunctionType(Ctx.VoidTy, ArgTys, Info);
+    KernelDecl->setType(FuncType);
+    KernelDecl->setParams(Params);
+
+    SemaRef.addSYCLKernel(KernelDecl);
+  }
+
+  bool enterStream(const CXXRecordDecl *RD, FieldDecl *FD, QualType Ty) final {
+    return enterStruct(RD, FD, Ty);
+  }
+
+  bool leaveStream(const CXXRecordDecl *RD, FieldDecl *FD, QualType Ty) final {
+    return leaveStruct(RD, FD, Ty);
+  }
+
+  bool enterStruct(const CXXRecordDecl *, FieldDecl *, QualType) final {
+    ++StructDepth;
+    return true;
+  }
+
+  bool leaveStruct(const CXXRecordDecl *, FieldDecl *, QualType) final {
+    --StructDepth;
+    return true;
+  }
+
+  bool enterStruct(const CXXRecordDecl *, const CXXBaseSpecifier &BS,
+                   QualType FieldTy) final {
+    ++StructDepth;
+    return true;
+  }
+
+  bool leaveStruct(const CXXRecordDecl *, const CXXBaseSpecifier &BS,
+                   QualType FieldTy) final {
+    --StructDepth;
+    return true;
+  }
+
+  bool handleSyclSpecialType(const CXXRecordDecl *, const CXXBaseSpecifier &BS,
+                             QualType FieldTy) final {
+    const auto *RecordDecl = FieldTy->getAsCXXRecordDecl();
+    assert(RecordDecl && "The type must be a RecordDecl");
+    CXXMethodDecl *InitMethod = getMethodByName(RecordDecl, InitMethodName);
+    assert(InitMethod && "The type must have the __init method");
+
+    // Don't do -1 here because we count on this to be the first parameter added
+    // (if any).
+    size_t ParamIndex = Params.size();
+    for (const ParmVarDecl *Param : InitMethod->parameters()) {
+      QualType ParamTy = Param->getType();
+      addParam(BS, ParamTy.getCanonicalType());
+    }
+    LastParamIndex = ParamIndex;
+    return true;
+  }
+
+  bool handleSyclSpecialType(FieldDecl *FD, QualType FieldTy) final {
+    return handleSpecialType(FD, FieldTy);
+  }
+
+  RecordDecl *wrapField(FieldDecl *Field, QualType FieldTy) {
+    RecordDecl *WrapperClass =
+        SemaRef.getASTContext().buildImplicitRecord("__wrapper_class");
+    WrapperClass->startDefinition();
+    Field = FieldDecl::Create(
+        SemaRef.getASTContext(), WrapperClass, SourceLocation(),
+        SourceLocation(), /*Id=*/nullptr, FieldTy,
+        SemaRef.getASTContext().getTrivialTypeSourceInfo(FieldTy,
+                                                         SourceLocation()),
+        /*BW=*/nullptr, /*Mutable=*/false, /*InitStyle=*/ICIS_NoInit);
+    Field->setAccess(AS_public);
+    WrapperClass->addDecl(Field);
+    WrapperClass->completeDefinition();
+    return WrapperClass;
+  };
+
+  bool handlePointerType(FieldDecl *FD, QualType FieldTy) final {
+    // USM allows to use raw pointers instead of buffers/accessors, but these
+    // pointers point to the specially allocated memory. For pointer fields we
+    // add a kernel argument with the same type as field but global address
+    // space, because OpenCL requires it.
+    QualType PointeeTy = FieldTy->getPointeeType();
+    Qualifiers Quals = PointeeTy.getQualifiers();
+    auto AS = Quals.getAddressSpace();
+    // Leave global_device and global_host address spaces as is to help FPGA
+    // device in memory allocations
+    if (AS != LangAS::opencl_global_device && AS != LangAS::opencl_global_host)
+      Quals.setAddressSpace(LangAS::opencl_global);
+    PointeeTy = SemaRef.getASTContext().getQualifiedType(
+        PointeeTy.getUnqualifiedType(), Quals);
+    QualType ModTy = SemaRef.getASTContext().getPointerType(PointeeTy);
+    // When the kernel is generated, struct type kernel arguments are
+    // decomposed; i.e. the parameters of the kernel are the fields of the
+    // struct, and not the struct itself. This causes an error in the backend
+    // when the struct field is a pointer, since non-USM pointers cannot be
+    // passed directly. To work around this issue, all pointers inside the
+    // struct are wrapped in a generated '__wrapper_class'.
+    if (StructDepth) {
+      RecordDecl *WrappedPointer = wrapField(FD, ModTy);
+      ModTy = SemaRef.getASTContext().getRecordType(WrappedPointer);
+    }
+
+    addParam(FD, ModTy);
+    return true;
+  }
+
+  bool handleSimpleArrayType(FieldDecl *FD, QualType FieldTy) final {
+    // Arrays are always wrapped in a struct since they cannot be passed
+    // directly.
+    RecordDecl *WrappedArray = wrapField(FD, FieldTy);
+    QualType ModTy = SemaRef.getASTContext().getRecordType(WrappedArray);
+    addParam(FD, ModTy);
+    return true;
+  }
+
+  bool handleScalarType(FieldDecl *FD, QualType FieldTy) final {
+    addParam(FD, FieldTy);
+    return true;
+  }
+
+  bool handleNonDecompStruct(const CXXRecordDecl *, FieldDecl *FD,
+                             QualType Ty) final {
+    addParam(FD, Ty);
+    return true;
+  }
+
+  bool handleNonDecompStruct(const CXXRecordDecl *Base,
+                             const CXXBaseSpecifier &BS, QualType Ty) final {
+    addParam(BS, Ty);
+    return true;
+  }
+
+  bool handleUnionType(FieldDecl *FD, QualType FieldTy) final {
+    return handleScalarType(FD, FieldTy);
+  }
+
+  bool handleSyclHalfType(FieldDecl *FD, QualType FieldTy) final {
+    addParam(FD, FieldTy);
+    return true;
+  }
+
+  void setBody(CompoundStmt *KB) { KernelDecl->setBody(KB); }
+
+  FunctionDecl *getKernelDecl() { return KernelDecl; }
+
+  llvm::ArrayRef<ParmVarDecl *> getParamVarDeclsForCurrentField() {
+    return ArrayRef<ParmVarDecl *>(std::begin(Params) + LastParamIndex,
+                                   std::end(Params));
+  }
+  using SyclKernelFieldHandler::handleSyclHalfType;
+};
+
+class SyclKernelBodyCreator : public SyclKernelFieldHandler {
+  SyclKernelDeclCreator &DeclCreator;
+  llvm::SmallVector<Stmt *, 16> BodyStmts;
+  llvm::SmallVector<InitListExpr *, 16> CollectionInitExprs;
+  llvm::SmallVector<Stmt *, 16> FinalizeStmts;
+  // This collection contains the information required to add/remove information
+  // about arrays as we enter them.  The InitializedEntity component is
+  // necessary for initializing child members.  uin64_t is the index of the
+  // current element being worked on, which is updated every time we visit
+  // nextElement.
+  llvm::SmallVector<std::pair<InitializedEntity, uint64_t>, 8> ArrayInfos;
+  VarDecl *KernelObjClone;
+  InitializedEntity VarEntity;
+  const CXXRecordDecl *KernelObj;
+  llvm::SmallVector<Expr *, 16> MemberExprBases;
+  FunctionDecl *KernelCallerFunc;
+  SourceLocation KernelCallerSrcLoc; // KernelCallerFunc source location.
+  // Contains a count of how many containers we're in.  This is used by the
+  // pointer-struct-wrapping code to ensure that we don't try to wrap
+  // non-top-level pointers.
+  uint64_t StructDepth = 0;
+
+  // Using the statements/init expressions that we've created, this generates
+  // the kernel body compound stmt. CompoundStmt needs to know its number of
+  // statements in advance to allocate it, so we cannot do this as we go along.
+  CompoundStmt *createKernelBody() {
+    assert(CollectionInitExprs.size() == 1 &&
+           "Should have been popped down to just the first one");
+    KernelObjClone->setInit(CollectionInitExprs.back());
+    Stmt *FunctionBody = KernelCallerFunc->getBody();
+
+    ParmVarDecl *KernelObjParam = *(KernelCallerFunc->param_begin());
+
+    // DeclRefExpr with valid source location but with decl which is not marked
+    // as used is invalid.
+    KernelObjClone->setIsUsed();
+    std::pair<DeclaratorDecl *, DeclaratorDecl *> MappingPair =
+        std::make_pair(KernelObjParam, KernelObjClone);
+
+    // Push the Kernel function scope to ensure the scope isn't empty
+    SemaRef.PushFunctionScope();
+    KernelBodyTransform KBT(MappingPair, SemaRef);
+    Stmt *NewBody = KBT.TransformStmt(FunctionBody).get();
+    BodyStmts.push_back(NewBody);
+
+    BodyStmts.insert(BodyStmts.end(), FinalizeStmts.begin(),
+                     FinalizeStmts.end());
+    return CompoundStmt::Create(SemaRef.getASTContext(), BodyStmts, {}, {});
+  }
+
+  // Creates a DeclRefExpr to the ParmVar that represents the current field.
+  Expr *createParamReferenceExpr() {
+    ParmVarDecl *KernelParameter =
+        DeclCreator.getParamVarDeclsForCurrentField()[0];
+
+    QualType ParamType = KernelParameter->getOriginalType();
+    Expr *DRE = SemaRef.BuildDeclRefExpr(KernelParameter, ParamType, VK_LValue,
+                                         KernelCallerSrcLoc);
+    return DRE;
+  }
+
+  // Creates a DeclRefExpr to the ParmVar that represents the current pointer
+  // field.
+  Expr *createPointerParamReferenceExpr(QualType PointerTy, bool Wrapped) {
+    ParmVarDecl *KernelParameter =
+        DeclCreator.getParamVarDeclsForCurrentField()[0];
+
+    QualType ParamType = KernelParameter->getOriginalType();
+    Expr *DRE = SemaRef.BuildDeclRefExpr(KernelParameter, ParamType, VK_LValue,
+                                         KernelCallerSrcLoc);
+
+    // Struct Type kernel arguments are decomposed. The pointer fields are
+    // then wrapped inside a compiler generated struct. Therefore when
+    // generating the initializers, we have to 'unwrap' the pointer.
+    if (Wrapped) {
+      CXXRecordDecl *WrapperStruct = ParamType->getAsCXXRecordDecl();
+      // Pointer field wrapped inside __wrapper_class
+      FieldDecl *Pointer = *(WrapperStruct->field_begin());
+      DRE = buildMemberExpr(DRE, Pointer);
+      ParamType = Pointer->getType();
+    }
+
+    DRE = ImplicitCastExpr::Create(SemaRef.Context, ParamType,
+                                   CK_LValueToRValue, DRE, /*BasePath=*/nullptr,
+                                   VK_PRValue, FPOptionsOverride());
+
+    if (PointerTy->getPointeeType().getAddressSpace() !=
+        ParamType->getPointeeType().getAddressSpace())
+      DRE = ImplicitCastExpr::Create(SemaRef.Context, PointerTy,
+                                     CK_AddressSpaceConversion, DRE, nullptr,
+                                     VK_PRValue, FPOptionsOverride());
+
+    return DRE;
+  }
+
+  Expr *createSimpleArrayParamReferenceExpr(QualType ArrayTy) {
+    ParmVarDecl *KernelParameter =
+        DeclCreator.getParamVarDeclsForCurrentField()[0];
+    QualType ParamType = KernelParameter->getOriginalType();
+    Expr *DRE = SemaRef.BuildDeclRefExpr(KernelParameter, ParamType, VK_LValue,
+                                         KernelCallerSrcLoc);
+
+    // Unwrap the array.
+    CXXRecordDecl *WrapperStruct = ParamType->getAsCXXRecordDecl();
+    FieldDecl *ArrayField = *(WrapperStruct->field_begin());
+    return buildMemberExpr(DRE, ArrayField);
+  }
+
+  // Returns 'true' if the thing we're visiting (Based on the FD/QualType pair)
+  // is an element of an array.  This will determine whether we do
+  // MemberExprBases in some cases or not, AND determines how we initialize
+  // values.
+  bool isArrayElement(const FieldDecl *FD, QualType Ty) const {
+    return !SemaRef.getASTContext().hasSameType(FD->getType(), Ty);
+  }
+
+  // Creates an initialized entity for a field/item. In the case where this is a
+  // field, returns a normal member initializer, if we're in a sub-array of a MD
+  // array, returns an element initializer.
+  InitializedEntity getFieldEntity(FieldDecl *FD, QualType Ty) {
+    if (isArrayElement(FD, Ty))
+      return InitializedEntity::InitializeElement(SemaRef.getASTContext(),
+                                                  ArrayInfos.back().second,
+                                                  ArrayInfos.back().first);
+    return InitializedEntity::InitializeMember(FD, &VarEntity);
+  }
+
+  void addFieldInit(FieldDecl *FD, QualType Ty, MultiExprArg ParamRef) {
+    InitializationKind InitKind =
+        InitializationKind::CreateCopy(KernelCallerSrcLoc, KernelCallerSrcLoc);
+    addFieldInit(FD, Ty, ParamRef, InitKind);
+  }
+
+  void addFieldInit(FieldDecl *FD, QualType Ty, MultiExprArg ParamRef,
+                    InitializationKind InitKind) {
+    addFieldInit(FD, Ty, ParamRef, InitKind, getFieldEntity(FD, Ty));
+  }
+
+  void addFieldInit(FieldDecl *FD, QualType Ty, MultiExprArg ParamRef,
+                    InitializationKind InitKind, InitializedEntity Entity) {
+    InitializationSequence InitSeq(SemaRef, Entity, InitKind, ParamRef);
+    ExprResult Init = InitSeq.Perform(SemaRef, Entity, InitKind, ParamRef);
+
+    InitListExpr *ParentILE = CollectionInitExprs.back();
+    ParentILE->updateInit(SemaRef.getASTContext(), ParentILE->getNumInits(),
+                          Init.get());
+  }
+
+  void addBaseInit(const CXXBaseSpecifier &BS, QualType Ty,
+                   InitializationKind InitKind) {
+    InitializedEntity Entity = InitializedEntity::InitializeBase(
+        SemaRef.Context, &BS, /*IsInheritedVirtualBase*/ false, &VarEntity);
+    InitializationSequence InitSeq(SemaRef, Entity, InitKind, None);
+    ExprResult Init = InitSeq.Perform(SemaRef, Entity, InitKind, None);
+
+    InitListExpr *ParentILE = CollectionInitExprs.back();
+    ParentILE->updateInit(SemaRef.getASTContext(), ParentILE->getNumInits(),
+                          Init.get());
+  }
+
+  void addSimpleBaseInit(const CXXBaseSpecifier &BS, QualType Ty) {
+    InitializationKind InitKind =
+        InitializationKind::CreateCopy(KernelCallerSrcLoc, KernelCallerSrcLoc);
+
+    InitializedEntity Entity = InitializedEntity::InitializeBase(
+        SemaRef.Context, &BS, /*IsInheritedVirtualBase*/ false, &VarEntity);
+
+    Expr *ParamRef = createParamReferenceExpr();
+    InitializationSequence InitSeq(SemaRef, Entity, InitKind, ParamRef);
+    ExprResult Init = InitSeq.Perform(SemaRef, Entity, InitKind, ParamRef);
+
+    InitListExpr *ParentILE = CollectionInitExprs.back();
+    ParentILE->updateInit(SemaRef.getASTContext(), ParentILE->getNumInits(),
+                          Init.get());
+  }
+
+  // Adds an initializer that handles a simple initialization of a field.
+  void addSimpleFieldInit(FieldDecl *FD, QualType Ty) {
+    Expr *ParamRef = createParamReferenceExpr();
+    addFieldInit(FD, Ty, ParamRef);
+  }
+
+  MemberExpr *buildMemberExpr(Expr *Base, ValueDecl *Member) {
+    DeclAccessPair MemberDAP = DeclAccessPair::make(Member, AS_none);
+    MemberExpr *Result = SemaRef.BuildMemberExpr(
+        Base, /*IsArrow */ false, KernelCallerSrcLoc, NestedNameSpecifierLoc(),
+        KernelCallerSrcLoc, Member, MemberDAP,
+        /*HadMultipleCandidates*/ false,
+        DeclarationNameInfo(Member->getDeclName(), KernelCallerSrcLoc),
+        Member->getType(), VK_LValue, OK_Ordinary);
+    return Result;
+  }
+
+  void addFieldMemberExpr(FieldDecl *FD, QualType Ty) {
+    if (!isArrayElement(FD, Ty))
+      MemberExprBases.push_back(buildMemberExpr(MemberExprBases.back(), FD));
+  }
+
+  void removeFieldMemberExpr(const FieldDecl *FD, QualType Ty) {
+    if (!isArrayElement(FD, Ty))
+      MemberExprBases.pop_back();
+  }
+
+  void createSpecialMethodCall(const CXXRecordDecl *RD, StringRef MethodName,
+                               SmallVectorImpl<Stmt *> &AddTo) {
+    CXXMethodDecl *Method = getMethodByName(RD, MethodName);
+    if (!Method)
+      return;
+
+    unsigned NumParams = Method->getNumParams();
+    llvm::SmallVector<Expr *, 4> ParamDREs(NumParams);
+    llvm::ArrayRef<ParmVarDecl *> KernelParameters =
+        DeclCreator.getParamVarDeclsForCurrentField();
+    for (size_t I = 0; I < NumParams; ++I) {
+      QualType ParamType = KernelParameters[I]->getOriginalType();
+      ParamDREs[I] = SemaRef.BuildDeclRefExpr(KernelParameters[I], ParamType,
+                                              VK_LValue, KernelCallerSrcLoc);
+    }
+
+    MemberExpr *MethodME = buildMemberExpr(MemberExprBases.back(), Method);
+
+    QualType ResultTy = Method->getReturnType();
+    ExprValueKind VK = Expr::getValueKindForType(ResultTy);
+    ResultTy = ResultTy.getNonLValueExprType(SemaRef.Context);
+    llvm::SmallVector<Expr *, 4> ParamStmts;
+    const auto *Proto = cast<FunctionProtoType>(Method->getType());
+    SemaRef.GatherArgumentsForCall(KernelCallerSrcLoc, Method, Proto, 0,
+                                   ParamDREs, ParamStmts);
+    // [kernel_obj or wrapper object].accessor.__init(_ValueType*,
+    // range<int>, range<int>, id<int>)
+    AddTo.push_back(CXXMemberCallExpr::Create(
+        SemaRef.Context, MethodME, ParamStmts, ResultTy, VK, KernelCallerSrcLoc,
+        FPOptionsOverride()));
+  }
+
+  // Creates an empty InitListExpr of the correct number of child-inits
+  // of this to append into.
+  void addCollectionInitListExpr(const CXXRecordDecl *RD) {
+    const ASTRecordLayout &Info =
+        SemaRef.getASTContext().getASTRecordLayout(RD);
+    uint64_t NumInitExprs = Info.getFieldCount() + RD->getNumBases();
+    addCollectionInitListExpr(QualType(RD->getTypeForDecl(), 0), NumInitExprs);
+  }
+
+  InitListExpr *createInitListExpr(const CXXRecordDecl *RD) {
+    const ASTRecordLayout &Info =
+        SemaRef.getASTContext().getASTRecordLayout(RD);
+    uint64_t NumInitExprs = Info.getFieldCount() + RD->getNumBases();
+    return createInitListExpr(QualType(RD->getTypeForDecl(), 0), NumInitExprs);
+  }
+
+  InitListExpr *createInitListExpr(QualType InitTy, uint64_t NumChildInits) {
+    InitListExpr *ILE = new (SemaRef.getASTContext()) InitListExpr(
+        SemaRef.getASTContext(), KernelCallerSrcLoc, {}, KernelCallerSrcLoc);
+    ILE->reserveInits(SemaRef.getASTContext(), NumChildInits);
+    ILE->setType(InitTy);
+
+    return ILE;
+  }
+
+  // Create an empty InitListExpr of the type/size for the rest of the visitor
+  // to append into.
+  void addCollectionInitListExpr(QualType InitTy, uint64_t NumChildInits) {
+
+    InitListExpr *ILE = createInitListExpr(InitTy, NumChildInits);
+    InitListExpr *ParentILE = CollectionInitExprs.back();
+    ParentILE->updateInit(SemaRef.getASTContext(), ParentILE->getNumInits(),
+                          ILE);
+
+    CollectionInitExprs.push_back(ILE);
+  }
+
+  // FIXME Avoid creation of kernel obj clone.
+  // See https://github.com/intel/llvm/issues/1544 for details.
+  static VarDecl *createKernelObjClone(ASTContext &Ctx, DeclContext *DC,
+                                       const CXXRecordDecl *KernelObj) {
+    TypeSourceInfo *TSInfo =
+        KernelObj->isLambda() ? KernelObj->getLambdaTypeInfo() : nullptr;
+    VarDecl *VD = VarDecl::Create(
+        Ctx, DC, KernelObj->getLocation(), KernelObj->getLocation(),
+        KernelObj->getIdentifier(), QualType(KernelObj->getTypeForDecl(), 0),
+        TSInfo, SC_None);
+
+    return VD;
+  }
+
+  const std::string &getInitMethodName() const { return InitMethodName; }
+
+  // Default inits the type, then calls the init-method in the body.
+  bool handleSpecialType(FieldDecl *FD, QualType Ty) {
+    addFieldInit(FD, Ty, None,
+                 InitializationKind::CreateDefault(KernelCallerSrcLoc));
+
+    addFieldMemberExpr(FD, Ty);
+
+    const auto *RecordDecl = Ty->getAsCXXRecordDecl();
+    createSpecialMethodCall(RecordDecl, getInitMethodName(), BodyStmts);
+    // A finalize-method is expected for special type such as stream.
+    CXXMethodDecl *FinalizeMethod =
+        getMethodByName(RecordDecl, FinalizeMethodName);
+    if (FinalizeMethod)
+      createSpecialMethodCall(RecordDecl, FinalizeMethodName, FinalizeStmts);
+
+    removeFieldMemberExpr(FD, Ty);
+
+    return true;
+  }
+
+  bool handleSpecialType(const CXXBaseSpecifier &BS, QualType Ty) {
+    const auto *RecordDecl = Ty->getAsCXXRecordDecl();
+    addBaseInit(BS, Ty, InitializationKind::CreateDefault(KernelCallerSrcLoc));
+    createSpecialMethodCall(RecordDecl, getInitMethodName(), BodyStmts);
+    return true;
+  }
+
+public:
+  static constexpr const bool VisitInsideSimpleContainers = false;
+  SyclKernelBodyCreator(Sema &S, SyclKernelDeclCreator &DC,
+                        const CXXRecordDecl *KernelObj,
+                        FunctionDecl *KernelCallerFunc)
+      : SyclKernelFieldHandler(S), DeclCreator(DC),
+        KernelObjClone(createKernelObjClone(S.getASTContext(),
+                                            DC.getKernelDecl(), KernelObj)),
+        VarEntity(InitializedEntity::InitializeVariable(KernelObjClone)),
+        KernelObj(KernelObj), KernelCallerFunc(KernelCallerFunc),
+        KernelCallerSrcLoc(KernelCallerFunc->getLocation()) {
+    CollectionInitExprs.push_back(createInitListExpr(KernelObj));
+
+    Stmt *DS = new (S.Context) DeclStmt(DeclGroupRef(KernelObjClone),
+                                        KernelCallerSrcLoc, KernelCallerSrcLoc);
+    BodyStmts.push_back(DS);
+    DeclRefExpr *KernelObjCloneRef = DeclRefExpr::Create(
+        S.Context, NestedNameSpecifierLoc(), KernelCallerSrcLoc, KernelObjClone,
+        false, DeclarationNameInfo(), QualType(KernelObj->getTypeForDecl(), 0),
+        VK_LValue);
+    MemberExprBases.push_back(KernelObjCloneRef);
+  }
+
+  ~SyclKernelBodyCreator() {
+    CompoundStmt *KernelBody = createKernelBody();
+    DeclCreator.setBody(KernelBody);
+  }
+
+  bool handleSyclSpecialType(FieldDecl *FD, QualType Ty) final {
+    return handleSpecialType(FD, Ty);
+  }
+
+  bool handleSyclSpecialType(const CXXRecordDecl *, const CXXBaseSpecifier &BS,
+                             QualType Ty) final {
+    return handleSpecialType(BS, Ty);
+  }
+
+  bool handleSyclSpecConstantType(FieldDecl *FD, QualType Ty) final {
+    return handleSpecialType(FD, Ty);
+  }
+
+  bool handleSyclHalfType(FieldDecl *FD, QualType Ty) final {
+    addSimpleFieldInit(FD, Ty);
+    return true;
+  }
+
+  bool handlePointerType(FieldDecl *FD, QualType FieldTy) final {
+    Expr *PointerRef =
+        createPointerParamReferenceExpr(FieldTy, StructDepth != 0);
+    addFieldInit(FD, FieldTy, PointerRef);
+    return true;
+  }
+
+  bool handleSimpleArrayType(FieldDecl *FD, QualType FieldTy) final {
+    Expr *ArrayRef = createSimpleArrayParamReferenceExpr(FieldTy);
+    InitializationKind InitKind = InitializationKind::CreateDirect({}, {}, {});
+
+    InitializedEntity Entity =
+        InitializedEntity::InitializeMember(FD, &VarEntity, /*Implicit*/ true);
+
+    addFieldInit(FD, FieldTy, ArrayRef, InitKind, Entity);
+    return true;
+  }
+
+  bool handleNonDecompStruct(const CXXRecordDecl *, FieldDecl *FD,
+                             QualType Ty) final {
+    addSimpleFieldInit(FD, Ty);
+    return true;
+  }
+
+  bool handleNonDecompStruct(const CXXRecordDecl *Base,
+                             const CXXBaseSpecifier &BS, QualType Ty) final {
+    addSimpleBaseInit(BS, Ty);
+    return true;
+  }
+
+  bool handleScalarType(FieldDecl *FD, QualType FieldTy) final {
+    addSimpleFieldInit(FD, FieldTy);
+    return true;
+  }
+
+  bool handleUnionType(FieldDecl *FD, QualType FieldTy) final {
+    addSimpleFieldInit(FD, FieldTy);
+    return true;
+  }
+
+  bool enterStream(const CXXRecordDecl *RD, FieldDecl *FD, QualType Ty) final {
+    ++StructDepth;
+    // Add a dummy init expression to catch the accessor initializers.
+    const auto *StreamDecl = Ty->getAsCXXRecordDecl();
+    CollectionInitExprs.push_back(createInitListExpr(StreamDecl));
+
+    addFieldMemberExpr(FD, Ty);
+    return true;
+  }
+
+  bool leaveStream(const CXXRecordDecl *RD, FieldDecl *FD, QualType Ty) final {
+    --StructDepth;
+    // Stream requires that its 'init' calls happen after its accessors init
+    // calls, so add them here instead.
+    const auto *StreamDecl = Ty->getAsCXXRecordDecl();
+
+    createSpecialMethodCall(StreamDecl, getInitMethodName(), BodyStmts);
+    createSpecialMethodCall(StreamDecl, FinalizeMethodName, FinalizeStmts);
+
+    removeFieldMemberExpr(FD, Ty);
+
+    CollectionInitExprs.pop_back();
+    return true;
+  }
+
+  bool enterStruct(const CXXRecordDecl *RD, FieldDecl *FD, QualType Ty) final {
+    ++StructDepth;
+    addCollectionInitListExpr(Ty->getAsCXXRecordDecl());
+
+    addFieldMemberExpr(FD, Ty);
+    return true;
+  }
+
+  bool leaveStruct(const CXXRecordDecl *, FieldDecl *FD, QualType Ty) final {
+    --StructDepth;
+    CollectionInitExprs.pop_back();
+
+    removeFieldMemberExpr(FD, Ty);
+    return true;
+  }
+
+  bool enterStruct(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS,
+                   QualType) final {
+    ++StructDepth;
+
+    CXXCastPath BasePath;
+    QualType DerivedTy(RD->getTypeForDecl(), 0);
+    QualType BaseTy = BS.getType();
+    SemaRef.CheckDerivedToBaseConversion(DerivedTy, BaseTy, KernelCallerSrcLoc,
+                                         SourceRange(), &BasePath,
+                                         /*IgnoreBaseAccess*/ true);
+    auto *Cast = ImplicitCastExpr::Create(
+        SemaRef.Context, BaseTy, CK_DerivedToBase, MemberExprBases.back(),
+        /* CXXCastPath=*/&BasePath, VK_LValue, FPOptionsOverride());
+    MemberExprBases.push_back(Cast);
+
+    addCollectionInitListExpr(BaseTy->getAsCXXRecordDecl());
+    return true;
+  }
+
+  bool leaveStruct(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS,
+                   QualType) final {
+    --StructDepth;
+    MemberExprBases.pop_back();
+    CollectionInitExprs.pop_back();
+    return true;
+  }
+
+  bool enterArray(FieldDecl *FD, QualType ArrayType,
+                  QualType ElementType) final {
+    uint64_t ArraySize = SemaRef.getASTContext()
+                             .getAsConstantArrayType(ArrayType)
+                             ->getSize()
+                             .getZExtValue();
+    addCollectionInitListExpr(ArrayType, ArraySize);
+    ArrayInfos.emplace_back(getFieldEntity(FD, ArrayType), 0);
+
+    // If this is the top-level array, we need to make a MemberExpr in addition
+    // to an array subscript.
+    addFieldMemberExpr(FD, ArrayType);
+    return true;
+  }
+
+  bool nextElement(QualType, uint64_t Index) final {
+    ArrayInfos.back().second = Index;
+
+    // Pop off the last member expr base.
+    if (Index != 0)
+      MemberExprBases.pop_back();
+
+    QualType SizeT = SemaRef.getASTContext().getSizeType();
+
+    llvm::APInt IndexVal{
+        static_cast<unsigned>(SemaRef.getASTContext().getTypeSize(SizeT)),
+        Index, SizeT->isSignedIntegerType()};
+
+    auto *IndexLiteral = IntegerLiteral::Create(
+        SemaRef.getASTContext(), IndexVal, SizeT, KernelCallerSrcLoc);
+
+    ExprResult IndexExpr = SemaRef.CreateBuiltinArraySubscriptExpr(
+        MemberExprBases.back(), KernelCallerSrcLoc, IndexLiteral,
+        KernelCallerSrcLoc);
+
+    assert(!IndexExpr.isInvalid());
+    MemberExprBases.push_back(IndexExpr.get());
+    return true;
+  }
+
+  bool leaveArray(FieldDecl *FD, QualType ArrayType,
+                  QualType ElementType) final {
+    CollectionInitExprs.pop_back();
+    ArrayInfos.pop_back();
+
+    assert(
+        !SemaRef.getASTContext().getAsConstantArrayType(ArrayType)->getSize() ==
+            0 &&
+        "Constant arrays must have at least 1 element");
+    // Remove the IndexExpr.
+    MemberExprBases.pop_back();
+
+    // Remove the field access expr as well.
+    removeFieldMemberExpr(FD, ArrayType);
+    return true;
+  }
+
+  using SyclKernelFieldHandler::handleSyclHalfType;
+};
+
+} // anonymous namespace
+
+static bool isZeroSizedArray(QualType Ty) {
+  if (const auto *CATy = dyn_cast<ConstantArrayType>(Ty))
+    return CATy->getSize() == 0;
+  return false;
+}
+
+static void checkSYCLType(Sema &S, QualType Ty, SourceRange Loc,
+                          llvm::DenseSet<QualType> Visited,
+                          SourceRange UsedAtLoc = SourceRange()) {
+  // Not all variable types are supported inside SYCL kernels,
+  // for example the quad type __float128 will cause errors in the
+  // SPIR-V translation phase.
+  // Here we check any potentially unsupported declaration and issue
+  // a deferred diagnostic, which will be emitted iff the declaration
+  // is discovered to reside in kernel code.
+  // The optional UsedAtLoc param is used when the SYCL usage is at a
+  // different location than the variable declaration and we need to
+  // inform the user of both, e.g. struct member usage vs declaration.
+
+  bool Emitting = false;
+
+  //--- check types ---
+
+  // zero length arrays
+  if (isZeroSizedArray(Ty)) {
+    S.SYCLDiagIfDeviceCode(Loc.getBegin(), diag::err_typecheck_zero_array_size);
+    Emitting = true;
+  }
+
+  // variable length arrays
+  if (Ty->isVariableArrayType()) {
+    S.SYCLDiagIfDeviceCode(Loc.getBegin(), diag::err_vla_unsupported);
+    Emitting = true;
+  }
+
+  // Sub-reference array or pointer, then proceed with that type.
+  while (Ty->isAnyPointerType() || Ty->isArrayType())
+    Ty = QualType{Ty->getPointeeOrArrayElementType(), 0};
+
+  if (Emitting && UsedAtLoc.isValid())
+    S.SYCLDiagIfDeviceCode(UsedAtLoc.getBegin(), diag::note_used_here);
+
+  //--- now recurse ---
+  // Pointers complicate recursion. Add this type to Visited.
+  // If already there, bail out.
+  if (!Visited.insert(Ty).second)
+    return;
+
+  if (const auto *ATy = dyn_cast<AttributedType>(Ty))
+    return checkSYCLType(S, ATy->getModifiedType(), Loc, Visited);
+
+  if (const auto *RD = Ty->getAsRecordDecl()) {
+    for (const auto &Field : RD->fields())
+      checkSYCLType(S, Field->getType(), Field->getSourceRange(), Visited, Loc);
+  } else if (const auto *FPTy = dyn_cast<FunctionProtoType>(Ty)) {
+    for (const auto &ParamTy : FPTy->param_types())
+      checkSYCLType(S, ParamTy, Loc, Visited);
+    checkSYCLType(S, FPTy->getReturnType(), Loc, Visited);
+  }
+}
+
+void Sema::checkSYCLDeviceVarDecl(VarDecl *Var) {
+  assert(getLangOpts().SYCLIsDevice &&
+         "Should only be called during SYCL compilation");
+  QualType Ty = Var->getType();
+  SourceRange Loc = Var->getLocation();
+  llvm::DenseSet<QualType> Visited;
+
+  checkSYCLType(*this, Ty, Loc, Visited);
+}
+
+// Generates the OpenCL kernel using KernelCallerFunc (kernel caller
+// function) defined is SYCL headers.
+// Generated OpenCL kernel contains the body of the kernel caller function,
+// receives OpenCL like parameters and additionally does some manipulation to
+// initialize captured lambda/functor fields with these parameters.
+// SYCL runtime marks kernel caller function with sycl_kernel attribute.
+// To be able to generate OpenCL kernel from KernelCallerFunc we put
+// the following requirements to the function which SYCL runtime can mark with
+// sycl_kernel attribute:
+//   - Must be template function with at least two template parameters.
+//     First parameter must represent "unique kernel name"
+//     Second parameter must be the function object type
+//   - Must have only one function parameter - function object.
+//
+// Example of kernel caller function:
+//   template <typename KernelName, typename KernelType/*, ...*/>
+//   __attribute__((sycl_kernel)) void kernel_caller_function(KernelType
+//                                                            KernelFuncObj) {
+//     KernelFuncObj();
+//   }
+//
+//
+void Sema::constructOpenCLKernel(FunctionDecl *KernelCallerFunc,
+                                 MangleContext &MC) {
+  // The first argument to the KernelCallerFunc is the lambda object.
+  const CXXRecordDecl *KernelObj = getKernelObjectType(KernelCallerFunc);
+  assert(KernelObj && "invalid kernel caller");
+
+  // Do not visit invalid kernel object.
+  if (KernelObj->isInvalidDecl())
+    return;
+
+  std::string KernelName = constructKernelName(*this, KernelCallerFunc, MC);
+  SyclKernelDeclCreator kernel_decl(*this, KernelName, KernelObj->getLocation(),
+                                    KernelCallerFunc->isInlined());
+  SyclKernelBodyCreator kernel_body(*this, kernel_decl, KernelObj,
+                                    KernelCallerFunc);
+
+  KernelObjVisitor Visitor{*this};
+  Visitor.VisitRecordBases(KernelObj, kernel_decl, kernel_body);
+  Visitor.VisitRecordFields(KernelObj, kernel_decl, kernel_body);
+}
 
 // -----------------------------------------------------------------------------
 // SYCL device specific diagnostics implementation
 // -----------------------------------------------------------------------------
 
 Sema::SemaDiagnosticBuilder Sema::SYCLDiagIfDeviceCode(SourceLocation Loc,
                                                        unsigned DiagID) {
   assert(getLangOpts().SYCLIsDevice &&
Sema::SemaDiagnosticBuilder Sema::SYCLDiagIfDeviceCode(SourceLocation Loc,
   return SemaDiagnosticBuilder(DiagKind, Loc, DiagID, FD, *this);
 }
 
 bool Sema::checkSYCLDeviceFunction(SourceLocation Loc, FunctionDecl *Callee) {
   assert(getLangOpts().SYCLIsDevice &&
          "Should only be called during SYCL compilation");
   assert(Callee && "Callee may not be null.");
 
-  // Errors in an unevaluated context don't need to be generated,
+  // Errors in unevaluated context don't need to be generated,
   // so we can safely skip them.
   if (isUnevaluatedContext() || isConstantEvaluated())
     return true;
 
+  FunctionDecl *Caller = dyn_cast<FunctionDecl>(getCurLexicalContext());
+
+  if (!Caller)
+    return true;
+
   SemaDiagnosticBuilder::Kind DiagKind = SemaDiagnosticBuilder::K_Nop;
 
   return DiagKind != SemaDiagnosticBuilder::K_Immediate &&
          DiagKind != SemaDiagnosticBuilder::K_ImmediateWithCallStack;
 }
+
+// -----------------------------------------------------------------------------
+// Utility class methods
+// -----------------------------------------------------------------------------
+bool Util::isSyclSpecialType(const QualType Ty) {
+  const CXXRecordDecl *RecTy = Ty->getAsCXXRecordDecl();
+  if (!RecTy)
+    return false;
+  return RecTy->hasAttr<SYCLSpecialClassAttr>();
+}
+
+bool Util::isSyclHalfType(const QualType Ty) {
+  const StringRef &Name = "half";
+  std::array<DeclContextDesc, 5> Scopes = {
+      Util::DeclContextDesc{clang::Decl::Kind::Namespace, "cl"},
+      Util::DeclContextDesc{clang::Decl::Kind::Namespace, "sycl"},
+      Util::DeclContextDesc{clang::Decl::Kind::Namespace, "detail"},
+      Util::DeclContextDesc{clang::Decl::Kind::Namespace, "half_impl"},
+      Util::DeclContextDesc{Decl::Kind::CXXRecord, Name}};
+  return matchQualifiedTypeName(Ty, Scopes);
+}
+
+bool Util::isSyclType(const QualType &Ty, StringRef Name, bool Tmpl) {
+  Decl::Kind ClassDeclKind =
+      Tmpl ? Decl::Kind::ClassTemplateSpecialization : Decl::Kind::CXXRecord;
+  std::array<DeclContextDesc, 3> Scopes = {
+      Util::DeclContextDesc{clang::Decl::Kind::Namespace, "cl"},
+      Util::DeclContextDesc{clang::Decl::Kind::Namespace, "sycl"},
+      Util::DeclContextDesc{ClassDeclKind, Name}};
+  return matchQualifiedTypeName(Ty, Scopes);
+}
+
+bool Util::matchContext(const DeclContext *Ctx,
+                        ArrayRef<Util::DeclContextDesc> Scopes) {
+  // The idea: check the declaration context chain starting from the item
+  // itself. At each step check the context is of expected kind
+  // (namespace) and name.
+  StringRef Name = "";
+
+  for (const auto &Scope : llvm::reverse(Scopes)) {
+    clang::Decl::Kind DK = Ctx->getDeclKind();
+    if (DK != Scope.first)
+      return false;
+
+    switch (DK) {
+    case clang::Decl::Kind::ClassTemplateSpecialization:
+      // ClassTemplateSpecializationDecl inherits from CXXRecordDecl
+    case clang::Decl::Kind::CXXRecord:
+      Name = cast<CXXRecordDecl>(Ctx)->getName();
+      break;
+    case clang::Decl::Kind::Namespace:
+      Name = cast<NamespaceDecl>(Ctx)->getName();
+      break;
+    default:
+      llvm_unreachable("matchContext: decl kind not supported");
+    }
+    if (Name != Scope.second)
+      return false;
+    Ctx = Ctx->getParent();
+  }
+  return Ctx->isTranslationUnit();
+}
+
+bool Util::matchQualifiedTypeName(const QualType &Ty,
+                                  ArrayRef<Util::DeclContextDesc> Scopes) {
+  const CXXRecordDecl *RecTy = Ty->getAsCXXRecordDecl();
+
+  if (!RecTy)
+    return false; // only classes/structs supported
+  const auto *Ctx = cast<DeclContext>(RecTy);
+  return Util::matchContext(Ctx, Scopes);
+}
+

clang/lib/Sema/SemaTemplateInstantiateDecl.cpp

 #include "clang/AST/ASTConsumer.h"
 #include "clang/AST/ASTContext.h"
 #include "clang/AST/ASTMutationListener.h"
 #include "clang/AST/DeclTemplate.h"
 #include "clang/AST/DeclVisitor.h"
 #include "clang/AST/DependentDiagnostic.h"
 #include "clang/AST/Expr.h"
 #include "clang/AST/ExprCXX.h"
+#include "clang/AST/Mangle.h"
 #include "clang/AST/PrettyDeclStackTrace.h"
 #include "clang/AST/TypeLoc.h"
 #include "clang/Basic/SourceManager.h"
 #include "clang/Basic/TargetInfo.h"
 #include "clang/Sema/Initialization.h"
 #include "clang/Sema/Lookup.h"
 #include "clang/Sema/ScopeInfo.h"
 #include "clang/Sema/SemaInternal.h"
if (ParentDC != D->getDeclContext()) {
     }
 
     D = Result;
   }
 
   return D;
 }
 
+static void processFunctionInstantiation(Sema &S,
+                                         SourceLocation PointOfInstantiation,
+                                         FunctionDecl *FD,
+                                         bool DefinitionRequired,
+                                         MangleContext &MC) {
+  S.InstantiateFunctionDefinition(/*FIXME:*/ PointOfInstantiation, FD, true,
+                                  DefinitionRequired, true);
+  if (!FD->isDefined())
+    return;
+  if (FD->hasAttr<SYCLKernelAttr>())
+    S.constructOpenCLKernel(FD, MC);
+  FD->setInstantiationIsPending(false);
+}
+
 /// Performs template instantiation for all implicit template
 /// instantiations we have seen until this point.
 void Sema::PerformPendingInstantiations(bool LocalOnly) {
+  std::unique_ptr<MangleContext> MangleCtx(
+      getASTContext().createMangleContext());
   std::deque<PendingImplicitInstantiation> delayedPCHInstantiations;
   while (!PendingLocalImplicitInstantiations.empty() ||
          (!LocalOnly && !PendingInstantiations.empty())) {
     PendingImplicitInstantiation Inst;
 
     if (PendingLocalImplicitInstantiations.empty()) {
       Inst = PendingInstantiations.front();
       PendingInstantiations.pop_front();
     } else {
       Inst = PendingLocalImplicitInstantiations.front();
       PendingLocalImplicitInstantiations.pop_front();
     }
 
     // Instantiate function definitions
     if (FunctionDecl *Function = dyn_cast<FunctionDecl>(Inst.first)) {
       bool DefinitionRequired = Function->getTemplateSpecializationKind() ==
                                 TSK_ExplicitInstantiationDefinition;
-      if (Function->isMultiVersion()) {
+      if (Function->isMultiVersion())
         getASTContext().forEachMultiversionedFunctionVersion(
-            Function, [this, Inst, DefinitionRequired](FunctionDecl *CurFD) {
-              InstantiateFunctionDefinition(/*FIXME:*/ Inst.second, CurFD, true,
-                                            DefinitionRequired, true);
-              if (CurFD->isDefined())
-                CurFD->setInstantiationIsPending(false);
+            Function, [this, Inst, DefinitionRequired,
+                       MangleCtx = move(MangleCtx)](FunctionDecl *CurFD) {
+              processFunctionInstantiation(*this, Inst.second, CurFD,
+                                           DefinitionRequired, *MangleCtx);
             });
-      } else {
-        InstantiateFunctionDefinition(/*FIXME:*/ Inst.second, Function, true,
-                                      DefinitionRequired, true);
-        if (Function->isDefined())
-          Function->setInstantiationIsPending(false);
-      }
+      else
+        processFunctionInstantiation(*this, Inst.second, Function,
+                                     DefinitionRequired, *MangleCtx);
       // Definition of a PCH-ed template declaration may be available only in the TU.
       if (!LocalOnly && LangOpts.PCHInstantiateTemplates &&
           TUKind == TU_Prefix && Function->instantiationIsPending())
         delayedPCHInstantiations.push_back(Inst);
       continue;
     }
 
     // Instantiate variable definitions

clang/lib/Sema/SemaType.cpp

if (!R.isUsable()) {
             << toString(ConstVal, 10) << ArraySize->getSourceRange();
         return QualType();
       }
 
       T = Context.getConstantArrayType(T, ConstVal, ArraySize, ASM, Quals);
     }
   }
 
-  if (T->isVariableArrayType() && !Context.getTargetInfo().isVLASupported()) {
+  if (T->isVariableArrayType() && !Context.getTargetInfo().isVLASupported() &&
+      !getLangOpts().SYCLIsDevice) {
     // CUDA device code and some other targets don't support VLAs.
     targetDiag(Loc, (getLangOpts().CUDA && getLangOpts().CUDAIsDevice)
                         ? diag::err_cuda_vla
                         : diag::err_vla_unsupported)
         << ((getLangOpts().CUDA && getLangOpts().CUDAIsDevice)
                 ? CurrentCUDATarget()
                 : CFT_InvalidTarget);
   }

clang/test/CodeGenSYCL/Inputs/sycl.hpp

+#pragma once
+
+namespace cl {
+namespace sycl {
+namespace access {
+
+enum class target {
+  global_buffer = 2014,
+  constant_buffer,
+  local,
+  image,
+  host_buffer,
+  host_image,
+  image_array
+};
+
+enum class mode {
+  read = 1024,
+  write,
+  read_write,
+  discard_write,
+  discard_read_write,
+  atomic
+};
+
+enum class placeholder {
+  false_t,
+  true_t
+};
+
+enum class address_space : int {
+  private_space = 0,
+  global_space,
+  constant_space,
+  local_space
+};
+} // namespace access
+
+template <int dim>
+struct id {
+  template <typename... T>
+  id(T... args) {} // fake constructor
+private:
+  // Some fake field added to see using of id arguments in the
+  // kernel wrapper
+  int Data;
+};
+
+template <int dim>
+struct range {
+  template <typename... T>
+  range(T... args) {} // fake constructor
+private:
+  // Some fake field added to see using of range arguments in the
+  // kernel wrapper
+  int Data;
+};
+
+template <int dim>
+struct _ImplT {
+  range<dim> AccessRange;
+  range<dim> MemRange;
+  id<dim> Offset;
+};
+
+template <typename dataT, int dimensions, access::mode accessmode,
+          access::target accessTarget = access::target::global_buffer,
+          access::placeholder isPlaceholder = access::placeholder::false_t>
+class __attribute__((sycl_special_class)) accessor {
+
+public:
+  void use(void) const {}
+  template <typename... T>
+  void use(T... args) {}
+  template <typename... T>
+  void use(T... args) const {}
+  _ImplT<dimensions> impl;
+
+private:
+  void __init(__attribute__((opencl_global)) dataT *Ptr,
+              range<dimensions> AccessRange,
+              range<dimensions> MemRange, id<dimensions> Offset) {}
+};
+
+} // namespace sycl
+} // namespace cl

clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp

+// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -emit-llvm %s -o - | FileCheck %s
+
+// This test checks that compiler generates correct kernel wrapper for basic
+// case.
+
+#include "Inputs/sycl.hpp"
+
+template <typename name, typename Func>
+__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
+  kernelFunc();
+}
+
+int main() {
+  cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write> accessorA;
+    kernel<class kernel_function>(
+      [=]() {
+        accessorA.use();
+      });
+  return 0;
+}
+
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_function
+// CHECK-SAME: i32 addrspace(1)* [[MEM_ARG:%[a-zA-Z0-9_]+]],
+// CHECK-SAME: %"struct.cl::sycl::range"* byval{{.*}}align 4 [[ACC_RANGE:%[a-zA-Z0-9_]+_1]],
+// CHECK-SAME: %"struct.cl::sycl::range"* byval{{.*}}align 4 [[MEM_RANGE:%[a-zA-Z0-9_]+_2]],
+// CHECK-SAME: %"struct.cl::sycl::id"* byval{{.*}}align 4 [[OFFSET:%[a-zA-Z0-9_]+]])
+// Check alloca for pointer argument
+// CHECK: [[MEM_ARG]].addr = alloca i32 addrspace(1)*
+// Check lambda object alloca
+// CHECK: [[ANONALLOCA:%[0-9]+]] = alloca %class.anon
+// CHECK: [[ARANGEA:%agg.tmp]] = alloca %"struct.cl::sycl::range"
+// CHECK: [[MRANGEA:%agg.tmp.*]] = alloca %"struct.cl::sycl::range"
+// CHECK: [[OIDA:%agg.tmp.*]] = alloca %"struct.cl::sycl::id"
+// CHECK: [[ANON:%[0-9]+]] = addrspacecast %class.anon* [[ANONALLOCA]] to %class.anon addrspace(4)*
+// CHECK: [[ARANGET:%agg.tmp.*]] = addrspacecast %"struct.cl::sycl::range"* [[ARANGEA]] to %"struct.cl::sycl::range" addrspace(4)*
+// CHECK: [[MRANGET:%agg.tmp.*]] = addrspacecast %"struct.cl::sycl::range"* [[MRANGEA]] to %"struct.cl::sycl::range" addrspace(4)*
+// CHECK: [[OIDT:%agg.tmp.*]] = addrspacecast %"struct.cl::sycl::id"* [[OIDA]] to %"struct.cl::sycl::id" addrspace(4)*
+//
+// Check store of kernel pointer argument to alloca
+// CHECK: store i32 addrspace(1)* [[MEM_ARG]], i32 addrspace(1)* addrspace(4)* [[MEM_ARG]].addr.ascast, align 8
+
+// Check for default constructor of accessor
+// CHECK: call spir_func {{.*}}accessor
+
+// Check accessor GEP
+// CHECK: [[ACCESSOR:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class.anon, %class.anon addrspace(4)* [[ANON]], i32 0, i32 0
+
+// Check load from kernel pointer argument alloca
+// CHECK: [[MEM_LOAD:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)* addrspace(4)* [[MEM_ARG]].addr.ascast
+
+// Check accessor __init method call
+// CHECK: [[ARANGE:%agg.tmp.*]] = addrspacecast %"struct.cl::sycl::range" addrspace(4)* [[ARANGET]] to %"struct.cl::sycl::range"*
+// CHECK: [[MRANGE:%agg.tmp.*]] = addrspacecast %"struct.cl::sycl::range" addrspace(4)* [[MRANGET]] to %"struct.cl::sycl::range"*
+// CHECK: [[OID:%agg.tmp.*]] = addrspacecast %"struct.cl::sycl::id" addrspace(4)* [[OIDT]] to %"struct.cl::sycl::id"*
+// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.cl::sycl::accessor" addrspace(4)* {{[^,]*}} [[ACCESSOR]], i32 addrspace(1)* [[MEM_LOAD]], %"struct.cl::sycl::range"* byval({{.*}}) align 4 [[ARANGE]], %"struct.cl::sycl::range"* byval({{.*}}) align 4 [[MRANGE]], %"struct.cl::sycl::id"* byval({{.*}}) align 4 [[OID]])
+
+// Check lambda "()" operator call
+// CHECK: call spir_func void @{{.*}}(%class.anon addrspace(4)* {{[^,]*}})

clang/test/CodeGenSYCL/device-functions.cpp

+// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
+
+// This test checks that the compiler emits functions called from a SYCL kernel
+// and doesn't emit functions not called from any SYCL kernel
+
+template <typename T>
+T bar(T arg);
+
+void foo() {
+  int a = 1 + 1 + bar(1);
+}
+
+template <typename T>
+T bar(T arg) {
+  return arg;
+}
+
+template <typename name, typename Func>
+__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) {
+  kernelFunc();
+}
+
+// Make sure that definitions for the types not used in SYCL kernels are not
+// emitted
+// CHECK-NOT: %struct.A
+// CHECK-NOT: @a = {{.*}} %struct.A
+struct A {
+  int x = 10;
+} a;
+
+int main() {
+  a.x = 8;
+  kernel_single_task<class test_kernel>([]() { foo(); });
+  return 0;
+}
+
+// baz is not called from the SYCL kernel, so it must not be emitted
+// CHECK-NOT: define {{.*}} @{{.*}}baz
+void baz() {}
+
+// CHECK-LABEL: define dso_local spir_kernel void @{{.*}}test_kernel
+// CHECK-LABEL: define internal spir_func void @_ZZ4mainENKUlvE_clEv(%class.anon addrspace(4)* {{[^,]*}} %this)
+// CHECK-LABEL: define {{.*}}spir_func void @_Z3foov()
+// CHECK-LABEL: define linkonce_odr spir_func i32 @_Z3barIiET_S0_(i32 %arg)

clang/test/CodeGenSYCL/unique_stable_name.cpp

-// RUN: %clang_cc1 -triple spir64-unknown-unknown-sycldevice -fsycl-is-device -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -triple x86_64-linux-pc  -fsycl-is-host -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
 // CHECK: @[[LAMBDA_KERNEL3:[^\w]+]] = private unnamed_addr constant [[LAMBDA_K3_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZ4mainEUlPZ4mainEUlvE_E_\00"
 // CHECK: @[[INT1:[^\w]+]] = private unnamed_addr constant [[INT_SIZE:\[[0-9]+ x i8\]]] c"_ZTSi\00"
 // CHECK: @[[STRING:[^\w]+]] = private unnamed_addr constant [[STRING_SIZE:\[[0-9]+ x i8\]]] c"_ZTSAppL_ZZ4mainE1jE_i\00",
 // CHECK: @[[INT2:[^\w]+]] = private unnamed_addr constant [[INT_SIZE]] c"_ZTSi\00"
 // CHECK: @[[LAMBDA_X:[^\w]+]] = private unnamed_addr constant [[LAMBDA_X_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZZ4mainENKUlvE0_clEvEUlvE_\00"
 // CHECK: @[[MACRO_X:[^\w]+]] = private unnamed_addr constant [[MACRO_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZZ4mainENKUlvE0_clEvEUlvE0_\00"
 // CHECK: @[[MACRO_Y:[^\w]+]] =  private unnamed_addr constant [[MACRO_SIZE]] c"_ZTSZZ4mainENKUlvE0_clEvEUlvE1_\00"
 // CHECK: @{{.*}} = private unnamed_addr constant [32 x i8] c"_ZTSZZ4mainENKUlvE0_clEvEUlvE2_\00", align 1
struct Derp {
   static constexpr const char str[] = "derp derp derp";
 };
 
 template <typename KernelName, typename KernelType>
 [[clang::sycl_kernel]] void kernel_single_task(KernelType kernelFunc) {
   kernelFunc();
 }
 
+template<typename KernelType>
+void unnamed_kernel_single_task(KernelType kernelFunc) {
+  kernel_single_task<KernelType>(kernelFunc);
+}
+
+template <typename KernelName, typename KernelType>
+void not_kernel_single_task(KernelType kernelFunc) {
+  kernelFunc();
+}
+
 int main() {
-  kernel_single_task<class kernel2>(func<Derp>);
-  // CHECK: call spir_func void @_Z18kernel_single_taskIZ4mainE7kernel2PFPKcvEEvT0_(i8 addrspace(4)* ()* @_Z4funcI4DerpEDTu33__builtin_sycl_unique_stable_nameDtsrT_3strEEEv)
+  not_kernel_single_task<class kernel2>(func<Derp>);
+  // CHECK: call void @_Z22not_kernel_single_taskIZ4mainE7kernel2PFPKcvEEvT0_(i8* ()* @_Z4funcI4DerpEDTu33__builtin_sycl_unique_stable_nameDtsrT_3strEEEv)
 
   auto l1 = []() { return 1; };
   auto l2 = [](decltype(l1) *l = nullptr) { return 2; };
-  kernel_single_task<class kernel3>(l2);
+  kernel_single_task<decltype(l2)>(l2);
   puts(__builtin_sycl_unique_stable_name(decltype(l2)));
-  // CHECK: call spir_func void @_Z18kernel_single_taskIZ4mainE7kernel3Z4mainEUlPZ4mainEUlvE_E_EvT0_
-  // CHECK: call spir_func void @puts(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[LAMBDA_K3_SIZE]], [[LAMBDA_K3_SIZE]]* @[[LAMBDA_KERNEL3]], i32 0, i32 0) to i8 addrspace(4)*))
+  // CHECK: call void @_Z18kernel_single_taskIZ4mainEUlPZ4mainEUlvE_E_S2_EvT0_
+  // CHECK: call void @puts(i8* getelementptr inbounds ([[LAMBDA_K3_SIZE]], [[LAMBDA_K3_SIZE]]* @[[LAMBDA_KERNEL3]], i32 0, i32 0))
 
   constexpr const char str[] = "lalala";
   static_assert(__builtin_strcmp(__builtin_sycl_unique_stable_name(decltype(str)), "_ZTSA7_Kc\0") == 0, "unexpected mangling");
 
   int i = 0;
   puts(__builtin_sycl_unique_stable_name(decltype(i++)));
-  // CHECK: call spir_func void @puts(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[INT_SIZE]], [[INT_SIZE]]* @[[INT1]], i32 0, i32 0) to i8 addrspace(4)*))
+  // CHECK: call void @puts(i8* getelementptr inbounds ([[INT_SIZE]], [[INT_SIZE]]* @[[INT1]], i32 0, i32 0))
 
   // FIXME: Ensure that j is incremented because VLAs are terrible.
   int j = 55;
   puts(__builtin_sycl_unique_stable_name(int[++j]));
-  // CHECK: call spir_func void @puts(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[STRING_SIZE]], [[STRING_SIZE]]* @[[STRING]], i32 0, i32 0) to i8 addrspace(4)*))
+  // CHECK: call void @puts(i8* getelementptr inbounds ([[STRING_SIZE]], [[STRING_SIZE]]* @[[STRING]], i32 0, i32 0))
 
-  // CHECK: define internal spir_func void @_Z18kernel_single_taskIZ4mainE7kernel2PFPKcvEEvT0_
-  // CHECK: declare spir_func i8 addrspace(4)* @_Z4funcI4DerpEDTu33__builtin_sycl_unique_stable_nameDtsrT_3strEEEv
-  // CHECK: define internal spir_func void @_Z18kernel_single_taskIZ4mainE7kernel3Z4mainEUlPZ4mainEUlvE_E_EvT0_
-  // CHECK: define internal spir_func void @_Z18kernel_single_taskIZ4mainE6kernelZ4mainEUlvE0_EvT0_
+  // CHECK: define internal void @_Z22not_kernel_single_taskIZ4mainE7kernel2PFPKcvEEvT0_
+  // CHECK: declare i8* @_Z4funcI4DerpEDTu33__builtin_sycl_unique_stable_nameDtsrT_3strEEEv
+  // CHECK: define internal void @_Z18kernel_single_taskIZ4mainEUlPZ4mainEUlvE_E_S2_EvT0_
+  // CHECK: define internal void @_Z18kernel_single_taskIZ4mainEUlvE0_S0_EvT0_
 
-  kernel_single_task<class kernel>(
+  unnamed_kernel_single_task(
       []() {
         puts(__builtin_sycl_unique_stable_name(int));
-        // CHECK: call spir_func void @puts(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[INT_SIZE]], [[INT_SIZE]]* @[[INT2]], i32 0, i32 0) to i8 addrspace(4)*))
+        // CHECK: call void @puts(i8* getelementptr inbounds ([[INT_SIZE]], [[INT_SIZE]]* @[[INT2]], i32 0, i32 0))
 
         auto x = []() {};
         puts(__builtin_sycl_unique_stable_name(decltype(x)));
-        // CHECK: call spir_func void @puts(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[LAMBDA_X_SIZE]], [[LAMBDA_X_SIZE]]* @[[LAMBDA_X]], i32 0, i32 0) to i8 addrspace(4)*))
+        // CHECK: call void @puts(i8* getelementptr inbounds ([[LAMBDA_X_SIZE]], [[LAMBDA_X_SIZE]]* @[[LAMBDA_X]], i32 0, i32 0))
 
         DEF_IN_MACRO();
-        // CHECK: call spir_func void @puts(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[MACRO_SIZE]], [[MACRO_SIZE]]* @[[MACRO_X]], i32 0, i32 0) to i8 addrspace(4)*))
-        // CHECK: call spir_func void @puts(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[MACRO_SIZE]], [[MACRO_SIZE]]* @[[MACRO_Y]], i32 0, i32 0) to i8 addrspace(4)*))
+        // CHECK: call void @puts(i8* getelementptr inbounds ([[MACRO_SIZE]], [[MACRO_SIZE]]* @[[MACRO_X]], i32 0, i32 0))
+        // CHECK: call void @puts(i8* getelementptr inbounds ([[MACRO_SIZE]], [[MACRO_SIZE]]* @[[MACRO_Y]], i32 0, i32 0))
 
         MACRO_CALLS_MACRO();
-        // CHECK: call spir_func void @puts(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[MACRO_MACRO_SIZE]], [[MACRO_MACRO_SIZE]]* @[[MACRO_MACRO_X]], i32 0, i32 0) to i8 addrspace(4)*))
-        // CHECK: call spir_func void @puts(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[MACRO_MACRO_SIZE]], [[MACRO_MACRO_SIZE]]* @[[MACRO_MACRO_Y]], i32 0, i32 0) to i8 addrspace(4)*))
+        // CHECK: call void @puts(i8* getelementptr inbounds ([[MACRO_MACRO_SIZE]], [[MACRO_MACRO_SIZE]]* @[[MACRO_MACRO_X]], i32 0, i32 0))
+        // CHECK: call void @puts(i8* getelementptr inbounds ([[MACRO_MACRO_SIZE]], [[MACRO_MACRO_SIZE]]* @[[MACRO_MACRO_Y]], i32 0, i32 0))
 
         template_param<int>();
-        // CHECK: call spir_func void @_Z14template_paramIiEvv
+        // CHECK: call void @_Z14template_paramIiEvv
 
         template_param<decltype(x)>();
-        // CHECK: call spir_func void @_Z14template_paramIZZ4mainENKUlvE0_clEvEUlvE_Evv
+        // CHECK: call void @_Z14template_paramIZZ4mainENKUlvE0_clEvEUlvE_Evv
 
         lambda_in_dependent_function<int>();
-        // CHECK: call spir_func void @_Z28lambda_in_dependent_functionIiEvv
+        // CHECK: call void @_Z28lambda_in_dependent_functionIiEvv
 
         lambda_in_dependent_function<decltype(x)>();
-        // CHECK: call spir_func void @_Z28lambda_in_dependent_functionIZZ4mainENKUlvE0_clEvEUlvE_Evv
+        // CHECK: call void @_Z28lambda_in_dependent_functionIZZ4mainENKUlvE0_clEvEUlvE_Evv
 
         lambda_no_dep<int, double>(3, 5.5);
-        // CHECK: call spir_func void @_Z13lambda_no_depIidEvT_T0_(i32 3, double 5.500000e+00)
+        // CHECK: call void @_Z13lambda_no_depIidEvT_T0_(i32 3, double 5.500000e+00)
 
         int a = 5;
         double b = 10.7;
         auto y = [](int a) { return a; };
         auto z = [](double b) { return b; };
         lambda_two_dep<decltype(y), decltype(z)>();
-        // CHECK: call spir_func void @_Z14lambda_two_depIZZ4mainENKUlvE0_clEvEUliE_ZZ4mainENKS0_clEvEUldE_Evv
+        // CHECK: call void @_Z14lambda_two_depIZZ4mainENKUlvE0_clEvEUliE_ZZ4mainENKS0_clEvEUldE_Evv
 
         lambda_two_dep<decltype(z), decltype(y)>();
-        // CHECK: call spir_func void @_Z14lambda_two_depIZZ4mainENKUlvE0_clEvEUldE_ZZ4mainENKS0_clEvEUliE_Evv
+        // CHECK: call void @_Z14lambda_two_depIZZ4mainENKUlvE0_clEvEUldE_ZZ4mainENKS0_clEvEUliE_Evv
       });
 }
 
-// CHECK: define linkonce_odr spir_func void @_Z14template_paramIiEvv
-// CHECK: call spir_func void @puts(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[INT_SIZE]], [[INT_SIZE]]* @[[INT3]], i32 0, i32 0) to i8 addrspace(4)*))
+// CHECK: define linkonce_odr void @_Z14template_paramIiEvv
+// CHECK: call void @puts(i8* getelementptr inbounds ([[INT_SIZE]], [[INT_SIZE]]* @[[INT3]], i32 0, i32 0))
 
-// CHECK: define internal spir_func void @_Z14template_paramIZZ4mainENKUlvE0_clEvEUlvE_Evv
-// CHECK: call spir_func void @puts(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[LAMBDA_SIZE]], [[LAMBDA_SIZE]]* @[[LAMBDA]], i32 0, i32 0) to i8 addrspace(4)*))
+// CHECK: define internal void @_Z14template_paramIZZ4mainENKUlvE0_clEvEUlvE_Evv
+// CHECK: call void @puts(i8* getelementptr inbounds ([[LAMBDA_SIZE]], [[LAMBDA_SIZE]]* @[[LAMBDA]], i32 0, i32 0))
 
-// CHECK: define linkonce_odr spir_func void @_Z28lambda_in_dependent_functionIiEvv
-// CHECK: call spir_func void @puts(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[DEP_INT_SIZE]], [[DEP_INT_SIZE]]* @[[LAMBDA_IN_DEP_INT]], i32 0, i32 0) to i8 addrspace(4)*))
+// CHECK: define linkonce_odr void @_Z28lambda_in_dependent_functionIiEvv
+// CHECK: call void @puts(i8* getelementptr inbounds ([[DEP_INT_SIZE]], [[DEP_INT_SIZE]]* @[[LAMBDA_IN_DEP_INT]], i32 0, i32 0))
 
-// CHECK: define internal spir_func void @_Z28lambda_in_dependent_functionIZZ4mainENKUlvE0_clEvEUlvE_Evv
-// CHECK: call spir_func void @puts(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[DEP_LAMBDA_SIZE]], [[DEP_LAMBDA_SIZE]]* @[[LAMBDA_IN_DEP_X]], i32 0, i32 0) to i8 addrspace(4)*))
+// CHECK: define internal void @_Z28lambda_in_dependent_functionIZZ4mainENKUlvE0_clEvEUlvE_Evv
+// CHECK: call void @puts(i8* getelementptr inbounds ([[DEP_LAMBDA_SIZE]], [[DEP_LAMBDA_SIZE]]* @[[LAMBDA_IN_DEP_X]], i32 0, i32 0))
 
-// CHECK: define linkonce_odr spir_func void @_Z13lambda_no_depIidEvT_T0_(i32 %a, double %b)
-// CHECK: call spir_func void @puts(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[NO_DEP_LAMBDA_SIZE]], [[NO_DEP_LAMBDA_SIZE]]* @[[LAMBDA_NO_DEP]], i32 0, i32 0) to i8 addrspace(4)*))
+// CHECK: define linkonce_odr void @_Z13lambda_no_depIidEvT_T0_(i32 %a, double %b)
+// CHECK: call void @puts(i8* getelementptr inbounds ([[NO_DEP_LAMBDA_SIZE]], [[NO_DEP_LAMBDA_SIZE]]* @[[LAMBDA_NO_DEP]], i32 0, i32 0))
 
-// CHECK: define internal spir_func void @_Z14lambda_two_depIZZ4mainENKUlvE0_clEvEUliE_ZZ4mainENKS0_clEvEUldE_Evv
-// CHECK: call spir_func void @puts(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[DEP_LAMBDA1_SIZE]], [[DEP_LAMBDA1_SIZE]]* @[[LAMBDA_TWO_DEP]], i32 0, i32 0) to i8 addrspace(4)*))
+// CHECK: define internal void @_Z14lambda_two_depIZZ4mainENKUlvE0_clEvEUliE_ZZ4mainENKS0_clEvEUldE_Evv
+// CHECK: call void @puts(i8* getelementptr inbounds ([[DEP_LAMBDA1_SIZE]], [[DEP_LAMBDA1_SIZE]]* @[[LAMBDA_TWO_DEP]], i32 0, i32 0))
 
-// CHECK: define internal spir_func void @_Z14lambda_two_depIZZ4mainENKUlvE0_clEvEUldE_ZZ4mainENKS0_clEvEUliE_Evv
-// CHECK: call spir_func void @puts(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[DEP_LAMBDA2_SIZE]], [[DEP_LAMBDA2_SIZE]]* @[[LAMBDA_TWO_DEP2]], i32 0, i32 0) to i8 addrspace(4)*))
+// CHECK: define internal void @_Z14lambda_two_depIZZ4mainENKUlvE0_clEvEUldE_ZZ4mainENKS0_clEvEUliE_Evv
+// CHECK: call void @puts(i8* getelementptr inbounds ([[DEP_LAMBDA2_SIZE]], [[DEP_LAMBDA2_SIZE]]* @[[LAMBDA_TWO_DEP2]], i32 0, i32 0))

clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp

-// RUN: %clang_cc1 -triple spir64-unknown-unknown-sycldevice -aux-triple x86_64-pc-windows-msvc -fsycl-is-device -disable-llvm-passes -fsycl-is-device -emit-llvm %s -o - | FileCheck %s
-// RUN: %clang_cc1 -triple x86_64-pc-windows-msvc -fsycl-is-device -disable-llvm-passes -fsycl-is-device -emit-llvm %s -o - | FileCheck %s
-
+// RUN: %clang_cc1 -triple spir64-unknown-unknown -aux-triple x86_64-pc-windows-msvc -fsycl-is-device -disable-llvm-passes -fsycl-is-device -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -triple x86_64-pc-windows-msvc -fsycl-is-device -disable-llvm-passes -fsycl-is-device -emit-llvm %s -o - | FileCheck %s --check-prefixes=WIN,CHECK
 
 template<typename KN, typename Func>
 __attribute__((sycl_kernel)) void kernel(Func F){
   F();
 }
 
 template<typename KN, typename Func>
 __attribute__((sycl_kernel)) void kernel2(Func F){

clang/test/Misc/pragma-attribute-supported-attributes-list.test

 // CHECK-NEXT: Pointer (SubjectMatchRule_record_not_is_union)
 // CHECK-NEXT: ReleaseHandle (SubjectMatchRule_variable_is_parameter)
 // CHECK-NEXT: RenderScriptKernel (SubjectMatchRule_function)
 // CHECK-NEXT: ReqdWorkGroupSize (SubjectMatchRule_function)
 // CHECK-NEXT: Restrict (SubjectMatchRule_function)
 // CHECK-NEXT: ReturnTypestate (SubjectMatchRule_function, SubjectMatchRule_variable_is_parameter)
 // CHECK-NEXT: ReturnsNonNull (SubjectMatchRule_objc_method, SubjectMatchRule_function)
 // CHECK-NEXT: ReturnsTwice (SubjectMatchRule_function)
+// CHECK-NEXT: SYCLSpecialClass (SubjectMatchRule_record)
 // CHECK-NEXT: ScopedLockable (SubjectMatchRule_record)
 // CHECK-NEXT: Section (SubjectMatchRule_function, SubjectMatchRule_variable_is_global, SubjectMatchRule_objc_method, SubjectMatchRule_objc_property)
 // CHECK-NEXT: SetTypestate (SubjectMatchRule_function_is_member)
 // CHECK-NEXT: SpeculativeLoadHardening (SubjectMatchRule_function, SubjectMatchRule_objc_method)
 // CHECK-NEXT: StandaloneDebug (SubjectMatchRule_record)
 // CHECK-NEXT: SwiftAsync (SubjectMatchRule_function, SubjectMatchRule_objc_method)
 // CHECK-NEXT: SwiftAsyncContext (SubjectMatchRule_variable_is_parameter)
 // CHECK-NEXT: SwiftAsyncError (SubjectMatchRule_function, SubjectMatchRule_objc_method)

clang/test/SemaSYCL/Inputs/sycl.hpp

+namespace cl {
+namespace sycl {
+namespace access {
+
+enum class target {
+  global_buffer = 2014,
+  constant_buffer,
+  local,
+  image,
+  host_buffer,
+  host_image,
+  image_array
+};
+
+enum class mode {
+  read = 1024,
+  write,
+  read_write,
+  discard_write,
+  discard_read_write,
+  atomic
+};
+
+enum class placeholder { false_t,
+                         true_t };
+
+enum class address_space : int {
+  private_space = 0,
+  global_space,
+  constant_space,
+  local_space
+};
+} // namespace access
+
+
+template <int dim>
+struct range {
+};
+
+template <int dim>
+struct id {
+};
+
+template <int dim>
+struct _ImplT {
+  range<dim> AccessRange;
+  range<dim> MemRange;
+  id<dim> Offset;
+};
+
+template <typename dataT, access::target accessTarget>
+struct DeviceValueType;
+
+template <typename dataT>
+struct DeviceValueType<dataT, access::target::global_buffer> {
+  using type = __attribute__((opencl_global)) dataT;
+};
+
+template <typename dataT>
+struct DeviceValueType<dataT, access::target::constant_buffer> {
+  using type = __attribute__((opencl_global)) const dataT;
+};
+
+template <typename dataT>
+struct DeviceValueType<dataT, access::target::local> {
+  using type = __attribute__((opencl_local)) dataT;
+};
+
+template <typename dataT, int dimensions, access::mode accessmode,
+          access::target accessTarget = access::target::global_buffer,
+          access::placeholder isPlaceholder = access::placeholder::false_t>
+class __attribute__((sycl_special_class)) accessor {
+public:
+  void use(void) const {}
+  void use(void *) const {}
+  _ImplT<dimensions> impl;
+
+private:
+  using PtrType = typename DeviceValueType<dataT, accessTarget>::type *;
+  void __init(PtrType Ptr, range<dimensions> AccessRange,
+              range<dimensions> MemRange, id<dimensions> Offset) {}
+  friend class stream;
+};
+
+} // namespace sycl
+} // namespace cl

clang/test/SemaSYCL/accessors-targets.cpp

+// RUN: %clang_cc1 -fsycl-is-device -ast-dump %s | FileCheck %s
+
+// This test checks that compiler generates correct kernel wrapper arguments for
+// different accessors targets.
+
+#include "Inputs/sycl.hpp"
+
+using namespace cl::sycl;
+
+template <typename name, typename Func>
+__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
+  kernelFunc();
+}
+
+int main() {
+
+  accessor<int, 1, access::mode::read_write,
+           access::target::local>
+      local_acc;
+  accessor<int, 1, access::mode::read_write,
+           access::target::global_buffer>
+      global_acc;
+  kernel<class use_local>(
+      [=]() {
+        local_acc.use();
+      });
+  kernel<class use_global>(
+      [=]() {
+        global_acc.use();
+      });
+}
+// CHECK: {{.*}}use_local{{.*}} 'void (__local int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)'
+// CHECK: {{.*}}use_global{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)'

clang/test/SemaSYCL/basic-kernel-wrapper.cpp

+// RUN: %clang_cc1 -fsycl-is-device -ast-dump %s | FileCheck %s
+
+// This test checks that compiler generates correct kernel wrapper for basic
+// case.
+
+#include "Inputs/sycl.hpp"
+
+template <typename Acc>
+struct AccWrapper { Acc accessor; };
+
+template <typename name, typename Func>
+__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
+  kernelFunc();
+}
+
+int main() {
+  cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write> acc;
+  kernel<class kernel_wrapper>(
+      [=]() {
+        acc.use();
+      });
+}
+
+// Check declaration of the kernel
+
+// CHECK: FunctionDecl {{.*}}kernel_wrapper{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)'
+
+// Check parameters of the kernel
+
+// CHECK: ParmVarDecl {{.*}} used [[_arg_Mem:[0-9a-zA-Z_]+]] '__global int *'
+// CHECK: ParmVarDecl {{.*}} used [[_arg_AccessRange:[0-9a-zA-Z_]+]] 'cl::sycl::range<1>'
+// CHECK: ParmVarDecl {{.*}} used [[_arg_MemRange:[0-9a-zA-Z_]+]] 'cl::sycl::range<1>'
+// CHECK: ParmVarDecl {{.*}} used [[_arg_Offset:[0-9a-zA-Z_]+]] 'cl::sycl::id<1>'
+
+// Check body of the kernel
+
+// Check lambda declaration inside the wrapper
+
+// CHECK: DeclStmt
+// CHECK-NEXT: VarDecl {{.*}} used '(lambda at {{.*}}basic-kernel-wrapper.cpp{{.*}})'
+
+// Check accessor initialization
+
+// CHECK: CXXMemberCallExpr {{.*}} 'void'
+// CHECK-NEXT: MemberExpr {{.*}} 'void ({{.*}}PtrType, range<1>, range<1>, id<1>)' lvalue .__init
+// CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write>':'cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t>' lvalue .
+// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}basic-kernel-wrapper.cpp{{.*}})' lvalue Var
+
+// CHECK-NEXT: ImplicitCastExpr {{.*}} <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr {{.*}} '__global int *' lvalue ParmVar {{.*}} '[[_arg_Mem]]' '__global int *'
+
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'range<1>':'cl::sycl::range<1>'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const cl::sycl::range<1>' lvalue <NoOp>
+// CHECK-NEXT: DeclRefExpr {{.*}} 'cl::sycl::range<1>' lvalue ParmVar {{.*}} '[[_arg_AccessRange]]' 'cl::sycl::range<1>'
+
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'range<1>':'cl::sycl::range<1>'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const cl::sycl::range<1>' lvalue <NoOp>
+// CHECK-NEXT: DeclRefExpr {{.*}} 'cl::sycl::range<1>' lvalue ParmVar {{.*}} '[[_arg_MemRange]]' 'cl::sycl::range<1>'
+
+// CHECK-NEXT: CXXConstructExpr {{.*}} 'id<1>':'cl::sycl::id<1>'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const cl::sycl::id<1>' lvalue <NoOp>
+// CHECK-NEXT: DeclRefExpr {{.*}} 'cl::sycl::id<1>' lvalue ParmVar {{.*}} '[[_arg_Offset]]' 'cl::sycl::id<1>'
+
+// Check that body of the kernel caller function is included into kernel
+
+// CHECK: CompoundStmt {{.*}}
+// CHECK-NEXT: CXXOperatorCallExpr {{.*}} 'void'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'void (*)() const' <FunctionToPointerDecay>
+// CHECK-NEXT: DeclRefExpr {{.*}} 'void () const' lvalue CXXMethod {{.*}} 'operator()' 'void () const'
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const (lambda at {{.*}}basic-kernel-wrapper.cpp{{.*}})' lvalue <NoOp>
+// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}basic-kernel-wrapper.cpp{{.*}})' lvalue Var
+
+// Check kernel wrapper attributes
+
+// CHECK: OpenCLKernelAttr {{.*}} Implicit
+// CHECK: AsmLabelAttr {{.*}} Implicit "{{.*}}kernel_wrapper{{.*}}"
+// CHECK: ArtificialAttr {{.*}} Implicit

clang/test/SemaSYCL/built-in-type-kernel-arg.cpp

+// RUN: %clang_cc1 -fsycl-is-device -ast-dump %s | FileCheck %s
+
+// This test checks that compiler generates correct initialization for arguments
+// that have struct or built-in type inside the OpenCL kernel
+
+#include "Inputs/sycl.hpp"
+
+template <typename name, typename Func>
+__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
+  kernelFunc();
+}
+
+struct test_struct {
+  int data;
+};
+
+void test(const int some_const) {
+  kernel<class kernel_const>(
+      [=]() {
+        int a = some_const;
+      });
+}
+
+int main() {
+  int data = 5;
+  test_struct s;
+  s.data = data;
+  kernel<class kernel_int>(
+      [=]() {
+        int kernel_data = data;
+      });
+  kernel<class kernel_struct>(
+      [=]() {
+        test_struct k_s;
+        k_s = s;
+      });
+  const int some_const = 10;
+  test(some_const);
+  return 0;
+}
+// Check kernel parameters
+// CHECK: FunctionDecl {{.*}}kernel_const{{.*}} 'void (const int)'
+// CHECK: ParmVarDecl {{.*}} used _arg_ 'const int'
+
+// Check that lambda field of const built-in type is initialized
+// CHECK: VarDecl {{.*}}'(lambda at {{.*}}built-in-type-kernel-arg.cpp{{.*}})'
+// CHECK-NEXT: InitListExpr
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr {{.*}} 'const int' lvalue ParmVar {{.*}} '_arg_' 'const int'
+
+// Check kernel parameters
+// CHECK: {{.*}}kernel_int{{.*}} 'void (int)'
+// CHECK: ParmVarDecl {{.*}} used _arg_ 'int'
+
+// Check that lambda field of built-in type is initialized
+// CHECK: VarDecl {{.*}}'(lambda at {{.*}}built-in-type-kernel-arg.cpp{{.*}})'
+// CHECK-NEXT: InitListExpr
+// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
+// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_' 'int'
+
+// Check kernel parameters
+// CHECK: {{.*}}kernel_struct{{.*}} 'void (test_struct)'
+// CHECK: ParmVarDecl {{.*}} used _arg_ 'test_struct'
+
+// Check that lambda field of struct type is initialized
+// CHECK: VarDecl {{.*}}'(lambda at {{.*}}built-in-type-kernel-arg.cpp{{.*}})'
+// CHECK-NEXT: InitListExpr
+// CHECK-NEXT: CXXConstructExpr {{.*}}'test_struct'{{.*}}void (const test_struct &)
+// CHECK-NEXT: ImplicitCastExpr {{.*}}'const test_struct' lvalue <NoOp>
+// CHECK-NEXT: DeclRefExpr {{.*}} 'test_struct' lvalue ParmVar {{.*}} '_arg_' 'test_struct'

clang/test/SemaSYCL/fake-accessors.cpp

+// RUN: %clang_cc1 -fsycl-is-device -Wno-unused-value -Wno-int-to-void-pointer-cast -ast-dump %s | FileCheck %s
+
+// This test checks that the compiler handles SYCL kenrel parameters of accessor
+// type correctly (i.e. sends a pointer as a separate parameter) and doesn't
+// apply this approach for user types even if they have the same name (e.g. 
+// `accessor` or `foo::cl::sycl::accessor`). We also verify that type aliasing
+// via typedef or using doesn't confuse the compiler.
+
+#include "Inputs/sycl.hpp"
+
+namespace foo {
+namespace cl {
+namespace sycl {
+class accessor {
+public:
+  int field;
+};
+} // namespace sycl
+} // namespace cl
+} // namespace foo
+
+class accessor {
+public:
+  int field;
+};
+
+typedef cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write,
+                           cl::sycl::access::target::global_buffer>
+    MyAccessorTD;
+
+using MyAccessorA = cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write,
+                                       cl::sycl::access::target::global_buffer>;
+
+template <typename name, typename Func>
+__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
+  kernelFunc();
+}
+
+int main() {
+  foo::cl::sycl::accessor acc = {1};
+  accessor acc1 = {1};
+
+  cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write> accessorA;
+  cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write> accessorB;
+  cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write> accessorC;
+    kernel<class fake_accessors>(
+        [=]() {
+          accessorA.use((void*)(acc.field + acc1.field));
+        });
+    kernel<class accessor_typedef>(
+        [=]() {
+          accessorB.use((void*)(acc.field + acc1.field));
+        });
+    kernel<class accessor_alias>(
+        [=]() {
+          accessorC.use((void*)(acc.field + acc1.field));
+        });
+    kernel<class user_types>(
+        [=]() {
+          acc.field + acc1.field;
+        });
+  return 0;
+}
+// CHECK: fake_accessors{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor)
+// CHECK: accessor_typedef{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor)
+// CHECK: accessor_alias{{.*}} 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor)
+// CHECK: user_types{{.*}} 'void (foo::cl::sycl::accessor, accessor)

clang/test/SemaSYCL/mangle-kernel.cpp

+// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -I %S/../Headers/Inputs/include/ -ast-dump %s | FileCheck %s --check-prefix=CHECK-64
+// RUN: %clang_cc1 -fsycl-is-device -triple spir-unknown-unknown-sycldevice -I %S/../Headers/Inputs/include/ -ast-dump %s | FileCheck %s --check-prefix=CHECK-32
+#include "Inputs/sycl.hpp"
+#include <stdlib.h>
+
+// This test checks that SYCL kernel name mangling matches Itanium ABI
+
+template <typename name, typename Func>
+__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
+  kernelFunc();
+}
+
+template <typename T>
+class SimpleVadd;
+
+int main() {
+  kernel<class SimpleVadd<int>>(
+      [=](){});
+
+  kernel<class SimpleVadd<double>>(
+      [=](){});
+
+  kernel<class SimpleVadd<size_t>>(
+      [=](){});
+  return 0;
+}
+
+// CHECK: _ZTS10SimpleVaddIiE
+// CHECK: _ZTS10SimpleVaddIdE
+// CHECK-64: _ZTS10SimpleVaddImE
+// CHECK-32: _ZTS10SimpleVaddIjE