From 2093f82e6b54c86f1a5a577aa456928b1e864472 Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Fri, 1 Aug 2025 16:50:26 -0400 Subject: [PATCH 1/2] [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 --- clang/lib/Sema/SemaSYCL.cpp | 4 +++ .../free_function_int_header_rtc_mode.cpp | 7 ++-- .../free_function_host_compiler.cpp | 32 +++++++++++++++++++ 3 files changed, 39 insertions(+), 4 deletions(-) create mode 100644 sycl/test/extensions/free_function_kernels/free_function_host_compiler.cpp diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index f983582ed2085..4a5f2aace1429 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/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() 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 48d5662e709aa..fae7be0120ba9 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; +} From 68d2d23163a4f9c3903f4c1b0256db415285cca5 Mon Sep 17 00:00:00 2001 From: Andrei Elovikov Date: Tue, 5 Aug 2025 07:32:52 -0700 Subject: [PATCH 2/2] [SYCL][TEST] Skip `free_function_host_compiler.cpp` when using `libcxx` (#19696) GNU `g++` host compiler doesn't work with libcxx without some voodoo magic. --- .../free_function_kernels/free_function_host_compiler.cpp | 1 + 1 file changed, 1 insertion(+) 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 index 296ee82184ad5..f468d3e2ffaa0 100644 --- a/sycl/test/extensions/free_function_kernels/free_function_host_compiler.cpp +++ b/sycl/test/extensions/free_function_kernels/free_function_host_compiler.cpp @@ -2,6 +2,7 @@ // RUN: %clangxx -fsycl -fno-sycl-unnamed-lambda %s // RUN: %clangxx -fsycl -fsycl-host-compiler=g++ %s // REQUIRES: linux +// UNSUPPORTED: libcxx // This test serves as a sanity check that compilation succeeds // for a simple free function kernel when using the host compiler