diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 07a6c25dced91..9c47fc88e2c06 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/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() diff --git a/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp b/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp index 29b697691f445..2ef457bfa4f9b 100644 --- a/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp +++ b/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 diff --git a/sycl/test/extensions/free_function_kernels/free_function_host_compiler.cpp b/sycl/test/extensions/free_function_kernels/free_function_host_compiler.cpp new file mode 100644 index 0000000000000..296ee82184ad5 --- /dev/null +++ b/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 +#include +#include + +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(q.get_context()); + sycl::kernel_id kID = + sycl::ext::oneapi::experimental::get_kernel_id(); + sycl::kernel krn = bundle.get_kernel(kID); + + q.submit([&](sycl::handler &cgh) { + sycl::nd_range<1> ndr; + cgh.parallel_for(ndr, krn); + }); + return 0; +}