float16.html



<!DOCTYPE html>
<!--[if IE 8]><html class="no-js lt-ie9" lang="en" > <![endif]-->
<!--[if gt IE 8]><!--> <html class="no-js" lang="en" > <!--<![endif]-->
<head>
  <meta charset="utf-8">
  
  <meta name="viewport" content="width=device-width, initial-scale=1.0">
  
  <title>Design Doc: float16 &mdash; PaddlePaddle  文档</title>
  

    <link rel="stylesheet" href="../_static/css/theme.css" type="text/css" />
  

        <link rel="index" title="索引"
              href="../genindex.html"/>
        <link rel="search" title="搜索" href="../search.html"/>
    <link rel="top" title="PaddlePaddle  文档" href="../index.html"/> 
<script>
var _hmt = _hmt || [];
(function() {
  var hm = document.createElement("script");
  hm.src = "//hm.baidu.com/hm.js?b9a314ab40d04d805655aab1deee08ba";
  var s = document.getElementsByTagName("script")[0]; 
  s.parentNode.insertBefore(hm, s);
})();
</script>


  <script src="../_static/js/modernizr.min.js"></script>

</head>

<body class="wy-body-for-nav" role="document">

  <div class="wy-grid-for-nav">

    
    <nav data-toggle="wy-nav-shift" class="wy-nav-side">
      <div class="wy-side-scroll">
        <div class="wy-side-nav-search">
          

            <a href="../index_cn.html" class="icon icon-home"> PaddlePaddle
          

          </a>

          
<div role="search">
  <form id="rtd-search-form" class="wy-form" action="../search.html" method="get">
    <input type="text" name="q" placeholder="Search docs" />
    <input type="hidden" name="check_keywords" value="yes" />
    <input type="hidden" name="area" value="default" />
  </form>
</div>

          
        </div>

        <div class="wy-menu wy-menu-vertical" data-spy="affix" role="navigation" aria-label="main navigation">
          
<nav class="doc-menu-vertical" role="navigation">

