如何防止将 thrust 的 device_vector 复制到设备

Question

所以我有一个帮助程序 class（创造性地命名为“BetterVector”），它被设计为在主机和设备之间来回传递，其大部分功能都可以从任何一方访问（一个重大缺陷device_vector)。但是，内核因非描述性分配错误而失败。

从堆栈跟踪来看，它似乎有时在复制构造函数上触发，有时在析构函数上触发，我不完全确定它为什么会改变。我认为是 device_vector 数据成员具有仅主机构造函数和解构函数，我使用 following post 来利用联合来防止调用这些函数，但问题仍然存在。如果您有任何建议，我们将不胜感激。

main.cu 测试文件：

#include <abstract/BetterVector.cuh>

struct thrust_functor {
    abstract::BetterVector<int> vector;

    explicit thrust_functor(const abstract::BetterVector<int> &vector) : vector(vector) {}

    __host__ void operator()(int i) {
        printf("Thrust functor index %d: %d\n", i, (int) vector[i]);
    }
};

__global__ void baseCudaPrint(abstract::BetterVector<int>* ptr) {
    const size_t i = blockIdx.x * blockDim.x + threadIdx.x;
    abstract::BetterVector<int> vector = *ptr;
    printf("Cuda kernel index %zu: %d\n", i, (int) vector[i]);
}


int main() {
    abstract::BetterVector<int> vector({1, 2, 3, 4});
    for (int i = 0; i < 4; i++) {
        printf("Host index %d: %d\n", i, (int) vector[i]);
    }
    printf("\n");

    abstract::BetterVector<int>* devVectorPtr;
    cudaMalloc(&devVectorPtr, sizeof(abstract::BetterVector<int>));
    cudaMemcpy(devVectorPtr, &vector, 1, cudaMemcpyHostToDevice);
    baseCudaPrint<<<1, vector.size()>>>(devVectorPtr);
    cudaDeviceSynchronize();
    cudaFree(devVectorPtr);
    printf("\n");

    thrust::counting_iterator<int> first(0);
    thrust::counting_iterator<int> last = first + vector.size();
    thrust::for_each(thrust::host, first, last, thrust_functor(vector));
    cudaDeviceSynchronize();
    printf("\n");
}

abstract/BetterVector.cuh:

#include <thrust/device_vector.h>
#include <thrust/device_ptr.h>
#include <thrust/functional.h>

namespace abstract {
template<typename T>
    struct equal_to : public thrust::unary_function<T, bool> {
        T lhs;

        __device__ __host__ explicit equal_to(T lhs) : lhs(lhs) {}

        __device__ __host__ bool operator()(T rhs) {
            return lhs == rhs;
        }
    };
template<typename T, typename VecType = thrust::device_vector<T>>
class BetterVector {
protected:
    typename VecType::pointer raw;
    size_t cachedSize;
    union {
        VecType vector;
    };

public:

    __host__ BetterVector() : vector(), raw(vector.data()), cachedSize(0) {}

    __host__ explicit BetterVector(size_t size) : vector(size), raw(vector.data()), cachedSize(size) {}

    __host__ explicit BetterVector(VecType vec) : vector(vec), raw(vector.data()), cachedSize(vec.size()) {}

    __host__ explicit BetterVector(std::vector<T> vec) : vector(vec), raw(vector.data()), cachedSize(vec.size()) {}

    __host__ __device__ BetterVector(const BetterVector &otherVec) :
#ifndef __CUDA_ARCH__
            vector(otherVec.vector),
#endif
            cachedSize(otherVec.cachedSize), raw(otherVec.raw) {}


    __host__ __device__ virtual ~BetterVector() {
#ifndef __CUDA_ARCH__
        vector.~VecType();
#endif
    }

    __host__ __device__ typename VecType::const_reference operator[](size_t index) const {
#ifdef __CUDA_ARCH__
        return raw[index];
#else
        return vector[index];
#endif
    }

    __host__ __device__ size_t size() const {
#ifdef __CUDA_ARCH__
        return cachedSize;
#else
        return vector.size();
#endif
    }
}

Answer 1

这里的中心问题似乎是，通过使用将项目放置在 union 中的技巧，这样构造函数和析构函数就不会被自动调用，您已经阻止了 vector 的正确初始化，并且您的构造函数没有完成那个。

对于测试代码的第一部分，通过 CUDA 内核调用，有一个对这个特定观察感兴趣的构造函数，这里：
```
__host__ explicit BetterVector(std::vector<T> vec) : vector(vec), raw(vector.data()), cachedSize(vec.size()) {}
```
我的说法是 vector(vec) 没有正确构建 vector。我怀疑这与 union 的使用有关，其中未调用定义的构造函数（并且可能使用了复制初始化程序，但这对我来说并不清楚）。

无论如何，我们可以使用 the link you provided 的线索来解决这个问题：

constructor can be called through so called "placement new"

如评论中所说，这个复制操作不可能是正确的，它只是复制1个字节：

cudaMemcpy(devVectorPtr, &vector, 1, cudaMemcpyHostToDevice);
                                  ^

printf的设备版本似乎不理解格式说明符%zu，我将其替换为%lu
这本身不是问题，但可能值得指出的是这行代码：
```
abstract::BetterVector<int> vector = *ptr;
```
在每个线程中产生一个单独的BetterVector对象，从传递给内核的对象初始化。

这一级别的“修复”将使您到达 main 代码在 CUDA 内核启动期间正确显示运行的地步。然而，此后的推力代码仍然存在我无法解决的问题。由于您的代码设计（在推力主机路径中使用 device_vector，即使它是主机函数，如果正常工作，对 for_each 的调用应该会“在后台”生成 3 个内核调用。很奇怪。 ) 无论如何我无法为你解决这个问题，但我可以说这 3 个内核调用每个都会触发对你的 __host__ __device__ 构造函数（以及相应的析构函数）的调用，这并不让我感到惊讶. Thrust 通过按值传递将一个 BetterVector 对象传递给每个内核启动，这样做会触发一个 constructor/destructor 序列以支持按值传递操作。因此，考虑到我们必须跳过重重障碍才能让之前的构造函数“工作”，因此该序列中可能存在问题。但我一直无法查明问题所在。

