diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td
--- a/clang/include/clang/Basic/Attr.td
+++ b/clang/include/clang/Basic/Attr.td
@@ -296,6 +296,7 @@
 def Borland : LangOpt<"Borland">;
 def CUDA : LangOpt<"CUDA">;
 def HIP : LangOpt<"HIP">;
+def SYCL : LangOpt<"SYCLIsDevice">;
 def COnly : LangOpt<"COnly", "!LangOpts.CPlusPlus">;
 def CPlusPlus : LangOpt<"CPlusPlus">;
 def OpenCL : LangOpt<"OpenCL">;
@@ -1021,6 +1022,20 @@
   let Documentation = [Undocumented];
 }
 
+def SYCLDevice : InheritableAttr {
+  let Spellings = [];
+  let Subjects = SubjectList<[Function, Var]>;
+  let LangOpts = [SYCL];
+  let Documentation = [Undocumented];
+}
+
+def SYCLKernel : InheritableAttr {
+  let Spellings = [Clang<"sycl_kernel">];
+  let Subjects = SubjectList<[Function]>;
+  let LangOpts = [SYCL];
+  let Documentation = [SYCLKernelDocs];
+}
+
 def C11NoReturn : InheritableAttr {
   let Spellings = [Keyword<"_Noreturn">];
   let Subjects = SubjectList<[Function], ErrorDiag>;
diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td
--- a/clang/include/clang/Basic/AttrDocs.td
+++ b/clang/include/clang/Basic/AttrDocs.td
@@ -253,6 +253,79 @@
   }];
 }
 
+def SYCLKernelDocs : Documentation {
+  let Category = DocCatFunction;
+  let Content = [{
+The ``sycl_kernel`` attribute specifies that a function will be used by the
+compiler to outline device code and to generate OpenCL kernel.
+Here is a code example of the SYCL program, which demonstrates compiler's
+outlining job:
+.. code-block:: c++
+
+  int foo(int x) { return ++x; }
+
+  using namespace cl::sycl;
+  queue Q;
+  buffer<int, 1> a(range<1>{1024});
+  Q.submit([&](handler& cgh) {
+    auto A = a.get_access<access::mode::write>(cgh);
+    cgh.parallel_for<init_a>(range<1>{1024}, [=](id<1> index) {
+      A[index] = index[0] * 2 + index[1] + foo(42);
+    });
+  }
+
+The lambda that is passed to the ``parallel_for`` is called a SYCL "kernel
+function". A SYCL "kernel function" defines entry point to the "device part"
+of the code. Compiler will traverse all symbols accessible from a
+"kernel function" and add them to the "device part" of the code. In this code
+example, compiler will add "foo" function to the "device part" of the code.
+More details about compilation of functions for device can be found in the
+SYCL 1.2.1 specification Section 6.4.
+To show to the compiler entry point to the "device part" of the code SYCL
+runtime can use the ``sycl_kernel`` attribute in the following way:
+.. code-block:: c++
+namespace cl {
+namespace sycl {
+class handler {
+  template <typename KernelName, typename KernelType/*, ...*/>
+  __attribute__((sycl_kernel)) void sycl_kernel_function(KernelType KernelFuncObj) {
+    // ...
+    KernelFuncObj();
+  }
+
+  template <typename KernelName, typename KernelType, int Dims>
+  void parallel_for(range<Dims> NumWorkItems, KernelType KernelFunc) {
+#ifdef __SYCL_DEVICE_ONLY__
+    sycl_kernel_function<KernelName, KernelType, Dims>(KernelFunc);
+#else
+    // Host implementation
+#endif
+  }
+};
+} // namespace sycl
+} // namespace cl
+
+The compiler will also generate OpenCL kernel using the function marked with the
+``sycl_kernel`` attribute.
+Here is the list of SYCL device compiler expectations with regard to the
+function marked with the ``sycl_kernel`` attribute:
+
+- Function template with at least two template parameters is expected. The compiler
+generates OpenCL kernel and uses first template parameter as unique name to the
+generated OpenCL kernel. Host application uses this unique name to invoke the
+OpenCL kernel generated for the ``sycl_kernel_function`` specialized by
+this name and second template parameter ``KernelType`` (which might be a lambda type).
+- Function must have at least one parameter. First parameter expected to be a
+function object type (named or unnamed i.e. lambda). Compiler uses function
+object type fields to generate OpenCL kernel parameters.
+- Function must return void. Compiler re-uses body of marked function to
+generate OpenCL kernel body and OpenCL kernel must return void.
+The ``sycl_kernel_function`` in the previous code sample meets these
+expectations.
+
+  }];
+}
+
 def C11NoReturnDocs : Documentation {
   let Category = DocCatFunction;
   let Content = [{
@@ -4194,4 +4267,4 @@
 not initialized on device side. It has internal linkage and is initialized by
 the initializer on host side.
   }];
-}
\ No newline at end of file
+}
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -9734,4 +9734,9 @@
   "%select{non-pointer|function pointer|void pointer}0 argument to "
   "'__builtin_launder' is not allowed">;
 
