OpenMP Target Offloading: GPU Acceleration via Pragmas — extending OpenMP directive-based parallelism to GPUs

OpenMP target offloading extends CPU-focused OpenMP directives to GPUs via pragmas specifying kernels and data movement, enabling GPU acceleration without rewriting code.

Target Construct and Data Mapping

#pragma omp target { ... } offloads code region to GPU. Map clause specifies data transfer: map(to:x) copies x from host to device, map(from:y) copies y device-to-host, map(tofrom:z) copies bidirectionally, map(alloc:w) allocates on device without initialization. map(delete:...) deallocates after region. Implicit data mapping (firstprivate, private) defaults to tofrom for scalars; arrays are private (not mapped). Data persistence across targets requires enter/exit data directives.

GPU Thread Hierarchy

teams distribute over GPU thread blocks. distribute parallelizes outer loop over teams. parallel for parallelizes inner loop over threads within team. Combined: #pragma omp target teams distribute parallel for { for (i=0; i

Simdization and Vector Operations

#pragma omp simd vectorizes loops within teams. Pragma combination: #pragma omp target teams distribute parallel for simd enables full GPU utilization (blocks, threads, vector lanes). simdlen clause specifies vector width; collapse clause fuses loop nests.

OpenMP 5.x Enhancements

OpenMP 5.0 (2018): requires clause enables unified shared memory (requires unified_shared_memory)—simplifying data movement, pointers accessible on both host/device. OpenMP 5.1 (2020): interop clause enables mixed CUDA/OpenMP code. OpenMP 5.2+ (2021+): memory spaces (omp_memspace_default, omp_memspace_large_cap) enable explicit memory hierarchy control; allocate directive binds arrays to specific spaces; interoperability with SYCL/HIP via vendor extensions.

Compiler Support

LLVM (Clang) and GCC support OpenMP offloading to NVIDIA/AMD/Intel GPUs via translation to CUDA/HIP. Performance depends on compiler maturity; LLVM Clang leads, GCC lags. Overhead: first target region incurs startup/compilation cost (~100 ms); subsequent regions execute kernels quickly.