2

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

#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 文件，因此在 cpp 文件中创建了一个_functor 对象并在内核函数中使用。问题是，在内核中执行 operator() 时，我发现了一些我无法仅查看代码来解释的随机行为。就好像我的结构有点损坏了。因此，在 a_functor 对象上调用 sizeof()，我发现：

  • CPU 代码（内核外的 .cpp 和 .cu）：64 字节

  • GPU 代码（内核内部）：68 字节

显然有某种不匹配毁了整个东西。更进一步，我跟踪了结构参数指针和结构本身之间的距离——试图检查产生的内存布局——这就是我发现的：

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）。

4

1 回答 1

3

从评论看来，您实际运行的代码与您发布的代码之间可能存在差异，因此如果没有人能够重现问题，很难给出比模糊答案更多的答案。也就是说，在 Windows 上，结构的布局和大小在 CPU 和 GPU 之间可能会有所不同，这些都记录在编程指南中：

在 Windows 上，与主机 Microsoft 编译器相比，CUDA 编译器可能会为满足以下任何条件的类类型 T 的 C++ 对象生成不同的内存布局：

  • T 具有虚函数或派生自具有虚函数的直接或间接基类；
  • T 有一个直接或间接的虚基类；
  • T 具有多个继承，具有多个直接或间接空基类。

这种对象的大小在主机和设备代码中也可能不同。只要类型 T 专门用于主机或设备代码，程序就应该可以正常工作。不要在主机和设备代码之间传递类型 T 的对象（例如，作为全局函数的参数或通过 cudaMemcpy*() 调用）。

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

于 2013-05-20T09:21:22.143 回答