OpenACC is not CUDA without the work. It is a compiler contract over loop scheduling, and it stops short exactly where the hard part of GPU performance begins: host-device data movement and memory layout. That distinction decides whether a directive-based port lands near a tuned kernel or plateaus somewhere you never measured.
The pragma model does deliver on speed of arrival. In our experience porting legacy Fortran and C++ solvers, OpenACC gets a first working GPU version running in days where an equivalent CUDA port takes weeks. What it does not do is restructure your arrays or rewrite your transfer pattern. The compiler will happily parallelise a loop nest whose inner stride guarantees uncoalesced access, and it will report success.
Where does OpenACC actually reach CUDA-level performance?
Directives close the gap on dense, regular loop nests with predictable access patterns and enough arithmetic per byte moved. Compute-bound kernels of that shape are the honest success case for the model.
The gap opens on memory-bound kernels. Once a workload needs explicit staging in shared memory, fused kernels to avoid round-tripping through global memory, asynchronous streams to overlap transfer with compute, or a layout change from array-of-structs to struct-of-arrays for coalesced access, more pragmas will not help. A directive-based version of a memory-bound kernel commonly lands well short of a tuned CUDA kernel on the same device, and the residual gap is closable only by data-layout and transfer changes.
The two are not mutually exclusive, which is the part teams miss. OpenACC and CUDA coexist in one binary — directives on the cold loops, explicit __global__ kernels on the hot ones, sharing device pointers. The decision is per kernel, not per codebase.
Per-kernel split: which model for which kernel
| Kernel property | Keep directive-based | Rewrite in CUDA | Evidence to check |
|---|---|---|---|
| Compute-bound, high arithmetic intensity | Yes | Rarely worth it | Achieved FLOP/s vs roofline |
| Memory-bound with poor coalescing | No | Yes | Achieved bandwidth vs device peak |
| Needs shared-memory staging or kernel fusion | No | Yes | Global-memory traffic per element |
| Runs once per timestep, small share of wall time | Yes | No | Profile share of total runtime |
| Transfer-dominated (host↔device per iteration) | Fix data regions first | Only after transfers are fixed | Transfer time vs kernel time |
The rule the table encodes: never rewrite a kernel in CUDA before a profiler has shown what limits it. NVIDIA Nsight Compute reporting achieved bandwidth well below device peak on a directive kernel is the signal that layout, not scheduling, is the constraint — and that is a measurement, not a guess. Treating OpenACC as a migration destination rather than an incremental step is where cost compounds, because the ceiling gets baked into the design before anyone has profiled it.
Portability is the other overclaim worth deflating. OpenACC and OpenMP target offload both compile for NVIDIA, AMD, and Intel GPUs, so the source moves. Tuning decisions do not: a layout and blocking choice tuned for one vendor’s memory hierarchy has to be re-examined on another. Source portability is real; performance portability is a separate piece of work.
This sits inside the broader memory-model argument we develop in our GPU acceleration and optimisation work, where the transfer pattern usually turns out to matter more than the kernel body.
If your directive-based port has stalled, the question we would ask first is not which framework to adopt — it is what achieved-bandwidth figure you have for the three kernels that dominate your runtime.