+// SYCL-specific diagnostics
+def warn_sycl_kernel_attribute_invalid : Warning<
+  "'sycl_kernel' attribute only applies to template funtions with special prototype, "
+  "please refer 'sycl_kernel' attribute documentation">, InGroup<IgnoredAttributes>;
+
 } // end of sema component.
diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/include/clang/Sema/Sema.h
@@ -11192,6 +11192,29 @@
     ConstructorDestructor,
     BuiltinFunction
   };
+
+private:
+  /// Contains function declarations to be added to the SYCL device code.
+  /// In SYCL, when we generate device code, we don't know which functions we
+  /// will emit before we emit sycl kernels, so we add device functions to this
+  /// array and handle it in separate way.
+  SmallVector<Decl *, 4> SyclDeviceFunctions;
+
+public:
+  /// This function adds the function declaration to the SYCL device code.
+  void addSyclDeviceFunc(Decl *D) {
+    D->addAttr(SYCLDeviceAttr::CreateImplicit(Context));
+    SyclDeviceFunctions.push_back(D);
+  }
+  /// Access to SYCL device function decls.
+  SmallVectorImpl<Decl *> &syclDeviceFuncs() { return SyclDeviceFunctions; }
+
+  /// Constructs an OpenCL kernel using the KernelCaller function and adds it to
+  /// the SYCL device code.
+  void constructOpenCLKernel(FunctionDecl *KernelCallerFunc);
+  /// This function marks all functions accessible from SYCL kernels with the
+  /// SYCL device attribute and adds them to the SYCL device code.
+  void markSYCLDevice(void);
 };
 
 /// RAII object that enters a new expression evaluation context.
diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp
--- a/clang/lib/CodeGen/CodeGenModule.cpp
+++ b/clang/lib/CodeGen/CodeGenModule.cpp
@@ -2405,6 +2405,11 @@
   if (Global->hasAttr<IFuncAttr>())
     return emitIFuncDefinition(GD);
 
+  // If this is SYCL device, only emit declarations marked with the SYCL device
+  // attribute.
+  if (LangOpts.SYCLIsDevice && !Global->hasAttr<SYCLDeviceAttr>())
+    return;
+
   // If this is a cpu_dispatch multiversion function, emit the resolver.
   if (Global->hasAttr<CPUDispatchAttr>())
     return emitCPUDispatchDefinition(GD);
@@ -2526,6 +2531,10 @@
     // The value must be emitted, but cannot be emitted eagerly.
     assert(!MayBeEmittedEagerly(Global));
     addDeferredDeclToEmit(GD);
