Блог

Я столкнулся с некоторыми проблемами при использовании ядра, которое использует некоторые структуры, которые я определил в 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;
}
  

Этот файл имеет размер 160 байт в C , но 152 байта в CUDA. Чтобы передать данные, я выделяю буфер на стороне CUDA и выполняю cudaMemcpy

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

что, я думаю, неверно, так как размер в CUDA и в C разный.

Как только я пытаюсь получить доступ к S::node0 , S::node1 или S::node3 в ядре, я получаю ошибку несогласованного доступа.

Итак, у меня есть три вопроса по этому вопросу:

• Почему размеры разные?
• Как я должен изменить код или выполнить копирование, чтобы это работало правильно?
• Должен ли я иметь структуру на стороне CUDA и выполнять специальную копию?

Редактировать: Благодаря принятому ответу я смог понять причину возникшей у меня проблемы. Eigen использует векторизацию, когда это возможно, и запрашивает для этого выравнивание на 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 байтов заполнения в CPU.

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

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

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

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 по этому поводу нет гарантии).