c++ 和 CUDA 中结构的不同大小

Question

我在使用内核时遇到了一些问题，该内核使用了我在 c++ 中定义的一些结构。 cuda-memcheck 给我的错误是对齐问题。

我尝试使用的结构包含一些指针，我猜这给我带来了问题。我已经在 .cu 文件和内核的主机函数中打印了 C++ 端和 CUDA 端的结构大小以控制台。这给出了不同的结果，这解释了我看到的问题，但我不确定它为什么会发生或如何解决它。

我使用的结构如下

struct Node {};
struct S
{
    Node *node0;
    Node *node1;
    Node *node2;
    double p0;
    double p1;
    double p2;
    double p3;

    Eigen::Matrix<double, 3, 2> f1;
    Eigen::Matrix<double, 3, 2> f2;
}

在 C++ 中它的大小为 160 字节，但在 CUDA 中为 152 字节。为了传输数据，我分配了一个 CUDA 端缓冲区并执行 cudaMemcpy

std::vector<S> someVector; // Consider it exists
S *ptr;
cudaMalloc(&ptr, sizeof(S) * someVector.size());
cudaMemcpy(ptr, someVector.data(), sizeof(S)*someVector.size(), cudaMemcpyHostToDevice);

我猜这是错误的，因为 CUDA 和 C++ 中的大小不同。

当我尝试在内核中访问 S::node0、S::node1 或 S::node3 时，我会收到未对齐的访问错误。

所以我对这个问题有三个问题：

• 为什么尺寸不同？
• 我应该如何更改代码或执行复制以使其正常工作？
• 我应该有一个 CUDA 侧结构并执行特殊复制吗？

编辑： 感谢接受的答案，我能够理解我遇到的问题的原因。 Eigen 尽可能使用 vectorizacion 并为此请求 16 字节对齐。当 Eigen 对象大小是 16 字节的倍数时，将启用矢量化。在我的特殊情况下，两个Eigen::Matrix<double, 3,2> 对矢量化有效。

但是，在 CUDA 中，Eigen 不要求 16 字节对齐。

由于我的结构有 4 个双精度数和 3 个指针，计算为 56 个字节，这不是 16 的倍数，因此在 CPU 中它必须添加 8 个填充字节，因此 Eigen 矩阵是 16 个字节对齐的。在 CUDA 中不会发生这种情况，因此大小不同。

我实现的解决方案是手动添加 8 个填充字节，因此 CPU 和 CUDA 中的结构相同。这解决了问题，并且不需要禁用矢量化。我发现另一个可行的解决方案是将Eigen::Matrix<double,3,2> 更改为2 Eigen::Matrix<double,3,1>。 Eigen::Matrix<double,3,1> 不满足向量化的要求，因此不需要在 CPU 中添加 8 个填充字节。

Answer 1

这种差异是由于 Eigen 在 C++ 和 CUDA 中请求内存对齐的方式。

在 C++ 中，S 与 16 字节对齐（您可以检查 alignof(S) == 16）。这是由于本征矩阵与 16 字节对齐，可能是因为使用了需要这种对齐的 SSE 寄存器。其余字段对齐为 8 字节（64 位指针和双精度）。

在Eigen/Core 头文件中，EIGEN_DONT_VECTORIZE 指令为 CUDA 启用。检查documentation时：

EIGEN_DONT_VECTORIZE - 定义时禁用显式矢量化。默认情况下未定义，除非 Eigen 的平台测试或用户定义 EIGEN_DONT_ALIGN 禁用对齐。

这基本上意味着 Eigen 矩阵在 CUDA 中没有特殊对齐，因此它们与元素类型对齐，在您的情况下为 double，导致矩阵对齐 8 字节，因此整个结构对齐。

解决它的最佳方法是强制对齐两种架构的结构。现在对 CUDA 不太流利，我认为您可以在 CUDA 中使用 __align__(16)（更多 here），并在 C++ 中使用 alignas(16)（since C++11）。如果您共享两种语言的声明，则可以定义一个宏以使用正确的运算符：

#ifdef __CUDACC__
# define MY_ALIGN(x) __align__(x)
#else
# define MY_ALIGN(x) alignas(x)
#endif

struct MY_ALIGN(16) S {
  // ...
};

无论如何，请注意此类低级副本，因为 Eigen 在 CUDA 中的实现可能与 C++ 中的不同（在 Eigen 的文档中对此没有任何保证）。