From 8e1f32b02fcfd217822de1626c6f016eb1137fd4 Mon Sep 17 00:00:00 2001 From: Dmitry Sidorov Date: Wed, 17 Apr 2019 13:17:25 +0300 Subject: [PATCH] [SYCL] Add support for discard access modes cl::sycl::discard_write and cl::sycl::discard_read_write are now supported. Do not proceed copying buffer from device to host and vice versa if any of mentioned modes are set. Signed-off-by: Dmitry Sidorov --- sycl/include/CL/sycl/detail/buffer_impl.hpp | 26 +++++++---- .../CL/sycl/detail/scheduler/requirements.h | 6 +-- sycl/test/basic_tests/accessor/accessor.cpp | 46 +++++++++++++++++++ 3 files changed, 67 insertions(+), 11 deletions(-) diff --git a/sycl/include/CL/sycl/detail/buffer_impl.hpp b/sycl/include/CL/sycl/detail/buffer_impl.hpp index 6d22ec91bb102..5da01ca38418d 100644 --- a/sycl/include/CL/sycl/detail/buffer_impl.hpp +++ b/sycl/include/CL/sycl/detail/buffer_impl.hpp @@ -262,7 +262,7 @@ template class buffer_impl { public: void moveMemoryTo(QueueImplPtr Queue, std::vector DepEvents, - EventImplPtr Event); + EventImplPtr Event, cl::sycl::access::mode Mode); void fill(QueueImplPtr Queue, std::vector DepEvents, EventImplPtr Event, const void *Pattern, size_t PatternSize, @@ -281,7 +281,7 @@ template class buffer_impl { bool isValidAccessToMem(cl::sycl::access::mode AccessMode); void allocate(QueueImplPtr Queue, std::vector DepEvents, - EventImplPtr Event, cl::sycl::access::mode mode); + EventImplPtr Event); cl_mem getOpenCLMem() const; @@ -409,7 +409,7 @@ void buffer_impl::copy( template void buffer_impl::moveMemoryTo( QueueImplPtr Queue, std::vector DepEvents, - EventImplPtr Event) { + EventImplPtr Event, cl::sycl::access::mode Mode) { ContextImplPtr Context = detail::getSyclObjImpl(Queue->get_context()); @@ -431,6 +431,10 @@ void buffer_impl::moveMemoryTo( // Copy from OCL device to host device. if (!OCLState.Queue->is_host() && Queue->is_host()) { + if (Mode == cl::sycl::access::mode::discard_write && + Mode == cl::sycl::access::mode::discard_read_write) + return; + const size_t ByteSize = get_size(); std::vector CLEvents = @@ -484,13 +488,18 @@ void buffer_impl::moveMemoryTo( std::vector CLEvents = detail::getOrWaitEvents(std::move(DepEvents), Context); + if (Mode == cl::sycl::access::mode::discard_write && + Mode == cl::sycl::access::mode::discard_read_write) + return; + cl_event &WriteBufEvent = Event->getHandleRef(); // Enqueue copying from host to new OCL buffer. Error = clEnqueueWriteBuffer(OCLState.Queue->getHandleRef(), OCLState.Mem, /*blocking_write=*/CL_FALSE, /*offset=*/0, - ByteSize, BufPtr, CLEvents.size(), CLEvents.data(), - &WriteBufEvent); // replace &WriteBufEvent to NULL + ByteSize, BufPtr, CLEvents.size(), + CLEvents.data(), /*replace &WriteBufEvent + to NULL*/ &WriteBufEvent); CHECK_OCL_CODE(Error); Event->setContextImpl(Context); @@ -507,8 +516,10 @@ buffer_impl::convertSycl2OCLMode(cl::sycl::access::mode mode) { case cl::sycl::access::mode::read: return CL_MEM_READ_ONLY; case cl::sycl::access::mode::write: + case cl::sycl::access::mode::discard_write: return CL_MEM_WRITE_ONLY; case cl::sycl::access::mode::read_write: + case cl::sycl::access::mode::discard_read_write: case cl::sycl::access::mode::atomic: return CL_MEM_READ_WRITE; default: @@ -534,8 +545,7 @@ bool buffer_impl::isValidAccessToMem( template void buffer_impl::allocate(QueueImplPtr Queue, std::vector DepEvents, - EventImplPtr Event, - cl::sycl::access::mode mode) { + EventImplPtr Event) { detail::waitEvents(DepEvents); @@ -556,7 +566,7 @@ void buffer_impl::allocate(QueueImplPtr Queue, cl_int Error; cl_mem Mem = - clCreateBuffer(Context->getHandleRef(), convertSycl2OCLMode(mode), + clCreateBuffer(Context->getHandleRef(), CL_MEM_READ_WRITE, ByteSize, nullptr, &Error); CHECK_OCL_CODE(Error); diff --git a/sycl/include/CL/sycl/detail/scheduler/requirements.h b/sycl/include/CL/sycl/detail/scheduler/requirements.h index 5bbf9af10fb49..f7fae8d6ac963 100644 --- a/sycl/include/CL/sycl/detail/scheduler/requirements.h +++ b/sycl/include/CL/sycl/detail/scheduler/requirements.h @@ -105,15 +105,15 @@ class BufferStorage : public BufferRequirement { void allocate(QueueImplPtr Queue, std::vector DepEvents, EventImplPtr Event) override { assert(m_Buffer != nullptr && "BufferStorage::m_Buffer is nullptr"); - m_Buffer->allocate(std::move(Queue), std::move(DepEvents), std::move(Event), - Mode); + m_Buffer->allocate(std::move(Queue), std::move(DepEvents), + std::move(Event)); } void moveMemoryTo(QueueImplPtr Queue, std::vector DepEvents, EventImplPtr Event) override { assert(m_Buffer != nullptr && "BufferStorage::m_Buffer is nullptr"); m_Buffer->moveMemoryTo(std::move(Queue), std::move(DepEvents), - std::move(Event)); + std::move(Event), Mode); } void fill(QueueImplPtr Queue, std::vector DepEvents, diff --git a/sycl/test/basic_tests/accessor/accessor.cpp b/sycl/test/basic_tests/accessor/accessor.cpp index 3af6e1f919778..c530730e3391c 100644 --- a/sycl/test/basic_tests/accessor/accessor.cpp +++ b/sycl/test/basic_tests/accessor/accessor.cpp @@ -174,4 +174,50 @@ int main() { } } } + + // Discard write accessor. + { + try { + sycl::queue Queue; + sycl::buffer buf(sycl::range<1>(3)); + + Queue.submit([&](sycl::handler& cgh) { + auto dev_acc = buf.get_access(cgh); + + cgh.parallel_for( + sycl::range<1>{3}, + [=](sycl::id<1> index) { dev_acc[index] = 42; }); + }); + + auto host_acc = buf.get_access(); + for (int i = 0; i != 3; ++i) + assert(host_acc[i] == 42); + + } catch (cl::sycl::exception e) { + std::cout << "SYCL exception caught: " << e.what(); + return 1; + } + } + + // Discard read-write accessor. + { + try { + sycl::queue Queue; + sycl::buffer buf(sycl::range<1>(3)); + + Queue.submit([&](sycl::handler& cgh) { + auto dev_acc = buf.get_access(cgh); + + cgh.parallel_for( + sycl::range<1>{3}, + [=](sycl::id<1> index) { dev_acc[index] = 42; }); + }); + + auto host_acc = + buf.get_access(); + } catch (cl::sycl::exception e) { + std::cout << "SYCL exception caught: " << e.what(); + return 1; + } + } }