| <!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Transitional//EN" "https://www.w3.org/TR/xhtml1/DTD/xhtml1-transitional.dtd"> |
| <html xmlns="http://www.w3.org/1999/xhtml"> |
| <head> |
| <meta http-equiv="Content-Type" content="text/xhtml;charset=UTF-8"/> |
| <meta http-equiv="X-UA-Compatible" content="IE=9"/> |
| <meta name="generator" content="Doxygen 1.8.17"/> |
| <meta name="viewport" content="width=device-width, initial-scale=1"/> |
| <title>mxnet: /work/mxnet/src/common/cuda/rtc/half-inl.h Source File</title> |
| <link href="tabs.css" rel="stylesheet" type="text/css"/> |
| <script type="text/javascript" src="jquery.js"></script> |
| <script type="text/javascript" src="dynsections.js"></script> |
| <link href="search/search.css" rel="stylesheet" type="text/css"/> |
| <script type="text/javascript" src="search/searchdata.js"></script> |
| <script type="text/javascript" src="search/search.js"></script> |
| <link href="doxygen.css" rel="stylesheet" type="text/css" /> |
| </head> |
| <body> |
| <div id="top"><!-- do not remove this div, it is closed by doxygen! --> |
| <div id="titlearea"> |
| <table cellspacing="0" cellpadding="0"> |
| <tbody> |
| <tr style="height: 56px;"> |
| <td id="projectalign" style="padding-left: 0.5em;"> |
| <div id="projectname">mxnet |
| </div> |
| </td> |
| </tr> |
| </tbody> |
| </table> |
| </div> |
| <!-- end header part --> |
| <!-- Generated by Doxygen 1.8.17 --> |
| <script type="text/javascript"> |
| /* @license magnet:?xt=urn:btih:cf05388f2679ee054f2beb29a391d25f4e673ac3&dn=gpl-2.0.txt GPL-v2 */ |
| var searchBox = new SearchBox("searchBox", "search",false,'Search'); |
| /* @license-end */ |
| </script> |
| <script type="text/javascript" src="menudata.js"></script> |
| <script type="text/javascript" src="menu.js"></script> |
| <script type="text/javascript"> |
| /* @license magnet:?xt=urn:btih:cf05388f2679ee054f2beb29a391d25f4e673ac3&dn=gpl-2.0.txt GPL-v2 */ |
| $(function() { |
| initMenu('',true,false,'search.php','Search'); |
| $(document).ready(function() { init_search(); }); |
| }); |
| /* @license-end */</script> |
| <div id="main-nav"></div> |
| <!-- window showing the filter options --> |
| <div id="MSearchSelectWindow" |
| onmouseover="return searchBox.OnSearchSelectShow()" |
| onmouseout="return searchBox.OnSearchSelectHide()" |
| onkeydown="return searchBox.OnSearchSelectKey(event)"> |
| </div> |
| |
| <!-- iframe showing the search results (closed by default) --> |
| <div id="MSearchResultsWindow"> |
| <iframe src="javascript:void(0)" frameborder="0" |
| name="MSearchResults" id="MSearchResults"> |
| </iframe> |
| </div> |
| |
| <div id="nav-path" class="navpath"> |
| <ul> |
| <li class="navelem"><a class="el" href="dir_68267d1309a1af8e8297ef4c3efbcdba.html">src</a></li><li class="navelem"><a class="el" href="dir_fdedb0aba14d44ce9d99bc100e026e6a.html">common</a></li><li class="navelem"><a class="el" href="dir_5d530576593496167de63f3f304bdbc7.html">cuda</a></li><li class="navelem"><a class="el" href="dir_24ef4685085d6522cb17f4d54bfb8994.html">rtc</a></li> </ul> |
| </div> |
| </div><!-- top --> |
| <div class="header"> |
| <div class="headertitle"> |
| <div class="title">half-inl.h</div> </div> |
| </div><!--header--> |
| <div class="contents"> |
| <a href="half-inl_8h.html">Go to the documentation of this file.</a><div class="fragment"><div class="line"><a name="l00001"></a><span class="lineno"> 1</span> <span class="comment">/*</span></div> |
| <div class="line"><a name="l00002"></a><span class="lineno"> 2</span> <span class="comment"> * Licensed to the Apache Software Foundation (ASF) under one</span></div> |
| <div class="line"><a name="l00003"></a><span class="lineno"> 3</span> <span class="comment"> * or more contributor license agreements. See the NOTICE file</span></div> |
| <div class="line"><a name="l00004"></a><span class="lineno"> 4</span> <span class="comment"> * distributed with this work for additional information</span></div> |
| <div class="line"><a name="l00005"></a><span class="lineno"> 5</span> <span class="comment"> * regarding copyright ownership. The ASF licenses this file</span></div> |
| <div class="line"><a name="l00006"></a><span class="lineno"> 6</span> <span class="comment"> * to you under the Apache License, Version 2.0 (the</span></div> |
| <div class="line"><a name="l00007"></a><span class="lineno"> 7</span> <span class="comment"> * "License"); you may not use this file except in compliance</span></div> |
| <div class="line"><a name="l00008"></a><span class="lineno"> 8</span> <span class="comment"> * with the License. You may obtain a copy of the License at</span></div> |
| <div class="line"><a name="l00009"></a><span class="lineno"> 9</span> <span class="comment"> *</span></div> |
| <div class="line"><a name="l00010"></a><span class="lineno"> 10</span> <span class="comment"> * http://www.apache.org/licenses/LICENSE-2.0</span></div> |
| <div class="line"><a name="l00011"></a><span class="lineno"> 11</span> <span class="comment"> *</span></div> |
| <div class="line"><a name="l00012"></a><span class="lineno"> 12</span> <span class="comment"> * Unless required by applicable law or agreed to in writing,</span></div> |
| <div class="line"><a name="l00013"></a><span class="lineno"> 13</span> <span class="comment"> * software distributed under the License is distributed on an</span></div> |
| <div class="line"><a name="l00014"></a><span class="lineno"> 14</span> <span class="comment"> * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY</span></div> |
| <div class="line"><a name="l00015"></a><span class="lineno"> 15</span> <span class="comment"> * KIND, either express or implied. See the License for the</span></div> |
| <div class="line"><a name="l00016"></a><span class="lineno"> 16</span> <span class="comment"> * specific language governing permissions and limitations</span></div> |
| <div class="line"><a name="l00017"></a><span class="lineno"> 17</span> <span class="comment"> * under the License.</span></div> |
| <div class="line"><a name="l00018"></a><span class="lineno"> 18</span> <span class="comment"> */</span></div> |
| <div class="line"><a name="l00019"></a><span class="lineno"> 19</span>  </div> |
| <div class="line"><a name="l00020"></a><span class="lineno"> 20</span> <span class="preprocessor">#ifndef MXNET_COMMON_CUDA_RTC_HALF_INL_H_</span></div> |
| <div class="line"><a name="l00021"></a><span class="lineno"> 21</span> <span class="preprocessor">#define MXNET_COMMON_CUDA_RTC_HALF_INL_H_</span></div> |
| <div class="line"><a name="l00022"></a><span class="lineno"> 22</span>  </div> |
| <div class="line"><a name="l00023"></a><span class="lineno"> 23</span> <span class="preprocessor">#if MXNET_USE_CUDA</span></div> |
| <div class="line"><a name="l00024"></a><span class="lineno"> 24</span>  </div> |
| <div class="line"><a name="l00025"></a><span class="lineno"> 25</span> <span class="keyword">namespace </span><a class="code" href="namespacemxnet.html">mxnet</a> {</div> |
| <div class="line"><a name="l00026"></a><span class="lineno"> 26</span> <span class="keyword">namespace </span>common {</div> |
| <div class="line"><a name="l00027"></a><span class="lineno"> 27</span> <span class="keyword">namespace </span>cuda {</div> |
| <div class="line"><a name="l00028"></a><span class="lineno"> 28</span> <span class="keyword">namespace </span>rtc {</div> |
| <div class="line"><a name="l00029"></a><span class="lineno"> 29</span>  </div> |
| <div class="line"><a name="l00030"></a><span class="lineno"><a class="line" href="namespacemxnet_1_1common_1_1cuda_1_1rtc.html#ad3ebfc4af0e3ca45f58ed000f71f2591"> 30</a></span> <span class="keyword">const</span> <span class="keywordtype">char</span> <a class="code" href="namespacemxnet_1_1common_1_1cuda_1_1rtc.html#ad3ebfc4af0e3ca45f58ed000f71f2591">fp16_support_string</a>[] = R<span class="stringliteral">"code(</span></div> |
| <div class="line"><a name="l00031"></a><span class="lineno"> 31</span> <span class="stringliteral">struct __align__(2) __half {</span></div> |
| <div class="line"><a name="l00032"></a><span class="lineno"> 32</span> <span class="stringliteral"> __host__ __device__ __half() : __x(0) { }</span></div> |
| <div class="line"><a name="l00033"></a><span class="lineno"> 33</span> <span class="stringliteral"> unsigned short __x;</span></div> |
| <div class="line"><a name="l00034"></a><span class="lineno"> 34</span> <span class="stringliteral">};</span></div> |
| <div class="line"><a name="l00035"></a><span class="lineno"> 35</span> <span class="stringliteral">/* Definitions of intrinsics */</span></div> |
| <div class="line"><a name="l00036"></a><span class="lineno"> 36</span> <span class="stringliteral">__device__ inline __half __float2half(const float f) {</span></div> |
| <div class="line"><a name="l00037"></a><span class="lineno"> 37</span> <span class="stringliteral"> __half val;</span></div> |
| <div class="line"><a name="l00038"></a><span class="lineno"> 38</span> <span class="stringliteral"> asm("{ cvt.rn.f16.f32 %0, %1;}\n" : "=h"(val.__x) : "f"(f));</span></div> |
| <div class="line"><a name="l00039"></a><span class="lineno"> 39</span> <span class="stringliteral"> return val;</span></div> |
| <div class="line"><a name="l00040"></a><span class="lineno"> 40</span> <span class="stringliteral">}</span></div> |
| <div class="line"><a name="l00041"></a><span class="lineno"> 41</span> <span class="stringliteral">__device__ inline float __half2float(const __half h) {</span></div> |
| <div class="line"><a name="l00042"></a><span class="lineno"> 42</span> <span class="stringliteral"> float val;</span></div> |
| <div class="line"><a name="l00043"></a><span class="lineno"> 43</span> <span class="stringliteral"> asm("{ cvt.f32.f16 %0, %1;}\n" : "=f"(val) : "h"(h.__x));</span></div> |
| <div class="line"><a name="l00044"></a><span class="lineno"> 44</span> <span class="stringliteral"> return val;</span></div> |
| <div class="line"><a name="l00045"></a><span class="lineno"> 45</span> <span class="stringliteral">}</span></div> |
| <div class="line"><a name="l00046"></a><span class="lineno"> 46</span> <span class="stringliteral"></span> </div> |
| <div class="line"><a name="l00047"></a><span class="lineno"> 47</span> <span class="stringliteral">typedef __half half;</span></div> |
| <div class="line"><a name="l00048"></a><span class="lineno"> 48</span> <span class="stringliteral"></span> </div> |
| <div class="line"><a name="l00049"></a><span class="lineno"> 49</span> <span class="stringliteral">template <typename DType></span></div> |
| <div class="line"><a name="l00050"></a><span class="lineno"> 50</span> <span class="stringliteral">struct AccType {</span></div> |
| <div class="line"><a name="l00051"></a><span class="lineno"> 51</span> <span class="stringliteral"> using type = DType;</span></div> |
| <div class="line"><a name="l00052"></a><span class="lineno"> 52</span> <span class="stringliteral"></span> </div> |
| <div class="line"><a name="l00053"></a><span class="lineno"> 53</span> <span class="stringliteral"> __device__ static inline type from(const DType& val) {</span></div> |
| <div class="line"><a name="l00054"></a><span class="lineno"> 54</span> <span class="stringliteral"> return val;</span></div> |
| <div class="line"><a name="l00055"></a><span class="lineno"> 55</span> <span class="stringliteral"> }</span></div> |
| <div class="line"><a name="l00056"></a><span class="lineno"> 56</span> <span class="stringliteral"></span> </div> |
| <div class="line"><a name="l00057"></a><span class="lineno"> 57</span> <span class="stringliteral"> __device__ static inline DType to(type val) {</span></div> |
| <div class="line"><a name="l00058"></a><span class="lineno"> 58</span> <span class="stringliteral"> return val;</span></div> |
| <div class="line"><a name="l00059"></a><span class="lineno"> 59</span> <span class="stringliteral"> }</span></div> |
| <div class="line"><a name="l00060"></a><span class="lineno"> 60</span> <span class="stringliteral"></span> </div> |
| <div class="line"><a name="l00061"></a><span class="lineno"> 61</span> <span class="stringliteral">};</span></div> |
| <div class="line"><a name="l00062"></a><span class="lineno"> 62</span> <span class="stringliteral"></span> </div> |
| <div class="line"><a name="l00063"></a><span class="lineno"> 63</span> <span class="stringliteral">template<></span></div> |
| <div class="line"><a name="l00064"></a><span class="lineno"> 64</span> <span class="stringliteral">struct AccType<half> {</span></div> |
| <div class="line"><a name="l00065"></a><span class="lineno"> 65</span> <span class="stringliteral"> using type = float;</span></div> |
| <div class="line"><a name="l00066"></a><span class="lineno"> 66</span> <span class="stringliteral"></span> </div> |
| <div class="line"><a name="l00067"></a><span class="lineno"> 67</span> <span class="stringliteral"> __device__ static inline type from(const half& val) {</span></div> |
| <div class="line"><a name="l00068"></a><span class="lineno"> 68</span> <span class="stringliteral"> return __half2float(val);</span></div> |
| <div class="line"><a name="l00069"></a><span class="lineno"> 69</span> <span class="stringliteral"> }</span></div> |
| <div class="line"><a name="l00070"></a><span class="lineno"> 70</span> <span class="stringliteral"></span> </div> |
| <div class="line"><a name="l00071"></a><span class="lineno"> 71</span> <span class="stringliteral"> __device__ static inline half to(type val) {</span></div> |
| <div class="line"><a name="l00072"></a><span class="lineno"> 72</span> <span class="stringliteral"> return __float2half(val);</span></div> |
| <div class="line"><a name="l00073"></a><span class="lineno"> 73</span> <span class="stringliteral"> }</span></div> |
| <div class="line"><a name="l00074"></a><span class="lineno"> 74</span> <span class="stringliteral">};</span></div> |
| <div class="line"><a name="l00075"></a><span class="lineno"> 75</span> <span class="stringliteral">)code";</span></div> |
| <div class="line"><a name="l00076"></a><span class="lineno"> 76</span> <span class="stringliteral"></span> </div> |
| <div class="line"><a name="l00077"></a><span class="lineno"> 77</span> <span class="stringliteral">} </span><span class="comment">// namespace rtc</span></div> |
| <div class="line"><a name="l00078"></a><span class="lineno"> 78</span> } <span class="comment">// namespace cuda</span></div> |
| <div class="line"><a name="l00079"></a><span class="lineno"> 79</span> } <span class="comment">// namespace common</span></div> |
| <div class="line"><a name="l00080"></a><span class="lineno"> 80</span> } <span class="comment">// namespace mxnet</span></div> |
| <div class="line"><a name="l00081"></a><span class="lineno"> 81</span>  </div> |
| <div class="line"><a name="l00082"></a><span class="lineno"> 82</span> <span class="preprocessor">#endif // MXNET_USE_CUDA</span></div> |
| <div class="line"><a name="l00083"></a><span class="lineno"> 83</span>  </div> |
| <div class="line"><a name="l00084"></a><span class="lineno"> 84</span> <span class="preprocessor">#endif // MXNET_COMMON_CUDA_RTC_HALF_INL_H_</span></div> |
| </div><!-- fragment --></div><!-- contents --> |
| <div class="ttc" id="anamespacemxnet_html"><div class="ttname"><a href="namespacemxnet.html">mxnet</a></div><div class="ttdoc">namespace of mxnet</div><div class="ttdef"><b>Definition:</b> api_registry.h:33</div></div> |
| <div class="ttc" id="anamespacemxnet_1_1common_1_1cuda_1_1rtc_html_ad3ebfc4af0e3ca45f58ed000f71f2591"><div class="ttname"><a href="namespacemxnet_1_1common_1_1cuda_1_1rtc.html#ad3ebfc4af0e3ca45f58ed000f71f2591">mxnet::common::cuda::rtc::fp16_support_string</a></div><div class="ttdeci">const char fp16_support_string[]</div><div class="ttdef"><b>Definition:</b> half-inl.h:30</div></div> |
| <!-- start footer part --> |
| <hr class="footer"/><address class="footer"><small> |
| Generated on Thu Jan 5 2023 03:47:40 for mxnet by  <a href="http://www.doxygen.org/index.html"> |
| <img class="footer" src="doxygen.png" alt="doxygen"/> |
| </a> 1.8.17 |
| </small></address> |
| </body> |
| </html> |
