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

Question

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

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

我使用的结构如下

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 个填充字节，因此特征矩阵是 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）。这是由于 Eigen 的矩阵与 16 字节对齐，可能是因为使用了需要这种对齐的 SSE 寄存器。您的其余字段对齐到 8 字节（64 位指针和双精度）。

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

EIGEN_DONT_VECTORIZE - disables explicit vectorization when defined. Not defined by default, unless alignment is disabled by Eigen's platform test or the user defining EIGEN_DONT_ALIGN.

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

解决它的最好方法是强制对齐两种架构的结构。现在对 CUDA 不是很流利，我认为你可以在 CUDA 中使用 __align__(16)（更多 here), and using alignas(16) in C++ (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 的文档中对此没有保证）。

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

Different sizes for a struct in cpp and CUDA

c++

cuda

alignment

eigen