无论如何，这是一个包含上述项目的示例：

$ cat t37.cu #include <thrust/device_vector.h> #include <thrust/device_ptr.h> #include <thrust/functional.h> namespace abstract { template<typename T> struct equal_to : public thrust::unary_function<T, bool> { T lhs; __device__ __host__ explicit equal_to(T lhs) : lhs(lhs) {} __device__ __host__ bool operator()(T rhs) { return lhs == rhs; } }; template<typename T, typename VecType = thrust::device_vector<T>> class BetterVector { protected: typename VecType::pointer raw; size_t cachedSize; union { VecType vector; }; public: __host__ BetterVector() : vector(), raw(vector.data()), cachedSize(0) {} __host__ explicit BetterVector(size_t size) : vector(size), raw(vector.data()), cachedSize(size) {} __host__ explicit BetterVector(VecType vec) : vector(vec), raw(vector.data()), cachedSize(vec.size()) {} // __host__ explicit BetterVector(std::vector<T> vec) : vector(vec), raw(vector.data()), cachedSize(vec.size()) {} __host__ explicit BetterVector(std::vector<T> vec) : cachedSize(vec.size()) { new (&vector) VecType(vec); raw = vector.data();} __host__ __device__ BetterVector(const BetterVector &otherVec) : #ifndef __CUDA_ARCH__ vector(otherVec.vector), #endif cachedSize(otherVec.cachedSize), raw(otherVec.raw) {} __host__ __device__ virtual ~BetterVector() { #ifndef __CUDA_ARCH__ vector.~VecType(); #endif } __host__ __device__ typename VecType::const_reference operator[](size_t index) const { #ifdef __CUDA_ARCH__ return raw[index]; #else return vector[index]; #endif } __host__ __device__ size_t size() const { #ifdef __CUDA_ARCH__ return cachedSize; #else return vector.size(); #endif } }; } struct thrust_functor { abstract::BetterVector<int> vector; explicit thrust_functor(const abstract::BetterVector<int> &vector) : vector(vector) {} __host__ void operator()(int i) { printf("Thrust functor index %d: %d\n", i, (int) vector[i]); } }; __global__ void baseCudaPrint(abstract::BetterVector<int>* ptr) { const size_t i = blockIdx.x * blockDim.x + threadIdx.x; abstract::BetterVector<int> vector = *ptr; printf("Cuda kernel index %lu: %d\n", i, (int) vector[i]); } int main() { // these indented lines mysteriously "fix" the thrust problems thrust::device_vector<int> x1(4,1); thrust::device_vector<int> x2(x1); // abstract::BetterVector<int> vector({1, 2, 3, 4}); for (int i = 0; i < 4; i++) { printf("Host index %d: %d\n", i, (int) vector[i]); } printf("\n"); abstract::BetterVector<int>* devVectorPtr; cudaMalloc(&devVectorPtr, sizeof(abstract::BetterVector<int>)); cudaMemcpy(devVectorPtr, &vector, sizeof(abstract::BetterVector<int>), cudaMemcpyHostToDevice); baseCudaPrint<<<1, vector.size()>>>(devVectorPtr); cudaDeviceSynchronize(); cudaFree(devVectorPtr); printf("\n"); thrust::counting_iterator<int> first(0); thrust::counting_iterator<int> last = first + vector.size(); thrust::for_each(thrust::host, first, last, thrust_functor(vector)); cudaDeviceSynchronize(); printf("\n"); } $ nvcc -std=c++14 t37.cu -o t37 -lineinfo -arch=sm_70 $ cuda-memcheck ./t37 ========= CUDA-MEMCHECK Host index 0: 1 Host index 1: 2 Host index 2: 3 Host index 3: 4 Cuda kernel index 0: 1 Cuda kernel index 1: 2 Cuda kernel index 2: 3 Cuda kernel index 3: 4 Thrust functor index 0: 1 Thrust functor index 1: 2 Thrust functor index 2: 3 Thrust functor index 3: 4 ========= ERROR SUMMARY: 0 errors $

我还会添加一个主观评论，我认为这种代码设计会很麻烦（以防万一还不清楚），我建议您考虑另一种“通用”向量的路径。仅举一个例子，您使用 thrust-provided [] 运算符允许通过主机代码访问的方法将非常慢。这将为以这种方式访问的每个项目调用单独的 cudaMemcpy 。无论如何，祝你好运！

如何防止将 thrust 的 device_vector 复制到设备

How to prevent the copy of thrust's device_vector to device

c++

cuda

thrust