[SYCL] Fix compile error for handler::copy with -fsycl-unnamed-lambda (#903)

Some `copy` methods in handler are implemented via `parallel_for` for
host device and hence, we need to have `KernelInfo` struct (provided by
integration header) to compile them.
This effectively creates cyclic dependency and we already have some kind
of a workaround for it for regular mode (without -fsycl-unnamed-lambda).

This patch applies the same solution for unnamed lambda mechanism.

Signed-off-by: Alexey Sachkov <[email protected]>

Author: AlexeySachkov

sycl/include/CL/sycl/detail/kernel_desc.hpp

@@ -63,7 +63,14 @@ template <class KernelNameType> struct KernelInfo {
   static constexpr const char *getName() { return ""; }
 };
 #else
-template <char...> struct KernelInfoData; // Should this have dummy impl?
+template <char...> struct KernelInfoData {
+  static constexpr unsigned getNumParams() { return 0; }
+  static const kernel_param_desc_t &getParamDesc(int Idx) {
+    static kernel_param_desc_t Dummy;
+    return Dummy;
+  }
+  static constexpr const char *getName() { return ""; }
+};
 
 // C++14 like index_sequence and make_index_sequence
 // not needed C++14 members (value_type, size) not implemented

sycl/test/regression/copy-with-unnamed-lambda.cpp

@@ -0,0 +1,59 @@
+// RUN: %clangxx -fsycl -fsycl-unnamed-lambda %s -o %t.out
+// The purpose of this test is to check that the following code can be
+// successfully compiled
+#include <CL/sycl.hpp>
+
+#include <iostream>
+
+int main() {
+  auto AsyncHandler = [](cl::sycl::exception_list EL) {
+    for (std::exception_ptr const &P : EL) {
+      try {
+        std::rethrow_exception(P);
+      } catch (std::exception const &E) {
+        std::cerr << "Caught async SYCL exception: " << E.what() << std::endl;
+      }
+    }
+  };
+
+  cl::sycl::queue Q(AsyncHandler);
+
+  constexpr size_t Size = 10;
+  const int ReferenceData[Size] = { 1, 2, 3, 4, 5, 6, 7, 8, 9, 10 };
+
+  cl::sycl::buffer<int> Buf(Size);
+
+  Q.submit([&](cl::sycl::handler &CGH) {
+    auto Acc = Buf.get_access<cl::sycl::access::mode::write>(CGH);
+    CGH.copy(ReferenceData, Acc);
+  });
+
+  Q.wait_and_throw();
+
+  auto Acc = Buf.get_access<cl::sycl::access::mode::read>();
+  for (size_t I = 0; I < Size; ++I) {
+    if (ReferenceData[I] != Acc[I]) {
+      std::cerr << "Incorrect result, got: " << Acc[I]
+                << ", expected: " << ReferenceData[I] << std::endl;
+      return 1;
+    }
+  }
+
+  int CopybackData[Size] = { 0 };
+  Q.submit([&](cl::sycl::handler &CGH) {
+    auto Acc = Buf.get_access<cl::sycl::access::mode::read>(CGH);
+    CGH.copy(Acc, CopybackData);
+  });
+
+  Q.wait_and_throw();
+
+  for (size_t I = 0; I < Size; ++I) {
+    if (ReferenceData[I] != CopybackData[I]) {
+      std::cerr << "Incorrect result, got: " << Acc[I]
+                << ", expected: " << ReferenceData[I] << std::endl;
+      return 1;
+    }
+  }
+
+  return 0;
+}