diff --git a/clang/include/clang/AST/GlobalDecl.h b/clang/include/clang/AST/GlobalDecl.h
index 97caff0198cb0..085b525cbd4a7 100644
--- a/clang/include/clang/AST/GlobalDecl.h
+++ b/clang/include/clang/AST/GlobalDecl.h
@@ -164,7 +164,8 @@ class GlobalDecl {
   }
 
   static KernelReferenceKind getDefaultKernelReference(const FunctionDecl *D) {
-    return (D->hasAttr<DeviceKernelAttr>() || D->getLangOpts().CUDAIsDevice)
+    return (D->hasAttr<DeviceKernelAttr>() || D->getLangOpts().CUDAIsDevice ||
+            D->hasAttr<SYCLKernelAttr>())
                ? KernelReferenceKind::Kernel
                : KernelReferenceKind::Stub;
   }
diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td
index 70a3f7a549372..450dd4874a3ef 100644
--- a/clang/include/clang/Basic/Attr.td
+++ b/clang/include/clang/Basic/Attr.td
@@ -1640,7 +1640,6 @@ def CUDAShared : InheritableAttr {
 }
 def : MutualExclusions<[CUDAConstant, CUDAShared, HIPManaged]>;
 
-<<<<<<< HEAD
 def GlobalStorageNonLocalVar : SubsetSubject<Var,
                                              [{S->hasGlobalStorage() &&
                                                !S->isLocalVarDeclOrParm()}],