<ul>
<li class="toctree-l1"><a class="reference internal" href="../getstarted/index_cn.html">新手入门</a><ul>
<li class="toctree-l2"><a class="reference internal" href="../getstarted/quickstart_cn.html">快速开始</a></li>
<li class="toctree-l2"><a class="reference internal" href="../getstarted/concepts/use_concepts_cn.html">基本使用概念</a></li>
</ul>
</li>
<li class="toctree-l1"><a class="reference internal" href="../build_and_install/index_cn.html">安装与编译</a><ul>
<li class="toctree-l2"><a class="reference internal" href="../build_and_install/pip_install_cn.html">使用pip安装</a></li>
<li class="toctree-l2"><a class="reference internal" href="../build_and_install/docker_install_cn.html">使用Docker安装运行</a></li>
<li class="toctree-l2"><a class="reference internal" href="../build_and_install/build_from_source_cn.html">从源码编译</a></li>
</ul>
</li>
<li class="toctree-l1"><a class="reference internal" href="../howto/index_cn.html">进阶使用</a><ul>
<li class="toctree-l2"><a class="reference internal" href="../howto/cmd_parameter/index_cn.html">命令行参数设置</a><ul>
<li class="toctree-l3"><a class="reference internal" href="../howto/cmd_parameter/use_case_cn.html">使用案例</a></li>
<li class="toctree-l3"><a class="reference internal" href="../howto/cmd_parameter/arguments_cn.html">参数概述</a></li>
<li class="toctree-l3"><a class="reference internal" href="../howto/cmd_parameter/detail_introduction_cn.html">细节描述</a></li>
</ul>
</li>
<li class="toctree-l2"><a class="reference internal" href="../howto/cluster/index_cn.html">分布式训练</a><ul>
<li class="toctree-l3"><a class="reference internal" href="../howto/cluster/preparations_cn.html">环境准备</a></li>
<li class="toctree-l3"><a class="reference internal" href="../howto/cluster/cmd_argument_cn.html">启动参数说明</a></li>
<li class="toctree-l3"><a class="reference internal" href="../howto/cluster/multi_cluster/index_cn.html">在不同集群中运行</a><ul>
<li class="toctree-l4"><a class="reference internal" href="../howto/cluster/multi_cluster/k8s_cn.html">Kubernetes单机训练</a></li>
<li class="toctree-l4"><a class="reference internal" href="../howto/cluster/multi_cluster/k8s_distributed_cn.html">Kubernetes分布式训练</a></li>
<li class="toctree-l4"><a class="reference internal" href="../howto/cluster/multi_cluster/openmpi_cn.html">在OpenMPI集群中启动训练</a></li>
<li class="toctree-l4"><a class="reference internal" href="../howto/cluster/multi_cluster/fabric_cn.html">使用fabric启动集群训练</a></li>
<li class="toctree-l4"><a class="reference internal" href="../howto/cluster/multi_cluster/k8s_aws_cn.html">Kubernetes on AWS</a></li>
</ul>
</li>
</ul>
</li>
<li class="toctree-l2"><a class="reference internal" href="../howto/capi/index_cn.html">C-API预测库</a><ul>
<li class="toctree-l3"><a class="reference internal" href="../howto/capi/compile_paddle_lib_cn.html">安装与编译C-API预测库</a></li>
<li class="toctree-l3"><a class="reference internal" href="../howto/capi/organization_of_the_inputs_cn.html">输入/输出数据组织</a></li>
<li class="toctree-l3"><a class="reference internal" href="../howto/capi/workflow_of_capi_cn.html">C-API使用流程</a></li>
</ul>
</li>
<li class="toctree-l2"><a class="reference internal" href="../howto/rnn/index_cn.html">RNN模型</a><ul>
<li class="toctree-l3"><a class="reference internal" href="../howto/rnn/rnn_config_cn.html">RNN配置</a></li>
<li class="toctree-l3"><a class="reference internal" href="../howto/rnn/recurrent_group_cn.html">Recurrent Group教程</a></li>
<li class="toctree-l3"><a class="reference internal" href="../howto/rnn/hierarchical_layer_cn.html">支持双层序列作为输入的Layer</a></li>
<li class="toctree-l3"><a class="reference internal" href="../howto/rnn/hrnn_rnn_api_compare_cn.html">单双层RNN API对比介绍</a></li>
</ul>
</li>
<li class="toctree-l2"><a class="reference internal" href="../howto/optimization/gpu_profiling_cn.html">GPU性能调优</a></li>
</ul>
</li>
<li class="toctree-l1"><a class="reference internal" href="../dev/index_cn.html">开发标准</a><ul>
<li class="toctree-l2"><a class="reference internal" href="../dev/contribute_to_paddle_cn.html">如何贡献代码</a></li>
<li class="toctree-l2"><a class="reference internal" href="../dev/write_docs_cn.html">如何贡献文档</a></li>
</ul>
</li>
<li class="toctree-l1"><a class="reference internal" href="../faq/index_cn.html">FAQ</a><ul>
<li class="toctree-l2"><a class="reference internal" href="../faq/build_and_install/index_cn.html">编译安装与单元测试</a></li>
<li class="toctree-l2"><a class="reference internal" href="../faq/model/index_cn.html">模型配置</a></li>
<li class="toctree-l2"><a class="reference internal" href="../faq/parameter/index_cn.html">参数设置</a></li>
<li class="toctree-l2"><a class="reference internal" href="../faq/local/index_cn.html">本地训练与预测</a></li>
<li class="toctree-l2"><a class="reference internal" href="../faq/cluster/index_cn.html">集群训练与预测</a></li>
</ul>
</li>
</ul>

</nav>

        </div>
      </div>
    </nav>

    <section data-toggle="wy-nav-shift" class="wy-nav-content-wrap">

      
      <nav class="wy-nav-top" role="navigation" aria-label="top navigation">
        <i data-toggle="wy-nav-top" class="fa fa-bars"></i>
        <a href="../index_cn.html">PaddlePaddle</a>
      </nav>


      <div class="wy-nav-content">
        <div class="rst-content">
          

<div role="navigation" aria-label="breadcrumbs navigation">
  <ul class="wy-breadcrumbs">
    <li><a href="../index_cn.html">Docs</a> &raquo;</li>
      
    <li>Design Doc: float16</li>
      <li class="wy-breadcrumbs-aside">
        
          
            <a href="../_sources/design/float16.md.txt" rel="nofollow"> View page source</a>
          
        
      </li>
  </ul>
  <hr/>
</div>
          <div role="main" class="document" itemscope="itemscope" itemtype="http://schema.org/Article">
           <div itemprop="articleBody">
            
  <div class="section" id="design-doc-float16">
