NVIDIA / NVIDIA/cccl

`thrust::mr::unsynchronized_pool_resource` fails to compile with `universal_host_pinned_memory_resource` Upstream

Open
#798 1 comment 0 reactions 1 assignee Claimed by @ericniebler View on GitHub
thrust
Dominant language
C++
Stars
2.5k
Forks
487
Avg merge
2d 7h
Merged PRs (30d)
296

Description

```
#include
#include
#include

using pinned_t = thrust::universal_host_pinned_memory_resource;
using pinned_pool_t = thrust::mr::unsynchronized_pool_resource;

// Constructs a resource that allocates pinned memory
pinned_t pinned_mr;

// Constructs a pool that uses the pinned resource as the upstream
pinned_pool_t pinned_pool{&pinned_mr};
```

Fails to compile: https://godbolt.org/z/KaovePse1

It complains about using a private type in the template arguments of a `__global__` function.

```
/thrust/system/cuda/detail/core/agent_launcher.h(933): error: A type that is defined inside a class and has private or protected access ("thrust::mr::unsynchronized_pool_resource::pool") cannot be used in the template argument type of a __global__ function template instantiation, unless the class is local to a __device__ or __global__ function
detected during:
instantiation of "thrust::cuda_cub::core::_kernel_agent" based on template
```

The private type in question is the [`unsynchornized_pool_resource::pool` type](https://github.com/NVIDIA/thrust/blob/583ab493073e6d2681e99fadea7ec5469f0fef3f/thrust/mr/pool.h#L191). Which is used as the value_type of a `host_vector`, [`m_pools`](https://github.com/NVIDIA/thrust/blob/583ab493073e6d2681e99fadea7ec5469f0fef3f/thrust/mr/pool.h#L197-L207). The compile error comes from calling [`resize`](https://github.com/NVIDIA/thrust/blob/583ab493073e6d2681e99fadea7ec5469f0fef3f/thrust/mr/pool.h#L116) on this vector which ultimately dispatches to `thrust::cuda_cub::__parallel_for::parallel_for`.

It is somewhat surprising to me that this ends up dispatching to the `cuda_cub` code path.

I've made a more minimal reproducer that shows the cause is isolated to using `universal_host_pinned_memory_resource` as the upstream resource for a `thrust::mr::allocator`.

https://godbolt.org/z/n64jr3GeY

I suspect the `universal_host_pinned_memory_resource` is causing dispatch to the `cuda_cub` path via ADL since `thrust::universal_host_pinned_memory_resource` is an alias to [`thrust::cuda::universal_host_pinned_memory_resource`](https://github.com/NVIDIA/thrust/blob/c72c7018a3a9e8faf6bbab31168ed9d86e1fb426/thrust/system/cuda/memory_resource.h#L122).

This may be intended behavior that using `universal_host_pinned_memory_resource` dispatches to the `cuda_cub` code paths since "universal" memory is expected to be device accessible. In which case, the easiest solution is probably to just move the [`pool`](https://github.com/NVIDIA/thrust/blob/583ab493073e6d2681e99fadea7ec5469f0fef3f/thrust/mr/pool.h#L191) nested private type to be non-private or just non-nested.

Contributor guide

Open the contributing guide

Assessment

This issue has not been assessed yet.

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.