4 ms·

I’ve never written any code for the GPU. Does anyone with more experience have an idea how similar writing GCN code for gcc is to standard CPU C/C++?

speps 8y ago
You'd usually write HLSL to target DXBC which is then compiled to GCN by the AMD driver. That's how it usually works in video games targeting DirectX 11 on PC. So to answer your question, you wouldn't really use C/C++ to target GCN. I think this is more akin to how CUDA works by using a dialect on top of C.
- llukas 8y ago
  CUDA supports most of C++14 and earlier: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#c-cplusplus-language-support https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.... With few CUDA specific extensions.
  - stochastic_monk 8y ago
    They even have special libraries, like thrust, for writing generic GPU code. Is it safe to say that this development for gcc doesn’t provide something comparable for GCN?
  - pjmlp 8y ago
    In fact newer CUDA generations have been designed to run C++ code. "CppCon 2017: Olivier Giroux "Designing (New) C++ Hardware”" https://www.youtube.com/watch?v=86seb-iZCnI https://www.youtube.com/watch?v=86seb-iZCnI
- shmerl 8y ago
  > You'd usually write HLSL to target DXBC which is then compiled to GCN by the AMD driver. Why not target SPIR-V for Vulkan driver?
Athas 8y ago
I've done a good bit of CUDA and OpenCL. I haven't dug into this GCC implementation, but my guess is that you're banned from using any features that require sophisticated runtime support (like exceptions or RTTI). Things like templates, which are a compile-time mechanism, should work fine. Also, you still have the usual nasty issues of the GPU having a separate memory space from the CPU. This is just a code generation backend; not some transparent GPU system model for C/C++.
atq2119 8y ago
So the funny thing is that gcc's backend for GCN is actually much closer to writing for standard CPUs than any other way of programming GPUs, because you directly program it at the wave-level. You see, the (not so) dirty secret of GCN is that from a very real perspective, the best way to think about it is that it's a bunch of CPU cores with extremely wide masked SIMD units. Almost all GPU programming languages obscure that fact, and their compilers do magic behind the scenes to make it appear as if you were programming individual threads without SIMD. GCC is the odd one out, which is really quite fascinating.
- mappu 8y ago
  Further reading: https://gcc.gnu.org/wiki/cauldron2018?action=AttachFile&do=get&target=amd-gcn-cauldron-presentation-2018.pdf https://gcc.gnu.org/wiki/cauldron2018?action=AttachFile&do=g... (e.g. slide 22)
- elcritch 8y ago
  Do the SIMD units have different memory blocks that while being controlled by one cpu core would make the super wide SIMD behave more like multiple threads with different memory units? Or at least act like it from a latency perspective?
  - atq2119 8y ago
    I guess you're talking about scatter / gather? All vector memory instructions in the GCN ISA are scatter / gather-type instructions, i.e. you provide one pointer (or buffer index, or texture coordinates) per SIMD lane. Docs say that the performance depends on the actual distribution of pointers, i.e. for best performance you should ensure that the pointers are consecutive (use SoA instead of AoS layouts, etc.), but I imagine that the drop in performance is somewhat gradual and graceful as your pointers become more scattered. I don't think the details are really documented though, and I wouldn't be surprised if those details changed between GCN generations.
    - justinclift 8y ago
      Sounds like it'd be a useful benchmark to do + publish for people's reference. May also lead to some tuning choices upstream too, depending on the results. :)
dragontamer 8y ago
I've written a few Reddit posts on the subject: https://www.reddit.com/r/ROCm/comments/adyu5k/amd_rocm_hcc_programming_introduction/ https://www.reddit.com/r/ROCm/comments/adyu5k/amd_rocm_hcc_p... https://www.reddit.com/r/gpgpu/comments/adyxpi/musings_on_vega_gcn_architecture/ https://www.reddit.com/r/gpgpu/comments/adyxpi/musings_on_ve... There's a 3rd story here: GCC's style is OpenMP 4.5. The big news is that the following kind of code is now possible: #pragma omp target map(a,b,c,d) { #pragma parallel for for(int i=0; i<1000000; i++){ a[i] = b[i] * c + d ; } } The above code will run on the GPU. GCC, CLang, and other languages are trying to support OpenMP 4.5's ability to offload code to GPUs. That's what this GCC9 news is about: AMD GPUs are gaining the ability to execute the above code, integrated into OpenMP Tasks and all that good stuff.
- twtw 8y ago
  Hasn't this worked via HSAIL for years?
  - vetrom 8y ago
    HSAIL depends on an ecosystem of non-free software components. Granted, it's an uphill battle to get the concert of moving pieces that would be needed for 3d/compute acceleration cards to work on a fully open source stack. That said, getting GCC to do a better job implementing code for those is a big step in that direction.

