diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp index f9f5feab38184..d3915e4a801d1 100644 --- a/sycl/include/CL/sycl/handler.hpp +++ b/sycl/include/CL/sycl/handler.hpp @@ -522,6 +522,79 @@ class __SYCL_EXPORT handler { return true; } +#ifndef __SYCL_DEVICE_ONLY__ + /// Copies the content of memory object accessed by Src into the memory + /// pointed by Dst. + /// + /// \param Src is a source SYCL accessor. + /// \param Dst is a pointer to destination memory. + template + detail::enable_if_t<(Dim > 0)> + copyAccToPtrHost(accessor Src, + TDst *Dst) { + range Range = Src.get_range(); + parallel_for> + (Range, [=](id Index) { + size_t LinearIndex = Index[0]; + for (int I = 1; I < Dim; ++I) + LinearIndex += Range[I] * Index[I]; + (reinterpret_cast(Dst))[LinearIndex] = Src[Index]; + }); + } + + /// Copies 1 element accessed by 0-dimensional accessor Src into the memory + /// pointed by Dst. + /// + /// \param Src is a source SYCL accessor. + /// \param Dst is a pointer to destination memory. + template + detail::enable_if_t + copyAccToPtrHost(accessor Src, + TDst *Dst) { + single_task> + ([=]() { + *Dst = readFromFirstAccElement(Src); + }); + } + + /// Copies the memory pointed by Src into the memory accessed by Dst. + /// + /// \param Src is a pointer to source memory. + /// \param Dst is a destination SYCL accessor. + template + detail::enable_if_t<(Dim > 0)> + copyPtrToAccHost(TDst *Src, + accessor Dst) { + range Range = Dst.get_range(); + parallel_for> + (Range, [=](id Index) { + size_t LinearIndex = Index[0]; + for (int I = 1; I < Dim; ++I) + LinearIndex += Range[I] * Index[I]; + Dst[Index] = (reinterpret_cast(Src))[LinearIndex]; + }); + } + + /// Copies 1 element pointed by Src to memory accessed by 0-dimensional + /// accessor Dst. + /// + /// \param Src is a pointer to source memory. + /// \param Dst is a destination SYCL accessor. + template + detail::enable_if_t + copyPtrToAccHost(TDst *Src, + accessor Dst) { + single_task> + ([=]() { + writeToFirstAccElement(Dst, *Src); + }); + } +#endif // __SYCL_DEVICE_ONLY__ + constexpr static bool isConstOrGlobal(access::target AccessTarget) { return AccessTarget == access::target::global_buffer || AccessTarget == access::target::constant_buffer; @@ -1206,7 +1279,7 @@ class __SYCL_EXPORT handler { // Explicit copy operations API - /// Copies the contents of memory object accessed by Src into the memory + /// Copies the content of memory object accessed by Src into the memory /// pointed by Dst. /// /// Source must have at least as many bytes as the range accessed by Dst. @@ -1228,7 +1301,7 @@ class __SYCL_EXPORT handler { copy(Src, RawDstPtr); } - /// Copies the contents of memory pointed by Src into the memory object + /// Copies the content of memory pointed by Src into the memory object /// accessed by Dst. /// /// Source must have at least as many bytes as the range accessed by Dst. @@ -1251,14 +1324,13 @@ class __SYCL_EXPORT handler { copy(RawSrcPtr, Dst); } - /// Copies the contents of memory object accessed by Src into the memory + /// Copies the content of memory object accessed by Src into the memory /// pointed by Dst. /// /// Source must have at least as many bytes as the range accessed by Dst. /// /// \param Src is a source SYCL accessor. /// \param Dst is a pointer to destination memory. - // TODO: support 0-dimensional and atomic accessors. template @@ -1270,17 +1342,8 @@ class __SYCL_EXPORT handler { #ifndef __SYCL_DEVICE_ONLY__ if (MIsHost) { // TODO: Temporary implementation for host. Should be handled by memory - // manger. - range Range = Src.get_range(); - parallel_for< class __copyAcc2Ptr< T_Src, T_Dst, Dims, AccessMode, - AccessTarget, IsPlaceholder>> - (Range, [=](id Index) { - size_t LinearIndex = Index[0]; - for (int I = 1; I < Dims; ++I) - LinearIndex += Range[I] * Index[I]; - ((T_Src *)Dst)[LinearIndex] = Src[Index]; - }); - + // manager. + copyAccToPtrHost(Src, Dst); return; } #endif @@ -1297,14 +1360,13 @@ class __SYCL_EXPORT handler { MAccStorage.push_back(std::move(AccImpl)); } - /// Copies the contents of memory pointed by Src into the memory object + /// Copies the content of memory pointed by Src into the memory object /// accessed by Dst. /// /// Source must have at least as many bytes as the range accessed by Dst. /// /// \param Src is a pointer to source memory. /// \param Dst is a destination SYCL accessor. - // TODO: support 0-dimensional and atomic accessors. template @@ -1317,17 +1379,8 @@ class __SYCL_EXPORT handler { #ifndef __SYCL_DEVICE_ONLY__ if (MIsHost) { // TODO: Temporary implementation for host. Should be handled by memory - // manger. - range Range = Dst.get_range(); - parallel_for< class __copyPtr2Acc< T_Src, T_Dst, Dims, AccessMode, - AccessTarget, IsPlaceholder>> - (Range, [=](id Index) { - size_t LinearIndex = Index[0]; - for (int I = 1; I < Dims; ++I) - LinearIndex += Range[I] * Index[I]; - - Dst[Index] = ((T_Dst *)Src)[LinearIndex]; - }); + // manager. + copyPtrToAccHost(Src, Dst); return; } #endif @@ -1344,7 +1397,7 @@ class __SYCL_EXPORT handler { MAccStorage.push_back(std::move(AccImpl)); } - /// Copies the contents of memory object accessed by Src to the memory + /// Copies the content of memory object accessed by Src to the memory /// object accessed by Dst. /// /// Dst must have at least as many bytes as the range accessed by Src. diff --git a/sycl/test/basic_tests/handler/handler_mem_op.cpp b/sycl/test/basic_tests/handler/handler_mem_op.cpp index cc51952431097..4fbd21a2072f3 100644 --- a/sycl/test/basic_tests/handler/handler_mem_op.cpp +++ b/sycl/test/basic_tests/handler/handler_mem_op.cpp @@ -251,6 +251,20 @@ template void test_copy_ptr_acc() { for (size_t I = 0; I < Size; ++I) { assert(Data[I] == Values[I]); } + + // Check copy from memory to 0-dimensional accessor. + T SrcValue = 99; + T DstValue = 0; + { + buffer DstBuf(&DstValue, range<1>(1)); + queue Queue; + Queue.submit([&](handler &Cgh) { + accessor + DstAcc(DstBuf, Cgh); + Cgh.copy(&SrcValue, DstAcc); + }); + } + assert(DstValue == 99); } template void test_copy_acc_ptr() { @@ -272,6 +286,38 @@ template void test_copy_acc_ptr() { for (size_t I = 0; I < Size; ++I) { assert(Data[I] == Values[I]); } + + // Check copy from 0-dimensional accessor to memory + T SrcValue = 99; + T DstValue = 0; + { + buffer SrcBuf(&SrcValue, range<1>(1)); + queue Queue; + Queue.submit([&](handler &Cgh) { + accessor + SrcAcc(SrcBuf, Cgh); + Cgh.copy(SrcAcc, &DstValue); + }); + } + assert(DstValue == 99); + + // Check copy from 0-dimensional placeholder accessor to memory + SrcValue = 77; + DstValue = 0; + { + buffer SrcBuf(&SrcValue, range<1>(1)); + accessor + SrcAcc(SrcBuf); + { + queue Queue; + Queue.submit([&](handler &Cgh) { + Cgh.require(SrcAcc); + Cgh.copy(SrcAcc, &DstValue); + }); + } + } + assert(DstValue == 77); } template void test_copy_shared_ptr_acc() {