Hacker News Reader

Comment by stochastic_monk | Hacker News Reader

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++?

View on HN

I've written a few Reddit posts on the subject:

https://www.reddit.com/r/ROCm/comments/adyu5k/amd_rocm_hcc_p...

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.

Hasn't this worked via HSAIL for years?

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.

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.

Further reading: https://gcc.gnu.org/wiki/cauldron2018?action=AttachFile&do=g...

(e.g. slide 22)

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?

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.

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. :)

From a memory perspective you can view it as multiple threads. Caches are a bit more complicated (each "cpu core"/streaming processor has it's own L1 cache, and cache coherency between L1 caches isn't nessesarily guaranteed). Memory bandwidth and latency are still bottlenecks.

The bigger implication of the SIMD design is that for any conditional all your "threads" should take the same route, otherwise you may end up computing both paths for every thread.

> 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.

This is not quite doing it justice. GPU threads allow for branching (which thus includes things like early return). They do loose performance because inactive branches are sleeping (although I didn't test this with Volta and onward, which is supposedly brings some improvements), but that's way better than just breaking vectorization outright. This makes it much easier to port some naive (or just cache-optimized) piece of numerics code to GPU rather than trying to vectorize it on CPU.

How's that different from masking (i.e per-line predication)?

the difference is that for masking you have to rely on the compiler to do-the-right-thing[TM] OR you have to program it yourself manually (which is quite annoying). With GPU threads the scheduling of branches is hardware supported, so it's much easier to map your scalar code to it as long as it's sufficiently data parallel.

That's the thing though: for AMD GCN, the compilers do have to do the right thing.

You can look at the actual hardware ISA, it's all documented e.g. here: https://rocm-documentation.readthedocs.io/en/latest/GCN_ISA_..., and both the HCC compiler and Mesa's OpenGL and Vulkan implementations are open-source and based on LLVM, where the backend is even upstream, so you can take a look for yourself.

Essentially everything those compilers do would be possible on a CPU with masked SIMD instructions (e.g. AVX512). In fact, that's basically how ISPC works. The one exception is that GCN has much more powerful scatter/gather support than CPUs do (and graphics-focused stuff like texture sampling hardware). But as far as the execution model is concerned, it's really the same.

I see, that's interesting. So essentially, AMD's hardware schedulers are not concerned with branching the same way Nvidia seems to be doing it?

Pascal does branching in a similar manner.

Volta / Turing does branching per-PC now, but it is opt-in optional. In any case, I expect Volta / Turing to prefer the old SIMD model of compute.

The only reason Volta / Turing supports per-PC branching is to handle some multi-thread communication issues. The execution units are still SIMD, so you take a heavy penalty going to the per-PC model.

PS: I haven't actually programmed Volta/Turing, its just my understanding from the white paper.

> How's that different from masking (i.e per-line predication)?

AMD GCN has hardware accelerated masking + stacks for this purpose. AMD's GPU is basically a combination of masking + fork/join commands, to split up wavefronts.

But its really just hardware accelerated masks. You sometimes need to store the masks + program counters in a stack to allow for divergent threads, and then pop the masks off and execute older commands... (Complex if / then / else branches where every thread takes a different set of branches).

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++.

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.

CUDA supports most of C++14 and earlier: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index....

With few CUDA specific extensions.

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?

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

> 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?

Reply on news.ycombinator.com