时间戳

首页 > 解决方案 > cpp 和 CUDA 中结构的不同大小

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

问题描述

我在使用内核时遇到了一些问题，该内核使用了我在 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在内核中，我就会得到一个未对齐的访问错误。

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

编辑： 感谢接受的答案，我能够理解我遇到的问题的原因。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 个填充字节。

标签： c++cudaalignmenteigen

解决方案

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

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

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

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

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

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

推荐阅读