Look ma, no CUDA! Programming GPUs with modern C++ and SYCL · Federico Ficarelli

[utterances-bot]

Look ma, no CUDA! Programming GPUs with modern C++ and SYCL · Federico Ficarelli

Volatile hacks.

https://nazavode.github.io/blog/sycl/

[timocafe]

It look's nice. But the million question is the memory layout.

• a CPU can accept more and less everything but a preference for SoA
• a GPU is SoA
• a FPGA ????

So yes multiple backend is good, but which memory layout ? SoA look like the obvious choice, but I have no idea of what need a FPGA

[nazavode]

@timocafe yep, your point is totally right. As all the SYCL folks I'm in touch with pointed out, they clearly avoided promising any kind of magic performance portability. It's an honest approach imho, we all know how hard it is across traditional CPUs, let alone completely different architectures. I'm not into FPGAs apart some very introductory training and they are really different beasts, so it should be no surprise that a radically different approach would be needed. As far I understood the underlying intent, the promise of SYCL is just to provide a common ecosystem in which writing different versions of the same kernel optimized for different execution backends should be as cheap and convenient as possible (no additional languages/tools/crappy compilers/so on).

[nazavode]

@timocafe a note about the memory layout, the abstraction provided here is quite low level, somewhere between std::span and the proposed (now postponed to C++23) std::mdspan: while the latter could in theory be extended to provide a uniform interface over different data layouts (e.g: aos/soa as you mentioned), sycl::buffer is just a glorified view over a stride-1 sequence of objects. In the (not so remote) case an FPGA kernel wants some weird layout, we should take care of it by hand as far I understand the API.

[cagnulein]

Ciao Federico, i've started playing with LLVM and SYCL just this month, and i've immediately saw a strange behaviour. Could you please have a look here intel/llvm#813 and give me your opinion?
Ciao e grazie!

[rcrovella]

regarding the section beginning with:

Data transfers are implicit

Unlike CUDA and OpenCL where explicit copies are required by the model,...

You may wish to read about CUDA unifed memory, which has been available in the CUDA programming model since CUDA 6.0 /2014/Kepler.