在CUDA中将类型转换为unsigned long long？

Posted 2023-04-15

【中文标题】在CUDA中将类型转换为unsigned long long？

【英文标题】：type casting to unsigned long long in CUDA?

【发布时间】：2013-08-21 04:36:35

【问题描述】：

基本上我想要的是一个类似 hiloint2uint64() 的函数，只需连接两个 32 位整数并将结果重新解释为 uint64。

我在 CUDA 中找不到任何可以做到这一点的函数，无论如何，是否有任何 ptx 代码可以进行这种类型转换？

【参考方案1】：

您可以像这样定义自己的函数：

__host__ __device__ unsigned long long int hiloint2uint64(int h, int l)

    int combined[] =  h, l ;

    return *reinterpret_cast<unsigned long long int*>(combined);

【讨论】：

C++ reinterpret_cast 是处理这种情况的正确方法。

【参考方案2】：

现在可能有点晚了，但最安全的方法可能是“手动”使用位移和或：

uint32_t ui_h = h;
uint32_t ui_l = l;
return (uint64_t(h)<<32)|(uint64_t(l));

请注意，其他答案中提出的其他解决方案并不安全，因为整数数组可能不是 8 字节对齐的（无论如何，移位一些位比内存读/写更快）

【参考方案3】：

使用 uint2（但将临时变量定义为 64 位值：unsigned long long int）而不是数组来确保对齐。 注意 l 和 h 的顺序。

__host__ __device__ __forceinline__ unsigned long long int hiloint2uint64(unsigned int h, unsigned int l)

    unsigned long long int result;
    uint2& src = *reinterpret_cast<uint2*>(&result);
    src.x = l;
    src.y = h;
    return result;

无论如何，CUDA 寄存器的大小都是 32 位。在最好的情况下，编译器不需要任何额外的代码。在最坏的情况下，它必须通过移动 32 位值来重新排序寄存器。

Godbolt 示例 https://godbolt.org/z/3r9WYK9e7 的优化程度。