[Lab 2] Expected slowdown from naive OpenACC no longer observed on modern hardware and compiler

[takuya4880]

Summary

Lab 2 (profiling_lab2.ipynb) is designed to show that a naive OpenACC parallelization
makes miniWeather slower on GPU than on CPU, due to excessive host↔device data movement.
However, on modern hardware and with current compiler versions this slowdown is no longer
observed — the GPU version is already significantly faster even without the manual
optimizations taught in later labs.

Observed Behavior

Tested on two systems:

System	CPU (serial)	GPU (naive OpenACC, Lab 2)
NVIDIA GB200 (NVLink-C2C unified memory)	~28 s	~860 ms
AMD CPU + NVIDIA A100	~28 s	similar speedup observed

Both results are the opposite of what the notebook predicts. This has been confirmed
with NVIDIA engineers as expected behavior on newer architectures.

Expected Behavior (per notebook)

The notebook states:

"You can see that the changes made actually slowed down the code and it runs slower
compared to the non-accelerated CPU-only version."

It also shows a repeating pattern of GPU idle time surrounded by PCIe data movement
(purple/teal boxes in Nsight Systems), which was the bottleneck on the hardware used
when the lab was originally authored.

Root Cause Analysis

Two factors likely contribute:

1. Hardware architecture improvements

• GB200 (Grace-Hopper / Blackwell): NVLink-C2C unified memory eliminates the
  PCIe data movement bottleneck entirely. Host↔device transfers that were costly on
  older PCIe-based systems are now near-zero cost, so the naive port benefits from
  GPU compute without paying the data movement penalty.
• A100: Even on PCIe-based A100 systems, faster HBM2e memory bandwidth significantly
  reduces the relative cost of data movement compared to older GPU generations.

2. Compiler intelligence improvements

• Modern NVHPC compiler versions are smarter about OpenACC data movement. Even from
  naive #pragma acc parallel directives, the compiler may now automatically apply
  more efficient data locality strategies (e.g., better managed memory prefetching
  hints, improved loop analysis for data lifetimes).
• The -gpu=ccnative flag compiles for the exact GPU architecture present at build time,
  potentially enabling additional hardware-specific optimizations.
• This means the "excessive data movement" visible in Nsight Systems on older compiler
  versions may be partially or fully optimized away at compile time on current NVHPC releases.

Impact

The educational goal of Lab 2 is to motivate students to optimize data movement by
first showing them the cost. If students on modern systems with current compilers see
a speedup instead, they may not understand why Labs 3+ (data movement optimization)
are necessary.

Suggested Remedies

1. Add a note on hardware and compiler dependency: Explain that the slowdown is
   architecture- and compiler-version-dependent. On unified-memory systems (GH200, GB200)
   or with newer NVHPC compilers, the naive port may already be fast.

2. Reframe the narrative around Nsight Systems' ability to reveal data movement
   patterns regardless of absolute cost, emphasizing that these patterns matter
   at scale (multi-GPU, very large problems).

Environment

• Hardware tested: NVIDIA GB200 (NVLink-C2C unified memory), NVIDIA A100 (AMD CPU host)
• Compiler: NVIDIA HPC SDK nvc++ 26.5-0 (64-bit, x86-64 Linux, target: znver2)
• Profiler: NVIDIA Nsight Systems
• Compiler flags: -acc -gpu=mem:managed,ccnative -Minfo=accel
  • -gpu=mem:managed: CUDA Managed Memory
  • -gpu=ccnative: compile for the native GPU architecture at build time

References

• Notebook: HPC_Profiler/_profiler/jupyter_notebook/profiling_lab2.ipynb