Hello friends,
I was working with some SYCL code and I tried the memory synchronization by using an host accessor on the host side, but it did not work. I don’t know why, because if I add a scope around the buffer creation and kernel function everything is fine.
This is the code:
cl::sycl::gpu_selector device;
cl::sycl::queue Q{device};
std::cout << "Running on "
<< Q.get_device().get_info<sycl::info::device::name>()
<< "\n";
constexpr int N = 10;
std::vector<float>vector(N);
for(int i=0; i<N; i++){
vector[i] = 100.0;
}
cl::sycl::buffer<float,1> buffVec(vector.data(),vector.size());
//WE DO NOT USE SCOPE
Q.submit([&](cl::sycl::handler &cgh){
//Create accesors to data
auto acc = buffVec.get_access<cl::sycl::access::mode::read_write>(cgh);
cgh.parallel_for<vec_mult>(cl::sycl::range<1>(N), [=](cl::sycl::id<1> index){
acc[index] *=2.0;
});
});
buffVec.get_access<cl::sycl::access::mode::read>(); //<--- Host accessor to Synchronize memory
for(int i=0; i<N; i++){
std::cout<<vector[i]<<std::endl;
}
Thanks for your support
Hi @nullpointer,
So the reason this isn’t doing what you expect is that by default host accessor don’t synchronize the data to the original host pointer as the buffers do on destruction. They simply make the data available on the host via the host accessor object, and this may be done via some internal allocation to the SYCL runtime.
There is a property which can be passed to a buffer object on construction called property::buffer::use_host_ptr which will prohibit the SYCL runtime from allocating memory and instead use the original host pointer. Using this property would give you the behaviour you are looking for, as it requires the accessor subscript operator to return the same address as the original host pointer.
However, it’s important to be careful of the scope of host accessors, especially when using this approach to access the data via the original host pointer. The data is only made available on the host and subsequently is only accessible via the original host pointer if using the property::buffer::use_host_ptr property, for the lifetime of the host accessor. This means once the object is destroyed the runtime is free to move the data asynchronously and therefore accessing the original pointer is undefined behaviour. So it’s best practice to keep the host accessor object alive for the duration of the access to the data, before releasing it back to the SYCL runtime.
I hope this helps, please let me know if you have any questions.
Gordon
1 Like
Hello Gordon,
Thanks for your answer, I really appreciate it. Now I know I was understanding the use of the host pointer accessor in a wrong way ( I saw this approach in a web page: Coding Games and Programming Challenges to Code Better). So, they say that the use of scopes “{}” will destruct the buffer and synchronize the memory… but what do you recommend?
Sorry for this inconvenience I’m new with SYCL
-Juan
Hello Juan,
Not at all, we’re very happy to help.
Thank you for sharing this link, I suspect this advice regarding synchronization with the host via host accessors is relying on an implementation detail of DPC++, so while that will work with DPC++, it may not be portable to other SYCL implementations. The advice regarding synchronization with the host via the buffer destructor is correct.
Generally we recommend that if you wish to read or modify the state of the data on the host between kernels to use host accessors, but use the host accessor object itself to read or modify elements of the data via the subscript operator, and if you wish to have the data copied back to the original pointer after the kernels are complete to use scope to destroy the buffer before accessing the original pointer directly.
Using the property::buffer::use_host_ptr as described above is also another alternative to using host accessors, though this approach does have implications for achieving pinned memory, for more details on this see Codeplay Developer - ComputeCpp CE - Guides - Memory Model
I hope this helps, please let me know if you have any questions.
Gordon