<span id="design-doc-float16"></span><h1>Design Doc: float16<a class="headerlink" href="#design-doc-float16" title="永久链接至标题">¶</a></h1>
<div class="section" id="why-float16">
<span id="why-float16"></span><h2>Why float16<a class="headerlink" href="#why-float16" title="永久链接至标题">¶</a></h2>
<p>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.</p>
<p>When high precision computation is not required, using float16 data type could potentially</p>
<ul class="simple">
<li>reduce storage space, memory bandwidth, and power usages;</li>
<li>increase the chance of data fitting into a smaller cache of lower latency;</li>
<li>provide arithmetic speed up if supported by hardware.</li>
</ul>
</div>
<div class="section" id="survey-of-current-float16-support">
<span id="survey-of-current-float16-support"></span><h2>Survey of current float16 support<a class="headerlink" href="#survey-of-current-float16-support" title="永久链接至标题">¶</a></h2>
<p>A brief survey of float16 support on different compilers, hardwares, and libraries can be found below. Interested readers can refer to <a class="reference external" href="https://github.com/PaddlePaddle/Paddle/issues/4853">link1</a> and <a class="reference external" href="https://github.com/Xreki/Xreki.github.io/blob/master/multi_data_types_in_dl_framework/ppt/float16_and_quantized_type.md">link2</a> for more info.</p>
<p>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 <code class="docutils literal"><span class="pre">__half</span></code> for cuda, <code class="docutils literal"><span class="pre">float16_t</span></code> for ARM, and <code class="docutils literal"><span class="pre">Eigen::half</span></code> for Eigen to make writing customized float16 kernels easier.</p>
<div class="section" id="compiler">
<span id="compiler"></span><h3>Compiler<a class="headerlink" href="#compiler" title="永久链接至标题">¶</a></h3>
<ul class="simple">
<li>nvcc supports <code class="docutils literal"><span class="pre">__half</span></code> data type after CUDA 7.5.</li>
<li><code class="docutils literal"><span class="pre">__fp16</span></code> or <code class="docutils literal"><span class="pre">float16_t</span></code> is supported as storage type for gcc &gt;= 6.1 and clang &gt;= 3.4.</li>
<li><code class="docutils literal"><span class="pre">__fp16</span></code> or <code class="docutils literal"><span class="pre">float16_t</span></code> is supported as arithmetic type for gcc &gt;= 7.1 and clang &gt;= 3.9.</li>
</ul>
</div>
<div class="section" id="hardware">
<span id="hardware"></span><h3>Hardware<a class="headerlink" href="#hardware" title="永久链接至标题">¶</a></h3>
<ul class="simple">
<li><code class="docutils literal"><span class="pre">__half</span></code> is supported on GPU with compute capability &gt;= 5.3.</li>
<li><code class="docutils literal"><span class="pre">__fp16</span></code> is supported as storage type for ARMv7-A, ARMv8-A, and above.</li>
<li><code class="docutils literal"><span class="pre">__fp16</span></code> 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).</li>
</ul>
</div>
<div class="section" id="libraries">
<span id="libraries"></span><h3>Libraries<a class="headerlink" href="#libraries" title="永久链接至标题">¶</a></h3>
<ul class="simple">
<li><a class="reference external" href="https://github.com/RLovelett/eigen">Eigen</a> &gt;= 3.3 supports float16 calculation on both GPU and CPU using the <code class="docutils literal"><span class="pre">Eigen::half</span></code> 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.</li>
<li><a class="reference external" href="https://github.com/ARM-software/ComputeLibrary">ARM compute library</a> &gt;= 17.02.01 supports NEON FP16 kernels (requires ARMv8.2-A CPU).</li>
</ul>
</div>
<div class="section" id="cuda-version-issue">
<span id="cuda-version-issue"></span><h3>CUDA version issue<a class="headerlink" href="#cuda-version-issue" title="永久链接至标题">¶</a></h3>
<p>There are currently three versions of CUDA that supports <code class="docutils literal"><span class="pre">__half</span></code> data type, namely, CUDA 7.5, 8.0, and 9.0.
CUDA 7.5 and 8.0 define <code class="docutils literal"><span class="pre">__half</span></code> as a simple struct that has a <code class="docutils literal"><span class="pre">uint16_t</span></code> data (see <a class="reference external" href="https://github.com/ptillet/isaac/blob/9212ab5a3ddbe48f30ef373f9c1fb546804c7a8c/include/isaac/external/CUDA/cuda_fp16.h"><code class="docutils literal"><span class="pre">cuda_fp16.h</span></code></a>) as follows:</p>
<div class="highlight-default"><div class="highlight"><pre><span></span><span class="n">typedef</span> <span class="n">struct</span> <span class="n">__align__</span><span class="p">(</span><span class="mi">2</span><span class="p">)</span> <span class="p">{</span>
   <span class="n">unsigned</span> <span class="n">short</span> <span class="n">x</span><span class="p">;</span>
<span class="p">}</span> <span class="n">__half</span><span class="p">;</span>

