Dask-cuda: LocalCudaCluster startup-time scales with >1 second per GPU

[beckernick]

Spinning up a LocalCudaCluster using all 16 GPUs on a DGX-2 seems to consistently take 20+ seconds. Given the time seems to scale loosely linearly with the number of GPUs used, I wonder if some of this seemingly serial processing might be able to be done in parallel.

from dask_cuda import LocalCUDACluster
from dask.distributed import Client

%%time
cluster = LocalCUDACluster()
client = Client(cluster)
CPU times: user 889 ms, sys: 442 ms, total: 1.33 s
Wall time: 22.4 s

%%time
​
cluster = LocalCUDACluster(
    n_workers=3,
    CUDA_VISIBLE_DEVICES="0,1,2",
)
client = Client(cluster)
CPU times: user 595 ms, sys: 179 ms, total: 774 ms
Wall time: 3.42 s

[pentschev]

Interesting, this seems to go down all the way to dask/distributed itself:

In [1]: from dask.distributed import Client

In [2]: from dask.distributed import LocalCluster

In [3]: %%time
   ...: cluster = LocalCluster(n_workers=160)
   ...: client = Client(cluster)
CPU times: user 2.81 s, sys: 844 ms, total: 3.65 s
Wall time: 9.98 s

In [4]: %%time
   ...: cluster = LocalCluster(n_workers=80)
   ...: client = Client(cluster)
CPU times: user 4.06 s, sys: 2.22 s, total: 6.28 s
Wall time: 5.8 s

In [5]: %%time
   ...: cluster = LocalCluster(n_workers=40)
   ...: client = Client(cluster)
CPU times: user 3.19 s, sys: 3.22 s, total: 6.4 s
Wall time: 3.86 s

It seems like part of it is indeed serial, but I don't know much about the internals, maybe @mrocklin or @jcrist could give us some insight.

I believe this is much more noticeable here because we create a CUDA context for each process which is indeed slow, and if that happens serially, then I'm not surprised we're seeing 20+ seconds for 16 GPUs.

[quasiben]

I think this is all serial and is happening in _correct_state_internal:

https://github.com/dask/distributed/blob/b60c4bfc5f46a9aa2f1aacbda21eac38abcc807e/distributed/deploy/spec.py#L304-L318

[jakirkham]

What brought you to that conclusion?

[jcrist]

Yeah, this loop could be parallelized using asyncio.gather instead of linearly waiting for each worker: https://github.com/dask/distributed/blob/b60c4bfc5f46a9aa2f1aacbda21eac38abcc807e/distributed/deploy/spec.py#L315-L317.

Something like this should work (missing support for exception handling).

await asyncio.gather(*workers, return_exceptions=True)

[quasiben]

What brought you to that conclusion?

@jakirkham I ran LocalCUDACluster through snakeviz and saw that the most time spent was in a call to threading.lock. This is obviously iis not really the problem and instead looked at dask-cuda and traced the last call in LocalCUDACluster

https://github.com/rapidsai/dask-cuda/blob/06ea368bcca4c6a2cd54262c36b2c1bdbb8b7dc2/dask_cuda/local_cuda_cluster.py#L121

That points to distributed and the for-loop in _correct_state_internal

[quasiben]

I think this startup time is impacting runs on a DGX2 with workers for each GPU (16). The default connection time out is 30s and when creating 16 workers the timeout can often be met and a restart of the worker is triggered. I've filed an issue for worker failures here: https://github.com/rapidsai/ucx-py/issues/376. For now I would suggest extending the timeout value while we prioritize this work

distributed:
  comm:
    timeouts:
      tcp: 50s              # time before calling an unresponsive connection dead

[quasiben]

@beckernick @randerzander can you try increasing timeouts and see if this prevents workers from dying in your applications ?

[jakirkham]

Submitted PR ( https://github.com/dask/distributed/pull/3402 ) upstream. Though it make still need more work.

[jakirkham]

Just wanted to add one note here, each worker acquires a CUDA context. This can take a bit of time.

In [1]: %time import numba.cuda; numba.cuda.current_context()                   
CPU times: user 381 ms, sys: 321 ms, total: 702 ms
Wall time: 1.01 s
Out[1]: <CUDA context c_void_p(94358057364640) of device 0>

[quasiben]

@mt-jones had some thought here as to why there is a slow down. I think it was that the cuda driver is going to create context serially and will be doing a somewhat costly look up each time a context is created. The fact that we are concurrently creating contexts doesn't help here and we are blocked by the driver

[quasiben]

I think we can close this issue as there is nothing to be done here. The cuda driver will serially process cuda context creation

[mtjrider]

@quasiben
Can we clarify exactly numba is doing when the context is created?
Is this what we do in Dask?

I've noticed there is always an initial memory allocation to establish that context. This memory allocation has a cost which is proportional to the number of GPUs on the system.

A note about runtime (primary) context initialization: device code is just-in-time compiled if necessary and loaded into device memory, but this initialization is not necessarily triggered immediately. It is triggered on the first material function-call to the runtime.

[mtjrider]

Regardless, I think this cost must be paid.

[quasiben]

@mt-jones i assume numba is calling the CUDA driver to create a context. @gmarkall do you have time to chime in here ? In dask-cuda we simply call out to numba to create the context:
https://github.com/rapidsai/dask-cuda/blob/091f0eecdda7d5f99b197c811dc9b73a5b620ed7/dask_cuda/initialize.py#L106-L110

[gmarkall]

@mt-jones i assume numba is calling the CUDA driver to create a context. @gmarkall do you have time to chime in here ? In dask-cuda we simply call out to numba to create the context:
https://github.com/rapidsai/dask-cuda/blob/091f0eecdda7d5f99b197c811dc9b73a5b620ed7/dask_cuda/initialize.py#L106-L110

Numba does use the driver to create a context - that does happen serially in the driver for each of the 16 contexts created, and does take while to create each one. The first Numba call (IIRC) will be to cuInit.

[jakirkham]

TBH I don't think it matters much whether Numba or something else does this. We have a similar issue with cuInit slowing down import cudf for example ( https://github.com/rapidsai/cudf/issues/627 ).

[mtjrider]

It’s not numba. It’s the process of loading device code ( among other things as pointed out in https://github.com/rapidsai/cudf/issues/627 )

The cost should be linear in the number of GPUs.
It’s not (from what I can see). There’s more going on. Thanks for referencing the issue @jakirkham. I’ll have a look when I get a moment.

[pentschev]

From internal discussions with CUDA developers, the scaling should not be linear, but quadratic:

It actually scales O(n^2). [...] You need to initialize P2P between GPUs. That's the O(n^2) component.

It can only be linear if you don't need P2P/IPC, but that requires workarounds that are somewhat complicated for users:

By the way, most of cuInit time is spent initializing UVM (80+%). You can try to workaround it by having a process running in the background that initializes CUDA and then idles forever. This will keep UVM initialized.

[jakirkham]

Thanks for adding this Peter! 😀

[jakirkham]

I'm curious what the next steps here should be. It sounds like a fix is out-of-our-hands (though would certainly be welcome 🙂). Should we close? Would it help if we documented this in some way? What else would be useful here?

[quasiben]

We could document the long startup times in a FAQ or something in the docs but I am also ok with closing

[github-actions[bot]]

This issue has been marked stale due to no recent activity in the past 30d. Please close this issue if no further response or action is needed. Otherwise, please respond with a comment indicating any updates or changes to the original issue and/or confirm this issue still needs to be addressed. This issue will be marked rotten if there is no activity in the next 60d.

[beckernick]

I think we can close.