I have been running my code on both non-unified memory setup (CPU + Hopper GPU) and an NVIDIA Grace Hopper system with Unified Memory. However, I am observing better performance on the non-unified Hopper GPU setup compared to Grace Hopper.
In my current approach, I allocate all memory on the CPU (using malloc) and rely on Grace Hopper’s Unified Memory to migrate data between the CPU and GPU on-demand via page faults. However, I suspect this might not be the most efficient strategy. Would it be better to systematically allocate memory based on expected first access? That is:
If the GPU is expected to access it first, allocate on the GPU.
If the CPU is expected to access it first, allocate on the CPU.
Since my codebase is quite large, I opted for the simple strategy of always allocating on the CPU to avoid extensive modifications. However, this seems to be negatively impacting performance on Grace Hopper compared to a traditional CPU + Hopper GPU setup.
Additionally, is malloc a better choice than cudaMallocManaged or vica-versa on Grace Hopper for better performance and memory management ?
I would appreciate any advice or insights from the community!
Even with very sophisticated driver prefetching heuristics, on-demand access with migration will never beat explicit bulk data copies or prefetches in terms of performance for large contiguous memory regions. This is the price for simplicity and ease of use. If the application’s access pattern is well defined and structured you should prefetch usingcudaMemPrefetchAsync
The GH200 whitepaper, however does not clarify what @tonywu93 accurately stated above. The performance gain is attributed to ATS when it should actually be NVLINK instead
The performance impact is projected in Figure 18 and highlights how applications transparently benefit from Grace Hopper features like ATS without any application-side changes.
Thanks for the reply. However, shouldn’t the native code of any project running on Grace Hopper be faster than on a traditional x86 CPU paired with a Hopper GPU, since ARM-based RISC architectures are generally considered faster than x86 for certain workloads, and the data transfer between CPU and GPU in Grace Hopper is also significantly faster?
That is a major misconception. There is no specific technical advantage (“speeds & feeds”) to the choice of ARM platforms. System-level performance is a function of overall system design and implementation specifics. From the data I have come across, the fastest x86-64 platforms beat NVIDIA’s integrated ARM64 solutions.
A reasonable indicator which host platforms deliver the highest system level performance is to check what host platforms are used as the basis of NVIDIA’s DGX brand of top-of-the-line systems. Since these are generally based on AMD and Intel platforms, I would assume that considerations other than purely technical ones enter into the choice. It is my understanding that NVIDIA’s latest Blackwell-based DGX models use Intel Xeon 6 processors (not sure of the model number).
What the use of ARM processor for the host does is make NVIDIA independent of CPU-based infrastructure from competitors, and in as far that NVIDIA is able to exercise microarchitectural choice, custom tailor the host platform to the most profitable GPU market, which is AI.
[Later:] This press release from Intel gives some details on the CPU in the latest DGX system:
The new Xeon 6 processors are available today, with one of the three currently serving as the host CPU for the NVIDIA DGX B300, the company’s latest generation of AI-accelerated systems. The NVIDIA DGX B300 integrates the Intel® Xeon® 6776P processor, which plays a vital role in managing, orchestrating and supporting the AI-accelerated system. With robust memory capacity and bandwidth, the Xeon 6776P supports the growing needs of AI models and datasets.
Thanks for your response. You’re right — the CPU itself isn’t necessarily faster in these systems.
However, in our workloads, the main performance bottleneck historically has been the CUDA APIs, due to excessive data transfers between the CPU and GPU.
NVLink-C2C offers significantly higher bandwidth and lower latency compared to traditional PCIe. Specifically:
NVLink-C2C can deliver up to 900 GB/s of coherent bandwidth between the Grace CPU and Hopper GPU.
In contrast, PCIe Gen4 provides up to 64 GB/s per x16 link (theoretical), and even PCIe Gen5 maxes out around 128 GB/s.
To summarize:
On Grace Hopper:
No need for explicit cudaMemcpy calls.
The memory is unified and coherent between the Grace CPU and Hopper GPU.
NVLink-C2C provides ultra-fast, direct memory access between CPU and GPU.
On x86 + Hopper:
Explicit cudaMemcpy calls are required to transfer data between CPU and GPU.
Communication uses PCIe, which is significantly slower than NVLink-C2C.
This often becomes a bottleneck in data-intensive workloads.
So again I am confused on why x86 + Hopper are performing better than Gracehopper ?
I cannot perform an analysis of an unknown workload relative to some of the differences between integrated and discrete systems you noted. The proper tool for such an analysis is a profiler. Have you given that a try?
Note that a unified physical memory is both a blessing and a curse. It is conducive to more efficient sharing of data between CPU and GPU. But a unified memory causes CPU and GPU to share the available bandwidth of one physical DRAM subsystem, so the bandwidth available to each is usually lower than in a system with separate system DRAM and GPU-attached DRAM.
Again, I do not know your workload, but outside of host/device transfers the performance of many CUDA API calls is mostly a function of single-threaded CPU performance. To my knowledge, high-end x86-64 CPUs still provide the best performance for that, with AMD EPYC leading the race. See SPEC CPU 2017 SPECspeed Integer results.
Could you please specify what profiling information you need? I’ve already profiled our workloads using Nsight Systems (note: our workloads are not similar to typical AI/ML workloads). As you correctly pointed out, our x86 CPU is an AMD EPYC 9474F 48-core processor, paired with the H100 NVL GPU.
In gracehopper, Each processor has its own dedicated memory (LPDDR5X for Grace, HBM3 for Hopper), but can access the other’s memory space via NVLink-C2C. So how is bandwidth a limitation, could you explain please ?
I’m trying to pinpoint whether our workloads are constrained by bandwidth or latency on both systems.
Are there any cuda hinting APIs you would suggest to improve the performance of our workloads ?
Right now we only use cudaMalloc on GPU, cudaMemcpy to transfer data and cudaFree
I cannot. I am not familiar with these systems or their architecture.
I do not need this information. I pointed you at the profilers (Nsight System, Nsight Compute) as tools for delivering the data that you will need to analyze the performance.
The work done by cudaMalloc() and cudaFree() is pretty much all single-threaded host-side activity. If these APIs are not used in accordance with best practices, they can take up a significant portion of application time. When you look at the time taken by these API calls in the profiler output, do you see a significant difference between the x86 system and the ARM system?
Depending on how your application is structured and what flavors of copy APIs are invoked, these data transfers may be handled concurrently with kernels running on the GPU and application code running on the host, which means their cost may not contribute, or may not contribute significantly, to the application’s end-to-end runtime. Look for overlap of these transfers with other activity in the timeline provided by the profiler.
The way I would tackle this kind of a performance comparison is to extract the timing for all individual components that contribute to end-to-end execution time and compare them one by one between the two systems. Look at the time taken for individual API calls, data transfers, and kernel invocations. Any salient differences? If that does not provide any useful leads, look at overlap/concurrency of activities. Any salient differences between the two system?
The moment one finds a significant difference between the two platforms one can then dive deeper to determine why that difference exists.
cudaMalloc and cudaFree are normally explicitly done outside of performance critical loops.
cudaMalloc is often replaced with the cudaMallocAsync variant, when working with streams.
As njuffa has explained, your individual timings for overlapping are important.
How long does your computation run? How much data has to be copied. How fast would the host-device transfer speed have to be to parallelize that fully? Is a parallelization possible in your application?
Besides benchmarking, that can also be calculated/estimated.