search

nazavode / nazavode.github.io

Sources for GitHub Pages user site, generated with gohugo.io
https://nazavode.github.io
1 stars 1 forks source link

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

Open utterances-bot opened 4 years ago

utterances-bot commented 4 years ago

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

Volatile hacks.

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

timocafe commented 4 years ago

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

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 commented 4 years ago

@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 commented 4 years ago

@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 commented 4 years ago

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 https://github.com/intel/llvm/issues/813 and give me your opinion? Ciao e grazie!

rcrovella commented 3 years ago

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.