AI/rocm-systems
2
0
Fork 0
Κώδικας Ζητήματα Pull Requests Δράσεις Πακέτα Έργα Κυκλοφορίες Wiki Δραστηριότητα

Update doxygen documentation

Περιήγηση Στο Πηγαίο
Αυτή η υποβολή περιλαμβάνεται σε:
Maneesh Gupta
2016-04-19 22:44:58 +05:30
γονέας 148799a371
υποβολή 26c6f9f861
90 αρχεία άλλαξαν με 338 προσθήκες και 343 διαγραφές
Λήψη Αρχείου Patch Λήψη Αρχείου Diff Ανάπτυξη όλων των αρχείων Σύμπτυξη όλων των αρχείων
docs/RuntimeAPI/html/hip__hcc_8h_source.html
+159 -156
Προβολή Αρχείου
@@ -4,7 +4,7 @@
<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.6"/>
<title>HIP: Heterogenous-computing Interface for Portability: /home/mangupta/hip_git/rel0.84.0/include/hcc_detail/hip_hcc.h Source File</title>
<title>HIP: Heterogenous-computing Interface for Portability: /home/mangupta/hip_git/release_0.84.00/include/hcc_detail/hip_hcc.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>
@@ -112,8 +112,8 @@ var searchBox = new SearchBox("searchBox", "search",false,'Search');
<div class="line"><a name="l00021"></a><span class="lineno"> 21</span>&#160;<span class="preprocessor"></span><span class="preprocessor">#define HIP_HCC_H</span></div>
<div class="line"><a name="l00022"></a><span class="lineno"> 22</span>&#160;<span class="preprocessor"></span></div>
<div class="line"><a name="l00023"></a><span class="lineno"> 23</span>&#160;<span class="preprocessor">#include &lt;hc.hpp&gt;</span></div>
<div class="line"><a name="l00024"></a><span class="lineno"> 24</span>&#160;<span class="preprocessor">#include &quot;hcc_detail/hip_util.h&quot;</span></div>
<div class="line"><a name="l00025"></a><span class="lineno"> 25</span>&#160;<span class="preprocessor">#include &quot;hcc_detail/staging_buffer.h&quot;</span></div>
<div class="line"><a name="l00024"></a><span class="lineno"> 24</span>&#160;<span class="preprocessor">#include &quot;hip/hcc_detail/hip_util.h&quot;</span></div>
<div class="line"><a name="l00025"></a><span class="lineno"> 25</span>&#160;<span class="preprocessor">#include &quot;hip/hcc_detail/staging_buffer.h&quot;</span></div>
<div class="line"><a name="l00026"></a><span class="lineno"> 26</span>&#160;</div>
<div class="line"><a name="l00027"></a><span class="lineno"> 27</span>&#160;<span class="preprocessor">#define HIP_HCC</span></div>
<div class="line"><a name="l00028"></a><span class="lineno"> 28</span>&#160;<span class="preprocessor"></span></div>
@@ -622,175 +622,178 @@ var searchBox = new SearchBox("searchBox", "search",false,'Search');
<div class="line"><a name="l00531"></a><span class="lineno"> 531</span>&#160; <span class="keywordtype">bool</span> removePeer(<a class="code" href="classihipDevice__t.html">ihipDevice_t</a> *peer);</div>
<div class="line"><a name="l00532"></a><span class="lineno"> 532</span>&#160; <span class="keywordtype">void</span> resetPeers(<a class="code" href="classihipDevice__t.html">ihipDevice_t</a> *thisDevice);</div>
<div class="line"><a name="l00533"></a><span class="lineno"> 533</span>&#160;</div>
<div class="line"><a name="l00534"></a><span class="lineno"> 534</span>&#160; uint32_t peerCnt()<span class="keyword"> const </span>{ <span class="keywordflow">return</span> _peerCnt; };</div>
<div class="line"><a name="l00535"></a><span class="lineno"> 535</span>&#160; hsa_agent_t *peerAgents()<span class="keyword"> const </span>{ <span class="keywordflow">return</span> _peerAgents; };</div>
<div class="line"><a name="l00534"></a><span class="lineno"> 534</span>&#160;</div>
<div class="line"><a name="l00535"></a><span class="lineno"> 535</span>&#160; <span class="keywordtype">void</span> addStream(<a class="code" href="classihipStream__t.html">ihipStream_t</a> *stream);</div>
<div class="line"><a name="l00536"></a><span class="lineno"> 536</span>&#160;</div>
<div class="line"><a name="l00537"></a><span class="lineno"> 537</span>&#160;</div>
<div class="line"><a name="l00538"></a><span class="lineno"> 538</span>&#160;<span class="keyword">private</span>:</div>
<div class="line"><a name="l00539"></a><span class="lineno"> 539</span>&#160; std::list&lt;ihipStream_t*&gt; _streams; <span class="comment">// streams associated with this device.</span></div>
<div class="line"><a name="l00540"></a><span class="lineno"> 540</span>&#160; ihipStream_t::SeqNum_t _stream_id;</div>
<div class="line"><a name="l00541"></a><span class="lineno"> 541</span>&#160;</div>
<div class="line"><a name="l00542"></a><span class="lineno"> 542</span>&#160; <span class="comment">// These reflect the currently Enabled set of peers for this GPU:</span></div>
<div class="line"><a name="l00543"></a><span class="lineno"> 543</span>&#160; std::list&lt;ihipDevice_t*&gt; _peers; <span class="comment">// list of enabled peer devices.</span></div>
<div class="line"><a name="l00544"></a><span class="lineno"> 544</span>&#160; uint32_t _peerCnt; <span class="comment">// number of enabled peers</span></div>
<div class="line"><a name="l00545"></a><span class="lineno"> 545</span>&#160; hsa_agent_t *_peerAgents; <span class="comment">// efficient packed array of enabled agents (to use for allocations.)</span></div>
<div class="line"><a name="l00546"></a><span class="lineno"> 546</span>&#160;<span class="keyword">private</span>:</div>
<div class="line"><a name="l00547"></a><span class="lineno"> 547</span>&#160; <span class="keywordtype">void</span> recomputePeerAgents();</div>
<div class="line"><a name="l00548"></a><span class="lineno"> 548</span>&#160;};</div>
<div class="line"><a name="l00549"></a><span class="lineno"> 549</span>&#160;</div>
<div class="line"><a name="l00550"></a><span class="lineno"> 550</span>&#160;<span class="comment">// Note Mutex selected based on DeviceMutex</span></div>
<div class="line"><a name="l00551"></a><span class="lineno"> 551</span>&#160;<span class="keyword">typedef</span> <a class="code" href="classihipDeviceCriticalBase__t.html">ihipDeviceCriticalBase_t&lt;DeviceMutex&gt;</a> <a class="code" href="classihipDeviceCriticalBase__t.html">ihipDeviceCritical_t</a>; </div>
<div class="line"><a name="l00537"></a><span class="lineno"> 537</span>&#160; uint32_t peerCnt()<span class="keyword"> const </span>{ <span class="keywordflow">return</span> _peerCnt; };</div>
<div class="line"><a name="l00538"></a><span class="lineno"> 538</span>&#160; hsa_agent_t *peerAgents()<span class="keyword"> const </span>{ <span class="keywordflow">return</span> _peerAgents; };</div>
<div class="line"><a name="l00539"></a><span class="lineno"> 539</span>&#160;</div>
<div class="line"><a name="l00540"></a><span class="lineno"> 540</span>&#160;</div>
<div class="line"><a name="l00541"></a><span class="lineno"> 541</span>&#160;<span class="keyword">private</span>:</div>
<div class="line"><a name="l00542"></a><span class="lineno"> 542</span>&#160; std::list&lt;ihipStream_t*&gt; _streams; <span class="comment">// streams associated with this device.</span></div>
<div class="line"><a name="l00543"></a><span class="lineno"> 543</span>&#160; ihipStream_t::SeqNum_t _stream_id;</div>
<div class="line"><a name="l00544"></a><span class="lineno"> 544</span>&#160;</div>
<div class="line"><a name="l00545"></a><span class="lineno"> 545</span>&#160; <span class="comment">// These reflect the currently Enabled set of peers for this GPU:</span></div>
<div class="line"><a name="l00546"></a><span class="lineno"> 546</span>&#160; std::list&lt;ihipDevice_t*&gt; _peers; <span class="comment">// list of enabled peer devices.</span></div>
<div class="line"><a name="l00547"></a><span class="lineno"> 547</span>&#160; uint32_t _peerCnt; <span class="comment">// number of enabled peers</span></div>
<div class="line"><a name="l00548"></a><span class="lineno"> 548</span>&#160; hsa_agent_t *_peerAgents; <span class="comment">// efficient packed array of enabled agents (to use for allocations.)</span></div>
<div class="line"><a name="l00549"></a><span class="lineno"> 549</span>&#160;<span class="keyword">private</span>:</div>
<div class="line"><a name="l00550"></a><span class="lineno"> 550</span>&#160; <span class="keywordtype">void</span> recomputePeerAgents();</div>
<div class="line"><a name="l00551"></a><span class="lineno"> 551</span>&#160;};</div>
<div class="line"><a name="l00552"></a><span class="lineno"> 552</span>&#160;</div>
<div class="line"><a name="l00553"></a><span class="lineno"> 553</span>&#160;<span class="comment">// This type is used by functions that need access to the critical device structures.</span></div>
<div class="line"><a name="l00554"></a><span class="lineno"> 554</span>&#160;<span class="keyword">typedef</span> <a class="code" href="classLockedAccessor.html">LockedAccessor&lt;ihipDeviceCritical_t&gt;</a> <a class="code" href="classLockedAccessor.html">LockedAccessor_DeviceCrit_t</a>;</div>
<div class="line"><a name="l00553"></a><span class="lineno"> 553</span>&#160;<span class="comment">// Note Mutex selected based on DeviceMutex</span></div>
<div class="line"><a name="l00554"></a><span class="lineno"> 554</span>&#160;<span class="keyword">typedef</span> <a class="code" href="classihipDeviceCriticalBase__t.html">ihipDeviceCriticalBase_t&lt;DeviceMutex&gt;</a> <a class="code" href="classihipDeviceCriticalBase__t.html">ihipDeviceCritical_t</a>; </div>
<div class="line"><a name="l00555"></a><span class="lineno"> 555</span>&#160;</div>
<div class="line"><a name="l00556"></a><span class="lineno"> 556</span>&#160;</div>
<div class="line"><a name="l00557"></a><span class="lineno"> 557</span>&#160;</div>
<div class="line"><a name="l00558"></a><span class="lineno"> 558</span>&#160;<span class="comment">//-------------------------------------------------------------------------------------------------</span></div>
<div class="line"><a name="l00559"></a><span class="lineno"> 559</span>&#160;<span class="comment">// Functions which read or write the critical data are named locked_.</span></div>
<div class="line"><a name="l00560"></a><span class="lineno"> 560</span>&#160;<span class="comment">// ihipDevice_t does not use recursive locks so the ihip implementation must avoid calling a locked_ function from within a locked_ function.</span></div>
<div class="line"><a name="l00561"></a><span class="lineno"> 561</span>&#160;<span class="comment">// External functions which call several locked_ functions will acquire and release the lock for each function. if this occurs in </span></div>
<div class="line"><a name="l00562"></a><span class="lineno"> 562</span>&#160;<span class="comment">// performance-sensitive code we may want to refactor by adding non-locked functions and creating a new locked_ member function to call them all.</span></div>
<div class="line"><a name="l00563"></a><span class="lineno"><a class="line" href="classihipDevice__t.html"> 563</a></span>&#160;<span class="keyword">class </span><a class="code" href="classihipDevice__t.html">ihipDevice_t</a></div>
<div class="line"><a name="l00564"></a><span class="lineno"> 564</span>&#160;{</div>
<div class="line"><a name="l00565"></a><span class="lineno"> 565</span>&#160;<span class="keyword">public</span>: <span class="comment">// Functions:</span></div>
<div class="line"><a name="l00566"></a><span class="lineno"> 566</span>&#160; <a class="code" href="classihipDevice__t.html">ihipDevice_t</a>() {}; <span class="comment">// note: calls constructor for _criticalData </span></div>
<div class="line"><a name="l00567"></a><span class="lineno"> 567</span>&#160; <span class="keywordtype">void</span> init(<span class="keywordtype">unsigned</span> device_index, <span class="keywordtype">unsigned</span> deviceCnt, hc::accelerator &amp;acc, <span class="keywordtype">unsigned</span> flags);</div>
<div class="line"><a name="l00568"></a><span class="lineno"> 568</span>&#160; ~<a class="code" href="classihipDevice__t.html">ihipDevice_t</a>();</div>
<div class="line"><a name="l00569"></a><span class="lineno"> 569</span>&#160;</div>
<div class="line"><a name="l00570"></a><span class="lineno"> 570</span>&#160; <span class="keywordtype">void</span> locked_addStream(<a class="code" href="classihipStream__t.html">ihipStream_t</a> *s);</div>
<div class="line"><a name="l00571"></a><span class="lineno"> 571</span>&#160; <span class="keywordtype">void</span> locked_removeStream(<a class="code" href="classihipStream__t.html">ihipStream_t</a> *s);</div>
<div class="line"><a name="l00572"></a><span class="lineno"> 572</span>&#160; <span class="keywordtype">void</span> locked_reset();</div>
<div class="line"><a name="l00573"></a><span class="lineno"> 573</span>&#160; <span class="keywordtype">void</span> locked_waitAllStreams();</div>
<div class="line"><a name="l00574"></a><span class="lineno"> 574</span>&#160; <span class="keywordtype">void</span> locked_syncDefaultStream(<span class="keywordtype">bool</span> waitOnSelf);</div>
<div class="line"><a name="l00575"></a><span class="lineno"> 575</span>&#160;</div>
<div class="line"><a name="l00576"></a><span class="lineno"> 576</span>&#160; <a class="code" href="classihipDeviceCriticalBase__t.html">ihipDeviceCritical_t</a> &amp;criticalData() { <span class="keywordflow">return</span> _criticalData; }; <span class="comment">// TODO, move private. Fix P2P.</span></div>
<div class="line"><a name="l00577"></a><span class="lineno"> 577</span>&#160;</div>
<div class="line"><a name="l00578"></a><span class="lineno"> 578</span>&#160;<span class="keyword">public</span>: <span class="comment">// Data, set at initialization:</span></div>
<div class="line"><a name="l00579"></a><span class="lineno"> 579</span>&#160; <span class="keywordtype">unsigned</span> _device_index; <span class="comment">// index into g_devices.</span></div>
<div class="line"><a name="l00556"></a><span class="lineno"> 556</span>&#160;<span class="comment">// This type is used by functions that need access to the critical device structures.</span></div>
<div class="line"><a name="l00557"></a><span class="lineno"> 557</span>&#160;<span class="keyword">typedef</span> <a class="code" href="classLockedAccessor.html">LockedAccessor&lt;ihipDeviceCritical_t&gt;</a> <a class="code" href="classLockedAccessor.html">LockedAccessor_DeviceCrit_t</a>;</div>
<div class="line"><a name="l00558"></a><span class="lineno"> 558</span>&#160;</div>
<div class="line"><a name="l00559"></a><span class="lineno"> 559</span>&#160;</div>
<div class="line"><a name="l00560"></a><span class="lineno"> 560</span>&#160;</div>
<div class="line"><a name="l00561"></a><span class="lineno"> 561</span>&#160;<span class="comment">//-------------------------------------------------------------------------------------------------</span></div>
<div class="line"><a name="l00562"></a><span class="lineno"> 562</span>&#160;<span class="comment">// Functions which read or write the critical data are named locked_.</span></div>
<div class="line"><a name="l00563"></a><span class="lineno"> 563</span>&#160;<span class="comment">// ihipDevice_t does not use recursive locks so the ihip implementation must avoid calling a locked_ function from within a locked_ function.</span></div>
<div class="line"><a name="l00564"></a><span class="lineno"> 564</span>&#160;<span class="comment">// External functions which call several locked_ functions will acquire and release the lock for each function. if this occurs in </span></div>
<div class="line"><a name="l00565"></a><span class="lineno"> 565</span>&#160;<span class="comment">// performance-sensitive code we may want to refactor by adding non-locked functions and creating a new locked_ member function to call them all.</span></div>
<div class="line"><a name="l00566"></a><span class="lineno"><a class="line" href="classihipDevice__t.html"> 566</a></span>&#160;<span class="keyword">class </span><a class="code" href="classihipDevice__t.html">ihipDevice_t</a></div>
<div class="line"><a name="l00567"></a><span class="lineno"> 567</span>&#160;{</div>
<div class="line"><a name="l00568"></a><span class="lineno"> 568</span>&#160;<span class="keyword">public</span>: <span class="comment">// Functions:</span></div>
<div class="line"><a name="l00569"></a><span class="lineno"> 569</span>&#160; <a class="code" href="classihipDevice__t.html">ihipDevice_t</a>() {}; <span class="comment">// note: calls constructor for _criticalData </span></div>
<div class="line"><a name="l00570"></a><span class="lineno"> 570</span>&#160; <span class="keywordtype">void</span> init(<span class="keywordtype">unsigned</span> device_index, <span class="keywordtype">unsigned</span> deviceCnt, hc::accelerator &amp;acc, <span class="keywordtype">unsigned</span> flags);</div>
<div class="line"><a name="l00571"></a><span class="lineno"> 571</span>&#160; ~<a class="code" href="classihipDevice__t.html">ihipDevice_t</a>();</div>
<div class="line"><a name="l00572"></a><span class="lineno"> 572</span>&#160;</div>
<div class="line"><a name="l00573"></a><span class="lineno"> 573</span>&#160; <span class="keywordtype">void</span> locked_addStream(<a class="code" href="classihipStream__t.html">ihipStream_t</a> *s);</div>
<div class="line"><a name="l00574"></a><span class="lineno"> 574</span>&#160; <span class="keywordtype">void</span> locked_removeStream(<a class="code" href="classihipStream__t.html">ihipStream_t</a> *s);</div>
<div class="line"><a name="l00575"></a><span class="lineno"> 575</span>&#160; <span class="keywordtype">void</span> locked_reset();</div>
<div class="line"><a name="l00576"></a><span class="lineno"> 576</span>&#160; <span class="keywordtype">void</span> locked_waitAllStreams();</div>
<div class="line"><a name="l00577"></a><span class="lineno"> 577</span>&#160; <span class="keywordtype">void</span> locked_syncDefaultStream(<span class="keywordtype">bool</span> waitOnSelf);</div>
<div class="line"><a name="l00578"></a><span class="lineno"> 578</span>&#160;</div>
<div class="line"><a name="l00579"></a><span class="lineno"> 579</span>&#160; <a class="code" href="classihipDeviceCriticalBase__t.html">ihipDeviceCritical_t</a> &amp;criticalData() { <span class="keywordflow">return</span> _criticalData; }; <span class="comment">// TODO, move private. Fix P2P.</span></div>
<div class="line"><a name="l00580"></a><span class="lineno"> 580</span>&#160;</div>
<div class="line"><a name="l00581"></a><span class="lineno"> 581</span>&#160; <a class="code" href="structhipDeviceProp__t.html">hipDeviceProp_t</a> _props; <span class="comment">// saved device properties.</span></div>
<div class="line"><a name="l00582"></a><span class="lineno"> 582</span>&#160; hc::accelerator _acc;</div>
<div class="line"><a name="l00583"></a><span class="lineno"> 583</span>&#160; hsa_agent_t _hsa_agent; <span class="comment">// hsa agent handle</span></div>
<div class="line"><a name="l00584"></a><span class="lineno"> 584</span>&#160;</div>
<div class="line"><a name="l00585"></a><span class="lineno"> 585</span>&#160; <span class="comment">// The NULL stream is used if no other stream is specified.</span></div>
<div class="line"><a name="l00586"></a><span class="lineno"> 586</span>&#160; <span class="comment">// NULL has special synchronization properties with other streams.</span></div>
<div class="line"><a name="l00587"></a><span class="lineno"> 587</span>&#160; <a class="code" href="classihipStream__t.html">ihipStream_t</a> *_default_stream;</div>
<div class="line"><a name="l00588"></a><span class="lineno"> 588</span>&#160;</div>
<div class="line"><a name="l00589"></a><span class="lineno"> 589</span>&#160;</div>
<div class="line"><a name="l00590"></a><span class="lineno"> 590</span>&#160; <span class="keywordtype">unsigned</span> _compute_units;</div>
<div class="line"><a name="l00581"></a><span class="lineno"> 581</span>&#160;<span class="keyword">public</span>: <span class="comment">// Data, set at initialization:</span></div>
<div class="line"><a name="l00582"></a><span class="lineno"> 582</span>&#160; <span class="keywordtype">unsigned</span> _device_index; <span class="comment">// index into g_devices.</span></div>
<div class="line"><a name="l00583"></a><span class="lineno"> 583</span>&#160;</div>
<div class="line"><a name="l00584"></a><span class="lineno"> 584</span>&#160; <a class="code" href="structhipDeviceProp__t.html">hipDeviceProp_t</a> _props; <span class="comment">// saved device properties.</span></div>
<div class="line"><a name="l00585"></a><span class="lineno"> 585</span>&#160; hc::accelerator _acc;</div>
<div class="line"><a name="l00586"></a><span class="lineno"> 586</span>&#160; hsa_agent_t _hsa_agent; <span class="comment">// hsa agent handle</span></div>
<div class="line"><a name="l00587"></a><span class="lineno"> 587</span>&#160;</div>
<div class="line"><a name="l00588"></a><span class="lineno"> 588</span>&#160; <span class="comment">// The NULL stream is used if no other stream is specified.</span></div>
<div class="line"><a name="l00589"></a><span class="lineno"> 589</span>&#160; <span class="comment">// NULL has special synchronization properties with other streams.</span></div>
<div class="line"><a name="l00590"></a><span class="lineno"> 590</span>&#160; <a class="code" href="classihipStream__t.html">ihipStream_t</a> *_default_stream;</div>
<div class="line"><a name="l00591"></a><span class="lineno"> 591</span>&#160;</div>
<div class="line"><a name="l00592"></a><span class="lineno"> 592</span>&#160; <a class="code" href="structStagingBuffer.html">StagingBuffer</a> *_staging_buffer[2]; <span class="comment">// one buffer for each direction.</span></div>
<div class="line"><a name="l00593"></a><span class="lineno"> 593</span>&#160;</div>
<div class="line"><a name="l00592"></a><span class="lineno"> 592</span>&#160;</div>
<div class="line"><a name="l00593"></a><span class="lineno"> 593</span>&#160; <span class="keywordtype">unsigned</span> _compute_units;</div>
<div class="line"><a name="l00594"></a><span class="lineno"> 594</span>&#160;</div>
<div class="line"><a name="l00595"></a><span class="lineno"> 595</span>&#160; <span class="keywordtype">unsigned</span> _device_flags;</div>
<div class="line"><a name="l00595"></a><span class="lineno"> 595</span>&#160; <a class="code" href="structStagingBuffer.html">StagingBuffer</a> *_staging_buffer[2]; <span class="comment">// one buffer for each direction.</span></div>
<div class="line"><a name="l00596"></a><span class="lineno"> 596</span>&#160;</div>
<div class="line"><a name="l00597"></a><span class="lineno"> 597</span>&#160;<span class="keyword">private</span>:</div>
<div class="line"><a name="l00598"></a><span class="lineno"> 598</span>&#160; <a class="code" href="group__GlobalDefs.html#gadf5010f6e140a53ecbdf949e73e87594">hipError_t</a> getProperties(<a class="code" href="structhipDeviceProp__t.html">hipDeviceProp_t</a>* prop);</div>
<div class="line"><a name="l00597"></a><span class="lineno"> 597</span>&#160;</div>
<div class="line"><a name="l00598"></a><span class="lineno"> 598</span>&#160; <span class="keywordtype">unsigned</span> _device_flags;</div>
<div class="line"><a name="l00599"></a><span class="lineno"> 599</span>&#160;</div>
<div class="line"><a name="l00600"></a><span class="lineno"> 600</span>&#160;<span class="keyword">private</span>: <span class="comment">// Critical data, protected with locked access:</span></div>
<div class="line"><a name="l00601"></a><span class="lineno"> 601</span>&#160; <span class="comment">// Members of _protected data MUST be accessed through the LockedAccessor.</span></div>
<div class="line"><a name="l00602"></a><span class="lineno"> 602</span>&#160; <span class="comment">// Search for LockedAccessor&lt;ihipDeviceCritical_t&gt; for examples; do not access _criticalData directly.</span></div>
<div class="line"><a name="l00603"></a><span class="lineno"> 603</span>&#160; <a class="code" href="classihipDeviceCriticalBase__t.html">ihipDeviceCritical_t</a> _criticalData;</div>
<div class="line"><a name="l00604"></a><span class="lineno"> 604</span>&#160;</div>
<div class="line"><a name="l00605"></a><span class="lineno"> 605</span>&#160;};</div>
<div class="line"><a name="l00606"></a><span class="lineno"> 606</span>&#160;</div>
<div class="line"><a name="l00600"></a><span class="lineno"> 600</span>&#160;<span class="keyword">private</span>:</div>
<div class="line"><a name="l00601"></a><span class="lineno"> 601</span>&#160; <a class="code" href="group__GlobalDefs.html#gadf5010f6e140a53ecbdf949e73e87594">hipError_t</a> getProperties(<a class="code" href="structhipDeviceProp__t.html">hipDeviceProp_t</a>* prop);</div>
<div class="line"><a name="l00602"></a><span class="lineno"> 602</span>&#160;</div>
<div class="line"><a name="l00603"></a><span class="lineno"> 603</span>&#160;<span class="keyword">private</span>: <span class="comment">// Critical data, protected with locked access:</span></div>
<div class="line"><a name="l00604"></a><span class="lineno"> 604</span>&#160; <span class="comment">// Members of _protected data MUST be accessed through the LockedAccessor.</span></div>
<div class="line"><a name="l00605"></a><span class="lineno"> 605</span>&#160; <span class="comment">// Search for LockedAccessor&lt;ihipDeviceCritical_t&gt; for examples; do not access _criticalData directly.</span></div>
<div class="line"><a name="l00606"></a><span class="lineno"> 606</span>&#160; <a class="code" href="classihipDeviceCriticalBase__t.html">ihipDeviceCritical_t</a> _criticalData;</div>
<div class="line"><a name="l00607"></a><span class="lineno"> 607</span>&#160;</div>
<div class="line"><a name="l00608"></a><span class="lineno"> 608</span>&#160;</div>
<div class="line"><a name="l00609"></a><span class="lineno"> 609</span>&#160;<span class="comment">// Global variable definition:</span></div>
<div class="line"><a name="l00610"></a><span class="lineno"> 610</span>&#160;<span class="keyword">extern</span> std::once_flag hip_initialized;</div>
<div class="line"><a name="l00611"></a><span class="lineno"> 611</span>&#160;<span class="keyword">extern</span> <a class="code" href="classihipDevice__t.html">ihipDevice_t</a> *g_devices; <span class="comment">// Array of all non-emulated (ie GPU) accelerators in the system.</span></div>
<div class="line"><a name="l00612"></a><span class="lineno"> 612</span>&#160;<span class="keyword">extern</span> <span class="keywordtype">bool</span> g_visible_device; <span class="comment">// Set the flag when HIP_VISIBLE_DEVICES is set</span></div>
<div class="line"><a name="l00613"></a><span class="lineno"> 613</span>&#160;<span class="keyword">extern</span> <span class="keywordtype">unsigned</span> g_deviceCnt;</div>
<div class="line"><a name="l00614"></a><span class="lineno"> 614</span>&#160;<span class="keyword">extern</span> std::vector&lt;int&gt; g_hip_visible_devices; <span class="comment">/* vector of integers that contains the visible device IDs */</span></div>
<div class="line"><a name="l00615"></a><span class="lineno"> 615</span>&#160;<span class="keyword">extern</span> hsa_agent_t g_cpu_agent ; <span class="comment">// the CPU agent.</span></div>
<div class="line"><a name="l00616"></a><span class="lineno"> 616</span>&#160;<span class="comment">//=================================================================================================</span></div>
<div class="line"><a name="l00617"></a><span class="lineno"> 617</span>&#160;<span class="keywordtype">void</span> ihipInit();</div>
<div class="line"><a name="l00618"></a><span class="lineno"> 618</span>&#160;<span class="keyword">const</span> <span class="keywordtype">char</span> *ihipErrorString(<a class="code" href="group__GlobalDefs.html#gadf5010f6e140a53ecbdf949e73e87594">hipError_t</a>);</div>
<div class="line"><a name="l00619"></a><span class="lineno"> 619</span>&#160;<a class="code" href="classihipDevice__t.html">ihipDevice_t</a> *ihipGetTlsDefaultDevice();</div>
<div class="line"><a name="l00620"></a><span class="lineno"> 620</span>&#160;<a class="code" href="classihipDevice__t.html">ihipDevice_t</a> *ihipGetDevice(<span class="keywordtype">int</span>);</div>
<div class="line"><a name="l00621"></a><span class="lineno"> 621</span>&#160;<span class="keywordtype">void</span> ihipSetTs(<a class="code" href="structhipEvent__t.html">hipEvent_t</a> e);</div>
<div class="line"><a name="l00622"></a><span class="lineno"> 622</span>&#160;</div>
<div class="line"><a name="l00623"></a><span class="lineno"> 623</span>&#160;<span class="keyword">template</span>&lt;<span class="keyword">typename</span> T&gt;</div>
<div class="line"><a name="l00624"></a><span class="lineno"> 624</span>&#160;hc::completion_future ihipMemcpyKernel(hipStream_t, T*, <span class="keyword">const</span> T*, <span class="keywordtype">size_t</span>);</div>
<div class="line"><a name="l00608"></a><span class="lineno"> 608</span>&#160;};</div>
<div class="line"><a name="l00609"></a><span class="lineno"> 609</span>&#160;</div>
<div class="line"><a name="l00610"></a><span class="lineno"> 610</span>&#160;</div>
<div class="line"><a name="l00611"></a><span class="lineno"> 611</span>&#160;</div>
<div class="line"><a name="l00612"></a><span class="lineno"> 612</span>&#160;<span class="comment">// Global variable definition:</span></div>
<div class="line"><a name="l00613"></a><span class="lineno"> 613</span>&#160;<span class="keyword">extern</span> std::once_flag hip_initialized;</div>
<div class="line"><a name="l00614"></a><span class="lineno"> 614</span>&#160;<span class="keyword">extern</span> <a class="code" href="classihipDevice__t.html">ihipDevice_t</a> *g_devices; <span class="comment">// Array of all non-emulated (ie GPU) accelerators in the system.</span></div>
<div class="line"><a name="l00615"></a><span class="lineno"> 615</span>&#160;<span class="keyword">extern</span> <span class="keywordtype">bool</span> g_visible_device; <span class="comment">// Set the flag when HIP_VISIBLE_DEVICES is set</span></div>
<div class="line"><a name="l00616"></a><span class="lineno"> 616</span>&#160;<span class="keyword">extern</span> <span class="keywordtype">unsigned</span> g_deviceCnt;</div>
<div class="line"><a name="l00617"></a><span class="lineno"> 617</span>&#160;<span class="keyword">extern</span> std::vector&lt;int&gt; g_hip_visible_devices; <span class="comment">/* vector of integers that contains the visible device IDs */</span></div>
<div class="line"><a name="l00618"></a><span class="lineno"> 618</span>&#160;<span class="keyword">extern</span> hsa_agent_t g_cpu_agent ; <span class="comment">// the CPU agent.</span></div>
<div class="line"><a name="l00619"></a><span class="lineno"> 619</span>&#160;<span class="comment">//=================================================================================================</span></div>
<div class="line"><a name="l00620"></a><span class="lineno"> 620</span>&#160;<span class="keywordtype">void</span> ihipInit();</div>
<div class="line"><a name="l00621"></a><span class="lineno"> 621</span>&#160;<span class="keyword">const</span> <span class="keywordtype">char</span> *ihipErrorString(<a class="code" href="group__GlobalDefs.html#gadf5010f6e140a53ecbdf949e73e87594">hipError_t</a>);</div>
<div class="line"><a name="l00622"></a><span class="lineno"> 622</span>&#160;<a class="code" href="classihipDevice__t.html">ihipDevice_t</a> *ihipGetTlsDefaultDevice();</div>
<div class="line"><a name="l00623"></a><span class="lineno"> 623</span>&#160;<a class="code" href="classihipDevice__t.html">ihipDevice_t</a> *ihipGetDevice(<span class="keywordtype">int</span>);</div>
<div class="line"><a name="l00624"></a><span class="lineno"> 624</span>&#160;<span class="keywordtype">void</span> ihipSetTs(<a class="code" href="structhipEvent__t.html">hipEvent_t</a> e);</div>
<div class="line"><a name="l00625"></a><span class="lineno"> 625</span>&#160;</div>
<div class="line"><a name="l00626"></a><span class="lineno"> 626</span>&#160;<span class="keyword">template</span>&lt;<span class="keyword">typename</span> T&gt;</div>
<div class="line"><a name="l00627"></a><span class="lineno"> 627</span>&#160;hc::completion_future ihipMemsetKernel(hipStream_t, T*, T, <span class="keywordtype">size_t</span>);</div>
<div class="line"><a name="l00627"></a><span class="lineno"> 627</span>&#160;hc::completion_future ihipMemcpyKernel(hipStream_t, T*, <span class="keyword">const</span> T*, <span class="keywordtype">size_t</span>);</div>
<div class="line"><a name="l00628"></a><span class="lineno"> 628</span>&#160;</div>
<div class="line"><a name="l00629"></a><span class="lineno"> 629</span>&#160;hipStream_t ihipSyncAndResolveStream(hipStream_t);</div>
<div class="line"><a name="l00630"></a><span class="lineno"> 630</span>&#160;<span class="keyword">template</span> &lt;<span class="keyword">typename</span> T&gt;</div>
<div class="line"><a name="l00629"></a><span class="lineno"> 629</span>&#160;<span class="keyword">template</span>&lt;<span class="keyword">typename</span> T&gt;</div>
<div class="line"><a name="l00630"></a><span class="lineno"> 630</span>&#160;hc::completion_future ihipMemsetKernel(hipStream_t, T*, T, <span class="keywordtype">size_t</span>);</div>
<div class="line"><a name="l00631"></a><span class="lineno"> 631</span>&#160;</div>
<div class="line"><a name="l00632"></a><span class="lineno"> 632</span>&#160;hc::completion_future</div>
<div class="line"><a name="l00633"></a><span class="lineno"> 633</span>&#160;ihipMemsetKernel(hipStream_t stream, T * ptr, T val, <span class="keywordtype">size_t</span> sizeBytes)</div>
<div class="line"><a name="l00634"></a><span class="lineno"> 634</span>&#160;{</div>
<div class="line"><a name="l00635"></a><span class="lineno"> 635</span>&#160; <span class="keywordtype">int</span> wg = std::min((<span class="keywordtype">unsigned</span>)8, stream-&gt;getDevice()-&gt;_compute_units);</div>
<div class="line"><a name="l00636"></a><span class="lineno"> 636</span>&#160; <span class="keyword">const</span> <span class="keywordtype">int</span> threads_per_wg = 256;</div>
<div class="line"><a name="l00637"></a><span class="lineno"> 637</span>&#160;</div>
<div class="line"><a name="l00638"></a><span class="lineno"> 638</span>&#160; <span class="keywordtype">int</span> threads = wg * threads_per_wg;</div>
<div class="line"><a name="l00639"></a><span class="lineno"> 639</span>&#160; <span class="keywordflow">if</span> (threads &gt; sizeBytes) {</div>
<div class="line"><a name="l00640"></a><span class="lineno"> 640</span>&#160; threads = ((sizeBytes + threads_per_wg - 1) / threads_per_wg) * threads_per_wg;</div>
<div class="line"><a name="l00641"></a><span class="lineno"> 641</span>&#160; }</div>
<div class="line"><a name="l00642"></a><span class="lineno"> 642</span>&#160;</div>
<div class="line"><a name="l00643"></a><span class="lineno"> 643</span>&#160;</div>
<div class="line"><a name="l00644"></a><span class="lineno"> 644</span>&#160; hc::extent&lt;1&gt; ext(threads);</div>
<div class="line"><a name="l00645"></a><span class="lineno"> 645</span>&#160; <span class="keyword">auto</span> ext_tile = ext.tile(threads_per_wg);</div>
<div class="line"><a name="l00632"></a><span class="lineno"> 632</span>&#160;hipStream_t ihipSyncAndResolveStream(hipStream_t);</div>
<div class="line"><a name="l00633"></a><span class="lineno"> 633</span>&#160;<span class="keyword">template</span> &lt;<span class="keyword">typename</span> T&gt;</div>
<div class="line"><a name="l00634"></a><span class="lineno"> 634</span>&#160;</div>
<div class="line"><a name="l00635"></a><span class="lineno"> 635</span>&#160;hc::completion_future</div>
<div class="line"><a name="l00636"></a><span class="lineno"> 636</span>&#160;ihipMemsetKernel(hipStream_t stream, T * ptr, T val, <span class="keywordtype">size_t</span> sizeBytes)</div>
<div class="line"><a name="l00637"></a><span class="lineno"> 637</span>&#160;{</div>
<div class="line"><a name="l00638"></a><span class="lineno"> 638</span>&#160; <span class="keywordtype">int</span> wg = std::min((<span class="keywordtype">unsigned</span>)8, stream-&gt;getDevice()-&gt;_compute_units);</div>
<div class="line"><a name="l00639"></a><span class="lineno"> 639</span>&#160; <span class="keyword">const</span> <span class="keywordtype">int</span> threads_per_wg = 256;</div>
<div class="line"><a name="l00640"></a><span class="lineno"> 640</span>&#160;</div>
<div class="line"><a name="l00641"></a><span class="lineno"> 641</span>&#160; <span class="keywordtype">int</span> threads = wg * threads_per_wg;</div>
<div class="line"><a name="l00642"></a><span class="lineno"> 642</span>&#160; <span class="keywordflow">if</span> (threads &gt; sizeBytes) {</div>
<div class="line"><a name="l00643"></a><span class="lineno"> 643</span>&#160; threads = ((sizeBytes + threads_per_wg - 1) / threads_per_wg) * threads_per_wg;</div>
<div class="line"><a name="l00644"></a><span class="lineno"> 644</span>&#160; }</div>
<div class="line"><a name="l00645"></a><span class="lineno"> 645</span>&#160;</div>
<div class="line"><a name="l00646"></a><span class="lineno"> 646</span>&#160;</div>
<div class="line"><a name="l00647"></a><span class="lineno"> 647</span>&#160; hc::completion_future cf =</div>
<div class="line"><a name="l00648"></a><span class="lineno"> 648</span>&#160; hc::parallel_for_each(</div>
<div class="line"><a name="l00649"></a><span class="lineno"> 649</span>&#160; stream-&gt;_av,</div>
<div class="line"><a name="l00650"></a><span class="lineno"> 650</span>&#160; ext_tile,</div>
<div class="line"><a name="l00651"></a><span class="lineno"> 651</span>&#160; [=] (hc::tiled_index&lt;1&gt; idx)</div>
<div class="line"><a name="l00652"></a><span class="lineno"> 652</span>&#160; __attribute__((hc))</div>
<div class="line"><a name="l00653"></a><span class="lineno"> 653</span>&#160; {</div>
<div class="line"><a name="l00654"></a><span class="lineno"> 654</span>&#160; <span class="keywordtype">int</span> offset = amp_get_global_id(0);</div>
<div class="line"><a name="l00655"></a><span class="lineno"> 655</span>&#160; <span class="comment">// TODO-HCC - change to hc_get_local_size()</span></div>
<div class="line"><a name="l00656"></a><span class="lineno"> 656</span>&#160; <span class="keywordtype">int</span> stride = amp_get_local_size(0) * hc_get_num_groups(0) ;</div>
<div class="line"><a name="l00657"></a><span class="lineno"> 657</span>&#160;</div>
<div class="line"><a name="l00658"></a><span class="lineno"> 658</span>&#160; <span class="keywordflow">for</span> (<span class="keywordtype">int</span> i=offset; i&lt;sizeBytes; i+=stride) {</div>
<div class="line"><a name="l00659"></a><span class="lineno"> 659</span>&#160; ptr[i] = val;</div>
<div class="line"><a name="l00660"></a><span class="lineno"> 660</span>&#160; }</div>
<div class="line"><a name="l00661"></a><span class="lineno"> 661</span>&#160; });</div>
<div class="line"><a name="l00662"></a><span class="lineno"> 662</span>&#160;</div>
<div class="line"><a name="l00663"></a><span class="lineno"> 663</span>&#160; <span class="keywordflow">return</span> cf;</div>
<div class="line"><a name="l00664"></a><span class="lineno"> 664</span>&#160;}</div>
<div class="line"><a name="l00647"></a><span class="lineno"> 647</span>&#160; hc::extent&lt;1&gt; ext(threads);</div>
<div class="line"><a name="l00648"></a><span class="lineno"> 648</span>&#160; <span class="keyword">auto</span> ext_tile = ext.tile(threads_per_wg);</div>
<div class="line"><a name="l00649"></a><span class="lineno"> 649</span>&#160;</div>
<div class="line"><a name="l00650"></a><span class="lineno"> 650</span>&#160; hc::completion_future cf =</div>
<div class="line"><a name="l00651"></a><span class="lineno"> 651</span>&#160; hc::parallel_for_each(</div>
<div class="line"><a name="l00652"></a><span class="lineno"> 652</span>&#160; stream-&gt;_av,</div>
<div class="line"><a name="l00653"></a><span class="lineno"> 653</span>&#160; ext_tile,</div>
<div class="line"><a name="l00654"></a><span class="lineno"> 654</span>&#160; [=] (hc::tiled_index&lt;1&gt; idx)</div>
<div class="line"><a name="l00655"></a><span class="lineno"> 655</span>&#160; __attribute__((hc))</div>
<div class="line"><a name="l00656"></a><span class="lineno"> 656</span>&#160; {</div>
<div class="line"><a name="l00657"></a><span class="lineno"> 657</span>&#160; <span class="keywordtype">int</span> offset = amp_get_global_id(0);</div>
<div class="line"><a name="l00658"></a><span class="lineno"> 658</span>&#160; <span class="comment">// TODO-HCC - change to hc_get_local_size()</span></div>
<div class="line"><a name="l00659"></a><span class="lineno"> 659</span>&#160; <span class="keywordtype">int</span> stride = amp_get_local_size(0) * hc_get_num_groups(0) ;</div>
<div class="line"><a name="l00660"></a><span class="lineno"> 660</span>&#160;</div>
<div class="line"><a name="l00661"></a><span class="lineno"> 661</span>&#160; <span class="keywordflow">for</span> (<span class="keywordtype">int</span> i=offset; i&lt;sizeBytes; i+=stride) {</div>
<div class="line"><a name="l00662"></a><span class="lineno"> 662</span>&#160; ptr[i] = val;</div>
<div class="line"><a name="l00663"></a><span class="lineno"> 663</span>&#160; }</div>
<div class="line"><a name="l00664"></a><span class="lineno"> 664</span>&#160; });</div>
<div class="line"><a name="l00665"></a><span class="lineno"> 665</span>&#160;</div>
<div class="line"><a name="l00666"></a><span class="lineno"> 666</span>&#160;<span class="keyword">template</span> &lt;<span class="keyword">typename</span> T&gt;</div>
<div class="line"><a name="l00667"></a><span class="lineno"> 667</span>&#160;hc::completion_future</div>
<div class="line"><a name="l00668"></a><span class="lineno"> 668</span>&#160;ihipMemcpyKernel(hipStream_t stream, T * c, <span class="keyword">const</span> T * a, <span class="keywordtype">size_t</span> sizeBytes)</div>
<div class="line"><a name="l00669"></a><span class="lineno"> 669</span>&#160;{</div>
<div class="line"><a name="l00670"></a><span class="lineno"> 670</span>&#160; <span class="keywordtype">int</span> wg = std::min((<span class="keywordtype">unsigned</span>)8, stream-&gt;getDevice()-&gt;_compute_units);</div>
<div class="line"><a name="l00671"></a><span class="lineno"> 671</span>&#160; <span class="keyword">const</span> <span class="keywordtype">int</span> threads_per_wg = 256;</div>
<div class="line"><a name="l00672"></a><span class="lineno"> 672</span>&#160;</div>
<div class="line"><a name="l00673"></a><span class="lineno"> 673</span>&#160; <span class="keywordtype">int</span> threads = wg * threads_per_wg;</div>
<div class="line"><a name="l00674"></a><span class="lineno"> 674</span>&#160; <span class="keywordflow">if</span> (threads &gt; sizeBytes) {</div>
<div class="line"><a name="l00675"></a><span class="lineno"> 675</span>&#160; threads = ((sizeBytes + threads_per_wg - 1) / threads_per_wg) * threads_per_wg;</div>
<div class="line"><a name="l00676"></a><span class="lineno"> 676</span>&#160; }</div>
<div class="line"><a name="l00677"></a><span class="lineno"> 677</span>&#160;</div>
<div class="line"><a name="l00678"></a><span class="lineno"> 678</span>&#160;</div>
<div class="line"><a name="l00679"></a><span class="lineno"> 679</span>&#160; hc::extent&lt;1&gt; ext(threads);</div>
<div class="line"><a name="l00680"></a><span class="lineno"> 680</span>&#160; <span class="keyword">auto</span> ext_tile = ext.tile(threads_per_wg);</div>
<div class="line"><a name="l00666"></a><span class="lineno"> 666</span>&#160; <span class="keywordflow">return</span> cf;</div>
<div class="line"><a name="l00667"></a><span class="lineno"> 667</span>&#160;}</div>
<div class="line"><a name="l00668"></a><span class="lineno"> 668</span>&#160;</div>
<div class="line"><a name="l00669"></a><span class="lineno"> 669</span>&#160;<span class="keyword">template</span> &lt;<span class="keyword">typename</span> T&gt;</div>
<div class="line"><a name="l00670"></a><span class="lineno"> 670</span>&#160;hc::completion_future</div>
<div class="line"><a name="l00671"></a><span class="lineno"> 671</span>&#160;ihipMemcpyKernel(hipStream_t stream, T * c, <span class="keyword">const</span> T * a, <span class="keywordtype">size_t</span> sizeBytes)</div>
<div class="line"><a name="l00672"></a><span class="lineno"> 672</span>&#160;{</div>
<div class="line"><a name="l00673"></a><span class="lineno"> 673</span>&#160; <span class="keywordtype">int</span> wg = std::min((<span class="keywordtype">unsigned</span>)8, stream-&gt;getDevice()-&gt;_compute_units);</div>
<div class="line"><a name="l00674"></a><span class="lineno"> 674</span>&#160; <span class="keyword">const</span> <span class="keywordtype">int</span> threads_per_wg = 256;</div>
<div class="line"><a name="l00675"></a><span class="lineno"> 675</span>&#160;</div>
<div class="line"><a name="l00676"></a><span class="lineno"> 676</span>&#160; <span class="keywordtype">int</span> threads = wg * threads_per_wg;</div>
<div class="line"><a name="l00677"></a><span class="lineno"> 677</span>&#160; <span class="keywordflow">if</span> (threads &gt; sizeBytes) {</div>
<div class="line"><a name="l00678"></a><span class="lineno"> 678</span>&#160; threads = ((sizeBytes + threads_per_wg - 1) / threads_per_wg) * threads_per_wg;</div>
<div class="line"><a name="l00679"></a><span class="lineno"> 679</span>&#160; }</div>
<div class="line"><a name="l00680"></a><span class="lineno"> 680</span>&#160;</div>
<div class="line"><a name="l00681"></a><span class="lineno"> 681</span>&#160;</div>
<div class="line"><a name="l00682"></a><span class="lineno"> 682</span>&#160; hc::completion_future cf =</div>
<div class="line"><a name="l00683"></a><span class="lineno"> 683</span>&#160; hc::parallel_for_each(</div>
<div class="line"><a name="l00684"></a><span class="lineno"> 684</span>&#160; stream-&gt;_av,</div>
<div class="line"><a name="l00685"></a><span class="lineno"> 685</span>&#160; ext_tile,</div>
<div class="line"><a name="l00686"></a><span class="lineno"> 686</span>&#160; [=] (hc::tiled_index&lt;1&gt; idx)</div>
<div class="line"><a name="l00687"></a><span class="lineno"> 687</span>&#160; __attribute__((hc))</div>
<div class="line"><a name="l00688"></a><span class="lineno"> 688</span>&#160; {</div>
<div class="line"><a name="l00689"></a><span class="lineno"> 689</span>&#160; <span class="keywordtype">int</span> offset = amp_get_global_id(0);</div>
<div class="line"><a name="l00690"></a><span class="lineno"> 690</span>&#160; <span class="comment">// TODO-HCC - change to hc_get_local_size()</span></div>
<div class="line"><a name="l00691"></a><span class="lineno"> 691</span>&#160; <span class="keywordtype">int</span> stride = amp_get_local_size(0) * hc_get_num_groups(0) ;</div>
<div class="line"><a name="l00692"></a><span class="lineno"> 692</span>&#160;</div>
<div class="line"><a name="l00693"></a><span class="lineno"> 693</span>&#160; <span class="keywordflow">for</span> (<span class="keywordtype">int</span> i=offset; i&lt;sizeBytes; i+=stride) {</div>
<div class="line"><a name="l00694"></a><span class="lineno"> 694</span>&#160; c[i] = a[i];</div>
<div class="line"><a name="l00695"></a><span class="lineno"> 695</span>&#160; }</div>
<div class="line"><a name="l00696"></a><span class="lineno"> 696</span>&#160; });</div>
<div class="line"><a name="l00697"></a><span class="lineno"> 697</span>&#160;</div>
<div class="line"><a name="l00698"></a><span class="lineno"> 698</span>&#160; <span class="keywordflow">return</span> cf;</div>
<div class="line"><a name="l00699"></a><span class="lineno"> 699</span>&#160;}</div>
<div class="line"><a name="l00682"></a><span class="lineno"> 682</span>&#160; hc::extent&lt;1&gt; ext(threads);</div>
<div class="line"><a name="l00683"></a><span class="lineno"> 683</span>&#160; <span class="keyword">auto</span> ext_tile = ext.tile(threads_per_wg);</div>
<div class="line"><a name="l00684"></a><span class="lineno"> 684</span>&#160;</div>
<div class="line"><a name="l00685"></a><span class="lineno"> 685</span>&#160; hc::completion_future cf =</div>
<div class="line"><a name="l00686"></a><span class="lineno"> 686</span>&#160; hc::parallel_for_each(</div>
<div class="line"><a name="l00687"></a><span class="lineno"> 687</span>&#160; stream-&gt;_av,</div>
<div class="line"><a name="l00688"></a><span class="lineno"> 688</span>&#160; ext_tile,</div>
<div class="line"><a name="l00689"></a><span class="lineno"> 689</span>&#160; [=] (hc::tiled_index&lt;1&gt; idx)</div>
<div class="line"><a name="l00690"></a><span class="lineno"> 690</span>&#160; __attribute__((hc))</div>
<div class="line"><a name="l00691"></a><span class="lineno"> 691</span>&#160; {</div>
<div class="line"><a name="l00692"></a><span class="lineno"> 692</span>&#160; <span class="keywordtype">int</span> offset = amp_get_global_id(0);</div>
<div class="line"><a name="l00693"></a><span class="lineno"> 693</span>&#160; <span class="comment">// TODO-HCC - change to hc_get_local_size()</span></div>
<div class="line"><a name="l00694"></a><span class="lineno"> 694</span>&#160; <span class="keywordtype">int</span> stride = amp_get_local_size(0) * hc_get_num_groups(0) ;</div>
<div class="line"><a name="l00695"></a><span class="lineno"> 695</span>&#160;</div>
<div class="line"><a name="l00696"></a><span class="lineno"> 696</span>&#160; <span class="keywordflow">for</span> (<span class="keywordtype">int</span> i=offset; i&lt;sizeBytes; i+=stride) {</div>
<div class="line"><a name="l00697"></a><span class="lineno"> 697</span>&#160; c[i] = a[i];</div>
<div class="line"><a name="l00698"></a><span class="lineno"> 698</span>&#160; }</div>
<div class="line"><a name="l00699"></a><span class="lineno"> 699</span>&#160; });</div>
<div class="line"><a name="l00700"></a><span class="lineno"> 700</span>&#160;</div>
<div class="line"><a name="l00701"></a><span class="lineno"> 701</span>&#160;<span class="preprocessor">#endif</span></div>
<div class="ttc" id="classihipDevice__t_html"><div class="ttname"><a href="classihipDevice__t.html">ihipDevice_t</a></div><div class="ttdef"><b>Definition:</b> hip_hcc.h:563</div></div>
<div class="line"><a name="l00701"></a><span class="lineno"> 701</span>&#160; <span class="keywordflow">return</span> cf;</div>
<div class="line"><a name="l00702"></a><span class="lineno"> 702</span>&#160;}</div>
<div class="line"><a name="l00703"></a><span class="lineno"> 703</span>&#160;</div>
<div class="line"><a name="l00704"></a><span class="lineno"> 704</span>&#160;<span class="preprocessor">#endif</span></div>
<div class="ttc" id="classihipDevice__t_html"><div class="ttname"><a href="classihipDevice__t.html">ihipDevice_t</a></div><div class="ttdef"><b>Definition:</b> hip_hcc.h:566</div></div>
<div class="ttc" id="structLockedBase_html"><div class="ttname"><a href="structLockedBase.html">LockedBase</a></div><div class="ttdef"><b>Definition:</b> hip_hcc.h:340</div></div>
<div class="ttc" id="classFakeMutex_html"><div class="ttname"><a href="classFakeMutex.html">FakeMutex</a></div><div class="ttdef"><b>Definition:</b> hip_hcc.h:279</div></div>
<div class="ttc" id="group__GlobalDefs_html_gadf5010f6e140a53ecbdf949e73e87594"><div class="ttname"><a href="group__GlobalDefs.html#gadf5010f6e140a53ecbdf949e73e87594">hipError_t</a></div><div class="ttdeci">hipError_t</div><div class="ttdef"><b>Definition:</b> hip_runtime_api.h:142</div></div>
@@ -807,7 +810,7 @@ var searchBox = new SearchBox("searchBox", "search",false,'Search');
</div><!-- fragment --></div><!-- contents -->
<!-- start footer part -->
<hr class="footer"/><address class="footer"><small>
Generated on Sat Apr 16 2016 14:52:09 for HIP: Heterogenous-computing Interface for Portability by &#160;<a href="http://www.doxygen.org/index.html">
Generated on Tue Apr 19 2016 22:44:02 for HIP: Heterogenous-computing Interface for Portability by &#160;<a href="http://www.doxygen.org/index.html">
<img class="footer" src="doxygen.png" alt="doxygen"/>
</a> 1.8.6
</small></address>