SIVARO
How to Use Multiple GPUs With OpenMP Offloading
I spent three weeks in early 2025 trying to get OpenMP offloading to scale across eight NVIDIA H100s for a client's LLM inference pipeline. The documentation was sparse. The examples were toy-sized. And every forum post I found was either from 2019 or written by someone who clearly never compiled the code they were sharing.
Here's what I learned the hard way.
What OpenMP Offloading Actually Is
OpenMP offloading is the target directive family that lets you write GPU code without leaving your C/C++ or Fortran source. You mark a region with #pragma omp target teams distribute parallel for, and the compiler generates device code, transfers data, and launches kernels. It's not CUDA. It's not HIP. It's a portable abstraction that compiles to whatever your compiler supports — NVPTX for NVIDIA, AMDGCN for AMD, SPIR-V for Intel.
Multi-GPU OpenMP offloading means taking that same abstraction and spreading your work across multiple devices. Same directives, new pragmas, and a set of runtime functions that let you query devices, set the active device, and manage per-device memory.
The key insight: OpenMP's multi-GPU model is device-centric, not execution-centric. You're not writing a single kernel that magically runs on many GPUs. You're writing logic that runs on the host, and you're dispatching work to specific devices.
The Architecture You Need to Understand
Most people think "multiple GPUs" means one big kernel split across devices. Wrong. That's not how OpenMP works, and honestly, it's not how most production systems work either.
OpenMP's model is simpler: you have a host (CPU) that orchestrates. Each GPU is a device. You use omp_get_num_devices() to count what's available, omp_set_default_device() to pick one, and omp_get_default_device() to ask which one is active.
The architecture looks like this:
Host CPU (OpenMP host threads)
├── Device 0: NVIDIA H100
├── Device 1: NVIDIA H100
└── Device 2: AMD MI300X
Data lives on specific devices. Computations map to specific devices. The host thread that encounters a target region assigns it to the default device — unless you tell it otherwise.
Here's the painful part I learned debugging a silent performance regression: implicit data mapping happens per device. If you write to a variable in a target region on Device 0, then access it on Device 1, you're looking at a host round-trip. The compiler won't warn you. The numbers will just be terrible.
Multi-GPU Programming Models That Work
Forget what the textbooks say. There are three patterns that actually work in production, and I've tested all three.
Model 1: Data-Parallel Slicing (The Simple One)
You split your data into N chunks (N = GPU count), map each chunk to a device, and run the same kernel on each. This is embarrassingly parallel, which means it's embarrassingly easy.
int num_devices = omp_get_num_devices();
int chunk_size = total_size / num_devices;
#pragma omp parallel for
for (int d = 0; d < num_devices; d++) {
omp_set_default_device(d);
#pragma omp target data map(to: input[chunk_size*d : chunk_size], \
coeffs[0 : coeffs_size]) \
map(from: output[chunk_size*d : chunk_size])
{
#pragma omp target teams distribute parallel for
for (int i = 0; i < chunk_size; i++) {
output[chunk_size*d + i] = process(input[chunk_size*d + i], coeffs);
}
}
}
Benchmark: I ran this on 4×A100s with a 400M-element array, 32-bit floats. Got 3.8× scaling. The 0.2 loss is host-side overhead from the target data setup and the thread synchronization in the outer parallel region.
Model 2: Pipeline Parallelism (The Sequential One)
For models where layer k+1 needs layer k's output, you can't slice the data. You slice the model. Each GPU owns a stage of the pipeline. This is how SIVARO deployed a 70B parameter model across 4×H100s in 2025 — not with Tensor Parallelism via NCCL, but with plain OpenMP offloading and the stages manually managed.
The trick is omp_get_initial_device() — you explicitly move intermediate results back to the host between stages:
omp_set_default_device(0);
#pragma omp target data map(to: input[0:N], params_0[0:P0]) \
map(from: stage1_out[0:N])
{
#pragma omp target teams distribute parallel for
for (int i = 0; i < N; i++)
stage1_out[i] = layer0_forward(input[i], params_0);
}
// Host boundary — unavoidable, but can be overlapped
omp_set_default_device(1);
#pragma omp target data map(to: stage1_out[0:N], params_1[0:P1]) \
map(from: stage2_out[0:N])
{
#pragma omp target teams distribute parallel for
for (int i = 0; i < N; i++)
stage2_out[i] = layer1_forward(stage1_out[i], params_1);
}
This works. It's also slower than a CUDA-based NCCL all-reduce approach by about 40% in my tests, because the host round-trip between stages adds latency. But it compiles anywhere, and for models with small activations (think < 100MB between stages), the hit is acceptable.
Model 3: Task-Based Dynamic Scheduling (The Flexible One)
This is the pattern I wish I'd known about earlier. OpenMP tasks (#pragma omp task) can be targeted to specific devices. The runtime handles load balancing — so if one GPU is busy, another picks up the next task.
#pragma omp parallel
{
#pragma omp single
for (int i = 0; i < num_tasks; i++) {
#pragma omp task device(i % omp_get_num_devices())
{
// Executed on device (i % num_devices)
process_batch(i);
}
}
}
The device() clause on the task directive is underdocumented. It's been in the spec since OpenMP 4.5, but GCC only added proper support in GCC 14 (released mid-2024), and Clang's support in LLVM 17 was buggy for certain datatypes. If you're on GCC 14+ or LLVM 18+, this works beautifully.
Data Mapping — Where Everyone Screws Up
The most common bug in OpenMP multi-GPU offloading isn't the compute. It's the data. Here's the rule:
Every device has its own memory space. The host must explicitly move data in and out. And data on Device 0 is NOT accessible on Device 1.
I can't count how many SIVARO clients showed me code where they did:
#pragma omp target teams distribute parallel for device(0)
for (int i = 0; i < N; i++) {
arr[i] = compute(arr[i]);
}
#pragma omp target teams distribute parallel for device(1)
for (int i = 0; i < N; i++) {
arr[i] = compute2(arr[i]); // arr on device 1? Nope. Not mapped.
}
The second region reads arr from host memory, copies it to Device 1, computes, copies back. Two implicit transfers. Per iteration of the outer host loop. Disaster.
The fix is explicit target data regions per device, and if you need data on multiple devices simultaneously, map(to: arr[0:N]) on each device's data region. The runtime tracks the mapping if you use omp_target_alloc and omp_target_memcpy explicitly, but those are verbose.
For the annotation-heavy style, this is the pattern I ended up settling on:
double *arr = (double*) malloc(N * sizeof(double));
// ... initialize on host ...
// Copy to ALL devices once
#pragma omp parallel num_threads(omp_get_num_devices())
{
int tid = omp_get_thread_num();
if (tid < omp_get_num_devices()) {
omp_set_default_device(tid);
#pragma omp target enter data map(to: arr[0:N])
}
}
That omp_target_enter data with map(to:) is the "persistent mapping" directive. The data stays on the device until you exit the data environment or call omp_target_exit_data. This is your friend.
Vendor-Specific Behavior You Need to Know (August 2026 Reality Check)
The spec is portable. The implementations are not.
-
NVIDIA (via Clang/LLVM): The most mature path.
-fopenmp -fopenmp-targets=nvptx64-nvidia-cudaworks, and recent LLVM 19/20 releases handle complex data types well. But you need CUDA toolkit installed even though you're not writing CUDA — the compiler usesnvptxtools for linking. -
AMD (via Clang): Requires
-fopenmp -fopenmp-targets=amdgcn-amd-amdhsa. Historically flaky with shared libraries. LLVM 18 fixed most issues. If you're on an MI300-series, expect driver-level weirdness withomxmemory pooling. -
Intel (via oneAPI/LLVM):
-fiopenmp -fopenmp-targets=spir64works on discrete Arc GPUs. Honestly, you probably aren't using this for serious HPC. -
NVIDIA (via NVHPC/NVC++): The commercial compiler handles multi-GPU better than open-source Clang for certain patterns, because the runtime is more mature. But you're locked into NVIDIA's ecosystem. We used NVHPC 24.7 at SIVARO in 2025, and it was stable until we wanted AMD GPUs for a cost optimization. Then we had to rewrite.
My recommendation: standardize on Clang/LLVM if you can. It's the only compiler that targets all three major vendors. GCC 15 finally added target GPU support for NVIDIA, but AMD support is still experimental as of mid-2026.
OpenMP Offloading Multi-GPU Example Code: A Complete Kernel
Let me give you a full example to make all of this tangible. This is the pattern I've used for batched inference across multiple GPUs.
#include <omp.h>
#include <stdio.h>
#include <stdlib.h>
#define NUM_ELEMENTS 1000000
#define BATCH_SIZE 250000
int main() {
int num_devices = omp_get_num_devices();
printf("Found %d devices
", num_devices);
if (num_devices < 2) {
printf("Need at least 2 GPUs
");
return 1;
}
// Host data
float *host_input = (float*) malloc(NUM_ELEMENTS * sizeof(float));
float *host_output = (float*) malloc(NUM_ELEMENTS * sizeof(float));
// Initialize
for (int i = 0; i < NUM_ELEMENTS; i++) {
host_input[i] = (float)i;
host_output[i] = 0.0f;
}
// Step 1: Partition indices per device
int chunk_size = NUM_ELEMENTS / num_devices;
int remaining = NUM_ELEMENTS % num_devices;
// Step 2: Offload per device using separate target regions
#pragma omp parallel num_threads(num_devices)
{
int device_id = omp_get_thread_num();
omp_set_default_device(device_id);
// Compute this device's slice
int start = device_id * chunk_size + (device_id < remaining ? device_id : remaining);
int count = chunk_size + (device_id < remaining ? 1 : 0);
// Map data explicitly
#pragma omp target data map(to: host_input[start : count]) \
map(from: host_output[start : count])
{
#pragma omp target teams distribute parallel for
for (int i = 0; i < count; i++) {
int global_idx = start + i;
// Your compute kernel
host_output[global_idx] = host_input[global_idx] * 2.0f + 1.0f;
}
}
}
// Verify on host
int errors = 0;
for (int i = 0; i < NUM_ELEMENTS; i++) {
float expected = host_input[i] * 2.0f + 1.0f;
if (host_output[i] != expected) errors++;
}
printf("Errors: %d
", errors);
free(host_input);
free(host_output);
return errors == 0 ? 0 : 1;
}
Compile with:
bash
clang -fopenmp -fopenmp-targets=nvptx64-nvidia-cuda -O3 -o multi_gpu multi_gpu.c
Run with:
bash
OMP_NUM_DEVICES=4 ./multi_gpu
The OMP_NUM_DEVICES env variable caps how many devices OpenMP sees. Useful for testing on a shared node.
Hidden Gotchas (From Production)
Gotcha 1: The Host Thread Default
OpenMP defaults to "any device" unless you specify. On some runtimes, omp_set_default_device(-1) means no device, and any target region compiles to host execution. I saw this corrupt data silently for two weeks. Always explicitly set the default device before entering a target region.
Gotcha 2: Memory Pinning
Implicit map operations are slow because they involve pageable memory transfers. For multi-GPU, where you're bouncing data between host and multiple devices, use omp_target_alloc for pinned memory:
float *dev_ptr = (float*) omp_target_alloc(N * sizeof(float), device_id);
This allocates device memory and returns a device pointer. Combine with omp_target_memcpy(dst, src, N * sizeof(float), dst_offset, src_offset, dst_device, src_device) for direct GPU-to-GPU copies. In my testing, this was 2.5× faster than implicit mapping for the same operation.
Gotcha 3: Thread Safety of Device Access
Each device can handle concurrent kernel launches from multiple host threads, but the runtime's device mapping isn't always thread-safe. I recommend one host thread per device (as in my example). It aligns with the natural parallelism and avoids race conditions in the runtime.
Gotcha 4: The nowait Clause
You can't use nowait with target regions that have data dependencies. Found that out during a training run that produced NaN gradients because Device 1 started reading data Device 0 hadn't finished writing. The spec technically allows it, but implementations are inconsistent.
When Multi-GPU OpenMP Is the Wrong Tool
Here's the honest part. If you need peak performance on a single node, and you're on NVIDIA only, just use CUDA + NCCL. You'll get 1.2–1.5× better performance because:
-
CUDA gives you direct control over peer-to-peer transfers. OpenMP offloading doesn't have a standard for GPU-to-GPU P2P — data goes through host memory unless you manually use
omp_target_memcpywith careful device IDs and hope the runtime uses the NVLink path. -
OpenMP's runtime adds overhead. Layer-level offloading for transformer models has 10–15% more launch latency than raw CUDA graphs with CUDA graphs capture. Field prediction from LLVM's own developers at the LLVM Developers Meeting in 2025 confirmed that the runtime overhead in OpenMP offloading is roughly 2–4x CUDA's launch overhead for small kernels.
-
Debugging is harder.
cuda-gdbjust works. OpenMP'sgdbsupport for GPU code is a patchwork of vendor extensions.
But if you need portability across NVIDIA and AMD (which is why most clients come to SIVARO), or you have existing Fortran/C++ OpenMP CPU code that you want to accelerate without a full rewrite, OpenMP offloading is worth it. You trade 20-30% performance for 80% less rewrite effort.
The Compiler Stack That's Actually Reliable
I've been testing compiler versions since 2022. Here's what I'd trust today:
- LLVM/Clang 19+: Best overall. Multi-target support (you can build one binary for NVIDIA and AMD). Good diagnostics. Active development.
- GCC 15: NVIDIA support is solid. AMD support improved but still spotty. Use it if you're already a GCC shop.
- NVHPC 25.x: Rock solid for NVIDIA, but it's a dead end for portability. If you're all-NVIDIA and want stability, this is your bet.
Avoid mixing compilers for the same binary. If you compile the kernels with Clang but link with GCC's libgomp runtime, you'll get runtime errors about inconsistent OpenMP versions. This wasted a full week for me earlier this year.
My Final Takeaway
Multiple GPUs with OpenMP offloading works if you respect three things:
-
Data placement is the programming model. Think in terms of device memory spaces, not global arrays.
-
One thread per device. Don't try to manage multiple devices from a single host thread unless you're using task-based parallelism — and even then, test carefully.
-
Measure before you optimize. I keep seeing people write elaborate hand-tuned multi-GPU code when a single GPU with a better algorithm would suffice. We tested a SIVARO client's digital signal processing workload in April 2026: single A100 with optimized kernels beat 4×V100s with naive multi-GPU slicing by 2.1×. More GPUs isn't automatically better.
The tooling has caught up. Clang 19+ handles multi-GPU offloading well, the official OpenMP examples in the 5.2 spec cover more multi-GPU patterns, and I've seen production deployments at hedge funds and robotics companies using it in 2026 without pulling their hair out.
Get the data mapping right. Make the host orchestration explicit. And don't be afraid to fall back to CUDA when the performance gap matters more than the portability benefit.
The spec is portable. The implementations are not. — Nishaant Dixit, August 29, 2026
Frequently Asked Questions
Q: Can I mix NVIDIA and AMD GPUs in the same OpenMP offloading program?
Yes, with LLVM/Clang 18+, you can compile with multiple -fopenmp-targets values and the runtime will dispatch to the appropriate device. But be careful: data mapping, peer-to-peer, and unified memory behavior differ. It works, but I'd shy away from mixing in production unless you have a compelling reason and build comprehensive tests.
Q: What's the difference between omp_set_default_device and the device() clause in target pragmas?
omp_set_default_device sets the device used for all subsequent targeting operations in the current thread's scope. The device() clause overrides this for a single construct. If you use neither, the default is implementation-defined — typically device 0. Always set it explicitly.
Q: Is there a performance difference between using one target data region with multiple target regions inside versus separate target data per device?
The first pattern (one target data region spanning multiple devices) doesn't make sense semantically — a target data region binds to one device. Separate target data per device is the right structure. The only alternative is using omp_target_alloc and omp_target_memcpy for manual control, which gives you flexibility but more code.
Q: Does OpenMP offloading support peer-to-peer GPU communication (NVLink or Infinity Fabric)?
The OpenMP 5.2 spec doesn't standardize P2P. In practice, with LLVM's libomptarget, omp_target_memcpy on NVIDIA systems can use cudaMemcpyPeer when the source and destination device pointers are known. It works, but performance isn't guaranteed and I've seen cases where it falls back to staging through host memory. Test on your specific hardware.
Q: What is OMP_TARGET_OFFLOAD and how does it affect multi-GPU code?
OMP_TARGET_OFFLOAD=DISABLED forces all target regions to execute on the host. It's a debug tool. OMP_TARGET_OFFLOAD=MANDATORY will error if target execution fails (e.g., no GPU found). When developing multi-GPU code on a single-GPU machine, set OMP_TARGET_OFFLOAD=MANDATORY to catch accidental host fallback.
Q: How do I compile a Fortran program with OpenMP offloading for multiple GPUs?
Same directives apply. Clang and Flang support most of the target directives in Fortran. The main difference is array syntax and numeric storage ordering. Fortran's column-major ordering can impact data mapping — be explicit with map(tofrom: arr(:)) over array sections.
Q: Is MPI + OpenMP offloading a good combination?
Yes. The classic pattern in the 2026 HPC world is MPI handles inter-node communication, while OpenMP offloading handles multi-GPU on each node. This hybrid MPI+OpenMP is the most common pattern in production HPC today. Use MPI rank 0 to orchestrate OpenMP device assignment. Just be aware that vendor-specific MPI implementations tie you to specific hardware — refer to the MPI Forum's OpenMP interop guidelines for theoretical details.
Q: What are the debugging tools for OpenMP offloading multi-GPU code?
For NVIDIA, use nsys and ncu to profile; both support OpenMP offloading. For AMD, use rocprof. The LLVM runtime supports LIBOMPTARGET_DEVICE_RTL_DEBUG=1 and LIBOMPTARGET_KERNEL_TRACE=1 environment variables for basic runtime diagnostics. And gdb has partial support with the omp device commands. None of this is as polished as CUDA tooling, but it's mature enough for production debugging.
Nishaant Dixit — Founder of SIVARO. Building data infrastructure and production AI systems since 2018. Built systems processing 200K events/sec.