From 1a20ff6fb4a0d7e99d82e17e5e4e9a92107cd15b Mon Sep 17 00:00:00 2001 From: "Skrobot, Daniel" Date: Tue, 3 Nov 2026 14:45:02 +0100 Subject: [PATCH 1/3] Throw kernel_not_supported for copies to/from unused device_image_scope device_global --- sycl/include/sycl/handler.hpp | 22 ++++++++++++++++++---- sycl/source/detail/memory_manager.cpp | 2 +- 2 files changed, 19 insertions(+), 5 deletions(-) diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index 1be6bce3d1e4e..4b65d4838ed76 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -2220,8 +2220,15 @@ class __SYCL_EXPORT handler { ext::oneapi::experimental::device_image_scope_key>(); if (!detail::isDeviceGlobalUsedInKernel(&Dest)) { - // If the corresponding device_global isn't used in any kernels, we fall - // back to doing the memory operation on host-only. + // A device_global with device_image_scope lives in a device image, so + // there is nothing to copy to if no kernel accesses it. + if constexpr (IsDeviceImageScoped) + throw sycl::exception( + make_error_code(errc::kernel_not_supported), + "No kernel accesses the device_global with device_image_scope."); + + // A device_global without device_image_scope that isn't used in any + // kernels falls back to doing the memory operation on host-only. memcpyToHostOnlyDeviceGlobal(&Dest, Src, sizeof(T), IsDeviceImageScoped, NumBytes, DestOffset); return; @@ -2257,8 +2264,15 @@ class __SYCL_EXPORT handler { ext::oneapi::experimental::device_image_scope_key>(); if (!detail::isDeviceGlobalUsedInKernel(&Src)) { - // If the corresponding device_global isn't used in any kernels, we fall - // back to doing the memory operation on host-only. + // A device_global with device_image_scope lives in a device image, so + // there is nothing to copy from if no kernel accesses it. + if constexpr (IsDeviceImageScoped) + throw sycl::exception( + make_error_code(errc::kernel_not_supported), + "No kernel accesses the device_global with device_image_scope."); + + // A device_global without device_image_scope that isn't used in any + // kernels falls back to doing the memory operation on host-only. memcpyFromHostOnlyDeviceGlobal(Dest, &Src, IsDeviceImageScoped, NumBytes, SrcOffset); return; diff --git a/sycl/source/detail/memory_manager.cpp b/sycl/source/detail/memory_manager.cpp index 6e5b9bea10420..9f5ab7fd1e410 100644 --- a/sycl/source/detail/memory_manager.cpp +++ b/sycl/source/detail/memory_manager.cpp @@ -1104,7 +1104,7 @@ getOrBuildProgramForDeviceGlobal(queue_impl &Queue, // If there are no kernels using the device_global we cannot proceed. if (DeviceGlobalEntry->MImageIdentifiers.size() == 0) - throw sycl::exception(make_error_code(errc::invalid), + throw sycl::exception(make_error_code(errc::kernel_not_supported), "No image exists with the device_global."); // Look for cached programs with the device_global. From 0ffd7d7b650e3303afc9e3a623d15139dcbe9529 Mon Sep 17 00:00:00 2001 From: "Skrobot, Daniel" Date: Tue, 3 Nov 2026 15:51:16 +0100 Subject: [PATCH 2/3] updated e2e test and added new cases in unittest for device global --- .../DeviceGlobal/device_global_unused_dis.cpp | 41 +++++++++++++--- sycl/unittests/Extensions/DeviceGlobal.cpp | 47 +++++++++++++++++++ 2 files changed, 81 insertions(+), 7 deletions(-) diff --git a/sycl/test-e2e/DeviceGlobal/device_global_unused_dis.cpp b/sycl/test-e2e/DeviceGlobal/device_global_unused_dis.cpp index 90e9e950e65ad..6efb290909a92 100644 --- a/sycl/test-e2e/DeviceGlobal/device_global_unused_dis.cpp +++ b/sycl/test-e2e/DeviceGlobal/device_global_unused_dis.cpp @@ -1,11 +1,38 @@ -// RUN: %{build} -fsycl-device-code-split=per_source -DUSE_DEVICE_IMAGE_SCOPE -o %t.out +// RUN: %{build} -DUSE_DEVICE_IMAGE_SCOPE -o %t.out // RUN: %{run} %t.out -// Tests that device_global with device_image_scope and no kernel uses can be -// copied to and from. -// NOTE: USE_DEVICE_IMAGE_SCOPE needs both kernels to be in the same image so -// we set -fsycl-device-code-split=per_source. +// Tests that copying to or from a device_global with device_image_scope that +// no kernel uses throws errc::kernel_not_supported. -#include "device_global_unused.hpp" +#include "common.hpp" -int main() { return test(); } +device_global memcpyDeviceGlobal; +device_global copyDeviceGlobal; + +namespace { +template bool throwsKernelNotSupported(FnT kFn) { + try { + kFn(); + } catch (const sycl::exception &e) { + return e.code() == errc::kernel_not_supported; + } + return false; +} +} // namespace + +int main() { + queue q; + int value = 42; + int failures = 0; + + failures += + !throwsKernelNotSupported([&] { q.memcpy(memcpyDeviceGlobal, &value); }); + failures += + !throwsKernelNotSupported([&] { q.memcpy(&value, memcpyDeviceGlobal); }); + failures += + !throwsKernelNotSupported([&] { q.copy(&value, copyDeviceGlobal); }); + failures += + !throwsKernelNotSupported([&] { q.copy(copyDeviceGlobal, &value); }); + + return failures; +} diff --git a/sycl/unittests/Extensions/DeviceGlobal.cpp b/sycl/unittests/Extensions/DeviceGlobal.cpp index 6aa8c32405830..d986805af6f61 100644 --- a/sycl/unittests/Extensions/DeviceGlobal.cpp +++ b/sycl/unittests/Extensions/DeviceGlobal.cpp @@ -34,6 +34,14 @@ sycl::ext::oneapi::experimental::device_global< sycl::ext::oneapi::experimental::device_image_scope))> DeviceGlobalImgScope; +sycl::ext::oneapi::experimental::device_global + DeviceGlobalUnused; +sycl::ext::oneapi::experimental::device_global< + DeviceGlobalElemType, + decltype(sycl::ext::oneapi::experimental::properties( + sycl::ext::oneapi::experimental::device_image_scope))> + DeviceGlobalImgScopeUnused; + namespace sycl { inline namespace _V1 { namespace detail { @@ -733,3 +741,42 @@ TEST_F(DeviceGlobalTest, DeviceGlobalImgScopeUseBeforeMemcpyFrom) { EXPECT_EQ(MockDeviceGlobalImgScopeMem[0], Vals[0]); EXPECT_EQ(MockDeviceGlobalImgScopeMem[1], Vals[1]); } + +TEST_F(DeviceGlobalTest, DeviceGlobalImgScopeUnusedCopyMemcpyThrows) { + int Vals[2] = {42, 1234}; + auto ExpectKernelNotSupported = [](auto &&KFn) { + try { + KFn(); + FAIL() << "Expected an exception"; + } catch (const sycl::exception &Excpt) { + EXPECT_EQ(Excpt.code(), sycl::errc::kernel_not_supported); + } + }; + + ExpectKernelNotSupported([&] { Q.copy(Vals, DeviceGlobalImgScopeUnused); }); + ExpectKernelNotSupported([&] { Q.copy(DeviceGlobalImgScopeUnused, Vals); }); + ExpectKernelNotSupported([&] { Q.memcpy(DeviceGlobalImgScopeUnused, Vals); }); + ExpectKernelNotSupported([&] { Q.memcpy(Vals, DeviceGlobalImgScopeUnused); }); + ExpectKernelNotSupported([&] { + Q.submit([&](sycl::handler &CGH) { + CGH.copy(Vals, DeviceGlobalImgScopeUnused); + }); + }); + ExpectKernelNotSupported([&] { + Q.submit([&](sycl::handler &CGH) { + CGH.copy(DeviceGlobalImgScopeUnused, Vals); + }); + }); +} + +// Without device_image_scope an unused device_global is still usable. +TEST_F(DeviceGlobalTest, DeviceGlobalUnusedCopyMemcpy) { + int WriteVals[2] = {42, 1234}; + int ReadVals[2] = {0, 0}; + + EXPECT_NO_THROW(Q.copy(WriteVals, DeviceGlobalUnused).wait()); + EXPECT_NO_THROW(Q.copy(DeviceGlobalUnused, ReadVals).wait()); + + EXPECT_EQ(ReadVals[0], WriteVals[0]); + EXPECT_EQ(ReadVals[1], WriteVals[1]); +} From e610ca86fcbb317d291ec98d11ce8245e60470e7 Mon Sep 17 00:00:00 2001 From: "Skrobot, Daniel" Date: Tue, 3 Nov 2026 17:55:55 +0100 Subject: [PATCH 3/3] apply feedback --- sycl/include/sycl/handler.hpp | 18 ++++++++++-------- .../DeviceGlobal/device_global_unused_dis.cpp | 2 +- sycl/unittests/Extensions/DeviceGlobal.cpp | 10 ++++++++++ 3 files changed, 21 insertions(+), 9 deletions(-) diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index 4b65d4838ed76..f4a89c9aa0195 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -2222,10 +2222,7 @@ class __SYCL_EXPORT handler { if (!detail::isDeviceGlobalUsedInKernel(&Dest)) { // A device_global with device_image_scope lives in a device image, so // there is nothing to copy to if no kernel accesses it. - if constexpr (IsDeviceImageScoped) - throw sycl::exception( - make_error_code(errc::kernel_not_supported), - "No kernel accesses the device_global with device_image_scope."); + throwIfDeviceGlobalNotUsedInKernel(); // A device_global without device_image_scope that isn't used in any // kernels falls back to doing the memory operation on host-only. @@ -2266,10 +2263,7 @@ class __SYCL_EXPORT handler { if (!detail::isDeviceGlobalUsedInKernel(&Src)) { // A device_global with device_image_scope lives in a device image, so // there is nothing to copy from if no kernel accesses it. - if constexpr (IsDeviceImageScoped) - throw sycl::exception( - make_error_code(errc::kernel_not_supported), - "No kernel accesses the device_global with device_image_scope."); + throwIfDeviceGlobalNotUsedInKernel(); // A device_global without device_image_scope that isn't used in any // kernels falls back to doing the memory operation on host-only. @@ -3058,6 +3052,14 @@ class __SYCL_EXPORT handler { void incrementArgShift(int Shift); + template + void throwIfDeviceGlobalNotUsedInKernel() const { + if constexpr (IsDeviceImageScoped) + throw sycl::exception( + make_error_code(errc::kernel_not_supported), + "No kernel accesses the device_global with device_image_scope."); + } + protected: /// Registers event dependencies in this command group. void depends_on(const detail::EventImplPtr &Event); diff --git a/sycl/test-e2e/DeviceGlobal/device_global_unused_dis.cpp b/sycl/test-e2e/DeviceGlobal/device_global_unused_dis.cpp index 6efb290909a92..986549009421f 100644 --- a/sycl/test-e2e/DeviceGlobal/device_global_unused_dis.cpp +++ b/sycl/test-e2e/DeviceGlobal/device_global_unused_dis.cpp @@ -13,7 +13,7 @@ namespace { template bool throwsKernelNotSupported(FnT kFn) { try { kFn(); - } catch (const sycl::exception &e) { + } catch (const exception &e) { return e.code() == errc::kernel_not_supported; } return false; diff --git a/sycl/unittests/Extensions/DeviceGlobal.cpp b/sycl/unittests/Extensions/DeviceGlobal.cpp index d986805af6f61..8afcdbe1d0025 100644 --- a/sycl/unittests/Extensions/DeviceGlobal.cpp +++ b/sycl/unittests/Extensions/DeviceGlobal.cpp @@ -767,6 +767,16 @@ TEST_F(DeviceGlobalTest, DeviceGlobalImgScopeUnusedCopyMemcpyThrows) { CGH.copy(DeviceGlobalImgScopeUnused, Vals); }); }); + ExpectKernelNotSupported([&] { + Q.submit([&](sycl::handler &CGH) { + CGH.memcpy(DeviceGlobalImgScopeUnused, Vals); + }); + }); + ExpectKernelNotSupported([&] { + Q.submit([&](sycl::handler &CGH) { + CGH.memcpy(Vals, DeviceGlobalImgScopeUnused); + }); + }); } // Without device_image_scope an unused device_global is still usable.