diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp
index 93fbf18a642f..c6d00cea785a 100644
--- a/clang/lib/CodeGen/CodeGenModule.cpp
+++ b/clang/lib/CodeGen/CodeGenModule.cpp
@@ -2592,40 +2592,39 @@ void CodeGenModule::GenKernelArgMetadata(llvm::Function *Fn,
       }
     }
 
-  bool IsEsimdFunction = FD && FD->hasAttr<SYCLSimdAttr>();
-
-  if (LangOpts.SYCLIsDevice && !IsEsimdFunction) {
-    Fn->setMetadata("kernel_arg_buffer_location",
-                    llvm::MDNode::get(VMContext, argSYCLBufferLocationAttr));
-    // Generate this metadata only if atleast one kernel argument is an
-    // accessor.
-    if (isKernelArgAnAccessor) {
-      Fn->setMetadata("kernel_arg_runtime_aligned",
-                      llvm::MDNode::get(VMContext, argSYCLAccessorPtrs));
-      Fn->setMetadata("kernel_arg_exclusive_ptr",
+  if (getLangOpts().SYCLIsDevice) {
+    if (FD && FD->hasAttr<SYCLSimdAttr>()) {
+      Fn->setMetadata("kernel_arg_accessor_ptr",
                       llvm::MDNode::get(VMContext, argSYCLAccessorPtrs));
-    }
-  } else {
-    if (getLangOpts().OpenCL || getLangOpts().SYCLIsDevice) {
-      Fn->setMetadata("kernel_arg_addr_space",
-                      llvm::MDNode::get(VMContext, addressQuals));
-      Fn->setMetadata("kernel_arg_access_qual",
-                      llvm::MDNode::get(VMContext, accessQuals));
-      Fn->setMetadata("kernel_arg_type",
-                      llvm::MDNode::get(VMContext, argTypeNames));
-      Fn->setMetadata("kernel_arg_base_type",
-                      llvm::MDNode::get(VMContext, argBaseTypeNames));
-      Fn->setMetadata("kernel_arg_type_qual",
-                      llvm::MDNode::get(VMContext, argTypeQuals));
-      if (IsEsimdFunction)
-        Fn->setMetadata("kernel_arg_accessor_ptr",
+    } else {
+      Fn->setMetadata("kernel_arg_buffer_location",
+                      llvm::MDNode::get(VMContext, argSYCLBufferLocationAttr));
+      // Generate this metadata only if at least one kernel argument is an
+      // accessor.
+      if (isKernelArgAnAccessor) {
+        Fn->setMetadata("kernel_arg_runtime_aligned",
                         llvm::MDNode::get(VMContext, argSYCLAccessorPtrs));
+        Fn->setMetadata("kernel_arg_exclusive_ptr",
+                        llvm::MDNode::get(VMContext, argSYCLAccessorPtrs));
+      }
     }
-    if (getCodeGenOpts().EmitOpenCLArgMetadata ||
-        getCodeGenOpts().HIPSaveKernelArgName)
-      Fn->setMetadata("kernel_arg_name",
-                      llvm::MDNode::get(VMContext, argNames));
   }
+
+  if (getLangOpts().OpenCL) {
+    Fn->setMetadata("kernel_arg_addr_space",
+                    llvm::MDNode::get(VMContext, addressQuals));
+    Fn->setMetadata("kernel_arg_access_qual",
+                    llvm::MDNode::get(VMContext, accessQuals));
+    Fn->setMetadata("kernel_arg_type",
+                    llvm::MDNode::get(VMContext, argTypeNames));
+    Fn->setMetadata("kernel_arg_base_type",
+                    llvm::MDNode::get(VMContext, argBaseTypeNames));
+    Fn->setMetadata("kernel_arg_type_qual",
+                    llvm::MDNode::get(VMContext, argTypeQuals));
+  }
+  if (getCodeGenOpts().EmitOpenCLArgMetadata ||
+      getCodeGenOpts().HIPSaveKernelArgName)
+    Fn->setMetadata("kernel_arg_name", llvm::MDNode::get(VMContext, argNames));
 }
 
 /// Determines whether the language options require us to model
diff --git a/clang/test/CodeGenSYCL/kernel-op-calls.cpp b/clang/test/CodeGenSYCL/kernel-op-calls.cpp
index 0c3a53586b46..aa5de944c129 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_addr_space !{{[0-9]+}} !kernel_arg_access_qual !{{[0-9]+}} !kernel_arg_type !{{[0-9]+}} !kernel_arg_base_type !{{[0-9]+}} !kernel_arg_type_qual !{{[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() {{.*}} !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]+}} {
       cgh.parallel_for(sycl::range<1>(10), EF);
     });
 
diff --git a/sycl/test/check_device_code/esimd/dae.cpp b/sycl/test/check_device_code/esimd/dae.cpp
index 2f0142395f4c..222a8628647a 100644
--- a/sycl/test/check_device_code/esimd/dae.cpp
+++ b/sycl/test/check_device_code/esimd/dae.cpp
@@ -15,10 +15,9 @@ __attribute__((sycl_kernel)) void my_kernel(Func kernelFunc) {
 
 SYCL_EXTERNAL SYCL_ESIMD_FUNCTION ESIMD_NOINLINE void callee(int x) {}
 
-// CHECK: define dso_local spir_kernel {{.*}} !kernel_arg_addr_space ![[#MD:]]
-// CHECK: !kernel_arg_access_qual ![[#MD]] !kernel_arg_type ![[#MD]] !kernel_arg_base_type ![[#MD]] !kernel_arg_type_qual ![[#MD]] !kernel_arg_accessor_ptr ![[#MD]]
+// CHECK: define dso_local spir_kernel {{.*}} !sycl_kernel_omit_args ![[#MD:]]
 SYCL_EXTERNAL void __attribute__((noinline)) caller(int x) {
   my_kernel<class kernel_abc>([=]() SYCL_ESIMD_KERNEL { callee(x); });
 }
 
-//CHECK: [[#MD]] = !{}
+//CHECK: [[#MD]] = !{i1 true}