Relation between at::Half and __half

[SakuraRiven]

Hi, we want to create atomicMax function with __half inputs by following pytorch_scatter/atomics.cuh at master. The detailed codes are

inline __device__ void operator()(at::Half *address, at::Half val) {       \
  unsigned int *address_as_ui =                                            \
      (unsigned int *)((char *)address - ((size_t)address & 2));           \
  unsigned int old = *address_as_ui;                                       \
  unsigned int assumed;                                                    \
                                                                           \
  do {                                                                     \
    assumed = old;                                                         \
    at::Half hsum;                                                         \
    hsum.x = (size_t)address & 2 ? (old >> 16) : (old & 0xffff);           \
    hsum = OP(hsum, val);                                                  \
    old = (size_t)address & 2 ? (old & 0xffff) | (hsum.x << 16)            \
                              : (old & 0xffff0000) | hsum.x;               \
    old = atomicCAS(address_as_ui, assumed, old);                          \
  } while (assumed != old);                                                \
}                                                                          \

When it comes to the __half type, however, we do not how to obtain the hsum.x. Are there some ways to convert as::Half and __half each other?

Also, how could we implement OP as max to deal with __half inputs?

Thanks : )