如何防止推力的 device_vector 复制到设备

如何解决如何防止推力的 device_vector 复制到设备

所以我有一个帮助类（创造性地命名为“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",(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",(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");
}

抽象/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),cachedSize(size) {}

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

    __host__ explicit BetterVector(std::vector<T> vec) : vector(vec),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
    }
}

解决方法

这里的中心问题似乎是，通过使用将项放置在 union 中的技巧，以便不会自动调用构造函数和析构函数，您阻止了 vector 的正确初始化，并且您的构造函数( s) 没有做到这一点。

1. 对于测试代码的第一部分，通过 CUDA 内核调用，这里有一个与此特定观察相关的构造函数：

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

   我的声明是 vector(vec) 没有正确构建 vector。我怀疑这与 union 的使用有关，其中未调用定义的构造函数（并且可能使用复制初始化程序，但我不清楚）。

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

可以通过所谓的“placement new”调用构造函数

2. 正如评论中提到的，这个复制操作不可能是正确的，它只是复制了 1 个字节：

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

3. printf 的设备版本似乎无法理解格式说明符 %zu，我将其替换为 %lu

4. 这本身不是问题，但可能值得指出的是这行代码：

   abstract::BetterVector<int> vector = *ptr;
   

   在每个线程中生成一个单独的BetterVector对象，从传递给内核的对象初始化。

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

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

$ 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(),cachedSize(0) {}

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

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

//    __host__ explicit BetterVector(std::vector<T> vec) : vector(vec),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",(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",(int) vector[i]);
    }
    printf("\n");

    abstract::BetterVector<int>* devVectorPtr;
    cudaMalloc(&devVectorPtr,sizeof(abstract::BetterVector<int>));
    cudaMemcpy(devVectorPtr,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
$

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