OpenMP Offload

Learning Objectives

  • Explain how OpenMP target directives offload computation to a GPU

  • Map data between host and device with map(to:), map(from:), and map(tofrom:)

  • Use the target teams distribute parallel for construct and explain how it maps to GPU hardware

  • Query for available devices at runtime and fall back to the host gracefully

Three Ways to Run the Triad on a GPU

Approach

Who writes the kernel

Who moves the data

Portability

CUDA

You (explicit __global__)

You (explicit cudaMemcpy)

NVIDIA only

OpenMP offload

Compiler (from your loop)

You (declarative map clauses)

Multi-vendor

std::par

Compiler

Compiler (unified memory)

Compiler-dependent

The Offload Execution Model

A GPU is organized as: many teams (CUDA: thread blocks), each containing many threads. OpenMP exposes this hierarchy:

#pragma omp target teams distribute parallel for
for (int i = 0; i < n; i++) {
  c[i] = a[i] + b[i];
}

Reading left to right:

  • target — execute on the default device (GPU). Host hands off to device.

  • teams — launch a league of teams (GPU thread blocks/SMs).

  • distribute — partition loop iterations across teams.

  • parallel for — within each team, run iterations on multiple threads.

Note

teams ≈ grid of thread blocks; threads from parallel for ≈ threads within a block.

Mapping Data Between Host and Device

#pragma omp target data map(to: a[0:n], b[0:n]) map(from: c[0:n])
{
  #pragma omp target teams distribute parallel for
  for (int i = 0; i < n; i++) {
    c[i] = a[i] + b[i];
  }
}

Clause

Meaning

CUDA analogue

map(to: a[0:n])

Copy host → device before region

cudaMemcpy(..., HostToDevice)

map(from: c[0:n])

Copy device → host after region

cudaMemcpy(..., DeviceToHost)

map(tofrom: x[0:n])

Copy both directions (default)

Both copies

map(alloc: t[0:n])

Allocate on device, no copy

cudaMalloc only

Warning

Defaulting everything to tofrom doubles your PCIe traffic. The triad only reads a and b and only writes c, so map(to: a, b) and map(from: c) move each array across the bus exactly once.

Device Discovery and Fallback

int num_devices = omp_get_num_devices();
printf("Number of available devices: %d\n", num_devices);

if (num_devices < 1) {
  printf("No GPU devices available, falling back to CPU execution\n");
} else {
  printf("Using GPU device 0\n");
  omp_set_default_device(0);
}

omp_get_num_devices() returns the number of attached accelerators. If no device is found, target regions still execute correctly on the host, so the same binary runs on a workstation without a GPU.

Build Flags

Compiler

CPU multithread

GPU offload (NVIDIA)

GCC

-fopenmp

-fopenmp -foffload=nvptx-none

Clang

-fopenmp

-fopenmp -fopenmp-targets=nvptx64

NVHPC (nvc++)

-mp

-mp=gpu

Warning

Stock GCC is usually built without the NVPTX offload plugin. On Debian/Ubuntu you need gcc-offload-nvptx. The NVIDIA HPC SDK (nvc++ -mp=gpu) is often the smoother path for NVIDIA hardware.

Running on ARC Clusters with GPU hardware

  1. Load compiler with offload support:

    module load nvhpc

  2. Request GPU in SLURM:

    #SBATCH --gres=gpu:1
#SBATCH --partition=<valid_gpu_partition>

  3. Compile and submit:

    nvc++ -mp=gpu -O3 stream_triad.cpp -o stream_triad
sbatch run.slurm

References