Разные размеры для структуры в cpp и 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-байтовое выравнивание для этого.Векторизация включается, когда размер собственного объекта кратен 16 байтам.В моем конкретном случае два Eigen::Matrix<double, 3,2> действительны для векторизации.

Однако в CUDA Эйген не запрашивает 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> не соответствует требованиям для векторизации, и поэтому нет необходимости добавлять 8 байтов заполнения в ЦП.

Answer 1

Такое различие связано с тем, как Eigen запрашивает выравнивание памяти в C ++ и CUDA.

В C ++ S выравнивается по 16 байтам (вы можете проверить, что alignof(S) == 16). Это происходит из-за матриц Эйгена, которые выровнены по 16 байтов, возможно, из-за использования регистров SSE, которые требуют такого выравнивания. Остальные ваши поля выровнены по 8 байтов (64-битные указатели и двойные).

В заголовочном файле Eigen/Core для CUDA включена директива EIGEN_DONT_VECTORIZE. При проверке документации :

EIGEN_DONT_VECTORIZE - отключает явную векторизацию, когда она определена. По умолчанию не определено, если только выравнивание не отключено тестом платформы Eigen или пользователем, определяющим EIGEN_DONT_ALIGN.

, что в основном означает, что собственные матрицы не имеют специального выравнивания в CUDA, поэтому они выровнены по типу элемента, double в вашем случае, что приводит к 8-байтовому выравниванию для матриц и, следовательно, для всей структуры.

Лучший способ ее решить - форсировать выравнивание структуры для обеих архитектур. Сейчас в CUDA не так свободно, я думаю, что вы можете сделать это с __align__(16) в CUDA (больше здесь ) и с использованием 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 об этом нет никаких гарантий).