From b45a7bb39551111cf507b480e3de34c5065895f2 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Thu, 26 Jun 2025 05:52:11 -0700 Subject: [PATCH 1/8] [SYCL] Remove linking workaround and test SYCL interlinking This commit removes a workaround for linking inline SYCL code with SYCLBIN. Additionally, it relaxes the requirement for dependency resolution when device images are looked up for anything other than executable state, avoiding errors when the symbol resolution is done manually. Signed-off-by: Larsen, Steffen --- sycl/source/detail/kernel_bundle_impl.hpp | 43 +--------- .../program_manager/program_manager.cpp | 5 +- .../SYCLBIN/Inputs/link_sycl_inline.cpp | 20 +++++ .../SYCLBIN/Inputs/link_sycl_inline.hpp | 79 +++++++++++++++++++ .../SYCLBIN/link_sycl_inline_input.cpp | 24 ++++++ .../SYCLBIN/link_sycl_inline_object.cpp | 24 ++++++ .../no_sycl_hpp_in_e2e_tests.cpp | 2 +- .../kernel-and-program/OutOfResources.cpp | 2 +- 8 files changed, 154 insertions(+), 45 deletions(-) create mode 100644 sycl/test-e2e/SYCLBIN/Inputs/link_sycl_inline.cpp create mode 100644 sycl/test-e2e/SYCLBIN/Inputs/link_sycl_inline.hpp create mode 100644 sycl/test-e2e/SYCLBIN/link_sycl_inline_input.cpp create mode 100644 sycl/test-e2e/SYCLBIN/link_sycl_inline_object.cpp diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index 67414a2ca197e..14892daf1764a 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -245,26 +245,6 @@ class kernel_bundle_impl // TODO: Unify with c'tor for sycl::compile and sycl::build by calling // sycl::join on vector of kernel_bundles - // Due to a bug in L0, specializations with conflicting IDs will overwrite - // each other when linked together, so to avoid this issue we link - // regular offline-compiled SYCL device images in separation. - // TODO: Remove when spec const overwriting issue has been fixed in L0. - std::vector OfflineDeviceImages; - std::unordered_set> - OfflineDeviceImageSet; - for (const kernel_bundle &ObjectBundle : - ObjectBundles) { - for (const DevImgPlainWithDeps &DeviceImageWithDeps : - getSyclObjImpl(ObjectBundle)->MDeviceImages) { - if (getSyclObjImpl(DeviceImageWithDeps.getMain())->getOriginMask() & - ImageOriginSYCLOffline) { - OfflineDeviceImages.push_back(&DeviceImageWithDeps); - for (const device_image_plain &DevImg : DeviceImageWithDeps) - OfflineDeviceImageSet.insert(getSyclObjImpl(DevImg)); - } - } - } - // Collect all unique images. std::vector DevImages; { @@ -273,9 +253,7 @@ class kernel_bundle_impl ObjectBundles) for (const device_image_plain &DevImg : getSyclObjImpl(ObjectBundle)->MUniqueDeviceImages) - if (OfflineDeviceImageSet.find(getSyclObjImpl(DevImg)) == - OfflineDeviceImageSet.end()) - DevImagesSet.insert(getSyclObjImpl(DevImg)); + DevImagesSet.insert(getSyclObjImpl(DevImg)); DevImages.reserve(DevImagesSet.size()); for (auto It = DevImagesSet.begin(); It != DevImagesSet.end();) DevImages.push_back(createSyclObjFromImpl( @@ -390,25 +368,6 @@ class kernel_bundle_impl // added. } - // ... And link the offline images in separation. (Workaround.) - for (const DevImgPlainWithDeps *DeviceImageWithDeps : OfflineDeviceImages) { - // Skip images which are not compatible with devices provided - if (std::none_of(MDevices.begin(), MDevices.end(), - [DeviceImageWithDeps](const device &Dev) { - return getSyclObjImpl(DeviceImageWithDeps->getMain()) - ->compatible_with_device(Dev); - })) - continue; - - std::vector LinkedResults = - detail::ProgramManager::getInstance().link( - DeviceImageWithDeps->getAll(), MDevices, PropList); - MDeviceImages.insert(MDeviceImages.end(), LinkedResults.begin(), - LinkedResults.end()); - MUniqueDeviceImages.insert(MUniqueDeviceImages.end(), - LinkedResults.begin(), LinkedResults.end()); - } - removeDuplicateImages(); for (const kernel_bundle &Bundle : ObjectBundles) { diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index 455810696260d..95d68c904a6a3 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -2599,7 +2599,10 @@ ProgramManager::getSYCLDeviceImagesWithCompatibleState( std::lock_guard KernelIDsGuard(m_KernelIDsMutex); ImgInfo.KernelIDs = m_BinImg2KernelIDs[BinImage]; } - ImgInfo.Deps = collectDeviceImageDeps(*BinImage, {DevImpl}); + ImgInfo.Deps = + collectDeviceImageDeps(*BinImage, {DevImpl}, + /*ErrorOnUnresolvableImport=*/TargetState == + bundle_state::executable); } const bundle_state ImgState = ImgInfo.State; const std::shared_ptr> &ImageKernelIDs = diff --git a/sycl/test-e2e/SYCLBIN/Inputs/link_sycl_inline.cpp b/sycl/test-e2e/SYCLBIN/Inputs/link_sycl_inline.cpp new file mode 100644 index 0000000000000..ba5989f51146a --- /dev/null +++ b/sycl/test-e2e/SYCLBIN/Inputs/link_sycl_inline.cpp @@ -0,0 +1,20 @@ +#include + +namespace syclext = sycl::ext::oneapi; +namespace syclexp = sycl::ext::oneapi::experimental; + +typedef void (*FuncPtrT)(size_t *); + +struct ArgsT { + size_t *Ptr; + FuncPtrT *FuncPtr; +}; + +SYCL_EXTERNAL size_t GetID() { + return syclext::this_work_item::get_nd_item<1>().get_global_id(); +} + +extern "C" SYCL_EXTERNAL SYCL_EXT_ONEAPI_FUNCTION_PROPERTY( + (syclexp::nd_range_kernel<1>)) void Kernel(ArgsT Args) { + (**Args.FuncPtr)(Args.Ptr); +} diff --git a/sycl/test-e2e/SYCLBIN/Inputs/link_sycl_inline.hpp b/sycl/test-e2e/SYCLBIN/Inputs/link_sycl_inline.hpp new file mode 100644 index 0000000000000..72336e20029f0 --- /dev/null +++ b/sycl/test-e2e/SYCLBIN/Inputs/link_sycl_inline.hpp @@ -0,0 +1,79 @@ +#include "common.hpp" + +#include + +namespace syclexp = sycl::ext::oneapi::experimental; + +typedef void (*FuncPtrT)(size_t *); + +struct ArgsT { + size_t *Ptr; + FuncPtrT *FuncPtr; +}; + +#ifdef __SYCL_DEVICE_ONLY__ +SYCL_EXTERNAL size_t GetID(); +#else +// Host-side code to avoid linker problems. Will never be called. +SYCL_EXTERNAL size_t GetID() { return 0; } +#endif + +SYCL_EXTERNAL +void Func(size_t *Ptr) { + size_t GlobalID = GetID(); + Ptr[GlobalID] = GlobalID; +} + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::single_task_kernel)) +void GetFuncPtr(ArgsT Args) { *Args.FuncPtr = Func; } + +constexpr size_t N = 32; + +int main(int argc, char *argv[]) { + assert(argc == 2); + + sycl::queue Q; + + int Failed = CommonLoadCheck(Q.get_context(), argv[1]); + +#if defined(SYCLBIN_INPUT_STATE) + auto SYCLBINInput = syclexp::get_kernel_bundle( + Q.get_context(), std::string{argv[1]}); + auto SYCLBINObj = sycl::compile(SYCLBINInput); +#elif defined(SYCLBIN_OBJECT_STATE) + auto SYCLBINObj = syclexp::get_kernel_bundle( + Q.get_context(), std::string{argv[1]}); +#else // defined(SYCLBIN_EXECUTABLE_STATE) +#error "Test does not work with executable state." +#endif + + auto KBObj = + syclexp::get_kernel_bundle( + Q.get_context()); + auto KBExe = sycl::link({KBObj, SYCLBINObj}); + + ArgsT Args{}; + Args.FuncPtr = sycl::malloc_shared(N, Q); + Args.Ptr = sycl::malloc_shared(N, Q); + + sycl::kernel GetFuncPtrKern = KBExe.ext_oneapi_get_kernel(); + Q.submit([&](sycl::handler &CGH) { + CGH.set_args(Args); + CGH.single_task(GetFuncPtrKern); + }).wait(); + + sycl::kernel Kern = KBExe.ext_oneapi_get_kernel("Kernel"); + Q.submit([&](sycl::handler &CGH) { + CGH.set_args(Args); + CGH.parallel_for(sycl::nd_range{{N}, {N}}, Kern); + }).wait(); + + for (size_t I = 0; I < N; ++I) { + if (Args.Ptr[I] != I) { + std::cout << Args.Ptr[I] << " != " << I << std::endl; + ++Failed; + } + } + + return Failed; +} diff --git a/sycl/test-e2e/SYCLBIN/link_sycl_inline_input.cpp b/sycl/test-e2e/SYCLBIN/link_sycl_inline_input.cpp new file mode 100644 index 0000000000000..4607c85e17493 --- /dev/null +++ b/sycl/test-e2e/SYCLBIN/link_sycl_inline_input.cpp @@ -0,0 +1,24 @@ +//==-------- link_sycl_inline_input.cpp --- SYCLBIN extension tests --------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +// REQUIRES: aspect-usm_shared_allocations + +// -- Test for linking between inline SYCL code and SYCLBIN code. + +// Due to the regression in https://github.com/intel/llvm/issues/18432 it will +// fail to build the SYCLBIN with nvptx targets. Once this is fixed, +// %{sycl_target_opts} should be added to the SYCLBIN generation run-line. +// REQUIRES: target-spir + +// RUN: %clangxx --offload-new-driver -fsyclbin=input -fsycl-allow-device-image-dependencies -Xclang -fsycl-allow-func-ptr %S/Inputs/link_sycl_inline.cpp -o %t.syclbin +// RUN: %{build} -fsycl-allow-device-image-dependencies -o %t.out +// RUN: %{l0_leak_check} %{run} %t.out %t.syclbin + +#define SYCLBIN_INPUT_STATE + +#include "Inputs/link_sycl_inline.hpp" diff --git a/sycl/test-e2e/SYCLBIN/link_sycl_inline_object.cpp b/sycl/test-e2e/SYCLBIN/link_sycl_inline_object.cpp new file mode 100644 index 0000000000000..d6ff02b6e2633 --- /dev/null +++ b/sycl/test-e2e/SYCLBIN/link_sycl_inline_object.cpp @@ -0,0 +1,24 @@ +//==-------- link_sycl_inline_object.cpp --- SYCLBIN extension tests --------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +// REQUIRES: aspect-usm_shared_allocations + +// -- Test for linking between inline SYCL code and SYCLBIN code. + +// Due to the regression in https://github.com/intel/llvm/issues/18432 it will +// fail to build the SYCLBIN with nvptx targets. Once this is fixed, +// %{sycl_target_opts} should be added to the SYCLBIN generation run-line. +// REQUIRES: target-spir + +// RUN: %clangxx --offload-new-driver -fsyclbin=object -fsycl-allow-device-image-dependencies -Xclang -fsycl-allow-func-ptr %S/Inputs/link_sycl_inline.cpp -o %t.syclbin +// RUN: %{build} -fsycl-allow-device-image-dependencies -o %t.out +// RUN: %{l0_leak_check} %{run} %t.out %t.syclbin + +#define SYCLBIN_OBJECT_STATE + +#include "Inputs/link_sycl_inline.hpp" diff --git a/sycl/test/e2e_test_requirements/no_sycl_hpp_in_e2e_tests.cpp b/sycl/test/e2e_test_requirements/no_sycl_hpp_in_e2e_tests.cpp index de2b939756ea0..5b97ae7a0cac9 100644 --- a/sycl/test/e2e_test_requirements/no_sycl_hpp_in_e2e_tests.cpp +++ b/sycl/test/e2e_test_requirements/no_sycl_hpp_in_e2e_tests.cpp @@ -6,7 +6,7 @@ // CHECK-DAG: README.md // CHECK-DAG: lit.cfg.py // -// CHECK-NUM-MATCHES: 25 +// CHECK-NUM-MATCHES: 26 // // This test verifies that `` isn't used in E2E tests. Instead, // fine-grained includes should used, see diff --git a/sycl/unittests/kernel-and-program/OutOfResources.cpp b/sycl/unittests/kernel-and-program/OutOfResources.cpp index 3ffe91503e881..cc9f18ae57b21 100644 --- a/sycl/unittests/kernel-and-program/OutOfResources.cpp +++ b/sycl/unittests/kernel-and-program/OutOfResources.cpp @@ -174,7 +174,7 @@ TEST_P(OutOfResourcesTestSuite, urProgramLink) { auto b3 = sycl::link({b1, b2}); EXPECT_FALSE(outOfResourcesToggle); // one restart due to out of resources, one link per each of b1 and b2. - EXPECT_EQ(nProgramLink, 3); + EXPECT_EQ(nProgramLink, 2); // no programs should be in the cache due to out of resources. { detail::KernelProgramCache::ProgramCache &Cache = From aa35a96347ecf228cb91a17799c5244841d84a8e Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Thu, 26 Jun 2025 07:13:53 -0700 Subject: [PATCH 2/8] Fix formatting Signed-off-by: Larsen, Steffen --- sycl/test-e2e/SYCLBIN/link_sycl_inline_object.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/sycl/test-e2e/SYCLBIN/link_sycl_inline_object.cpp b/sycl/test-e2e/SYCLBIN/link_sycl_inline_object.cpp index d6ff02b6e2633..655815adeecf6 100644 --- a/sycl/test-e2e/SYCLBIN/link_sycl_inline_object.cpp +++ b/sycl/test-e2e/SYCLBIN/link_sycl_inline_object.cpp @@ -1,4 +1,5 @@ -//==-------- link_sycl_inline_object.cpp --- SYCLBIN extension tests --------==// +//==-------- link_sycl_inline_object.cpp --- SYCLBIN extension tests +//--------==// // // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. // See https://llvm.org/LICENSE.txt for license information. From 7a36ba111f957f17b0acd6add18aa80d844a5c39 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Thu, 26 Jun 2025 07:52:09 -0700 Subject: [PATCH 3/8] Free your memory! Signed-off-by: Larsen, Steffen --- sycl/test-e2e/SYCLBIN/Inputs/link_sycl_inline.hpp | 3 +++ 1 file changed, 3 insertions(+) diff --git a/sycl/test-e2e/SYCLBIN/Inputs/link_sycl_inline.hpp b/sycl/test-e2e/SYCLBIN/Inputs/link_sycl_inline.hpp index 72336e20029f0..3dbbb9a070d02 100644 --- a/sycl/test-e2e/SYCLBIN/Inputs/link_sycl_inline.hpp +++ b/sycl/test-e2e/SYCLBIN/Inputs/link_sycl_inline.hpp @@ -75,5 +75,8 @@ int main(int argc, char *argv[]) { } } + sycl::free(Args.FuncPtr, Q); + sycl::free(Args.Ptr, Q); + return Failed; } From d87e098e2ad7986af2c77967fb3d07a88fa20fac Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Fri, 27 Jun 2025 04:09:46 -0700 Subject: [PATCH 4/8] Relax linking restriction instead Signed-off-by: Larsen, Steffen --- sycl/source/detail/kernel_bundle_impl.hpp | 49 ++++++++++++++++++++++- 1 file changed, 48 insertions(+), 1 deletion(-) diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index 14892daf1764a..ef2ae1ce9e6a1 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -245,6 +245,31 @@ class kernel_bundle_impl // TODO: Unify with c'tor for sycl::compile and sycl::build by calling // sycl::join on vector of kernel_bundles + // Due to a bug in L0, specializations with conflicting IDs will overwrite + // each other when linked together, so to avoid this issue we link + // images with specialization constants in separation. + // TODO: Remove when spec const overwriting issue has been fixed in L0. + std::vector ImagesWithSpecConsts; + std::unordered_set> + ImagesWithSpecConstsSet; + for (const kernel_bundle &ObjectBundle : + ObjectBundles) { + for (const DevImgPlainWithDeps &DeviceImageWithDeps : + getSyclObjImpl(ObjectBundle)->MDeviceImages) { + if (std::none_of(DeviceImageWithDeps.begin(), DeviceImageWithDeps.end(), + [](const device_image_plain &DevImg) { + const RTDeviceBinaryImage *BinImg = + getSyclObjImpl(DevImg)->get_bin_image_ref(); + return BinImg && BinImg->getSpecConstants().size(); + })) + continue; + + ImagesWithSpecConsts.push_back(&DeviceImageWithDeps); + for (const device_image_plain &DevImg : DeviceImageWithDeps) + ImagesWithSpecConstsSet.insert(getSyclObjImpl(DevImg)); + } + } + // Collect all unique images. std::vector DevImages; { @@ -253,7 +278,9 @@ class kernel_bundle_impl ObjectBundles) for (const device_image_plain &DevImg : getSyclObjImpl(ObjectBundle)->MUniqueDeviceImages) - DevImagesSet.insert(getSyclObjImpl(DevImg)); + if (ImagesWithSpecConstsSet.find(getSyclObjImpl(DevImg)) == + ImagesWithSpecConstsSet.end()) + DevImagesSet.insert(getSyclObjImpl(DevImg)); DevImages.reserve(DevImagesSet.size()); for (auto It = DevImagesSet.begin(); It != DevImagesSet.end();) DevImages.push_back(createSyclObjFromImpl( @@ -368,6 +395,26 @@ class kernel_bundle_impl // added. } + // ... And link the offline images in separation. (Workaround.) + for (const DevImgPlainWithDeps *DeviceImageWithDeps : + ImagesWithSpecConsts) { + // Skip images which are not compatible with devices provided + if (std::none_of(MDevices.begin(), MDevices.end(), + [DeviceImageWithDeps](const device &Dev) { + return getSyclObjImpl(DeviceImageWithDeps->getMain()) + ->compatible_with_device(Dev); + })) + continue; + + std::vector LinkedResults = + detail::ProgramManager::getInstance().link( + DeviceImageWithDeps->getAll(), MDevices, PropList); + MDeviceImages.insert(MDeviceImages.end(), LinkedResults.begin(), + LinkedResults.end()); + MUniqueDeviceImages.insert(MUniqueDeviceImages.end(), + LinkedResults.begin(), LinkedResults.end()); + } + removeDuplicateImages(); for (const kernel_bundle &Bundle : ObjectBundles) { From 1fc0e36e2ebaad7bdf752b607bce7382003bf345 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Mon, 30 Jun 2025 07:53:28 -0700 Subject: [PATCH 5/8] XFAIL on OCL CPU Signed-off-by: Larsen, Steffen --- sycl/test-e2e/SYCLBIN/link_sycl_inline_input.cpp | 3 +++ sycl/test-e2e/SYCLBIN/link_sycl_inline_object.cpp | 3 +++ 2 files changed, 6 insertions(+) diff --git a/sycl/test-e2e/SYCLBIN/link_sycl_inline_input.cpp b/sycl/test-e2e/SYCLBIN/link_sycl_inline_input.cpp index 4607c85e17493..616ef877fe30f 100644 --- a/sycl/test-e2e/SYCLBIN/link_sycl_inline_input.cpp +++ b/sycl/test-e2e/SYCLBIN/link_sycl_inline_input.cpp @@ -15,6 +15,9 @@ // %{sycl_target_opts} should be added to the SYCLBIN generation run-line. // REQUIRES: target-spir +// XFAIL: opencl && cpu +// XFAIL-TRACKER: CMPLRLLVM-68800 + // RUN: %clangxx --offload-new-driver -fsyclbin=input -fsycl-allow-device-image-dependencies -Xclang -fsycl-allow-func-ptr %S/Inputs/link_sycl_inline.cpp -o %t.syclbin // RUN: %{build} -fsycl-allow-device-image-dependencies -o %t.out // RUN: %{l0_leak_check} %{run} %t.out %t.syclbin diff --git a/sycl/test-e2e/SYCLBIN/link_sycl_inline_object.cpp b/sycl/test-e2e/SYCLBIN/link_sycl_inline_object.cpp index 655815adeecf6..e20d2cfcf6e99 100644 --- a/sycl/test-e2e/SYCLBIN/link_sycl_inline_object.cpp +++ b/sycl/test-e2e/SYCLBIN/link_sycl_inline_object.cpp @@ -16,6 +16,9 @@ // %{sycl_target_opts} should be added to the SYCLBIN generation run-line. // REQUIRES: target-spir +// XFAIL: opencl && cpu +// XFAIL-TRACKER: CMPLRLLVM-68800 + // RUN: %clangxx --offload-new-driver -fsyclbin=object -fsycl-allow-device-image-dependencies -Xclang -fsycl-allow-func-ptr %S/Inputs/link_sycl_inline.cpp -o %t.syclbin // RUN: %{build} -fsycl-allow-device-image-dependencies -o %t.out // RUN: %{l0_leak_check} %{run} %t.out %t.syclbin From fdf07fc64294db3ce04b57fec99b681d01167992 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Tue, 1 Jul 2025 02:02:47 -0700 Subject: [PATCH 6/8] Update requirement comment Signed-off-by: Larsen, Steffen --- sycl/test-e2e/SYCLBIN/link_sycl_inline_input.cpp | 7 ++++--- sycl/test-e2e/SYCLBIN/link_sycl_inline_object.cpp | 7 ++++--- 2 files changed, 8 insertions(+), 6 deletions(-) diff --git a/sycl/test-e2e/SYCLBIN/link_sycl_inline_input.cpp b/sycl/test-e2e/SYCLBIN/link_sycl_inline_input.cpp index 616ef877fe30f..88ebbcf0aef32 100644 --- a/sycl/test-e2e/SYCLBIN/link_sycl_inline_input.cpp +++ b/sycl/test-e2e/SYCLBIN/link_sycl_inline_input.cpp @@ -10,9 +10,10 @@ // -- Test for linking between inline SYCL code and SYCLBIN code. -// Due to the regression in https://github.com/intel/llvm/issues/18432 it will -// fail to build the SYCLBIN with nvptx targets. Once this is fixed, -// %{sycl_target_opts} should be added to the SYCLBIN generation run-line. +// ptxas currently fails to compile images with unresolved symbols. Disable for +// other targets than SPIR-V until this has been resolved. (CMPLRLLVM-68810) +// Note: %{sycl_target_opts} should be added to the SYCLBIN compilation lines +// once fixed. // REQUIRES: target-spir // XFAIL: opencl && cpu diff --git a/sycl/test-e2e/SYCLBIN/link_sycl_inline_object.cpp b/sycl/test-e2e/SYCLBIN/link_sycl_inline_object.cpp index e20d2cfcf6e99..b53bf6344bbbc 100644 --- a/sycl/test-e2e/SYCLBIN/link_sycl_inline_object.cpp +++ b/sycl/test-e2e/SYCLBIN/link_sycl_inline_object.cpp @@ -11,9 +11,10 @@ // -- Test for linking between inline SYCL code and SYCLBIN code. -// Due to the regression in https://github.com/intel/llvm/issues/18432 it will -// fail to build the SYCLBIN with nvptx targets. Once this is fixed, -// %{sycl_target_opts} should be added to the SYCLBIN generation run-line. +// ptxas currently fails to compile images with unresolved symbols. Disable for +// other targets than SPIR-V until this has been resolved. (CMPLRLLVM-68810) +// Note: %{sycl_target_opts} should be added to the SYCLBIN compilation lines +// once fixed. // REQUIRES: target-spir // XFAIL: opencl && cpu From bc1c6ac38f2b7c61cef7ef0000fa52b31d73ea12 Mon Sep 17 00:00:00 2001 From: Steffen Larsen Date: Wed, 2 Jul 2025 11:16:07 +0200 Subject: [PATCH 7/8] Apply suggestions from code review --- sycl/test-e2e/SYCLBIN/link_sycl_inline_input.cpp | 3 +++ sycl/test-e2e/SYCLBIN/link_sycl_inline_object.cpp | 3 +++ 2 files changed, 6 insertions(+) diff --git a/sycl/test-e2e/SYCLBIN/link_sycl_inline_input.cpp b/sycl/test-e2e/SYCLBIN/link_sycl_inline_input.cpp index 88ebbcf0aef32..0848730386cc8 100644 --- a/sycl/test-e2e/SYCLBIN/link_sycl_inline_input.cpp +++ b/sycl/test-e2e/SYCLBIN/link_sycl_inline_input.cpp @@ -19,6 +19,9 @@ // XFAIL: opencl && cpu // XFAIL-TRACKER: CMPLRLLVM-68800 +// XFAIL: linux && arch-intel_gpu_bmg_g21 +// XFAIL-TRACKER: https://github.com/intel/llvm/issues/19258 + // RUN: %clangxx --offload-new-driver -fsyclbin=input -fsycl-allow-device-image-dependencies -Xclang -fsycl-allow-func-ptr %S/Inputs/link_sycl_inline.cpp -o %t.syclbin // RUN: %{build} -fsycl-allow-device-image-dependencies -o %t.out // RUN: %{l0_leak_check} %{run} %t.out %t.syclbin diff --git a/sycl/test-e2e/SYCLBIN/link_sycl_inline_object.cpp b/sycl/test-e2e/SYCLBIN/link_sycl_inline_object.cpp index b53bf6344bbbc..4aace113f4fd1 100644 --- a/sycl/test-e2e/SYCLBIN/link_sycl_inline_object.cpp +++ b/sycl/test-e2e/SYCLBIN/link_sycl_inline_object.cpp @@ -20,6 +20,9 @@ // XFAIL: opencl && cpu // XFAIL-TRACKER: CMPLRLLVM-68800 +// XFAIL: linux && arch-intel_gpu_bmg_g21 +// XFAIL-TRACKER: https://github.com/intel/llvm/issues/19258 + // RUN: %clangxx --offload-new-driver -fsyclbin=object -fsycl-allow-device-image-dependencies -Xclang -fsycl-allow-func-ptr %S/Inputs/link_sycl_inline.cpp -o %t.syclbin // RUN: %{build} -fsycl-allow-device-image-dependencies -o %t.out // RUN: %{l0_leak_check} %{run} %t.out %t.syclbin From 2c97cb8385a527d7cef390543626ae428e38667a Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Wed, 2 Jul 2025 03:34:02 -0700 Subject: [PATCH 8/8] Fix sycl/sycl.hpp count Signed-off-by: Larsen, Steffen --- sycl/test/e2e_test_requirements/no_sycl_hpp_in_e2e_tests.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test/e2e_test_requirements/no_sycl_hpp_in_e2e_tests.cpp b/sycl/test/e2e_test_requirements/no_sycl_hpp_in_e2e_tests.cpp index 5b97ae7a0cac9..13f0fc51bc037 100644 --- a/sycl/test/e2e_test_requirements/no_sycl_hpp_in_e2e_tests.cpp +++ b/sycl/test/e2e_test_requirements/no_sycl_hpp_in_e2e_tests.cpp @@ -6,7 +6,7 @@ // CHECK-DAG: README.md // CHECK-DAG: lit.cfg.py // -// CHECK-NUM-MATCHES: 26 +// CHECK-NUM-MATCHES: 27 // // This test verifies that `` isn't used in E2E tests. Instead, // fine-grained includes should used, see