Why is the execution time of cudaMalloc so variable (when using hotspot benchmark from Rodinia Benchmark suite)?

[abhishekghos]

I experimented with MPS on the hotspot CUDA application from the Rodinia-3.1 benchmark suite. A clone of that is maintained in my Git Hub.

(with MPS Enabled on the Linux System)

cuda/hotspot$ make
cuda/hotspot$ sudo nsys profile -o MPS-2-1 hotspot_launch.sh 2

The naming convention of the nsys output file used by me is as follows:
MPS-<no_of_invocation_to_hotspot>-<interaction number>
So, the above MPS-2-1 means launched hotspot 2 times under MPS for the first time. MPS-2-4 means launched hotspot 2 times under MPS for the fourth time. (The nsys report files are present in rodinia-3.1-hotspot/cuda/hotspot/MPS_NSYS_REPORTS)

But, I noticed a huge variability in the cudaMalloc time.

For instance:

1. First Iteration:
   

   
   
   The first cudaMalloc takes 184.398 ms, while the second cudaMalloc takes 172.541 ms

2. Second Iteration:
   

   
   
   The first cudaMalloc takes 58.08 ms, while the second cudaMalloc takes 168.570 ms

3. Third Iteration:
   

   
   
   The first cudaMalloc takes 168.320 ms, while the second cudaMalloc takes 60.238 ms

And so on… Why is there this amount of variability?

---

Now I disable MPS and run the same experiments and profile them.
(with MPS disabled on the Linux System)

cuda/hotspot$ sudo nsys profile -o NO-MPS-2-1 hotspot_launch.sh 2

The naming convention of the nsys output file used by me is as follows:
NO-MPS-<no_of_invocation_to_hotspot>-<interaction number>
So, the above NO-MPS-2-1 means launched hotspot 2 times under no MPS (time-slicing) for the first time. (The nsys report files are present in rodinia-3.1-hotspot/cuda/hotspot/NO_MPS_NSYS_REPORTS)

Now the time of cudaMalloc seems to be double of what was reported with MPS enabled.

1. First Iteration:
   

   
   
   The first cudaMalloc takes 239.444 ms, while the second cudaMalloc takes 223.165 ms.

2. Second Iteration:
   

   
   
   The first cudaMalloc takes 253.195 ms, while the second cudaMalloc takes 244.965 ms.

[Robert_Crovella]

Possible reasons for variability in cudaMalloc time duration:

1. The first cudaMalloc call in an application may be also one that incurs CUDA initialization overhead.
2. cudaMalloc is a synchronizing operation. That means it cannot complete until all previous device activity in that context is idle. So viewed from an API perspective, if there is ongoing device activity (kernel execution) when the cudaMalloc call is issued, it will wait there (from an API perspective) until the device is idle.
3. In a multithreaded situation (and perhaps MPS multi-process), the CUDA runtime may institute locks to serialize certain activity. The acquisition cost of these locks will vary depending on instance specifics.

There may be other reasons for this as well.

For these reasons, canonical advice is to do all device memory allocations early in your application, and keep them out of concurrent work issuance loops/situations.

As an alternative when this is unavoidable, the cuda memory pool system can be investigated.

[njuffa]

It is typical that one observes a wide distribution of execution times for allocation requests to dynamic memory allocators. Most allocators, including the CUDA allocator, comprise multiple layers, and typically the cost of an allocation request depends on how deep the allocator has to reach into this hierarchy. This in turn depends on the size of an allocation request, attributes of the allocation request (if any), and the present internal state of the allocator (e.g. fragmentation).

In order to further comment on this specific case one would need to know details of the cudaMalloc() calls associated with these profiles. The cluster of intense “OS runtime libraries” activity associated with each cudaMalloc() call shown suggests to me that these are large allocation causing the allocator to dip down to the deepest (most expensive) allocator layer. However, without knowing the internal details of the CUDA allocator (as far as I know, not documented by NVIDIA, it may well change over time) I am not particular confident in this diagnosis.