diff --git a/sycl/source/detail/accessor_impl.hpp b/sycl/source/detail/accessor_impl.hpp index 643247549e414..8eb3bbd05d305 100644 --- a/sycl/source/detail/accessor_impl.hpp +++ b/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; diff --git a/sycl/source/detail/scheduler/commands.cpp b/sycl/source/detail/scheduler/commands.cpp index e71a8d2878db8..c59c6bbe1c413 100644 --- a/sycl/source/detail/scheduler/commands.cpp +++ b/sycl/source/detail/scheduler/commands.cpp @@ -1737,11 +1737,14 @@ ur_result_t MemCpyCommandHost::enqueueImp() { std::vector 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; diff --git a/sycl/source/detail/scheduler/graph_builder.cpp b/sycl/source/detail/scheduler/graph_builder.cpp index 046d3f25d066e..9f6cd3300775d 100644 --- a/sycl/source/detail/scheduler/graph_builder.cpp +++ b/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()) { diff --git a/sycl/test-e2e/Basic/accessor/ranged_no_init.cpp b/sycl/test-e2e/Basic/accessor/ranged_no_init.cpp new file mode 100644 index 0000000000000..d64989140d3be --- /dev/null +++ b/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 + +#include +#include +#include + +int main() { + constexpr std::size_t n = 16, lo = 4, hi = 8; + + int observed[n]; + { + sycl::buffer 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 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; +}