+  } else if (LangOpts.SYCLIsDevice) {
+    // SYCL kernels can be templated and not called from anywhere in the
+    // module but should be emitted.
+    addDeferredDeclToEmit(GD);
   } else {
     // Otherwise, remember that we saw a deferred decl with this name.  The
     // first use of the mangled name will cause it to move into
diff --git a/clang/lib/Parse/ParseAST.cpp b/clang/lib/Parse/ParseAST.cpp
--- a/clang/lib/Parse/ParseAST.cpp
+++ b/clang/lib/Parse/ParseAST.cpp
@@ -168,6 +168,12 @@
   for (Decl *D : S.WeakTopLevelDecls())
     Consumer->HandleTopLevelDecl(DeclGroupRef(D));
 
+  if (S.getLangOpts().SYCLIsDevice) {
+    for (Decl *D : S.syclDeviceFuncs()) {
+      Consumer->HandleTopLevelDecl(DeclGroupRef(D));
+    }
+  }
+
   Consumer->HandleTranslationUnit(S.getASTContext());
 
   // Finalize the template instantiation observer chain.
diff --git a/clang/lib/Sema/CMakeLists.txt b/clang/lib/Sema/CMakeLists.txt
--- a/clang/lib/Sema/CMakeLists.txt
+++ b/clang/lib/Sema/CMakeLists.txt
@@ -56,6 +56,7 @@
   SemaStmt.cpp
   SemaStmtAsm.cpp
   SemaStmtAttr.cpp
+  SemaSYCL.cpp
   SemaTemplate.cpp
   SemaTemplateDeduction.cpp
   SemaTemplateInstantiate.cpp
diff --git a/clang/lib/Sema/Sema.cpp b/clang/lib/Sema/Sema.cpp
--- a/clang/lib/Sema/Sema.cpp
+++ b/clang/lib/Sema/Sema.cpp
@@ -907,6 +907,9 @@
     PerformPendingInstantiations();
   }
 
+  if (getLangOpts().SYCLIsDevice)
+    markSYCLDevice();
+
   assert(LateParsedInstantiations.empty() &&
          "end of TU template instantiation should not create more "
          "late-parsed templates");
diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -6444,6 +6444,55 @@
       AL.getRange(), S.Context, AL.getAttributeSpellingListIndex()));
 }
 
