Question

的开发人员，在半精度算术方面遇到了困难。

在库中的某个时候，我需要在 constexpr 函数中将 a 转换IntType为相关的FloatType（相同的位数）（以评估多项式）。

我有这个功能（更大的类型aerobus::i16::val<x>）：

template<typename valueType>
static constexpr valueType get() { return static_cast<valueType>(x); }

在一个封闭的类型中。

当 valueType 是float或double且 x 是int32_t或时，它可以正常工作，但当 valueType 是且 x 是int64_t时，则不行。__halfint16_t

nvcc 12.6.r12.6 生成以下错误：

../src/aerobus.h:576:55: note: ‘static constexpr valueType aerobus::i16::val<x>::get() [with valueType = __half; short int x = 64]’ is not usable as a ‘constexpr’ function because:
  576 |             static constexpr INLINED DEVICE valueType get() { return (valueType)x; }
      |                                                       ^~~
../src/aerobus.h:576:55: error: call to non-‘constexpr’ function ‘__half::__half(short int)’
In file included from cuda_fp16.cpp:1:
/usr/local/cuda-12.6/bin/../targets/x86_64-linux/include/cuda_fp16.h:4652:25: note: ‘__half::__half(short int)’ declared here
 4652 |     __CUDA_HOSTDEVICE__ __half(const short val) { __x = __short2half_rn(val).__x; }

我该如何解决这个问题？我是否被迫实现从 int16_t 到 __half 的 constexpr 转换函数？

看起来是这样的，如果没有constexpr现成的版本。您可以使用它std::is_constant_evaluated来避免在运行时进行低效的转换。 — 
@passerby 谢谢。我希望它存在于 cuda 标头的某个地方。 — 
@PasserBy：您认为什么是“现成的”？请参阅我的回答和那里的链接… —

Accepted Answer

cuda_fp16.hpp 标头中没有任何内容包含 constexpr，因此您必须自己完成这项艰苦的工作。以下内容为我编译：

static constexpr unsigned short my_internal_float2half(const float f, unsigned int& sign, unsigned int& remainder) {
    unsigned int x;
    unsigned int u;
    unsigned int result;
    x = std::bit_cast<int>(f); //c++20 
    ////(void)std::memcpy(&x, &f, sizeof(f)); //not constexpr
    u = (x & 0x7fffffffU);
    sign = ((x >> 16U) & 0x8000U);
    // NaN/+Inf/-Inf
    if (u >= 0x7f800000U) {
        remainder = 0U;
        result = ((u == 0x7f800000U) ? (sign | 0x7c00U) : 0x7fffU);
    } else if (u > 0x477fefffU) { // Overflows
        remainder = 0x80000000U;
        result = (sign | 0x7bffU);
    } else if (u >= 0x38800000U) { // Normal numbers
        remainder = u << 19U;
        u -= 0x38000000U;
        result = (sign | (u >> 13U));
    } else if (u < 0x33000001U) { // +0/-0
        remainder = u;
        result = sign;
    } else { // Denormal numbers
        const unsigned int exponent = u >> 23U;
        const unsigned int shift = 0x7eU - exponent;
        unsigned int mantissa = (u & 0x7fffffU);
        mantissa |= 0x800000U;
        remainder = mantissa << (32U - shift);
        result = (sign | (mantissa >> shift));
        result &= 0x0000FFFFU;
    }
    return static_cast<unsigned short>(result);
}

static constexpr __half my_float2half_rn(const float a) {
    __half val;
    __half_raw r;
    unsigned int sign = 0U;
    unsigned int remainder = 0U;
    r.x = my_internal_float2half(a, sign, remainder);
    if ((remainder > 0x80000000U) || ((remainder == 0x80000000U) && ((r.x & 0x1U) != 0U))) {
        r.x++;
    }
    val = std::bit_cast<__half>(r); //allowed, see operator= for __raw_half -> __half
    return val;
}

static constexpr __half my_int2half_rn(const int i) {
    __half h;
        // double-rounding is not a problem here: if integer
        // has more than 24 bits, it is already too large to
        // be represented in half precision, and result will
        // be infinity.
        const float  f = static_cast<float>(i);
    h = my_float2half_rn(f);
    return h;
}

__device__ consteval __half convert_int16_to_half(int16_t i) {
    return my_float2half_rn(static_cast<float>(i));
}

#ifdef INSIDE_OPS_CLASS_DEF
//in your class header
template<typename VT>
static constexpr VT get() { 
    static_assert(std::is_same<int16_t, decltype(this->x)>::value);
    if constexpr (std::is_same<VT, __half>::value) {
        return convert_int16_to_half(this->x);
    } else {
        return static_cast<valueType>(x); 
    }
}
#endif

请注意，在之外consteval，您将需要使用内置版本，因为它们使用单个汇编语句，让 GPU 进行转换。

请注意，如果您担心效率，那么您真的不应该只处理单个__half，而是成对处理它们。

Answer 2

如果您可以分离主机端和设备端实现（您可以 – 至少使用#ifdef __CUDA_ARCH__等） – 并且您的主机端编译器支持 C++23，那么您可以按照通常的方式将整数转换为半精度类型。示例：

#include <stdfloat>

constexpr std::float16_t magic(std::int16_t x)
{
    return static_cast<std::float16_t>(x);
}

这 (GodBolt)。当然，您需要将其重新解释std::float16_t为__half，它不是 constexpr，但您可以将其留到运行时，而不是编译时。

其他可以尝试的事情：

使用此转换的现有、已建立的实现，可以将其标记为 constexpr。Johan的自己的答案；但您也可以考虑，该库正是用于这些目的并且相当受欢迎。，。int2half
如果你能够以 constexpr 方式向上转换为浮点数（同样，这需要检查，例如 clang、MSVC），则可以使用编译器的 constexpr int 到 float 转换
向 NVIDIA 提交有关此问题的错误。这不是一个立即的解决方案，但来自用户的压力可能会让他们为我们提供 constexpr fp16 函数。

我没有库的这一部分运行时选项。但我肯定会向 NVIDIA 提交错误报告。 — 
@RegisPortalez：如果库的这一部分在运行时没有使用，那么也许你应该放弃cuda_fp16.hpp并使用 C++ 标准std::float16_t或编译器的内置 FP16 类型（_Float16例如在 GCC 中） — 
FP16 库的问题在于它是 c++11。您需要重写它，与 cuda_fp16 标头相同。 —

c++ – 如何以 constexpr 方式将整数转换为 CUDA 的 __half FP16 类型？ – 堆栈内存溢出

最佳答案
2

最佳答案 2

最佳答案
2