visual-c++ - CPU 和 GPU 代码与 CUDA 之间的内存布局不匹配

Question

我正在经历一个非常奇怪的情况。我有这个模板结构:

#ifdef __CUDACC__
#define __HOSTDEVICE __host__ __device__
#else
#define __HOSTDEVICE
#endif

template <typename T>
struct matrix
{
    T* ptr;
    int col_size, row_size;
    int stride;
    // some host & device methods
};

struct dummy1 {};
struct dummy2 : dummy1 {};

template <typename T>
struct a_functor : dummy2
{
    matriz<T> help_m;
    matrix<T> x, y;
    T *x_ptr, *y_ptr;
    int bsx, ind_thr;
    __HOSTDEVICE void operator()(T* __x, T* __y)
    {
        // functor code
    }
};

我已经构造了我的代码以分离 cpp 和 cu 文件，因此 a_functor 对象在 cpp 文件中创建并在内核函数中使用。问题是，在内核中执行 operator() 时，我发现了一些仅查看代码无法解释的随机行为。就像我的结构有点损坏。因此，在 a_functor 对象上调用 sizeof()，我发现:

CPU 代码(内核外的 .cpp 和 .cu):64 字节

GPU 代码(内核内部):68 字节

显然有某种不匹配破坏了整个东西。更进一步，我跟踪了 struct 参数指针和 struct 本身之间的距离 - 尝试检查生成的内存布局 - 这是我发现的:

a_functor foo;
// CPU
(char*)(&foo.help_m)    - (char*)(&foo)   = 0
(char*)(&foo.x)         - (char*)(&foo)   = 16
(char*)(&foo.y)         - (char*)(&foo)   = 32
(char*)(&foo.x_ptr)     - (char*)(&foo)   = 48
(char*)(&foo.y_ptr)     - (char*)(&foo)   = 52
(char*)(&foo.bsx)       - (char*)(&foo)   = 56
(char*)(&foo.ind_thr)   - (char*)(&foo)   = 60

// GPU - inside a_functor::operator(), in-kernel
(char*)(&this->help_m)  - (char*)(this)   = 4
(char*)(&this->x)       - (char*)(this)   = 20
(char*)(&this->y)       - (char*)(this)   = 36
(char*)(&this->x_ptr)   - (char*)(this)   = 52
(char*)(&this->y_ptr)   - (char*)(this)   = 56
(char*)(&this->bsx)     - (char*)(this)   = 60
(char*)(&this->ind_thr) - (char*)(this)   = 64

我真的不明白为什么 nvcc 为我的结构生成了这个内存布局(那 4 个字节应该是/做什么!？!)。我认为这可能是一个对齐问题，我尝试明确对齐 a_functor，但我不能，因为它是在内核中按值传递的

template <typename T, typename Str>
__global__ void mykernel(Str foo, T* src, T*dst);

当我尝试编译时，我得到

错误:无法将显式对齐太大的参数传递给 全局 win32平台上的例程

所以，为了解决这个奇怪的情况(......我确实认为这是一个 nvcc 错误)，我该怎么办？我唯一能想到的就是调整对齐并通过指针将我的结构传递给内核以避免上述错误。但是，我真的想知道:为什么内存布局不匹配？!真的没意思。。。

更多信息:我正在使用 Visual Studio 2008，在 Windows XP 32 位平台上使用 MSVC 进行编译。我安装了最新的 CUDA Toolkit 5.0.35。我的卡是 GeForce GTX 570(计算能力 2.0)。

Answer 1

从评论看来，您实际运行的代码和您发布的代码之间可能存在差异，因此在没有人能够重现问题的情况下，很难给出模糊的答案。也就是说，在 Windows 上，有些情况下结构的布局和大小在 CPU 和 GPU 之间可能不同，这些是 documented在编程指南中:

On Windows, the CUDA compiler may produce a different memory layout, compared to the host Microsoft compiler, for a C++ object of class type T that satisfies any of the following conditions:

• T has virtual functions or derives from a direct or indirect base class that has virtual functions;
• T has a direct or indirect virtual base class;
• T has multiple inheritance with more than one direct or indirect empty base class.

The size for such an object may also be different in host and device code. As long as type T is used exclusively in host or device code, the program should work correctly. Do not pass objects of type T between host and device code (e.g., as arguments to global functions or through cudaMemcpy*() calls).

第三种情况可能适用于您有一个空基类的情况，您在实际代码中是否有多重继承？