Fix Thrust/CUB Linkage Issues

[gevtushenko]

Description

closes #166

Using CUB/Thrust in shared libraries is a known source of issues. For a while, the solution to these issues consisted of wrapping CUB/Thrust namespaces in THRUST_CUB_WRAPPED_NAMESPACE macro so that different shared libraries have different symbol names. This solution has poor discoverability, since issues present themselves in forms of segmentation faults, hangs, wrong results, etc. To eliminate the symbol visibility issues on our end, the following design is proposed:

1. Hide kernel launchers: it’s important that kernel launchers like the Thrust triple_chevron always reside in the same library as the CUB/Thrust host API using these kernel launchers.

2. Hide all kernels: it’s important that CUB/Thrust kernels always reside in the same library as the CUB/Thrust host API using these kernels.

3. Incorporate GPU architectures into symbol names: it’s important that kernels compiled for a given GPU architecture are always used by the CUB/Thrust host API compiled for that architecture.

Applying mentioned recommendations to the CCCL codebase involves the following changes:

1. thrust::cuda_cub::launcher::triple_chevron is annotated with _LIBCUDACXX_HIDDEN

2. instead of annotating kernels with __global__ we introduce {CUB,THRUST}_DETAIL_KERNEL_ATTRIBUTES macro that’s equivalent to __global__ _LIBCUDACXX_HIDDEN and annotate every CUB/Thrust kernel with it

3. Thrust symbols are placed inside an inline namespace containing the set of GPU architectures for which the TU is being compiled when CUDA is used as the device system.

Potentially breaking changes:

• In Linux builds, CUB/Thrust kernels used to have default visibility. This PR changes that, so direct references to CUB/Thrust kernels in shared libraries will be broken. Since CUB/Thrust kernels are considered an implementation detail, this isn't a blocker for the PR.

• Thrust ABI is broken when the set of GPU architectures doesn’t match. If a library exposes a function taking, say, device vector and is compiled for SM80, any code compiled for different architectures won’t be able to link against that library. The behavior can be opted out by defining THRUST_DISABLE_ABI_NAMESPACE or providing wrapped namespace.

Checklist

• New or existing tests cover these changes.
• The documentation is up to date with these changes.

[miscco]

Minor nits

[ahendriksen]

This is a really clean PR with great documentation! I added some questions and one documentation suggestion.