Kokkos: OpenMPTarget: Why a separate execution space? (Question)

[mhoemmen]

Hi all! I had a chat yesterday with @crtrott about why OpenMPTarget needs to be a compile-time option and can't just be a run-time option. (@crtrott explained that the "target" processor may have a different instruction set than the "host" processor, so the compiler may need to generate different code.) I get that, but I'm curious why it's actually a different execution space, as opposed to a policy argument. Is the idea that each execution space corresponds to one code target? How would that correspond to use of target with the device(which) clause? (I'm guessing that most use of device(which) requires homogeneous devices, so that the which argument doesn't affect code generation.)

[crtrott]

You still have heterogenous systems where the more "CPU" like resources are used with non-target OpenMP while the more "accelerator" like resources are utilized via OpenMP-Target regions. This is already the case on IBM+NVIDIA GPUs, where the parallel host execution space is OpenMP and the device execution space is OpenMP with target directives. (Obviously the preferred way right now on IBM+NVIDIA is still OpenMP + CUDA, not OpenMP + OpenMPTarget).

[mhoemmen]

@crtrott Right, I get that bit -- it just seems like the Kokkos back-end should be the same for both options, so you could have just one execution space implementation. Is the reason that OpenMPTarget needs to store some information on the GPU using those funny pragmas that stash allocations there for later use? Anyway, I just mainly wanted to record the question.

[crtrott]

It actually is two different code paths, basically in its most simple form:

#pragma omp parallel for
for(int i=0; i<N; i++)

vs

#pragma omp target team distribute parallel for [team_size, ...]
for(int i=0; i<N; i++)

And that means it is two different compilation and two different execution spaces.

[ibaned]

It looks like this was answered. If not, please reopen.