在没有内存移动的情况下交换 CUDA 推力设备向量答案

【问题标题】：Swapping CUDA Thrust device vectors without memory movements在没有内存移动的情况下交换 CUDA 推力设备向量
【发布时间】：2017-10-05 05:19:28
【问题描述】：

如果我有两个 cudaMalloced 数组，我可以通过简单地交换相关指针来交换它们而无需内存移动。

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

我的问题是：有没有办法在不移动内存的情况下交换 CUDA Thrust device_vectors？

【问题讨论】：

thrust::vector 类有一个contiguous_storage 类型的成员，用于存储向量内容。在内部交换向量时，仅交换 contiguous_storage 的 begin() 迭代器、size 和 allocator。因此不涉及数据的内存副本。您可以在文件contiguous_storage.inl 内的swap 成员函数中检查这一点。
在赋值运算符的情况下，如果你看vector_base::operator=的代码，它使用了assign函数，它似乎执行了向量内容的完整内存复制。
@sgarizvi 感谢您的 cmets。实际上，这与@talonmies 在下面的 cmets 中指出的反对意见相同。然而，奇怪的是我在时间轴中找不到内存副本。也许thrust 使用内核来执行复制？
@sgarizvi 我已经在我的主代码中用swap 替换了显式副本，这对时间有好处。至少，swap 似乎比复制快。
我认为这是向量类相对较新的特性。在过去，我很确定交换使用复制分配并触发内存复制。

标签： cuda thrust

【解决方案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

【讨论】：

【解决方案2】：

我不知道。

没有公开的构造函数采用现有的device_ptr，并且device_vector 中的底层基向量是私有的，因此没有办法自己潜入并执行指针交换。这些是我能想到的在不触发标准复制构造函数的情况下完成这项工作的唯一方法。

编辑以添加此答案似乎是错误的。似乎最近（可能在推力 1.6 前后）的变化已经实现了一个内部指针交换交换机制，可以通过device_vector.swap() 调用。这绕过了swap() 的常用复制构造函数，并且不会触发内存传输 .

【讨论】：

如果您不知道，那么有 99.99% 的可能性是不可能的 :-) 一如既往地感谢您。
考虑一下，您也许可以通过破解一个自定义的分配器类来做一些事情，该分配器类返回另一个设备向量的内存。但是你还有很多其他问题可能无法解决
只是一个问题：d_b.swap(d_a) 是否暗示记忆运动？
我已经有一段时间没有看到这个了，但我 99% 确信推力中的所有交换都使用复制构造函数，并且会触发内存移动
我不能 -- AFAIK 这是swap github.com/thrust/thrust/blob/master/thrust/detail/swap.h 只是temp=a; a=b; b=temp;。这应该会导致内存移动