OpenCL USM/SVM requirement when using only buffers

[pjaaskel]

Hi!

Do I read it right, that USM/SVM support is required from the OpenCL device even if the application would not use USM allocations, but only buffer-based accesses?

BR,
Pekka

[illuhad]

Hi Pekka, this is correct.

The background is the following: SYCL has a massive problem in that there are zero interop guarantees in the spec between buffers and USM pointers. This means that e.g. a SYCL library developer has to effectively implement their library twice: Once with a buffer interface and once with a USM interface because they cannot know which model the client code will use, and buffers and USM pointers cannot talk to each other. Obviously this can massively drive up the cost of implementing SYCL libraries and I believe it is one of the major issues holding back the SYCL ecosystem.
Because of this, we have decided that buffers and USM pointers should be universally interoperable as an extension in our implementation. This is achieved by implementing buffers on top of USM device allocations, since this is the only way we can reliably turn a buffer into a USM pointer and vice versa.

If this is an issue for pocl, it might be possible to reinterpret USM pointers as index into an array of cl_mem or similar when used in buffers at the expense of the extension, but I'd prefer if we could avoid such hacks of course :-)

[illuhad]

I think the problem you describe is not a SYCL problem as such, but just that the world is not ready yet for apps to use raw unified SVM/USM pointers everywhere (preferably system-shared or FG system SVM!). Having them supported widely would make the library writer's issues go away. I've been hoping this dream to come reality since HSAF started to push the concept, but seems I have to wait a little bit longer :D Meanwhile I think any portable heterogeneous code base shall use buffers if it's not strictly necessary to pass around raw pointers.

I'm not sure. SYCL knows different kinds of USM pointers - I am talking only about device USM allocations, not shared USM allocations. USM device allocations are basically like cudaMalloc. You get a pointer that can exclusively be dereferenced only on the target device. To transfer data, explicit queue::memcpy() calls are needed. I think that such pointers are in fact widely supported at least across the common (Intel/NVIDIA/AMD) hardware. The only hardware that might have a problem here is FPGAs, but those are kind of niche anyway.

[pjaaskel]

Yep, device allocations are of course easier, but still not trivial due to the need to have the unified address space (the range has to be preserved from the host VM). In comparison to cl_mem, that is.

[illuhad]

That's true. But as I said, basically all common hardware supports this. And this is also one of the important building blocks to enable more complex data structures.

Almost every CUDA app that gets migrated to SYCL uses USM device allocations, because porting cudaMalloc to sycl::malloc_device is very straight-forward compared to migrating to buffer which can require substantial refactoring. And the CUDA-to-SYCL migration tools also use USM. Because of this, and because of the additional control that USM device allocations allow over the buffer model, pretty much every SYCL HPC app uses USM. So, I see USM support as crucial in the SYCL ecosystem. And none of all these USM apps can talk to code that uses SYCL buffers.

I understand that USM can be difficult to support for something like PoCL-R, but if there's a way to support it, it might be worth the effort. The future of SYCL without USM is uncertain as more apps move to USM.

[pjaaskel]

Yep, I agree. Minimal (CG SVM or device-USM) is top on my priority list to add to PoCL-R. Should be doable.

[illuhad]

Sounds good. Let me know if you need anything from our side to enable PoCL-R with our SYCL implementation :)