<span class="n">typedef</span> <span class="n">__half</span> <span class="n">half</span><span class="p">;</span>
</pre></div>
</div>
<p>This struct does not define any overloaded arithmetic operators. So you have to directly use <code class="docutils literal"><span class="pre">__hadd</span></code> instead of <code class="docutils literal"><span class="pre">+</span></code> to correctly add two half types:</p>
<div class="highlight-default"><div class="highlight"><pre><span></span><span class="n">__global__</span> <span class="n">void</span> <span class="n">Add</span><span class="p">()</span> <span class="p">{</span>
  <span class="n">half</span> <span class="n">a</span><span class="p">,</span> <span class="n">b</span><span class="p">,</span> <span class="n">c</span><span class="p">;</span>
  <span class="n">c</span> <span class="o">=</span> <span class="n">__hadd</span><span class="p">(</span><span class="n">a</span><span class="p">,</span> <span class="n">b</span><span class="p">);</span> <span class="o">//</span> <span class="n">correct</span>
  <span class="n">c</span> <span class="o">=</span> <span class="n">a</span> <span class="o">+</span> <span class="n">b</span><span class="p">;</span> <span class="o">//</span> <span class="n">compiler</span> <span class="n">error</span><span class="p">:</span> <span class="n">no</span> <span class="n">operator</span> <span class="s2">&quot;+&quot;</span> <span class="n">matches</span> <span class="n">these</span> <span class="n">operands</span>
<span class="p">}</span>
</pre></div>
</div>
<p>CUDA 9.0 provides a major update to the half data type. The related code can be found in the updated <a class="reference external" href="https://github.com/ptillet/isaac/blob/master/include/isaac/external/CUDA/cuda_fp16.h"><code class="docutils literal"><span class="pre">cuda_fp16.h</span></code></a> and the newly added <a class="reference external" href="https://github.com/ptillet/isaac/blob/master/include/isaac/external/CUDA/cuda_fp16.hpp"><code class="docutils literal"><span class="pre">cuda_fp16.hpp</span></code></a>.</p>
<p>Essentially, CUDA 9.0 renames the original <code class="docutils literal"><span class="pre">__half</span></code> type in 7.5 and 8.0 as <code class="docutils literal"><span class="pre">__half_raw</span></code>, and defines a new <code class="docutils literal"><span class="pre">__half</span></code> class type that has constructors, conversion operators, and also provides overloaded arithmetic operators such as follows:</p>
<div class="highlight-default"><div class="highlight"><pre><span></span><span class="n">typedef</span> <span class="n">struct</span> <span class="n">__CUDA_ALIGN__</span><span class="p">(</span><span class="mi">2</span><span class="p">)</span> <span class="p">{</span>
    <span class="n">unsigned</span> <span class="n">short</span> <span class="n">x</span><span class="p">;</span>
<span class="p">}</span> <span class="n">__half_raw</span><span class="p">;</span>


<span class="n">struct</span> <span class="n">__CUDA_ALIGN__</span><span class="p">(</span><span class="mi">2</span><span class="p">)</span> <span class="n">__half</span> <span class="p">{</span>
<span class="n">protected</span><span class="p">:</span>
    <span class="n">unsigned</span> <span class="n">short</span> <span class="n">__x</span><span class="p">;</span>
<span class="n">public</span><span class="p">:</span>
    <span class="o">//</span> <span class="n">constructors</span> <span class="ow">and</span> <span class="n">conversion</span> <span class="n">operators</span> <span class="n">from</span><span class="o">/</span><span class="n">to</span> 
    <span class="o">//</span> <span class="n">__half_raw</span> <span class="ow">and</span> <span class="n">other</span> <span class="n">built</span><span class="o">-</span><span class="ow">in</span> <span class="n">data</span> <span class="n">types</span>
<span class="p">}</span>

<span class="n">typedef</span> <span class="n">__half</span> <span class="n">half</span><span class="p">;</span>

<span class="n">__device__</span> <span class="n">__forceinline__</span> 
<span class="n">__half</span> <span class="n">operator</span><span class="o">+</span><span class="p">(</span><span class="n">const</span> <span class="n">__half</span> <span class="o">&amp;</span><span class="n">lh</span><span class="p">,</span> <span class="n">const</span> <span class="n">__half</span> <span class="o">&amp;</span><span class="n">rh</span><span class="p">)</span> <span class="p">{</span> 
    <span class="k">return</span> <span class="n">__hadd</span><span class="p">(</span><span class="n">lh</span><span class="p">,</span> <span class="n">rh</span><span class="p">);</span> 
<span class="p">}</span>

