Skip to content

P
Paddle

PaddlePaddle / Paddle
大约 1 年 前同步成功

通知 2299
Star 20931
Fork 5422
P
Paddle
切换分支/标签
查找文件 普通视图历史永久链接Permalink
float16.html 26.7 KB
Newer Older
T
Deploy to GitHub Pages: 4d422156d42ee21e11656937401cae0081e3c1a5  
Travis CI 已提交
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35 36 37 38 39 40 41 42 43 44 45 46 47 48 49 50 51 52 53 54 55 56 57 58 59 60 61 62 63 64 65 66 67 68 69 70 71 72 73 74 75 76 77 78 79 80 81 82 83 84 85 86 87 88 89 90 91 92 93 94 95 96 97 98 99 100 101 102 103 104 105 106 107 108 109 110 111 


<!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"/> 

  <link rel="stylesheet" href="https://cdn.jsdelivr.net/perfect-scrollbar/0.6.14/css/perfect-scrollbar.min.css" type="text/css" />
  <link rel="stylesheet" href="../_static/css/override.css" type="text/css" />
  <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">

  
  <header class="site-header">
    <div class="site-logo">
      <a href="/"><img src="../_static/images/PP_w.png"></a>
    </div>
    <div class="site-nav-links">
      <div class="site-menu">
        <a class="fork-on-github" href="https://github.com/PaddlePaddle/Paddle" target="_blank"><i class="fa fa-github"></i>Fork me on Github</a>
        <div class="language-switcher dropdown">
          <a type="button" data-toggle="dropdown">
            <span>English</span>
            <i class="fa fa-angle-up"></i>
            <i class="fa fa-angle-down"></i>
          </a>
          <ul class="dropdown-menu">
            <li><a href="/doc_cn">中文</a></li>
            <li><a href="/doc">English</a></li>
          </ul>
        </div>
        <ul class="site-page-links">
          <li><a href="/">Home</a></li>
        </ul>
      </div>
      <div class="doc-module">
        
        <ul>
<li class="toctree-l1"><a class="reference internal" href="../getstarted/index_cn.html">新手入门</a></li>
<li class="toctree-l1"><a class="reference internal" href="../howto/index_cn.html">进阶指南</a></li>
<li class="toctree-l1"><a class="reference internal" href="../api/index_cn.html">API</a></li>
<li class="toctree-l1"><a class="reference internal" href="../faq/index_cn.html">FAQ</a></li>
<li class="toctree-l1"><a class="reference internal" href="../mobile/index_cn.html">MOBILE</a></li>
</ul>

        
<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>
  </header>
  
  <div class="main-content-wrap">

    
    <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/build_and_install/index_cn.html">安装与编译</a><ul>