+static void handleSYCLKernelAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
+
+  // The 'sycl_device' attribute applies only to functions
+  auto *FD = dyn_cast<FunctionDecl>(D);
+  if (!FD) {
+    S.Diag(AL.getLoc(), diag::warn_sycl_kernel_attribute_invalid);
+    return;
+  }
+
+  FunctionTemplateDecl *FT = FD->getDescribedFunctionTemplate();
+
+  // Function template is expected
+  if (!FT) {
+    S.Diag(AL.getLoc(), diag::warn_sycl_kernel_attribute_invalid);
+    return;
+  }
+
+  // Function must have at least two template parameters
+  TemplateParameterList *TL = FT->getTemplateParameters();
+  if (TL->size() < 2) {
+    S.Diag(AL.getLoc(), diag::warn_sycl_kernel_attribute_invalid);
+    return;
+  }
+
+  // Template parameters must be a typenames
+  for (unsigned I = 0; I < 2; ++I) {
+    NamedDecl* TParam = TL->getParam(I);
+    if (!isa<TemplateTypeParmDecl>(TParam)) {
+      S.Diag(AL.getLoc(), diag::warn_sycl_kernel_attribute_invalid);
+      return;
+    }
+  }
+
+  // Function must have at least one argument
+  if (getFunctionOrMethodNumParams(D) < 1) {
+    S.Diag(AL.getLoc(), diag::warn_sycl_kernel_attribute_invalid);
+    return;
+  }
+
+  // Function must return void
+  QualType RetTy = getFunctionOrMethodResultType(D);
+  if (!RetTy->isVoidType()) {
+    S.Diag(AL.getLoc(), diag::warn_sycl_kernel_attribute_invalid);
+    return;
+  }
+
+  handleSimpleAttribute<SYCLKernelAttr>(S, D, AL);
+}
+
 static void handleDestroyAttr(Sema &S, Decl *D, const ParsedAttr &A) {
   if (!cast<VarDecl>(D)->hasGlobalStorage()) {
     S.Diag(D->getLocation(), diag::err_destroy_attr_on_non_static_var)
@@ -6767,6 +6816,9 @@
   case ParsedAttr::AT_Flatten:
     handleSimpleAttribute<FlattenAttr>(S, D, AL);
     break;
+  case ParsedAttr::AT_SYCLKernel:
+    handleSYCLKernelAttr(S, D, AL);
+    break;
   case ParsedAttr::AT_Format:
     handleFormatAttr(S, D, AL);
     break;
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
new file mode 100644
--- /dev/null
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -0,0 +1,74 @@
+//===- 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 "clang/AST/RecursiveASTVisitor.h"
+#include "clang/Sema/Sema.h"
+
+using namespace clang;
+
+class MarkDeviceFunction : public RecursiveASTVisitor<MarkDeviceFunction> {
+public:
+  MarkDeviceFunction(Sema &S)
+      : RecursiveASTVisitor<MarkDeviceFunction>(), SemaRef(S) {}
+
+  bool VisitCallExpr(CallExpr *E) {
+    if (FunctionDecl *Callee = E->getDirectCallee()) {
+      Callee = Callee->getCanonicalDecl();
+      // Remember that all SYCL kernel functions have deferred
+      // instantiation as template functions. It means that
+      // all functions used by kernel have already been parsed and have
+      // definitions.
+      if (FunctionDecl *Def = Callee->getDefinition()) {
+        if (!Def->hasAttr<SYCLDeviceAttr>()) {
+          SemaRef.addSyclDeviceFunc(Def);
+          this->TraverseStmt(Def->getBody());
+        }
+      }
+    }
+    return true;
+  }
+
+  bool VisitCXXConstructExpr(CXXConstructExpr *E) {
+    CXXConstructorDecl *Ctor = E->getConstructor();
+
+    if (FunctionDecl *Def = Ctor->getDefinition())
+      SemaRef.addSyclDeviceFunc(Def);
+
+    const CXXRecordDecl *ConstructedType = Ctor->getParent();
+    if (ConstructedType->hasUserDeclaredDestructor()) {
+      CXXDestructorDecl *Dtor = ConstructedType->getDestructor();
+
+      if (FunctionDecl *Def = Dtor->getDefinition())
+        SemaRef.addSyclDeviceFunc(Def);
+    }
+    return true;
+  }
+
+private:
+  Sema &SemaRef;
+};
+
+void Sema::constructOpenCLKernel(FunctionDecl *KernelCallerFunc) {
+  addSyclDeviceFunc(KernelCallerFunc);
+}
+
+void Sema::markSYCLDevice(void) {
+  // Let's mark all called functions with the SYCL Device attribute.
+  MarkDeviceFunction Marker(*this);
+  for (const auto &Elt : syclDeviceFuncs()) {
+    if (auto *Func = dyn_cast<FunctionDecl>(Elt)) {
+      if (FunctionDecl *Def = Func->getDefinition()) {
+        if (!Def->hasAttr<SYCLDeviceAttr>())
+          addSyclDeviceFunc(Def);
+        Marker.TraverseStmt(Def->getBody());
+      }
+    }
+  }
+}
diff --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
--- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
+++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
@@ -5532,14 +5532,30 @@
             Function, [this, Inst, DefinitionRequired](FunctionDecl *CurFD) {
               InstantiateFunctionDefinition(/*FIXME:*/ Inst.second, CurFD, true,
                                             DefinitionRequired, true);
-              if (CurFD->isDefined())
+              if (CurFD->isDefined()) {
+                // Because all SYCL kernel functions are template functions -
+                // they have deferred instantination. We need bodies of these
+                // functions so we are checking for the SYCL kernel attribute
+                // after instantination.
+                if (getLangOpts().SYCLIsDevice &&
+                    CurFD->hasAttr<SYCLKernelAttr>())
+                  constructOpenCLKernel(CurFD);
                 CurFD->setInstantiationIsPending(false);
+              }
             });
       } else {
         InstantiateFunctionDefinition(/*FIXME:*/ Inst.second, Function, true,
                                       DefinitionRequired, true);
-        if (Function->isDefined())
+        if (Function->isDefined()) {
+          // Because all SYCL kernel functions are template functions - they
+          // have deferred instantination. We need bodies of these functions
+          // so we are checking for the SYCL kernel attribute after
+          // instantination.
+          if (getLangOpts().SYCLIsDevice &&
+              Function->hasAttr<SYCLKernelAttr>())
+            constructOpenCLKernel(Function);
           Function->setInstantiationIsPending(false);
+        }
       }
       continue;
     }
