[BUG]: Suboptimal swap performance on universal vectors

[gevtushenko]

Is this a duplicate?

• I confirmed there appear to be no duplicate issues for this bug and that I agree to the Code of Conduct

Type of Bug

Performance

Component

Thrust

Describe the bug

Swap of universal vectors is unnecessary slow and occupies more memory than it should.
Device and host vectors have swap overloads in thrust:: namespace.
These overloads call efficient .swap member function:

cccl/thrust/thrust/device_vector.h

Lines 539 to 543 in 80031e2

	template <typename T, typename Alloc>
	void swap(device_vector<T, Alloc>& a, device_vector<T, Alloc>& b)
	{
	a.swap(b);
	}

Universal and universal host pinned vectors are missing such an overload and the vector_base they alias to has swap overload in thrust::details, so it’s unreachable.

How to Reproduce

template <template <typename> class Vector>
void swap_time(const char* str) {
    Vector<char> a(1 << 30, 'a');
    Vector<char> b(1 << 30, 'b');

    auto begin = std::chrono::high_resolution_clock::now();
    thrust::swap(a, b);
    // a.swap(b);
    auto end = std::chrono::high_resolution_clock::now();
    std::cout << str << " swap time: " << std::chrono::duration<double>(end - begin).count() << " s" << std::endl;
    std::cout << "a: " << static_cast<char>(a[0]) << "; "
              << "b: " << static_cast<char>(b[0]) << std::endl;
}

int main() {
    swap_time<thrust::device_vector>("device");
    swap_time<thrust::host_vector>("host");
    swap_time<thrust::universal_vector>("universal");
    swap_time<thrust::universal_host_pinned_vector>("universal_host_pinned");
}

Expected behavior

• Swap of universal host pinned vector is 30 million times faster
• Swap of universal vector is 4 million times faster

Reproduction link

No response

Operating System

No response

nvidia-smi output

No response

NVCC version

No response

[bernhardmgruber]

While I agree that we should fix the performance issue, I have troubles with how we swap vectors. The established practice in C++ on how to swap any two objects is:

using std::swap;
swap(a, b);

You make std::swap available and then ADL-call swap, so if a better (more specialized) implementation is available, it will be chosen. Alternatively, you can use using cuda::std::swap;. Also using thrust::swap; works here. All these versions don't show the performance problem.

Only calling thrust::swap(a, b); directly is slow for universal vector, because the semantics are: "use the thrust swap algorithm" and not "find the best swap algorithm, while [std|cuda::std|thrust]::swap is available as fallback".

[bernhardmgruber]

There are two ways to address this.

1. Turn the aliases for thrust::universal_vector etc. into dedicated types and add specializations inside namespace thrust. All swap optimizations would only work for thrust::swap. All other swaps (from cuda::std or std) would do 3 moves.
2. Make thrust::swap an alias to cuda::std::swap and add a bunch of fixes. All ways to swap would now take the optimal path. This requires some fixes to libcu++ as well, including this one: [BUG]: cuda::std::swap cannot ADL-two-step-swap a type from std #2984.

I strongly prefer 2.

[bernhardmgruber]

I implemented solution 2. in #2985, which is independent of #2984.