Cuda 类型双关 - memcpy 与 UB union

Question

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

Cuda 中推荐的做法是什么？它总是会不幸地在调试版本中调用 memcpy() 吗？

Answer 1

所以我对此很好奇，并在编译器资源管理器中查看了一下 https://godbolt.org/z/Cv5ozC 看来cuda执行了完整的memcpy。没有一般性的推荐，但我们可以使用 std::bit_cast 从 c++20 标准中学习一些东西。我认为实施 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