Cuda类型转换 - memcpy vs UB联合

Question

在非Cuda C++代码中，目前建议使用通过memcpy进行类型切换而不是通过union进行UB。尽管这可能会在Debug版本中导致性能问题，并且事实上为了获得更好的Release版本性能，我有时不得不采用UB方法。

在Cuda中推荐的做法是什么？在Debug版本中是否总是不幸地调用memcpy()？

- iam

---

你能展示一个memcpy方式的例子，在发布版本中会导致性能下降吗？ - Richard Hodges

不，因为它在我不再能访问的旧公司代码库中。我认为它是在一个矢量类库中，用的是 ~VS2013（在这种情况下没有使用Cuda来明确）。但人们希望在VS2017中的发布模式下，现在应该可以了 - 但我不确定。 - iam

我认为在现代编译器中，memcpy路线将是最有效和可靠的。类型转换可能会导致未定义行为，甚至可能导致代码根本无法工作 - 特别是在优化时。 - Richard Hodges

在C++中确实如此，但在CUDA C++中可能并非如此，因为由于GPU寄存器的限制，它可能会执行“意外”的操作，例如由于能够通过union或memcpy重新解释事物而导致内存溢出。 因此，32位寄存器大小类型之间的pun转可能具有优势。 - iam

Answer 1

所以我对此很好奇，看了一下编译器探索器https://godbolt.org/z/Cv5ozC。似乎cuda做了完整的memcpy。没有通用建议，但我们可以从c++20标准中学到一些东西，即std::bit_cast。我认为采用std::bit_cast的实现并将其内部改为reinterpret 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

- esdanol

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

- Antonio Sánchez