T
Deploy to GitHub Pages: 0c0ff2828ccedb51db23290d6df9e4c83839d6af  
Travis CI 已提交
112 113 
<li class="toctree-l3"><a class="reference internal" href="../getstarted/build_and_install/pip_install_cn.html">使用pip安装</a></li>
<li class="toctree-l3"><a class="reference internal" href="../getstarted/build_and_install/docker_install_cn.html">使用Docker安装运行</a></li>
T
Deploy to GitHub Pages: 458ffbf4fe7fe6da5e2ea133c027d560815cc3cc  
Travis CI 已提交
114 
<li class="toctree-l3"><a class="reference internal" href="../howto/dev/build_cn.html">用Docker编译和测试PaddlePaddle</a></li>
T
Deploy to GitHub Pages: 0c0ff2828ccedb51db23290d6df9e4c83839d6af  
Travis CI 已提交
115 
<li class="toctree-l3"><a class="reference internal" href="../getstarted/build_and_install/build_from_source_cn.html">从源码编译</a></li>
T
Deploy to GitHub Pages: 4d422156d42ee21e11656937401cae0081e3c1a5  
Travis CI 已提交
116 117 118 119 120 121 122 123 124 125 126 127 
</ul>
</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="../howto/index_cn.html">进阶指南</a><ul>
<li class="toctree-l2"><a class="reference internal" href="../howto/usage/cmd_parameter/index_cn.html">设置命令行参数</a><ul>
<li class="toctree-l3"><a class="reference internal" href="../howto/usage/cmd_parameter/use_case_cn.html">使用案例</a></li>
<li class="toctree-l3"><a class="reference internal" href="../howto/usage/cmd_parameter/arguments_cn.html">参数概述</a></li>
<li class="toctree-l3"><a class="reference internal" href="../howto/usage/cmd_parameter/detail_introduction_cn.html">细节描述</a></li>
</ul>
</li>
T
Deploy to GitHub Pages: 33e074dab85d23bccdb1876df6490a7a5514e395  
Travis CI 已提交
128 129 130 131 132 133 
<li class="toctree-l2"><a class="reference internal" href="../howto/usage/cluster/cluster_train_cn.html">分布式训练</a><ul>
<li class="toctree-l3"><a class="reference internal" href="../howto/usage/cluster/fabric_cn.html">fabric集群</a></li>
<li class="toctree-l3"><a class="reference internal" href="../howto/usage/cluster/openmpi_cn.html">openmpi集群</a></li>
<li class="toctree-l3"><a class="reference internal" href="../howto/usage/cluster/k8s_cn.html">kubernetes单机</a></li>
<li class="toctree-l3"><a class="reference internal" href="../howto/usage/cluster/k8s_distributed_cn.html">kubernetes distributed分布式</a></li>
<li class="toctree-l3"><a class="reference internal" href="../howto/usage/cluster/k8s_aws_cn.html">AWS上运行kubernetes集群训练</a></li>
T
Deploy to GitHub Pages: 16fd9f18e304043a1035ad237884a93be354804c  
Travis CI 已提交
134 135 
</ul>
</li>
T
Deploy to GitHub Pages: 70ee79c2567d5128a9d527b444fcee30f83c02a5  
Travis CI 已提交
136 
<li class="toctree-l2"><a class="reference internal" href="../howto/dev/contribute_to_paddle_cn.html">如何贡献代码</a></li>
T
Deploy to GitHub Pages: 4d422156d42ee21e11656937401cae0081e3c1a5  
Travis CI 已提交
137 138 139 140 141 142 143 144 145 146 147 148 149 150 151 152 153 154 155 156 157 158 
<li class="toctree-l2"><a class="reference internal" href="../howto/dev/write_docs_cn.html">如何贡献/修改文档</a></li>
<li class="toctree-l2"><a class="reference internal" href="../howto/deep_model/rnn/index_cn.html">RNN相关模型</a><ul>
<li class="toctree-l3"><a class="reference internal" href="../howto/deep_model/rnn/rnn_config_cn.html">RNN配置</a></li>
<li class="toctree-l3"><a class="reference internal" href="../howto/deep_model/rnn/recurrent_group_cn.html">Recurrent Group教程</a></li>
<li class="toctree-l3"><a class="reference internal" href="../howto/deep_model/rnn/hierarchical_layer_cn.html">支持双层序列作为输入的Layer</a></li>
<li class="toctree-l3"><a class="reference internal" href="../howto/deep_model/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="../api/index_cn.html">API</a><ul>
<li class="toctree-l2"><a class="reference internal" href="../api/v2/model_configs.html">模型配置</a><ul>
<li class="toctree-l3"><a class="reference internal" href="../api/v2/config/activation.html">Activation</a></li>
<li class="toctree-l3"><a class="reference internal" href="../api/v2/config/layer.html">Layers</a></li>
<li class="toctree-l3"><a class="reference internal" href="../api/v2/config/evaluators.html">Evaluators</a></li>
<li class="toctree-l3"><a class="reference internal" href="../api/v2/config/optimizer.html">Optimizer</a></li>
<li class="toctree-l3"><a class="reference internal" href="../api/v2/config/pooling.html">Pooling</a></li>
<li class="toctree-l3"><a class="reference internal" href="../api/v2/config/networks.html">Networks</a></li>
<li class="toctree-l3"><a class="reference internal" href="../api/v2/config/attr.html">Parameter Attribute</a></li>
</ul>
</li>
T
Deploy to GitHub Pages: d4d8f742a3bc4366bf410bc9ff29c0c32666bf60  
Travis CI 已提交
159 160 161 162 163 164 
<li class="toctree-l2"><a class="reference internal" href="../api/v2/data.html">数据访问</a><ul>
<li class="toctree-l3"><a class="reference internal" href="../api/v2/data/data_reader.html">Data Reader Interface</a></li>
<li class="toctree-l3"><a class="reference internal" href="../api/v2/data/image.html">Image Interface</a></li>
<li class="toctree-l3"><a class="reference internal" href="../api/v2/data/dataset.html">Dataset</a></li>
</ul>
</li>
T
Deploy to GitHub Pages: 4d422156d42ee21e11656937401cae0081e3c1a5  
Travis CI 已提交
165 
<li class="toctree-l2"><a class="reference internal" href="../api/v2/run_logic.html">训练与应用</a></li>
T
Deploy to GitHub Pages: 5ac72d95014e453c13a389248d2d2361a45a3402  
Travis CI 已提交
166 167 168 169 170 171 172 173 174 175 176 177 178 
<li class="toctree-l2"><a class="reference internal" href="../api/v2/fluid.html">Fluid</a><ul>
<li class="toctree-l3"><a class="reference internal" href="../api/v2/fluid/layers.html">Layers</a></li>
<li class="toctree-l3"><a class="reference internal" href="../api/v2/fluid/data_feeder.html">DataFeeder</a></li>
<li class="toctree-l3"><a class="reference internal" href="../api/v2/fluid/executor.html">Executor</a></li>
<li class="toctree-l3"><a class="reference internal" href="../api/v2/fluid/initializer.html">Initializer</a></li>
<li class="toctree-l3"><a class="reference internal" href="../api/v2/fluid/evaluator.html">Evaluator</a></li>
<li class="toctree-l3"><a class="reference internal" href="../api/v2/fluid/nets.html">Nets</a></li>
<li class="toctree-l3"><a class="reference internal" href="../api/v2/fluid/optimizer.html">Optimizer</a></li>
<li class="toctree-l3"><a class="reference internal" href="../api/v2/fluid/param_attr.html">ParamAttr</a></li>
<li class="toctree-l3"><a class="reference internal" href="../api/v2/fluid/profiler.html">Profiler</a></li>
<li class="toctree-l3"><a class="reference internal" href="../api/v2/fluid/regularizer.html">Regularizer</a></li>
</ul>
</li>
T
Deploy to GitHub Pages: 4d422156d42ee21e11656937401cae0081e3c1a5  
Travis CI 已提交
179 180 181 182 183 184 185 186 187 188 189 
</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>
<li class="toctree-l1"><a class="reference internal" href="../mobile/index_cn.html">MOBILE</a><ul>
T
Deploy to GitHub Pages: c808fbbfcbaaf5c08f6254bfdb860f5dac76a627  
Travis CI 已提交
190 191 192 
<li class="toctree-l2"><a class="reference internal" href="../mobile/cross_compiling_for_android_cn.html">Android平台编译指南</a></li>
<li class="toctree-l2"><a class="reference internal" href="../mobile/cross_compiling_for_ios_cn.html">iOS平台编译指南</a></li>
<li class="toctree-l2"><a class="reference internal" href="../mobile/cross_compiling_for_raspberry_cn.html">Raspberry Pi平台编译指南</a></li>
T
Deploy to GitHub Pages: 4d422156d42ee21e11656937401cae0081e3c1a5  
Travis CI 已提交
193 194 195 196 197 198 199 200 201 202 203 204 205 206 207 208 209 210 211 212 213 214 215 216 217 218 219 220 221 222 223 224 225 226 227 228 229 230 231 232 233 234 235 236 237 238 239 240 241 242 243 244 245 246 247 248 249 250 251 252 253 254 255 256 257 258 259 260 261 262 
</ul>
</li>
</ul>

        
    </nav>
    
    <section class="doc-content-wrap">

      

 







