[SYCL] Free function host compilation bugfix

clang/lib/Sema/SemaSYCL.cpp

@@ -6947,6 +6947,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;
+}