Remove ComputeCpp-specific code from SYCL Vptr
diff --git a/Eigen/src/Core/arch/SYCL/SyclMemoryModel.h b/Eigen/src/Core/arch/SYCL/SyclMemoryModel.h
index 2b96587..54eedfa 100644
--- a/Eigen/src/Core/arch/SYCL/SyclMemoryModel.h
+++ b/Eigen/src/Core/arch/SYCL/SyclMemoryModel.h
@@ -141,7 +141,7 @@
/* basic type for all buffers
*/
- using buffer_t = cl::sycl::buffer_mem;
+ using buffer_t = cl::sycl::buffer<buffer_data_type_t>;
/**
* Node that stores information about a device allocation.
@@ -237,17 +237,14 @@
template <typename buffer_data_type = buffer_data_type_t>
cl::sycl::buffer<buffer_data_type, 1> get_buffer(
const virtual_pointer_t ptr) {
- using sycl_buffer_t = cl::sycl::buffer<buffer_data_type, 1>;
- // get_node() returns a `buffer_mem`, so we need to cast it to a `buffer<>`.
- // We can do this without the `buffer_mem` being a pointer, as we
- // only declare member variables in the base class (`buffer_mem`) and not in
- // the child class (`buffer<>).
auto node = get_node(ptr);
+ auto& map_node = node->second;
eigen_assert(node->first == ptr || node->first < ptr);
- eigen_assert(ptr < static_cast<virtual_pointer_t>(node->second.m_size +
+ eigen_assert(ptr < static_cast<virtual_pointer_t>(map_node.m_size +
node->first));
- return *(static_cast<sycl_buffer_t *>(&node->second.m_buffer));
+ return map_node.m_buffer.reinterpret<buffer_data_type>(
+ cl::sycl::range<1>{map_node.m_size / sizeof(buffer_data_type)});
}
/**
@@ -429,8 +426,11 @@
template <class BufferT>
virtual_pointer_t add_pointer_impl(BufferT b) {
virtual_pointer_t retVal = nullptr;
- size_t bufSize = b.get_count();
- pMapNode_t p{b, bufSize, false};
+ size_t bufSize = b.get_count() * sizeof(buffer_data_type_t);
+ auto byte_buffer =
+ b.template reinterpret<buffer_data_type_t>(cl::sycl::range<1>{bufSize});
+ pMapNode_t p{byte_buffer, bufSize, false};
+
// If this is the first pointer:
if (m_pointerMap.empty()) {
virtual_pointer_t initialVal{m_baseAddress};
