[SYCL] Fix kernel::get_native (#3655)

This patch fixes compilation errors with kernel::get_native usage and improves test coverage for this method

Author: Alexander Batashev

sycl/include/CL/sycl/backend/level_zero.hpp

@@ -57,7 +57,7 @@ struct interop<backend::level_zero,
 };
 
 namespace detail {
-template <> class BackendReturn<backend::level_zero, kernel> {
+template <> struct BackendReturn<backend::level_zero, kernel> {
   using type = ze_kernel_handle_t;
 };
 

sycl/include/CL/sycl/kernel.hpp

@@ -177,7 +177,8 @@ class __SYCL_EXPORT kernel {
   typename backend_traits<Backend>::template return_type<kernel>
   get_native() const {
     return detail::pi::cast<
-        backend_traits<Backend>::template return_type<kernel>>(getNativeImpl());
+        typename backend_traits<Backend>::template return_type<kernel>>(
+        getNativeImpl());
   }
 
 private:

sycl/test/on-device/basic_tests/interop/get_native_ocl.cpp

@@ -0,0 +1,36 @@
+// REQUIRES: opencl, opencl_dev_kit
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %opencl_options %s -o %t.ocl.out
+// RUN: %t.ocl.out
+
+#include <CL/cl.h>
+
+#include <CL/sycl/backend/opencl.hpp>
+#include <sycl/sycl.hpp>
+
+constexpr auto BE = sycl::backend::opencl;
+
+class TestKernel;
+
+int main() {
+  sycl::queue Q;
+
+  if (0) {
+    Q.submit([](sycl::handler &CGH) { CGH.single_task<TestKernel>([] {}); });
+  }
+
+  sycl::kernel_id KernelID = sycl::get_kernel_id<TestKernel>();
+
+  sycl::kernel_bundle KernelBundle =
+      sycl::get_kernel_bundle<sycl::bundle_state::executable>(Q.get_context());
+
+  sycl::kernel Kernel = KernelBundle.get_kernel(KernelID);
+
+  cl_kernel Handle = Kernel.get_native<BE>();
+
+  size_t Size = 0;
+  cl_int Err =
+      clGetKernelInfo(Handle, CL_KERNEL_FUNCTION_NAME, 0, nullptr, &Size);
+  assert(Err == CL_SUCCESS);
+
+  return 0;
+}

sycl/test/on-device/basic_tests/interop/get_native_ze.cpp

@@ -0,0 +1,35 @@
+// REQUIRES: level_zero, level_zero_dev_kit
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %level_zero_options %s -o %t.ze.out
+// RUN: %t.ze.out
+
+#include <level_zero/ze_api.h>
+
+#include <CL/sycl/backend/level_zero.hpp>
+#include <sycl/sycl.hpp>
+
+constexpr auto BE = sycl::backend::level_zero;
+
+class TestKernel;
+
+int main() {
+  sycl::queue Q;
+
+  if (0) {
+    Q.submit([](sycl::handler &CGH) { CGH.single_task<TestKernel>([] {}); });
+  }
+
+  sycl::kernel_id KernelID = sycl::get_kernel_id<TestKernel>();
+
+  sycl::kernel_bundle KernelBundle =
+      sycl::get_kernel_bundle<sycl::bundle_state::executable>(Q.get_context());
+
+  sycl::kernel Kernel = KernelBundle.get_kernel(KernelID);
+
+  ze_kernel_handle_t Handle = Kernel.get_native<BE>();
+
+  ze_kernel_properties_t KernelProperties;
+  ze_result_t Err = zeKernelGetProperties(Handle, &KernelProperties);
+  assert(Err == ZE_RESULT_SUCCESS);
+
+  return 0;
+}