Half base type #1706
Half base type #1706
Conversation
265cb47 to
b0c488b
Compare
8bd8d1c to
2ab1acd
Compare
|
@MarcelKoch I have made the half without types.hpp dependence such that we can include it directly without circular dependence and the instantiate function have the definition of half now |
7eaf547 to
d4f3893
Compare
f924ff5 to
5a7ae0c
Compare
| @@ -17,7 +17,15 @@ | |||
| #include "utils.hpp" | |||
|
|
|||
|
|
|||
| struct __half; | |||
There was a problem hiding this comment.
I don't think there's a way around the issue, but we should be aware that __half is a reserved identifier according to the C++ standard, so we should technically not be defining anything with that name. Also there might be a tiny potential of name mangling issues if this is ever handled as a class instead of a struct with MSVC.
| half x = create_from_bits("0" "01111111" "00000000000000000000000"); | ||
|
|
||
| ASSERT_EQ(get_bits(x), get_bits("0" "01111" "0000000000")); |
There was a problem hiding this comment.
FYI, since C++14 we can use binary literals 0b01111111
|
|
||
| template <> | ||
| struct culibs_type_impl<std::complex<half>> { | ||
| using type = __half2; |
There was a problem hiding this comment.
Just for my understanding, __half2 is a vector type derived from __half?
There was a problem hiding this comment.
you are right. only one function __hcmadd does complex operation, but the other are all vector type operation.
There was a problem hiding this comment.
I have check cusparse they use __half2 to hold complex<__half>, too. https://docs.nvidia.com/cuda/cusparse/#cudadatatype-t
| # GDB doesn't seem to consider the user-defined conversion in its Value.cast, | ||
| # so we need to call the conversion operator explicitly | ||
| address = hex(val.address) | ||
| self.float_val = gdb.parse_and_eval(f"reinterpret_cast<gko::half*>({address})->operator float()") |
There was a problem hiding this comment.
If that function is ever not exported, we may need to implement an xfunction for it
There was a problem hiding this comment.
I am not familar with that. This is done by @MarcelKoch
could you implement that?
There was a problem hiding this comment.
We likely don't need it now, just for future reference
8cc9759 to
0fa6dc9
Compare
… half to another test
Co-authored-by: Marcel Koch <marcel.koch@kit.edu>
Co-authored-by: Marcel Koch <marcel.koch@kit.edu>
This PR implements the half precision for matrices and components Related PR: #1706
This pr moves the half out of extended_float.hpp with some modification and half is available in the public interface.
Currently, we can still mark some conversion/operation with GKO_ATTRIBUTE and GKO_INLINE.
However, I think it is good to ensure we only use the vendor's impl on device side.
To achieve it, Jacobi needs to use __half not half now.
This also removes the undefined behavior from
reinterpret_casttostd::memcpyone question: should we mark the half precision in experimental namespace?in usual interface from discussion