diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td
index c6a9961e33454..ca111375d5501 100644
--- a/clang/include/clang/Basic/Attr.td
+++ b/clang/include/clang/Basic/Attr.td
@@ -1650,7 +1650,7 @@ def DeviceKernel : DeclOrTypeAttr {
     // list, but here we have the same spelling with unscores and without,
     // so handle that case manually.
       return A.getAttributeSpellingListIndex() == Keyword_kernel ||
-      A.getAttrName()->getName() == "kernel";
+      (A.getAttrName() && A.getAttrName()->getName() == "kernel");
     }
     static inline bool isOpenCLSpelling(const AttributeCommonInfo* A) {
     if (!A) return false;
diff --git a/clang/lib/Basic/Targets/X86.h b/clang/lib/Basic/Targets/X86.h
index 6f8a2365be256..5fbd545a8d079 100644
--- a/clang/lib/Basic/Targets/X86.h
+++ b/clang/lib/Basic/Targets/X86.h
@@ -796,7 +796,7 @@ class LLVM_LIBRARY_VISIBILITY X86_64TargetInfo : public X86TargetInfo {
     case CC_X86RegCall:
       return CCCR_OK;
     case CC_DeviceKernel:
-      return IsOpenCL ? CCCR_OK : CCCR_Warning;
+      return IsOpenCL || IsSYCLDevice ? CCCR_OK : CCCR_Warning;
     default:
       return CCCR_Warning;
     }
@@ -842,10 +842,12 @@ class LLVM_LIBRARY_VISIBILITY X86_64TargetInfo : public X86TargetInfo {
   void adjust(DiagnosticsEngine &Diags, LangOptions &Opts) override {
     TargetInfo::adjust(Diags, Opts);
     IsOpenCL = Opts.OpenCL;
+    IsSYCLDevice = Opts.SYCLIsDevice;
   }
 
 private:
   bool IsOpenCL = false;
+  bool IsSYCLDevice = false;
 };
 
 // x86-64 UEFI target
diff --git a/clang/lib/CodeGen/CGCall.cpp b/clang/lib/CodeGen/CGCall.cpp
index 37adc21d89ae3..6e96762678cad 100644
--- a/clang/lib/CodeGen/CGCall.cpp
+++ b/clang/lib/CodeGen/CGCall.cpp
@@ -94,6 +94,10 @@ unsigned CodeGenTypes::ClangCallConvToLLVMCallConv(CallingConv CC) {
       return llvm::CallingConv::AMDGPU_KERNEL;
     if (CGM.getTriple().isNVPTX())
       return llvm::CallingConv::PTX_Kernel;
+    if (CGM.getLangOpts().SYCLIsNativeCPU)
+      return CGM.getTarget().getDefaultCallingConv();
+    if (CGM.getLangOpts().SYCLIsDevice)
+      return CGM.getTarget().getDefaultCallingConv();
     llvm_unreachable("Unknown kernel calling convention");
   }
   case CC_PreserveMost:
@@ -2628,8 +2632,8 @@ void CodeGenModule::ConstructAttributeList(StringRef Name,
                                  NumElemsParam);
     }
 
