Six Easy Pieces — the OpenACC way
Background and Overview
When it comes to looking at modernization options for legacy C (or C++) code running on edge devices, there are uncommon asks to “leverage mult-core capabilities”, “leverage GPU capabilities”, “rewrite for C++17 or C++20” among a host of other requirements. Sometimes some basic tinkering with available hardware (in this case the Jetson AGX Orin that is NOT a discrete PCIe CUDA (basically an Integrated Ampere GPU (Compute Capability 8.7)) can reveal options that are simpler than a brute-force rewrite to CUDA — Enter OpenACC (https://www.openacc.org/)
Even though OpenACC was originally designed with discrete NVIDIA GPUs in mind, its ability to automatically parallelize loops and control data movement through simple pragmas makes it extremely valuable even when running purely on CPU-only ARM systems like the Jetson Orin or Raspberry Pi 5 with ARM64 Linux.
How do I get access to OpenACC that can be validated on a GPU?
NVIDIA’s HPC (High Performance Computing) SDK can be used for basic experiments to validate this idea. Although NVC++ (the C++ compiler provided by the HPC SDK) does NOT officially support Jetson AGX Orin, NVIDIA has a release that can be used on an Ubuntu Linux Arm Server. This is good enough since L4T (Linux for Tegra) is part of NVIDIA’s JetPack SDK and provides a development environment for projects that require both ARM-based processing and powerful GPU acceleration. Note that even basic std::execution::par constructs of C++17 won’t be supported on this GPU. So we can safely exclude that from our trivial benchmarking exercise.
What is being benchmarked here?
Pick a compute-intensive mathematical operation (could be matrix multiplication that is most common with GPUs — imagine one thread assigned to “multiply” each element of a 2048x2048 matrix and the parallelism involved. It’s going to be obviously fast compared to a typical for loop that you would otherwise use in a non-CUDA construct with for loops etc.). Here we’ve taken an even simpler mathematical operation (complex enough to demonstrate benefits of the ACC & CUDA approach against the traditional CPU based sequential computation). So what exactly is the math operation here:
“Compute the sum of sin(i) \* cos(i) for about 100 million values of i (ranging from 0 to a 100 million)”. It is not difficult to imagine the code to look something like this:
const int N = 100000000; …
for (int i = 0; i < N; i++) { sum += std::sin(i) \* std::cos(i); }
We look at the latency associated with the above “intensive” compute as follows:
- Without any OpenACC (#pragma acc parallel loop reduction(+:sum)) directive
- With the OpenACC directive
- Refactoring the above code to just use brute-force CUDA Runtime
How do we expect the “compute” workflow / execution to happen?
Something like the figure below:
The typical optimization / modernization options to make a code GPU aware
Show me the code
Well…here you go: https://github.com/Quest1Codes/nvidia-jetson-agx-orin
Check for availability of the OpenACC runtime
\# cd /mnt/quest1/hpc-sdk/nvidia-jetson-agx-orin/acc \# nvc++ -acc -gpu=cc87 checkacc.cpp -o checkacc \# ./checkacc OpenACC supported. Version: 201711
Check latency when using only CPU. No optimization or use of GPU. Just serial computation
\# ./test\_acc —acc=off OpenACC unused or not supported Sum \= 0.447924 Execution Time: 1905.23 ms
Check the impact of the pragma directive — An interim option for migration
\# cd /mnt/quest1/hpc-sdk/nvidia-jetson-agx-orin/acc \# nvc++ -acc -gpu=cc87 TrigonometricOps\acc.cpp -o test\acc \# ./test\_acc —acc=on Offloading to GPU using OpenACC : \[201711\] Sum = 0.447924 Execution Time: 369.992 ms
Check the impact on latency due to direct usage of CUDA Runtime — The Re-write option
\# cd /mnt/quest1/hpc-sdk/nvidia-jetson-agx-orin/cuda \# nvcc -arch=sm\87 -o cuda\trig TrigonometricOps.cu -lm \# ./cuda\_trig Total sum \= 0.447924 CUDA Execution Time: 187.655 ms
The Trivial Benchmark Summary
Obvious takeaways
- Clearly CUDA delivers the highest performance, thanks to explicit control over thread hierarchy, memory access patterns, and kernel launches.
- OpenACC provides a great migration path — significantly reducing latency with minimal code changes, acting as the “middle ground” between CPU-only and full CUDA rewrites.
- CPU-only sequential code is easy to write but becomes the bottleneck for compute-heavy mathematical workloads.
Six Easy Pieces — Some practical use-cases
_for modernizing traditional C (or C++ code) with OpenACC pragma decorations_
1\. Accelerating Sensor Data Processing on Robotics Platforms (Jetson / ROS2)
- Robotic workloads like LiDAR point cloud filtering, sliding window smoothing, or Kalman filter updates often involve large independent loop operations.
- These can be accelerated using:
#pragma acc parallel loop for (int i = 0; i < N; i++) { filtered\[i\] = raw\[i\] \* gain; }
Implication: Immediate speedup on ARM multi-core without writing pthreads / OpenMP manually.
2\. Fast Image Processing Without OpenCV GPU Modules
Suppose you’re doing Gaussian blur / Sobel / thresholding in plain C++ (or CUDA is overkill or unavailable):
#pragma acc parallel loop collapse(2) for(int y=0; y<H; y++) { for(int x=0; x<W; x++) { output\[y\]\[x\] = (input\[y\]\[x\] + input\[y\]\[x+1\] + input\[y\]\[x\-1\]) / 3.0f; } }
3\. Scientific & Numerical Code Modernization (Fast Refactoring Without Rewrites)
Legacy simulation codes, DSP pipelines, or numerical kernels in finance, RF, or control systems can be sped up using incremental annotations instead of architectural rewrites.
Implication: Ideal scenario: Legacy engineering code → “Just add pragmas” and benchmark.
4\. Time-Critical Loops in ML Inference Pre/Post-Processing
- Even when using TensorRT or PyTorch on Jetson, custom tensor transformations and normalization loops can be bottlenecks.
- Instead of re-writing in CUDA, annotate (example shown below):
#pragma acc loop vector for(int i = 0; i < N; i++) x\[i\] = (x\[i\] - mean) / var;
5\. Parallel CSV / JSON Parsing or Compression on Edge Gateways
Even string and file parsing can benefit from loop-level auto-vectorization + parallelization.
6\. Accelerating a 2D Matrix Multiplication with #pragma acc kernels
Suppose you have a classic CPU-only matrix multiplication:
Concluding Remarks & Thoughts
Even though NVIDIA’s full-featured HPC SDK does not officially target embedded GPUs like the Jetson Orin AGX, the underlying OpenACC approach remains an invaluable tool — not for GPU offload alone, but as a way to unlock multi-core performance on ARM CPUs with almost no code restructuring. It sits in a sweet spot between raw CUDA and standard C++, allowing incremental performance tuning through lightweight pragmas.
In an era where even edge devices must run heavy computation, OpenACC offers a rare gift: scalability without complexity, through simple and highly effective pragma decorations.