From d0622abdcb2d863fa2ecc806248a2b1bf4ad9fa3 Mon Sep 17 00:00:00 2001 From: "Klochkov, Denis" Date: Thu, 6 Aug 2026 15:50:20 +0200 Subject: [PATCH 1/5] [SYCL] do no emit free function body into host object file --- clang/lib/Sema/SemaSYCL.cpp | 10 ++- .../host_object_no_bloat.cpp | 75 +++++++++++++++++++ 2 files changed, 83 insertions(+), 2 deletions(-) create mode 100644 sycl/test/extensions/free_function_kernels/host_object_no_bloat.cpp diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 62a610342fec0..128eeb00fb25d 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -4494,8 +4494,14 @@ class FreeFunctionKernelBodyCreator : public SyclKernelFieldHandler { FreeFunctionSrcLoc(FF->getLocation()) {} ~FreeFunctionKernelBodyCreator() { - CompoundStmt *KernelBody = createFreeFunctionKernelBody(); - DeclCreator.setBody(KernelBody); + // For host compilation, skip creating the wrapper body to reduce object + // size. The wrapper (__sycl_kernel_*) won't be called on host, so leaving + // it as a declaration-only significantly reduces host object bloat (~3x + // size issue). + if (SemaSYCLRef.SemaRef.getLangOpts().SYCLIsDevice) { + CompoundStmt *KernelBody = createFreeFunctionKernelBody(); + DeclCreator.setBody(KernelBody); + } } bool handleSyclSpecialType(FieldDecl *FD, QualType FieldTy) final { diff --git a/sycl/test/extensions/free_function_kernels/host_object_no_bloat.cpp b/sycl/test/extensions/free_function_kernels/host_object_no_bloat.cpp new file mode 100644 index 0000000000000..bb79418da7f12 --- /dev/null +++ b/sycl/test/extensions/free_function_kernels/host_object_no_bloat.cpp @@ -0,0 +1,75 @@ +// This test verifies that free function kernels don't bloat host object files. +// The __sycl_kernel_* wrapper should not have a body during host compilation. +// This addresses CMPLRLLVM-77222 where free function kernels caused ~3x bloat. +// +// Compile to object file (bundled) +// RUN: %clangxx -fsycl -fsycl-targets=spir64 -c %s -o %t.o +// +// Extract host object +// RUN: clang-offload-bundler -type=o -targets=host-x86_64-unknown-linux-gnu -input=%t.o -output=%t_host.o -unbundle +// +// Dump symbols to a file and verify wrapper is NOT defined +// RUN: llvm-nm %t_host.o > %t_host_syms.txt +// +// The wrapper function __sycl_kernel_* should not have type 'T' (defined text). +// If it appears with 'T', the fix has regressed. +// RUN: FileCheck %s --input-file %t_host_syms.txt --check-prefix=CHECK-HOST +// +// CHECK-HOST-NOT: {{[0-9a-f]+ T .*__sycl_kernel.*test_kernel}} + +#include +#include +#include +#include +#include + +namespace syclexp = sycl::ext::oneapi::experimental; + +// Free function kernel with nd_range_kernel property +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) +void test_kernel(sycl::nd_item<1> item, int *data, int n) { + auto id = item.get_global_id(0); + if (id < n) { + data[id] = data[id] * 2 + 1; + } +} + +// Single task free function kernel +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::single_task_kernel)) +void init_kernel(int *ptr, int value) { + *ptr = value; +} + +int main() { + sycl::queue q; + const int N = 64; + + // Allocate USM memory + int *data = sycl::malloc_shared(N, q); + + // Initialize data using single_task free function + syclexp::single_task(q, syclexp::kernel_function_s{}, data, 42); + + // Get kernel bundle for free function kernels + auto kb = sycl::get_kernel_bundle( + q.get_context(), {q.get_device()}, + {sycl::get_kernel_id>()}); + + // Launch free function kernel using nd_launch with kernel bundle + q.submit([&](sycl::handler &h) { + h.use_kernel_bundle(kb); + sycl::nd_range<1> range{N, 64}; + syclexp::nd_launch(h, range, syclexp::kernel_function_s{}, + data, N); + }); + + // Alternative: direct queue submission + sycl::nd_range<1> range{N, 64}; + syclexp::nd_launch(q, range, syclexp::kernel_function_s{}, + data, N); + + q.wait(); + + sycl::free(data, q); + return 0; +} From e1d11367cee652fdbc60508d2da3a6800cafe55b Mon Sep 17 00:00:00 2001 From: "Klochkov, Denis" Date: Thu, 6 Aug 2026 17:15:49 +0200 Subject: [PATCH 2/5] [SYCL][E2E] limit test to run on windows --- .../extensions/free_function_kernels/host_object_no_bloat.cpp | 1 + 1 file changed, 1 insertion(+) diff --git a/sycl/test/extensions/free_function_kernels/host_object_no_bloat.cpp b/sycl/test/extensions/free_function_kernels/host_object_no_bloat.cpp index bb79418da7f12..5a965d4570c7f 100644 --- a/sycl/test/extensions/free_function_kernels/host_object_no_bloat.cpp +++ b/sycl/test/extensions/free_function_kernels/host_object_no_bloat.cpp @@ -3,6 +3,7 @@ // This addresses CMPLRLLVM-77222 where free function kernels caused ~3x bloat. // // Compile to object file (bundled) +// REQUIRES: linux // RUN: %clangxx -fsycl -fsycl-targets=spir64 -c %s -o %t.o // // Extract host object From ddd9282961168d3bb7efc13579ee425947544df4 Mon Sep 17 00:00:00 2001 From: "Klochkov, Denis" Date: Thu, 6 Aug 2026 17:50:32 +0200 Subject: [PATCH 3/5] [SYCL][E2E] add sanity check --- .../free_function_kernels/host_object_no_bloat.cpp | 9 +++++++++ 1 file changed, 9 insertions(+) diff --git a/sycl/test/extensions/free_function_kernels/host_object_no_bloat.cpp b/sycl/test/extensions/free_function_kernels/host_object_no_bloat.cpp index 5a965d4570c7f..5d9ec699870a4 100644 --- a/sycl/test/extensions/free_function_kernels/host_object_no_bloat.cpp +++ b/sycl/test/extensions/free_function_kernels/host_object_no_bloat.cpp @@ -12,6 +12,15 @@ // Dump symbols to a file and verify wrapper is NOT defined // RUN: llvm-nm %t_host.o > %t_host_syms.txt // +// Sanity check that the dump has real content, so the CHECK-NOT below can't +// pass vacuously (e.g. because of an empty/broken symbol dump). The original +// kernel functions keep their host definitions; only the __sycl_kernel_* +// wrapper's body is elided. +// RUN: FileCheck %s --input-file %t_host_syms.txt --check-prefix=CHECK-HOST-SANITY +// +// CHECK-HOST-SANITY-DAG: {{[0-9a-f]+ T _Z11test_kernel}} +// CHECK-HOST-SANITY-DAG: {{[0-9a-f]+ T _Z11init_kernel}} +// // The wrapper function __sycl_kernel_* should not have type 'T' (defined text). // If it appears with 'T', the fix has regressed. // RUN: FileCheck %s --input-file %t_host_syms.txt --check-prefix=CHECK-HOST From 7132baa3373f1728283464cc4a74dfcecc90f80d Mon Sep 17 00:00:00 2001 From: "Klochkov, Denis" Date: Thu, 6 Aug 2026 17:51:08 +0200 Subject: [PATCH 4/5] [SYCL][E2E] fix formatting --- .../free_function_kernels/host_object_no_bloat.cpp | 8 +++----- 1 file changed, 3 insertions(+), 5 deletions(-) diff --git a/sycl/test/extensions/free_function_kernels/host_object_no_bloat.cpp b/sycl/test/extensions/free_function_kernels/host_object_no_bloat.cpp index 5d9ec699870a4..99fa44f063385 100644 --- a/sycl/test/extensions/free_function_kernels/host_object_no_bloat.cpp +++ b/sycl/test/extensions/free_function_kernels/host_object_no_bloat.cpp @@ -46,9 +46,7 @@ void test_kernel(sycl::nd_item<1> item, int *data, int n) { // Single task free function kernel SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::single_task_kernel)) -void init_kernel(int *ptr, int value) { - *ptr = value; -} +void init_kernel(int *ptr, int value) { *ptr = value; } int main() { sycl::queue q; @@ -75,8 +73,8 @@ int main() { // Alternative: direct queue submission sycl::nd_range<1> range{N, 64}; - syclexp::nd_launch(q, range, syclexp::kernel_function_s{}, - data, N); + syclexp::nd_launch(q, range, syclexp::kernel_function_s{}, data, + N); q.wait(); From d6beec7feaef45f456149850a21537415c007496 Mon Sep 17 00:00:00 2001 From: "Klochkov, Denis" Date: Thu, 6 Aug 2026 19:38:30 +0200 Subject: [PATCH 5/5] [SYCL][E2E] call free function from host directly --- .../host_object_no_bloat.cpp | 25 +++++++++++++++++++ 1 file changed, 25 insertions(+) diff --git a/sycl/test/extensions/free_function_kernels/host_object_no_bloat.cpp b/sycl/test/extensions/free_function_kernels/host_object_no_bloat.cpp index 99fa44f063385..ddba98b94c35c 100644 --- a/sycl/test/extensions/free_function_kernels/host_object_no_bloat.cpp +++ b/sycl/test/extensions/free_function_kernels/host_object_no_bloat.cpp @@ -26,6 +26,13 @@ // RUN: FileCheck %s --input-file %t_host_syms.txt --check-prefix=CHECK-HOST // // CHECK-HOST-NOT: {{[0-9a-f]+ T .*__sycl_kernel.*test_kernel}} +// +// Build to verify the free function kernels still compile and link +// correctly. Actual device execution is covered by the e2e test suite; +// sycl/test tests are compile-only and must not depend on a real device. +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t2.out + +#include #include #include @@ -49,6 +56,24 @@ SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::single_task_kernel)) void init_kernel(int *ptr, int value) { *ptr = value; } int main() { + // Call test_kernel directly as a plain host function to confirm it + // keeps a real, callable definition on the host (not just the device-only + // __sycl_kernel_* wrapper). + { + sycl::id<1> zero{0}; + sycl::range<1> extent{1}; + auto group = sycl::detail::Builder::createGroup<1>(extent, extent, zero); + auto globalItem = + sycl::detail::Builder::createItem<1, true>(extent, zero, zero); + auto localItem = sycl::detail::Builder::createItem<1, false>(extent, zero); + auto item = + sycl::detail::Builder::createNDItem<1>(globalItem, localItem, group); + + int hostData = 20; + test_kernel(item, &hostData, 1); + assert(hostData == 41); + } + sycl::queue q; const int N = 64;