Identify and list features of existing extended floating point wrapper types

[jrhemstad]

@cliffburdick said it best:

CUDA does not have standardized classes/operators available on both host and device to support these operations, so we define it here.

As it happens, there appears to be a lot of projects and people who feel the same:

• cutlass::bfloat16_t
• cutlass::half_t
• cute::bfloat (alias to cutlass::bfloat16_t)
• cute::half (alias to cutlass::half_t)
• matx::matxFp16
• matx::matxBf16
• cub::half_t
• cub::bfloat16_t
• c10::BFloat16
• c10::Half
• Eigen::BFloat16
• Eigen::Half
• cuSPARSE made some internal wrappers too

Things to look into:

• RAPIDS?

[weinbe2]

Kokkos has Kokkos::Experimental::half_t and Kokkos::Experimental::bhalf_t, where it'll properly wrap the data types if a given target supports it... and otherwise it's a thin wrapper around float.

It has a few levels of abstraction (as Kokkos does), but the core abstract implementation is here: https://github.com/kokkos/kokkos/blob/master/core/src/impl/Kokkos_Half_FloatingPointWrapper.hpp ; and, as an example, the way it'll picks a proper type for CUDA if a given version/arch supports it is here: https://github.com/kokkos/kokkos/blob/master/core/src/Cuda/Kokkos_Cuda_Half_Impl_Type.hpp

[achirkin]

In rapidsai/raft ANN benchmarks, we don't have a wrapper; instead, we have a uint16_t stub for CPU-only builds. Also across raft, the lack of <type_traits> support (i.e. something like std::is_floating_point) has been an annoying issue, which still prevents parts of raft from supporting half.

[jrhemstad]

In rapidsai/raft ANN benchmarks, we don't have a wrapper; instead, we have a uint16_t stub for CPU-only builds. Also across raft, the lack of <type_traits> support (i.e. something like std::is_floating_point) has been an annoying issue, which still prevents parts of raft from supporting half.

Thanks @achirkin, this is really helpful feedback. We've heard a lot of similar things, and so we're working on figuring out how to address all of these problems and more :)

[miscco]

I bricked my workstation for a few hours so here is a comparison. I tried to be diligent, but it is hard to find every function:

half

