SYCL Explicit Data Movement from Image to Host

[ScottS3]

I'm having trouble trying to figure out how to do explicit data transfer to/from host when using sycl::image. I'm familiar with sycl::handler.copy when using buffers:

float* host_ptr;
sycl::buffer<float, 1> buf;
sycl::queue& q
q.submit([&](sycl::handler& cgh) {
  auto acc = buf.get_access<sycl::access::mode::read>(cgh);
  cgh.copy(acc, host_ptr);
});

q.wait_and_throw();

but I can't quite find the equivalent code when dealing with sycl::image. If I pass an image accessor to .copy it complains that the accessor doesn't have a subscipt operator. Does anyone know the correct method for this?

My previous OpenCL code had clEnqueueReadImage and clEnqueueWriteImage for this sort of data image transfer. Is there a sycl equivalent?

note: I'm aware of the implicit copy back to host on sycl::image deconstruction, but I want to know the explicit solution, rather than the implicit one.

[VaishnaviV_Intel]

Hi,

In my example there are three sycl::images. I called two of them "textures" and one "theImage". These are just names. The first two are uploaded exactly once, but used by every operation in the loop. The "theImage” is uploaded each time.

Thus, as I understand it, this demonstrates what you are asking after: how to use the sycl::image class but such that is NOT constantly re-uploaded to the device. The "textures" in my sample demonstrate this. They are used repeatedly in every iteration of the loop, but uploaded only once.

You said that you also "... and we download them if/when we require that info host side". For that you can use a host accessor. Or just wait for the image to go out of scope (if possible).

The queue/handler.copy() operations works on USM pointers or accessors. It is not called directly on a buffer or image object. But it should work with an accessor gotten from an image object. So it is available.

Thanks & Regards,

Vankudothu Vaishnavi. 

[ScottS3]

Hi,

"They are used repeatedly in every iteration of the loop, but uploaded only once."

Yes I understand that, but in that example the 'upload' can only happen with the creation of an image, or 'download' with the destruction of an image.  I'm trying to figure out how to upload/download without reallocating/unallocating an image i.e. uploading to an existing device image without recreation, or downloading without destroying a device image.

We usually create/allocate an image once and then potentially upload to it and download from it without recreating/destroying it constantly like in the example given.  Hence the inquiry on how to manually manage and transfer device image data.

"But it should work with an accessor gotten from an image object. So it is available."

I would love to see a working example of copy with images.  I've tried a few times to get copy working with images but I've never managed to get it compiling or working.

Kind Regards
Scott

[VaishnaviV_Intel]

Hi,

>>I'm trying to figure out how to upload/download without reallocating/unallocating an image i.e. uploading to an existing device image without recreation, or downloading without destroying a device image.

Accessors allow you to both read and write the image data without destroying the image. And accessor can be used on the device in kernel code (we call those device accessors) or outside a kernel in the host code ( host accessors). When your code uses an accessor SYCL will automatically figure out whether the data needs to be uploaded or downloaded for you. So if you create an image and then use an accessor to read it on the device, SYCL will figure out that that data needs to be uploaded and enqueue the upload appropriately. If, after that, a different kernel were to be submitted to the device that also creates an accessor and accesses the image, SYCL would figure out that it does NOT need to copy the image data a second time. But, if you put a host accessor that accesses the data between those two kernels, SYCL will realize that the data will have to be downloaded (so the host accessor has the right data) and then uploaded again after (for the second kernel to have the right data). All of this happens automatically. And it takes read vs write into consideration ( if the first kernel only reads, and never writes, then there is no need to download the data ).

Hope you got it. So by using accessors and kernels you can use a single image. That might be something like this:

sycl::image myImage;
for (i =0 : N ){
    // 1. submit kernel with device accessor
   //  2. make host accessor and do operation on host (copy data somewhere)
}

The data will be automatically uploaded/downloaded as needed before steps 1 and 2. And, again, if everything is just reading, then there would be only one image upload , not one per loop iteration.

>> would love to see a working example of copy with images. I've tried a few times to get copy working with images, but I've never managed to get it compiling or working.

You don't need to worry too much about how long the image stays where. So, Please try to write your code in a way that suits your use case.

If you still face any issues, please let us know.

Thanks & Regards,

Vankudothu Vaishnavi.

[ScottS3]

Hi,

"When your code uses an accessor SYCL will automatically figure out whether the data needs to be uploaded or downloaded for you."

Ah so accessors are what SYCL uses to organise data transfer for you.  Not knowing exactly when image data gets uploaded or downloaded explicitly will probably block usage of SYCL for us, which is a shame.

"But, if you put a host accessor that accesses the data between those two kernels"

I'm assuming this would be the functions we wrote earlier that used host accessors.  Still a little unsure how SYCL would know to upload/download data when it doesn't have host memory associated with it.  I guess we would use the host accessor, but they seemed very slow compared to manual upload/download kernels we wrote earlier.

"would love to see a working example of copy with images"

I'm still curious to see if there is a working example of this?  I've never been able to get it compiling?  Do you have an example?

Regards

Scott

[VaishnaviV_Intel]

Hi,

Thank you for your patience. The issue raised by you will be resolved with new image APIs and that has been targeted to be fixed in the Intel oneAPI 2024.2 version. Please download it and let us know if this resolves your issue. We will be closing this thread from our side, If the issue still persists with the new release then create a new thread for us to investigate.

Please refer to the following release notes for the updates,

https://www.intel.com/content/www/us/en/developer/articles/release-notes/intel-oneapi-dpc-c-compiler-release-notes.html

Thanks & Regards,

Vankudothu Vaishnavi.