AMD GCN Radeon Support in GCC 9

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.

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

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.

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.

> C++ is explicitly unsupported

Given that C++ now has a defined memory model, it might very well be that it's now impossible to implement C++ in a standard conforming way on exotic architectures.

This appears to be a GCN backend intended for use with OpenMP and friends, so not really comparable to cuda (more comparable maybe to the llvm nvptx backend).

Why CUDA is faster than OpenCL?

OpenMP 4.5 provides the #pragma target data, target enter data, target exit data, and target update constructs to provide for the memory-management to-and-from the GPU.

omp_target_alloc, omp_target_memcpy, omp_target-associate_ptr provide some other utility functions that you'd expect for these use cases. omp_target_memcpy can copy from host-to-device, device-to-host, or even device-to-device.

From the perspective of high performance computing (e.g. achieve a significant fraction of what the hardware is capable of for a given problem) I've found GPU programming always more accessible since the introduction of CUDA - assuming that the problem is something for which GPU makes sense in the first place (e.g. throughput is dominating over latency).

That is to say, I'd much rather optimize GPU kernels for things like sequential access, register use and programmable caches than make it vectorize on the various versions of AVX, treat multi-core parallelization separately and then fit everything into the fastest CPU cache possible just because the CPU's memory bandwidth is so damn slow.

GPU Programming is easier than AVX Programming on the CPU. That's why CUDA has become more popular.