diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 8f45e45dc3f8f..bd2b58d809529 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -4684,6 +4684,10 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
 
   for (const KernelDesc &K : KernelDescs) {
     const size_t N = K.Params.size();
+    PresumedLoc PLoc = S.Context.getSourceManager().getPresumedLoc(
+        S.Context.getSourceManager()
+            .getExpansionRange(K.KernelLocation)
+            .getEnd());
     if (K.IsUnnamedKernel) {
       O << "template <> struct KernelInfoData<";
       OutputStableNameInChars(O, K.StableName);
@@ -4707,6 +4711,44 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
     O << "  __SYCL_DLL_LOCAL\n";
     O << "  static constexpr bool isESIMD() { return " << K.IsESIMDKernel
       << "; }\n";
+    O << "  __SYCL_DLL_LOCAL\n";
+    O << "  static constexpr const char* getFileName() {\n";
+    O << "#ifndef NDEBUG\n";
+    O << "    return \""
+      << std::string(PLoc.getFilename())
+             .substr(std::string(PLoc.getFilename()).find_last_of("/\\") + 1);
+    O << "\";\n";
+    O << "#else\n";
+    O << "    return \"\";\n";
+    O << "#endif\n";
+    O << "  }\n";
+    O << "  __SYCL_DLL_LOCAL\n";
+    O << "  static constexpr const char* getFunctionName() {\n";
+    O << "#ifndef NDEBUG\n";
+    O << "    return \"";
+    SYCLKernelNameTypePrinter Printer(O, Policy);
+    Printer.Visit(K.NameType);
+    O << "\";\n";
+    O << "#else\n";
+    O << "    return \"\";\n";
+    O << "#endif\n";
+    O << "  }\n";
+    O << "  __SYCL_DLL_LOCAL\n";
+    O << "  static constexpr unsigned getLineNumber() {\n";
+    O << "#ifndef NDEBUG\n";
+    O << "    return " << PLoc.getLine() << ";\n";
+    O << "#else\n";
+    O << "    return 0;\n";
+    O << "#endif\n";
+    O << "  }\n";
+    O << "  __SYCL_DLL_LOCAL\n";
+    O << "  static constexpr unsigned getColumnNumber() {\n";
+    O << "#ifndef NDEBUG\n";
+    O << "    return " << PLoc.getColumn() << ";\n";
+    O << "#else\n";
+    O << "    return 0;\n";
+    O << "#endif\n";
+    O << "  }\n";
     O << "};\n";
     CurStart += N;
   }
diff --git a/clang/test/CodeGenSYCL/code_location.cpp b/clang/test/CodeGenSYCL/code_location.cpp
new file mode 100644
index 0000000000000..9cb9bcda33ccc
--- /dev/null
+++ b/clang/test/CodeGenSYCL/code_location.cpp
@@ -0,0 +1,308 @@
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem -sycl-std=2020 -fsycl-int-header=%t.h %s -o %t.out
+// RUN: FileCheck -input-file=%t.h %s
+// RUN: %clang_cc1 -fsycl-is-host -sycl-std=2020 %s | FileCheck -input-file=%t.h %s
+
+#include "Inputs/sycl.hpp"
+
+// Check that meaningful information is returned when NDEBUG is not defined
+// and empty strings and 0s are emitted when it is.
+int test1() {
+  cl::sycl::queue q;
+  q.submit([&](cl::sycl::handler &h) { h.single_task([] {}); });
+  q.submit([&](cl::sycl::handler &h) { h.single_task<class KernelName>([]() {}); });
+  return 0;
+}
+// CHECK: template <> struct KernelInfoData<'_', 'Z', 'T', 'S', 'Z', 'Z', '5', 't', 'e', 's', 't', '1', 'v', 'E', 'N', 'K', 'U', 'l', 'R', 'N', '2', 'c', 'l', '4', 's', 'y', 'c', 'l', '7', 'h', 'a', 'n', 'd', 'l', 'e', 'r', 'E', 'E', '_', 'c', 'l', 'E', 'S', '2', '_', 'E', 'U', 'l', 'v', 'E', '_'> {
+// CHECK:   static constexpr const char* getFileName() {
+// CHECK: #ifndef NDEBUG
+// CHECK:     return "code_location.cpp";
+// CHECK: #else
+// CHECK:     return "";
+// CHECK: #endif
+// CHECK:   }
+// CHECK:   static constexpr const char* getFunctionName() {
+// CHECK: #ifndef NDEBUG
+// CHECK:     return "";
+// CHECK: #else
+// CHECK:     return "";
+// CHECK: #endif
+// CHECK:   }
+// CHECK:   static constexpr unsigned getLineNumber() {
+// CHECK: #ifndef NDEBUG
+// CHECK:     return 11;
+// CHECK: #else
+// CHECK:     return 0;
+// CHECK: #endif
+// CHECK:   }
+// CHECK:   static constexpr unsigned getColumnNumber() {
+// CHECK: #ifndef NDEBUG
+// CHECK:     return 54;
+// CHECK: #else
+// CHECK:     return 0;
+// CHECK: #endif
+// CHECK:   }
+// CHECK: };
+
+// CHECK: template <> struct KernelInfo<KernelName> {
+// CHECK:   static constexpr const char* getFileName() {
+// CHECK: #ifndef NDEBUG
+// CHECK:     return "code_location.cpp";
+// CHECK: #else
+// CHECK:     return "";
+// CHECK: #endif
+// CHECK:   }
+// CHECK:   __SYCL_DLL_LOCAL
+// CHECK:   static constexpr const char* getFunctionName() {
+// CHECK: #ifndef NDEBUG
+// CHECK:     return "KernelName";
+// CHECK: #else
+// CHECK:     return "";
+// CHECK: #endif
+// CHECK:   }
+// CHECK:   __SYCL_DLL_LOCAL
+// CHECK:   static constexpr unsigned getLineNumber() {
+// CHECK: #ifndef NDEBUG
+// CHECK:     return 12;
+// CHECK: #else
+// CHECK:     return 0;
+// CHECK: #endif
+// CHECK:   }
+// CHECK:   __SYCL_DLL_LOCAL
+// CHECK:   static constexpr unsigned getColumnNumber() {
+// CHECK: #ifndef NDEBUG
+// CHECK:     return 72;
+// CHECK: #else
+// CHECK:     return 0;
+// CHECK: #endif
+// CHECK:   }
+// CHECK: };
+
+// Check that the right name and location is returned when
+// lambda and kernel name are defined on different lines
+class KernelName2;
+int test2() {
+  cl::sycl::queue q;
+  q.submit([&](cl::sycl::handler &h) { h.single_task<KernelName2>(
+                                           [] { int i = 2; }); });
+  return 0;
+}
+// CHECK: template <> struct KernelInfo<::KernelName2> {
+// CHECK:   static constexpr const char* getFileName() {
+// CHECK: #ifndef NDEBUG
+// CHECK:     return "code_location.cpp";
+// CHECK: #else
+// CHECK:     return "";
+// CHECK: #endif
+// CHECK:   }
+// CHECK:   static constexpr const char* getFunctionName() {
+// CHECK: #ifndef NDEBUG
+// CHECK:     return "::KernelName2";
+// CHECK: #else
+// CHECK:     return "";
+// CHECK: #endif
+// CHECK:   }
+// CHECK:   static constexpr unsigned getLineNumber() {
+// CHECK: #ifndef NDEBUG
+// CHECK:     return 86;
+// CHECK: #else
+// CHECK:     return 0;
+// CHECK: #endif
+// CHECK:   }
+// CHECK:   static constexpr unsigned getColumnNumber() {
+// CHECK: #ifndef NDEBUG
+// CHECK:     return 44;
+// CHECK: #else
+// CHECK:     return 0;
+// CHECK: #endif
+// CHECK:   }
+// CHECK: };
+
+// Check that fully qualified name is returned
+template <typename T> class KernelName3;
+int test3() {
+  cl::sycl::queue q;
+  q.submit([&](cl::sycl::handler &h) { h.single_task<KernelName3<KernelName2>>(
+                                           [] { int i = 3; }); });
+  return 0;
+}
+// CHECK: template <> struct KernelInfo<::KernelName3<::KernelName2>> {
+// CHECK:   static constexpr const char* getFileName() {
+// CHECK: #ifndef NDEBUG
+// CHECK:     return "code_location.cpp";
+// CHECK: #else
+// CHECK:     return "";
+// CHECK: #endif
+// CHECK:   }
+// CHECK:   static constexpr const char* getFunctionName() {
+// CHECK: #ifndef NDEBUG
+// CHECK:     return "::KernelName3<::KernelName2>";
+// CHECK: #else
+// CHECK:     return "";
+// CHECK: #endif
+// CHECK:   }
+// CHECK:   static constexpr unsigned getLineNumber() {
+// CHECK: #ifndef NDEBUG
+// CHECK:     return 125;
+// CHECK: #else
+// CHECK:     return 0;
+// CHECK: #endif
+// CHECK:   }
+// CHECK:   static constexpr unsigned getColumnNumber() {
+// CHECK: #ifndef NDEBUG
+// CHECK:     return 44;
+// CHECK: #else
+// CHECK:     return 0;
+// CHECK: #endif
+// CHECK:   }
+// CHECK: };
+
+// Check that the location information returned is that of l4
+auto l4 = []() { return 4; };
+int test4() {
+  cl::sycl::queue q;
+  q.submit([=](cl::sycl::handler &h) { h.single_task<class KernelName4>(l4); });
+  return 0;
+}
+// CHECK: template <> struct KernelInfo<KernelName4> {
+// CHECK:   static constexpr const char* getFileName() {
+// CHECK: #ifndef NDEBUG
+// CHECK:     return "code_location.cpp";
+// CHECK: #else
+// CHECK:     return "";
+// CHECK: #endif
+// CHECK:   }
+// CHECK:   static constexpr const char* getFunctionName() {
+// CHECK: #ifndef NDEBUG
+// CHECK:     return "KernelName4";
+// CHECK: #else
+// CHECK:     return "";
+// CHECK: #endif
+// CHECK:   }
+// CHECK:   static constexpr unsigned getLineNumber() {
+// CHECK: #ifndef NDEBUG
+// CHECK:     return 160;
+// CHECK: #else
+// CHECK:     return 0;
+// CHECK: #endif
+// CHECK:   }
+// CHECK:   static constexpr unsigned getColumnNumber() {
+// CHECK: #ifndef NDEBUG
+// CHECK:     return 11;
+// CHECK: #else
+// CHECK:     return 0;
+// CHECK: #endif
+// CHECK:   }
+// CHECK: };
+
+// Check that fully qualified name is returned when unnamed lambda
+// kernel is enclosed in a namespace
+namespace NS {
+int test5() {
+  cl::sycl::queue q;
+  q.submit([=](cl::sycl::handler &h) { h.single_task([] {}); });
+  q.submit([=](cl::sycl::handler &h) { h.single_task<class KernelName5>([] {}); });
+  return 0;
+}
+} // namespace NS
+// CHECK: template <> struct KernelInfoData<'_', 'Z', 'T', 'S', 'Z', 'Z', 'N', '2', 'N', 'S', '5', 't', 'e', 's', 't', '5', 'E', 'v', 'E', 'N', 'K', 'U', 'l', 'R', 'N', '2', 'c', 'l', '4', 's', 'y', 'c', 'l', '7', 'h', 'a', 'n', 'd', 'l', 'e', 'r', 'E', 'E', '_', 'c', 'l', 'E', 'S', '3', '_', 'E', 'U', 'l', 'v', 'E', '_'> {
+// CHECK:   static constexpr const char* getFileName() {
+// CHECK: #ifndef NDEBUG
+// CHECK:     return "code_location.cpp";
+// CHECK: #else
+// CHECK:     return "";
+// CHECK: #endif
+// CHECK:   }
+// CHECK:   static constexpr const char* getFunctionName() {
+// CHECK: #ifndef NDEBUG
+// CHECK:     return "NS::";
+// CHECK: #else
+// CHECK:     return "";
+// CHECK: #endif
+// CHECK:   }
+// CHECK:   static constexpr unsigned getLineNumber() {
+// CHECK: #ifndef NDEBUG
+// CHECK:     return 202;
+// CHECK: #else
+// CHECK:     return 0;
+// CHECK: #endif
+// CHECK:   }
+// CHECK:   static constexpr unsigned getColumnNumber() {
+// CHECK: #ifndef NDEBUG
+// CHECK:     return 54;
+// CHECK: #else
+// CHECK:     return 0;
+// CHECK: #endif
+// CHECK:   }
+// CHECK: };
+// CHECK: template <> struct KernelInfo<NS::KernelName5> {
+// CHECK:   static constexpr const char* getFileName() {
+// CHECK: #ifndef NDEBUG
+// CHECK:     return "code_location.cpp";
+// CHECK: #else
+// CHECK:     return "";
+// CHECK: #endif
+// CHECK:   }
+// CHECK:   static constexpr const char* getFunctionName() {
+// CHECK: #ifndef NDEBUG
+// CHECK:     return "NS::KernelName5";
+// CHECK: #else
+// CHECK:     return "";
+// CHECK: #endif
+// CHECK:   }
+// CHECK:   static constexpr unsigned getLineNumber() {
+// CHECK: #ifndef NDEBUG
+// CHECK:     return 203;
+// CHECK: #else
+// CHECK:     return 0;
+// CHECK: #endif
+// CHECK:   }
+// CHECK:   static constexpr unsigned getColumnNumber() {
+// CHECK: #ifndef NDEBUG
+// CHECK:     return 73;
+// CHECK: #else
+// CHECK:     return 0;
+// CHECK: #endif
+// CHECK:   }
+// CHECK: };
+
+// Check that the location information returned is that of the Functor
+struct Functor {
+  void operator()() const {
+  }
+};
+int test6() {
+  Functor F;
+  cl::sycl::queue q;
+  q.submit([=](cl::sycl::handler &h) { h.single_task<class KernelName6>(F); });
+  return 0;
+}
+// CHECK: template <> struct KernelInfo<KernelName6> {
+// CHECK:   static constexpr const char* getFileName() {
+// CHECK: #ifndef NDEBUG
+// CHECK:     return "code_location.cpp";
+// CHECK: #else
+// CHECK:     return "";
+// CHECK: #endif
+// CHECK:   }
+// CHECK:   static constexpr const char* getFunctionName() {
+// CHECK: #ifndef NDEBUG
+// CHECK:     return "KernelName6";
+// CHECK: #else
+// CHECK:     return "";
+// CHECK: #endif
+// CHECK:   }
+// CHECK:   static constexpr unsigned getLineNumber() {
+// CHECK: #ifndef NDEBUG
+// CHECK:     return 269;
+// CHECK: #else
+// CHECK:     return 0;
+// CHECK: #endif
+// CHECK:   }
+// CHECK:   static constexpr unsigned getColumnNumber() {
+// CHECK: #ifndef NDEBUG
+// CHECK:     return 8;
+// CHECK: #else
+// CHECK:     return 0;
+// CHECK: #endif
+// CHECK:   }
+// CHECK: };
diff --git a/sycl/include/CL/sycl/detail/kernel_desc.hpp b/sycl/include/CL/sycl/detail/kernel_desc.hpp
index e3daedc6ad701..3545631083a70 100644
--- a/sycl/include/CL/sycl/detail/kernel_desc.hpp
+++ b/sycl/include/CL/sycl/detail/kernel_desc.hpp
@@ -75,6 +75,10 @@ template <class KernelNameType> struct KernelInfo {
   }
   static constexpr const char *getName() { return ""; }
   static constexpr bool isESIMD() { return 0; }
+  static constexpr const char *getFileName() { return ""; }
+  static constexpr const char *getFunctionName() { return ""; }
+  static constexpr unsigned getLineNumber() { return 0; }
+  static constexpr unsigned getColumnNumber() { return 0; }
 };
 #else
 template <char...> struct KernelInfoData {
@@ -85,6 +89,10 @@ template <char...> struct KernelInfoData {
   }
   static constexpr const char *getName() { return ""; }
   static constexpr bool isESIMD() { return 0; }
+  static constexpr const char *getFileName() { return ""; }
+  static constexpr const char *getFunctionName() { return ""; }
+  static constexpr unsigned getLineNumber() { return 0; }
+  static constexpr unsigned getColumnNumber() { return 0; }
 };
 
 // C++14 like index_sequence and make_index_sequence
@@ -123,6 +131,10 @@ template <class KernelNameType> struct KernelInfo {
   }
   static constexpr const char *getName() { return SubKernelInfo::getName(); }
   static constexpr bool isESIMD() { return SubKernelInfo::isESIMD(); }
+  static constexpr const char *getFileName() { return ""; }
+  static constexpr const char *getFunctionName() { return ""; }
+  static constexpr unsigned getLineNumber() { return 0; }
+  static constexpr unsigned getColumnNumber() { return 0; }
 };
 #endif //__SYCL_UNNAMED_LAMBDA__