<span class="o">//</span> <span class="n">Other</span> <span class="n">overloaded</span> <span class="n">operators</span>
</pre></div>
</div>
<p>This new design makes <code class="docutils literal"><span class="pre">c</span> <span class="pre">=</span> <span class="pre">a</span> <span class="pre">+</span> <span class="pre">b</span></code> work correctly for CUDA half data type.</p>
</div>
</div>
<div class="section" id="implementation">
<span id="implementation"></span><h2>Implementation<a class="headerlink" href="#implementation" title="永久链接至标题">¶</a></h2>
<p>The float16 class holds a 16-bit <code class="docutils literal"><span class="pre">uint16_t</span></code> data internally.</p>
<div class="highlight-default"><div class="highlight"><pre><span></span><span class="n">struct</span> <span class="n">float16</span> <span class="p">{</span>
  <span class="n">uint16_t</span> <span class="n">x</span><span class="p">;</span>
<span class="p">};</span>
</pre></div>
</div>
<p>float16 supports the following features:</p>
<ul class="simple">
<li>constructors / assignment operators that take input from primitive data types including bool, integers of various length, float, and double.</li>
<li>constructors / assignment operators that take input from <code class="docutils literal"><span class="pre">__half</span></code> on cuda, <code class="docutils literal"><span class="pre">float16_t</span></code> on ARM, and <code class="docutils literal"><span class="pre">Eigen::half</span></code> on Eigen.</li>
<li>conversion operators to primitive data types and half precision data types on cuda, ARM and Eigen.</li>
<li>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.</li>
</ul>
<p>To support the above features, two fundamental conversion functions are provided:</p>
<div class="highlight-default"><div class="highlight"><pre><span></span><span class="n">float16</span> <span class="n">float_to_half_rn</span><span class="p">(</span><span class="nb">float</span> <span class="n">f</span><span class="p">);</span>  <span class="o">//</span> <span class="n">convert</span> <span class="n">to</span> <span class="n">half</span> <span class="n">precision</span> <span class="ow">in</span> <span class="nb">round</span><span class="o">-</span><span class="n">to</span><span class="o">-</span><span class="n">nearest</span><span class="o">-</span><span class="n">even</span> <span class="n">mode</span>
<span class="nb">float</span> <span class="n">half_to_float</span><span class="p">(</span><span class="n">float16</span> <span class="n">h</span><span class="p">);</span>
</pre></div>
</div>
<p>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.</p>
</div>
<div class="section" id="to-do">
<span id="to-do"></span><h2>To do<a class="headerlink" href="#to-do" title="永久链接至标题">¶</a></h2>
<p>After float16 class is available, some of the future items are below:</p>
<ul class="simple">
<li>Update pybind/tensor_py.h to bind c++ float16 with numpy float16.</li>
<li>Modify <code class="docutils literal"><span class="pre">GetKernelType()</span></code> method in <code class="docutils literal"><span class="pre">framework/operator.h</span></code> to make it compatible with float16.</li>
<li>Create a type-casting operator that can convert the data type in tensor between float16 and other types.</li>
</ul>
</div>
</div>


           </div>
          </div>
          <footer>
  

  <hr/>

  <div role="contentinfo">
    <p>
        &copy; Copyright 2016, PaddlePaddle developers.

    </p>
  </div>
  Built with <a href="http://sphinx-doc.org/">Sphinx</a> using a <a href="https://github.com/snide/sphinx_rtd_theme">theme</a> provided by <a href="https://readthedocs.org">Read the Docs</a>. 

</footer>

        </div>
      </div>

    </section>

  </div>
  

    <script type="text/javascript">
        var DOCUMENTATION_OPTIONS = {
            URL_ROOT:'../',
            VERSION:'',
            COLLAPSE_INDEX:false,
            FILE_SUFFIX:'.html',
            HAS_SOURCE:  true
        };
    </script>
      <script type="text/javascript" src="../_static/jquery.js"></script>
      <script type="text/javascript" src="../_static/underscore.js"></script>
      <script type="text/javascript" src="../_static/doctools.js"></script>
      <script type="text/javascript" src="../_static/translations.js"></script>
      <script type="text/javascript" src="https://cdn.bootcss.com/mathjax/2.7.0/MathJax.js"></script>

  
    <script type="text/javascript" src="../_static/js/theme.js"></script>
  

  <script type="text/javascript">
      jQuery(function () {
          SphinxRtdTheme.StickyNav.enable();
      });
  </script>
   

</body>
</html>