[FEA]: What to do with `cuda::device` namespace

[davebayer]

Is this a duplicate?

• [x] I confirmed there appear to be no duplicate issues for this request and that I agree to the Code of Conduct

Area

libcu++

Is your feature request related to a problem? Please describe.

From libcu++ documentation:

• cuda:: / <cuda/*>: Conforming extensions to the Standard Library that work in host and device code.
• cuda::device / <cuda/device/*>: Conforming extensions to the Standard Library that work only in device code.

(Unfortunately?) this is not true. There are features (like cuda::for_each_canceled_block) which are device only and features (like cuda::device::warp_shuffle) which are not in <cuda/device/*> path.

Describe the solution you'd like

My idea is to make cuda::device an inline namespace and put everything to <cuda/*> path (as it currently is).

We could put the device only features from cuda to cuda::device namespace without the need to rewrite any of the existing code. It would be a breaking change, so I would suggest to import the symbols from cuda to cuda::device namespace for now.

The cuda::device namespace would still keep it's importance, the user could e. g. import all of the features from cuda::device and be sure only __device__ symbols are available, like:

namespace cd = cuda::device;

void kernel(...)
{
  auto res1 = cd::warp_shuffle(...);
  // auto res2 = cd::some_host_only_function_from_namespace_cuda(...); // not found
}

It also brings me to an idea that we could import __host__ __device__ features from cuda namespace to cuda::device.

In the future we could do something similar with cuda::host namespace where would be all of the __host__ only features.

One more idea: we could move cuda::ptx to cuda::device::ptx (which makes sense), it wouldn't require changes in user's code (but would be a breaking change.

Describe alternatives you've considered

We can put everything in cuda namespace and create aliases in cuda::device namespace, which would also make sense. But that would mean cuda::device wouldn't be inline namespace anymore.

Maybe it would be a better solution?

But I don't know, maybe it would be better to just get rid of the cuda::device namespace completely and have everything in cuda namespace.

Additional context

No response

[pciolkosz]

I agree with the solution. From internal discussion cuda::device namespace was added at a time where everything in cuda namespace was host/device and we needed a new namespace for device-only APIs. But now we have a bunch of host-only runtime APIs along with the migration of APIs from Thust/CUB into cuda and the cuda::device namespace does not provide the semantic benefit anymore.

We should stop putting new APIs into cuda::device namespace and provide aliases to all cuda::device APIs in cuda namespace (at proposed in the description).

[miscco]

I believe organization is a crucial part of a library. Maybe the cuda::device is too wide, but I definitely want to see other namespaces used to better group features and not cram everything into cuda::