diff --git a/clang/test/CodeGenSYCL/device-functions.cpp b/clang/test/CodeGenSYCL/device-functions.cpp
new file mode 100644
--- /dev/null
+++ b/clang/test/CodeGenSYCL/device-functions.cpp
@@ -0,0 +1,29 @@
+// RUN: %clang_cc1 -triple spir64-unknown-unknown -std=c++11 -fsycl-is-device -S -emit-llvm %s -o - | FileCheck %s
+
+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(Func kernelFunc) {
+  kernelFunc();
+}
+
+int main() {
+  kernel_single_task<class fake_kernel>([]() { foo(); });
+  return 0;
+}
+// CHECK: define spir_func void @{{.*}}foo
+// CHECK: define linkonce_odr spir_func i32 @{{.*}}bar
+// CHECK: define internal spir_func void @{{.*}}kernel_single_task
+// FIXME: Next function is lambda () operator. spir_func calling convention
+// is missed for C++ methods.
+// CHECK: define internal void @"_ZZ4mainENK3$_0clEv"(%class.anon* %this)
diff --git a/clang/test/Misc/pragma-attribute-supported-attributes-list.test b/clang/test/Misc/pragma-attribute-supported-attributes-list.test
--- a/clang/test/Misc/pragma-attribute-supported-attributes-list.test
+++ b/clang/test/Misc/pragma-attribute-supported-attributes-list.test
@@ -126,6 +126,7 @@
 // CHECK-NEXT: ReturnTypestate (SubjectMatchRule_function, SubjectMatchRule_variable_is_parameter)
 // CHECK-NEXT: ReturnsNonNull (SubjectMatchRule_objc_method, SubjectMatchRule_function)
 // CHECK-NEXT: ReturnsTwice (SubjectMatchRule_function)
+// CHECK-NEXT: SYCLKernel (SubjectMatchRule_function)
 // 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)
