Cuda type punning - memcpy против UB union

Question

В не-Cuda C++-коде в настоящее время предлагается практика, заключающаяся в том, что вместо объединения UB следует использовать штампование типов через memcpy. Несмотря на то, что это может вызывать проблемы с производительностью в сборках Debug, и тот факт, что мне пришлось несколько раз обращаться к корню UB для повышения производительности в сборках Release.

Какова рекомендуемая практика в Cuda? И будет ли он всегда вызывать memcpy() в сборках Debug?

Answer 1

Мне было любопытно об этом, и я заглянул в проводник компилятора https://godbolt.org/z/Cv5ozC Похоже, что cuda полностью выполняет memcpy. Общих рекомендаций нет, но мы можем кое-что узнать из стандарта C++20 с помощью std::bit_cast. Я думаю, что реализация std::bit_cast и изменение кишок для переинтерпретации приведений - лучший путь вперед. Это все еще неопределенное поведение, но вы компилируете для одной архитектуры, поэтому оно, по крайней мере, будет постоянно неопределенным. И это добавляет много защиты от идиотов, которую допускает неопределенный маршрут поведения.

template <class To, class From, class Res = typename std::enable_if<
    (sizeof(To) == sizeof(From)) && 
    (alignof(To) == alignof(From)) &&
    std::is_trivially_copyable<From>::value &&
    std::is_trivially_copyable<To>::value,
    To>::type>
__device__ Res& bit_cast(From& src) noexcept {
    return *reinterpret_cast<To*>(&src);
}

По-прежнему нет хорошего способа привести float[4] к float4, потому что система типов не может представить, что указатель на начало массива удовлетворяет требованиям выравнивания float4.

Если вам нужна дополнительная информация, я разобрался с этим, посмотрев доклад на CPPCon2019: https://www.youtube.com/watch?v=_qzMpk-22cc

Answer 2

Если вы сначала загрузите src в регистры, затем memcpy опущено, поэтому я думаю, что подход memcpy по-прежнему является предпочтительным способом избежать UB.

      template <class To, class From>
__device__  To bit_cast(From& src) noexcept {
    To tgt;
    From staged = src;
    memcpy(&tgt, &staged, sizeof(To));
    return tgt;
}

Модифицированная версия примера @ esdanol, которая показывает тот же сгенерированный код: https://cuda.godbolt.org/z/4f9n1ndEW