[SYCL] Throw kernel_not_supported for copies to/from unused device_image_scope device_global

sycl/include/sycl/handler.hpp

@@ -2220,8 +2220,12 @@ 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.
+      throwIfDeviceGlobalNotUsedInKernel<IsDeviceImageScoped>();
+
+      // 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 +2261,12 @@ 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.
+      throwIfDeviceGlobalNotUsedInKernel<IsDeviceImageScoped>();
+
+      // 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;
@@ -3044,6 +3052,14 @@ class __SYCL_EXPORT handler {
 
   void incrementArgShift(int Shift);
 
+  template <bool IsDeviceImageScoped>
+  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);

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.

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<int, TestProperties> memcpyDeviceGlobal;
+device_global<int, TestProperties> copyDeviceGlobal;
+
+namespace {
+template <typename FnT> bool throwsKernelNotSupported(FnT kFn) {
+  try {
+    kFn();
+  } catch (const 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;
+}

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<DeviceGlobalElemType>
+    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,52 @@ 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);
+    });
+  });
+  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.
+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]);
+}