Skip to content

Latest commit

 

History

History
343 lines (289 loc) · 13.9 KB

float16_and_quantized_type.md

File metadata and controls

343 lines (289 loc) · 13.9 KB
Raw

Float16 and Quantized Int8 Type

Yiqun Liu

2017-06-27

Part I, float16 - FP16, half

IEEE 754存储格式

  • float

    • 符号位： 1位
    • 指数位： 8位，范围 $2^{-126}$ ~ $2^{127}$
    • 尾数位：23位 50%

  • half

    • 符号位： 1位
    • 指数位： 5位，范围 $2^{-14}$ ~ $2^{15}$
    • 尾数位：10位 52%

基本类型支持

  • caffe2
    • 类型定义caffe2/caffe2/core/types.h 
      namespace caffe2 {
typedef struct CAFFE2_ALIGNED(2) __f16 { uint16_t x; } float16;
}  // namespace caffe2
    • 类型转换 caffe2/caffe2/utils/conversions.h 
      inline float16 cpu_float2half_rn(float f) {
  float16 ret;
  ...
  exponent = ((u >> 23) & 0xff);
  mantissa = (u & 0x7fffff);
  ...
  ret.x = (sign | (exponent << 10) | mantissa);
  return ret;
}
inline float cpu_half2float(float16 h) {
  unsigned sign = ((h.x >> 15) & 1);
  unsigned exponent = ((h.x >> 10) & 0x1f);
  unsigned mantissa = ((h.x & 0x3ff) << 13);
  ...
  unsigned i = ((sign << 31) | (exponent << 23) | mantissa);
  float ret;
  memcpy(&ret, &i, sizeof(i));
  return ret;
}

基本类型支持

  • CUDA include/cuda_fp16.h
    • CUDA 7.5后 文档
    • 类型定义 
      typedef struct __align__(2) {
  unsigned short x;
} __half;
typedef struct __align__(4) {
  unsigned int x;
} __half2;
    • 类型转换 
      __device__ __half __float2half(const float a);
__device__ float __half2float(const __half a);
__device__ __half2 __floats2half2_rn(const float a, const float b);
__device__ float2 __half22float2(const __half2 a);
    • 计算函数 
      __device__ __half __hadd(const __half a, const __half b);
__device__ __half2 __hadd2(const __half2 a, const __half2 b);

基本类型支持

half为什么比float快？

  • 硬件支持-NVIDIA GPU, Mixed-Precision Programming with CUDA 8
    • half2类型，两路向量半精度融合乘加指令(HFMA2)，一条指令操作两个half数据 
      __device__ __half2 __hfma2(const __half2 a, const __half2 b, const __half2 c); // a * b + c
    • 使用示例 
      __global__ void haxpy(int n, half a, const half *x, half *y) {
  int start = threadIdx.x + blockDim.x * blockIdx.x;
  int stride = blockDim.x * gridDim.x;
#if __CUDA_ARCH__ >= 530
  int n2 = n/2;
  half2 *x2 = (half2*)x, *y2 = (half2*)y;
  for (int i = start; i < n2; i+= stride) 
    y2[i] = __hfma2(__halves2half2(a, a), x2[i], y2[i]);
  // first thread handles singleton for odd arrays
  if (start == 0 && (n%2)) y[n-1] = __hfma(a, x[n-1], y[n-1]);   
#else
  for (int i = start; i < n; i+= stride)
    y[i] = __float2half(__half2float(a) * __half2float(x[i]) + __half2float(y[i]));
#endif
}
    • Eigen内部会自动使用half2类型计算 Eigen/src/Core/arch/CUDA/PacketMathHalf.h

half为什么比float快？

  • 计算库支持-NVIDIA GPU
    • cuBLAS
      • cublasHgemm，使用FP16计算，并且作为输入输出
      • cublasSgemmEx，使用FP32计算，输入数据可以是FP32、FP16或INT8，输出数据可以是FP32或FP16
    • cuDNN
      • 5.0支持FP16卷积前向计算，5.1支持FP16卷积后向计算
    • TensorRT
      • v1有FP16 inference卷积支持
    • 其他支持FP16 gemm的计算库：nervanagpu，openai-gemm

half为什么比float快？

  • 硬件支持-armv8 CPU
    • 基本数据类型float16_t
    • 向量数据类型float16x8_t
    • 函数支持

深度学习系统中的应用

  • caffe2
    • if分支控制不同数据类型的计算 caffe2/operators/fully_connected_op_gpu.cc 
      template <> bool FullyConnectedOp<CUDAContext>::RunOnDevice() {
  if (Input(0).IsType<float>()) {
    return DoRunWithType<
      float, // X
      float, // W
      float, // B
      float, // Y
      float>(); // Math
  } else if (Input(0).IsType<float16>()) {
    return DoRunWithType<
      float16, // X
      float16, // W
      float16, // B
      float16, // Y
      float>(); // Math
  } else {
    CAFFE_THROW("Unsupported type");
  }
  return false;
}
    • 只有部分Op支持float16
      • FullyConnectedOp, CudnnConvOp, CudnnSpatialBNOp, CuDNNPoolOp, CuDNNReluOp, CuDNNDropoutOp, MaxPoolWithIndexOp, WeightdSumOp, SumOp, CUDAAddOp

深度学习系统中的应用

PartII, Quantized int8 - Fixed point

基本类型支持

基本类型支持

  • 类型转换方法一
    • 对于 $Y = W * X$
    • 若 $\tilde{W} = scale_W * W$，且 $\tilde{X} = scale_X * X$
    • 那么 $Y = \tilde{W} * \tilde{X} / (scale_W * scale_X)$
    • 取 $scale_W = 128 / max(|W_{ij}|)$，那么 $-127 &lt; \tilde{W}_{ij} &lt; 128$
  • 类型转换方法二 60%
    • 令 $\tilde{Y}=\tilde{W} * \tilde{X}$，则 $Y_{ij}=\tilde{Y} / {scale_W_i * scale_X_j}$

基本类型支持

  • 类型转换方法-Google
    • On the efficient representation and execution of deep acoustic models
    • 使用QUInt8量化，希望量化到的范围[0, 255] $$ \tilde{V} = \frac{255}{V_{max}-V_{min}} * (V - V_{min})$$
    • 令 $Q=\frac{255}{V_{max}-V_{min}}$，恢复 $$ V = \tilde{V}*Q^{-1} + V_{min}$$
    • 作用于$Y=WX$，可得 $$ Y=\tilde{W}\tilde{X}*Q_W^{-1}*Q_X^{-1} + \tilde{W}*Q_W^{-1}*X_{min} + \tilde{X}*Q_X^{-1}*W_{min}+W_{min}*X_{min}$$

为什么更快？

  • 硬件支持-NVIDIA GPU， Mixed-Precision Programming with CUDA 8
    • include/sm_61_intrinsics.h 70%
    • 计算函数
    __device__ int __dp4a(int srcA, int srcB, int c);
__device__ int __dp4a(char4 srcA, char4 srcB, int c);

为什么更快？

  • cuBLAS
    • CUDA 8.0引入新接口，cublasGemm，支持INT8计算
  • cuDNN
    • v6增加INT8 inference支持
  • TensorRT
    • v2将会增加INT8 inference支持

为什么更快？

  • 硬件支持-armv7a CPU
    • 提供基于int8x8_t的计算 
      int16x8_t vmlal_s8(int16x8_t a, int8x8_t b, int8x8_t c);
    • 计算库

深度学习系统中的应用

深度学习系统中的应用

Thank You!