不需要内存移动的交换CUDA Thrust设备向量

Question

如果我有两个使用cudaMalloc分配的数组，我可以通过交换相关指针而无需内存移动来交换它们。

如果我有两个CUDA Thrust设备向量，例如d_a和d_b，我可以使用第三个临时向量d_c来交换它们，但这将需要内存移动。

我的问题是：是否有一种方法可以在不进行内存移动的情况下交换CUDA Thrust设备向量？

- Vitality

---

thrust::vector 类有一个 contiguous_storage 类型的成员，用于存储向量内容。当向量交换时，内部仅交换 contiguous_storage 的 begin() 迭代器、size 和 allocator。因此，不涉及数据的内存复制。您可以在文件 contiguous_storage.inl 中的 swap 成员函数中检查这一点。 - sgarizvi

在赋值运算符的情况下，如果您查看vector_base::operator=的代码，它使用了assign函数，该函数似乎执行向量内容的完整内存复制。 - sgarizvi

@sgarizvi 谢谢你的评论。实际上，这正是@talonmies在下面的评论中指出的同样问题。然而，奇怪的是我在时间线中找不到内存拷贝的记录。也许thrust使用了一个内核来执行拷贝操作？ - Vitality

我认为这是向量类的一个相对较新的特性。在旧时代，我很确定swap使用了复制赋值，并触发了内存复制。 - talonmies

@RobertCrovella，你的评论非常贴合这篇帖子。如果你有使用向量引用的替代解决方案，我相信提供一个带例子的答案对整个社区都会有益处。还有一个问题：device_vector.swap()在底层实际上是交换向量引用吗？ - Vitality

Answer 1

看起来device_vector.swap()避免了内存移动。

实际上，考虑以下代码：

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>

#include <thrust\device_vector.h>

void printDeviceVector(thrust::device_vector<int> &d_a) {

    for (int k = 0; k < d_a.size(); k++) {

        int temp = d_a[k];
        printf("%i\n", temp);

    }

}

int main()
{
    const int N = 10;

    thrust::device_vector<int> d_a(N, 1);
    thrust::device_vector<int> d_b(N, 2);

    // --- Original
    printf("Original device vector d_a\n");
    printDeviceVector(d_a);
    printf("Original device vector d_b\n");
    printDeviceVector(d_b);

    d_b.swap(d_a);

    // --- Original
    printf("Final device vector d_a\n");
    printDeviceVector(d_a);
    printf("Final device vector d_b\n");
    printDeviceVector(d_b);

    d_a.clear();
    thrust::device_vector<int>().swap(d_a); 
    d_b.clear();
    thrust::device_vector<int>().swap(d_b);

    cudaDeviceReset();

    return 0;
}

使用

    d_b.swap(d_a);

如果我们对它进行分析，我们会在时间轴上看不到任何设备之间的内存移动：

另一方面, 如果我们将 d_b.swap(d_a) 更改为

d_b = d_a;

然后，在时间轴上出现设备到设备的移动：

最后，d_b.swap(d_a)比d_b = d_a更有利于时序。对于N = 33554432，时间如下：

d_b.swap(d_a)     0.001152ms
d_b = d_a         3.181824ms

- Vitality

Answer 2

据我所知没有这种方法。

没有暴露接受现有device_ptr的构造函数，而device_vector中的基础向量是私有的，因此没有办法深入其中并执行指针交换。这些可能是在不触发标准复制构造函数的情况下使其工作的唯一方式。

---

编辑添加：似乎这个答案是错误的。最近（大约在Thrust 1.6左右）的更改实现了一个内部的指针交换交换机制，可以通过device_vector.swap()调用。这会绕过swap()的通常复制构造函数惯用语，并且不会触发内存传输。

- talonmies

---

如果你不知道，那么99.99%的可能性是不可能的 :-) 谢谢，一如既往。 - Vitality

我已经使用了 d_b.swap(d_a)，并且在时间轴上没有看到设备之间的复制（请参见编辑）。 - Vitality

我已经有一段时间没有看过这个了，但我99%确定thrust中的所有交换都使用了复制构造函数，并且这会触发内存移动。 - talonmies

我不行 - 据我所知这是 swap https://github.com/thrust/thrust/blob/master/thrust/detail/swap.h 只需要 temp=a; a=b; b=temp;。那应该会导致内存移动。 - talonmies

从时间上看，d_b.swap(d_a)明显比d_b = d_a更有优势。 - Vitality