@@ -1667,13 +1666,13 @@ def SYCLGlobalVar : InheritableAttr {
   let LangOpts = [SYCLIsDevice];
   // Only used internally by the SYCL implementation
   let Documentation = [Undocumented];
-=======
+}
+
 def SYCLKernel : InheritableAttr {
   let Spellings = [Clang<"sycl_kernel">];
   let Subjects = SubjectList<[FunctionTmpl]>;
-  let LangOpts = [SYCLDevice];
+  let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
   let Documentation = [SYCLKernelDocs];
->>>>>>> 1db148cc946eb95fefd5399766e379fc030eef78
 }
 
 def DeviceKernel : DeclOrTypeAttr {
diff --git a/clang/lib/AST/Decl.cpp b/clang/lib/AST/Decl.cpp
index b663653101139..234dcd59de537 100644
--- a/clang/lib/AST/Decl.cpp
+++ b/clang/lib/AST/Decl.cpp
@@ -5569,7 +5569,9 @@ FunctionDecl *FunctionDecl::CreateDeserialized(ASTContext &C, GlobalDeclID ID) {
 
 bool FunctionDecl::isReferenceableKernel() const {
   return hasAttr<CUDAGlobalAttr>() ||
-         DeviceKernelAttr::isOpenCLSpelling(getAttr<DeviceKernelAttr>()) || hasAttr<DeviceKernelAttr>();
+  return hasAttr<CUDAGlobalAttr>() ||
+         DeviceKernelAttr::isOpenCLSpelling(getAttr<DeviceKernelAttr>()) ||
+         hasAttr<SYCLKernelAttr>();
 }
 
 BlockDecl *BlockDecl::Create(ASTContext &C, DeclContext *DC, SourceLocation L) {
diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp
index 669b2e25cd9a4..d881f4ea65e0f 100644
--- a/clang/lib/CodeGen/CodeGenFunction.cpp
+++ b/clang/lib/CodeGen/CodeGenFunction.cpp
@@ -772,7 +772,7 @@ void CodeGenFunction::EmitKernelMetadata(const FunctionDecl *FD,
   }
 
   bool IsKernelOrDevice =
-      FD->hasAttr<DeviceKernelAttr>() || FD->hasAttr<SYCLDeviceAttr>();
+      FD->hasAttr<SYCLKernelAttr>() || FD->hasAttr<SYCLDeviceAttr>();
   const IntelReqdSubGroupSizeAttr *ReqSubGroup =
       FD->getAttr<IntelReqdSubGroupSizeAttr>();
 
diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp
index 427dd8290dad6..ae8d30217573b 100644
--- a/clang/lib/CodeGen/CodeGenModule.cpp
+++ b/clang/lib/CodeGen/CodeGenModule.cpp
@@ -6953,7 +6953,7 @@ CodeGenModule::getLLVMLinkageForDeclarator(const DeclaratorDecl *D,
   // with the SYCL_EXTERNAL macro. For any function or variable that does not
   // have this, linkonce_odr suffices. If -fno-sycl-rdc is passed, we know there
   // is only one translation unit and can so mark them internal.
-  if (getLangOpts().SYCLIsDevice && !D->hasAttr<DeviceKernelAttr>() &&
+  if (getLangOpts().SYCLIsDevice && !D->hasAttr<SYCLKernelAttr>() &&
       !D->hasAttr<SYCLDeviceAttr>() && !D->hasAttr<SYCLExternalAttr>() &&
       !SemaSYCL::isTypeDecoratedWithDeclAttribute<SYCLDeviceGlobalAttr>(
           D->getType()))
diff --git a/clang/lib/Sema/SemaChecking.cpp b/clang/lib/Sema/SemaChecking.cpp
index 8c9aff93f0031..458b88633249a 100644
--- a/clang/lib/Sema/SemaChecking.cpp
+++ b/clang/lib/Sema/SemaChecking.cpp
@@ -4408,7 +4408,7 @@ void Sema::checkCall(NamedDecl *FDecl, const FunctionProtoType *Proto,
   if (FD)
     diagnoseArgDependentDiagnoseIfAttrs(FD, ThisArg, Args, Loc);
 
-  if (FD && FD->hasAttr<DeviceKernelAttr>() && getLangOpts().isSYCL())
+  if (FD && FD->hasAttr<SYCLKernelAttr>())
     SYCL().CheckSYCLKernelCall(FD, Args);
 
   // Diagnose variadic calls in SYCL.
diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index 54c5fe49b7efd..833037e8adfee 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -20957,7 +20957,7 @@ Sema::DeviceDiagnosticReason Sema::getEmissionReason(const FunctionDecl *FD) {
   // FIXME: This should really be a bitwise-or of the language modes.
   if (FD->hasAttr<SYCLSimdAttr>())
     return Sema::DeviceDiagnosticReason::Esimd;
-  if (FD->hasAttr<SYCLDeviceAttr>() || FD->hasAttr<DeviceKernelAttr>())
+  if (FD->hasAttr<SYCLDeviceAttr>() || FD->hasAttr<SYCLKernelAttr>())
     return getLangOpts().SYCLCUDACompat
                ? Sema::DeviceDiagnosticReason::SyclCudaCompat
                : Sema::DeviceDiagnosticReason::Sycl;
@@ -20978,20 +20978,12 @@ Sema::FunctionEmissionStatus Sema::getEmissionStatus(const FunctionDecl *FD,
                                                      bool Final) {
   assert(FD && "Expected non-null FunctionDecl");
 
-<<<<<<< HEAD
-=======
-  // SYCL functions can be template, so we check if they have appropriate
-  // attribute prior to checking if it is a template.
-  if (LangOpts.SYCLIsDevice && FD->hasAttr<SYCLKernelAttr>())
-    return FunctionEmissionStatus::Emitted;
-
->>>>>>> 1db148cc946eb95fefd5399766e379fc030eef78
   // Templates are emitted when they're instantiated.
   if (FD->isDependentContext())
     return FunctionEmissionStatus::TemplateDiscarded;
 
   if (LangOpts.SYCLIsDevice &&
-      (FD->hasAttr<SYCLDeviceAttr>() || FD->hasAttr<DeviceKernelAttr>()) &&
+      (FD->hasAttr<SYCLDeviceAttr>() || FD->hasAttr<SYCLKernelAttr>()) &&
       !FD->hasAttr<ArtificialAttr>())
     return FunctionEmissionStatus::Emitted;
 
@@ -21077,7 +21069,7 @@ Sema::FunctionEmissionStatus Sema::getEmissionStatus(const FunctionDecl *FD,
   }
 
   if (getLangOpts().SYCLIsDevice) {
-    if (!FD->hasAttr<SYCLDeviceAttr>() && !FD->hasAttr<DeviceKernelAttr>())
+    if (!FD->hasAttr<SYCLDeviceAttr>() && !FD->hasAttr<SYCLKernelAttr>())
       return FunctionEmissionStatus::Unknown;
 
     // Check whether this function is externally visible -- if so, it's
diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index 30359dd76e45b..aeba952bf2e06 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -5266,20 +5266,7 @@ static void handleCallConvAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
 static void handleDeviceKernelAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
   const auto *FD = dyn_cast_or_null<FunctionDecl>(D);
   bool IsFunctionTemplate = FD && FD->getDescribedFunctionTemplate();
-<<<<<<< HEAD
-  if (S.getLangOpts().isSYCL()) {
-    if (!IsFunctionTemplate) {
-      S.Diag(AL.getLoc(), diag::warn_attribute_wrong_decl_type_str)
-          << AL << AL.isRegularKeywordAttribute() << "function templates";
-    } else {
-      S.SYCL().handleKernelAttr(D, AL);
-    }
-  } else if (DeviceKernelAttr::isSYCLSpelling(AL)) {
-    S.Diag(AL.getLoc(), diag::warn_attribute_ignored) << AL;
-  } else if (S.getASTContext().getTargetInfo().getTriple().isNVPTX()) {
-=======
   if (S.getASTContext().getTargetInfo().getTriple().isNVPTX()) {
->>>>>>> 1db148cc946eb95fefd5399766e379fc030eef78
     handleGlobalAttr(S, D, AL);
   } else {
     // OpenCL C++ will throw a more specific error.
@@ -7184,13 +7171,11 @@ ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D, const ParsedAttr &AL,
   case ParsedAttr::AT_EnumExtensibility:
     handleEnumExtensibilityAttr(S, D, AL);
     break;
-<<<<<<< HEAD
   case ParsedAttr::AT_SYCLSimd:
     handleSimpleAttribute<SYCLSimdAttr>(S, D, AL);
-=======
+    break;
   case ParsedAttr::AT_SYCLKernel:
     S.SYCL().handleKernelAttr(D, AL);
->>>>>>> 1db148cc946eb95fefd5399766e379fc030eef78
     break;
   case ParsedAttr::AT_SYCLExternal:
     handleSimpleAttribute<SYCLExternalAttr>(S, D, AL);
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 0e568e4f76186..8f55303f9284b 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -911,7 +911,7 @@ class SingleDeviceFunctionTracker {
     // a SYCLKernel or SYCLDevice attribute on it, add it to the set of
     // routines potentially reachable on device. This is to diagnose such
     // cases later in finalizeSYCLDelayedAnalysis().
-    if (!CurrentDecl->isDefined() && !CurrentDecl->hasAttr<DeviceKernelAttr>() &&
+    if (!CurrentDecl->isDefined() && !CurrentDecl->hasAttr<SYCLKernelAttr>() &&
         !CurrentDecl->hasAttr<SYCLDeviceAttr>())
       Parent.SemaSYCLRef.addFDToReachableFromSyclDevice(CurrentDecl,
                                                         CallStack.back());
@@ -971,7 +971,7 @@ class SingleDeviceFunctionTracker {
     if (isSYCLKernelBodyFunction(CurrentDecl)) {
       // This is a direct callee of the kernel.
       if (CallStack.size() == 1 &&
-          CallStack.back()->hasAttr<DeviceKernelAttr>()) {
+          CallStack.back()->hasAttr<SYCLKernelAttr>()) {
         assert(!KernelBody && "inconsistent call graph - only one kernel body "
                               "function can be called");
         KernelBody = CurrentDecl;
@@ -3009,7 +3009,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
     // to TransformStmt in replaceWithLocalClone can diagnose something that got
     // diagnosed on the actual kernel.
     KernelDecl->addAttr(
-        DeviceKernelAttr::CreateImplicit(SemaSYCLRef.getASTContext()));
+        SYCLKernelAttr::CreateImplicit(SemaSYCLRef.getASTContext()));
 
     SemaSYCLRef.addSyclDeviceDecl(KernelDecl);
   }
@@ -6093,7 +6093,7 @@ void SemaSYCL::finalizeSYCLDelayedAnalysis(const FunctionDecl *Caller,
     return;
 
   // If Callee has a SYCL attribute, no diagnostic needed.
-  if (Callee->hasAttr<SYCLDeviceAttr>() || Callee->hasAttr<DeviceKernelAttr>())
+  if (Callee->hasAttr<SYCLDeviceAttr>() || Callee->hasAttr<SYCLKernelAttr>())
     return;
 
   // If Callee has a CUDA device attribute, no diagnostic needed.
@@ -7716,24 +7716,6 @@ void SemaSYCL::performSYCLDelayedAttributesAnalaysis(const FunctionDecl *FD) {
            diag::warn_sycl_incorrect_use_attribute_non_kernel_function)
           << KernelAttr;
   }
-<<<<<<< HEAD
-=======
-
-  // Function must have at least one argument.
-  if (getFunctionOrMethodNumParams(D) != 1) {
-    Diag(FT->getLocation(), diag::warn_sycl_kernel_num_of_function_params);
-    return;
-  }
-
-  // Function must return void.
-  QualType RetTy = getFunctionOrMethodResultType(D);
-  if (!RetTy->isVoidType()) {
-    Diag(FT->getLocation(), diag::warn_sycl_kernel_return_type);
-    return;
-  }
-
-  handleSimpleAttribute<SYCLKernelAttr>(*this, D, AL);
->>>>>>> 1db148cc946eb95fefd5399766e379fc030eef78
 }
 
 void SemaSYCL::handleKernelEntryPointAttr(Decl *D, const ParsedAttr &AL) {
diff --git a/clang/lib/Sema/SemaSYCLDeclAttr.cpp b/clang/lib/Sema/SemaSYCLDeclAttr.cpp
index 58ee6e81534be..83d8ba5bb7528 100644
--- a/clang/lib/Sema/SemaSYCLDeclAttr.cpp
+++ b/clang/lib/Sema/SemaSYCLDeclAttr.cpp
@@ -55,7 +55,7 @@ void SemaSYCL::handleKernelAttr(Decl *D, const ParsedAttr &AL) {
     return;
   }
 
-  handleSimpleAttribute<DeviceKernelAttr>(*this, D, AL);
+  handleSimpleAttribute<SYCLKernelAttr>(*this, D, AL);
 }
 
 // Returns a DupArgResult value; Same means the args have the same value,
diff --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
index 24579c20d81e4..86e417df70127 100644
--- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
+++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
@@ -7659,7 +7659,7 @@ static void processFunctionInstantiation(Sema &S,
                                   DefinitionRequired, AtEndOfTU);
   if (!FD->isDefined())
     return;
-  if (S.LangOpts.SYCLIsDevice && FD->hasAttr<DeviceKernelAttr>())
+  if (S.LangOpts.SYCLIsDevice && FD->hasAttr<SYCLKernelAttr>())
     S.SYCL().ConstructOpenCLKernel(FD, MC);
   FD->setInstantiationIsPending(false);
 }
diff --git a/clang/lib/Sema/SemaType.cpp b/clang/lib/Sema/SemaType.cpp
index 61edb435a8adf..a39df3ceef008 100644
--- a/clang/lib/Sema/SemaType.cpp
+++ b/clang/lib/Sema/SemaType.cpp
@@ -3797,12 +3797,7 @@ static CallingConv getCCForDeclaratorChunk(
     }
   }
   for (const ParsedAttr &AL : D.getDeclSpec().getAttributes()) {
-<<<<<<< HEAD
-    if (AL.getKind() == ParsedAttr::AT_DeviceKernel &&
-	!DeviceKernelAttr::isSYCLSpelling(AL)) {
-=======
     if (AL.getKind() == ParsedAttr::AT_DeviceKernel) {
->>>>>>> 1db148cc946eb95fefd5399766e379fc030eef78
       CC = CC_DeviceKernel;
       break;
     }
diff --git a/clang/test/SemaSYCL/sycl-esimd-ast.cpp b/clang/test/SemaSYCL/sycl-esimd-ast.cpp
index 7d0a1e65708fd..8c3e01aa529a8 100644
--- a/clang/test/SemaSYCL/sycl-esimd-ast.cpp
+++ b/clang/test/SemaSYCL/sycl-esimd-ast.cpp
@@ -14,7 +14,7 @@ int main() {
   deviceQueue.submit([&](sycl::handler &h) {
     // CHECK-LABEL: FunctionDecl {{.*}}test_kernel1
     // CHECK:       SYCLSimdAttr {{.*}} Implicit
-    // CHECK-NEXT:  DeviceKernelAttr {{.*}} Implicit
+    // CHECK-NEXT:  SYCLKernelAttr {{.*}} Implicit
     // CHECK-NEXT:  AsmLabelAttr {{.*}} Implicit
     // CHECK-NEXT:  SYCLSimdAttr {{.*}}
     h.single_task<class test_kernel1>(
@@ -22,7 +22,7 @@ int main() {
 
     // CHECK-LABEL: FunctionDecl {{.*}}test_kernel2
     // CHECK:       SYCLSimdAttr {{.*}} Implicit
-    // CHECK-NEXT:  DeviceKernelAttr {{.*}} Implicit
+    // CHECK-NEXT:  SYCLKernelAttr {{.*}} Implicit
     // CHECK-NEXT:  AsmLabelAttr {{.*}} Implicit
     // CHECK-NEXT:  SYCLSimdAttr {{.*}}
     h.single_task<class test_kernel2>(