1. Scope

Ten names: std::par, Kokkos, Level Zero, OpenMP offload, OpenCL, HIP, HSA, SYCL, CUDA Runtime, CUDA Driver, are frequently discussed as if they were alternatives to one another. They are not. They occupy five distinct levels of a single tower, and several of them are implemented in terms of others.

2. The Five Levels

L5: std::par is shorthand for the ISO C++17 execution policies std::execution::par and par_unseq applied to the standard algorithms (for_each, transform, reduce, and the C++20 ranges forms). There is no device concept in the language, no explicit data movement, and no kernel. The compiler infers everything. Cleanest code available; least controllable.

L4: Portability layers. Two mechanisms, one purpose: write once, run on any vendor. Kokkos, RAJA, and Alpaka are C++ template libraries that resolve to a backend at compile time; Kokkos additionally supplies multidimensional Views with selectable memory layout, which is the feature that actually delivers performance portability and which no lower layer provides. OpenMP target offload does the same job with #pragma omp target directives lowered by the compiler into libomptarget calls. Directives suit large Fortran and C codebases; templates suit modern C++.

OpenACC and OpenMP offload deserve separating. OpenACC came first (2011, from Cray, PGI, NVIDIA, and CAPS) and predates OpenMP 4.0’s target (2013); it was in part a proving ground whose ideas were later absorbed. Its distinguishing feature is the descriptive kernels construct, which hands a loop nest to the compiler and lets the compiler perform dependence analysis and decide the mapping, as opposed to the prescriptive parallel loop (and all of OpenMP target), where the programmer dictates the decomposition. OpenACC’s parallelism levels are gang, worker, and vector, corresponding roughly to OpenMP’s teams, threads, and SIMD, and on NVIDIA hardware to block, warp, and lane. Its data model – copyin, copyout, create, present, unstructured enter data/exit data, and update —- is close enough to OpenMP’s map clauses that mechanical translation between the two is feasible, which matters below. Asynchronous queues (async(n), wait) map onto CUDA streams.

The convergence is real but incomplete: OpenMP 5.0 added a descriptive loop construct explicitly borrowing from OpenACC’s kernels, and a long-discussed merger of the two specifications has never happened. The practical difference today is implementation breadth, not expressiveness. For many people, OpenACC is dieing and tend to use OpenMP.

L3: Programming models. What people actually write. The CUDA Runtime API (cudaMalloc, «<»>) is NVIDIA’s convenience layer. HIP is AMD’s deliberate near-clone of it, close enough that migrates code by textual cuda*->hip* substitution. OpenCL is the veteran cross-vendor option: C host API, separate-source kernels, portable to CPUs, GPUs, and FPGAs, but verbose and second-class on NVIDIA. SYCL is Khronos’s modern answer, single-source C++ with lambdas, defined over a pluggable backend model.

L2: Driver APIs. The CUDA Driver API (libcuda.so, cu*), AMD’s HSA/ROCr (libhsa-runtime64.so, hsa_*), and Intel’s Level Zero (libze_loader.so, ze*) are exact analogues. Each exposes explicit contexts, devices, command queues, memory allocation, and module loading. None cares what language produced the kernel; each consumes a binary format (cubin/PTX, AMD code objects, SPIR-V). These are targets for compilers and runtime libraries, not usually written by hand.

L1: Kernel-mode drivers. nvidia.ko, amdgpu + KFD, i915/xe.

3. Layer Skeleton

            ╔═══════════════════════════════════════════════════════════════╗
 L5         ║                std::par   (ISO C++ algorithms)                ║
 language   ╚═══════════════════════════════════════════════════════════════╝

            ╔═════════════════════════╗ ╔═══════════════════════════════════╗
 L4         ║  Kokkos / RAJA / Alpaka ║ ║  OpenACC   │   OpenMP target      ║
 portability║  (C++ templates)        ║ ║  (directives, partly descriptive) ║
            ╚═════════════════════════╝ ╚═══════════════════════════════════╝

            ┌───────────────┬────────────────┬───────────────┬──────────────┐
 L3         │ CUDA Runtime  │      HIP       │     SYCL      │    OpenCL    │
 model      │  (cudaMalloc) │   (hipMalloc)  │  DPC++ / ACpp │              │
            └───────────────┴────────────────┴───────────────┴──────────────┘

            ┌───────────────┬────────────────┬───────────────┐
 L2         │  CUDA Driver  │   HSA / ROCr   │  Level Zero   │
 driver API │   libcuda     │ libhsa-runtime │  libze_loader │
            └───────────────┴────────────────┴───────────────┘

            ┌───────────────┬────────────────┬───────────────┐
 L1         │   nvidia.ko   │  amdgpu + KFD  │   i915 / xe   │
 kernel     └───────────────┴────────────────┴───────────────┘
                 NVIDIA           AMD              Intel

OpenCL sits at L3 in no single vendor column: its ICD loader binds to any of the three at run time. OpenACC sits at L4 but, like OpenMP target, its edges descend straight to L2, it has no L3 layer beneath it.

4. Vendor Mapping

Tier NVIDIA AMD Intel Cross-vendor
L5 language — — — std::par
L4 portability — — — Kokkos, RAJA, OpenACC, OpenMP offload
L3 model CUDA Runtime HIP (uses SYCL) OpenCL, SYCL
L2 driver API CUDA Driver HSA / ROCr Level Zero —
L1 kernel driver nvidia.ko amdgpu + KFD i915 / xe —

OpenACC is nominally cross-vendor but is the least evenly supported entry in that column.

5. Kokkos

Kokkos picks at compile time via templates rather than at run time via dlopen:

Kokkos ─┬─► CUDA Runtime API   (Kokkos::Cuda)
        ├─► HIP                (Kokkos::HIP)
        ├─► SYCL               (Kokkos::SYCL) ─► UR ─► Level Zero/…
        ├─► OpenMP target      (Kokkos::Experimental::OpenMPTarget)
        ├─► OpenACC            (Kokkos::Experimental::OpenACC, added 4.0)
        └─► OpenMP/threads/HPX (host)