From 602f8c27f5307f1da966df2fc26745ecd0e78fc9 Mon Sep 17 00:00:00 2001 From: Mehdi Goli Date: Fri, 20 Jan 2017 18:23:20 +0000 Subject: Reverting back to the previous TensorDeviceSycl.h as the total number of buffer is not enough for tensorflow. --- .../Eigen/CXX11/src/Tensor/TensorDeviceSycl.h | 83 +++++++++++----------- 1 file changed, 41 insertions(+), 42 deletions(-) (limited to 'unsupported/Eigen/CXX11/src/Tensor/TensorDeviceSycl.h') diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorDeviceSycl.h b/unsupported/Eigen/CXX11/src/Tensor/TensorDeviceSycl.h index a30090714..722a5d894 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorDeviceSycl.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorDeviceSycl.h @@ -15,16 +15,13 @@ #if defined(EIGEN_USE_SYCL) && !defined(EIGEN_CXX11_TENSOR_TENSOR_DEVICE_SYCL_H) #define EIGEN_CXX11_TENSOR_TENSOR_DEVICE_SYCL_H -#include "TensorSyclLegacyPointer.h" - namespace Eigen { #define ConvertToActualTypeSycl(Scalar, buf_acc) reinterpret_cast::pointer_t>((&(*buf_acc.get_pointer()))) template class MemCopyFunctor { public: - MemCopyFunctor(read_accessor src_acc, write_accessor dst_acc, size_t rng, size_t i, size_t offset) - : m_src_acc(src_acc), m_dst_acc(dst_acc), m_rng(rng), m_i(i), m_offset(offset) {} + MemCopyFunctor(read_accessor src_acc, write_accessor dst_acc, size_t rng, size_t i, size_t offset) : m_src_acc(src_acc), m_dst_acc(dst_acc), m_rng(rng), m_i(i), m_offset(offset) {} void operator()(cl::sycl::nd_item<1> itemID) { auto src_ptr = ConvertToActualTypeSycl(Scalar, m_src_acc); @@ -55,7 +52,6 @@ namespace Eigen { }; - EIGEN_STRONG_INLINE auto get_sycl_supported_devices()->decltype(cl::sycl::device::get_devices()){ auto devices = cl::sycl::device::get_devices(); std::vector::iterator it =devices.begin(); @@ -78,10 +74,11 @@ struct QueueInterface { bool exception_caught_ = false; mutable std::mutex mutex_; + /// std::map is the container used to make sure that we create only one buffer /// per pointer. The lifespan of the buffer now depends on the lifespan of SyclDevice. /// If a non-read-only pointer is needed to be accessed on the host we should manually deallocate it. - //mutable std::map> buffer_map; + mutable std::map> buffer_map; /// sycl queue mutable cl::sycl::queue m_queue; /// creating device by using cl::sycl::selector or cl::sycl::device both are the same and can be captured through dev_Selector typename @@ -119,42 +116,45 @@ m_queue(cl::sycl::queue(s, [&](cl::sycl::exception_list l) { /// use this pointer as a key in our buffer_map and we make sure that we dedicate only one buffer only for this pointer. /// The device pointer would be deleted by calling deallocate function. EIGEN_STRONG_INLINE void* allocate(size_t num_bytes) const { + auto buf = cl::sycl::buffer(cl::sycl::range<1>(num_bytes)); + auto ptr =buf.get_access().get_pointer(); + buf.set_final_data(nullptr); std::lock_guard lock(mutex_); - return codeplay::legacy::malloc(num_bytes); + buffer_map.insert(std::pair>(static_cast(ptr),buf)); + return static_cast(ptr); } /// This is used to deallocate the device pointer. p is used as a key inside /// the map to find the device buffer and delete it. EIGEN_STRONG_INLINE void deallocate(void *p) const { std::lock_guard lock(mutex_); - return codeplay::legacy::free(p); + auto it = buffer_map.find(static_cast(p)); + if (it != buffer_map.end()) { + buffer_map.erase(it); + } } EIGEN_STRONG_INLINE void deallocate_all() const { std::lock_guard lock(mutex_); - codeplay::legacy::clear(); + buffer_map.clear(); } - EIGEN_STRONG_INLINE codeplay::legacy::PointerMapper& pointerMapper() const { + EIGEN_STRONG_INLINE std::map>::iterator find_buffer(const void* ptr) const { std::lock_guard lock(mutex_); - return codeplay::legacy::getPointerMapper(); - } - - EIGEN_STRONG_INLINE cl::sycl::buffer get_buffer(void* ptr) const { - std::lock_guard lock(mutex_); - return pointerMapper().get_buffer(pointerMapper().get_buffer_id(ptr)); - } - - EIGEN_STRONG_INLINE size_t get_buffer_offset(void* ptr) const { - std::lock_guard lock(mutex_); - return pointerMapper().get_offset(ptr); + auto it1 = buffer_map.find(static_cast(ptr)); + if (it1 != buffer_map.end()){ + return it1; + } + else{ + for(std::map>::iterator it=buffer_map.begin(); it!=buffer_map.end(); ++it){ + auto size = it->second.get_size(); + if((it->first < (static_cast(ptr))) && ((static_cast(ptr)) < (it->first + size)) ) return it; + } + } + std::cerr << "No sycl buffer found. Make sure that you have allocated memory for your buffer by calling malloc-ed function."<< std::endl; + abort(); } - /*EIGEN_STRONG_INLINE void* get_buffer_id(void* ptr) const { - std::lock_guard lock(mutex_); - return static_cast(pointerMapper().get_buffer_id(ptr)); - }*/ - // This function checks if the runtime recorded an error for the // underlying stream device. EIGEN_STRONG_INLINE bool ok() const { @@ -165,7 +165,7 @@ m_queue(cl::sycl::queue(s, [&](cl::sycl::exception_list l) { } // destructor - ~QueueInterface() { codeplay::legacy::clear(); } + ~QueueInterface() { buffer_map.clear(); } }; struct SyclDevice { @@ -183,11 +183,10 @@ struct SyclDevice { } /// Accessing the created sycl device buffer for the device pointer - EIGEN_STRONG_INLINE cl::sycl::buffer get_sycl_buffer(const void * ptr) const { - return m_queue_stream->get_buffer(const_cast(ptr)); + EIGEN_STRONG_INLINE cl::sycl::buffer& get_sycl_buffer(const void * ptr) const { + return m_queue_stream->find_buffer(ptr)->second; } - /// This is used to prepare the number of threads and also the number of threads per block for sycl kernels template EIGEN_STRONG_INLINE void parallel_for_setup(Index n, Index &tileSize, Index &rng, Index &GRange) const { @@ -274,8 +273,6 @@ struct SyclDevice { if (xMode != 0) GRange0 += static_cast(tileSize0 - xMode); } } - - /// allocate device memory EIGEN_STRONG_INLINE void *allocate(size_t num_bytes) const { return m_queue_stream->allocate(num_bytes); @@ -290,15 +287,17 @@ struct SyclDevice { /// the memcpy function template EIGEN_STRONG_INLINE void memcpy(void *dst, const Index *src, size_t n) const { - auto offset= m_queue_stream->get_buffer_offset((void*)src); - auto i= m_queue_stream->get_buffer_offset(dst); + auto it1 = m_queue_stream->find_buffer((void*)src); + auto it2 = m_queue_stream->find_buffer(dst); + auto offset= (static_cast(static_cast(src))) - it1->first; + auto i= (static_cast(dst)) - it2->first; offset/=sizeof(Index); i/=sizeof(Index); size_t rng, GRange, tileSize; parallel_for_setup(n/sizeof(Index), tileSize, rng, GRange); sycl_queue().submit([&](cl::sycl::handler &cgh) { - auto src_acc =get_sycl_accessor(cgh, src); - auto dst_acc =get_sycl_accessor(cgh, dst); + auto src_acc =it1->second.template get_access(cgh); + auto dst_acc =it2->second.template get_access(cgh); typedef decltype(src_acc) read_accessor; typedef decltype(dst_acc) write_accessor; cgh.parallel_for(cl::sycl::nd_range<1>(cl::sycl::range<1>(GRange), cl::sycl::range<1>(tileSize)), MemCopyFunctor(src_acc, dst_acc, rng, i, offset)); @@ -311,11 +310,10 @@ struct SyclDevice { /// on it. Using a discard_write accessor guarantees that we do not bring back the current value of the /// buffer to host. Then we use the memcpy to copy the data to the host accessor. The first time that /// this buffer is accessed, the data will be copied to the device. - template EIGEN_STRONG_INLINE void memcpyHostToDevice(T *dst, const T *src, size_t n) const { + template EIGEN_STRONG_INLINE void memcpyHostToDevice(Index *dst, const Index *src, size_t n) const { auto host_acc= get_sycl_buffer(dst). template get_access(); ::memcpy(host_acc.get_pointer(), src, n); } - /// The memcpyDeviceToHost is used to copy the data from host to device. Here, in order to avoid double copying the data. We create a sycl /// buffer with map_allocator for the destination pointer with a discard_write accessor on it. The lifespan of the buffer is bound to the /// lifespan of the memcpyDeviceToHost function. We create a kernel to copy the data, from the device- only source buffer to the destination @@ -323,14 +321,15 @@ struct SyclDevice { /// would be available on the dst pointer using fast copy technique (map_allocator). In this case we can make sure that we copy the data back /// to the cpu only once per function call. template EIGEN_STRONG_INLINE void memcpyDeviceToHost(void *dst, const Index *src, size_t n) const { - auto offset =m_queue_stream->get_buffer_offset((void *)src); + auto it = m_queue_stream->find_buffer(src); + auto offset =static_cast(static_cast(src))- it->first; offset/=sizeof(Index); size_t rng, GRange, tileSize; parallel_for_setup(n/sizeof(Index), tileSize, rng, GRange); // Assuming that the dst is the start of the destination pointer auto dest_buf = cl::sycl::buffer >(static_cast(dst), cl::sycl::range<1>(n)); sycl_queue().submit([&](cl::sycl::handler &cgh) { - auto src_acc= get_sycl_accessor(cgh, src); + auto src_acc= it->second.template get_access(cgh); auto dst_acc =dest_buf.template get_access(cgh); typedef decltype(src_acc) read_accessor; typedef decltype(dst_acc) write_accessor; @@ -344,8 +343,7 @@ struct SyclDevice { EIGEN_STRONG_INLINE void memset(void *data, int c, size_t n) const { size_t rng, GRange, tileSize; parallel_for_setup(n, tileSize, rng, GRange); - auto buf =get_sycl_buffer(static_cast(static_cast(data))); - sycl_queue().submit(memsetCghFunctor(buf,rng, GRange, tileSize, c )); + sycl_queue().submit(memsetCghFunctor(get_sycl_buffer(static_cast(static_cast(data))),rng, GRange, tileSize, c )); synchronize(); } @@ -411,6 +409,7 @@ struct SyclDevice { }; + } // end namespace Eigen #endif // EIGEN_CXX11_TENSOR_TENSOR_DEVICE_SYCL_H -- cgit v1.2.3