diff --git a/clang/test/SemaSYCL/device-attributes-on-non-sycl.cpp b/clang/test/SemaSYCL/device-attributes-on-non-sycl.cpp
new file mode 100644
--- /dev/null
+++ b/clang/test/SemaSYCL/device-attributes-on-non-sycl.cpp
@@ -0,0 +1,14 @@
+// RUN: %clang_cc1 -std=c++11 -fsyntax-only -fsycl-is-device -verify %s
+// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify -x c++ %s
+
+#ifndef __SYCL_DEVICE_ONLY__
+// expected-warning@+7 {{'sycl_kernel' attribute ignored}}
+// expected-warning@+8 {{'sycl_kernel' attribute ignored}}
+#else
+// expected-no-diagnostics
+#endif
+
+template <typename T, typename A, int B>
+__attribute__((sycl_kernel)) void foo(T P);
+template <typename T, typename A, int B>
+[[clang::sycl_kernel]] void foo1(T P);
diff --git a/clang/test/SemaSYCL/device-attributes.cpp b/clang/test/SemaSYCL/device-attributes.cpp
new file mode 100644
--- /dev/null
+++ b/clang/test/SemaSYCL/device-attributes.cpp
@@ -0,0 +1,41 @@
+// RUN: %clang_cc1 -std=c++11 -fsyntax-only -fsycl-is-device -verify %s
+
+[[clang::sycl_kernel]] int gv2 = 0; // expected-warning {{'sycl_kernel' attribute only applies to functions}}
+__attribute__((sycl_kernel)) int gv3 = 0; // expected-warning {{'sycl_kernel' attribute only applies to functions}}
+
+__attribute__((sycl_kernel(1))) void foo(); // expected-error {{'sycl_kernel' attribute takes no arguments}}
+[[clang::sycl_kernel(1)]] void foo2(); // expected-error {{'sycl_kernel' attribute takes no arguments}}
+
+// Only template functions
+__attribute__((sycl_kernel)) void foo(); // expected-warning {{'sycl_kernel' only applies to template funtions with special prototype, please refer 'sycl_kernel' attribute documentation}}
+[[clang::sycl_kernel]] void foo1(); // expected-warning {{'sycl_kernel' only applies to template funtions with special prototype, please refer 'sycl_kernel' attribute documentation}}
+
+// At least two template parameters
+template <typename T>
+__attribute__((sycl_kernel)) void foo(T P); // expected-warning {{'sycl_kernel' only applies to template funtions with special prototype, please refer 'sycl_kernel' attribute documentation}}
+template <typename T>
+[[clang::sycl_kernel]] void foo1(T P); // expected-warning {{'sycl_kernel' only applies to template funtions with special prototype, please refer 'sycl_kernel' attribute documentation}}
+
+// Both first two template parameters must be a typenames
+template <typename T, int A>
+__attribute__((sycl_kernel)) void foo(T P); // expected-warning {{'sycl_kernel' only applies to template funtions with special prototype, please refer 'sycl_kernel' attribute documentation}}
+template <typename T, int A>
+[[clang::sycl_kernel]] void foo1(T P); // expected-warning {{'sycl_kernel' only applies to template funtions with special prototype, please refer 'sycl_kernel' attribute documentation}}
+
+// Must return void
+template <typename T, typename A>
+__attribute__((sycl_kernel)) int foo(T P); // expected-warning {{'sycl_kernel' only applies to template funtions with special prototype, please refer 'sycl_kernel' attribute documentation}}
+template <typename T, typename A>
+[[clang::sycl_kernel]] int foo1(T P); // expected-warning {{'sycl_kernel' only applies to template funtions with special prototype, please refer 'sycl_kernel' attribute documentation}}
+
+// Must take at least one argument
+template <typename T, typename A>
+__attribute__((sycl_kernel)) void foo(); // expected-warning {{'sycl_kernel' only applies to template funtions with special prototype, please refer 'sycl_kernel' attribute documentation}}
+template <typename T, typename A>
+[[clang::sycl_kernel]] void foo1(); // expected-warning {{'sycl_kernel' only applies to template funtions with special prototype, please refer 'sycl_kernel' attribute documentation}}
+
+// No diagnosticts
+template <typename T, typename A, int B>
+__attribute__((sycl_kernel)) void foo(T P);
+template <typename T, typename A, int B>
+[[clang::sycl_kernel]] void foo1(T P);
diff --git a/clang/test/SemaSYCL/device-code-outlining.cpp b/clang/test/SemaSYCL/device-code-outlining.cpp
new file mode 100644
--- /dev/null
+++ b/clang/test/SemaSYCL/device-code-outlining.cpp
@@ -0,0 +1,37 @@
+// RUN: %clang_cc1 -std=c++11 -fsycl-is-device -ast-dump %s | FileCheck %s
+
+template <typename T>
+T bar(T arg);
+// CHECK: FunctionTemplateDecl {{.*}} bar
+// CHECK: SYCLDeviceAttr {{.*}} Implicit
+
+void foo() {
+  int a = 1 + 1 + bar(1);
+}
+// CHECK: FunctionDecl {{.*}} foo
+// CHECK: SYCLDeviceAttr {{.*}} Implicit
+
+template <typename T>
+T bar(T arg) {
+  return arg;
+}
+
+template <typename name, typename Func>
+__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {
+  kernelFunc();
+}
+// CHECK: FunctionTemplateDecl {{.*}} kernel_single_task
+// CHECK: SYCLDeviceAttr {{.*}} Implicit
+
+void host_foo() {
+  int b = 0;
+}
+// CHECK: FunctionDecl {{.*}} host_foo
+// CHECK-NOT: SYCLDeviceAttr
+// CHECK: FunctionDecl {{.*}} main
+
+int main() {
+  kernel_single_task<class fake_kernel>([]() { foo(); });
+  host_foo();
+  return 0;
+}