2016-04-16 14:55:10 +05:30
<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Transitional//EN" "http://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.6" />
2016-04-19 22:44:58 +05:30
< title > HIP: Heterogenous-computing Interface for Portability: /home/mangupta/hip_git/release_0.84.00/include/hcc_detail/hip_hcc.h Source File</ title >
2016-04-16 14:55:10 +05:30
< 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/search.js" ></ script >
< script type = "text/javascript" >
$ ( document ). ready ( function () { searchBox . OnSelectItem ( 0 ); });
</ 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 style = "padding-left: 0.5em;" >
< div id = "projectname" > HIP: Heterogenous-computing Interface for Portability
</ div >
</ td >
</ tr >
</ tbody >
</ table >
</ div >
<!-- end header part -->
<!-- Generated by Doxygen 1.8.6 -->
< script type = "text/javascript" >
var searchBox = new SearchBox ( "searchBox" , "search" , false , 'Search' );
</ script >
< div id = "navrow1" class = "tabs" >
< ul class = "tablist" >
< li >< a href = "index.html" >< span > Main  Page</ span ></ a ></ li >
< li >< a href = "pages.html" >< span > Related  Pages</ span ></ a ></ li >
< li >< a href = "modules.html" >< span > Modules</ span ></ a ></ li >
< li >< a href = "annotated.html" >< span > Classes</ span ></ a ></ li >
< li class = "current" >< a href = "files.html" >< span > Files</ span ></ a ></ li >
< li >
< div id = "MSearchBox" class = "MSearchBoxInactive" >
< span class = "left" >
< img id = "MSearchSelect" src = "search/mag_sel.png"
onmouseover = "return searchBox.OnSearchSelectShow()"
onmouseout = "return searchBox.OnSearchSelectHide()"
alt = "" />
< input type = "text" id = "MSearchField" value = "Search" accesskey = "S"
onfocus = "searchBox.OnSearchFieldFocus(true)"
onblur = "searchBox.OnSearchFieldFocus(false)"
onkeyup = "searchBox.OnSearchFieldChange(event)" />
</ span >< span class = "right" >
< a id = "MSearchClose" href = "javascript:searchBox.CloseResultsWindow()" >< img id = "MSearchCloseImg" border = "0" src = "search/close.png" alt = "" /></ a >
</ span >
</ div >
</ li >
</ ul >
</ div >
< div id = "navrow2" class = "tabs2" >
< ul class = "tablist" >
< li >< a href = "files.html" >< span > File  List</ span ></ a ></ li >
< li >< a href = "globals.html" >< span > File  Members</ span ></ a ></ li >
</ ul >
</ div >
<!-- window showing the filter options -->
< div id = "MSearchSelectWindow"
onmouseover = "return searchBox.OnSearchSelectShow()"
onmouseout = "return searchBox.OnSearchSelectHide()"
onkeydown = "return searchBox.OnSearchSelectKey(event)" >
< a class = "SelectItem" href = "javascript:void(0)" onclick = "searchBox.OnSelectItem(0)" >< span class = "SelectionMark" >   </ span > All</ a >< a class = "SelectItem" href = "javascript:void(0)" onclick = "searchBox.OnSelectItem(1)" >< span class = "SelectionMark" >   </ span > Classes</ a >< a class = "SelectItem" href = "javascript:void(0)" onclick = "searchBox.OnSelectItem(2)" >< span class = "SelectionMark" >   </ span > Files</ a >< a class = "SelectItem" href = "javascript:void(0)" onclick = "searchBox.OnSelectItem(3)" >< span class = "SelectionMark" >   </ span > Functions</ a >< a class = "SelectItem" href = "javascript:void(0)" onclick = "searchBox.OnSelectItem(4)" >< span class = "SelectionMark" >   </ span > Variables</ a >< a class = "SelectItem" href = "javascript:void(0)" onclick = "searchBox.OnSelectItem(5)" >< span class = "SelectionMark" >   </ span > Typedefs</ a >< a class = "SelectItem" href = "javascript:void(0)" onclick = "searchBox.OnSelectItem(6)" >< span class = "SelectionMark" >   </ span > Enumerations</ a >< a class = "SelectItem" href = "javascript:void(0)" onclick = "searchBox.OnSelectItem(7)" >< span class = "SelectionMark" >   </ span > Enumerator</ a >< a class = "SelectItem" href = "javascript:void(0)" onclick = "searchBox.OnSelectItem(8)" >< span class = "SelectionMark" >   </ span > Macros</ a >< a class = "SelectItem" href = "javascript:void(0)" onclick = "searchBox.OnSelectItem(9)" >< span class = "SelectionMark" >   </ span > Groups</ a >< a class = "SelectItem" href = "javascript:void(0)" onclick = "searchBox.OnSelectItem(10)" >< span class = "SelectionMark" >   </ span > Pages</ a ></ 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_d44c64559bbebec7f509842c48db8b23.html" > include</ a ></ li >< li class = "navelem" >< a class = "el" href = "dir_6d8604cb65fa6b83549668eb0ce09cac.html" > hcc_detail</ a ></ li > </ ul >
</ div >
</ div > <!-- top -->
< div class = "header" >
< div class = "headertitle" >
< div class = "title" > hip_hcc.h</ div > </ div >
</ div > <!--header-->
< div class = "contents" >
< 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" > Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.</ span ></ div >
< div class = "line" >< a name = "l00003" ></ a >< span class = "lineno" > 3</ span >   < span class = "comment" > Permission is hereby granted, free of charge, to any person obtaining a copy</ span ></ div >
< div class = "line" >< a name = "l00004" ></ a >< span class = "lineno" > 4</ span >   < span class = "comment" > of this software and associated documentation files (the " Software" ), to deal</ span ></ div >
< div class = "line" >< a name = "l00005" ></ a >< span class = "lineno" > 5</ span >   < span class = "comment" > in the Software without restriction, including without limitation the rights</ span ></ div >
< div class = "line" >< a name = "l00006" ></ a >< span class = "lineno" > 6</ span >   < span class = "comment" > to use, copy, modify, merge, publish, distribute, sublicense, and/or sell</ span ></ div >
< div class = "line" >< a name = "l00007" ></ a >< span class = "lineno" > 7</ span >   < span class = "comment" > copies of the Software, and to permit persons to whom the Software is</ span ></ div >
< div class = "line" >< a name = "l00008" ></ a >< span class = "lineno" > 8</ span >   < span class = "comment" > furnished to do so, subject to the following conditions:</ span ></ div >
< div class = "line" >< a name = "l00009" ></ a >< span class = "lineno" > 9</ span >   < span class = "comment" > The above copyright notice and this permission notice shall be included in</ span ></ div >
< div class = "line" >< a name = "l00010" ></ a >< span class = "lineno" > 10</ span >   < span class = "comment" > all copies or substantial portions of the Software.</ span ></ div >
< div class = "line" >< a name = "l00011" ></ a >< span class = "lineno" > 11</ span >   < span class = "comment" > THE SOFTWARE IS PROVIDED " AS IS" , WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR</ span ></ div >
< div class = "line" >< a name = "l00012" ></ a >< span class = "lineno" > 12</ span >   < span class = "comment" > IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,</ span ></ div >
< div class = "line" >< a name = "l00013" ></ a >< span class = "lineno" > 13</ span >   < span class = "comment" > FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE</ span ></ div >
< div class = "line" >< a name = "l00014" ></ a >< span class = "lineno" > 14</ span >   < span class = "comment" > AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER</ span ></ div >
< div class = "line" >< a name = "l00015" ></ a >< span class = "lineno" > 15</ span >   < span class = "comment" > LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,</ span ></ div >
< div class = "line" >< a name = "l00016" ></ a >< span class = "lineno" > 16</ span >   < span class = "comment" > OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN</ span ></ div >
< div class = "line" >< a name = "l00017" ></ a >< span class = "lineno" > 17</ span >   < span class = "comment" > THE SOFTWARE.</ 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 HIP_HCC_H</ span ></ div >
< div class = "line" >< a name = "l00021" ></ a >< span class = "lineno" > 21</ span >   < 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 >   < span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00023" ></ a >< span class = "lineno" > 23</ span >   < span class = "preprocessor" > #include < hc.hpp> </ span ></ div >
2016-04-19 22:44:58 +05:30
< div class = "line" >< a name = "l00024" ></ a >< span class = "lineno" > 24</ span >   < span class = "preprocessor" > #include " hip/hcc_detail/hip_util.h" </ span ></ div >
< div class = "line" >< a name = "l00025" ></ a >< span class = "lineno" > 25</ span >   < span class = "preprocessor" > #include " hip/hcc_detail/staging_buffer.h" </ span ></ div >
2016-04-16 14:55:10 +05:30
< div class = "line" >< a name = "l00026" ></ a >< span class = "lineno" > 26</ span >   </ div >
< div class = "line" >< a name = "l00027" ></ a >< span class = "lineno" > 27</ span >   < span class = "preprocessor" > #define HIP_HCC</ span ></ div >
< div class = "line" >< a name = "l00028" ></ a >< span class = "lineno" > 28</ span >   < span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00029" ></ a >< span class = "lineno" > 29</ span >   < span class = "preprocessor" > #if defined(__HCC__) && (__hcc_workweek__ < 1502)</ span ></ div >
< div class = "line" >< a name = "l00030" ></ a >< span class = "lineno" > 30</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #error(" This version of HIP requires a newer version of HCC." );</ span ></ div >
< div class = "line" >< a name = "l00031" ></ a >< span class = "lineno" > 31</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #endif</ span ></ div >
< div class = "line" >< a name = "l00032" ></ a >< span class = "lineno" > 32</ span >   < span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00033" ></ a >< span class = "lineno" > 33</ span >   < span class = "comment" > // #define USE_MEMCPYTOSYMBOL</ span ></ div >
< div class = "line" >< a name = "l00034" ></ a >< span class = "lineno" > 34</ span >   < span class = "comment" > //</ span ></ div >
< div class = "line" >< a name = "l00035" ></ a >< span class = "lineno" > 35</ span >   < span class = "comment" > //Use the new HCC accelerator_view::copy instead of am_copy</ span ></ div >
< div class = "line" >< a name = "l00036" ></ a >< span class = "lineno" > 36</ span >   < span class = "preprocessor" > #define USE_AV_COPY 0</ span ></ div >
< div class = "line" >< a name = "l00037" ></ a >< span class = "lineno" > 37</ span >   < span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00038" ></ a >< span class = "lineno" > 38</ span >   < span class = "comment" > // Compile peer-to-peer support.</ span ></ div >
< div class = "line" >< a name = "l00039" ></ a >< span class = "lineno" > 39</ span >   < span class = "comment" > // > = 2 : use HCC hc:accelerator::get_is_peer</ span ></ div >
< div class = "line" >< a name = "l00040" ></ a >< span class = "lineno" > 40</ span >   < span class = "comment" > // > = 3 : use hc::am_memtracker_update_peers(...)</ span ></ div >
< div class = "line" >< a name = "l00041" ></ a >< span class = "lineno" > 41</ span >   < span class = "preprocessor" > #define USE_PEER_TO_PEER 0</ span ></ div >
< div class = "line" >< a name = "l00042" ></ a >< span class = "lineno" > 42</ span >   < span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00043" ></ a >< span class = "lineno" > 43</ span >   < span class = "comment" > // Use new lock API in HCC:</ span ></ div >
< div class = "line" >< a name = "l00044" ></ a >< span class = "lineno" > 44</ span >   < span class = "preprocessor" > #define USE_HCC_LOCK 0</ span ></ div >
< div class = "line" >< a name = "l00045" ></ a >< span class = "lineno" > 45</ span >   < span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00046" ></ a >< span class = "lineno" > 46</ span >   < span class = "comment" > //#define INLINE static inline</ span ></ div >
< div class = "line" >< a name = "l00047" ></ a >< span class = "lineno" > 47</ span >   </ div >
< div class = "line" >< a name = "l00048" ></ a >< span class = "lineno" > 48</ span >   < span class = "comment" > //---</ span ></ div >
< div class = "line" >< a name = "l00049" ></ a >< span class = "lineno" > 49</ span >   < span class = "comment" > // Environment variables:</ span ></ div >
< div class = "line" >< a name = "l00050" ></ a >< span class = "lineno" > 50</ span >   </ div >
< div class = "line" >< a name = "l00051" ></ a >< span class = "lineno" > 51</ span >   < span class = "comment" > // Intended to distinguish whether an environment variable should be visible only in debug mode, or in debug+release.</ span ></ div >
< div class = "line" >< a name = "l00052" ></ a >< span class = "lineno" > 52</ span >   < span class = "comment" > //static const int debug = 0;</ span ></ div >
< div class = "line" >< a name = "l00053" ></ a >< span class = "lineno" > 53</ span >   < span class = "keyword" > extern</ span > < span class = "keyword" > const</ span > < span class = "keywordtype" > int</ span > release;</ div >
< div class = "line" >< a name = "l00054" ></ a >< span class = "lineno" > 54</ span >   </ div >
< div class = "line" >< a name = "l00055" ></ a >< span class = "lineno" > 55</ span >   < span class = "keyword" > extern</ span > < span class = "keywordtype" > int</ span > HIP_LAUNCH_BLOCKING;</ div >
< div class = "line" >< a name = "l00056" ></ a >< span class = "lineno" > 56</ span >   </ div >
< div class = "line" >< a name = "l00057" ></ a >< span class = "lineno" > 57</ span >   < span class = "keyword" > extern</ span > < span class = "keywordtype" > int</ span > HIP_PRINT_ENV;</ div >
< div class = "line" >< a name = "l00058" ></ a >< span class = "lineno" > 58</ span >   < span class = "keyword" > extern</ span > < span class = "keywordtype" > int</ span > HIP_ATP_MARKER;</ div >
< div class = "line" >< a name = "l00059" ></ a >< span class = "lineno" > 59</ span >   < span class = "comment" > //extern int HIP_TRACE_API;</ span ></ div >
< div class = "line" >< a name = "l00060" ></ a >< span class = "lineno" > 60</ span >   < span class = "keyword" > extern</ span > < span class = "keywordtype" > int</ span > HIP_ATP;</ div >
< div class = "line" >< a name = "l00061" ></ a >< span class = "lineno" > 61</ span >   < span class = "keyword" > extern</ span > < span class = "keywordtype" > int</ span > HIP_DB;</ div >
< div class = "line" >< a name = "l00062" ></ a >< span class = "lineno" > 62</ span >   < span class = "keyword" > extern</ span > < span class = "keywordtype" > int</ span > HIP_STAGING_SIZE; < span class = "comment" > /* size of staging buffers, in KB */</ span ></ div >
< div class = "line" >< a name = "l00063" ></ a >< span class = "lineno" > 63</ span >   < span class = "keyword" > extern</ span > < span class = "keywordtype" > int</ span > HIP_STAGING_BUFFERS; < span class = "comment" > // TODO - remove, two buffers should be enough.</ span ></ div >
< div class = "line" >< a name = "l00064" ></ a >< span class = "lineno" > 64</ span >   < span class = "keyword" > extern</ span > < span class = "keywordtype" > int</ span > HIP_PININPLACE;</ div >
< div class = "line" >< a name = "l00065" ></ a >< span class = "lineno" > 65</ span >   < span class = "keyword" > extern</ span > < span class = "keywordtype" > int</ span > HIP_STREAM_SIGNALS; < span class = "comment" > /* number of signals to allocate at stream creation */</ span ></ div >
< div class = "line" >< a name = "l00066" ></ a >< span class = "lineno" > 66</ span >   < span class = "keyword" > extern</ span > < span class = "keywordtype" > int</ span > HIP_VISIBLE_DEVICES; < span class = "comment" > /* Contains a comma-separated sequence of GPU identifiers */</ span ></ div >
< div class = "line" >< a name = "l00067" ></ a >< span class = "lineno" > 67</ span >   </ div >
< div class = "line" >< a name = "l00068" ></ a >< span class = "lineno" > 68</ span >   </ div >
< div class = "line" >< a name = "l00069" ></ a >< span class = "lineno" > 69</ span >   < span class = "comment" > //---</ span ></ div >
< div class = "line" >< a name = "l00070" ></ a >< span class = "lineno" > 70</ span >   < span class = "comment" > // Chicken bits for disabling functionality to work around potential issues:</ span ></ div >
< div class = "line" >< a name = "l00071" ></ a >< span class = "lineno" > 71</ span >   < span class = "keyword" > extern</ span > < span class = "keywordtype" > int</ span > HIP_DISABLE_HW_KERNEL_DEP;</ div >
< div class = "line" >< a name = "l00072" ></ a >< span class = "lineno" > 72</ span >   < span class = "keyword" > extern</ span > < span class = "keywordtype" > int</ span > HIP_DISABLE_HW_COPY_DEP;</ div >
< div class = "line" >< a name = "l00073" ></ a >< span class = "lineno" > 73</ span >   </ div >
< div class = "line" >< a name = "l00074" ></ a >< span class = "lineno" > 74</ span >   < span class = "keyword" > extern</ span > thread_local < span class = "keywordtype" > int</ span > tls_defaultDevice;</ div >
< div class = "line" >< a name = "l00075" ></ a >< span class = "lineno" > 75</ span >   < span class = "keyword" > extern</ span > thread_local < a class = "code" href = "group__GlobalDefs.html#gadf5010f6e140a53ecbdf949e73e87594" > hipError_t</ a > tls_lastHipError;</ div >
< div class = "line" >< a name = "l00076" ></ a >< span class = "lineno" > 76</ span >   < span class = "keyword" > class </ span >< a class = "code" href = "classihipStream__t.html" > ihipStream_t</ a > ;</ div >
< div class = "line" >< a name = "l00077" ></ a >< span class = "lineno" > 77</ span >   < span class = "keyword" > class </ span >< a class = "code" href = "classihipDevice__t.html" > ihipDevice_t</ a > ;</ div >
< div class = "line" >< a name = "l00078" ></ a >< span class = "lineno" > 78</ span >   </ div >
< div class = "line" >< a name = "l00079" ></ a >< span class = "lineno" > 79</ span >   </ div >
< div class = "line" >< a name = "l00080" ></ a >< span class = "lineno" > 80</ span >   < span class = "comment" > // Color defs for debug messages:</ span ></ div >
< div class = "line" >< a name = "l00081" ></ a >< span class = "lineno" > 81</ span >   < span class = "preprocessor" > #define KNRM " \x1B[0m" </ span ></ div >
< div class = "line" >< a name = "l00082" ></ a >< span class = "lineno" > 82</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #define KRED " \x1B[31m" </ span ></ div >
< div class = "line" >< a name = "l00083" ></ a >< span class = "lineno" > 83</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #define KGRN " \x1B[32m" </ span ></ div >
< div class = "line" >< a name = "l00084" ></ a >< span class = "lineno" > 84</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #define KYEL " \x1B[33m" </ span ></ div >
< div class = "line" >< a name = "l00085" ></ a >< span class = "lineno" > 85</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #define KBLU " \x1B[34m" </ span ></ div >
< div class = "line" >< a name = "l00086" ></ a >< span class = "lineno" > 86</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #define KMAG " \x1B[35m" </ span ></ div >
< div class = "line" >< a name = "l00087" ></ a >< span class = "lineno" > 87</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #define KCYN " \x1B[36m" </ span ></ div >
< div class = "line" >< a name = "l00088" ></ a >< span class = "lineno" > 88</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #define KWHT " \x1B[37m" </ span ></ div >
< div class = "line" >< a name = "l00089" ></ a >< span class = "lineno" > 89</ span >   < span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00090" ></ a >< span class = "lineno" > 90</ span >   < span class = "preprocessor" > #define API_COLOR KGRN</ span ></ div >
< div class = "line" >< a name = "l00091" ></ a >< span class = "lineno" > 91</ span >   < span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00092" ></ a >< span class = "lineno" > 92</ span >   </ div >
< div class = "line" >< a name = "l00093" ></ a >< span class = "lineno" > 93</ span >   < span class = "preprocessor" > #define HIP_HCC </ span ></ div >
< div class = "line" >< a name = "l00094" ></ a >< span class = "lineno" > 94</ span >   < span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00095" ></ a >< span class = "lineno" > 95</ span >   < span class = "comment" > // If set, thread-safety is enforced on all stream functions.</ span ></ div >
< div class = "line" >< a name = "l00096" ></ a >< span class = "lineno" > 96</ span >   < span class = "comment" > // Stream functions will acquire a mutex before entering critical sections.</ span ></ div >
< div class = "line" >< a name = "l00097" ></ a >< span class = "lineno" > 97</ span >   < span class = "preprocessor" > #define STREAM_THREAD_SAFE 1</ span ></ div >
< div class = "line" >< a name = "l00098" ></ a >< span class = "lineno" > 98</ span >   < span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00099" ></ a >< span class = "lineno" > 99</ span >   </ div >
< div class = "line" >< a name = "l00100" ></ a >< span class = "lineno" > 100</ span >   < span class = "preprocessor" > #define DEVICE_THREAD_SAFE 1</ span ></ div >
< div class = "line" >< a name = "l00101" ></ a >< span class = "lineno" > 101</ span >   < span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00102" ></ a >< span class = "lineno" > 102</ span >   < span class = "comment" > // If FORCE_COPY_DEP=1 , HIP runtime will add </ span ></ div >
< div class = "line" >< a name = "l00103" ></ a >< span class = "lineno" > 103</ span >   < span class = "comment" > // synchronization for copy commands in the same stream, regardless of command type.</ span ></ div >
< div class = "line" >< a name = "l00104" ></ a >< span class = "lineno" > 104</ span >   < span class = "comment" > // If FORCE_COPY_DEP=0 data copies of the same kind (H2H, H2D, D2H, D2D) are assumed to be implicitly ordered.</ span ></ div >
< div class = "line" >< a name = "l00105" ></ a >< span class = "lineno" > 105</ span >   < span class = "comment" > // ROCR runtime implementation currently provides this guarantee when using SDMA queues but not </ span ></ div >
< div class = "line" >< a name = "l00106" ></ a >< span class = "lineno" > 106</ span >   < span class = "comment" > // when using shader queues. </ span ></ div >
< div class = "line" >< a name = "l00107" ></ a >< span class = "lineno" > 107</ span >   < span class = "comment" > // TODO - measure if this matters for performance, in particular for back-to-back small copies.</ span ></ div >
< div class = "line" >< a name = "l00108" ></ a >< span class = "lineno" > 108</ span >   < span class = "comment" > // If not, we can simplify the copy dependency tracking by collapsing to a single Copy type, and always forcing dependencies for copy commands.</ span ></ div >
< div class = "line" >< a name = "l00109" ></ a >< span class = "lineno" > 109</ span >   < span class = "preprocessor" > #define FORCE_SAMEDIR_COPY_DEP 1</ span ></ div >
< div class = "line" >< a name = "l00110" ></ a >< span class = "lineno" > 110</ span >   < span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00111" ></ a >< span class = "lineno" > 111</ span >   </ div >
< div class = "line" >< a name = "l00112" ></ a >< span class = "lineno" > 112</ span >   < span class = "comment" > // Compile debug trace mode - this prints debug messages to stderr when env var HIP_DB is set.</ span ></ div >
< div class = "line" >< a name = "l00113" ></ a >< span class = "lineno" > 113</ span >   < span class = "comment" > // May be set to 0 to remove debug if checks - possible code size and performance difference?</ span ></ div >
< div class = "line" >< a name = "l00114" ></ a >< span class = "lineno" > 114</ span >   < span class = "preprocessor" > #define COMPILE_HIP_DB 1</ span ></ div >
< div class = "line" >< a name = "l00115" ></ a >< span class = "lineno" > 115</ span >   < span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00116" ></ a >< span class = "lineno" > 116</ span >   </ div >
< div class = "line" >< a name = "l00117" ></ a >< span class = "lineno" > 117</ span >   < span class = "comment" > // Compile HIP tracing capability.</ span ></ div >
< div class = "line" >< a name = "l00118" ></ a >< span class = "lineno" > 118</ span >   < span class = "comment" > // 0x1 = print a string at function entry with arguments.</ span ></ div >
< div class = "line" >< a name = "l00119" ></ a >< span class = "lineno" > 119</ span >   < span class = "comment" > // 0x2 = prints a simple message with function name + return code when function exits.</ span ></ div >
< div class = "line" >< a name = "l00120" ></ a >< span class = "lineno" > 120</ span >   < span class = "comment" > // 0x3 = print both.</ span ></ div >
< div class = "line" >< a name = "l00121" ></ a >< span class = "lineno" > 121</ span >   < span class = "comment" > // Must be enabled at runtime with HIP_TRACE_API</ span ></ div >
< div class = "line" >< a name = "l00122" ></ a >< span class = "lineno" > 122</ span >   < span class = "preprocessor" > #define COMPILE_HIP_TRACE_API 0x3 </ span ></ div >
< div class = "line" >< a name = "l00123" ></ a >< span class = "lineno" > 123</ span >   < span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00124" ></ a >< span class = "lineno" > 124</ span >   </ div >
< div class = "line" >< a name = "l00125" ></ a >< span class = "lineno" > 125</ span >   < span class = "comment" > // Compile code that generates trace markers for CodeXL ATP at HIP function begin/end.</ span ></ div >
< div class = "line" >< a name = "l00126" ></ a >< span class = "lineno" > 126</ span >   < span class = "comment" > // ATP is standard CodeXL format that includes timestamps for kernels, HSA RT APIs, and HIP APIs.</ span ></ div >
< div class = "line" >< a name = "l00127" ></ a >< span class = "lineno" > 127</ span >   < span class = "preprocessor" > #ifndef COMPILE_HIP_ATP_MARKER</ span ></ div >
< div class = "line" >< a name = "l00128" ></ a >< span class = "lineno" > 128</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #define COMPILE_HIP_ATP_MARKER 0</ span ></ div >
< div class = "line" >< a name = "l00129" ></ a >< span class = "lineno" > 129</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #endif</ span ></ div >
< div class = "line" >< a name = "l00130" ></ a >< span class = "lineno" > 130</ span >   < span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00131" ></ a >< span class = "lineno" > 131</ span >   </ div >
< div class = "line" >< a name = "l00132" ></ a >< span class = "lineno" > 132</ span >   < span class = "comment" > // #include CPP files to produce one object file</ span ></ div >
< div class = "line" >< a name = "l00133" ></ a >< span class = "lineno" > 133</ span >   < span class = "preprocessor" > #define ONE_OBJECT_FILE 0</ span ></ div >
< div class = "line" >< a name = "l00134" ></ a >< span class = "lineno" > 134</ span >   < span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00135" ></ a >< span class = "lineno" > 135</ span >   </ div >
< div class = "line" >< a name = "l00136" ></ a >< span class = "lineno" > 136</ span >   < span class = "comment" > // Compile support for trace markers that are displayed on CodeXL GUI at start/stop of each function boundary.</ span ></ div >
< div class = "line" >< a name = "l00137" ></ a >< span class = "lineno" > 137</ span >   < span class = "comment" > // TODO - currently we print the trace message at the beginning. if we waited, we could also include return codes, and any values returned</ span ></ div >
< div class = "line" >< a name = "l00138" ></ a >< span class = "lineno" > 138</ span >   < span class = "comment" > // through ptr-to-args (ie the pointers allocated by hipMalloc).</ span ></ div >
< div class = "line" >< a name = "l00139" ></ a >< span class = "lineno" > 139</ span >   < span class = "preprocessor" > #if COMPILE_HIP_ATP_MARKER</ span ></ div >
< div class = "line" >< a name = "l00140" ></ a >< span class = "lineno" > 140</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #include " AMDTActivityLogger.h" </ span ></ div >
< div class = "line" >< a name = "l00141" ></ a >< span class = "lineno" > 141</ span >   < span class = "preprocessor" > #define SCOPED_MARKER(markerName,group,userString) amdtScopedMarker(markerName, group, userString)</ span ></ div >
< div class = "line" >< a name = "l00142" ></ a >< span class = "lineno" > 142</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #else </ span ></ div >
< div class = "line" >< a name = "l00143" ></ a >< span class = "lineno" > 143</ span >   < span class = "preprocessor" ></ span >< span class = "comment" > // Swallow scoped markers:</ span ></ div >
< div class = "line" >< a name = "l00144" ></ a >< span class = "lineno" > 144</ span >   < span class = "preprocessor" > #define SCOPED_MARKER(markerName,group,userString) </ span ></ div >
< div class = "line" >< a name = "l00145" ></ a >< span class = "lineno" > 145</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #endif</ span ></ div >
< div class = "line" >< a name = "l00146" ></ a >< span class = "lineno" > 146</ span >   < span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00147" ></ a >< span class = "lineno" > 147</ span >   </ div >
< div class = "line" >< a name = "l00148" ></ a >< span class = "lineno" > 148</ span >   < span class = "preprocessor" > #if COMPILE_HIP_ATP_MARKER || (COMPILE_HIP_TRACE_API & 0x1) </ span ></ div >
< div class = "line" >< a name = "l00149" ></ a >< span class = "lineno" > 149</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #define API_TRACE(...)\</ span ></ div >
< div class = "line" >< a name = "l00150" ></ a >< span class = "lineno" > 150</ span >   < span class = "preprocessor" > {\</ span ></ div >
< div class = "line" >< a name = "l00151" ></ a >< span class = "lineno" > 151</ span >   < span class = "preprocessor" > if (HIP_ATP_MARKER || (COMPILE_HIP_DB && HIP_TRACE_API)) {\</ span ></ div >
< div class = "line" >< a name = "l00152" ></ a >< span class = "lineno" > 152</ span >   < span class = "preprocessor" > std::string s = std::string(__func__) + " (" + ToString(__VA_ARGS__) + ' )' ;\</ span ></ div >
< div class = "line" >< a name = "l00153" ></ a >< span class = "lineno" > 153</ span >   < span class = "preprocessor" > if (COMPILE_HIP_DB && HIP_TRACE_API) {\</ span ></ div >
< div class = "line" >< a name = "l00154" ></ a >< span class = "lineno" > 154</ span >   < span class = "preprocessor" > fprintf (stderr, API_COLOR "<< hip-api: %s\n" KNRM, s.c_str());\</ span ></ div >
< div class = "line" >< a name = "l00155" ></ a >< span class = "lineno" > 155</ span >   < span class = "preprocessor" > }\</ span ></ div >
< div class = "line" >< a name = "l00156" ></ a >< span class = "lineno" > 156</ span >   < span class = "preprocessor" > SCOPED_MARKER(s.c_str(), " HIP" , NULL);\</ span ></ div >
< div class = "line" >< a name = "l00157" ></ a >< span class = "lineno" > 157</ span >   < span class = "preprocessor" > }\</ span ></ div >
< div class = "line" >< a name = "l00158" ></ a >< span class = "lineno" > 158</ span >   < span class = "preprocessor" > }</ span ></ div >
< div class = "line" >< a name = "l00159" ></ a >< span class = "lineno" > 159</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #else</ span ></ div >
< div class = "line" >< a name = "l00160" ></ a >< span class = "lineno" > 160</ span >   < span class = "preprocessor" ></ span >< span class = "comment" > // Swallow API_TRACE</ span ></ div >
< div class = "line" >< a name = "l00161" ></ a >< span class = "lineno" > 161</ span >   < span class = "preprocessor" > #define API_TRACE(...)</ span ></ div >
< div class = "line" >< a name = "l00162" ></ a >< span class = "lineno" > 162</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #endif</ span ></ div >
< div class = "line" >< a name = "l00163" ></ a >< span class = "lineno" > 163</ span >   < span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00164" ></ a >< span class = "lineno" > 164</ span >   </ div >
< div class = "line" >< a name = "l00165" ></ a >< span class = "lineno" > 165</ span >   </ div >
< div class = "line" >< a name = "l00166" ></ a >< span class = "lineno" > 166</ span >   < span class = "comment" > // This macro should be called at the beginning of every HIP API.</ span ></ div >
< div class = "line" >< a name = "l00167" ></ a >< span class = "lineno" > 167</ span >   < span class = "comment" > // It initialies the hip runtime (exactly once), and</ span ></ div >
< div class = "line" >< a name = "l00168" ></ a >< span class = "lineno" > 168</ span >   < span class = "comment" > // generate trace string that can be output to stderr or to ATP file.</ span ></ div >
< div class = "line" >< a name = "l00169" ></ a >< span class = "lineno" > 169</ span >   < span class = "preprocessor" > #define HIP_INIT_API(...) \</ span ></ div >
< div class = "line" >< a name = "l00170" ></ a >< span class = "lineno" > 170</ span >   < span class = "preprocessor" > std::call_once(hip_initialized, ihipInit);\</ span ></ div >
< div class = "line" >< a name = "l00171" ></ a >< span class = "lineno" > 171</ span >   < span class = "preprocessor" > API_TRACE(__VA_ARGS__);</ span ></ div >
< div class = "line" >< a name = "l00172" ></ a >< span class = "lineno" > 172</ span >   < span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00173" ></ a >< span class = "lineno" > 173</ span >   < span class = "preprocessor" > #define ihipLogStatus(_hip_status) \</ span ></ div >
< div class = "line" >< a name = "l00174" ></ a >< span class = "lineno" > 174</ span >   < span class = "preprocessor" > ({\</ span ></ div >
< div class = "line" >< a name = "l00175" ></ a >< span class = "lineno" > 175</ span >   < span class = "preprocessor" > hipError_t _local_hip_status = _hip_status; </ span >< span class = "comment" > /*local copy so _hip_status only evaluated once*/</ span >< span class = "preprocessor" > \</ span ></ div >
< div class = "line" >< a name = "l00176" ></ a >< span class = "lineno" > 176</ span >   < span class = "preprocessor" > tls_lastHipError = _local_hip_status;\</ span ></ div >
< div class = "line" >< a name = "l00177" ></ a >< span class = "lineno" > 177</ span >   < span class = "preprocessor" > \</ span ></ div >
< div class = "line" >< a name = "l00178" ></ a >< span class = "lineno" > 178</ span >   < span class = "preprocessor" > if ((COMPILE_HIP_TRACE_API & 0x2) && HIP_TRACE_API) {\</ span ></ div >
< div class = "line" >< a name = "l00179" ></ a >< span class = "lineno" > 179</ span >   < span class = "preprocessor" > fprintf(stderr, " %ship-api: %-30s ret=%2d (%s)>> \n" KNRM, (_local_hip_status == 0) ? API_COLOR:KRED, __func__, _local_hip_status, ihipErrorString(_local_hip_status));\</ span ></ div >
< div class = "line" >< a name = "l00180" ></ a >< span class = "lineno" > 180</ span >   < span class = "preprocessor" > }\</ span ></ div >
< div class = "line" >< a name = "l00181" ></ a >< span class = "lineno" > 181</ span >   < span class = "preprocessor" > _local_hip_status;\</ span ></ div >
< div class = "line" >< a name = "l00182" ></ a >< span class = "lineno" > 182</ span >   < span class = "preprocessor" > })</ span ></ div >
< div class = "line" >< a name = "l00183" ></ a >< span class = "lineno" > 183</ span >   < span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00184" ></ a >< span class = "lineno" > 184</ span >   </ div >
< div class = "line" >< a name = "l00185" ></ a >< span class = "lineno" > 185</ span >   </ div >
< div class = "line" >< a name = "l00186" ></ a >< span class = "lineno" > 186</ span >   </ div >
< div class = "line" >< a name = "l00187" ></ a >< span class = "lineno" > 187</ span >   < span class = "comment" > //---</ span ></ div >
< div class = "line" >< a name = "l00188" ></ a >< span class = "lineno" > 188</ span >   < span class = "comment" > //HIP_DB Debug flags:</ span ></ div >
< div class = "line" >< a name = "l00189" ></ a >< span class = "lineno" > 189</ span >   < span class = "preprocessor" > #define DB_API 0 </ span >< span class = "comment" > /* 0x01 - shortcut to enable HIP_TRACE_API on single switch */</ span >< span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00190" ></ a >< span class = "lineno" > 190</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #define DB_SYNC 1 </ span >< span class = "comment" > /* 0x02 - trace synchronization pieces */</ span >< span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00191" ></ a >< span class = "lineno" > 191</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #define DB_MEM 2 </ span >< span class = "comment" > /* 0x04 - trace memory allocation / deallocation */</ span >< span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00192" ></ a >< span class = "lineno" > 192</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #define DB_COPY1 3 </ span >< span class = "comment" > /* 0x08 - trace memory copy commands. . */</ span >< span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00193" ></ a >< span class = "lineno" > 193</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #define DB_SIGNAL 4 </ span >< span class = "comment" > /* 0x10 - trace signal pool commands */</ span >< span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00194" ></ a >< span class = "lineno" > 194</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #define DB_COPY2 5 </ span >< span class = "comment" > /* 0x20 - trace memory copy commands. Detailed. */</ span >< span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00195" ></ a >< span class = "lineno" > 195</ span >   < span class = "preprocessor" ></ span >< span class = "comment" > // When adding a new debug flag, also add to the char name table below.</ span ></ div >
< div class = "line" >< a name = "l00196" ></ a >< span class = "lineno" > 196</ span >   </ div >
< div class = "line" >< a name = "l00197" ></ a >< span class = "lineno" > 197</ span >   < span class = "keyword" > static</ span > < span class = "keyword" > const</ span > < span class = "keywordtype" > char</ span > *dbName [] =</ div >
< div class = "line" >< a name = "l00198" ></ a >< span class = "lineno" > 198</ span >   {</ div >
< div class = "line" >< a name = "l00199" ></ a >< span class = "lineno" > 199</ span >   KNRM < span class = "stringliteral" > " hip-api" </ span > , < span class = "comment" > // not used, </ span ></ div >
< div class = "line" >< a name = "l00200" ></ a >< span class = "lineno" > 200</ span >   KYEL < span class = "stringliteral" > " hip-sync" </ span > ,</ div >
< div class = "line" >< a name = "l00201" ></ a >< span class = "lineno" > 201</ span >   KCYN < span class = "stringliteral" > " hip-mem" </ span > ,</ div >
< div class = "line" >< a name = "l00202" ></ a >< span class = "lineno" > 202</ span >   KMAG < span class = "stringliteral" > " hip-copy1" </ span > ,</ div >
< div class = "line" >< a name = "l00203" ></ a >< span class = "lineno" > 203</ span >   KRED < span class = "stringliteral" > " hip-signal" </ span > ,</ div >
< div class = "line" >< a name = "l00204" ></ a >< span class = "lineno" > 204</ span >   KNRM < span class = "stringliteral" > " hip-copy2" </ span > ,</ div >
< div class = "line" >< a name = "l00205" ></ a >< span class = "lineno" > 205</ span >   };</ div >
< div class = "line" >< a name = "l00206" ></ a >< span class = "lineno" > 206</ span >   </ div >
< div class = "line" >< a name = "l00207" ></ a >< span class = "lineno" > 207</ span >   < span class = "preprocessor" > #if COMPILE_HIP_DB</ span ></ div >
< div class = "line" >< a name = "l00208" ></ a >< span class = "lineno" > 208</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #define tprintf(trace_level, ...) {\</ span ></ div >
< div class = "line" >< a name = "l00209" ></ a >< span class = "lineno" > 209</ span >   < span class = "preprocessor" > if (HIP_DB & (1<< (trace_level))) {\</ span ></ div >
< div class = "line" >< a name = "l00210" ></ a >< span class = "lineno" > 210</ span >   < span class = "preprocessor" > fprintf (stderr, " %s:" , dbName[trace_level]); \</ span ></ div >
< div class = "line" >< a name = "l00211" ></ a >< span class = "lineno" > 211</ span >   < span class = "preprocessor" > fprintf (stderr, __VA_ARGS__);\</ span ></ div >
< div class = "line" >< a name = "l00212" ></ a >< span class = "lineno" > 212</ span >   < span class = "preprocessor" > fprintf (stderr, " %s" , KNRM); \</ span ></ div >
< div class = "line" >< a name = "l00213" ></ a >< span class = "lineno" > 213</ span >   < span class = "preprocessor" > }\</ span ></ div >
< div class = "line" >< a name = "l00214" ></ a >< span class = "lineno" > 214</ span >   < span class = "preprocessor" > }</ span ></ div >
< div class = "line" >< a name = "l00215" ></ a >< span class = "lineno" > 215</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #else </ span ></ div >
< div class = "line" >< a name = "l00216" ></ a >< span class = "lineno" > 216</ span >   < span class = "preprocessor" ></ span >< span class = "comment" > /* Compile to empty code */</ span ></ div >
< div class = "line" >< a name = "l00217" ></ a >< span class = "lineno" > 217</ span >   < span class = "preprocessor" > #define tprintf(trace_level, ...) </ span ></ div >
< div class = "line" >< a name = "l00218" ></ a >< span class = "lineno" > 218</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #endif</ span ></ div >
< div class = "line" >< a name = "l00219" ></ a >< span class = "lineno" > 219</ span >   < span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00220" ></ a >< span class = "lineno" >< a class = "line" href = "classihipException.html" > 220</ a ></ span >   < span class = "keyword" > class </ span >< a class = "code" href = "classihipException.html" > ihipException</ a > : < span class = "keyword" > public</ span > std::exception</ div >
< div class = "line" >< a name = "l00221" ></ a >< span class = "lineno" > 221</ span >   {</ div >
< div class = "line" >< a name = "l00222" ></ a >< span class = "lineno" > 222</ span >   < span class = "keyword" > public</ span > :</ div >
< div class = "line" >< a name = "l00223" ></ a >< span class = "lineno" > 223</ span >   < a class = "code" href = "classihipException.html" > ihipException</ a > (< a class = "code" href = "group__GlobalDefs.html#gadf5010f6e140a53ecbdf949e73e87594" > hipError_t</ a > e) : _code(e) {};</ div >
< div class = "line" >< a name = "l00224" ></ a >< span class = "lineno" > 224</ span >   </ div >
< div class = "line" >< a name = "l00225" ></ a >< span class = "lineno" > 225</ span >   < a class = "code" href = "group__GlobalDefs.html#gadf5010f6e140a53ecbdf949e73e87594" > hipError_t</ a > _code; </ div >
< div class = "line" >< a name = "l00226" ></ a >< span class = "lineno" > 226</ span >   };</ div >
< div class = "line" >< a name = "l00227" ></ a >< span class = "lineno" > 227</ span >   </ div >
< div class = "line" >< a name = "l00228" ></ a >< span class = "lineno" > 228</ span >   </ div >
< div class = "line" >< a name = "l00229" ></ a >< span class = "lineno" > 229</ span >   < span class = "preprocessor" > #ifdef __cplusplus</ span ></ div >
< div class = "line" >< a name = "l00230" ></ a >< span class = "lineno" > 230</ span >   < span class = "preprocessor" ></ span >< span class = "keyword" > extern</ span > < span class = "stringliteral" > " C" </ span > {</ div >
< div class = "line" >< a name = "l00231" ></ a >< span class = "lineno" > 231</ span >   < span class = "preprocessor" > #endif</ span ></ div >
< div class = "line" >< a name = "l00232" ></ a >< span class = "lineno" > 232</ span >   < span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00233" ></ a >< span class = "lineno" > 233</ span >   < span class = "keyword" > typedef</ span > < span class = "keyword" > class </ span >< a class = "code" href = "classihipStream__t.html" > ihipStream_t</ a > * hipStream_t;</ div >
< div class = "line" >< a name = "l00234" ></ a >< span class = "lineno" > 234</ span >   < span class = "comment" > //typedef struct hipEvent_t {</ span ></ div >
< div class = "line" >< a name = "l00235" ></ a >< span class = "lineno" > 235</ span >   < span class = "comment" > // struct ihipEvent_t *_handle;</ span ></ div >
< div class = "line" >< a name = "l00236" ></ a >< span class = "lineno" > 236</ span >   < span class = "comment" > //} hipEvent_t;</ span ></ div >
< div class = "line" >< a name = "l00237" ></ a >< span class = "lineno" > 237</ span >   </ div >
< div class = "line" >< a name = "l00238" ></ a >< span class = "lineno" > 238</ span >   < span class = "preprocessor" > #ifdef __cplusplus</ span ></ div >
< div class = "line" >< a name = "l00239" ></ a >< span class = "lineno" > 239</ span >   < span class = "preprocessor" ></ span > }</ div >
< div class = "line" >< a name = "l00240" ></ a >< span class = "lineno" > 240</ span >   < span class = "preprocessor" > #endif</ span ></ div >
< div class = "line" >< a name = "l00241" ></ a >< span class = "lineno" > 241</ span >   < span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00242" ></ a >< span class = "lineno" > 242</ span >   < span class = "keyword" > const</ span > hipStream_t hipStreamNull = 0x0;</ div >
< div class = "line" >< a name = "l00243" ></ a >< span class = "lineno" > 243</ span >   </ div >
< div class = "line" >< a name = "l00244" ></ a >< span class = "lineno" > 244</ span >   </ div >
< div class = "line" >< a name = "l00245" ></ a >< span class = "lineno" > 245</ span >   < span class = "keyword" > enum</ span > ihipCommand_t {</ div >
< div class = "line" >< a name = "l00246" ></ a >< span class = "lineno" > 246</ span >   ihipCommandCopyH2H,</ div >
< div class = "line" >< a name = "l00247" ></ a >< span class = "lineno" > 247</ span >   ihipCommandCopyH2D, </ div >
< div class = "line" >< a name = "l00248" ></ a >< span class = "lineno" > 248</ span >   ihipCommandCopyD2H,</ div >
< div class = "line" >< a name = "l00249" ></ a >< span class = "lineno" > 249</ span >   ihipCommandCopyD2D,</ div >
< div class = "line" >< a name = "l00250" ></ a >< span class = "lineno" > 250</ span >   ihipCommandKernel,</ div >
< div class = "line" >< a name = "l00251" ></ a >< span class = "lineno" > 251</ span >   };</ div >
< div class = "line" >< a name = "l00252" ></ a >< span class = "lineno" > 252</ span >   </ div >
< div class = "line" >< a name = "l00253" ></ a >< span class = "lineno" > 253</ span >   < span class = "keyword" > static</ span > < span class = "keyword" > const</ span > < span class = "keywordtype" > char</ span > * ihipCommandName[] = {</ div >
< div class = "line" >< a name = "l00254" ></ a >< span class = "lineno" > 254</ span >   < span class = "stringliteral" > " CopyH2H" </ span > , < span class = "stringliteral" > " CopyH2D" </ span > , < span class = "stringliteral" > " CopyD2H" </ span > , < span class = "stringliteral" > " CopyD2D" </ span > , < span class = "stringliteral" > " Kernel" </ span ></ div >
< div class = "line" >< a name = "l00255" ></ a >< span class = "lineno" > 255</ span >   };</ div >
< div class = "line" >< a name = "l00256" ></ a >< span class = "lineno" > 256</ span >   </ div >
< div class = "line" >< a name = "l00257" ></ a >< span class = "lineno" > 257</ span >   </ div >
< div class = "line" >< a name = "l00258" ></ a >< span class = "lineno" > 258</ span >   </ div >
< div class = "line" >< a name = "l00259" ></ a >< span class = "lineno" > 259</ span >   < span class = "keyword" > typedef</ span > uint64_t SIGSEQNUM;</ div >
< div class = "line" >< a name = "l00260" ></ a >< span class = "lineno" > 260</ span >   </ div >
< div class = "line" >< a name = "l00261" ></ a >< span class = "lineno" > 261</ span >   < span class = "comment" > //---</ span ></ div >
< div class = "line" >< a name = "l00262" ></ a >< span class = "lineno" > 262</ span >   < span class = "comment" > // Small wrapper around signals.</ span ></ div >
< div class = "line" >< a name = "l00263" ></ a >< span class = "lineno" > 263</ span >   < span class = "comment" > // Designed to be used from stream.</ span ></ div >
< div class = "line" >< a name = "l00264" ></ a >< span class = "lineno" > 264</ span >   < span class = "comment" > // TODO-someday refactor this class so it can be stored in a vector<> </ span ></ div >
< div class = "line" >< a name = "l00265" ></ a >< span class = "lineno" > 265</ span >   < span class = "comment" > // we already store the index here so we can use for garbage collection.</ span ></ div >
< div class = "line" >< a name = "l00266" ></ a >< span class = "lineno" >< a class = "line" href = "structihipSignal__t.html" > 266</ a ></ span >   < span class = "keyword" > struct </ span >< a class = "code" href = "structihipSignal__t.html" > ihipSignal_t</ a > {</ div >
< div class = "line" >< a name = "l00267" ></ a >< span class = "lineno" > 267</ span >   hsa_signal_t _hsa_signal; < span class = "comment" > // hsa signal handle</ span ></ div >
< div class = "line" >< a name = "l00268" ></ a >< span class = "lineno" > 268</ span >   < span class = "keywordtype" > int</ span > _index; < span class = "comment" > // Index in pool, used for garbage collection.</ span ></ div >
< div class = "line" >< a name = "l00269" ></ a >< span class = "lineno" > 269</ span >   SIGSEQNUM _sig_id; < span class = "comment" > // unique sequentially increasing ID.</ span ></ div >
< div class = "line" >< a name = "l00270" ></ a >< span class = "lineno" > 270</ span >   </ div >
< div class = "line" >< a name = "l00271" ></ a >< span class = "lineno" > 271</ span >   < a class = "code" href = "structihipSignal__t.html" > ihipSignal_t</ a > ();</ div >
< div class = "line" >< a name = "l00272" ></ a >< span class = "lineno" > 272</ span >   ~< a class = "code" href = "structihipSignal__t.html" > ihipSignal_t</ a > ();</ div >
< div class = "line" >< a name = "l00273" ></ a >< span class = "lineno" > 273</ span >   </ div >
< div class = "line" >< a name = "l00274" ></ a >< span class = "lineno" > 274</ span >   < span class = "keywordtype" > void</ span > release();</ div >
< div class = "line" >< a name = "l00275" ></ a >< span class = "lineno" > 275</ span >   };</ div >
< div class = "line" >< a name = "l00276" ></ a >< span class = "lineno" > 276</ span >   </ div >
< div class = "line" >< a name = "l00277" ></ a >< span class = "lineno" > 277</ span >   </ div >
< div class = "line" >< a name = "l00278" ></ a >< span class = "lineno" > 278</ span >   < span class = "comment" > // Used to remove lock, for performance or stimulating bugs.</ span ></ div >
< div class = "line" >< a name = "l00279" ></ a >< span class = "lineno" >< a class = "line" href = "classFakeMutex.html" > 279</ a ></ span >   < span class = "keyword" > class </ span >< a class = "code" href = "classFakeMutex.html" > FakeMutex</ a ></ div >
< div class = "line" >< a name = "l00280" ></ a >< span class = "lineno" > 280</ span >   {</ div >
< div class = "line" >< a name = "l00281" ></ a >< span class = "lineno" > 281</ span >   < span class = "keyword" > public</ span > :</ div >
< div class = "line" >< a name = "l00282" ></ a >< span class = "lineno" > 282</ span >   < span class = "keywordtype" > void</ span > lock() { }</ div >
< div class = "line" >< a name = "l00283" ></ a >< span class = "lineno" > 283</ span >   < span class = "keywordtype" > bool</ span > try_lock() {< span class = "keywordflow" > return</ span > < span class = "keyword" > true</ span > ; }</ div >
< div class = "line" >< a name = "l00284" ></ a >< span class = "lineno" > 284</ span >   < span class = "keywordtype" > void</ span > unlock() { }</ div >
< div class = "line" >< a name = "l00285" ></ a >< span class = "lineno" > 285</ span >   };</ div >
< div class = "line" >< a name = "l00286" ></ a >< span class = "lineno" > 286</ span >   </ div >
< div class = "line" >< a name = "l00287" ></ a >< span class = "lineno" > 287</ span >   </ div >
< div class = "line" >< a name = "l00288" ></ a >< span class = "lineno" > 288</ span >   < span class = "preprocessor" > #if STREAM_THREAD_SAFE</ span ></ div >
< div class = "line" >< a name = "l00289" ></ a >< span class = "lineno" > 289</ span >   < span class = "preprocessor" ></ span >< span class = "keyword" > typedef</ span > std::mutex StreamMutex;</ div >
< div class = "line" >< a name = "l00290" ></ a >< span class = "lineno" > 290</ span >   < span class = "preprocessor" > #else</ span ></ div >
< div class = "line" >< a name = "l00291" ></ a >< span class = "lineno" > 291</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #warning " Stream thread-safe disabled" </ span ></ div >
< div class = "line" >< a name = "l00292" ></ a >< span class = "lineno" > 292</ span >   < span class = "preprocessor" ></ span >< span class = "keyword" > typedef</ span > < a class = "code" href = "classFakeMutex.html" > FakeMutex</ a > StreamMutex;</ div >
< div class = "line" >< a name = "l00293" ></ a >< span class = "lineno" > 293</ span >   < span class = "preprocessor" > #endif</ span ></ div >
< div class = "line" >< a name = "l00294" ></ a >< span class = "lineno" > 294</ span >   < span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00295" ></ a >< span class = "lineno" > 295</ span >   < span class = "preprocessor" > #if DEVICE_THREAD_SAFE</ span ></ div >
< div class = "line" >< a name = "l00296" ></ a >< span class = "lineno" > 296</ span >   < span class = "preprocessor" ></ span >< span class = "keyword" > typedef</ span > std::mutex DeviceMutex;</ div >
< div class = "line" >< a name = "l00297" ></ a >< span class = "lineno" > 297</ span >   < span class = "preprocessor" > #else</ span ></ div >
< div class = "line" >< a name = "l00298" ></ a >< span class = "lineno" > 298</ span >   < span class = "preprocessor" ></ span >< span class = "keyword" > typedef</ span > < a class = "code" href = "classFakeMutex.html" > FakeMutex</ a > DeviceMutex;</ div >
< div class = "line" >< a name = "l00299" ></ a >< span class = "lineno" > 299</ span >   < span class = "preprocessor" > #warning " Device thread-safe disabled" </ span ></ div >
< div class = "line" >< a name = "l00300" ></ a >< span class = "lineno" > 300</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #endif</ span ></ div >
< div class = "line" >< a name = "l00301" ></ a >< span class = "lineno" > 301</ span >   < span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00302" ></ a >< span class = "lineno" > 302</ span >   < span class = "comment" > //</ span ></ div >
< div class = "line" >< a name = "l00303" ></ a >< span class = "lineno" > 303</ span >   < span class = "comment" > //---</ span ></ div >
< div class = "line" >< a name = "l00304" ></ a >< span class = "lineno" > 304</ span >   < span class = "comment" > // Protects access to the member _data with a lock acquired on contruction/destruction.</ span ></ div >
< div class = "line" >< a name = "l00305" ></ a >< span class = "lineno" > 305</ span >   < span class = "comment" > // T must contain a _mutex field which meets the BasicLockable requirements (lock/unlock)</ span ></ div >
< div class = "line" >< a name = "l00306" ></ a >< span class = "lineno" > 306</ span >   < span class = "keyword" > template</ span > < < span class = "keyword" > typename</ span > T> </ div >
< div class = "line" >< a name = "l00307" ></ a >< span class = "lineno" >< a class = "line" href = "classLockedAccessor.html" > 307</ a ></ span >   < span class = "keyword" > class </ span >< a class = "code" href = "classLockedAccessor.html" > LockedAccessor</ a ></ div >
< div class = "line" >< a name = "l00308" ></ a >< span class = "lineno" > 308</ span >   {</ div >
< div class = "line" >< a name = "l00309" ></ a >< span class = "lineno" > 309</ span >   < span class = "keyword" > public</ span > :</ div >
< div class = "line" >< a name = "l00310" ></ a >< span class = "lineno" > 310</ span >   < a class = "code" href = "classLockedAccessor.html" > LockedAccessor</ a > (T & criticalData, < span class = "keywordtype" > bool</ span > autoUnlock=< span class = "keyword" > true</ span > ) : </ div >
< div class = "line" >< a name = "l00311" ></ a >< span class = "lineno" > 311</ span >   _criticalData(& criticalData),</ div >
< div class = "line" >< a name = "l00312" ></ a >< span class = "lineno" > 312</ span >   _autoUnlock(autoUnlock)</ div >
< div class = "line" >< a name = "l00313" ></ a >< span class = "lineno" > 313</ span >   </ div >
< div class = "line" >< a name = "l00314" ></ a >< span class = "lineno" > 314</ span >   {</ div >
< div class = "line" >< a name = "l00315" ></ a >< span class = "lineno" > 315</ span >   _criticalData-> _mutex.lock();</ div >
< div class = "line" >< a name = "l00316" ></ a >< span class = "lineno" > 316</ span >   };</ div >
< div class = "line" >< a name = "l00317" ></ a >< span class = "lineno" > 317</ span >   </ div >
< div class = "line" >< a name = "l00318" ></ a >< span class = "lineno" > 318</ span >   ~< a class = "code" href = "classLockedAccessor.html" > LockedAccessor</ a > () </ div >
< div class = "line" >< a name = "l00319" ></ a >< span class = "lineno" > 319</ span >   {</ div >
< div class = "line" >< a name = "l00320" ></ a >< span class = "lineno" > 320</ span >   < span class = "keywordflow" > if</ span > (_autoUnlock) {</ div >
< div class = "line" >< a name = "l00321" ></ a >< span class = "lineno" > 321</ span >   _criticalData-> _mutex.unlock();</ div >
< div class = "line" >< a name = "l00322" ></ a >< span class = "lineno" > 322</ span >   }</ div >
< div class = "line" >< a name = "l00323" ></ a >< span class = "lineno" > 323</ span >   }</ div >
< div class = "line" >< a name = "l00324" ></ a >< span class = "lineno" > 324</ span >   </ div >
< div class = "line" >< a name = "l00325" ></ a >< span class = "lineno" > 325</ span >   < span class = "keywordtype" > void</ span > unlock() </ div >
< div class = "line" >< a name = "l00326" ></ a >< span class = "lineno" > 326</ span >   {</ div >
< div class = "line" >< a name = "l00327" ></ a >< span class = "lineno" > 327</ span >   _criticalData-> _mutex.unlock();</ div >
< div class = "line" >< a name = "l00328" ></ a >< span class = "lineno" > 328</ span >   }</ div >
< div class = "line" >< a name = "l00329" ></ a >< span class = "lineno" > 329</ span >   </ div >
< div class = "line" >< a name = "l00330" ></ a >< span class = "lineno" > 330</ span >   < span class = "comment" > // Syntactic sugar so -> can be used to get the underlying type.</ span ></ div >
< div class = "line" >< a name = "l00331" ></ a >< span class = "lineno" > 331</ span >   T *operator-> () { < span class = "keywordflow" > return</ span > _criticalData; };</ div >
< div class = "line" >< a name = "l00332" ></ a >< span class = "lineno" > 332</ span >   </ div >
< div class = "line" >< a name = "l00333" ></ a >< span class = "lineno" > 333</ span >   < span class = "keyword" > private</ span > :</ div >
< div class = "line" >< a name = "l00334" ></ a >< span class = "lineno" > 334</ span >   T *_criticalData;</ div >
< div class = "line" >< a name = "l00335" ></ a >< span class = "lineno" > 335</ span >   < span class = "keywordtype" > bool</ span > _autoUnlock;</ div >
< div class = "line" >< a name = "l00336" ></ a >< span class = "lineno" > 336</ span >   };</ div >
< div class = "line" >< a name = "l00337" ></ a >< span class = "lineno" > 337</ span >   </ div >
< div class = "line" >< a name = "l00338" ></ a >< span class = "lineno" > 338</ span >   </ div >
< div class = "line" >< a name = "l00339" ></ a >< span class = "lineno" > 339</ span >   < span class = "keyword" > template</ span > < < span class = "keyword" > typename</ span > MUTEX_TYPE> </ div >
< div class = "line" >< a name = "l00340" ></ a >< span class = "lineno" >< a class = "line" href = "structLockedBase.html" > 340</ a ></ span >   < span class = "keyword" > struct </ span >< a class = "code" href = "structLockedBase.html" > LockedBase</ a > {</ div >
< div class = "line" >< a name = "l00341" ></ a >< span class = "lineno" > 341</ span >   </ div >
< div class = "line" >< a name = "l00342" ></ a >< span class = "lineno" > 342</ span >   < span class = "comment" > // Experts-only interface for explicit locking. </ span ></ div >
< div class = "line" >< a name = "l00343" ></ a >< span class = "lineno" > 343</ span >   < span class = "comment" > // Most uses should use the lock-accessor.</ span ></ div >
< div class = "line" >< a name = "l00344" ></ a >< span class = "lineno" > 344</ span >   < span class = "keywordtype" > void</ span > lock() { _mutex.lock(); }</ div >
< div class = "line" >< a name = "l00345" ></ a >< span class = "lineno" > 345</ span >   < span class = "keywordtype" > void</ span > unlock() { _mutex.unlock(); }</ div >
< div class = "line" >< a name = "l00346" ></ a >< span class = "lineno" > 346</ span >   </ div >
< div class = "line" >< a name = "l00347" ></ a >< span class = "lineno" > 347</ span >   MUTEX_TYPE _mutex;</ div >
< div class = "line" >< a name = "l00348" ></ a >< span class = "lineno" > 348</ span >   };</ div >
< div class = "line" >< a name = "l00349" ></ a >< span class = "lineno" > 349</ span >   </ div >
< div class = "line" >< a name = "l00350" ></ a >< span class = "lineno" > 350</ span >   </ div >
< div class = "line" >< a name = "l00351" ></ a >< span class = "lineno" > 351</ span >   < span class = "keyword" > template</ span > < < span class = "keyword" > typename</ span > MUTEX_TYPE> </ div >
< div class = "line" >< a name = "l00352" ></ a >< span class = "lineno" >< a class = "line" href = "classihipStreamCriticalBase__t.html" > 352</ a ></ span >   < span class = "keyword" > class </ span >< a class = "code" href = "classihipStreamCriticalBase__t.html" > ihipStreamCriticalBase_t</ a > : < span class = "keyword" > public</ span > < a class = "code" href = "structLockedBase.html" > LockedBase</ a > < MUTEX_TYPE> </ div >
< div class = "line" >< a name = "l00353" ></ a >< span class = "lineno" > 353</ span >   {</ div >
< div class = "line" >< a name = "l00354" ></ a >< span class = "lineno" > 354</ span >   < span class = "keyword" > public</ span > :</ div >
< div class = "line" >< a name = "l00355" ></ a >< span class = "lineno" > 355</ span >   < a class = "code" href = "classihipStreamCriticalBase__t.html" > ihipStreamCriticalBase_t</ a > () :</ div >
< div class = "line" >< a name = "l00356" ></ a >< span class = "lineno" > 356</ span >   _last_command_type(ihipCommandCopyH2H),</ div >
< div class = "line" >< a name = "l00357" ></ a >< span class = "lineno" > 357</ span >   _last_copy_signal(NULL),</ div >
< div class = "line" >< a name = "l00358" ></ a >< span class = "lineno" > 358</ span >   _signalCursor(0),</ div >
< div class = "line" >< a name = "l00359" ></ a >< span class = "lineno" > 359</ span >   _oldest_live_sig_id(1),</ div >
< div class = "line" >< a name = "l00360" ></ a >< span class = "lineno" > 360</ span >   _stream_sig_id(0)</ div >
< div class = "line" >< a name = "l00361" ></ a >< span class = "lineno" > 361</ span >   {</ div >
< div class = "line" >< a name = "l00362" ></ a >< span class = "lineno" > 362</ span >   _signalPool.resize(HIP_STREAM_SIGNALS > 0 ? HIP_STREAM_SIGNALS : 1);</ div >
< div class = "line" >< a name = "l00363" ></ a >< span class = "lineno" > 363</ span >   };</ div >
< div class = "line" >< a name = "l00364" ></ a >< span class = "lineno" > 364</ span >   </ div >
< div class = "line" >< a name = "l00365" ></ a >< span class = "lineno" > 365</ span >   ~< a class = "code" href = "classihipStreamCriticalBase__t.html" > ihipStreamCriticalBase_t</ a > () {</ div >
< div class = "line" >< a name = "l00366" ></ a >< span class = "lineno" > 366</ span >   _signalPool.clear();</ div >
< div class = "line" >< a name = "l00367" ></ a >< span class = "lineno" > 367</ span >   }</ div >
< div class = "line" >< a name = "l00368" ></ a >< span class = "lineno" > 368</ span >   </ div >
< div class = "line" >< a name = "l00369" ></ a >< span class = "lineno" > 369</ span >   < a class = "code" href = "classihipStreamCriticalBase__t.html" > ihipStreamCriticalBase_t< StreamMutex> </ a > * mlock() { < a class = "code" href = "structLockedBase.html" > LockedBase< MUTEX_TYPE> ::lock</ a > (); < span class = "keywordflow" > return</ span > < span class = "keyword" > this</ span > ;};</ div >
< div class = "line" >< a name = "l00370" ></ a >< span class = "lineno" > 370</ span >   </ div >
< div class = "line" >< a name = "l00371" ></ a >< span class = "lineno" > 371</ span >   </ div >
< div class = "line" >< a name = "l00372" ></ a >< span class = "lineno" > 372</ span >   < span class = "keyword" > public</ span > :</ div >
< div class = "line" >< a name = "l00373" ></ a >< span class = "lineno" > 373</ span >   < span class = "comment" > // Critical Data:</ span ></ div >
< div class = "line" >< a name = "l00374" ></ a >< span class = "lineno" > 374</ span >   ihipCommand_t _last_command_type; < span class = "comment" > // type of the last command</ span ></ div >
< div class = "line" >< a name = "l00375" ></ a >< span class = "lineno" > 375</ span >   </ div >
< div class = "line" >< a name = "l00376" ></ a >< span class = "lineno" > 376</ span >   < span class = "comment" > // signal of last copy command sent to the stream.</ span ></ div >
< div class = "line" >< a name = "l00377" ></ a >< span class = "lineno" > 377</ span >   < span class = "comment" > // May be NULL, indicating the previous command has completley finished and future commands don' t need to create a dependency.</ span ></ div >
< div class = "line" >< a name = "l00378" ></ a >< span class = "lineno" > 378</ span >   < span class = "comment" > // Copy can be either H2D or D2H.</ span ></ div >
< div class = "line" >< a name = "l00379" ></ a >< span class = "lineno" > 379</ span >   < a class = "code" href = "structihipSignal__t.html" > ihipSignal_t</ a > *_last_copy_signal;</ div >
< div class = "line" >< a name = "l00380" ></ a >< span class = "lineno" > 380</ span >   </ div >
< div class = "line" >< a name = "l00381" ></ a >< span class = "lineno" > 381</ span >   hc::completion_future _last_kernel_future; < span class = "comment" > // Completion future of last kernel command sent to GPU.</ span ></ div >
< div class = "line" >< a name = "l00382" ></ a >< span class = "lineno" > 382</ span >   </ div >
< div class = "line" >< a name = "l00383" ></ a >< span class = "lineno" > 383</ span >   < span class = "comment" > // Signal pool:</ span ></ div >
< div class = "line" >< a name = "l00384" ></ a >< span class = "lineno" > 384</ span >   < span class = "keywordtype" > int</ span > _signalCursor;</ div >
< div class = "line" >< a name = "l00385" ></ a >< span class = "lineno" > 385</ span >   SIGSEQNUM _oldest_live_sig_id; < span class = "comment" > // oldest live seq_id, anything < this can be allocated.</ span ></ div >
< div class = "line" >< a name = "l00386" ></ a >< span class = "lineno" > 386</ span >   std::deque< ihipSignal_t> _signalPool; < span class = "comment" > // Pool of signals for use by this stream.</ span ></ div >
< div class = "line" >< a name = "l00387" ></ a >< span class = "lineno" > 387</ span >   </ div >
< div class = "line" >< a name = "l00388" ></ a >< span class = "lineno" > 388</ span >   </ div >
< div class = "line" >< a name = "l00389" ></ a >< span class = "lineno" > 389</ span >   SIGSEQNUM _stream_sig_id; < span class = "comment" > // Monotonically increasing unique signal id.</ span ></ div >
< div class = "line" >< a name = "l00390" ></ a >< span class = "lineno" > 390</ span >   };</ div >
< div class = "line" >< a name = "l00391" ></ a >< span class = "lineno" > 391</ span >   </ div >
< div class = "line" >< a name = "l00392" ></ a >< span class = "lineno" > 392</ span >   </ div >
< div class = "line" >< a name = "l00393" ></ a >< span class = "lineno" > 393</ span >   < span class = "keyword" > typedef</ span > < a class = "code" href = "classihipStreamCriticalBase__t.html" > ihipStreamCriticalBase_t< StreamMutex> </ a > < a class = "code" href = "classihipStreamCriticalBase__t.html" > ihipStreamCritical_t</ a > ; </ div >
< div class = "line" >< a name = "l00394" ></ a >< span class = "lineno" > 394</ span >   < span class = "keyword" > typedef</ span > < a class = "code" href = "classLockedAccessor.html" > LockedAccessor< ihipStreamCritical_t> </ a > < a class = "code" href = "classLockedAccessor.html" > LockedAccessor_StreamCrit_t</ a > ;</ div >
< div class = "line" >< a name = "l00395" ></ a >< span class = "lineno" > 395</ span >   </ div >
< div class = "line" >< a name = "l00396" ></ a >< span class = "lineno" > 396</ span >   </ div >
< div class = "line" >< a name = "l00397" ></ a >< span class = "lineno" > 397</ span >   </ div >
< div class = "line" >< a name = "l00398" ></ a >< span class = "lineno" > 398</ span >   < span class = "comment" > // Internal stream structure.</ span ></ div >
< div class = "line" >< a name = "l00399" ></ a >< span class = "lineno" >< a class = "line" href = "classihipStream__t.html" > 399</ a ></ span >   < span class = "keyword" > class </ span >< a class = "code" href = "classihipStream__t.html" > ihipStream_t</ a > {</ div >
< div class = "line" >< a name = "l00400" ></ a >< span class = "lineno" > 400</ span >   < span class = "keyword" > public</ span > :</ div >
< div class = "line" >< a name = "l00401" ></ a >< span class = "lineno" > 401</ span >   < span class = "keyword" > typedef</ span > uint64_t SeqNum_t ;</ div >
< div class = "line" >< a name = "l00402" ></ a >< span class = "lineno" > 402</ span >   </ div >
< div class = "line" >< a name = "l00403" ></ a >< span class = "lineno" > 403</ span >   < a class = "code" href = "classihipStream__t.html" > ihipStream_t</ a > (< span class = "keywordtype" > unsigned</ span > device_index, hc::accelerator_view av, < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > int</ span > flags);</ div >
< div class = "line" >< a name = "l00404" ></ a >< span class = "lineno" > 404</ span >   ~< a class = "code" href = "classihipStream__t.html" > ihipStream_t</ a > ();</ div >
< div class = "line" >< a name = "l00405" ></ a >< span class = "lineno" > 405</ span >   </ div >
< div class = "line" >< a name = "l00406" ></ a >< span class = "lineno" > 406</ span >   < span class = "comment" > // kind is hipMemcpyKind</ span ></ div >
< div class = "line" >< a name = "l00407" ></ a >< span class = "lineno" > 407</ span >   < span class = "keywordtype" > void</ span > copySync (< a class = "code" href = "classLockedAccessor.html" > LockedAccessor_StreamCrit_t</ a > & crit, < span class = "keywordtype" > void</ span > * dst, < span class = "keyword" > const</ span > < span class = "keywordtype" > void</ span > * src, < span class = "keywordtype" > size_t</ span > sizeBytes, < span class = "keywordtype" > unsigned</ span > kind);</ div >
< div class = "line" >< a name = "l00408" ></ a >< span class = "lineno" > 408</ span >   < span class = "keywordtype" > void</ span > locked_copySync (< span class = "keywordtype" > void</ span > * dst, < span class = "keyword" > const</ span > < span class = "keywordtype" > void</ span > * src, < span class = "keywordtype" > size_t</ span > sizeBytes, < span class = "keywordtype" > unsigned</ span > kind);</ div >
< div class = "line" >< a name = "l00409" ></ a >< span class = "lineno" > 409</ span >   </ div >
< div class = "line" >< a name = "l00410" ></ a >< span class = "lineno" > 410</ span >   < span class = "keywordtype" > void</ span > copyAsync(< span class = "keywordtype" > void</ span > * dst, < span class = "keyword" > const</ span > < span class = "keywordtype" > void</ span > * src, < span class = "keywordtype" > size_t</ span > sizeBytes, < span class = "keywordtype" > unsigned</ span > kind);</ div >
< div class = "line" >< a name = "l00411" ></ a >< span class = "lineno" > 411</ span >   </ div >
< div class = "line" >< a name = "l00412" ></ a >< span class = "lineno" > 412</ span >   < span class = "comment" > //---</ span ></ div >
< div class = "line" >< a name = "l00413" ></ a >< span class = "lineno" > 413</ span >   < span class = "comment" > // Thread-safe accessors - these acquire / release mutex:</ span ></ div >
< div class = "line" >< a name = "l00414" ></ a >< span class = "lineno" > 414</ span >   < span class = "keywordtype" > bool</ span > lockopen_preKernelCommand();</ div >
< div class = "line" >< a name = "l00415" ></ a >< span class = "lineno" > 415</ span >   < span class = "keywordtype" > void</ span > lockclose_postKernelCommand(hc::completion_future & kernel_future);</ div >
< div class = "line" >< a name = "l00416" ></ a >< span class = "lineno" > 416</ span >   </ div >
< div class = "line" >< a name = "l00417" ></ a >< span class = "lineno" > 417</ span >   < span class = "keywordtype" > int</ span > preCopyCommand(< a class = "code" href = "classLockedAccessor.html" > LockedAccessor_StreamCrit_t</ a > & crit, < a class = "code" href = "structihipSignal__t.html" > ihipSignal_t</ a > *lastCopy, hsa_signal_t *waitSignal, ihipCommand_t copyType);</ div >
< div class = "line" >< a name = "l00418" ></ a >< span class = "lineno" > 418</ span >   </ div >
< div class = "line" >< a name = "l00419" ></ a >< span class = "lineno" > 419</ span >   < span class = "keywordtype" > void</ span > locked_reclaimSignals(SIGSEQNUM sigNum);</ div >
< div class = "line" >< a name = "l00420" ></ a >< span class = "lineno" > 420</ span >   < span class = "keywordtype" > void</ span > locked_wait(< span class = "keywordtype" > bool</ span > assertQueueEmpty=< span class = "keyword" > false</ span > );</ div >
< div class = "line" >< a name = "l00421" ></ a >< span class = "lineno" > 421</ span >   SIGSEQNUM locked_lastCopySeqId() {< a class = "code" href = "classLockedAccessor.html" > LockedAccessor_StreamCrit_t</ a > crit(_criticalData); < span class = "keywordflow" > return</ span > lastCopySeqId(crit); };</ div >
< div class = "line" >< a name = "l00422" ></ a >< span class = "lineno" > 422</ span >   </ div >
< div class = "line" >< a name = "l00423" ></ a >< span class = "lineno" > 423</ span >   < span class = "comment" > // Use this if we already have the stream critical data mutex:</ span ></ div >
< div class = "line" >< a name = "l00424" ></ a >< span class = "lineno" > 424</ span >   < span class = "keywordtype" > void</ span > wait(< a class = "code" href = "classLockedAccessor.html" > LockedAccessor_StreamCrit_t</ a > & crit, < span class = "keywordtype" > bool</ span > assertQueueEmpty=< span class = "keyword" > false</ span > );</ div >
< div class = "line" >< a name = "l00425" ></ a >< span class = "lineno" > 425</ span >   </ div >
< div class = "line" >< a name = "l00426" ></ a >< span class = "lineno" > 426</ span >   </ div >
< div class = "line" >< a name = "l00427" ></ a >< span class = "lineno" > 427</ span >   </ div >
< div class = "line" >< a name = "l00428" ></ a >< span class = "lineno" > 428</ span >   < span class = "comment" > // Non-threadsafe accessors - must be protected by high-level stream lock with accessor passed to function.</ span ></ div >
< div class = "line" >< a name = "l00429" ></ a >< span class = "lineno" > 429</ span >   SIGSEQNUM lastCopySeqId (< a class = "code" href = "classLockedAccessor.html" > LockedAccessor_StreamCrit_t</ a > & crit) { < span class = "keywordflow" > return</ span > crit-> _last_copy_signal ? crit-> _last_copy_signal-> _sig_id : 0; };</ div >
< div class = "line" >< a name = "l00430" ></ a >< span class = "lineno" > 430</ span >   < a class = "code" href = "structihipSignal__t.html" > ihipSignal_t</ a > * allocSignal (< a class = "code" href = "classLockedAccessor.html" > LockedAccessor_StreamCrit_t</ a > & crit);</ div >
< div class = "line" >< a name = "l00431" ></ a >< span class = "lineno" > 431</ span >   </ div >
< div class = "line" >< a name = "l00432" ></ a >< span class = "lineno" > 432</ span >   </ div >
< div class = "line" >< a name = "l00433" ></ a >< span class = "lineno" > 433</ span >   < span class = "comment" > //-- Non-racy accessors:</ span ></ div >
< div class = "line" >< a name = "l00434" ></ a >< span class = "lineno" > 434</ span >   < span class = "comment" > // These functions access fields set at initialization time and are non-racy (so do not acquire mutex)</ span ></ div >
< div class = "line" >< a name = "l00435" ></ a >< span class = "lineno" > 435</ span >   < a class = "code" href = "classihipDevice__t.html" > ihipDevice_t</ a > * getDevice() < span class = "keyword" > const</ span > ;</ div >
< div class = "line" >< a name = "l00436" ></ a >< span class = "lineno" > 436</ span >   </ div >
< div class = "line" >< a name = "l00437" ></ a >< span class = "lineno" > 437</ span >   </ div >
< div class = "line" >< a name = "l00438" ></ a >< span class = "lineno" > 438</ span >   < span class = "keyword" > public</ span > :</ div >
< div class = "line" >< a name = "l00439" ></ a >< span class = "lineno" > 439</ span >   < span class = "comment" > //---</ span ></ div >
< div class = "line" >< a name = "l00440" ></ a >< span class = "lineno" > 440</ span >   < span class = "comment" > //Public member vars - these are set at initialization and never change:</ span ></ div >
< div class = "line" >< a name = "l00441" ></ a >< span class = "lineno" > 441</ span >   SeqNum_t _id; < span class = "comment" > // monotonic sequence ID</ span ></ div >
< div class = "line" >< a name = "l00442" ></ a >< span class = "lineno" > 442</ span >   hc::accelerator_view _av;</ div >
< div class = "line" >< a name = "l00443" ></ a >< span class = "lineno" > 443</ span >   < span class = "keywordtype" > unsigned</ span > _flags;</ div >
< div class = "line" >< a name = "l00444" ></ a >< span class = "lineno" > 444</ span >   </ div >
< div class = "line" >< a name = "l00445" ></ a >< span class = "lineno" > 445</ span >   < span class = "keyword" > private</ span > : < span class = "comment" > // Critical Data. THis MUST be accessed through LockedAccessor_StreamCrit_t</ span ></ div >
< div class = "line" >< a name = "l00446" ></ a >< span class = "lineno" > 446</ span >   < a class = "code" href = "classihipStreamCriticalBase__t.html" > ihipStreamCritical_t</ a > _criticalData;</ div >
< div class = "line" >< a name = "l00447" ></ a >< span class = "lineno" > 447</ span >   </ div >
< div class = "line" >< a name = "l00448" ></ a >< span class = "lineno" > 448</ span >   < span class = "keyword" > private</ span > :</ div >
< div class = "line" >< a name = "l00449" ></ a >< span class = "lineno" > 449</ span >   < span class = "keywordtype" > void</ span > enqueueBarrier(hsa_queue_t* queue, < a class = "code" href = "structihipSignal__t.html" > ihipSignal_t</ a > *depSignal);</ div >
< div class = "line" >< a name = "l00450" ></ a >< span class = "lineno" > 450</ span >   < span class = "keywordtype" > void</ span > waitCopy(< a class = "code" href = "classLockedAccessor.html" > LockedAccessor_StreamCrit_t</ a > & crit, < a class = "code" href = "structihipSignal__t.html" > ihipSignal_t</ a > *signal);</ div >
< div class = "line" >< a name = "l00451" ></ a >< span class = "lineno" > 451</ span >   </ div >
< div class = "line" >< a name = "l00452" ></ a >< span class = "lineno" > 452</ span >   < span class = "comment" > // The unsigned return is hipMemcpyKind</ span ></ div >
< div class = "line" >< a name = "l00453" ></ a >< span class = "lineno" > 453</ span >   < span class = "keywordtype" > unsigned</ span > resolveMemcpyDirection(< span class = "keywordtype" > bool</ span > srcInDeviceMem, < span class = "keywordtype" > bool</ span > dstInDeviceMem);</ div >
< div class = "line" >< a name = "l00454" ></ a >< span class = "lineno" > 454</ span >   < span class = "keywordtype" > void</ span > setCopyAgents(< span class = "keywordtype" > unsigned</ span > kind, ihipCommand_t *commandType, hsa_agent_t *srcAgent, hsa_agent_t *dstAgent);</ div >
< div class = "line" >< a name = "l00455" ></ a >< span class = "lineno" > 455</ span >   </ div >
< div class = "line" >< a name = "l00456" ></ a >< span class = "lineno" > 456</ span >   < span class = "keywordtype" > unsigned</ span > _device_index; < span class = "comment" > // index into the g_device array </ span ></ div >
< div class = "line" >< a name = "l00457" ></ a >< span class = "lineno" > 457</ span >   </ div >
< div class = "line" >< a name = "l00458" ></ a >< span class = "lineno" > 458</ span >   < span class = "keyword" > friend</ span > std::ostream& operator<< (std::ostream& os, < span class = "keyword" > const</ span > < a class = "code" href = "classihipStream__t.html" > ihipStream_t</ a > & s);</ div >
< div class = "line" >< a name = "l00459" ></ a >< span class = "lineno" > 459</ span >   };</ div >
< div class = "line" >< a name = "l00460" ></ a >< span class = "lineno" > 460</ span >   </ div >
< div class = "line" >< a name = "l00461" ></ a >< span class = "lineno" > 461</ span >   </ div >
< div class = "line" >< a name = "l00462" ></ a >< span class = "lineno" > 462</ span >   < span class = "keyword" > inline</ span > std::ostream& operator<< (std::ostream& os, < span class = "keyword" > const</ span > < a class = "code" href = "classihipStream__t.html" > ihipStream_t</ a > & s)</ div >
< div class = "line" >< a name = "l00463" ></ a >< span class = "lineno" > 463</ span >   {</ div >
< div class = "line" >< a name = "l00464" ></ a >< span class = "lineno" > 464</ span >   os << < span class = "stringliteral" > " stream#" </ span > ;</ div >
< div class = "line" >< a name = "l00465" ></ a >< span class = "lineno" > 465</ span >   os << s._device_index;</ div >
< div class = "line" >< a name = "l00466" ></ a >< span class = "lineno" > 466</ span >   os << < span class = "charliteral" > ' .' </ span > ;</ div >
< div class = "line" >< a name = "l00467" ></ a >< span class = "lineno" > 467</ span >   os << s._id;</ div >
< div class = "line" >< a name = "l00468" ></ a >< span class = "lineno" > 468</ span >   < span class = "keywordflow" > return</ span > os;</ div >
< div class = "line" >< a name = "l00469" ></ a >< span class = "lineno" > 469</ span >   }</ div >
< div class = "line" >< a name = "l00470" ></ a >< span class = "lineno" > 470</ span >   </ div >
< div class = "line" >< a name = "l00471" ></ a >< span class = "lineno" > 471</ span >   </ div >
< div class = "line" >< a name = "l00472" ></ a >< span class = "lineno" > 472</ span >   < span class = "comment" > //----</ span ></ div >
< div class = "line" >< a name = "l00473" ></ a >< span class = "lineno" > 473</ span >   < span class = "comment" > // Internal event structure:</ span ></ div >
< div class = "line" >< a name = "l00474" ></ a >< span class = "lineno" > 474</ span >   < span class = "keyword" > enum</ span > hipEventStatus_t {</ div >
< div class = "line" >< a name = "l00475" ></ a >< span class = "lineno" > 475</ span >   hipEventStatusUnitialized = 0, < span class = "comment" > // event is unutilized, must be " Created" before use.</ span ></ div >
< div class = "line" >< a name = "l00476" ></ a >< span class = "lineno" > 476</ span >   hipEventStatusCreated = 1,</ div >
< div class = "line" >< a name = "l00477" ></ a >< span class = "lineno" > 477</ span >   hipEventStatusRecording = 2, < span class = "comment" > // event has been enqueued to record something.</ span ></ div >
< div class = "line" >< a name = "l00478" ></ a >< span class = "lineno" > 478</ span >   hipEventStatusRecorded = 3, < span class = "comment" > // event has been recorded - timestamps are valid.</ span ></ div >
< div class = "line" >< a name = "l00479" ></ a >< span class = "lineno" > 479</ span >   } ;</ div >
< div class = "line" >< a name = "l00480" ></ a >< span class = "lineno" > 480</ span >   </ div >
< div class = "line" >< a name = "l00481" ></ a >< span class = "lineno" > 481</ span >   </ div >
< div class = "line" >< a name = "l00482" ></ a >< span class = "lineno" > 482</ span >   < span class = "comment" > // internal hip event structure.</ span ></ div >
< div class = "line" >< a name = "l00483" ></ a >< span class = "lineno" >< a class = "line" href = "structihipEvent__t.html" > 483</ a ></ span >   < span class = "keyword" > struct </ span >< a class = "code" href = "structihipEvent__t.html" > ihipEvent_t</ a > {</ div >
< div class = "line" >< a name = "l00484" ></ a >< span class = "lineno" > 484</ span >   hipEventStatus_t _state;</ div >
< div class = "line" >< a name = "l00485" ></ a >< span class = "lineno" > 485</ span >   </ div >
< div class = "line" >< a name = "l00486" ></ a >< span class = "lineno" > 486</ span >   hipStream_t _stream; < span class = "comment" > // Stream where the event is recorded, or NULL if all streams.</ span ></ div >
< div class = "line" >< a name = "l00487" ></ a >< span class = "lineno" > 487</ span >   < span class = "keywordtype" > unsigned</ span > _flags;</ div >
< div class = "line" >< a name = "l00488" ></ a >< span class = "lineno" > 488</ span >   </ div >
< div class = "line" >< a name = "l00489" ></ a >< span class = "lineno" > 489</ span >   hc::completion_future _marker;</ div >
< div class = "line" >< a name = "l00490" ></ a >< span class = "lineno" > 490</ span >   uint64_t _timestamp; < span class = "comment" > // store timestamp, may be set on host or by marker.</ span ></ div >
< div class = "line" >< a name = "l00491" ></ a >< span class = "lineno" > 491</ span >   </ div >
< div class = "line" >< a name = "l00492" ></ a >< span class = "lineno" > 492</ span >   SIGSEQNUM _copy_seq_id;</ div >
< div class = "line" >< a name = "l00493" ></ a >< span class = "lineno" > 493</ span >   } ;</ div >
< div class = "line" >< a name = "l00494" ></ a >< span class = "lineno" > 494</ span >   </ div >
< div class = "line" >< a name = "l00495" ></ a >< span class = "lineno" > 495</ span >   </ div >
< div class = "line" >< a name = "l00496" ></ a >< span class = "lineno" > 496</ span >   </ div >
< div class = "line" >< a name = "l00497" ></ a >< span class = "lineno" > 497</ span >   </ div >
< div class = "line" >< a name = "l00498" ></ a >< span class = "lineno" > 498</ span >   </ div >
< div class = "line" >< a name = "l00499" ></ a >< span class = "lineno" > 499</ span >   < span class = "comment" > //---</ span ></ div >
< div class = "line" >< a name = "l00500" ></ a >< span class = "lineno" > 500</ span >   < span class = "comment" > // Data that must be protected with thread-safe access</ span ></ div >
< div class = "line" >< a name = "l00501" ></ a >< span class = "lineno" > 501</ span >   < span class = "comment" > // All members are private - this class must be accessed through friend LockedAccessor which </ span ></ div >
< div class = "line" >< a name = "l00502" ></ a >< span class = "lineno" > 502</ span >   < span class = "comment" > // will lock the mutex on construction and unlock on destruction.</ span ></ div >
< div class = "line" >< a name = "l00503" ></ a >< span class = "lineno" > 503</ span >   < span class = "comment" > //</ span ></ div >
< div class = "line" >< a name = "l00504" ></ a >< span class = "lineno" > 504</ span >   < span class = "comment" > // MUTEX_TYPE is template argument so can easily convert to FakeMutex for performance or stress testing.</ span ></ div >
< div class = "line" >< a name = "l00505" ></ a >< span class = "lineno" > 505</ span >   < span class = "keyword" > template</ span > < < span class = "keyword" > class</ span > MUTEX_TYPE> </ div >
< div class = "line" >< a name = "l00506" ></ a >< span class = "lineno" >< a class = "line" href = "classihipDeviceCriticalBase__t.html" > 506</ a ></ span >   < span class = "keyword" > class </ span >< a class = "code" href = "classihipDeviceCriticalBase__t.html" > ihipDeviceCriticalBase_t</ a > : < a class = "code" href = "structLockedBase.html" > LockedBase</ a > < MUTEX_TYPE> </ div >
< div class = "line" >< a name = "l00507" ></ a >< span class = "lineno" > 507</ span >   {</ div >
< div class = "line" >< a name = "l00508" ></ a >< span class = "lineno" > 508</ span >   < span class = "keyword" > public</ span > :</ div >
< div class = "line" >< a name = "l00509" ></ a >< span class = "lineno" > 509</ span >   < a class = "code" href = "classihipDeviceCriticalBase__t.html" > ihipDeviceCriticalBase_t</ a > () : _stream_id(0), _peerAgents(< span class = "keyword" > nullptr</ span > ) {};</ div >
< div class = "line" >< a name = "l00510" ></ a >< span class = "lineno" > 510</ span >   </ div >
< div class = "line" >< a name = "l00511" ></ a >< span class = "lineno" > 511</ span >   < span class = "keywordtype" > void</ span > init(< span class = "keywordtype" > unsigned</ span > deviceCnt) {</ div >
< div class = "line" >< a name = "l00512" ></ a >< span class = "lineno" > 512</ span >   assert(_peerAgents == < span class = "keyword" > nullptr</ span > );</ div >
< div class = "line" >< a name = "l00513" ></ a >< span class = "lineno" > 513</ span >   _peerAgents = < span class = "keyword" > new</ span > hsa_agent_t[deviceCnt];</ div >
< div class = "line" >< a name = "l00514" ></ a >< span class = "lineno" > 514</ span >   };</ div >
< div class = "line" >< a name = "l00515" ></ a >< span class = "lineno" > 515</ span >   </ div >
< div class = "line" >< a name = "l00516" ></ a >< span class = "lineno" > 516</ span >   ~< a class = "code" href = "classihipDeviceCriticalBase__t.html" > ihipDeviceCriticalBase_t</ a > () {</ div >
< div class = "line" >< a name = "l00517" ></ a >< span class = "lineno" > 517</ span >   < span class = "keywordflow" > if</ span > (_peerAgents != < span class = "keyword" > nullptr</ span > ) {</ div >
< div class = "line" >< a name = "l00518" ></ a >< span class = "lineno" > 518</ span >   < span class = "keyword" > delete</ span > _peerAgents;</ div >
< div class = "line" >< a name = "l00519" ></ a >< span class = "lineno" > 519</ span >   _peerAgents = < span class = "keyword" > nullptr</ span > ;</ div >
< div class = "line" >< a name = "l00520" ></ a >< span class = "lineno" > 520</ span >   }</ div >
< div class = "line" >< a name = "l00521" ></ a >< span class = "lineno" > 521</ span >   }</ div >
< div class = "line" >< a name = "l00522" ></ a >< span class = "lineno" > 522</ span >   < span class = "keyword" > friend</ span > < span class = "keyword" > class </ span >< a class = "code" href = "classLockedAccessor.html" > LockedAccessor</ a > < ihipDeviceCriticalBase_t> ;</ div >
< div class = "line" >< a name = "l00523" ></ a >< span class = "lineno" > 523</ span >   </ div >
< div class = "line" >< a name = "l00524" ></ a >< span class = "lineno" > 524</ span >   std::list< ihipStream_t*> & streams() { < span class = "keywordflow" > return</ span > _streams; };</ div >
< div class = "line" >< a name = "l00525" ></ a >< span class = "lineno" > 525</ span >   < span class = "keyword" > const</ span > std::list< ihipStream_t*> & const_streams()< span class = "keyword" > const </ span > { < span class = "keywordflow" > return</ span > _streams; };</ div >
< div class = "line" >< a name = "l00526" ></ a >< span class = "lineno" > 526</ span >   </ div >
< div class = "line" >< a name = "l00527" ></ a >< span class = "lineno" > 527</ span >   < span class = "comment" > // " Allocate" a stream ID:</ span ></ div >
< div class = "line" >< a name = "l00528" ></ a >< span class = "lineno" > 528</ span >   ihipStream_t::SeqNum_t incStreamId() { < span class = "keywordflow" > return</ span > _stream_id++; };</ div >
< div class = "line" >< a name = "l00529" ></ a >< span class = "lineno" > 529</ span >   </ div >
< div class = "line" >< a name = "l00530" ></ a >< span class = "lineno" > 530</ span >   < span class = "keywordtype" > bool</ span > addPeer(< a class = "code" href = "classihipDevice__t.html" > ihipDevice_t</ a > *peer);</ div >
< div class = "line" >< a name = "l00531" ></ a >< span class = "lineno" > 531</ span >   < 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 >   < 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 >   </ div >
2016-04-19 22:44:58 +05:30
< div class = "line" >< a name = "l00534" ></ a >< span class = "lineno" > 534</ span >   </ div >
< div class = "line" >< a name = "l00535" ></ a >< span class = "lineno" > 535</ span >   < span class = "keywordtype" > void</ span > addStream(< a class = "code" href = "classihipStream__t.html" > ihipStream_t</ a > *stream);</ div >
2016-04-16 14:55:10 +05:30
< div class = "line" >< a name = "l00536" ></ a >< span class = "lineno" > 536</ span >   </ div >
2016-04-19 22:44:58 +05:30
< div class = "line" >< a name = "l00537" ></ a >< span class = "lineno" > 537</ span >   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 >   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 >   </ div >
< div class = "line" >< a name = "l00540" ></ a >< span class = "lineno" > 540</ span >   </ div >
< div class = "line" >< a name = "l00541" ></ a >< span class = "lineno" > 541</ span >   < span class = "keyword" > private</ span > :</ div >
< div class = "line" >< a name = "l00542" ></ a >< span class = "lineno" > 542</ span >   std::list< ihipStream_t*> _streams; < span class = "comment" > // streams associated with this device.</ span ></ div >
< div class = "line" >< a name = "l00543" ></ a >< span class = "lineno" > 543</ span >   ihipStream_t::SeqNum_t _stream_id;</ div >
< div class = "line" >< a name = "l00544" ></ a >< span class = "lineno" > 544</ span >   </ div >
< div class = "line" >< a name = "l00545" ></ a >< span class = "lineno" > 545</ span >   < 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 >   std::list< ihipDevice_t*> _peers; < span class = "comment" > // list of enabled peer devices.</ span ></ div >
< div class = "line" >< a name = "l00547" ></ a >< span class = "lineno" > 547</ span >   uint32_t _peerCnt; < span class = "comment" > // number of enabled peers</ span ></ div >
< div class = "line" >< a name = "l00548" ></ a >< span class = "lineno" > 548</ span >   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 >   < span class = "keyword" > private</ span > :</ div >
< div class = "line" >< a name = "l00550" ></ a >< span class = "lineno" > 550</ span >   < span class = "keywordtype" > void</ span > recomputePeerAgents();</ div >
< div class = "line" >< a name = "l00551" ></ a >< span class = "lineno" > 551</ span >   };</ div >
2016-04-16 14:55:10 +05:30
< div class = "line" >< a name = "l00552" ></ a >< span class = "lineno" > 552</ span >   </ div >
2016-04-19 22:44:58 +05:30
< div class = "line" >< a name = "l00553" ></ a >< span class = "lineno" > 553</ span >   < span class = "comment" > // Note Mutex selected based on DeviceMutex</ span ></ div >
< div class = "line" >< a name = "l00554" ></ a >< span class = "lineno" > 554</ span >   < span class = "keyword" > typedef</ span > < a class = "code" href = "classihipDeviceCriticalBase__t.html" > ihipDeviceCriticalBase_t< DeviceMutex> </ a > < a class = "code" href = "classihipDeviceCriticalBase__t.html" > ihipDeviceCritical_t</ a > ; </ div >
2016-04-16 14:55:10 +05:30
< div class = "line" >< a name = "l00555" ></ a >< span class = "lineno" > 555</ span >   </ div >
2016-04-19 22:44:58 +05:30
< div class = "line" >< a name = "l00556" ></ a >< span class = "lineno" > 556</ span >   < 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 >   < span class = "keyword" > typedef</ span > < a class = "code" href = "classLockedAccessor.html" > LockedAccessor< ihipDeviceCritical_t> </ a > < a class = "code" href = "classLockedAccessor.html" > LockedAccessor_DeviceCrit_t</ a > ;</ div >
< div class = "line" >< a name = "l00558" ></ a >< span class = "lineno" > 558</ span >   </ div >
< div class = "line" >< a name = "l00559" ></ a >< span class = "lineno" > 559</ span >   </ div >
< div class = "line" >< a name = "l00560" ></ a >< span class = "lineno" > 560</ span >   </ div >
< div class = "line" >< a name = "l00561" ></ a >< span class = "lineno" > 561</ span >   < span class = "comment" > //-------------------------------------------------------------------------------------------------</ span ></ div >
< div class = "line" >< a name = "l00562" ></ a >< span class = "lineno" > 562</ span >   < 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 >   < 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 >   < 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 >   < 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 >   < 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 >   {</ div >
< div class = "line" >< a name = "l00568" ></ a >< span class = "lineno" > 568</ span >   < span class = "keyword" > public</ span > : < span class = "comment" > // Functions:</ span ></ div >
< div class = "line" >< a name = "l00569" ></ a >< span class = "lineno" > 569</ span >   < 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 >   < span class = "keywordtype" > void</ span > init(< span class = "keywordtype" > unsigned</ span > device_index, < span class = "keywordtype" > unsigned</ span > deviceCnt, hc::accelerator & acc, < span class = "keywordtype" > unsigned</ span > flags);</ div >
< div class = "line" >< a name = "l00571" ></ a >< span class = "lineno" > 571</ span >   ~< a class = "code" href = "classihipDevice__t.html" > ihipDevice_t</ a > ();</ div >
< div class = "line" >< a name = "l00572" ></ a >< span class = "lineno" > 572</ span >   </ div >
< div class = "line" >< a name = "l00573" ></ a >< span class = "lineno" > 573</ span >   < 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 >   < 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 >   < span class = "keywordtype" > void</ span > locked_reset();</ div >
< div class = "line" >< a name = "l00576" ></ a >< span class = "lineno" > 576</ span >   < span class = "keywordtype" > void</ span > locked_waitAllStreams();</ div >
< div class = "line" >< a name = "l00577" ></ a >< span class = "lineno" > 577</ span >   < 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 >   </ div >
< div class = "line" >< a name = "l00579" ></ a >< span class = "lineno" > 579</ span >   < a class = "code" href = "classihipDeviceCriticalBase__t.html" > ihipDeviceCritical_t</ a > & criticalData() { < span class = "keywordflow" > return</ span > _criticalData; }; < span class = "comment" > // TODO, move private. Fix P2P.</ span ></ div >
2016-04-16 14:55:10 +05:30
< div class = "line" >< a name = "l00580" ></ a >< span class = "lineno" > 580</ span >   </ div >
2016-04-19 22:44:58 +05:30
< div class = "line" >< a name = "l00581" ></ a >< span class = "lineno" > 581</ span >   < 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 >   < 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 >   </ div >
< div class = "line" >< a name = "l00584" ></ a >< span class = "lineno" > 584</ span >   < 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 >   hc::accelerator _acc;</ div >
< div class = "line" >< a name = "l00586" ></ a >< span class = "lineno" > 586</ span >   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 >   </ div >
< div class = "line" >< a name = "l00588" ></ a >< span class = "lineno" > 588</ span >   < 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 >   < 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 >   < a class = "code" href = "classihipStream__t.html" > ihipStream_t</ a > *_default_stream;</ div >
2016-04-16 14:55:10 +05:30
< div class = "line" >< a name = "l00591" ></ a >< span class = "lineno" > 591</ span >   </ div >
2016-04-19 22:44:58 +05:30
< div class = "line" >< a name = "l00592" ></ a >< span class = "lineno" > 592</ span >   </ div >
< div class = "line" >< a name = "l00593" ></ a >< span class = "lineno" > 593</ span >   < span class = "keywordtype" > unsigned</ span > _compute_units;</ div >
2016-04-16 14:55:10 +05:30
< div class = "line" >< a name = "l00594" ></ a >< span class = "lineno" > 594</ span >   </ div >
2016-04-19 22:44:58 +05:30
< div class = "line" >< a name = "l00595" ></ a >< span class = "lineno" > 595</ span >   < a class = "code" href = "structStagingBuffer.html" > StagingBuffer</ a > *_staging_buffer[2]; < span class = "comment" > // one buffer for each direction.</ span ></ div >
2016-04-16 14:55:10 +05:30
< div class = "line" >< a name = "l00596" ></ a >< span class = "lineno" > 596</ span >   </ div >
2016-04-19 22:44:58 +05:30
< div class = "line" >< a name = "l00597" ></ a >< span class = "lineno" > 597</ span >   </ div >
< div class = "line" >< a name = "l00598" ></ a >< span class = "lineno" > 598</ span >   < span class = "keywordtype" > unsigned</ span > _device_flags;</ div >
2016-04-16 14:55:10 +05:30
< div class = "line" >< a name = "l00599" ></ a >< span class = "lineno" > 599</ span >   </ div >
2016-04-19 22:44:58 +05:30
< div class = "line" >< a name = "l00600" ></ a >< span class = "lineno" > 600</ span >   < span class = "keyword" > private</ span > :</ div >
< div class = "line" >< a name = "l00601" ></ a >< span class = "lineno" > 601</ span >   < 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 >   </ div >
< div class = "line" >< a name = "l00603" ></ a >< span class = "lineno" > 603</ span >   < 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 >   < 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 >   < span class = "comment" > // Search for LockedAccessor< ihipDeviceCritical_t> for examples; do not access _criticalData directly.</ span ></ div >
< div class = "line" >< a name = "l00606" ></ a >< span class = "lineno" > 606</ span >   < a class = "code" href = "classihipDeviceCriticalBase__t.html" > ihipDeviceCritical_t</ a > _criticalData;</ div >
2016-04-16 14:55:10 +05:30
< div class = "line" >< a name = "l00607" ></ a >< span class = "lineno" > 607</ span >   </ div >
2016-04-19 22:44:58 +05:30
< div class = "line" >< a name = "l00608" ></ a >< span class = "lineno" > 608</ span >   };</ div >
< div class = "line" >< a name = "l00609" ></ a >< span class = "lineno" > 609</ span >   </ div >
< div class = "line" >< a name = "l00610" ></ a >< span class = "lineno" > 610</ span >   </ div >
< div class = "line" >< a name = "l00611" ></ a >< span class = "lineno" > 611</ span >   </ div >
< div class = "line" >< a name = "l00612" ></ a >< span class = "lineno" > 612</ span >   < span class = "comment" > // Global variable definition:</ span ></ div >
< div class = "line" >< a name = "l00613" ></ a >< span class = "lineno" > 613</ span >   < span class = "keyword" > extern</ span > std::once_flag hip_initialized;</ div >
< div class = "line" >< a name = "l00614" ></ a >< span class = "lineno" > 614</ span >   < 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 >   < 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 >   < 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 >   < span class = "keyword" > extern</ span > std::vector< int> 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 >   < 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 >   < span class = "comment" > //=================================================================================================</ span ></ div >
< div class = "line" >< a name = "l00620" ></ a >< span class = "lineno" > 620</ span >   < span class = "keywordtype" > void</ span > ihipInit();</ div >
< div class = "line" >< a name = "l00621" ></ a >< span class = "lineno" > 621</ span >   < 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 >   < a class = "code" href = "classihipDevice__t.html" > ihipDevice_t</ a > *ihipGetTlsDefaultDevice();</ div >
< div class = "line" >< a name = "l00623" ></ a >< span class = "lineno" > 623</ span >   < 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 >   < span class = "keywordtype" > void</ span > ihipSetTs(< a class = "code" href = "structhipEvent__t.html" > hipEvent_t</ a > e);</ div >
2016-04-16 14:55:10 +05:30
< div class = "line" >< a name = "l00625" ></ a >< span class = "lineno" > 625</ span >   </ div >
< div class = "line" >< a name = "l00626" ></ a >< span class = "lineno" > 626</ span >   < span class = "keyword" > template</ span > < < span class = "keyword" > typename</ span > T> </ div >
2016-04-19 22:44:58 +05:30
< div class = "line" >< a name = "l00627" ></ a >< span class = "lineno" > 627</ span >   hc::completion_future ihipMemcpyKernel(hipStream_t, T*, < span class = "keyword" > const</ span > T*, < span class = "keywordtype" > size_t</ span > );</ div >
2016-04-16 14:55:10 +05:30
< div class = "line" >< a name = "l00628" ></ a >< span class = "lineno" > 628</ span >   </ div >
2016-04-19 22:44:58 +05:30
< div class = "line" >< a name = "l00629" ></ a >< span class = "lineno" > 629</ span >   < span class = "keyword" > template</ span > < < span class = "keyword" > typename</ span > T> </ div >
< div class = "line" >< a name = "l00630" ></ a >< span class = "lineno" > 630</ span >   hc::completion_future ihipMemsetKernel(hipStream_t, T*, T, < span class = "keywordtype" > size_t</ span > );</ div >
2016-04-16 14:55:10 +05:30
< div class = "line" >< a name = "l00631" ></ a >< span class = "lineno" > 631</ span >   </ div >
2016-04-19 22:44:58 +05:30
< div class = "line" >< a name = "l00632" ></ a >< span class = "lineno" > 632</ span >   hipStream_t ihipSyncAndResolveStream(hipStream_t);</ div >
< div class = "line" >< a name = "l00633" ></ a >< span class = "lineno" > 633</ span >   < span class = "keyword" > template</ span > < < span class = "keyword" > typename</ span > T> </ div >
< div class = "line" >< a name = "l00634" ></ a >< span class = "lineno" > 634</ span >   </ div >
< div class = "line" >< a name = "l00635" ></ a >< span class = "lineno" > 635</ span >   hc::completion_future</ div >
< div class = "line" >< a name = "l00636" ></ a >< span class = "lineno" > 636</ span >   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 >   {</ div >
< div class = "line" >< a name = "l00638" ></ a >< span class = "lineno" > 638</ span >   < span class = "keywordtype" > int</ span > wg = std::min((< span class = "keywordtype" > unsigned</ span > )8, stream-> getDevice()-> _compute_units);</ div >
< div class = "line" >< a name = "l00639" ></ a >< span class = "lineno" > 639</ span >   < 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 >   </ div >
< div class = "line" >< a name = "l00641" ></ a >< span class = "lineno" > 641</ span >   < span class = "keywordtype" > int</ span > threads = wg * threads_per_wg;</ div >
< div class = "line" >< a name = "l00642" ></ a >< span class = "lineno" > 642</ span >   < span class = "keywordflow" > if</ span > (threads > sizeBytes) {</ div >
< div class = "line" >< a name = "l00643" ></ a >< span class = "lineno" > 643</ span >   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 >   }</ div >
< div class = "line" >< a name = "l00645" ></ a >< span class = "lineno" > 645</ span >   </ div >
2016-04-16 14:55:10 +05:30
< div class = "line" >< a name = "l00646" ></ a >< span class = "lineno" > 646</ span >   </ div >
2016-04-19 22:44:58 +05:30
< div class = "line" >< a name = "l00647" ></ a >< span class = "lineno" > 647</ span >   hc::extent< 1> ext(threads);</ div >
< div class = "line" >< a name = "l00648" ></ a >< span class = "lineno" > 648</ span >   < 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 >   </ div >
< div class = "line" >< a name = "l00650" ></ a >< span class = "lineno" > 650</ span >   hc::completion_future cf =</ div >
< div class = "line" >< a name = "l00651" ></ a >< span class = "lineno" > 651</ span >   hc::parallel_for_each(</ div >
< div class = "line" >< a name = "l00652" ></ a >< span class = "lineno" > 652</ span >   stream-> _av,</ div >
< div class = "line" >< a name = "l00653" ></ a >< span class = "lineno" > 653</ span >   ext_tile,</ div >
< div class = "line" >< a name = "l00654" ></ a >< span class = "lineno" > 654</ span >   [=] (hc::tiled_index< 1> idx)</ div >
< div class = "line" >< a name = "l00655" ></ a >< span class = "lineno" > 655</ span >   __attribute__((hc))</ div >
< div class = "line" >< a name = "l00656" ></ a >< span class = "lineno" > 656</ span >   {</ div >
< div class = "line" >< a name = "l00657" ></ a >< span class = "lineno" > 657</ span >   < span class = "keywordtype" > int</ span > offset = amp_get_global_id(0);</ div >
< div class = "line" >< a name = "l00658" ></ a >< span class = "lineno" > 658</ span >   < 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 >   < 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 >   </ div >
< div class = "line" >< a name = "l00661" ></ a >< span class = "lineno" > 661</ span >   < span class = "keywordflow" > for</ span > (< span class = "keywordtype" > int</ span > i=offset; i< sizeBytes; i+=stride) {</ div >
< div class = "line" >< a name = "l00662" ></ a >< span class = "lineno" > 662</ span >   ptr[i] = val;</ div >
< div class = "line" >< a name = "l00663" ></ a >< span class = "lineno" > 663</ span >   }</ div >
< div class = "line" >< a name = "l00664" ></ a >< span class = "lineno" > 664</ span >   });</ div >
2016-04-16 14:55:10 +05:30
< div class = "line" >< a name = "l00665" ></ a >< span class = "lineno" > 665</ span >   </ div >
2016-04-19 22:44:58 +05:30
< div class = "line" >< a name = "l00666" ></ a >< span class = "lineno" > 666</ span >   < span class = "keywordflow" > return</ span > cf;</ div >
< div class = "line" >< a name = "l00667" ></ a >< span class = "lineno" > 667</ span >   }</ div >
< div class = "line" >< a name = "l00668" ></ a >< span class = "lineno" > 668</ span >   </ div >
< div class = "line" >< a name = "l00669" ></ a >< span class = "lineno" > 669</ span >   < span class = "keyword" > template</ span > < < span class = "keyword" > typename</ span > T> </ div >
< div class = "line" >< a name = "l00670" ></ a >< span class = "lineno" > 670</ span >   hc::completion_future</ div >
< div class = "line" >< a name = "l00671" ></ a >< span class = "lineno" > 671</ span >   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 >   {</ div >
< div class = "line" >< a name = "l00673" ></ a >< span class = "lineno" > 673</ span >   < span class = "keywordtype" > int</ span > wg = std::min((< span class = "keywordtype" > unsigned</ span > )8, stream-> getDevice()-> _compute_units);</ div >
< div class = "line" >< a name = "l00674" ></ a >< span class = "lineno" > 674</ span >   < 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 >   </ div >
< div class = "line" >< a name = "l00676" ></ a >< span class = "lineno" > 676</ span >   < span class = "keywordtype" > int</ span > threads = wg * threads_per_wg;</ div >
< div class = "line" >< a name = "l00677" ></ a >< span class = "lineno" > 677</ span >   < span class = "keywordflow" > if</ span > (threads > sizeBytes) {</ div >
< div class = "line" >< a name = "l00678" ></ a >< span class = "lineno" > 678</ span >   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 >   }</ div >
< div class = "line" >< a name = "l00680" ></ a >< span class = "lineno" > 680</ span >   </ div >
2016-04-16 14:55:10 +05:30
< div class = "line" >< a name = "l00681" ></ a >< span class = "lineno" > 681</ span >   </ div >
2016-04-19 22:44:58 +05:30
< div class = "line" >< a name = "l00682" ></ a >< span class = "lineno" > 682</ span >   hc::extent< 1> ext(threads);</ div >
< div class = "line" >< a name = "l00683" ></ a >< span class = "lineno" > 683</ span >   < 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 >   </ div >
< div class = "line" >< a name = "l00685" ></ a >< span class = "lineno" > 685</ span >   hc::completion_future cf =</ div >
< div class = "line" >< a name = "l00686" ></ a >< span class = "lineno" > 686</ span >   hc::parallel_for_each(</ div >
< div class = "line" >< a name = "l00687" ></ a >< span class = "lineno" > 687</ span >   stream-> _av,</ div >
< div class = "line" >< a name = "l00688" ></ a >< span class = "lineno" > 688</ span >   ext_tile,</ div >
< div class = "line" >< a name = "l00689" ></ a >< span class = "lineno" > 689</ span >   [=] (hc::tiled_index< 1> idx)</ div >
< div class = "line" >< a name = "l00690" ></ a >< span class = "lineno" > 690</ span >   __attribute__((hc))</ div >
< div class = "line" >< a name = "l00691" ></ a >< span class = "lineno" > 691</ span >   {</ div >
< div class = "line" >< a name = "l00692" ></ a >< span class = "lineno" > 692</ span >   < span class = "keywordtype" > int</ span > offset = amp_get_global_id(0);</ div >
< div class = "line" >< a name = "l00693" ></ a >< span class = "lineno" > 693</ span >   < 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 >   < 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 >   </ div >
< div class = "line" >< a name = "l00696" ></ a >< span class = "lineno" > 696</ span >   < span class = "keywordflow" > for</ span > (< span class = "keywordtype" > int</ span > i=offset; i< sizeBytes; i+=stride) {</ div >
< div class = "line" >< a name = "l00697" ></ a >< span class = "lineno" > 697</ span >   c[i] = a[i];</ div >
< div class = "line" >< a name = "l00698" ></ a >< span class = "lineno" > 698</ span >   }</ div >
< div class = "line" >< a name = "l00699" ></ a >< span class = "lineno" > 699</ span >   });</ div >
2016-04-16 14:55:10 +05:30
< div class = "line" >< a name = "l00700" ></ a >< span class = "lineno" > 700</ span >   </ div >
2016-04-19 22:44:58 +05:30
< div class = "line" >< a name = "l00701" ></ a >< span class = "lineno" > 701</ span >   < span class = "keywordflow" > return</ span > cf;</ div >
< div class = "line" >< a name = "l00702" ></ a >< span class = "lineno" > 702</ span >   }</ div >
< div class = "line" >< a name = "l00703" ></ a >< span class = "lineno" > 703</ span >   </ div >
< div class = "line" >< a name = "l00704" ></ a >< span class = "lineno" > 704</ span >   < 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 >
2016-04-16 14:55:10 +05:30
< 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 >
< div class = "ttc" id = "structhipEvent__t_html" >< div class = "ttname" >< a href = "structhipEvent__t.html" > hipEvent_t</ a ></ div >< div class = "ttdef" >< b > Definition:</ b > hip_runtime_api.h:47</ div ></ div >
< div class = "ttc" id = "classihipDeviceCriticalBase__t_html" >< div class = "ttname" >< a href = "classihipDeviceCriticalBase__t.html" > ihipDeviceCriticalBase_t</ a ></ div >< div class = "ttdef" >< b > Definition:</ b > hip_hcc.h:506</ div ></ div >
< div class = "ttc" id = "structihipSignal__t_html" >< div class = "ttname" >< a href = "structihipSignal__t.html" > ihipSignal_t</ a ></ div >< div class = "ttdef" >< b > Definition:</ b > hip_hcc.h:266</ div ></ div >
< div class = "ttc" id = "structhipDeviceProp__t_html" >< div class = "ttname" >< a href = "structhipDeviceProp__t.html" > hipDeviceProp_t</ a ></ div >< div class = "ttdef" >< b > Definition:</ b > hip_runtime_api.h:74</ div ></ div >
< div class = "ttc" id = "structStagingBuffer_html" >< div class = "ttname" >< a href = "structStagingBuffer.html" > StagingBuffer</ a ></ div >< div class = "ttdef" >< b > Definition:</ b > staging_buffer.h:40</ div ></ div >
< div class = "ttc" id = "structihipEvent__t_html" >< div class = "ttname" >< a href = "structihipEvent__t.html" > ihipEvent_t</ a ></ div >< div class = "ttdef" >< b > Definition:</ b > hip_hcc.h:483</ div ></ div >
< div class = "ttc" id = "classihipException_html" >< div class = "ttname" >< a href = "classihipException.html" > ihipException</ a ></ div >< div class = "ttdef" >< b > Definition:</ b > hip_hcc.h:220</ div ></ div >
< div class = "ttc" id = "classihipStream__t_html" >< div class = "ttname" >< a href = "classihipStream__t.html" > ihipStream_t</ a ></ div >< div class = "ttdef" >< b > Definition:</ b > hip_hcc.h:399</ div ></ div >
< div class = "ttc" id = "classihipStreamCriticalBase__t_html" >< div class = "ttname" >< a href = "classihipStreamCriticalBase__t.html" > ihipStreamCriticalBase_t</ a ></ div >< div class = "ttdef" >< b > Definition:</ b > hip_hcc.h:352</ div ></ div >
< div class = "ttc" id = "classLockedAccessor_html" >< div class = "ttname" >< a href = "classLockedAccessor.html" > LockedAccessor</ a ></ div >< div class = "ttdef" >< b > Definition:</ b > hip_hcc.h:307</ div ></ div >
</ div > <!-- fragment --> </ div > <!-- contents -->
<!-- start footer part -->
< hr class = "footer" />< address class = "footer" >< small >
2016-04-19 22:44:58 +05:30
Generated on Tue Apr 19 2016 22:44:02 for HIP: Heterogenous-computing Interface for Portability by   < a href = "http://www.doxygen.org/index.html" >
2016-04-16 14:55:10 +05:30
< img class = "footer" src = "doxygen.png" alt = "doxygen" />
</ a > 1.8.6
</ small ></ address >
</ body >
</ html >
