Design Doc: float16

Why float16

Half precision (float16) is a binary floating-point format that occupies 16 bits in memory. float16 is half the size of traditional 32-bit single precision format (float) and has lower precision and smaller range.

When high precision computation is not required, using float16 data type could potentially

  • reduce storage space, memory bandwidth, and power usages;
  • increase the chance of data fitting into a smaller cache of lower latency;
  • provide arithmetic speed up if supported by hardware.

Survey of current float16 support

A brief survey of float16 support on different compilers, hardwares, and libraries can be found below. Interested readers can refer to link1 and link2 for more info.

The goal of float16 is to serve as a key for the executor to find and run the correct version of compute method specialized for float16 in operator kernel. It should be compatible with various natively supported float16 implementations including __half for cuda, float16_t for ARM, and Eigen::half for Eigen to make writing customized float16 kernels easier.

Compiler

  • nvcc supports __half data type after CUDA 7.5.
  • __fp16 or float16_t is supported as storage type for gcc >= 6.1 and clang >= 3.4.
  • __fp16 or float16_t is supported as arithmetic type for gcc >= 7.1 and clang >= 3.9.

Hardware

  • __half is supported on GPU with compute capability >= 5.3.
  • __fp16 is supported as storage type for ARMv7-A, ARMv8-A, and above.
  • __fp16 is supported as arithmetic type after ARMv8.2-A (currently, the only microarchitecture implementing ARMv8.2-A is ARM Cortex-A75, which is announced in May 2017. There seems to be no application processors currently available on market that adopts this architecture. It is reported that Qualcomm Snapdragon 845 uses Cortex-A75 design and will be available in mobile devices in early 2018).

Libraries

  • Eigen >= 3.3 supports float16 calculation on both GPU and CPU using the Eigen::half class. It is mostly useful for Nvidia GPUs because of the overloaded arithmetic operators using cuda intrinsics. It falls back to using software emulation on CPU for calculation and there is no special treatment to ARM processors.
  • ARM compute library >= 17.02.01 supports NEON FP16 kernels (requires ARMv8.2-A CPU).

CUDA version issue

There are currently three versions of CUDA that supports __half data type, namely, CUDA 7.5, 8.0, and 9.0. CUDA 7.5 and 8.0 define __half as a simple struct that has a uint16_t data (see cuda_fp16.h) as follows:

typedef struct __align__(2) {
   unsigned short x;
} __half;

typedef __half half;

This struct does not define any overloaded arithmetic operators. So you have to directly use __hadd instead of + to correctly add two half types:

__global__ void Add() {
  half a, b, c;
  c = __hadd(a, b); // correct
  c = a + b; // compiler error: no operator "+" matches these operands
}

CUDA 9.0 provides a major update to the half data type. The related code can be found in the updated cuda_fp16.h and the newly added cuda_fp16.hpp.

Essentially, CUDA 9.0 renames the original __half type in 7.5 and 8.0 as __half_raw, and defines a new __half class type that has constructors, conversion operators, and also provides overloaded arithmetic operators such as follows:

typedef struct __CUDA_ALIGN__(2) {
    unsigned short x;
} __half_raw;


struct __CUDA_ALIGN__(2) __half {
protected:
    unsigned short __x;
public:
    // constructors and conversion operators from/to 
    // __half_raw and other built-in data types
}

typedef __half half;

__device__ __forceinline__ 
__half operator+(const __half &lh, const __half &rh) { 
    return __hadd(lh, rh); 
}

// Other overloaded operators

This new design makes c = a + b work correctly for CUDA half data type.

Implementation

The float16 class holds a 16-bit uint16_t data internally.

struct float16 {
  uint16_t x;
};

float16 supports the following features:

  • constructors / assignment operators that take input from primitive data types including bool, integers of various length, float, and double.
  • constructors / assignment operators that take input from __half on cuda, float16_t on ARM, and Eigen::half on Eigen.
  • conversion operators to primitive data types and half precision data types on cuda, ARM and Eigen.
  • overloaded arithmetic operators for cuda, arm, and non-arm cpu, respectively. These operators will take advantage of the cuda and ARM intrinsics on the corresponding hardware.

To support the above features, two fundamental conversion functions are provided:

float16 float_to_half_rn(float f);  // convert to half precision in round-to-nearest-even mode
float half_to_float(float16 h);

which provides one-to-one conversion between float32 and float16. These twos functions will do different conversion routines based on the current hardware. CUDA/ARM instrinsics will be used when the corresonding hardware is available. If the hardware or compiler level does not support float32 to float16 conversion, software emulation will be performed to do the conversion.

To do

After float16 class is available, some of the future items are below:

  • Update pybind/tensor_py.h to bind c++ float16 with numpy float16.
  • Modify GetKernelType() method in framework/operator.h to make it compatible with float16.
  • Create a type-casting operator that can convert the data type in tensor between float16 and other types.