-    if (DeviceKernelAttr::isOpenCLSpelling(
-            TargetDecl->getAttr<DeviceKernelAttr>()) &&
+    if (TargetDecl->hasAttr<DeviceKernelAttr>() &&
+        (getLangOpts().OpenCL || getLangOpts().SYCLIsDevice) &&
         CallingConv != CallingConv::CC_C &&
         CallingConv != CallingConv::CC_SpirFunction) {
       // Check CallingConv to avoid adding uniform-work-group-size attribute to
@@ -3013,8 +3017,9 @@ void CodeGenModule::ConstructAttributeList(StringRef Name,
     // > data type, the OpenCL compiler can assume that the pointee is always
     // > appropriately aligned as required by the data type.
     if (TargetDecl &&
-        DeviceKernelAttr::isOpenCLSpelling(
-            TargetDecl->getAttr<DeviceKernelAttr>()) &&
+
+        (TargetDecl->hasAttr<DeviceKernelAttr>() &&
+         (getLangOpts().OpenCL || getLangOpts().SYCLIsDevice)) &&
         ParamType->isPointerType()) {
       QualType PTy = ParamType->getPointeeType();
       if (!PTy->isIncompleteType() && PTy->isConstantSizeType()) {
diff --git a/clang/lib/Sema/SemaChecking.cpp b/clang/lib/Sema/SemaChecking.cpp
index 325f3307f47c5..5fb158df21600 100644
--- a/clang/lib/Sema/SemaChecking.cpp
+++ b/clang/lib/Sema/SemaChecking.cpp
@@ -4014,7 +4014,7 @@ void Sema::checkCall(NamedDecl *FDecl, const FunctionProtoType *Proto,
   if (FD)
     diagnoseArgDependentDiagnoseIfAttrs(FD, ThisArg, Args, Loc);
 
-  if (FD && FD->hasAttr<DeviceKernelAttr>())
+  if (FD && FD->hasAttr<DeviceKernelAttr>() && getLangOpts().isSYCL())
     SYCL().CheckSYCLKernelCall(FD, Args);
 
   // Diagnose variadic calls in SYCL.
diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index b89f3f1acce53..dc56c354b43c9 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -20762,7 +20762,8 @@ Sema::FunctionEmissionStatus Sema::getEmissionStatus(const FunctionDecl *FD,
     return FunctionEmissionStatus::TemplateDiscarded;
 
   if (LangOpts.SYCLIsDevice &&
-      (FD->hasAttr<SYCLDeviceAttr>() || FD->hasAttr<DeviceKernelAttr>()))
+      (FD->hasAttr<SYCLDeviceAttr>() || FD->hasAttr<DeviceKernelAttr>()) &&
+      !FD->hasAttr<ArtificialAttr>())
     return FunctionEmissionStatus::Emitted;
 
   // Check whether this function is an externally visible definition.
diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index f09a1ce036c7d..aeeb3a50e6d8d 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -5355,7 +5355,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();
-  if (S.getLangOpts().SYCLIsDevice) {
+  if (S.getLangOpts().isSYCL()) {
     if (!IsFunctionTemplate) {
       S.Diag(AL.getLoc(), diag::warn_attribute_wrong_decl_type_str)
           << AL << AL.isRegularKeywordAttribute() << "function templates";
diff --git a/clang/lib/Sema/SemaType.cpp b/clang/lib/Sema/SemaType.cpp
index e2b65e72a9f5d..f28e45bf001c4 100644
--- a/clang/lib/Sema/SemaType.cpp
+++ b/clang/lib/Sema/SemaType.cpp
@@ -3785,12 +3785,11 @@ static CallingConv getCCForDeclaratorChunk(
       }
     }
   }
-  if (!S.getLangOpts().isSYCL()) {
-    for (const ParsedAttr &AL : D.getDeclSpec().getAttributes()) {
-      if (AL.getKind() == ParsedAttr::AT_DeviceKernel) {
-        CC = CC_DeviceKernel;
-        break;
-      }
+  for (const ParsedAttr &AL : D.getDeclSpec().getAttributes()) {
+    if (AL.getKind() == ParsedAttr::AT_DeviceKernel &&
+	!DeviceKernelAttr::isSYCLSpelling(AL)) {
+      CC = CC_DeviceKernel;
+      break;
     }
   }
   return CC;
diff --git a/clang/test/CodeGenSYCL/kernel-op-calls.cpp b/clang/test/CodeGenSYCL/kernel-op-calls.cpp
index aa5de944c1299..1dee50ba4b3ad 100644
--- a/clang/test/CodeGenSYCL/kernel-op-calls.cpp
+++ b/clang/test/CodeGenSYCL/kernel-op-calls.cpp
@@ -53,7 +53,7 @@ int main() {
 
   Q.submit([&](sycl::handler& cgh) {
       ESIMDFunctor EF;
-      // CHECK: define dso_local spir_kernel void @_ZTS12ESIMDFunctor() {{.*}} !intel_reqd_sub_group_size !{{[0-9]+}} !work_group_size_hint !{{[0-9]+}} !kernel_arg_accessor_ptr !{{[0-9]+}} !sycl_explicit_simd !{{[0-9]+}} !sycl_fixed_targets !{{[0-9]+}} {
+      // CHECK: define dso_local spir_kernel void @_ZTS12ESIMDFunctor() {{.*}} !sycl_explicit_simd !{{[0-9]+}} !intel_reqd_sub_group_size !{{[0-9]+}} !work_group_size_hint !{{[0-9]+}} !kernel_arg_accessor_ptr !{{[0-9]+}} !sycl_fixed_targets !{{[0-9]+}} {
       cgh.parallel_for(sycl::range<1>(10), EF);
     });
 
diff --git a/clang/test/CodeGenSYCL/usm-int-header.cpp b/clang/test/CodeGenSYCL/usm-int-header.cpp
index 080772b4b4a60..d99b689bf6531 100644
--- a/clang/test/CodeGenSYCL/usm-int-header.cpp
+++ b/clang/test/CodeGenSYCL/usm-int-header.cpp
@@ -30,4 +30,4 @@ int main() {
     });
 }
 
-// CHECK: FunctionDecl {{.*}}usm_test{{.*}} 'void (__global int *, __global float *)'
+// CHECK: FunctionDecl {{.*}}usm_test{{.*}} 'void (__global int *, __global float *)  __attribute__((device_kernel))'
diff --git a/clang/test/SemaOpenCL/sampled_image_overload.cl b/clang/test/SemaOpenCL/sampled_image_overload.cl
index 4d4a722018a29..b14026e6eabd5 100644
--- a/clang/test/SemaOpenCL/sampled_image_overload.cl
+++ b/clang/test/SemaOpenCL/sampled_image_overload.cl
@@ -3,7 +3,7 @@
 void __attribute__((overloadable)) foo(__ocl_sampled_image1d_ro_t);
 void __attribute__((overloadable)) foo(__ocl_sampled_image2d_ro_t);
 
-// CHECK: FunctionDecl {{.*}} <{{.*}}> line:{{.*}} ker 'void (__private __ocl_sampled_image1d_ro_t, __private __ocl_sampled_image2d_ro_t)'
+// CHECK: FunctionDecl {{.*}} <{{.*}}> line:{{.*}} ker 'void (__private __ocl_sampled_image1d_ro_t, __private __ocl_sampled_image2d_ro_t) __attribute__((device_kernel))'
 void kernel ker(__ocl_sampled_image1d_ro_t src1, __ocl_sampled_image2d_ro_t src2) {
   // CHECK: CallExpr
   // CHECK-NEXT: ImplicitCastExpr {{.*}} <{{.*}}> 'void (*)(__private __ocl_sampled_image1d_ro_t)'
diff --git a/clang/test/SemaSYCL/accessor_inheritance.cpp b/clang/test/SemaSYCL/accessor_inheritance.cpp
index 8d83bbf5b2738..efbd84d0906a6 100644
--- a/clang/test/SemaSYCL/accessor_inheritance.cpp
+++ b/clang/test/SemaSYCL/accessor_inheritance.cpp
@@ -30,7 +30,7 @@ int main() {
 }
 
 // Check kernel parameters
-// CHECK: FunctionDecl {{.*}}kernel{{.*}} 'void (int, int, __global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>, __global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>, int)'
+// CHECK: FunctionDecl {{.*}}kernel{{.*}} 'void (int, int, __global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>, __global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>, int) __attribute__((device_kernel))'
 // CHECK: ParmVarDecl{{.*}} used _arg_A 'int'
 // CHECK: ParmVarDecl{{.*}} used _arg_B 'int'
 // CHECK: ParmVarDecl{{.*}} used _arg_AccField '__global char *'
diff --git a/clang/test/SemaSYCL/accessors-targets-image.cpp b/clang/test/SemaSYCL/accessors-targets-image.cpp
index f608e66c107cd..4fcc401360400 100644
--- a/clang/test/SemaSYCL/accessors-targets-image.cpp
+++ b/clang/test/SemaSYCL/accessors-targets-image.cpp
@@ -72,12 +72,12 @@ int main() {
   });
 }
 
-// CHECK: {{.*}}use_image1d_r 'void (__read_only image1d_t)'
-// CHECK: {{.*}}use_image2d_r 'void (__read_only image2d_t)'
-// CHECK: {{.*}}use_image3d_r 'void (__read_only image3d_t)'
-// CHECK: {{.*}}use_image1d_w 'void (__write_only image1d_t)'
-// CHECK: {{.*}}use_image2d_w 'void (__write_only image2d_t)'
-// CHECK: {{.*}}use_image3d_w 'void (__write_only image3d_t)'
+// CHECK: {{.*}}use_image1d_r 'void (__read_only image1d_t) __attribute__((device_kernel))'
+// CHECK: {{.*}}use_image2d_r 'void (__read_only image2d_t) __attribute__((device_kernel))'
+// CHECK: {{.*}}use_image3d_r 'void (__read_only image3d_t) __attribute__((device_kernel))'
+// CHECK: {{.*}}use_image1d_w 'void (__write_only image1d_t) __attribute__((device_kernel))'
+// CHECK: {{.*}}use_image2d_w 'void (__write_only image2d_t) __attribute__((device_kernel))'
+// CHECK: {{.*}}use_image3d_w 'void (__write_only image3d_t) __attribute__((device_kernel))'
 
 // TODO: SYCL specific fail - analyze and enable
 // XFAIL: target=x86_64-pc-windows-msvc
diff --git a/clang/test/SemaSYCL/accessors-targets.cpp b/clang/test/SemaSYCL/accessors-targets.cpp
index 6a06fc86687cf..4ab164754e1f7 100644
--- a/clang/test/SemaSYCL/accessors-targets.cpp
+++ b/clang/test/SemaSYCL/accessors-targets.cpp
@@ -51,7 +51,7 @@ int main() {
         });
   });
 }
-// CHECK: {{.*}}use_local_dep{{.*}} 'void (__local int *, sycl::range<1>, sycl::range<1>, sycl::id<1>)'
-// CHECK: {{.*}}use_local{{.*}} 'void (__local int *, sycl::range<1>, sycl::range<1>, sycl::id<1>)'
-// CHECK: {{.*}}use_global{{.*}} 'void (__global int *, sycl::range<1>, sycl::range<1>, sycl::id<1>)'
-// CHECK: {{.*}}use_constant{{.*}} 'void (__constant int *, sycl::range<1>, sycl::range<1>, sycl::id<1>)'
+// CHECK: {{.*}}use_local_dep{{.*}} 'void (__local int *, sycl::range<1>, sycl::range<1>, sycl::id<1>) __attribute__((device_kernel))'
+// CHECK: {{.*}}use_local{{.*}} 'void (__local int *, sycl::range<1>, sycl::range<1>, sycl::id<1>) __attribute__((device_kernel))'
+// CHECK: {{.*}}use_global{{.*}} 'void (__global int *, sycl::range<1>, sycl::range<1>, sycl::id<1>) __attribute__((device_kernel))'
+// CHECK: {{.*}}use_constant{{.*}} 'void (__constant int *, sycl::range<1>, sycl::range<1>, sycl::id<1>) __attribute__((device_kernel))'
diff --git a/clang/test/SemaSYCL/array-kernel-param.cpp b/clang/test/SemaSYCL/array-kernel-param.cpp
index d9b37c88b28dd..1f6a3b807d0b6 100644
--- a/clang/test/SemaSYCL/array-kernel-param.cpp
+++ b/clang/test/SemaSYCL/array-kernel-param.cpp
@@ -114,7 +114,7 @@ int main() {
 }
 
 // Check Kernel_Accessor parameters
-// CHECK: FunctionDecl {{.*}}Kernel_Accessor{{.*}} 'void (__global int *, sycl::range<1>, sycl::range<1>, sycl::id<1>, __global int *, sycl::range<1>, sycl::range<1>, sycl::id<1>)'
+// CHECK: FunctionDecl {{.*}}Kernel_Accessor{{.*}} 'void (__global int *, sycl::range<1>, sycl::range<1>, sycl::id<1>, __global int *, sycl::range<1>, sycl::range<1>, sycl::id<1>) __attribute__((device_kernel))'
 // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ReadWriteAccessor '__global int *'
 // CHECK-NEXT: SYCLAccessorPtrAttr
 // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ReadWriteAccessor 'sycl::range<1>'
@@ -131,7 +131,7 @@ int main() {
 // CHECK-NEXT: MemberExpr {{.*}}__init
 
 // Check Kernel_Array parameters
-// CHECK: FunctionDecl {{.*}}Kernel_Array{{.*}} 'void (__wrapper_class)'
+// CHECK: FunctionDecl {{.*}}Kernel_Array{{.*}} 'void (__wrapper_class) __attribute__((device_kernel))'
 // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_Array '__wrapper_class'
 // Check Kernel_Array inits
 // CHECK-NEXT: CompoundStmt
@@ -150,7 +150,7 @@ int main() {
 // CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_Array' '__wrapper_class'
 
 // Check Kernel_Array_Ptrs parameters
-// CHECK: FunctionDecl {{.*}}Kernel_Array_Ptrs{{.*}} 'void (__wrapper_class)'
+// CHECK: FunctionDecl {{.*}}Kernel_Array_Ptrs{{.*}} 'void (__wrapper_class) __attribute__((device_kernel))'
 // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ArrayOfPointers '__wrapper_class'
 // Check Kernel_Array_Ptrs inits
 // CHECK-NEXT: CompoundStmt
@@ -180,7 +180,7 @@ int main() {
 // CHECK-NEXT: IntegerLiteral {{.*}} 1
 
 // Check Kernel_StructAccArray parameters
-// CHECK: FunctionDecl {{.*}}Kernel_StructAccArray{{.*}} 'void (__global int *, sycl::range<1>, sycl::range<1>, sycl::id<1>, __global int *, sycl::range<1>, sycl::range<1>, sycl::id<1>)'
+// CHECK: FunctionDecl {{.*}}Kernel_StructAccArray{{.*}} 'void (__global int *, sycl::range<1>, sycl::range<1>, sycl::id<1>, __global int *, sycl::range<1>, sycl::range<1>, sycl::id<1>) __attribute__((device_kernel))'
 // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc '__global int *'
 // CHECK-NEXT: SYCLAccessorPtrAttr
 // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'sycl::range<1>'
@@ -207,7 +207,7 @@ int main() {
 // CHECK-NEXT: MemberExpr {{.*}}__init
 
 // Check Kernel_TemplatedStructArray parameters
-// CHECK: FunctionDecl {{.*}}Kernel_TemplatedStructArray{{.*}} 'void (S<int>)'
+// CHECK: FunctionDecl {{.*}}Kernel_TemplatedStructArray{{.*}} 'void (S<int>) __attribute__((device_kernel))'
 // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_s 'S<int>'
 // CHECK-NEXT: CompoundStmt
 // CHECK-NEXT: DeclStmt
@@ -218,7 +218,7 @@ int main() {
 // CHECK-NEXT: DeclRefExpr {{.*}} 'S<int>' lvalue ParmVar {{.*}} '_arg_s' 'S<int>'
 
 // Check Kernel_Array_2D parameters
-// CHECK: FunctionDecl {{.*}}Kernel_Array_2D{{.*}} 'void (__wrapper_class)'
+// CHECK: FunctionDecl {{.*}}Kernel_Array_2D{{.*}} 'void (__wrapper_class) __attribute__((device_kernel))'
 // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_array_2D '__wrapper_class'
 // Check Kernel_Array_2D inits
 // CHECK-NEXT: CompoundStmt
@@ -250,7 +250,7 @@ int main() {
 // CHECK-NEXT: ArrayInitIndexExpr {{.*}} 'unsigned
 
 // Check Kernel_NonDecomposedStruct parameters.
-// CHECK: FunctionDecl {{.*}}Kernel_NonDecomposedStruct{{.*}} 'void (__wrapper_class)'
+// CHECK: FunctionDecl {{.*}}Kernel_NonDecomposedStruct{{.*}} 'void (__wrapper_class) __attribute__((device_kernel))'
 // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_NonDecompStructArray '__wrapper_class'
 // Check Kernel_NonDecomposedStruct inits
 // CHECK-NEXT: CompoundStmt
@@ -271,7 +271,7 @@ int main() {
 // CHECK-NEXT: ArrayInitIndexExpr {{.*}} 'unsigned
 
 // Check Kernel_StructWithPointers parameters.
-// CHECK: FunctionDecl {{.*}}Kernel_StructWithPointers{{.*}} 'void (__wrapper_class)'
+// CHECK: FunctionDecl {{.*}}Kernel_StructWithPointers{{.*}} 'void (__wrapper_class) __attribute__((device_kernel))'
 // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_StructWithPointersArray '__wrapper_class'
 // Check Kernel_StructWithPointers inits
 // CHECK-NEXT: CompoundStmt
@@ -303,7 +303,7 @@ int main() {
 // CHECK-NEXT: IntegerLiteral {{.*}} 1
 
 // Check Kernel_Array_Ptrs_2D parameters
-// CHECK: FunctionDecl {{.*}}Kernel_Array_Ptrs_2D{{.*}} 'void (__wrapper_class, __wrapper_class)'
+// CHECK: FunctionDecl {{.*}}Kernel_Array_Ptrs_2D{{.*}} 'void (__wrapper_class, __wrapper_class) __attribute__((device_kernel))'
 // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ArrayOfPointers_2D '__wrapper_class'
 // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ArrayOfPointers '__wrapper_class'
 
diff --git a/clang/test/SemaSYCL/basic-kernel-wrapper.cpp b/clang/test/SemaSYCL/basic-kernel-wrapper.cpp
index 808b8f2d9d37d..f66c5f04b4b7c 100644
--- a/clang/test/SemaSYCL/basic-kernel-wrapper.cpp
+++ b/clang/test/SemaSYCL/basic-kernel-wrapper.cpp
@@ -20,7 +20,7 @@ int main() {
 
 // Check declaration of the kernel
 
-// CHECK: FunctionDecl {{.*}}kernel_wrapper{{.*}} 'void (__global int *, sycl::range<1>, sycl::range<1>, sycl::id<1>)'
+// CHECK: FunctionDecl {{.*}}kernel_wrapper{{.*}} 'void (__global int *, sycl::range<1>, sycl::range<1>, sycl::id<1>) __attribute__((device_kernel))'
 
 // Check parameters of the kernel
 
diff --git a/clang/test/SemaSYCL/binding_decl_lambda_nullptr.cpp b/clang/test/SemaSYCL/binding_decl_lambda_nullptr.cpp
index 9aca54588f388..f3c73bba2495f 100644
--- a/clang/test/SemaSYCL/binding_decl_lambda_nullptr.cpp
+++ b/clang/test/SemaSYCL/binding_decl_lambda_nullptr.cpp
@@ -18,5 +18,5 @@ void foo() {
   h.single_task<class C>(Lambda);
 }
 
-// CHECK: FunctionDecl {{.*}}foo{{.*}} 'void (int)'
+// CHECK: FunctionDecl {{.*}}foo{{.*}} 'void (int) __attribute__((device_kernel))'
 // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_bind_x 'int'
diff --git a/clang/test/SemaSYCL/built-in-type-kernel-arg.cpp b/clang/test/SemaSYCL/built-in-type-kernel-arg.cpp
index 5f3da924e982d..86cc871701e95 100644
--- a/clang/test/SemaSYCL/built-in-type-kernel-arg.cpp
+++ b/clang/test/SemaSYCL/built-in-type-kernel-arg.cpp
@@ -77,7 +77,7 @@ int main() {
   return 0;
 }
 // Check kernel parameters
-// CHECK: FunctionDecl {{.*}}kernel_const{{.*}} 'void (const int)'
+// CHECK: FunctionDecl {{.*}}kernel_const{{.*}} 'void (const int) __attribute__((device_kernel))'
 // CHECK: ParmVarDecl {{.*}} used _arg_some_const 'const int'
 
 // Check that lambda field of const built-in type is initialized
@@ -87,7 +87,7 @@ int main() {
 // CHECK-NEXT: DeclRefExpr {{.*}} 'const int' lvalue ParmVar {{.*}} '_arg_some_const' 'const int'
 
 // Check kernel parameters
-// CHECK: {{.*}}kernel_int{{.*}} 'void (int)'
+// CHECK: {{.*}}kernel_int{{.*}} 'void (int) __attribute__((device_kernel))'
 // CHECK: ParmVarDecl {{.*}} used _arg_data 'int'
 
 // Check that lambda field of built-in type is initialized
@@ -97,7 +97,7 @@ int main() {
 // CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_data' 'int'
 
 // Check kernel parameters
-// CHECK: {{.*}}kernel_struct{{.*}} 'void (__generated_test_struct)'
+// CHECK: {{.*}}kernel_struct{{.*}} 'void (__generated_test_struct) __attribute__((device_kernel))'
 // CHECK: ParmVarDecl {{.*}} used _arg_s '__generated_test_struct'
 
 // Check that lambda field of struct type is initialized
@@ -111,7 +111,7 @@ int main() {
 // CHECK-NEXT: DeclRefExpr {{.*}} '__generated_test_struct' lvalue ParmVar {{.*}} '_arg_s'
 
 // Check kernel parameters
-// CHECK: {{.*}}kernel_pointer{{.*}} 'void (__global int *, __global int *, __wrapper_class)'
+// CHECK: {{.*}}kernel_pointer{{.*}} 'void (__global int *, __global int *, __wrapper_class) __attribute__((device_kernel))'
 // CHECK: ParmVarDecl {{.*}} used _arg_new_data_addr '__global int *'
 // CHECK: ParmVarDecl {{.*}} used _arg_data_addr '__global int *'
 // CHECK: ParmVarDecl {{.*}} used _arg_ptr_array '__wrapper_class'
@@ -145,7 +145,7 @@ int main() {
 // CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ptr_array'
 // CHECK-NEXT: IntegerLiteral {{.*}} 1
 
-// CHECK: FunctionDecl {{.*}}kernel_nns{{.*}} 'void (__generated_test_struct_simple)'
+// CHECK: FunctionDecl {{.*}}kernel_nns{{.*}} 'void (__generated_test_struct_simple) __attribute__((device_kernel))'
 // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_tds '__generated_test_struct_simple'
 
 // CHECK: VarDecl {{.*}} used __SYCLKernel
diff --git a/clang/test/SemaSYCL/decomposition.cpp b/clang/test/SemaSYCL/decomposition.cpp
index f4699c3bd9ee8..91366ebfcde44 100644
--- a/clang/test/SemaSYCL/decomposition.cpp
+++ b/clang/test/SemaSYCL/decomposition.cpp
@@ -90,32 +90,32 @@ int main() {
   myQueue.submit([&](sycl::handler &h) {
     h.single_task<class NonDecomposed>([=]() { return SimpleStruct.i + ArrayOfSimpleStruct[0].i + NonDecompStruct.i + ArrayOfNonDecompStruct[0].i; });
   });
-  // CHECK: FunctionDecl {{.*}}NonDecomposed{{.*}} 'void (StructNonDecomposed, __wrapper_class, StructWithNonDecomposedStruct, __wrapper_class)'
+  // CHECK: FunctionDecl {{.*}}NonDecomposed{{.*}} 'void (StructNonDecomposed, __wrapper_class, StructWithNonDecomposedStruct, __wrapper_class) __attribute__((device_kernel))'
 
   {
     StructWithArray<StructWithAccessor> t1;
     myQueue.submit([&](sycl::handler &h) {
       h.single_task<class Acc1>([=]() { return t1.i; });
     });
-    // CHECK: FunctionDecl {{.*}}Acc1{{.*}} 'void (__global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>, __wrapper_class, __global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>, __wrapper_class, __global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>, __wrapper_class, StructNonDecomposed, int)'
+    // CHECK: FunctionDecl {{.*}}Acc1{{.*}} 'void (__global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>, __wrapper_class, __global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>, __wrapper_class, __global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>, __wrapper_class, StructNonDecomposed, int) __attribute__((device_kernel))'
 
     DerivedStruct<StructWithAccessor> t2;
     myQueue.submit([&](sycl::handler &h) {
       h.single_task<class Acc2>([=]() { return t2.i; });
     });
-    // CHECK: FunctionDecl {{.*}}Acc2{{.*}} 'void (__global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>, __wrapper_class, StructNonDecomposed, int)'
+    // CHECK: FunctionDecl {{.*}}Acc2{{.*}} 'void (__global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>, __wrapper_class, StructNonDecomposed, int) __attribute__((device_kernel))'
 
     StructWithArray<StructInheritedAccessor> t3;
     myQueue.submit([&](sycl::handler &h) {
       h.single_task<class Acc3>([=]() { return t3.i; });
     });
-    // CHECK: FunctionDecl {{.*}}Acc3{{.*}} 'void (__global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>, int, __global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>, int, __global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>, int, StructNonDecomposed, int)'
+    // CHECK: FunctionDecl {{.*}}Acc3{{.*}} 'void (__global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>, int, __global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>, int, __global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>, int, StructNonDecomposed, int) __attribute__((device_kernel))'
 
     DerivedStruct<StructInheritedAccessor> t4;
     myQueue.submit([&](sycl::handler &h) {
       h.single_task<class Acc4>([=]() { return t4.i; });
     });
-    // CHECK: FunctionDecl {{.*}}Acc4{{.*}} 'void (__global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>, int, StructNonDecomposed, int)'
+    // CHECK: FunctionDecl {{.*}}Acc4{{.*}} 'void (__global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>, int, StructNonDecomposed, int) __attribute__((device_kernel))'
   }
 
   {
@@ -123,13 +123,13 @@ int main() {
     myQueue.submit([&](sycl::handler &h) {
       h.single_task<class Sampl1>([=]() { return t1.i; });
     });
-    // CHECK: FunctionDecl {{.*}}Sampl1{{.*}} 'void (sampler_t, sampler_t, sampler_t, StructNonDecomposed, int)'
+    // CHECK: FunctionDecl {{.*}}Sampl1{{.*}} 'void (sampler_t, sampler_t, sampler_t, StructNonDecomposed, int) __attribute__((device_kernel))'
 
     DerivedStruct<StructWithSampler> t2;
     myQueue.submit([&](sycl::handler &h) {
       h.single_task<class Sampl2>([=]() { return t2.i; });
     });
-    // CHECK: FunctionDecl {{.*}}Sampl2{{.*}} 'void (sampler_t, StructNonDecomposed, int)'
+    // CHECK: FunctionDecl {{.*}}Sampl2{{.*}} 'void (sampler_t, StructNonDecomposed, int) __attribute__((device_kernel))'
   }
 
   {
@@ -137,12 +137,12 @@ int main() {
     myQueue.submit([&](sycl::handler &h) {
       h.single_task<class Stream1>([=]() { return t1.i; });
     });
-    // CHECK: FunctionDecl {{.*}}Stream1{{.*}} 'void (__global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>, int, __global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>, int, __global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>, int, StructNonDecomposed, int)'
+    // CHECK: FunctionDecl {{.*}}Stream1{{.*}} 'void (__global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>, int, __global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>, int, __global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>, int, StructNonDecomposed, int) __attribute__((device_kernel))'
     DerivedStruct<StructWithStream> t2;
     myQueue.submit([&](sycl::handler &h) {
       h.single_task<class Stream2>([=]() { return t2.i; });
     });
-    // CHECK: FunctionDecl {{.*}}Stream2{{.*}} 'void (__global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>, int, StructNonDecomposed, int)'
+    // CHECK: FunctionDecl {{.*}}Stream2{{.*}} 'void (__global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>, int, StructNonDecomposed, int) __attribute__((device_kernel))'
   }
 
   {
@@ -150,13 +150,13 @@ int main() {
     myQueue.submit([&](sycl::handler &h) {
       h.single_task<class Half1>([=]() { return t1.i; });
     });
-    // CHECK: FunctionDecl {{.*}}Half1{{.*}} 'void (StructWithArray<StructWithHalf>)'
+    // CHECK: FunctionDecl {{.*}}Half1{{.*}} 'void (StructWithArray<StructWithHalf>) __attribute__((device_kernel))'
 
     DerivedStruct<StructWithHalf> t2;
     myQueue.submit([&](sycl::handler &h) {
       h.single_task<class Half2>([=]() { return t2.i; });
     });
-    // CHECK: FunctionDecl {{.*}}Half2{{.*}} 'void (DerivedStruct<StructWithHalf>)'
+    // CHECK: FunctionDecl {{.*}}Half2{{.*}} 'void (DerivedStruct<StructWithHalf>) __attribute__((device_kernel))'
   }
 
   {
@@ -164,25 +164,25 @@ int main() {
     myQueue.submit([&](sycl::handler &h) {
       h.single_task<class Pointer>([=]() { return SimpleStructWithPtr.i; });
     });
-    // CHECK: FunctionDecl {{.*}}Pointer{{.*}} 'void (__generated_StructWithPtr)'
+    // CHECK: FunctionDecl {{.*}}Pointer{{.*}} 'void (__generated_StructWithPtr) __attribute__((device_kernel))'
 
     Nested::TDStrWithPTR TDStructWithPtr;
     myQueue.submit([&](sycl::handler &h) {
       h.single_task<class TDStr>([=]() { return TDStructWithPtr.i; });
     });
-    // CHECK: FunctionDecl {{.*}}TDStr{{.*}} 'void (__generated_StructWithPtr)'
+    // CHECK: FunctionDecl {{.*}}TDStr{{.*}} 'void (__generated_StructWithPtr) __attribute__((device_kernel))'
 
     StructWithArray<StructWithPtr> t1;
     myQueue.submit([&](sycl::handler &h) {
       h.single_task<class NestedArrayOfStructWithPointer>([=]() { return t1.i; });
     });
-    // CHECK: FunctionDecl {{.*}}NestedArrayOfStructWithPointer{{.*}} 'void (__generated_StructWithArray)'
+    // CHECK: FunctionDecl {{.*}}NestedArrayOfStructWithPointer{{.*}} 'void (__generated_StructWithArray) __attribute__((device_kernel))'
 
     DerivedStruct<StructWithPtr> t2;
     myQueue.submit([&](sycl::handler &h) {
       h.single_task<class PointerInBase>([=]() { return t2.i; });
     });
-    // CHECK: FunctionDecl {{.*}}PointerInBase{{.*}} 'void (__generated_DerivedStruct)'
+    // CHECK: FunctionDecl {{.*}}PointerInBase{{.*}} 'void (__generated_DerivedStruct) __attribute__((device_kernel))'
   }
 
   {
@@ -190,18 +190,18 @@ int main() {
     myQueue.submit([&](sycl::handler &h) {
       h.single_task<class NonTrivial>([=]() { return NonTrivialStructWithPtr.i;});
     });
-    // CHECK: FunctionDecl {{.*}}NonTrivial{{.*}} 'void (__generated_NonTrivialType)'
+    // CHECK: FunctionDecl {{.*}}NonTrivial{{.*}} 'void (__generated_NonTrivialType) __attribute__((device_kernel))'
 
     NonTrivialType NonTrivialTypeArray[2]{0,0};
     myQueue.submit([&](sycl::handler &h) {
       h.single_task<class ArrayOfNonTrivialStruct>([=]() { return NonTrivialTypeArray[0].i;});
     });
-    // CHECK: FunctionDecl {{.*}}ArrayOfNonTrivialStruct{{.*}} 'void (__wrapper_class)'
+    // CHECK: FunctionDecl {{.*}}ArrayOfNonTrivialStruct{{.*}} 'void (__wrapper_class) __attribute__((device_kernel))'
 
     NonTrivialDerived NonTrivialDerivedStructWithPtr(10);
     myQueue.submit([&](sycl::handler &h) {
       h.single_task<class NonTrivialStructInBase>([=]() { return NonTrivialDerivedStructWithPtr.i;});
     });
-    // CHECK: FunctionDecl {{.*}}NonTrivialStructInBase{{.*}} 'void (__generated_NonTrivialDerived)'
+    // CHECK: FunctionDecl {{.*}}NonTrivialStructInBase{{.*}} 'void (__generated_NonTrivialDerived) __attribute__((device_kernel))'
   }
 }
diff --git a/clang/test/SemaSYCL/free_function_array_kernel_param.cpp b/clang/test/SemaSYCL/free_function_array_kernel_param.cpp
index cbb62eb59fa89..322e495dcabcb 100755
--- a/clang/test/SemaSYCL/free_function_array_kernel_param.cpp
+++ b/clang/test/SemaSYCL/free_function_array_kernel_param.cpp
@@ -26,7 +26,7 @@ void ff_6(KArgWithPtrArray<ArrSize> KArg) {
 
 template void ff_6(KArgWithPtrArray<TestArrSize> KArg);
 
-// CHECK: FunctionDecl {{.*}}__sycl_kernel{{.*}}'void (__generated_KArgWithPtrArray)'
+// CHECK: FunctionDecl {{.*}}__sycl_kernel{{.*}}'void (__generated_KArgWithPtrArray) __attribute__((device_kernel))'
 // CHECK-NEXT: ParmVarDecl {{.*}} used __arg_KArg '__generated_KArgWithPtrArray'
 // CHECK-NEXT: CompoundStmt
 // CHECK-NEXT: CallExpr
diff --git a/clang/test/SemaSYCL/free_function_kernel_params.cpp b/clang/test/SemaSYCL/free_function_kernel_params.cpp
index e0f1e0665d551..3f6e27acd8843 100644
--- a/clang/test/SemaSYCL/free_function_kernel_params.cpp
+++ b/clang/test/SemaSYCL/free_function_kernel_params.cpp
@@ -12,7 +12,7 @@ void ff_2(int *ptr, int start, int end) {
   for (int i = start; i <= end; i++)
     ptr[i] = start;
 }
-// CHECK: FunctionDecl {{.*}}__sycl_kernel_{{.*}} 'void (__global int *, int, int)'
+// CHECK: FunctionDecl {{.*}}__sycl_kernel_{{.*}} 'void (__global int *, int, int) __attribute__((device_kernel))'
 // CHECK-NEXT: ParmVarDecl {{.*}} __arg_ptr '__global int *'
 // CHECK-NEXT: ParmVarDecl {{.*}} __arg_start 'int'
 // CHECK-NEXT: ParmVarDecl {{.*}} __arg_end 'int'
@@ -41,7 +41,7 @@ __attribute__((sycl_device))
 // Explicit instantiation with "int*"
 template void ff_3(int* ptr, int start, int end);
 
-// CHECK: FunctionDecl {{.*}}__sycl_kernel_{{.*}} 'void (__global int *, int, int)'
+// CHECK: FunctionDecl {{.*}}__sycl_kernel_{{.*}} 'void (__global int *, int, int) __attribute__((device_kernel))'
 // CHECK-NEXT: ParmVarDecl {{.*}} __arg_ptr '__global int *'
 // CHECK-NEXT: ParmVarDecl {{.*}} __arg_start 'int'
 // CHECK-NEXT: ParmVarDecl {{.*}} __arg_end 'int'
@@ -90,7 +90,7 @@ __attribute__((sycl_device))
 [[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 0)]]
 void ff_4(NoPointers S1, Pointers S2, Agg S3) {
 }
-// CHECK: FunctionDecl {{.*}}__sycl_kernel{{.*}}'void (NoPointers, __generated_Pointers, __generated_Agg)'
+// CHECK: FunctionDecl {{.*}}__sycl_kernel{{.*}}'void (NoPointers, __generated_Pointers, __generated_Agg) __attribute__((device_kernel))'
 // CHECK-NEXT: ParmVarDecl {{.*}} used __arg_S1 'NoPointers'
 // CHECK-NEXT: ParmVarDecl {{.*}} used __arg_S2 '__generated_Pointers'
 // CHECK-NEXT: ParmVarDecl {{.*}} used __arg_S3 '__generated_Agg'
@@ -118,7 +118,7 @@ __attribute__((sycl_device))
 [[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 0)]]
 void ff_5(Agg1 S1, Derived S2, Derived1 S3) {
 }
-// CHECK: FunctionDecl {{.*}}__sycl_kernel{{.*}}'void (Agg1, __generated_Derived, __generated_Derived1)'
+// CHECK: FunctionDecl {{.*}}__sycl_kernel{{.*}}'void (Agg1, __generated_Derived, __generated_Derived1) __attribute__((device_kernel))'
 // CHECK-NEXT: ParmVarDecl {{.*}} used __arg_S1 'Agg1'
 // CHECK-NEXT: ParmVarDecl {{.*}} used __arg_S2 '__generated_Derived'
 // CHECK-NEXT: ParmVarDecl {{.*}} used __arg_S3 '__generated_Derived1'
@@ -150,7 +150,7 @@ __attribute__((sycl_device))
 
 // Explicit instantiation.
 template void ff_6(Agg S1, Derived1 S2, int);
-// CHECK: FunctionDecl {{.*}}__sycl_kernel{{.*}}'void (__generated_Agg, __generated_Derived1, int)'
+// CHECK: FunctionDecl {{.*}}__sycl_kernel{{.*}}'void (__generated_Agg, __generated_Derived1, int) __attribute__((device_kernel))'
 // CHECK-NEXT: ParmVarDecl {{.*}} used __arg_S1 '__generated_Agg'
 // CHECK-NEXT: ParmVarDecl {{.*}} used __arg_S2 '__generated_Derived1'
 // CHECK-NEXT: ParmVarDecl {{.*}} used __arg_end 'int'
@@ -177,7 +177,7 @@ __attribute__((sycl_device))
 [[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 0)]]
 void ff_7(sycl::work_group_memory<int> mem) {
 }
-// CHECK: FunctionDecl {{.*}}__sycl_kernel{{.*}}'void (__local int *)'
+// CHECK: FunctionDecl {{.*}}__sycl_kernel{{.*}}'void (__local int *) __attribute__((device_kernel))'
 // CHECK-NEXT: ParmVarDecl {{.*}} used __arg_Ptr '__local int *'
 // CHECK-NEXT: CompoundStmt
 // CHECK-NEXT: DeclStmt
@@ -197,7 +197,7 @@ __attribute__((sycl_device))
 [[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 0)]]
 void ff_8(sycl::dynamic_work_group_memory<int> DynMem) {
 }
-// CHECK: FunctionDecl {{.*}}__sycl_kernel{{.*}}'void (__local int *)'
+// CHECK: FunctionDecl {{.*}}__sycl_kernel{{.*}}'void (__local int *) __attribute__((device_kernel))'
 // CHECK-NEXT: ParmVarDecl {{.*}} used __arg_Ptr '__local int *'
 // CHECK-NEXT: CompoundStmt
 // CHECK-NEXT: DeclStmt
@@ -266,7 +266,7 @@ void ff_10(sycl::sampler S) {
 
 // CHECK: FunctionDecl {{.*}}'void (sycl::sampler)'
 // CHECK: ParmVarDecl {{.*}}S 'sycl::sampler'
-// CHECK: FunctionDecl {{.*}}'void (sampler_t)'
+// CHECK: FunctionDecl {{.*}}'void (sampler_t) __attribute__((device_kernel))'
 // CHECK: ParmVarDecl {{.*}}__arg_Sampler 'sampler_t'
 // CHECK: CXXMemberCallExpr {{.*}}'void'
 // CHECK-NEXT: MemberExpr {{.*}}.__init
@@ -279,7 +279,7 @@ void ff_11(sycl::stream str) {
 
 // CHECK: FunctionDecl {{.*}}'void (sycl::stream)'
 // CHECK: ParmVarDecl {{.*}}str 'sycl::stream'
-// CHECK: FunctionDecl {{.*}}'void (__global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>, int)'
+// CHECK: FunctionDecl {{.*}}'void (__global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>, int) __attribute__((device_kernel))'
 // CHECK: ParmVarDecl {{.*}}__arg_Ptr '__global char *'
 // CHECK: ParmVarDecl {{.*}}__arg_AccessRange 'sycl::range<1>'
 // CHECK: ParmVarDecl {{.*}}__arg_MemRange 'sycl::range<1>'
@@ -296,7 +296,7 @@ void ff_12(sycl::ext::oneapi::experimental::annotated_arg<int> arg) {
 
 // CHECK: FunctionDecl {{.*}}'void (sycl::ext::oneapi::experimental::annotated_arg<int>)'
 // CHECK: ParmVarDecl {{.*}}arg 'sycl::ext::oneapi::experimental::annotated_arg<int>'
-// CHECK: FunctionDecl {{.*}}'void (int)'
+// CHECK: FunctionDecl {{.*}}'void (int) __attribute__((device_kernel))'
 // CHECK: ParmVarDecl {{.*}}__arg__obj 'int'
 // CHECK: CXXMemberCallExpr {{.*}}
 // CHECK-NEXT: MemberExpr {{.*}}.__init
@@ -309,7 +309,7 @@ void ff_13(sycl::ext::oneapi::experimental::annotated_ptr<int> ptr) {
 
 // CHECK: FunctionDecl {{.*}}'void (sycl::ext::oneapi::experimental::annotated_ptr<int>)'
 // CHECK: ParmVarDecl {{.*}}ptr 'sycl::ext::oneapi::experimental::annotated_ptr<int>'
-// CHECK: FunctionDecl {{.*}}'void (int *)'
+// CHECK: FunctionDecl {{.*}}'void (int *) __attribute__((device_kernel))'
 // CHECK: ParmVarDecl {{.*}}__arg__obj 'int *'
 // CHECK: CXXMemberCallExpr {{.*}}
 // CHECK-NEXT: MemberExpr {{.*}}.__init
diff --git a/clang/test/SemaSYCL/half-kernel-arg.cpp b/clang/test/SemaSYCL/half-kernel-arg.cpp
index b2a301782d2e0..e0a0cec82ac0e 100644
--- a/clang/test/SemaSYCL/half-kernel-arg.cpp
+++ b/clang/test/SemaSYCL/half-kernel-arg.cpp
@@ -17,7 +17,7 @@ int main() {
   });
 }
 
-// CHECK: {{.*}}kernel_half{{.*}} 'void (sycl::half)'
+// CHECK: {{.*}}kernel_half{{.*}} 'void (sycl::half) __attribute__((device_kernel))'
 // CHECK: ParmVarDecl {{.*}} used _arg_HostHalf 'sycl::half':'sycl::detail::half_impl::half'
 // // Check that lambda field of half type is initialized
 // CHECK: VarDecl {{.*}}'(lambda at {{.*}}'
diff --git a/clang/test/SemaSYCL/intel-max-global-work-dim-device-ast.cpp b/clang/test/SemaSYCL/intel-max-global-work-dim-device-ast.cpp
index b8a171c832896..d1c0884e309f2 100644
--- a/clang/test/SemaSYCL/intel-max-global-work-dim-device-ast.cpp
+++ b/clang/test/SemaSYCL/intel-max-global-work-dim-device-ast.cpp
@@ -300,7 +300,7 @@ int main() {
 
     // Ignore duplicate attribute with same argument value.
     h.single_task<class test_kernell2>(
-        // CHECK-LABEL: FunctionDecl {{.*}}test_kernell2 'void ()'
+        // CHECK-LABEL: FunctionDecl {{.*}}test_kernell2 'void () __attribute__((device_kernel))'
         // CHECK:       SYCLIntelMaxGlobalWorkDimAttr
         // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
         // CHECK-NEXT:  value: Int 3
diff --git a/clang/test/SemaSYCL/kernel-handler.cpp b/clang/test/SemaSYCL/kernel-handler.cpp
index ec9644a3bec24..979adadbd47ed 100644
--- a/clang/test/SemaSYCL/kernel-handler.cpp
+++ b/clang/test/SemaSYCL/kernel-handler.cpp
@@ -28,7 +28,7 @@ int main() {
 }
 
 // Check test_kernel_handler parameters
-// NONATIVESUPPORT: FunctionDecl {{.*}}test_kernel_handler{{.*}} 'void ((lambda at {{.*}}kernel-handler.cpp{{.*}}), __global char *)'
+// NONATIVESUPPORT: FunctionDecl {{.*}}test_kernel_handler{{.*}} 'void ((lambda at {{.*}}kernel-handler.cpp{{.*}}), __global char *) __attribute__((device_kernel))'
 // NONATIVESUPPORT-NEXT: ParmVarDecl {{.*}} used _arg__sycl_functor '(lambda at {{.*}}'
 // NONATIVESUPPORT-NEXT: ParmVarDecl {{.*}} used  _arg__specialization_constants_buffer '__global char *'
 
@@ -57,7 +57,7 @@ int main() {
 // NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} 'kernel_handler':'sycl::kernel_handler' lvalue Var {{.*}} 'kh' 'kernel_handler':'sycl::kernel_handler'
 
 // Check test_pfwg_kernel_handler parameters
-// NONATIVESUPPORT: FunctionDecl {{.*}}test_pfwg_kernel_handler{{.*}} 'void ((lambda at {{.*}}kernel-handler.cpp{{.*}}), __global char *)'
+// NONATIVESUPPORT: FunctionDecl {{.*}}test_pfwg_kernel_handler{{.*}} 'void ((lambda at {{.*}}kernel-handler.cpp{{.*}}), __global char *) __attribute__((device_kernel))'
 // NONATIVESUPPORT-NEXT: ParmVarDecl {{.*}} used  _arg__sycl_functor '(lambda at {{.*}}kernel-handler.cpp{{.*}})'
 // NONATIVESUPPORT-NEXT: SYCLScopeAttr {{.*}} Implicit WorkGroup
 // NONATIVESUPPORT-NEXT: ParmVarDecl {{.*}} used  _arg__specialization_constants_buffer '__global char *'
@@ -91,7 +91,7 @@ int main() {
 // Test AST for default SPIR architecture
 
 // Check test_kernel_handler parameters
-// NATIVESUPPORT: FunctionDecl {{.*}}test_kernel_handler{{.*}} 'void (int, __global char *)'
+// NATIVESUPPORT: FunctionDecl {{.*}}test_kernel_handler{{.*}} 'void (int, __global char *) __attribute__((device_kernel))'
 // NATIVESUPPORT-NEXT: ParmVarDecl {{.*}} used _arg_a 'int'
 // NATIVESUPPORT-NEXT: ParmVarDecl {{.*}} used  _arg__specialization_constants_buffer '__global char *'
 
diff --git a/clang/test/SemaSYCL/kernel_functor_location.cpp b/clang/test/SemaSYCL/kernel_functor_location.cpp
index f5111630e74d5..00422ec025e2b 100644
--- a/clang/test/SemaSYCL/kernel_functor_location.cpp
+++ b/clang/test/SemaSYCL/kernel_functor_location.cpp
@@ -11,7 +11,7 @@ struct Functor {
   void operator()() const {}
 };
 
-// CHECK: FunctionDecl {{.*}} _ZTS7Functor 'void ()'
+// CHECK: FunctionDecl {{.*}} _ZTS7Functor 'void () __attribute__((device_kernel))'
 // CHECK-NEXT: |-CompoundStmt {{.*}} <{{.*}}line{{.*}}>
 
 int main() {
diff --git a/clang/test/SemaSYCL/no-decomp.cpp b/clang/test/SemaSYCL/no-decomp.cpp
index 5f61574935c8b..b9fc897e5f91d 100644
--- a/clang/test/SemaSYCL/no-decomp.cpp
+++ b/clang/test/SemaSYCL/no-decomp.cpp
@@ -37,20 +37,20 @@ int main() {
     wrapping_acc acc;
     cgh.single_task(acc);
   });
-  // ALL: FunctionDecl {{.*}} _ZTS12wrapping_acc 'void (__wrapper_class, __global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>)'
+  // ALL: FunctionDecl {{.*}} _ZTS12wrapping_acc 'void (__wrapper_class, __global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>) __attribute__((device_kernel))'
 
   q.submit([&](sycl::handler &cgh) {
     pointer_wrap ptr;
     cgh.single_task(ptr);
   });
-  // NODECOMP: FunctionDecl {{.*}} _ZTS12pointer_wrap 'void (pointer_wrap)'
-  // DECOMP: FunctionDecl {{.*}} _ZTS12pointer_wrap 'void (__global int *)'
+  // NODECOMP: FunctionDecl {{.*}} _ZTS12pointer_wrap 'void (pointer_wrap) __attribute__((device_kernel))'
+  // DECOMP: FunctionDecl {{.*}} _ZTS12pointer_wrap 'void (__global int *) __attribute__((device_kernel))'
 
   q.submit([&](sycl::handler &cgh) {
     empty e;
     cgh.single_task(e);
   });
-  // ALL: FunctionDecl {{.*}} _ZTS5empty 'void ()'
+  // ALL: FunctionDecl {{.*}} _ZTS5empty 'void () __attribute__((device_kernel))'
 
   return 0;
 }
diff --git a/clang/test/SemaSYCL/sampler.cpp b/clang/test/SemaSYCL/sampler.cpp
index 5d812f8ae8012..83dfcea06c312 100644
--- a/clang/test/SemaSYCL/sampler.cpp
+++ b/clang/test/SemaSYCL/sampler.cpp
@@ -20,7 +20,7 @@ int main() {
 }
 
 // Check declaration of the test kernel
-// CHECK: FunctionDecl {{.*}}SamplerLambda{{.*}} 'void (sampler_t)'
+// CHECK: FunctionDecl {{.*}}SamplerLambda{{.*}} 'void (sampler_t) __attribute__((device_kernel))'
 //
 // Check parameters of the test kernel
 // CHECK: ParmVarDecl {{.*}} used [[_arg_sampler:[0-9a-zA-Z_]+]] 'sampler_t'
diff --git a/clang/test/SemaSYCL/union-kernel-param.cpp b/clang/test/SemaSYCL/union-kernel-param.cpp
index 312c0b8f05009..4fbc4aaa56502 100644
--- a/clang/test/SemaSYCL/union-kernel-param.cpp
+++ b/clang/test/SemaSYCL/union-kernel-param.cpp
@@ -24,7 +24,7 @@ int main() {
 }
 
 // Check kernel parameters
-// CHECK: FunctionDecl {{.*}}kernel{{.*}} 'void (MyUnion)'
+// CHECK: FunctionDecl {{.*}}kernel{{.*}} 'void (MyUnion) __attribute__((device_kernel))'
 // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_accel 'MyUnion'
 
 // Check kernel inits
diff --git a/clang/test/SemaSYCL/union-kernel-param1.cpp b/clang/test/SemaSYCL/union-kernel-param1.cpp
index aef01a26b140a..c3567fb40cc8a 100644
--- a/clang/test/SemaSYCL/union-kernel-param1.cpp
+++ b/clang/test/SemaSYCL/union-kernel-param1.cpp
@@ -33,7 +33,7 @@ int main() {
 }
 
 // Check kernel_A parameters
-// CHECK: FunctionDecl {{.*}}kernel_A{{.*}} 'void (union union_acc_t)'
+// CHECK: FunctionDecl {{.*}}kernel_A{{.*}} 'void (union union_acc_t) __attribute__((device_kernel))'
 // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_union_acc 'union union_acc_t':'union_acc_t'
 
 // Check kernel_A inits
@@ -46,7 +46,7 @@ int main() {
 // CHECK: DeclRefExpr {{.*}} 'union union_acc_t':'union_acc_t' lvalue ParmVar {{.*}} '_arg_union_acc' 'union union_acc_t':'union_acc_t'
 
 // Check kernel_B parameters
-// CHECK: FunctionDecl {{.*}}kernel_B{{.*}} 'void (S<int>)'
+// CHECK: FunctionDecl {{.*}}kernel_B{{.*}} 'void (S<int>) __attribute__((device_kernel))'
 // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_s 'S<int>'
 
 // Check kernel_B inits
diff --git a/clang/test/SemaSYCL/union-kernel-param2.cpp b/clang/test/SemaSYCL/union-kernel-param2.cpp
index 920d73603e6c1..590f220f161cf 100644
--- a/clang/test/SemaSYCL/union-kernel-param2.cpp
+++ b/clang/test/SemaSYCL/union-kernel-param2.cpp
@@ -59,7 +59,7 @@ int main() {
 }
 
 // Check kernel_A parameters
-// CHECK: FunctionDecl {{.*}}kernel_A{{.*}} 'void (union MyUnion)'
+// CHECK: FunctionDecl {{.*}}kernel_A{{.*}} 'void (union MyUnion) __attribute__((device_kernel))'
 // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_union_mem 'union MyUnion':'MyUnion'
 
 // Check kernel_A inits
@@ -72,7 +72,7 @@ int main() {
 // CHECK-NEXT: DeclRefExpr {{.*}} 'union MyUnion':'MyUnion' lvalue ParmVar {{.*}} '_arg_union_mem' 'union MyUnion':'MyUnion'
 
 // Check kernel_B parameters
-// CHECK: FunctionDecl {{.*}}kernel_B{{.*}} 'void (union MyUnion, __global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>)'
+// CHECK: FunctionDecl {{.*}}kernel_B{{.*}} 'void (union MyUnion, __global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>) __attribute__((device_kernel))'
 // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_union_mem 'union MyUnion':'MyStruct::MyUnion'
 // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_AccField '__global char *'
 // CHECK: ParmVarDecl {{.*}} used _arg_AccField 'sycl::range<1>'
@@ -98,7 +98,7 @@ int main() {
 // CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}union-kernel-param2.cpp:48:9)' lvalue Var {{.*}} '__SYCLKernel' '(lambda at {{.*}}union-kernel-param2.cpp:48:9)'
 
 // Check kernel_C parameters
-// CHECK: FunctionDecl {{.*}}kernel_C{{.*}} 'void (__generated_MyStructWithPtr)'
+// CHECK: FunctionDecl {{.*}}kernel_C{{.*}} 'void (__generated_MyStructWithPtr) __attribute__((device_kernel))'
 // CHECK-NEXT: ParmVarDecl {{.*}} used _arg_structWithPtr_mem '__generated_MyStructWithPtr'
 
 // Check kernel_C inits
diff --git a/clang/test/SemaSYCL/work_group_memory_inheritance.cpp b/clang/test/SemaSYCL/work_group_memory_inheritance.cpp
index 6f1f6badbdc59..1216254726a7e 100644
--- a/clang/test/SemaSYCL/work_group_memory_inheritance.cpp
+++ b/clang/test/SemaSYCL/work_group_memory_inheritance.cpp
@@ -20,7 +20,7 @@ int main() {
   return 0;
 }
 
-// CHECK: FunctionDecl {{.*}}kernel{{.*}} 'void (__local int *)'
+// CHECK: FunctionDecl {{.*}}kernel{{.*}} 'void (__local int *) __attribute__((device_kernel))'
 // CHECK-NEXT: ParmVarDecl {{.*}}used _arg__base '__local int *'
 // CHECK-NEXT: CompoundStmt {{.*}}
 // CHECK-NEXT: DeclStmt {{.*}}
diff --git a/clang/test/SemaSYCL/wrapped-accessor.cpp b/clang/test/SemaSYCL/wrapped-accessor.cpp
index ba2bb59cc79c6..87165a8e11263 100644
--- a/clang/test/SemaSYCL/wrapped-accessor.cpp
+++ b/clang/test/SemaSYCL/wrapped-accessor.cpp
@@ -25,7 +25,7 @@ int main() {
 }
 
 // Check declaration of the kernel
-// CHECK: wrapped_access{{.*}} 'void (__global int *, sycl::range<1>, sycl::range<1>, sycl::id<1>)'
+// CHECK: wrapped_access{{.*}} 'void (__global int *, sycl::range<1>, sycl::range<1>, sycl::id<1>) __attribute__((device_kernel))'
 
 // Check parameters of the kernel
 // CHECK: ParmVarDecl {{.*}} used _arg_accessor '__global int *'
diff --git a/llvm/test/Transforms/SPIRITTAnnotations/itt_instrumentation_debug.ll b/llvm/test/Transforms/SPIRITTAnnotations/itt_instrumentation_debug.ll
index 5917cd0aba910..c8907fee2895e 100644
--- a/llvm/test/Transforms/SPIRITTAnnotations/itt_instrumentation_debug.ll
+++ b/llvm/test/Transforms/SPIRITTAnnotations/itt_instrumentation_debug.ll
@@ -90,7 +90,7 @@ attributes #2 = { convergent }
 !11 = !{i32 3, i32 200000}
 !12 = !{i32 2, i32 2}
 !13 = distinct !DISubprogram(name: "test", scope: !6, file: !6, line: 4, type: !14, scopeLine: 4, flags: DIFlagPrototyped, spFlags: DISPFlagDefinition, unit: !2, retainedNodes: !4)
-!14 = !DISubroutineType(cc: DW_CC_LLVM_OpenCLKernel, types: !15)
+!14 = !DISubroutineType(cc: DW_CC_LLVM_DeviceKernel, types: !15)
 !15 = !{null}
 !16 = !DILocation(line: 5, column: 3, scope: !13)
 !17 = !DILocation(line: 6, column: 3, scope: !13)