库达型双关语 - memcpy vs UB 联盟

Cuda type punning - memcpy vs UB union

本文关键字：vs UB 联盟 memcpy 达型 双关语      更新时间：2023-10-16

在非 Cuda C++ 代码中，当前建议的做法是应该使用通过 memcpy 的类型双关语，而不是通过联合进行 UB。 尽管它可能会导致调试版本中的性能问题，并且我不得不几次使用 UB 根才能在发布版本中获得更好的性能。

库达的推荐做法是什么？ 它会总是不幸地在调试版本中调用 memcpy（） 吗？

所以我对此很好奇，并在编译器资源管理器中查看了一下https://godbolt.org/z/Cv5ozC似乎库达做了完整的记忆。没有一般的建议，但我们可以从 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

如果您先将 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

相关文章：