[SYCL] no_init property limited to range of effected accessor.

sycl/source/detail/accessor_impl.hpp

@@ -114,6 +114,15 @@ class AccessorImplHost {
 
   void *&MData = MAccData.MData;
 
+  // True iff this accessor covers the whole underlying memory object.
+  // discard_* elision of host<->device transfers is only safe when this holds
+  // — for a ranged accessor, no_init applies only to elements within the
+  // accessor's range; elements outside must be preserved (SYCL 2020 §4.7.6.4).
+  bool isFullMemoryAccess() const {
+    return !MIsSubBuffer && MOffset == id<3>{0, 0, 0} &&
+           MAccessRange == MMemoryRange;
+  }
+
   Command *MBlockedCmd = nullptr;
 
   bool PerWI = false;

sycl/source/detail/scheduler/commands.cpp

@@ -1737,11 +1737,14 @@ ur_result_t MemCpyCommandHost::enqueueImp() {
   std::vector<ur_event_handle_t> RawEvents = getUrEvents(EventImpls);
 
   ur_event_handle_t UREvent = nullptr;
-  // Omit copying if mode is discard one.
+  // Omit copying if mode is discard one and the accessor covers the full
+  // memory object; a ranged discard accessor must still preserve elements
+  // outside its range (SYCL 2020 §4.7.6.4).
   // TODO: Handle this at the graph building time by, for example, creating
   // empty node instead of memcpy.
-  if (MDstReq.MAccessMode == access::mode::discard_read_write ||
-      MDstReq.MAccessMode == access::mode::discard_write) {
+  if ((MDstReq.MAccessMode == access::mode::discard_read_write ||
+       MDstReq.MAccessMode == access::mode::discard_write) &&
+      MDstReq.isFullMemoryAccess()) {
     Command::waitForEvents(Queue, EventImpls, UREvent);
 
     return UR_RESULT_SUCCESS;

sycl/source/detail/scheduler/graph_builder.cpp

@@ -385,8 +385,9 @@ Scheduler::GraphBuilder::insertMemoryMove(MemObjRecord *Record,
     NewCmd = insertMapUnmapForLinkedCmds(AllocaCmdSrc, AllocaCmdDst, MapMode);
     Record->MHostAccess = MapMode;
   } else {
-    if ((Req->MAccessMode == access::mode::discard_write) ||
-        (Req->MAccessMode == access::mode::discard_read_write)) {
+    if (((Req->MAccessMode == access::mode::discard_write) ||
+         (Req->MAccessMode == access::mode::discard_read_write)) &&
+        Req->isFullMemoryAccess()) {
       Record->setCurContext(Context);
       return nullptr;
     } else {
@@ -728,9 +729,11 @@ AllocaCommandBase *Scheduler::GraphBuilder::getOrCreateAllocaForReq(
       // might need to also create a host alloca right away in order to perform
       // the initial memory write.
       if (Record->MAllocaCommands.empty()) {
-        if (!HostUnifiedMemory &&
-            Req->MAccessMode != access::mode::discard_write &&
-            Req->MAccessMode != access::mode::discard_read_write) {
+        const bool IsFullDiscard =
+            (Req->MAccessMode == access::mode::discard_write ||
+             Req->MAccessMode == access::mode::discard_read_write) &&
+            Req->isFullMemoryAccess();
+        if (!HostUnifiedMemory && !IsFullDiscard) {
           // There's no need to make a host allocation if the buffer is not
           // initialized with user data.
           if (MemObj->hasUserDataPtr()) {

sycl/test-e2e/Basic/accessor/ranged_no_init.cpp

@@ -0,0 +1,64 @@
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+// RUN: env SYCL_HOST_UNIFIED_MEMORY=0 %{run} %t.out
+
+// Regression test for a scheduler bug where sycl::property::no_init on a
+// ranged accessor caused elements OUTSIDE the accessor's range to be
+// discarded (zeroed) on devices without host-unified memory.
+//
+// SYCL 2020 §4.7.6.4 (Table 24, property::no_init):
+//   "If this is a ranged accessor, this applies only to the elements within
+//    the accessor's range. The values of unwritten elements outside of this
+//    range are preserved."
+
+#include <sycl/detail/core.hpp>
+
+#include <cstddef>
+#include <cstdlib>
+#include <iostream>
+
+int main() {
+  constexpr std::size_t n = 16, lo = 4, hi = 8;
+
+  int observed[n];
+  {
+    sycl::buffer<int, 1> buf{sycl::range<1>(n)};
+
+    {
+      sycl::host_accessor h{buf, sycl::write_only};
+      for (std::size_t i = 0; i < n; ++i)
+        h[i] = 42;
+    }
+
+    sycl::queue{}
+        .submit([&](sycl::handler &cgh) {
+          sycl::accessor<int, 1, sycl::access_mode::write> a{
+              buf, cgh, sycl::range<1>(hi - lo), sycl::id<1>(lo),
+              sycl::property_list{sycl::no_init}};
+          cgh.parallel_for(sycl::range<1>(hi - lo),
+                           [=](sycl::id<1> i) { a[i] = 7; });
+        })
+        .wait_and_throw();
+
+    sycl::host_accessor h{buf, sycl::read_only};
+    for (std::size_t i = 0; i < n; ++i)
+      observed[i] = h[i];
+  }
+
+  int failures = 0;
+  for (std::size_t i = 0; i < n; ++i) {
+    const int expected = (i >= lo && i < hi) ? 7 : 42;
+    if (observed[i] != expected) {
+      std::cerr << "  index " << i << ": got " << observed[i] << ", expected "
+                << expected << "\n";
+      ++failures;
+    }
+  }
+  if (failures) {
+    std::cerr << failures
+              << " element(s) outside the ranged no_init accessor were not "
+                 "preserved.\n";
+    return 1;
+  }
+  return 0;
+}