[SYCL] Free function host compilation bugfix (#19541)

The `KernelInfoData` struct is specialized in the integration header for
free function kernels and results in compilation errors when using other
host compilers to compile code for example while using
`-fsycl-host-compiler=g++` . This PR removes its generation for free
function kernels in the integration header altogether since it is not
used anywhere.

For more info about the bug tackled by this PR: 

`KernelInfoData` struct is specialized in the integration header for
unnamed kernels such as free function kernels with the primary
definition located in `kernel_desc.hpp`. The primary definition however
seems to be conditional upon the fact that the unnamed lambda extension
must be enabled which is not the case with other host compilers and we
end up with an error of specialization without a primary template.
Furthermore setting `-fsycl-unnamed-lambda` doesn't do the trick as
apparently it is not allowed to be used together with
`-fsycl-host-compiler`. Since it seems unreasonable to disallow using
other host compilers with free function kernels, this seemed one of the
easier solutions at the moment.

Another approach would be to potentially remove the ifdefs in
kernel_desc.hpp that conditionally include/exclude `KernelInfo `and
`KernelInfoData` and have them both available unconditionally. Tagging
@bader for some input about this alternative approach.

---------

Co-authored-by: Tom Honermann <[email protected]>

Author: lbushi25

clang/lib/Sema/SemaSYCL.cpp

@@ -6980,6 +6980,10 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
 
   for (const KernelDesc &K : KernelDescs) {
     const size_t N = K.Params.size();
+    if (S.isFreeFunction(K.SyclKernel)) {
+      CurStart += N;
+      continue;
+    }
     PresumedLoc PLoc = S.getASTContext().getSourceManager().getPresumedLoc(
         S.getASTContext()
             .getSourceManager()

clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp

@@ -44,10 +44,9 @@ int main(){
 // CHECK-NEXT:    "{{.*}}__sycl_kernel_free_function_nd_rangePiii",
 // CHECK-NEXT:    "{{.*}}Kernel_Function",
 
-
-// CHECK: static constexpr const char* getName() { return "{{.*}}__sycl_kernel_free_function_singlePiii"; }
-// CHECK: static constexpr const char* getName() { return "{{.*}}__sycl_kernel_free_function_nd_rangePiii"; }
-// CHECK: static constexpr const char* getName() { return "{{.*}}Kernel_Function"; }
+// CHECK: _Z34__sycl_kernel_free_function_singlePiii
+// CHECK: _Z36__sycl_kernel_free_function_nd_rangePiii
+// CHECK: _ZTSZ4mainE15Kernel_Function
 
 // CHECK-RTC-NOT: free_function_single_kernel
 // CHECK-RTC-NOT: free_function_nd_range

sycl/test/extensions/free_function_kernels/free_function_host_compiler.cpp

@@ -0,0 +1,32 @@
+// RUN: %clangxx -fsycl %s
+// RUN: %clangxx -fsycl -fno-sycl-unnamed-lambda %s
+// RUN: %clangxx -fsycl -fsycl-host-compiler=g++ %s
+// REQUIRES: linux
+
+// This test serves as a sanity check that compilation succeeds
+// for a simple free function kernel when using the host compiler
+// flag with unnamed lambda extension enabled(default) and disabled.
+
+#include <sycl/detail/core.hpp>
+#include <sycl/ext/oneapi/experimental/free_function_traits.hpp>
+#include <sycl/kernel_bundle.hpp>
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(
+    (sycl::ext::oneapi::experimental::nd_range_kernel<1>))
+void kernel() {}
+
+int main() {
+  sycl::queue q;
+
+  sycl::kernel_bundle bundle =
+      sycl::get_kernel_bundle<sycl::bundle_state::executable>(q.get_context());
+  sycl::kernel_id kID =
+      sycl::ext::oneapi::experimental::get_kernel_id<kernel>();
+  sycl::kernel krn = bundle.get_kernel(kID);
+
+  q.submit([&](sycl::handler &cgh) {
+    sycl::nd_range<1> ndr;
+    cgh.parallel_for(ndr, krn);
+  });
+  return 0;
+}