speps 8y ago

You'd usually write HLSL to target DXBC which is then compiled to GCN by the AMD driver. That's how it usually works in video games targeting DirectX 11 on PC. So to answer your question, you wouldn't really use C/C++ to target GCN. I think this is more akin to how CUDA works by using a dialect on top of C.

llukas 8y ago
CUDA supports most of C++14 and earlier: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#c-cplusplus-language-support https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.... With few CUDA specific extensions.
- stochastic_monk 8y ago
  They even have special libraries, like thrust, for writing generic GPU code. Is it safe to say that this development for gcc doesn’t provide something comparable for GCN?
- pjmlp 8y ago
  In fact newer CUDA generations have been designed to run C++ code. "CppCon 2017: Olivier Giroux "Designing (New) C++ Hardware”" https://www.youtube.com/watch?v=86seb-iZCnI https://www.youtube.com/watch?v=86seb-iZCnI
shmerl 8y ago
> You'd usually write HLSL to target DXBC which is then compiled to GCN by the AMD driver. Why not target SPIR-V for Vulkan driver?

Athas 8y ago

I've done a good bit of CUDA and OpenCL. I haven't dug into this GCC implementation, but my guess is that you're banned from using any features that require sophisticated runtime support (like exceptions or RTTI). Things like templates, which are a compile-time mechanism, should work fine. Also, you still have the usual nasty issues of the GPU having a separate memory space from the CPU. This is just a code generation backend; not some transparent GPU system model for C/C++.

atq2119 8y ago

So the funny thing is that gcc's backend for GCN is actually much closer to writing for standard CPUs than any other way of programming GPUs, because you directly program it at the wave-level. You see, the (not so) dirty secret of GCN is that from a very real perspective, the best way to think about it is that it's a bunch of CPU cores with extremely wide masked SIMD units. Almost all GPU programming languages obscure that fact, and their compilers do magic behind the scenes to make it appear as if you were programming individual threads without SIMD. GCC is the odd one out, which is really quite fascinating.

mappu 8y ago
Further reading: https://gcc.gnu.org/wiki/cauldron2018?action=AttachFile&do=get&target=amd-gcn-cauldron-presentation-2018.pdf https://gcc.gnu.org/wiki/cauldron2018?action=AttachFile&do=g... (e.g. slide 22)
elcritch 8y ago
Do the SIMD units have different memory blocks that while being controlled by one cpu core would make the super wide SIMD behave more like multiple threads with different memory units? Or at least act like it from a latency perspective?
- atq2119 8y ago
  I guess you're talking about scatter / gather? All vector memory instructions in the GCN ISA are scatter / gather-type instructions, i.e. you provide one pointer (or buffer index, or texture coordinates) per SIMD lane. Docs say that the performance depends on the actual distribution of pointers, i.e. for best performance you should ensure that the pointers are consecutive (use SoA instead of AoS layouts, etc.), but I imagine that the drop in performance is somewhat gradual and graceful as your pointers become more scattered. I don't think the details are really documented though, and I wouldn't be surprised if those details changed between GCN generations.
  - justinclift 8y ago
    Sounds like it'd be a useful benchmark to do + publish for people's reference. May also lead to some tuning choices upstream too, depending on the results. :)

dragontamer 8y ago

I've written a few Reddit posts on the subject: https://www.reddit.com/r/ROCm/comments/adyu5k/amd_rocm_hcc_programming_introduction/ https://www.reddit.com/r/ROCm/comments/adyu5k/amd_rocm_hcc_p... https://www.reddit.com/r/gpgpu/comments/adyxpi/musings_on_vega_gcn_architecture/ https://www.reddit.com/r/gpgpu/comments/adyxpi/musings_on_ve... There's a 3rd story here: GCC's style is OpenMP 4.5. The big news is that the following kind of code is now possible: #pragma omp target map(a,b,c,d) { #pragma parallel for for(int i=0; i<1000000; i++){ a[i] = b[i] * c + d ; } } The above code will run on the GPU. GCC, CLang, and other languages are trying to support OpenMP 4.5's ability to offload code to GPUs. That's what this GCC9 news is about: AMD GPUs are gaining the ability to execute the above code, integrated into OpenMP Tasks and all that good stuff.

twtw 8y ago
Hasn't this worked via HSAIL for years?
- vetrom 8y ago
  HSAIL depends on an ecosystem of non-free software components. Granted, it's an uphill battle to get the concert of moving pieces that would be needed for 3d/compute acceleration cards to work on a fully open source stack. That said, getting GCC to do a better job implementing code for those is a big step in that direction.