API	Kokkos	MatX	CUB	pytorch	Eigen	cutlass
cast_to_half	✔️	❌	✔️	✔️	✔️	✔️
cast_from_half	✔️	✔️	✔️	✔️	✔️	✔️
conversion to float	✔️	✔️	✔️	✔️	✔️	✔️
int constructor	✔️	✔️	✔️	✔️	✔️	✔️
converting constructor	✔️	✔️	✔️	✔️	✔️	✔️
converting assignment	✔️	✔️	❌	❌	❌	❌
---	---	---	---	---	---	---
operator+= half	✔️	✔️	✔️	✔️	✔️	✔️
operator+= float	✔️	✔️	❌	✔️	❌	❌
operator-= half	✔️	✔️	✔️	✔️	✔️	✔️
operator-= float	✔️	✔️	❌	✔️	❌	❌
operator*= half	✔️	✔️	✔️	✔️	✔️	✔️
operator*= float	✔️	✔️	❌	✔️	❌	❌
operator/= half	✔️	✔️	✔️	✔️	✔️	✔️
operator/= float	✔️	✔️	❌	✔️	❌	❌
operator+ half, half	✔️	✔️	✔️	✔️	✔️	✔️
operator+ T, half	✔️	✔️	❌	✔️	❌	❌
operator- half, half	✔️	✔️	✔️	✔️	✔️	✔️
operator- T, half	✔️	✔️	❌	✔️	❌	❌
operator* half, half	✔️	✔️	✔️	✔️	✔️	✔️
operator* T, half	✔️	✔️	❌	✔️	❌	❌
operator/ half, half	✔️	✔️	✔️	✔️	✔️	✔️
operator/ T, half	✔️	✔️	❌	✔️	❌	❌
unary+	✔️	❌	❌	❌	❌	❌
unary-	✔️	❌	❌	❌	❌	✔️
---	---	---	---	---	---	---
operator== half	✔️	✔️	✔️	❌	✔️	✔️
operator== float	❌	✔️	❌	❌	❌	❌
operator!= half	✔️	✔️	✔️	❌	✔️	✔️
operator!= float	❌	✔️	❌	❌	❌	❌
operator< half	✔️	✔️	✔️	❌	✔️	✔️
operator< float	✔️	❌	❌	❌	❌	❌
operator> half	✔️	✔️	✔️	❌	✔️	✔️
operator> float	✔️	❌	❌	❌	❌	❌
operator<= half	✔️	✔️	✔️	❌	✔️	✔️
operator<= float	✔️	❌	❌	❌	❌	❌
operator>= half	✔️	✔️	✔️	❌	✔️	✔️
operator>= float	✔️	❌	❌	❌	❌	❌
---	---	---	---	---	---	---
operatorOR	✔️	❌	❌	❌	❌	❌
operator&&	✔️	❌	❌	❌	❌	❌
operator++	✔️	✔️	✔️	❌	✔️	✔️
operator++(int)	✔️	✔️	✔️	❌	✔️	✔️
operator--	✔️	❌	❌	❌	✔️	✔️
operator--(int)	✔️	❌	❌	❌	✔️	✔️
---	---	---	---	---	---	---
abs	✔️	✔️	❌	❌	✔️	❌
fmod	✔️	❌	❌	❌	✔️	❌
fmax	✔️	❌	❌	❌	❌	❌
fmin	✔️	❌	❌	❌	❌	❌
fdim	✔️	❌	❌	❌	❌	❌
exp	✔️	✔️	❌	❌	✔️	❌
exp2	✔️	❌	❌	❌	❌	❌
expm1	✔️	❌	❌	❌	✔️	❌
log	✔️	✔️	❌	❌	✔️	❌
log1p	✔️	❌	❌	❌	✔️	❌
log2	✔️	✔️	❌	❌	✔️	❌
log10	✔️	✔️	❌	❌	✔️	❌
pow	✔️	✔️	❌	❌	✔️	❌
sqrt	✔️	✔️	❌	❌	✔️	✔️
cbrt	✔️	❌	❌	❌	❌	❌
hypot	✔️	❌	❌	❌	❌	❌
erf	✔️	❌	❌	❌	❌	❌
erfc	✔️	❌	❌	❌	❌	❌
tgamma	✔️	❌	❌	❌	❌	❌
lgamma	✔️	❌	❌	❌	❌	❌
signbit	❌	❌	❌	❌	❌	✔️
fpclassify	❌	❌	❌	❌	❌	✔️
copysign	❌	❌	❌	❌	❌	✔️
isinf	✔️	✔️	❌	❌	✔️	✔️
isnan	✔️	❌	❌	❌	✔️	✔️
isfinite	✔️	❌	❌	❌	✔️	✔️
isnormal	❌	❌	❌	❌	❌	✔️
ceil	✔️	✔️	❌	❌	✔️	❌
floor	✔️	✔️	❌	❌	✔️	❌
round	✔️	✔️	❌	❌	✔️	❌
trunc	✔️	❌	❌	❌	❌	❌
rint	❌	❌	❌	❌	✔️	❌
sin	✔️	✔️	❌	❌	✔️	❌
cos	✔️	✔️	❌	❌	✔️	❌
tan	✔️	✔️	❌	❌	✔️	❌
asin	✔️	✔️	❌	❌	✔️	❌
acos	✔️	✔️	❌	❌	✔️	❌
atan	✔️	✔️	❌	❌	✔️	❌
atan2	✔️	✔️	❌	❌	✔️	❌
sinh	✔️	✔️	❌	❌	❌	❌
cosh	✔️	✔️	❌	❌	❌	❌
tanh	✔️	✔️	❌	❌	✔️	❌
asinh	✔️	✔️	❌	❌	❌	❌
acosh	✔️	✔️	❌	❌	❌	❌
atanh	✔️	✔️	❌	❌	✔️	❌
numeric_limits	✔️	❌	✔️	✔️	✔️	✔️

[jrhemstad]

Awesome work @miscco. I think the only missing piece is to understand what each of these solutions does for atomic operations. Do they overload things like atomicAdd? Something else?

@cliffburdick was there anything you recall having to add for the matx types that you don't see reflected in the table above?

[cliffburdick]

Awesome work @miscco. I think the only missing piece is to understand what each of these solutions does for atomic operations. Do they overload things like atomicAdd? Something else?

@cliffburdick was there anything you recall having to add for the matx types that you don't see reflected in the table above?

Not sure why I didn't see this. I think this is comprehensive, but one of the main issues was also having types that are compatible with their complex counterparts.