的开发人员,在半精度算术方面遇到了困难。
在库中的某个时候,我需要在 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
时,则不行。__half
int16_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 转换函数?
3
最佳答案
2
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
,而是成对处理它们。
1
-
哇,太棒了。我会尝试一下。
–
|
如果您可以分离主机端和设备端实现(您可以 – 至少使用#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 函数。
4
-
我没有库的这一部分运行时选项。但我肯定会向 NVIDIA 提交错误报告。
– -
@RegisPortalez:如果库的这一部分在运行时没有使用,那么也许你应该放弃
cuda_fp16.hpp
并使用 C++ 标准std::float16_t
或编译器的内置 FP16 类型(_Float16
例如在 GCC 中)
– -
1FP16 库的问题在于它是 c++11。您需要重写它,与 cuda_fp16 标头相同。
–
-
向 NVIDIA 提交了一个错误:
–
|
constexpr
现成的版本。您可以使用它std::is_constant_evaluated
来避免在运行时进行低效的转换。–
–
–
|