linux - CUDA 多设备问题，thrust::system_error

Question

我开发了一种使用推力 的算法。我的办公室计算机有一张支持 CUDA 的卡，其架构为:

--- General information about Device 0 Name: Quadro 2000 Compute Capability: 2.1 Clock Rate: 1251000 kHz Device Overlap: Enabled Kernel Execution Timeout: Disabled

在这台机器上，我的算法运行没有错误。但是，当尝试生成 device_vector 时，实验室机器上的干净构建会抛出讨厌的 thrust::system::system_error。两台机器都运行 RedHat 6 并且配置相同，除了多个图形卡。该实验室机器包含三个支持 CUDA 的卡，具有以下架构:

--- General information about Device 0 Name: Tesla C2050 Compute Capability: 2.0 Clock Rate: 1147000 kHz Device Overlap: Enabled Kernel Execution Timeout: Disabled

--- General information about Device 1 Name: Quadro 2000 Compute Capability: 2.1 Clock Rate: 1251000 kHz Device Overlap: Enabled Kernel Execution Timeout: Disabled

--- General information about Device 2 Name: Quadro 2000 Compute Capability: 2.1 Clock Rate: 1251000 kHz Device Overlap: Enabled Kernel Execution Timeout: Enabled`

我知道 thrust 需要针对目标架构进行编译才能工作。因此，我将 CUDA 设备设置为 1。但是，错误仍然存​​在。

作为调试措施，我在 device_vector 分配之前放置了一个 cudaGetDevice() 调用。设备被正确地声明为 1。

int device;
CUDA_CHECK_RETURN(cudaGetDevice(&device), __FILE__, __LINE__);
std::cout << "Operating on device " << device << std::endl; // <-- device 1

// copy the turns to the runtime
thrust::device_vector<MalfunctionTurn> d_turns = turns; // <-- error here

我已无计可施，无法调试它。有没有人见过这样的错误？更值得注意的是，cudaSetDevice() 是否存在我不知道的限制？我很担心，因为不同机器上的两张相同的卡不能运行相同的代码。

提前致谢。

---

编辑

编译命令行:nvcc -rdc=true -arch=sm_21 -O3 file

这是一个重现错误的最小示例:

#define DEVICE __device__
#define HOST __host__

#include <thrust/host_vector.h>
#include <thrust/device_vector.h>

template <typename T, std::size_t N>
class Container {
public:

    DEVICE HOST
    Container() {

    }

private:
    T data[N];
};

typedef Container<double, 7> double7;

template <std::size_t N = 10 >
class History {
public:

    DEVICE HOST
    History() {

    }

    DEVICE HOST
    virtual ~History() {

    }

private:

    double7 history[N];
};

int main() {

    try {

        thrust::host_vector<History<> > histories(1);
        thrust::device_vector<History<> > d_histories = histories;
    } catch (const thrust::system_error &) {
        std::cerr << "boo boo" << std::endl;
    }

    return 0;
}

Answer 1

据我所知，您的代码(在幕后)违反了 usage of classes with virtual functions 上的 CUDA 限制。 :

It is not allowed to pass as an argument to a __global__ function an object of a class with virtual functions.

如果我拿走你的代码并删除:

#include <JARSS.h>

并将其替换为:

#define HOST   __host__
#define DEVICE __device__

我可以编译它。然而，在引擎盖下，推力执行这条线:

    thrust::device_vector<History<> > d_histories = histories;

通过启动一个内核，该内核将主机上的对象作为参数复制到设备。 (例如，您可以使用 nvprof 验证内核启动。)这是 thrust 的常见行为。问题在于那些具有虚拟析构函数的对象不能以这种方式复制。

与您的问题所述相反，此代码不应在任何 CUDA 机器上正确运行。

您可以通过注释掉虚拟析构函数来“修复”您的代码。如果您需要多态行为，对实际代码的修复可能比这更复杂。应该仍然可以使用带有推力的多态对象行为，请参阅 this answer一个有效的例子。