<div role="navigation" aria-label="breadcrumbs navigation">
  <ul class="wy-breadcrumbs">
      
    <li>Design Doc: float16</li>
  </ul>
</div>
      
      <div class="wy-nav-content" id="doc-content">
        <div class="rst-content">
          <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>
T
Deploy to GitHub Pages: 35572355c2261c493aa782ba1255971f4dfa385e  
Travis CI 已提交
263 264 265 266 267 268 269 270 271 272 273 274 275 276 277 278 279 280 281 282 283 284 285 286 287 288 289 290 291 292 293 294 295 296 297 298 299 300 301 302 303 304 305 306 307 308 
<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>
T
Deploy to GitHub Pages: 4d422156d42ee21e11656937401cae0081e3c1a5  
Travis CI 已提交
309 310 311 312 313 314 315 316 317 318 319 320 321 322 323 324 325 326 327 328 329 330 331 332 333 334 335 336 
</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>
T
Deploy to GitHub Pages: 91b724821cae55543d73f2bc79bef13f6c16fc36  
Travis CI 已提交
337 
<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>
T
Deploy to GitHub Pages: 4d422156d42ee21e11656937401cae0081e3c1a5  
Travis CI 已提交
338 339 340 341 342 343 344 345 346 347 348 349 350 351 352 353 354 355 356 357 358 359 360 361 362 363 364 365 366 367 368 369 370 371 372 373 374 375 376 377 378 379 380 381 382 383 384 385 386 387 388 389 390 391 392 393 394 395 396 397 398 399 400 
<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,
            SOURCELINK_SUFFIX: ".txt",
        };
    </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 src="https://maxcdn.bootstrapcdn.com/bootstrap/3.3.7/js/bootstrap.min.js" integrity="sha384-Tc5IQib027qvyjSMfHjOMaLkfuWVxZxUPnCJA7l2mCWNIpG9mGCD8wGNIcPD7Txa" crossorigin="anonymous"></script>
  <script src="https://cdn.jsdelivr.net/perfect-scrollbar/0.6.14/js/perfect-scrollbar.jquery.min.js"></script>
  <script src="../_static/js/paddle_doc_init.js"></script> 

</body>
</html>