在非Cuda C++代码中,目前建议使用通过memcpy进行类型切换而不是通过union进行UB。尽管这可能会在Debug版本中导致性能问题,并且事实上为了获得更好的Release版本性能,我有时不得不采用UB方法。
在Cuda中推荐的做法是什么?在Debug版本中是否总是不幸地调用memcpy()?
在Cuda中推荐的做法是什么?在Debug版本中是否总是不幸地调用memcpy()?
所以我对此很好奇,看了一下编译器探索器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
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