Add a CUDA stream pool

[harrism]

Closes #613. Adds a simple pool class which contains a circular buffer of cuda_stream objects with a get_stream method to get a view of one of the streams. Also adds tests and a simple benchmark.

Benchmark results. First line is getting a stream from the stream pool. Second is using the RAII rmm::cuda_stream() class which calls cudaStreamCreate() and cudaStreamDestroy().

---------------------------------------------------------------------------------
Benchmark                       Time             CPU   Iterations UserCounters...
---------------------------------------------------------------------------------
BM_StreamPoolGetStream      0.240 us        0.240 us      2918230 items_per_second=4.16296M/s
BM_CudaStreamClass           3.01 us         2.90 us       249073 items_per_second=344.521k/s

TODO in a follow up:

• Cython/Python bindings
• Stream priorities and non-blocking streams, once support is added to rmm::cuda_stream().

[kkraus14]

cmake / conda lgtm

[rongou]

Should this be private?

[harrism]

It's a default parameter to a public constructor, so I assumed it must be public...

[rongou]

The constructor has access to private fields. There is no reason for any caller to change this default, right?

[jrhemstad]

From what I remember, Cython doesn't really work well with default args, so they usually need access to the definition of what is being used as the default.

[harrism]

This is constexpr, so a caller cannot change it. I made it public because I thought it might be useful for users to be able to access the default size at run time.

[mzient]

This makes this pool unusable for DALI - unwittingly submitting new (and potentially urgent) work on a stream that already has hundreds of milliseconds worth of work scheduled is not an option.

[harrism]

You do know that there are hardware limits to concurrency, right? See table 15 in the CUDA programming guide. The current high end GPUs have at most 128 concurrent execution contexts, so using more than 128 streams is redundant -- you will get false dependencies even with more streams.

The size of the pool is configurable. Set it to 128 if you want 128 streams.

We can also discuss what the best default size is before we merge this. I chose 16 because that is the minimum maximum concurrency of GPUs that RAPIDS supports today. I didn't want to create hundreds of streams if the user doesn't need them.

We can also consider a pool that can grow, with a parameterized maximum size.

[jrhemstad]

As we make the pool more complex, we can also have a notion of normal streams and high priority streams, e.g., get_stream and get_priority_stream.

[mzient]

I guess that we have a very different idea of a stream pool then.

My idea (which I'm going to pursue in DALI) was to have an pool where you take streams from the pool and explicitly return them. That's more or less equivalent to create/destroy, but without the ugly handle reuse.

The pool would have two lists:

• ready list
• free list

When a client requests a stream, it's taken from the ready list. When it's returned to the pool, it's placed in the free list. Then there are several options:

1. Add a callback to the stream, so that it's returned to the ready list when all work is complete.
2. Call cudaStreamQuery periodically in a worker thread and collect ready streams.
3. Call cudaStreamQuery on the streams in the free list when there's no ready stream and return the ready ones to the ready list (this is the least preferred option, as it introduces a potentially large and unpredictable overhead).

In any case, if there are no ready streams and there are no ready streams, a new stream is created and returned.

[harrism]

I did consider a design with get_stream and return_stream, but I think this is early optimization. The present design is extremely simple, and will likely provide good performance for many use cases. If and when we find that the performance is insufficient, we can add complexity as needed. Certainly, a design that requires worker threads is much more complex, and I will not add that complexity to RMM without a real-world benchmark to drive it.