Extend device_scalar to optionally use pinned bounce buffer

[vuule]

Description

Depends on #16945

Added cudf::detail::device_scalar, derived from rmm::device_scalar. The new class overrides function members that perform copies between host and device. New implementation uses a cudf::detail::host_vector as a bounce buffer to avoid performing a pageable copy.

Replaced rmm::device_scalar with cudf::detail::device_scalar across libcudf.

Checklist

• I am familiar with the Contributing Guidelines.
• New or existing tests cover these changes.
• The documentation is up to date with these changes.

[vuule]

bonus feature from having a bounce buffer - we don't need to worry about the value lifetime. rmm::device_scalar prohibits passing an rvalue here, but we don't need to.

[davidwendt]

Why the global change? Could we still use rmm::device_scalar when not using a pinned bounce buffer?

[vuule]

Why the global change? Could we still use rmm::device_scalar when not using a pinned bounce buffer?

(discussed offline, posting here for viz)
We're trying to avoid any pageable copies because they cause copy engine contention in multi-threaded use cases. Using pinned memory in the bounce_buffer would make this a pinned copy, which is slightly better in general. In addition, some users can choose to perform small copies using a kernel to further avoid the copy engine in their multi-threaded applications.

[vyasr]

The new class seems fine, but do we really need to use it everywhere? Does this mean that any usage of rmm::device_scalar in cudf is now forbidden because it could introduce unexpected performance overheads? If so, should we include a pre-commit hook or something to that effect to enforce that? Also, if this implication is correct @davidwendt may want to weigh in.

The title says "optionally use pinned bounce buffer", but this usage looks unconditional. Is the optionality encoded in the host vector, or is it simply no longer optional?

[vuule]

The title says "optionally use pinned bounce buffer", but this usage looks unconditional. Is the optionality encoded in the host vector, or is it simply no longer optional?

The use of a bounce buffer is unconditional, but host_vector, which is used as a bounce buffer, optionally uses pinned memory.
There's no perf impact when the bounce buffer is pageable. I'll evaluate impact from pinned memory once we're closer to eliminating all "unconditional" pageable memory use (for memory that ends up on the GPU, or copied to from the GPU).

[vuule]

The new class seems fine, but do we really need to use it everywhere? Does this mean that any usage of rmm::device_scalar in cudf is now forbidden because it could introduce unexpected performance overheads?

Pretty much. My aim is to stop using rmm::device_scalar outside of public APIs. Same goes for std::vector and thrust::host_vector (if copied to/from GPU). I don't know if we can automate this.

[vyasr]

Approving based on some extensive offline discussion. We're going to move frward with this and see how this type works in libcudf, then consider upstreaming it to rmm if it's generalizable, and if it's not then looking into developing clear guidelines for when it should be used in cudf.

[davidwendt]

Curious why the move ctor required code but this did not?

[vuule]

Default implementations should be fine in both cases. Compiled fine on 12.5 🤷
I suspect it's an 11.8 compiler bug, but really didn't want to dig into it, with a handy workaround available.

[vuule]

/merge