2016-01-26 20:14:33 -06:00
<!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-02-18 20:43:03 -06:00
< title > HIP: Heterogenous-computing Interface for Portability: /home/bensander/HIP.public/include/hcc_detail/hip_runtime.h Source File</ title >
2016-01-26 20:14:33 -06:00
< 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_runtime.h</ div > </ div >
</ div > <!--header-->
< div class = "contents" >
< a href = "hcc__detail_2hip__runtime_8h.html" > Go to the documentation of this file.</ a >< div class = "fragment" >< div class = "line" >< a name = "l00001" ></ a >< span class = "lineno" > 1</ span >   < span class = "comment" > /*</ span ></ div >
< div class = "line" >< a name = "l00002" ></ a >< span class = "lineno" > 2</ span >   < span class = "comment" > 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" ></ span ></ div >
< div class = "line" >< a name = "l00004" ></ a >< span class = "lineno" > 4</ span >   < span class = "comment" > Permission is hereby granted, free of charge, to any person obtaining a copy</ span ></ div >
< div class = "line" >< a name = "l00005" ></ a >< span class = "lineno" > 5</ span >   < span class = "comment" > of this software and associated documentation files (the " Software" ), to deal</ span ></ div >
< div class = "line" >< a name = "l00006" ></ a >< span class = "lineno" > 6</ span >   < span class = "comment" > in the Software without restriction, including without limitation the rights</ span ></ div >
< div class = "line" >< a name = "l00007" ></ a >< span class = "lineno" > 7</ span >   < span class = "comment" > to use, copy, modify, merge, publish, distribute, sublicense, and/or sell</ span ></ div >
< div class = "line" >< a name = "l00008" ></ a >< span class = "lineno" > 8</ span >   < span class = "comment" > copies of the Software, and to permit persons to whom the Software is</ span ></ div >
< div class = "line" >< a name = "l00009" ></ a >< span class = "lineno" > 9</ span >   < span class = "comment" > furnished to do so, subject to the following conditions:</ span ></ div >
< div class = "line" >< a name = "l00010" ></ a >< span class = "lineno" > 10</ span >   < span class = "comment" ></ span ></ div >
< div class = "line" >< a name = "l00011" ></ a >< span class = "lineno" > 11</ span >   < span class = "comment" > The above copyright notice and this permission notice shall be included in</ span ></ div >
< div class = "line" >< a name = "l00012" ></ a >< span class = "lineno" > 12</ span >   < span class = "comment" > all copies or substantial portions of the Software.</ span ></ div >
< div class = "line" >< a name = "l00013" ></ a >< span class = "lineno" > 13</ span >   < span class = "comment" ></ span ></ div >
< div class = "line" >< a name = "l00014" ></ a >< span class = "lineno" > 14</ span >   < span class = "comment" > THE SOFTWARE IS PROVIDED " AS IS" , WITHOUT WARRANTY OF ANY KIND, EXPRESS OR</ span ></ div >
< div class = "line" >< a name = "l00015" ></ a >< span class = "lineno" > 15</ span >   < span class = "comment" > IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,</ span ></ div >
< div class = "line" >< a name = "l00016" ></ a >< span class = "lineno" > 16</ span >   < span class = "comment" > FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE</ span ></ div >
< div class = "line" >< a name = "l00017" ></ a >< span class = "lineno" > 17</ span >   < span class = "comment" > AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER</ span ></ div >
< div class = "line" >< a name = "l00018" ></ a >< span class = "lineno" > 18</ span >   < span class = "comment" > LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,</ span ></ div >
< div class = "line" >< a name = "l00019" ></ a >< span class = "lineno" > 19</ span >   < span class = "comment" > OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN</ span ></ div >
< div class = "line" >< a name = "l00020" ></ a >< span class = "lineno" > 20</ span >   < span class = "comment" > THE SOFTWARE.</ span ></ div >
< div class = "line" >< a name = "l00021" ></ a >< span class = "lineno" > 21</ span >   < span class = "comment" > */</ span ></ div >
< div class = "line" >< a name = "l00027" ></ a >< span class = "lineno" > 27</ span >   < span class = "preprocessor" > #pragma once</ 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 = "comment" > //---</ span ></ div >
< div class = "line" >< a name = "l00030" ></ a >< span class = "lineno" > 30</ span >   < span class = "comment" > // Top part of file can be compiled with any compiler</ span ></ div >
< div class = "line" >< a name = "l00031" ></ a >< span class = "lineno" > 31</ span >   </ div >
< div class = "line" >< a name = "l00032" ></ a >< span class = "lineno" > 32</ span >   </ div >
< div class = "line" >< a name = "l00033" ></ a >< span class = "lineno" > 33</ span >   < span class = "preprocessor" > #include < cstring> </ span ></ div >
< div class = "line" >< a name = "l00034" ></ a >< span class = "lineno" > 34</ span >   < span class = "preprocessor" > #include < cmath> </ span ></ div >
< div class = "line" >< a name = "l00035" ></ a >< span class = "lineno" > 35</ span >   < span class = "preprocessor" > #include < string.h> </ span ></ div >
< div class = "line" >< a name = "l00036" ></ a >< span class = "lineno" > 36</ span >   < span class = "preprocessor" > #include < stddef.h> </ span ></ div >
< div class = "line" >< a name = "l00037" ></ a >< span class = "lineno" > 37</ span >   </ div >
< div class = "line" >< a name = "l00038" ></ a >< span class = "lineno" > 38</ span >   </ div >
< div class = "line" >< a name = "l00039" ></ a >< span class = "lineno" > 39</ span >   < span class = "preprocessor" > #define CUDA_SUCCESS hipSuccess</ span ></ div >
< div class = "line" >< a name = "l00040" ></ a >< span class = "lineno" > 40</ span >   < span class = "preprocessor" ></ span ></ div >
2016-02-18 20:43:03 -06:00
< div class = "line" >< a name = "l00041" ></ a >< span class = "lineno" > 41</ span >   < span class = "preprocessor" > #include < < a class = "code" href = "hcc__detail_2hip__runtime__api_8h.html" > hip_runtime_api.h</ a > > </ span ></ div >
2016-01-26 20:14:33 -06:00
< div class = "line" >< a name = "l00042" ></ a >< span class = "lineno" > 42</ span >   </ div >
< div class = "line" >< a name = "l00043" ></ a >< span class = "lineno" > 43</ span >   < span class = "comment" > //---</ span ></ div >
< div class = "line" >< a name = "l00044" ></ a >< span class = "lineno" > 44</ span >   < span class = "comment" > // Remainder of this file only compiles with HCC</ span ></ div >
< div class = "line" >< a name = "l00045" ></ a >< span class = "lineno" > 45</ span >   < span class = "preprocessor" > #ifdef __HCC__</ span ></ div >
< div class = "line" >< a name = "l00046" ></ a >< span class = "lineno" > 46</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #include < hc.hpp> </ span ></ div >
< div class = "line" >< a name = "l00047" ></ a >< span class = "lineno" > 47</ span >   < span class = "preprocessor" > #include < grid_launch.h> </ span ></ div >
< div class = "line" >< a name = "l00048" ></ a >< span class = "lineno" > 48</ span >   </ div >
< div class = "line" >< a name = "l00049" ></ a >< span class = "lineno" > 49</ span >   < span class = "comment" > //TODO-HCC-GL - change this to typedef.</ span ></ div >
< div class = "line" >< a name = "l00050" ></ a >< span class = "lineno" > 50</ span >   < span class = "comment" > //typedef grid_launch_parm hipLaunchParm ;</ span ></ div >
< div class = "line" >< a name = "l00051" ></ a >< span class = "lineno" > 51</ span >   < span class = "preprocessor" > #define hipLaunchParm grid_launch_parm</ span ></ div >
< div class = "line" >< a name = "l00052" ></ a >< span class = "lineno" > 52</ span >   < span class = "preprocessor" ></ span ></ div >
2016-02-18 20:43:03 -06:00
< div class = "line" >< a name = "l00053" ></ a >< span class = "lineno" > 53</ span >   < span class = "preprocessor" > #include < < a class = "code" href = "hip__texture_8h.html" > hcc_detail/hip_texture.h</ a > > </ span ></ div >
< div class = "line" >< a name = "l00054" ></ a >< span class = "lineno" > 54</ span >   < span class = "preprocessor" > #include < < a class = "code" href = "host__defines_8h.html" > hcc_detail/host_defines.h</ a > > </ span ></ div >
2016-01-26 20:14:33 -06:00
< div class = "line" >< a name = "l00055" ></ a >< span class = "lineno" > 55</ span >   </ div >
< div class = "line" >< a name = "l00056" ></ a >< span class = "lineno" > 56</ span >   < span class = "comment" > // TODO-HCC remove old definitions ; ~1602 hcc supports __HCC_ACCELERATOR__ define.</ span ></ div >
< div class = "line" >< a name = "l00057" ></ a >< span class = "lineno" > 57</ span >   < span class = "preprocessor" > #if defined (__KALMAR_ACCELERATOR__) && not defined (__HCC_ACCELERATOR__)</ span ></ div >
< div class = "line" >< a name = "l00058" ></ a >< span class = "lineno" > 58</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #define __HCC_ACCELERATOR__ __KALMAR_ACCELERATOR__</ span ></ div >
< div class = "line" >< a name = "l00059" ></ a >< span class = "lineno" > 59</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #endif</ span ></ div >
< div class = "line" >< a name = "l00060" ></ a >< span class = "lineno" > 60</ span >   < span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00061" ></ a >< span class = "lineno" > 61</ span >   < span class = "comment" > // Feature tests:</ span ></ div >
< div class = "line" >< a name = "l00062" ></ a >< span class = "lineno" > 62</ span >   < span class = "preprocessor" > #if defined(__HCC_ACCELERATOR__) and (__HCC_ACCELERATOR__ != 0)</ span ></ div >
< div class = "line" >< a name = "l00063" ></ a >< span class = "lineno" > 63</ span >   < span class = "preprocessor" ></ span >< span class = "comment" > // Device compile and not host compile:</ span ></ div >
< div class = "line" >< a name = "l00064" ></ a >< span class = "lineno" > 64</ span >   </ div >
2016-02-18 20:43:03 -06:00
< div class = "line" >< a name = "l00065" ></ a >< span class = "lineno" > 65</ span >   < span class = "comment" > //TODO-HCC enable __HIP_ARCH_HAS_ATOMICS__ when HCC supports these.</ span ></ div >
< div class = "line" >< a name = "l00066" ></ a >< span class = "lineno" > 66</ span >   < span class = "comment" > // 32-bit Atomics:</ span ></ div >
< div class = "line" >< a name = "l00067" ></ a >< span class = "lineno" > 67</ span >   < span class = "preprocessor" > #define __HIP_ARCH_HAS_GLOBAL_INT32_ATOMICS__ (1)</ span ></ div >
< div class = "line" >< a name = "l00068" ></ a >< span class = "lineno" > 68</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #define __HIP_ARCH_HAS_GLOBAL_FLOAT_ATOMIC_EXCH__ (1)</ span ></ div >
< div class = "line" >< a name = "l00069" ></ a >< span class = "lineno" > 69</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #define __HIP_ARCH_HAS_SHARED_INT32_ATOMICS__ (0)</ span ></ div >
< div class = "line" >< a name = "l00070" ></ a >< span class = "lineno" > 70</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #define __HIP_ARCH_HAS_SHARED_FLOAT_ATOMIC_EXCH__ (0)</ span ></ div >
< div class = "line" >< a name = "l00071" ></ a >< span class = "lineno" > 71</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #define __HIP_ARCH_HAS_FLOAT_ATOMIC_ADD__ (0)</ span ></ div >
< div class = "line" >< a name = "l00072" ></ a >< span class = "lineno" > 72</ span >   < span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00073" ></ a >< span class = "lineno" > 73</ span >   < span class = "comment" > // 64-bit Atomics:</ span ></ div >
< div class = "line" >< a name = "l00074" ></ a >< span class = "lineno" > 74</ span >   < span class = "preprocessor" > #define __HIP_ARCH_HAS_GLOBAL_INT64_ATOMICS__ (1)</ span ></ div >
< div class = "line" >< a name = "l00075" ></ a >< span class = "lineno" > 75</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #define __HIP_ARCH_HAS_SHARED_INT64_ATOMICS__ (0)</ span ></ div >
< div class = "line" >< a name = "l00076" ></ a >< span class = "lineno" > 76</ span >   < span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00077" ></ a >< span class = "lineno" > 77</ span >   < span class = "comment" > // Doubles</ span ></ div >
< div class = "line" >< a name = "l00078" ></ a >< span class = "lineno" > 78</ span >   < span class = "preprocessor" > #define __HIP_ARCH_HAS_DOUBLES__ (1)</ span ></ div >
2016-01-26 20:14:33 -06:00
< div class = "line" >< a name = "l00079" ></ a >< span class = "lineno" > 79</ span >   < span class = "preprocessor" ></ span ></ div >
2016-02-18 20:43:03 -06:00
< div class = "line" >< a name = "l00080" ></ a >< span class = "lineno" > 80</ span >   < span class = "comment" > //warp cross-lane operations:</ span ></ div >
< div class = "line" >< a name = "l00081" ></ a >< span class = "lineno" > 81</ span >   < span class = "preprocessor" > #define __HIP_ARCH_HAS_WARP_VOTE__ (1)</ span ></ div >
< div class = "line" >< a name = "l00082" ></ a >< span class = "lineno" > 82</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #define __HIP_ARCH_HAS_WARP_BALLOT__ (1)</ span ></ div >
< div class = "line" >< a name = "l00083" ></ a >< span class = "lineno" > 83</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #define __HIP_ARCH_HAS_WARP_SHUFFLE__ (1)</ span ></ div >
< div class = "line" >< a name = "l00084" ></ a >< span class = "lineno" > 84</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #define __HIP_ARCH_HAS_WARP_FUNNEL_SHIFT__ (0)</ span ></ div >
< div class = "line" >< a name = "l00085" ></ a >< span class = "lineno" > 85</ span >   < span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00086" ></ a >< span class = "lineno" > 86</ span >   < span class = "comment" > //sync</ span ></ div >
< div class = "line" >< a name = "l00087" ></ a >< span class = "lineno" > 87</ span >   < span class = "preprocessor" > #define __HIP_ARCH_HAS_THREAD_FENCE_SYSTEM__ (0)</ span ></ div >
< div class = "line" >< a name = "l00088" ></ a >< span class = "lineno" > 88</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #define __HIP_ARCH_HAS_SYNC_THREAD_EXT__ (0)</ 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 = "comment" > // misc</ span ></ div >
< div class = "line" >< a name = "l00091" ></ a >< span class = "lineno" > 91</ span >   < span class = "preprocessor" > #define __HIP_ARCH_HAS_SURFACE_FUNCS__ (0)</ span ></ div >
< div class = "line" >< a name = "l00092" ></ a >< span class = "lineno" > 92</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #define __HIP_ARCH_HAS_3DGRID__ (1)</ span ></ div >
< div class = "line" >< a name = "l00093" ></ a >< span class = "lineno" > 93</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #define __HIP_ARCH_HAS_DYNAMIC_PARALLEL__ (0)</ 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 = "preprocessor" > #endif</ span ></ div >
< div class = "line" >< a name = "l00096" ></ a >< span class = "lineno" > 96</ span >   < span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00097" ></ a >< span class = "lineno" > 97</ span >   </ div >
< div class = "line" >< a name = "l00098" ></ a >< span class = "lineno" > 98</ 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 >   </ div >
< div class = "line" >< a name = "l00101" ></ a >< span class = "lineno" > 101</ span >   < span class = "comment" > //TODO-HCC this is currently ignored by HCC target of HIP</ span ></ div >
< div class = "line" >< a name = "l00102" ></ a >< span class = "lineno" > 102</ span >   < span class = "preprocessor" > #define __launch_bounds__(requiredMaxThreadsPerBlock, minBlocksPerMultiprocessor)</ span ></ div >
2016-01-26 20:14:33 -06:00
< div class = "line" >< a name = "l00103" ></ a >< span class = "lineno" > 103</ span >   < span class = "preprocessor" ></ span ></ div >
2016-02-18 20:43:03 -06:00
< div class = "line" >< a name = "l00104" ></ a >< span class = "lineno" > 104</ span >   < span class = "comment" > // Detect if we are compiling C++ mode or C mode</ span ></ div >
< div class = "line" >< a name = "l00105" ></ a >< span class = "lineno" > 105</ span >   < span class = "preprocessor" > #if defined(__cplusplus)</ span ></ div >
< div class = "line" >< a name = "l00106" ></ a >< span class = "lineno" > 106</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #define __HCC_CPP__</ span ></ div >
< div class = "line" >< a name = "l00107" ></ a >< span class = "lineno" > 107</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #elif defined(__STDC_VERSION__)</ span ></ div >
< div class = "line" >< a name = "l00108" ></ a >< span class = "lineno" > 108</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #define __HCC_C__</ span ></ div >
< div class = "line" >< a name = "l00109" ></ a >< span class = "lineno" > 109</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #endif</ span ></ div >
2016-01-26 20:14:33 -06:00
< div class = "line" >< a name = "l00110" ></ a >< span class = "lineno" > 110</ span >   < span class = "preprocessor" ></ span ></ div >
2016-02-18 20:43:03 -06:00
< div class = "line" >< a name = "l00111" ></ a >< span class = "lineno" > 111</ span >   < span class = "preprocessor" > #define clock_t long long int</ span ></ div >
< div class = "line" >< a name = "l00112" ></ a >< span class = "lineno" > 112</ span >   < span class = "preprocessor" ></ span > __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > int</ span > clock64() { < span class = "keywordflow" > return</ span > (< span class = "keywordtype" > long</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > int</ span > )hc::__clock_u64(); };</ div >
< div class = "line" >< a name = "l00113" ></ a >< span class = "lineno" > 113</ span >   __device__ < span class = "keyword" > inline</ span > clock_t clock() { < span class = "keywordflow" > return</ span > (clock_t)hc::__clock_u64(); };</ div >
< div class = "line" >< a name = "l00114" ></ a >< span class = "lineno" > 114</ span >   </ div >
< div class = "line" >< a name = "l00115" ></ a >< span class = "lineno" > 115</ span >   < span class = "comment" > //atomicAdd()</ span ></ div >
< div class = "line" >< a name = "l00116" ></ a >< span class = "lineno" > 116</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > int</ span > atomicAdd(< span class = "keywordtype" > int</ span > * address, < span class = "keywordtype" > int</ span > val)</ div >
< div class = "line" >< a name = "l00117" ></ a >< span class = "lineno" > 117</ span >   {</ div >
< div class = "line" >< a name = "l00118" ></ a >< span class = "lineno" > 118</ span >   < span class = "keywordflow" > return</ span > hc::atomic_fetch_add(address,val);</ div >
< div class = "line" >< a name = "l00119" ></ a >< span class = "lineno" > 119</ span >   }</ div >
< div class = "line" >< a name = "l00120" ></ a >< span class = "lineno" > 120</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > int</ span > atomicAdd(< span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > int</ span > * address,</ div >
< div class = "line" >< a name = "l00121" ></ a >< span class = "lineno" > 121</ span >   < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > int</ span > val)</ div >
< div class = "line" >< a name = "l00122" ></ a >< span class = "lineno" > 122</ span >   {</ div >
< div class = "line" >< a name = "l00123" ></ a >< span class = "lineno" > 123</ span >   < span class = "keywordflow" > return</ span > hc::atomic_fetch_add(address,val);</ 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 >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > int</ span > atomicAdd(< span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > int</ span > * address,</ div >
< div class = "line" >< a name = "l00126" ></ a >< span class = "lineno" > 126</ span >   < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > int</ span > val)</ div >
< div class = "line" >< a name = "l00127" ></ a >< span class = "lineno" > 127</ span >   {</ div >
< div class = "line" >< a name = "l00128" ></ a >< span class = "lineno" > 128</ span >   < span class = "keywordflow" > return</ span > (< span class = "keywordtype" > long</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > int</ span > )hc::atomic_fetch_add((uint64_t*)address,(uint64_t)val);</ div >
< div class = "line" >< a name = "l00129" ></ a >< span class = "lineno" > 129</ span >   }</ div >
< div class = "line" >< a name = "l00130" ></ a >< span class = "lineno" > 130</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > float</ span > atomicAdd(< span class = "keywordtype" > float</ span > * address, < span class = "keywordtype" > float</ span > val)</ 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 = "keywordflow" > return</ span > hc::atomic_fetch_add(address,val);</ div >
< div class = "line" >< a name = "l00133" ></ a >< span class = "lineno" > 133</ span >   }</ div >
< div class = "line" >< a name = "l00134" ></ a >< span class = "lineno" > 134</ span >   </ div >
< div class = "line" >< a name = "l00135" ></ a >< span class = "lineno" > 135</ span >   < span class = "comment" > //atomicSub()</ span ></ div >
< div class = "line" >< a name = "l00136" ></ a >< span class = "lineno" > 136</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > int</ span > atomicSub(< span class = "keywordtype" > int</ span > * address, < span class = "keywordtype" > int</ span > val)</ div >
< div class = "line" >< a name = "l00137" ></ a >< span class = "lineno" > 137</ span >   {</ div >
< div class = "line" >< a name = "l00138" ></ a >< span class = "lineno" > 138</ span >   < span class = "keywordflow" > return</ span > hc::atomic_fetch_sub(address,val);</ div >
< div class = "line" >< a name = "l00139" ></ a >< span class = "lineno" > 139</ span >   }</ div >
< div class = "line" >< a name = "l00140" ></ a >< span class = "lineno" > 140</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > int</ span > atomicSub(< span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > int</ span > * address,</ div >
< div class = "line" >< a name = "l00141" ></ a >< span class = "lineno" > 141</ span >   < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > int</ span > val)</ div >
< div class = "line" >< a name = "l00142" ></ a >< span class = "lineno" > 142</ span >   {</ div >
< div class = "line" >< a name = "l00143" ></ a >< span class = "lineno" > 143</ span >   < span class = "keywordflow" > return</ span > hc::atomic_fetch_sub(address,val);</ div >
< div class = "line" >< a name = "l00144" ></ a >< span class = "lineno" > 144</ span >   }</ div >
< div class = "line" >< a name = "l00145" ></ a >< span class = "lineno" > 145</ span >   </ div >
< div class = "line" >< a name = "l00146" ></ a >< span class = "lineno" > 146</ span >   < span class = "comment" > //atomicExch()</ span ></ div >
< div class = "line" >< a name = "l00147" ></ a >< span class = "lineno" > 147</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > int</ span > atomicExch(< span class = "keywordtype" > int</ span > * address, < span class = "keywordtype" > int</ span > val)</ div >
< div class = "line" >< a name = "l00148" ></ a >< span class = "lineno" > 148</ span >   {</ div >
< div class = "line" >< a name = "l00149" ></ a >< span class = "lineno" > 149</ span >   < span class = "keywordflow" > return</ span > hc::atomic_exchange(address,val);</ div >
< div class = "line" >< a name = "l00150" ></ a >< span class = "lineno" > 150</ span >   }</ div >
< div class = "line" >< a name = "l00151" ></ a >< span class = "lineno" > 151</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > int</ span > atomicExch(< span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > int</ span > * address,</ div >
< div class = "line" >< a name = "l00152" ></ a >< span class = "lineno" > 152</ span >   < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > int</ span > val)</ div >
< div class = "line" >< a name = "l00153" ></ a >< span class = "lineno" > 153</ span >   {</ div >
< div class = "line" >< a name = "l00154" ></ a >< span class = "lineno" > 154</ span >   < span class = "keywordflow" > return</ span > hc::atomic_exchange(address,val);</ div >
< div class = "line" >< a name = "l00155" ></ a >< span class = "lineno" > 155</ span >   }</ div >
< div class = "line" >< a name = "l00156" ></ a >< span class = "lineno" > 156</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > int</ span > atomicExch(< span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > int</ span > * address,</ div >
< div class = "line" >< a name = "l00157" ></ a >< span class = "lineno" > 157</ span >   < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > int</ span > val)</ div >
< div class = "line" >< a name = "l00158" ></ a >< span class = "lineno" > 158</ span >   {</ div >
< div class = "line" >< a name = "l00159" ></ a >< span class = "lineno" > 159</ span >   < span class = "keywordflow" > return</ span > (< span class = "keywordtype" > long</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > int</ span > )hc::atomic_exchange((uint64_t*)address,(uint64_t)val);</ div >
< div class = "line" >< a name = "l00160" ></ a >< span class = "lineno" > 160</ span >   }</ div >
< div class = "line" >< a name = "l00161" ></ a >< span class = "lineno" > 161</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > float</ span > atomicExch(< span class = "keywordtype" > float</ span > * address, < span class = "keywordtype" > float</ span > val)</ div >
< div class = "line" >< a name = "l00162" ></ a >< span class = "lineno" > 162</ span >   {</ div >
< div class = "line" >< a name = "l00163" ></ a >< span class = "lineno" > 163</ span >   < span class = "keywordflow" > return</ span > hc::atomic_exchange(address,val);</ 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" > //atomicMin()</ span ></ div >
< div class = "line" >< a name = "l00167" ></ a >< span class = "lineno" > 167</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > int</ span > atomicMin(< span class = "keywordtype" > int</ span > * address, < span class = "keywordtype" > int</ span > val)</ div >
< div class = "line" >< a name = "l00168" ></ a >< span class = "lineno" > 168</ span >   {</ div >
< div class = "line" >< a name = "l00169" ></ a >< span class = "lineno" > 169</ span >   < span class = "keywordflow" > return</ span > hc::atomic_fetch_min(address,val);</ div >
< div class = "line" >< a name = "l00170" ></ a >< span class = "lineno" > 170</ span >   }</ div >
< div class = "line" >< a name = "l00171" ></ a >< span class = "lineno" > 171</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > int</ span > atomicMin(< span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > int</ span > * address,</ div >
< div class = "line" >< a name = "l00172" ></ a >< span class = "lineno" > 172</ span >   < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > int</ span > val)</ div >
< div class = "line" >< a name = "l00173" ></ a >< span class = "lineno" > 173</ span >   {</ div >
< div class = "line" >< a name = "l00174" ></ a >< span class = "lineno" > 174</ span >   < span class = "keywordflow" > return</ span > hc::atomic_fetch_min(address,val);</ div >
< div class = "line" >< a name = "l00175" ></ a >< span class = "lineno" > 175</ span >   }</ div >
< div class = "line" >< a name = "l00176" ></ a >< span class = "lineno" > 176</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > int</ span > atomicMin(< span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > int</ span > * address,</ div >
< div class = "line" >< a name = "l00177" ></ a >< span class = "lineno" > 177</ span >   < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > int</ span > val)</ div >
< div class = "line" >< a name = "l00178" ></ a >< span class = "lineno" > 178</ span >   {</ div >
< div class = "line" >< a name = "l00179" ></ a >< span class = "lineno" > 179</ span >   < span class = "keywordflow" > return</ span > (< span class = "keywordtype" > long</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > int</ span > )hc::atomic_fetch_min((uint64_t*)address,(uint64_t)val);</ div >
< div class = "line" >< a name = "l00180" ></ a >< span class = "lineno" > 180</ span >   }</ div >
< div class = "line" >< a name = "l00181" ></ a >< span class = "lineno" > 181</ span >   </ div >
< div class = "line" >< a name = "l00182" ></ a >< span class = "lineno" > 182</ span >   < span class = "comment" > //atomicMax()</ span ></ div >
< div class = "line" >< a name = "l00183" ></ a >< span class = "lineno" > 183</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > int</ span > atomicMax(< span class = "keywordtype" > int</ span > * address, < span class = "keywordtype" > int</ span > val)</ 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 >   < span class = "keywordflow" > return</ span > hc::atomic_fetch_max(address,val);</ 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 >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > int</ span > atomicMax(< span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > int</ span > * address,</ div >
< div class = "line" >< a name = "l00188" ></ a >< span class = "lineno" > 188</ span >   < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > int</ span > val)</ div >
< div class = "line" >< a name = "l00189" ></ a >< span class = "lineno" > 189</ span >   {</ div >
< div class = "line" >< a name = "l00190" ></ a >< span class = "lineno" > 190</ span >   < span class = "keywordflow" > return</ span > hc::atomic_fetch_max(address,val);</ div >
< div class = "line" >< a name = "l00191" ></ a >< span class = "lineno" > 191</ span >   }</ div >
< div class = "line" >< a name = "l00192" ></ a >< span class = "lineno" > 192</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > int</ span > atomicMax(< span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > int</ span > * address,</ div >
< div class = "line" >< a name = "l00193" ></ a >< span class = "lineno" > 193</ span >   < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > int</ span > val)</ div >
< div class = "line" >< a name = "l00194" ></ a >< span class = "lineno" > 194</ span >   {</ div >
< div class = "line" >< a name = "l00195" ></ a >< span class = "lineno" > 195</ span >   < span class = "keywordflow" > return</ span > (< span class = "keywordtype" > long</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > int</ span > )hc::atomic_fetch_max((uint64_t*)address,(uint64_t)val);</ 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 >   </ div >
< div class = "line" >< a name = "l00198" ></ a >< span class = "lineno" > 198</ span >   < span class = "comment" > //atomicCAS()</ span ></ div >
< div class = "line" >< a name = "l00199" ></ a >< span class = "lineno" > 199</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > int</ span > atomicCAS(< span class = "keywordtype" > int</ span > * address, < span class = "keywordtype" > int</ span > compare, < span class = "keywordtype" > int</ span > val)</ div >
< div class = "line" >< a name = "l00200" ></ a >< span class = "lineno" > 200</ span >   {</ div >
< div class = "line" >< a name = "l00201" ></ a >< span class = "lineno" > 201</ span >   hc::atomic_compare_exchange(address,& compare,val);</ div >
< div class = "line" >< a name = "l00202" ></ a >< span class = "lineno" > 202</ span >   < span class = "keywordflow" > return</ span > *address;</ div >
2016-01-26 20:14:33 -06:00
< div class = "line" >< a name = "l00203" ></ a >< span class = "lineno" > 203</ span >   }</ div >
2016-02-18 20:43:03 -06:00
< div class = "line" >< a name = "l00204" ></ a >< span class = "lineno" > 204</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > int</ span > atomicCAS(< span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > int</ span > * address,</ div >
< div class = "line" >< a name = "l00205" ></ a >< span class = "lineno" > 205</ span >   < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > int</ span > compare,</ div >
< div class = "line" >< a name = "l00206" ></ a >< span class = "lineno" > 206</ span >   < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > int</ span > val)</ div >
2016-01-26 20:14:33 -06:00
< div class = "line" >< a name = "l00207" ></ a >< span class = "lineno" > 207</ span >   {</ div >
2016-02-18 20:43:03 -06:00
< div class = "line" >< a name = "l00208" ></ a >< span class = "lineno" > 208</ span >   hc::atomic_compare_exchange(address,& compare,val);</ div >
< div class = "line" >< a name = "l00209" ></ a >< span class = "lineno" > 209</ span >   < span class = "keywordflow" > return</ span > *address;</ div >
< div class = "line" >< a name = "l00210" ></ a >< span class = "lineno" > 210</ span >   }</ div >
< div class = "line" >< a name = "l00211" ></ a >< span class = "lineno" > 211</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > int</ span > atomicCAS(< span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > int</ span > * address,</ div >
< div class = "line" >< a name = "l00212" ></ a >< span class = "lineno" > 212</ span >   < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > int</ span > compare,</ div >
< div class = "line" >< a name = "l00213" ></ a >< span class = "lineno" > 213</ span >   < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > int</ span > val)</ div >
< div class = "line" >< a name = "l00214" ></ a >< span class = "lineno" > 214</ span >   {</ div >
< div class = "line" >< a name = "l00215" ></ a >< span class = "lineno" > 215</ span >   hc::atomic_compare_exchange((uint64_t*)address,(uint64_t*)& compare,(uint64_t)val);</ div >
< div class = "line" >< a name = "l00216" ></ a >< span class = "lineno" > 216</ span >   < span class = "keywordflow" > return</ span > *address;</ div >
< div class = "line" >< a name = "l00217" ></ a >< span class = "lineno" > 217</ span >   }</ div >
< div class = "line" >< a name = "l00218" ></ a >< span class = "lineno" > 218</ span >   </ div >
< div class = "line" >< a name = "l00219" ></ a >< span class = "lineno" > 219</ span >   < span class = "comment" > //atomicAnd()</ span ></ div >
< div class = "line" >< a name = "l00220" ></ a >< span class = "lineno" > 220</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > int</ span > atomicAnd(< span class = "keywordtype" > int</ span > * address, < span class = "keywordtype" > int</ span > val)</ 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 = "keywordflow" > return</ span > hc::atomic_fetch_and(address,val);</ div >
< div class = "line" >< a name = "l00223" ></ a >< span class = "lineno" > 223</ span >   }</ div >
< div class = "line" >< a name = "l00224" ></ a >< span class = "lineno" > 224</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > int</ span > atomicAnd(< span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > int</ span > * address,</ div >
2016-01-26 20:14:33 -06:00
< div class = "line" >< a name = "l00225" ></ a >< span class = "lineno" > 225</ span >   < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > int</ span > val)</ div >
< div class = "line" >< a name = "l00226" ></ a >< span class = "lineno" > 226</ span >   {</ div >
2016-02-18 20:43:03 -06:00
< div class = "line" >< a name = "l00227" ></ a >< span class = "lineno" > 227</ span >   < span class = "keywordflow" > return</ span > hc::atomic_fetch_and(address,val);</ 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 >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > int</ span > atomicAnd(< span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > int</ span > * address,</ div >
< div class = "line" >< a name = "l00230" ></ a >< span class = "lineno" > 230</ span >   < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > int</ span > val)</ div >
< div class = "line" >< a name = "l00231" ></ a >< span class = "lineno" > 231</ span >   {</ div >
< div class = "line" >< a name = "l00232" ></ a >< span class = "lineno" > 232</ span >   < span class = "keywordflow" > return</ span > (< span class = "keywordtype" > long</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > int</ span > )hc::atomic_fetch_and((uint64_t*)address,(uint64_t)val);</ div >
< div class = "line" >< a name = "l00233" ></ a >< span class = "lineno" > 233</ span >   }</ div >
< div class = "line" >< a name = "l00234" ></ a >< span class = "lineno" > 234</ span >   </ div >
< div class = "line" >< a name = "l00235" ></ a >< span class = "lineno" > 235</ span >   < span class = "comment" > //atomicOr()</ span ></ div >
< div class = "line" >< a name = "l00236" ></ a >< span class = "lineno" > 236</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > int</ span > atomicOr(< span class = "keywordtype" > int</ span > * address, < span class = "keywordtype" > int</ span > val)</ 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 = "keywordflow" > return</ span > hc::atomic_fetch_or(address,val);</ div >
< div class = "line" >< a name = "l00239" ></ a >< span class = "lineno" > 239</ span >   }</ div >
< div class = "line" >< a name = "l00240" ></ a >< span class = "lineno" > 240</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > int</ span > atomicOr(< span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > int</ span > * address,</ div >
< div class = "line" >< a name = "l00241" ></ a >< span class = "lineno" > 241</ span >   < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > int</ span > val)</ div >
< div class = "line" >< a name = "l00242" ></ a >< span class = "lineno" > 242</ span >   {</ div >
< div class = "line" >< a name = "l00243" ></ a >< span class = "lineno" > 243</ span >   < span class = "keywordflow" > return</ span > hc::atomic_fetch_or(address,val);</ 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 >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > int</ span > atomicOr(< span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > int</ span > * address,</ div >
< div class = "line" >< a name = "l00246" ></ a >< span class = "lineno" > 246</ span >   < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > int</ span > val)</ div >
< div class = "line" >< a name = "l00247" ></ a >< span class = "lineno" > 247</ span >   {</ div >
< div class = "line" >< a name = "l00248" ></ a >< span class = "lineno" > 248</ span >   < span class = "keywordflow" > return</ span > (< span class = "keywordtype" > long</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > int</ span > )hc::atomic_fetch_or((uint64_t*)address,(uint64_t)val);</ div >
< div class = "line" >< a name = "l00249" ></ a >< span class = "lineno" > 249</ span >   }</ div >
< div class = "line" >< a name = "l00250" ></ a >< span class = "lineno" > 250</ span >   </ div >
< div class = "line" >< a name = "l00251" ></ a >< span class = "lineno" > 251</ span >   < span class = "comment" > //atomicXor()</ span ></ div >
< div class = "line" >< a name = "l00252" ></ a >< span class = "lineno" > 252</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > int</ span > atomicXor(< span class = "keywordtype" > int</ span > * address, < span class = "keywordtype" > int</ span > val)</ div >
< div class = "line" >< a name = "l00253" ></ a >< span class = "lineno" > 253</ span >   {</ div >
< div class = "line" >< a name = "l00254" ></ a >< span class = "lineno" > 254</ span >   < span class = "keywordflow" > return</ span > hc::atomic_fetch_xor(address,val);</ 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 >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > int</ span > atomicXor(< span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > int</ span > * address,</ div >
< div class = "line" >< a name = "l00257" ></ a >< span class = "lineno" > 257</ span >   < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > int</ span > val)</ 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 = "keywordflow" > return</ span > hc::atomic_fetch_xor(address,val);</ 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 >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > int</ span > atomicXor(< span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > int</ span > * address,</ div >
< div class = "line" >< a name = "l00262" ></ a >< span class = "lineno" > 262</ span >   < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > int</ span > val)</ div >
< div class = "line" >< a name = "l00263" ></ a >< span class = "lineno" > 263</ span >   {</ div >
< div class = "line" >< a name = "l00264" ></ a >< span class = "lineno" > 264</ span >   < span class = "keywordflow" > return</ span > (< span class = "keywordtype" > long</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > int</ span > )hc::atomic_fetch_xor((uint64_t*)address,(uint64_t)val);</ div >
< div class = "line" >< a name = "l00265" ></ a >< span class = "lineno" > 265</ span >   }</ div >
< div class = "line" >< a name = "l00266" ></ a >< span class = "lineno" > 266</ span >   </ div >
< div class = "line" >< a name = "l00267" ></ a >< span class = "lineno" > 267</ span >   < span class = "preprocessor" > #include < hc.hpp> </ span ></ div >
< div class = "line" >< a name = "l00268" ></ a >< span class = "lineno" > 268</ span >   < span class = "comment" > // integer intrinsic function __poc __clz __ffs __brev</ span ></ div >
< div class = "line" >< a name = "l00269" ></ a >< span class = "lineno" > 269</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > int</ span > __popc( < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > int</ span > input)</ 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 >   < span class = "keywordflow" > return</ span > hc::__popcount_u32_b32( input);</ div >
< div class = "line" >< a name = "l00272" ></ a >< span class = "lineno" > 272</ span >   }</ 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 >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > int</ span > __popcll( < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > int</ span > input)</ 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 >   < span class = "keywordflow" > return</ span > hc::__popcount_u32_b64(input);</ 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 >   </ div >
< div class = "line" >< a name = "l00279" ></ a >< span class = "lineno" > 279</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > int</ span > __clz(< span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > int</ span > input)</ 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 = "keywordflow" > return</ span > hc::__firstbit_u32_u32( input);</ div >
< div class = "line" >< a name = "l00282" ></ a >< span class = "lineno" > 282</ span >   }</ div >
< div class = "line" >< a name = "l00283" ></ a >< span class = "lineno" > 283</ span >   </ div >
< div class = "line" >< a name = "l00284" ></ a >< span class = "lineno" > 284</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > int</ span > __clzll(< span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > int</ span > input)</ 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 >   < span class = "keywordflow" > return</ span > hc::__firstbit_u32_u64( input);</ 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 >   </ div >
< div class = "line" >< a name = "l00289" ></ a >< span class = "lineno" > 289</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > int</ span > __clz(< span class = "keywordtype" > int</ span > input)</ div >
2016-01-26 20:14:33 -06:00
< div class = "line" >< a name = "l00290" ></ a >< span class = "lineno" > 290</ span >   {</ div >
2016-02-18 20:43:03 -06:00
< div class = "line" >< a name = "l00291" ></ a >< span class = "lineno" > 291</ span >   < span class = "keywordflow" > return</ span > hc::__firstbit_u32_s32( input);</ div >
2016-01-26 20:14:33 -06:00
< div class = "line" >< a name = "l00292" ></ a >< span class = "lineno" > 292</ span >   }</ div >
< div class = "line" >< a name = "l00293" ></ a >< span class = "lineno" > 293</ span >   </ div >
2016-02-18 20:43:03 -06:00
< div class = "line" >< a name = "l00294" ></ a >< span class = "lineno" > 294</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > int</ span > __clzll(< span class = "keywordtype" > long</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > int</ span > input)</ div >
2016-01-26 20:14:33 -06:00
< div class = "line" >< a name = "l00295" ></ a >< span class = "lineno" > 295</ span >   {</ div >
2016-02-18 20:43:03 -06:00
< div class = "line" >< a name = "l00296" ></ a >< span class = "lineno" > 296</ span >   < span class = "keywordflow" > return</ span > hc::__firstbit_u32_s64( input);</ div >
2016-01-26 20:14:33 -06:00
< div class = "line" >< a name = "l00297" ></ a >< span class = "lineno" > 297</ span >   }</ div >
< div class = "line" >< a name = "l00298" ></ a >< span class = "lineno" > 298</ span >   </ div >
2016-02-18 20:43:03 -06:00
< div class = "line" >< a name = "l00299" ></ a >< span class = "lineno" > 299</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > int</ span > __ffs(< span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > int</ span > input)</ div >
2016-01-26 20:14:33 -06:00
< div class = "line" >< a name = "l00300" ></ a >< span class = "lineno" > 300</ span >   {</ div >
2016-02-18 20:43:03 -06:00
< div class = "line" >< a name = "l00301" ></ a >< span class = "lineno" > 301</ span >   < span class = "keywordflow" > return</ span > hc::__lastbit_u32_u32( input)+1;</ div >
2016-01-26 20:14:33 -06:00
< div class = "line" >< a name = "l00302" ></ a >< span class = "lineno" > 302</ span >   }</ div >
< div class = "line" >< a name = "l00303" ></ a >< span class = "lineno" > 303</ span >   </ div >
2016-02-18 20:43:03 -06:00
< div class = "line" >< a name = "l00304" ></ a >< span class = "lineno" > 304</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > int</ span > __ffsll(< span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > int</ span > input)</ div >
2016-01-26 20:14:33 -06:00
< div class = "line" >< a name = "l00305" ></ a >< span class = "lineno" > 305</ span >   {</ div >
2016-02-18 20:43:03 -06:00
< div class = "line" >< a name = "l00306" ></ a >< span class = "lineno" > 306</ span >   < span class = "keywordflow" > return</ span > hc::__lastbit_u32_u64( input)+1;</ div >
2016-01-26 20:14:33 -06:00
< div class = "line" >< a name = "l00307" ></ a >< span class = "lineno" > 307</ span >   }</ div >
< div class = "line" >< a name = "l00308" ></ a >< span class = "lineno" > 308</ span >   </ div >
2016-02-18 20:43:03 -06:00
< div class = "line" >< a name = "l00309" ></ a >< span class = "lineno" > 309</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > int</ span > __ffs(< span class = "keywordtype" > int</ span > input)</ div >
2016-01-26 20:14:33 -06:00
< div class = "line" >< a name = "l00310" ></ a >< span class = "lineno" > 310</ span >   {</ div >
2016-02-18 20:43:03 -06:00
< div class = "line" >< a name = "l00311" ></ a >< span class = "lineno" > 311</ span >   < span class = "keywordflow" > return</ span > hc::__lastbit_u32_s32( input)+1;</ div >
2016-01-26 20:14:33 -06:00
< div class = "line" >< a name = "l00312" ></ a >< span class = "lineno" > 312</ span >   }</ div >
< div class = "line" >< a name = "l00313" ></ a >< span class = "lineno" > 313</ span >   </ div >
2016-02-18 20:43:03 -06:00
< div class = "line" >< a name = "l00314" ></ a >< span class = "lineno" > 314</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > int</ span > __ffsll(< span class = "keywordtype" > long</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > int</ span > input)</ div >
2016-01-26 20:14:33 -06:00
< div class = "line" >< a name = "l00315" ></ a >< span class = "lineno" > 315</ span >   {</ div >
2016-02-18 20:43:03 -06:00
< div class = "line" >< a name = "l00316" ></ a >< span class = "lineno" > 316</ span >   < span class = "keywordflow" > return</ span > hc::__lastbit_u32_s64( input)+1;</ div >
2016-01-26 20:14:33 -06:00
< 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 >   </ div >
2016-02-18 20:43:03 -06:00
< div class = "line" >< a name = "l00319" ></ a >< span class = "lineno" > 319</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > int</ span > __brev( < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > int</ span > input)</ div >
2016-01-26 20:14:33 -06:00
< div class = "line" >< a name = "l00320" ></ a >< span class = "lineno" > 320</ span >   {</ div >
2016-02-18 20:43:03 -06:00
< div class = "line" >< a name = "l00321" ></ a >< span class = "lineno" > 321</ span >   < span class = "keywordflow" > return</ span > hc::__bitrev_b32( input);</ div >
2016-01-26 20:14:33 -06:00
< 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 >
2016-02-18 20:43:03 -06:00
< div class = "line" >< a name = "l00324" ></ a >< span class = "lineno" > 324</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > int</ span > __brevll( < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > int</ span > input)</ div >
2016-01-26 20:14:33 -06:00
< div class = "line" >< a name = "l00325" ></ a >< span class = "lineno" > 325</ span >   {</ div >
2016-02-18 20:43:03 -06:00
< div class = "line" >< a name = "l00326" ></ a >< span class = "lineno" > 326</ span >   < span class = "keywordflow" > return</ span > hc::__bitrev_b64( input);</ div >
2016-01-26 20:14:33 -06:00
< div class = "line" >< a name = "l00327" ></ a >< span class = "lineno" > 327</ span >   }</ div >
< div class = "line" >< a name = "l00328" ></ a >< span class = "lineno" > 328</ span >   </ div >
2016-02-18 20:43:03 -06:00
< div class = "line" >< a name = "l00329" ></ a >< span class = "lineno" > 329</ span >   < span class = "comment" > // warp vote function __all __any __ballot</ span ></ div >
< div class = "line" >< a name = "l00330" ></ a >< span class = "lineno" > 330</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > int</ span > __all( < span class = "keywordtype" > int</ span > input)</ div >
< div class = "line" >< a name = "l00331" ></ a >< span class = "lineno" > 331</ span >   {</ div >
< div class = "line" >< a name = "l00332" ></ a >< span class = "lineno" > 332</ span >   < span class = "keywordflow" > return</ span > hc::__all( input);</ div >
< div class = "line" >< a name = "l00333" ></ a >< span class = "lineno" > 333</ span >   }</ div >
< div class = "line" >< a name = "l00334" ></ a >< span class = "lineno" > 334</ span >   </ div >
< div class = "line" >< a name = "l00335" ></ a >< span class = "lineno" > 335</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > int</ span > __any( < span class = "keywordtype" > int</ span > input)</ 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 >   < span class = "keywordflow" > if</ span > ( hc::__any( input)!=0) < span class = "keywordflow" > return</ span > 1;</ div >
< div class = "line" >< a name = "l00338" ></ a >< span class = "lineno" > 338</ span >   < span class = "keywordflow" > else</ span > < span class = "keywordflow" > return</ span > 0;</ div >
< div class = "line" >< a name = "l00339" ></ a >< span class = "lineno" > 339</ span >   }</ div >
2016-01-26 20:14:33 -06:00
< div class = "line" >< a name = "l00340" ></ a >< span class = "lineno" > 340</ span >   </ div >
2016-02-18 20:43:03 -06:00
< div class = "line" >< a name = "l00341" ></ a >< span class = "lineno" > 341</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > long</ span > < span class = "keywordtype" > int</ span > __ballot( < span class = "keywordtype" > int</ span > input)</ div >
2016-01-26 20:14:33 -06:00
< div class = "line" >< a name = "l00342" ></ a >< span class = "lineno" > 342</ span >   {</ div >
2016-02-18 20:43:03 -06:00
< div class = "line" >< a name = "l00343" ></ a >< span class = "lineno" > 343</ span >   < span class = "keywordflow" > return</ span > hc::__ballot( input);</ div >
2016-01-26 20:14:33 -06:00
< div class = "line" >< a name = "l00344" ></ a >< span class = "lineno" > 344</ span >   }</ div >
< div class = "line" >< a name = "l00345" ></ a >< span class = "lineno" > 345</ span >   </ div >
2016-02-18 20:43:03 -06:00
< div class = "line" >< a name = "l00346" ></ a >< span class = "lineno" > 346</ span >   < span class = "comment" > // warp shuffle functions</ span ></ div >
< div class = "line" >< a name = "l00347" ></ a >< span class = "lineno" > 347</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > int</ span > __shfl(< span class = "keywordtype" > int</ span > input, < span class = "keywordtype" > int</ span > lane, < span class = "keywordtype" > int</ span > width)</ div >
2016-01-26 20:14:33 -06:00
< div class = "line" >< a name = "l00348" ></ a >< span class = "lineno" > 348</ span >   {</ div >
2016-02-18 20:43:03 -06:00
< div class = "line" >< a name = "l00349" ></ a >< span class = "lineno" > 349</ span >   < span class = "keywordflow" > return</ span > hc::__shfl(input,lane,width);</ div >
2016-01-26 20:14:33 -06:00
< 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 >   </ div >
2016-02-18 20:43:03 -06:00
< div class = "line" >< a name = "l00352" ></ a >< span class = "lineno" > 352</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > int</ span > __shfl_up(< span class = "keywordtype" > int</ span > input, < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > int</ span > lane_delta, < span class = "keywordtype" > int</ span > width)</ 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 = "keywordflow" > return</ span > hc::__shfl_up(input,lane_delta,width);</ div >
< div class = "line" >< a name = "l00355" ></ a >< span class = "lineno" > 355</ span >   }</ div >
< div class = "line" >< a name = "l00356" ></ a >< span class = "lineno" > 356</ span >   </ div >
< div class = "line" >< a name = "l00357" ></ a >< span class = "lineno" > 357</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > int</ span > __shfl_down(< span class = "keywordtype" > int</ span > input, < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > int</ span > lane_delta, < span class = "keywordtype" > int</ span > width)</ div >
< div class = "line" >< a name = "l00358" ></ a >< span class = "lineno" > 358</ span >   {</ div >
< div class = "line" >< a name = "l00359" ></ a >< span class = "lineno" > 359</ span >   < span class = "keywordflow" > return</ span > hc::__shfl_down(input,lane_delta,width);</ div >
< div class = "line" >< a name = "l00360" ></ a >< span class = "lineno" > 360</ span >   }</ div >
2016-01-26 20:14:33 -06:00
< div class = "line" >< a name = "l00361" ></ a >< span class = "lineno" > 361</ span >   </ div >
2016-02-18 20:43:03 -06:00
< div class = "line" >< a name = "l00362" ></ a >< span class = "lineno" > 362</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > int</ span > __shfl_xor(< span class = "keywordtype" > int</ span > input, < span class = "keywordtype" > int</ span > lane_mask, < span class = "keywordtype" > int</ span > width)</ 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 >   < span class = "keywordflow" > return</ span > hc::__shfl_xor(input,lane_mask,width);</ div >
< div class = "line" >< a name = "l00365" ></ a >< span class = "lineno" > 365</ span >   }</ div >
2016-01-26 20:14:33 -06:00
< div class = "line" >< a name = "l00366" ></ a >< span class = "lineno" > 366</ span >   </ div >
2016-02-18 20:43:03 -06:00
< div class = "line" >< a name = "l00367" ></ a >< span class = "lineno" > 367</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > float</ span > __shfl(< span class = "keywordtype" > float</ span > input, < span class = "keywordtype" > int</ span > lane, < span class = "keywordtype" > int</ span > width)</ 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 >   < span class = "keywordflow" > return</ span > hc::__shfl(input,lane,width);</ 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 >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > float</ span > __shfl_up(< span class = "keywordtype" > float</ span > input, < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > int</ span > lane_delta, < span class = "keywordtype" > int</ span > width)</ div >
< div class = "line" >< a name = "l00373" ></ a >< span class = "lineno" > 373</ span >   {</ div >
< div class = "line" >< a name = "l00374" ></ a >< span class = "lineno" > 374</ span >   < span class = "keywordflow" > return</ span > hc::__shfl_up(input,lane_delta,width);</ 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 >   </ div >
< div class = "line" >< a name = "l00377" ></ a >< span class = "lineno" > 377</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > float</ span > __shfl_down(< span class = "keywordtype" > float</ span > input, < span class = "keywordtype" > unsigned</ span > < span class = "keywordtype" > int</ span > lane_delta, < span class = "keywordtype" > int</ span > width)</ div >
< div class = "line" >< a name = "l00378" ></ a >< span class = "lineno" > 378</ span >   {</ div >
< div class = "line" >< a name = "l00379" ></ a >< span class = "lineno" > 379</ span >   < span class = "keywordflow" > return</ span > hc::__shfl_down(input,lane_delta,width);</ div >
< div class = "line" >< a name = "l00380" ></ a >< span class = "lineno" > 380</ span >   }</ div >
2016-01-26 20:14:33 -06:00
< div class = "line" >< a name = "l00381" ></ a >< span class = "lineno" > 381</ span >   </ div >
2016-02-18 20:43:03 -06:00
< div class = "line" >< a name = "l00382" ></ a >< span class = "lineno" > 382</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > float</ span > __shfl_xor(< span class = "keywordtype" > float</ span > input, < span class = "keywordtype" > int</ span > lane_mask, < span class = "keywordtype" > int</ span > width)</ div >
< div class = "line" >< a name = "l00383" ></ a >< span class = "lineno" > 383</ span >   {</ div >
< div class = "line" >< a name = "l00384" ></ a >< span class = "lineno" > 384</ span >   < span class = "keywordflow" > return</ span > hc::__shfl_xor(input,lane_mask,width);</ div >
< div class = "line" >< a name = "l00385" ></ a >< span class = "lineno" > 385</ span >   }</ div >
< div class = "line" >< a name = "l00386" ></ a >< span class = "lineno" > 386</ 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 >   < span class = "preprocessor" > #include < hc_math.hpp> </ span ></ div >
< div class = "line" >< a name = "l00389" ></ a >< span class = "lineno" > 389</ span >   < span class = "comment" > // TODO: Choose whether default is precise math or fast math based on compilation flag.</ span ></ div >
< div class = "line" >< a name = "l00390" ></ a >< span class = "lineno" > 390</ span >   < span class = "preprocessor" > #ifdef __HCC_ACCELERATOR__</ span ></ div >
< div class = "line" >< a name = "l00391" ></ a >< span class = "lineno" > 391</ span >   < span class = "preprocessor" ></ span >< span class = "keyword" > using namespace </ span > hc::precise_math;</ div >
< div class = "line" >< a name = "l00392" ></ a >< span class = "lineno" > 392</ span >   < span class = "preprocessor" > #endif</ span ></ div >
< div class = "line" >< a name = "l00393" ></ a >< span class = "lineno" > 393</ span >   < span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00394" ></ a >< span class = "lineno" > 394</ span >   < span class = "comment" > //TODO: Undo this once min/max functions are supported by hc</ span ></ div >
< div class = "line" >< a name = "l00395" ></ a >< span class = "lineno" > 395</ span >   < span class = "keyword" > inline</ span > < span class = "keywordtype" > int</ span > min(< span class = "keywordtype" > int</ span > arg1, < span class = "keywordtype" > int</ span > arg2) __attribute((hc,cpu)) { \</ div >
< div class = "line" >< a name = "l00396" ></ a >< span class = "lineno" > 396</ span >   return (< span class = "keywordtype" > int</ span > )(hc::precise_math::fmin((< span class = "keywordtype" > float</ span > )arg1, (< span class = "keywordtype" > float</ span > )arg2));}</ div >
< div class = "line" >< a name = "l00397" ></ a >< span class = "lineno" > 397</ span >   < span class = "keyword" > inline</ span > < span class = "keywordtype" > int</ span > max(< span class = "keywordtype" > int</ span > arg1, < span class = "keywordtype" > int</ span > arg2) __attribute((hc,cpu)) { \</ div >
< div class = "line" >< a name = "l00398" ></ a >< span class = "lineno" > 398</ span >   return (< span class = "keywordtype" > int</ span > )(hc::precise_math::fmax((< span class = "keywordtype" > float</ span > )arg1, (< span class = "keywordtype" > float</ span > )arg2));}</ div >
< div class = "line" >< a name = "l00399" ></ a >< span class = "lineno" > 399</ span >   </ div >
< div class = "line" >< a name = "l00400" ></ a >< span class = "lineno" > 400</ span >   </ div >
< div class = "line" >< a name = "l00401" ></ a >< span class = "lineno" > 401</ span >   < span class = "comment" > //TODO - add a couple fast math operations here, the set here will grow :</ span ></ div >
< div class = "line" >< a name = "l00402" ></ a >< span class = "lineno" > 402</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > float</ span > __cosf(< span class = "keywordtype" > float</ span > x) {< span class = "keywordflow" > return</ span > hc::fast_math::cosf(x); };</ div >
< div class = "line" >< a name = "l00403" ></ a >< span class = "lineno" > 403</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > float</ span > __expf(< span class = "keywordtype" > float</ span > x) {< span class = "keywordflow" > return</ span > hc::fast_math::expf(x); };</ div >
< div class = "line" >< a name = "l00404" ></ a >< span class = "lineno" > 404</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > float</ span > __frsqrt_rn(< span class = "keywordtype" > float</ span > x) {< span class = "keywordflow" > return</ span > hc::fast_math::rsqrt(x); };</ div >
< div class = "line" >< a name = "l00405" ></ a >< span class = "lineno" > 405</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > float</ span > __fsqrt_rd(< span class = "keywordtype" > float</ span > x) {< span class = "keywordflow" > return</ span > hc::fast_math::sqrt(x); };</ div >
< div class = "line" >< a name = "l00406" ></ a >< span class = "lineno" > 406</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > float</ span > __fsqrt_rn(< span class = "keywordtype" > float</ span > x) {< span class = "keywordflow" > return</ span > hc::fast_math::sqrt(x); };</ div >
< div class = "line" >< a name = "l00407" ></ a >< span class = "lineno" > 407</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > float</ span > __fsqrt_ru(< span class = "keywordtype" > float</ span > x) {< span class = "keywordflow" > return</ span > hc::fast_math::sqrt(x); };</ div >
< div class = "line" >< a name = "l00408" ></ a >< span class = "lineno" > 408</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > float</ span > __fsqrt_rz(< span class = "keywordtype" > float</ span > x) {< span class = "keywordflow" > return</ span > hc::fast_math::sqrt(x); };</ div >
< div class = "line" >< a name = "l00409" ></ a >< span class = "lineno" > 409</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > float</ span > __log10f(< span class = "keywordtype" > float</ span > x) {< span class = "keywordflow" > return</ span > hc::fast_math::log10f(x); };</ div >
< div class = "line" >< a name = "l00410" ></ a >< span class = "lineno" > 410</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > float</ span > __log2f(< span class = "keywordtype" > float</ span > x) {< span class = "keywordflow" > return</ span > hc::fast_math::log2f(x); };</ div >
< div class = "line" >< a name = "l00411" ></ a >< span class = "lineno" > 411</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > float</ span > __logf(< span class = "keywordtype" > float</ span > x) {< span class = "keywordflow" > return</ span > hc::fast_math::logf(x); };</ div >
< div class = "line" >< a name = "l00412" ></ a >< span class = "lineno" > 412</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > float</ span > __powf(< span class = "keywordtype" > float</ span > base, < span class = "keywordtype" > float</ span > exponent) {< span class = "keywordflow" > return</ span > hc::fast_math::powf(base, exponent); };</ div >
< div class = "line" >< a name = "l00413" ></ a >< span class = "lineno" > 413</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > void</ span > __sincosf(< span class = "keywordtype" > float</ span > x, < span class = "keywordtype" > float</ span > *s, < span class = "keywordtype" > float</ span > *c) {< span class = "keywordflow" > return</ span > hc::fast_math::sincosf(x, s, c); };</ div >
< div class = "line" >< a name = "l00414" ></ a >< span class = "lineno" > 414</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > float</ span > __sinf(< span class = "keywordtype" > float</ span > x) {< span class = "keywordflow" > return</ span > hc::fast_math::sinf(x); };</ div >
< div class = "line" >< a name = "l00415" ></ a >< span class = "lineno" > 415</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > float</ span > __tanf(< span class = "keywordtype" > float</ span > x) {< span class = "keywordflow" > return</ span > hc::fast_math::tanf(x); };</ div >
< div class = "line" >< a name = "l00416" ></ a >< span class = "lineno" > 416</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > float</ span > __dsqrt_rd(< span class = "keywordtype" > double</ span > x) {< span class = "keywordflow" > return</ span > hc::fast_math::sqrt(x); };</ div >
< div class = "line" >< a name = "l00417" ></ a >< span class = "lineno" > 417</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > float</ span > __dsqrt_rn(< span class = "keywordtype" > double</ span > x) {< span class = "keywordflow" > return</ span > hc::fast_math::sqrt(x); };</ div >
< div class = "line" >< a name = "l00418" ></ a >< span class = "lineno" > 418</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > float</ span > __dsqrt_ru(< span class = "keywordtype" > double</ span > x) {< span class = "keywordflow" > return</ span > hc::fast_math::sqrt(x); };</ div >
< div class = "line" >< a name = "l00419" ></ a >< span class = "lineno" > 419</ span >   __device__ < span class = "keyword" > inline</ span > < span class = "keywordtype" > float</ span > __dsqrt_rz(< span class = "keywordtype" > double</ span > x) {< span class = "keywordflow" > return</ span > hc::fast_math::sqrt(x); };</ div >
2016-01-26 20:14:33 -06:00
< div class = "line" >< a name = "l00420" ></ a >< span class = "lineno" > 420</ span >   </ div >
2016-02-18 20:43:03 -06:00
< div class = "line" >< a name = "l00424" ></ a >< span class = "lineno" >< a class = "line" href = "hcc__detail_2hip__runtime_8h.html#a48f5f9da77c5fab1fbcf0205bb347d89" > 424</ a ></ span >   < span class = "preprocessor" > #define hipThreadIdx_x (amp_get_local_id(2))</ span ></ div >
< div class = "line" >< a name = "l00425" ></ a >< span class = "lineno" > 425</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #define hipThreadIdx_y (amp_get_local_id(1))</ span ></ div >
< div class = "line" >< a name = "l00426" ></ a >< span class = "lineno" > 426</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #define hipThreadIdx_z (amp_get_local_id(0))</ span ></ div >
< div class = "line" >< a name = "l00427" ></ a >< span class = "lineno" > 427</ span >   < span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00428" ></ a >< span class = "lineno" > 428</ span >   < span class = "preprocessor" > #define hipBlockIdx_x (hc_get_group_id(2))</ span ></ div >
< div class = "line" >< a name = "l00429" ></ a >< span class = "lineno" > 429</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #define hipBlockIdx_y (hc_get_group_id(1))</ span ></ div >
< div class = "line" >< a name = "l00430" ></ a >< span class = "lineno" > 430</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #define hipBlockIdx_z (hc_get_group_id(0))</ span ></ div >
< div class = "line" >< a name = "l00431" ></ a >< span class = "lineno" > 431</ span >   < span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00432" ></ a >< span class = "lineno" > 432</ span >   < span class = "preprocessor" > #define hipBlockDim_x (amp_get_local_size(2))</ span ></ div >
< div class = "line" >< a name = "l00433" ></ a >< span class = "lineno" > 433</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #define hipBlockDim_y (amp_get_local_size(1))</ span ></ div >
< div class = "line" >< a name = "l00434" ></ a >< span class = "lineno" > 434</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #define hipBlockDim_z (amp_get_local_size(0))</ span ></ div >
< div class = "line" >< a name = "l00435" ></ a >< span class = "lineno" > 435</ span >   < span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00436" ></ a >< span class = "lineno" > 436</ span >   < span class = "preprocessor" > #define hipGridDim_x (hc_get_num_groups(2))</ span ></ div >
< div class = "line" >< a name = "l00437" ></ a >< span class = "lineno" > 437</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #define hipGridDim_y (hc_get_num_groups(1))</ span ></ div >
< div class = "line" >< a name = "l00438" ></ a >< span class = "lineno" > 438</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #define hipGridDim_z (hc_get_num_groups(0))</ span ></ div >
< div class = "line" >< a name = "l00439" ></ a >< span class = "lineno" > 439</ span >   < span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00440" ></ a >< span class = "lineno" > 440</ span >   </ div >
< div class = "line" >< a name = "l00441" ></ a >< span class = "lineno" > 441</ span >   < span class = "keyword" > extern</ span > < span class = "keywordtype" > int</ span > warpSize ;</ div >
< div class = "line" >< a name = "l00442" ></ a >< span class = "lineno" > 442</ span >   </ div >
< div class = "line" >< a name = "l00443" ></ a >< span class = "lineno" > 443</ span >   </ div >
< div class = "line" >< a name = "l00444" ></ a >< span class = "lineno" > 444</ span >   < span class = "preprocessor" > #define __syncthreads() hc_barrier(CLK_LOCAL_MEM_FENCE)</ span ></ div >
< div class = "line" >< a name = "l00445" ></ a >< span class = "lineno" > 445</ span >   < span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00446" ></ a >< span class = "lineno" > 446</ span >   </ div >
< div class = "line" >< a name = "l00447" ></ a >< span class = "lineno" > 447</ span >   < span class = "preprocessor" > #if 0</ span ></ div >
< div class = "line" >< a name = "l00448" ></ a >< span class = "lineno" > 448</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #define KALMAR_PFE_BEGIN() \</ span ></ div >
< div class = "line" >< a name = "l00449" ></ a >< span class = "lineno" > 449</ span >   < span class = "preprocessor" > hc::extent< 3> ext(lp.gridDim.x, lp.gridDim.y, lp.gridDim.z);\</ span ></ div >
< div class = "line" >< a name = "l00450" ></ a >< span class = "lineno" > 450</ span >   < span class = "preprocessor" > auto __hipExtTile = ext.tile(lp.groupDim.x, lp.groupDim.y, lp.groupDim.z);\</ span ></ div >
< div class = "line" >< a name = "l00451" ></ a >< span class = "lineno" > 451</ span >   < span class = "preprocessor" > __hipExtTile.set_dynamic_group_segment_size(lp.groupMemBytes);\</ span ></ div >
< div class = "line" >< a name = "l00452" ></ a >< span class = "lineno" > 452</ span >   < span class = "preprocessor" > \</ span ></ div >
< div class = "line" >< a name = "l00453" ></ a >< span class = "lineno" > 453</ span >   < span class = "preprocessor" > hc::completion_future cf = hc::parallel_for_each (\</ span ></ div >
< div class = "line" >< a name = "l00454" ></ a >< span class = "lineno" > 454</ span >   < span class = "preprocessor" > *lp.av,\</ span ></ div >
< div class = "line" >< a name = "l00455" ></ a >< span class = "lineno" > 455</ span >   < span class = "preprocessor" > __hipExtTile,\</ span ></ div >
< div class = "line" >< a name = "l00456" ></ a >< span class = "lineno" > 456</ span >   < span class = "preprocessor" > [=] (hc::tiled_index< 3> __hipIdx) mutable [[hc]]</ span ></ div >
< div class = "line" >< a name = "l00457" ></ a >< span class = "lineno" > 457</ span >   < span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00458" ></ a >< span class = "lineno" > 458</ span >   </ 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 >   < span class = "preprocessor" > #define KALMAR_PFE_END \</ span ></ div >
< div class = "line" >< a name = "l00461" ></ a >< span class = "lineno" > 461</ span >   < span class = "preprocessor" > ); \</ span ></ div >
< div class = "line" >< a name = "l00462" ></ a >< span class = "lineno" > 462</ span >   < span class = "preprocessor" > if (HIP_LAUNCH_BLOCKING) {\</ span ></ div >
< div class = "line" >< a name = "l00463" ></ a >< span class = "lineno" > 463</ span >   < span class = "preprocessor" > if (HIP_TRACE_API) {\</ span ></ div >
< div class = "line" >< a name = "l00464" ></ a >< span class = "lineno" > 464</ span >   < span class = "preprocessor" > fprintf(stderr, " hiptrace1: HIP_LAUNCH_BLOCKING ...\n" );\</ span ></ div >
< div class = "line" >< a name = "l00465" ></ a >< span class = "lineno" > 465</ span >   < span class = "preprocessor" > }\</ span ></ div >
< div class = "line" >< a name = "l00466" ></ a >< span class = "lineno" > 466</ span >   < span class = "preprocessor" > cf.wait(); \</ span ></ div >
< div class = "line" >< a name = "l00467" ></ a >< span class = "lineno" > 467</ span >   < span class = "preprocessor" > if (HIP_TRACE_API) {\</ span ></ div >
< div class = "line" >< a name = "l00468" ></ a >< span class = "lineno" > 468</ span >   < span class = "preprocessor" > fprintf(stderr, " hiptrace1: ...completed.\n" );\</ span ></ div >
< div class = "line" >< a name = "l00469" ></ a >< span class = "lineno" > 469</ span >   < span class = "preprocessor" > }\</ span ></ div >
< div class = "line" >< a name = "l00470" ></ a >< span class = "lineno" > 470</ span >   < span class = "preprocessor" > }</ span ></ div >
< div class = "line" >< a name = "l00471" ></ a >< span class = "lineno" > 471</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #endif</ span ></ div >
< div class = "line" >< a name = "l00472" ></ a >< span class = "lineno" > 472</ span >   < span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00473" ></ a >< span class = "lineno" > 473</ span >   </ div >
< div class = "line" >< a name = "l00474" ></ a >< span class = "lineno" > 474</ span >   </ div >
< div class = "line" >< a name = "l00475" ></ a >< span class = "lineno" > 475</ span >   < span class = "preprocessor" > #define HIP_KERNEL_NAME(...) __VA_ARGS__</ span ></ div >
< div class = "line" >< a name = "l00476" ></ a >< span class = "lineno" > 476</ span >   < span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00477" ></ a >< span class = "lineno" > 477</ span >   </ div >
< div class = "line" >< a name = "l00478" ></ a >< span class = "lineno" > 478</ span >   < span class = "preprocessor" > #ifdef __HCC_CPP__</ span ></ div >
< div class = "line" >< a name = "l00479" ></ a >< span class = "lineno" > 479</ span >   < span class = "preprocessor" ></ span > hc::accelerator_view *ihipLaunchKernel(< a class = "code" href = "structihipStream__t.html" > hipStream_t</ a > stream);</ 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 >   < span class = "preprocessor" > #if not defined(DISABLE_GRID_LAUNCH)</ span ></ div >
< div class = "line" >< a name = "l00482" ></ a >< span class = "lineno" > 482</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #define hipLaunchKernel(_kernelName, _numBlocks3D, _blockDim3D, _groupMemBytes, _stream, ...) \</ span ></ div >
< div class = "line" >< a name = "l00483" ></ a >< span class = "lineno" > 483</ span >   < span class = "preprocessor" > do {\</ span ></ div >
< div class = "line" >< a name = "l00484" ></ a >< span class = "lineno" > 484</ span >   < span class = "preprocessor" > grid_launch_parm lp;\</ span ></ div >
< div class = "line" >< a name = "l00485" ></ a >< span class = "lineno" > 485</ span >   < span class = "preprocessor" > lp.gridDim.x = _numBlocks3D.x; \</ span ></ div >
< div class = "line" >< a name = "l00486" ></ a >< span class = "lineno" > 486</ span >   < span class = "preprocessor" > lp.gridDim.y = _numBlocks3D.y; \</ span ></ div >
< div class = "line" >< a name = "l00487" ></ a >< span class = "lineno" > 487</ span >   < span class = "preprocessor" > lp.gridDim.z = _numBlocks3D.z; \</ span ></ div >
< div class = "line" >< a name = "l00488" ></ a >< span class = "lineno" > 488</ span >   < span class = "preprocessor" > lp.groupDim.x = _blockDim3D.x; \</ span ></ div >
< div class = "line" >< a name = "l00489" ></ a >< span class = "lineno" > 489</ span >   < span class = "preprocessor" > lp.groupDim.y = _blockDim3D.y; \</ span ></ div >
< div class = "line" >< a name = "l00490" ></ a >< span class = "lineno" > 490</ span >   < span class = "preprocessor" > lp.groupDim.z = _blockDim3D.z; \</ span ></ div >
< div class = "line" >< a name = "l00491" ></ a >< span class = "lineno" > 491</ span >   < span class = "preprocessor" > lp.groupMemBytes = _groupMemBytes;\</ span ></ div >
< div class = "line" >< a name = "l00492" ></ a >< span class = "lineno" > 492</ span >   < span class = "preprocessor" > hc::completion_future cf;\</ span ></ div >
< div class = "line" >< a name = "l00493" ></ a >< span class = "lineno" > 493</ span >   < span class = "preprocessor" > lp.cf = & cf; \</ span ></ div >
< div class = "line" >< a name = "l00494" ></ a >< span class = "lineno" > 494</ span >   < span class = "preprocessor" > lp.av = (ihipLaunchKernel(_stream)); \</ span ></ div >
< div class = "line" >< a name = "l00495" ></ a >< span class = "lineno" > 495</ span >   < span class = "preprocessor" > if (HIP_TRACE_API) {\</ span ></ div >
< div class = "line" >< a name = "l00496" ></ a >< span class = "lineno" > 496</ span >   < span class = "preprocessor" > fprintf(stderr, " hiptrace1: launch ' %s' gridDim:[%d.%d.%d] groupDim:[%d.%d.%d] groupMem:+%d stream=%p\n" , \</ span ></ div >
< div class = "line" >< a name = "l00497" ></ a >< span class = "lineno" > 497</ span >   < span class = "preprocessor" > #_kernelName, lp.gridDim.z, lp.gridDim.y, lp.gridDim.x, lp.groupDim.z, lp.groupDim.y, lp.groupDim.x, lp.groupMemBytes, (void*)(_stream));\</ span ></ div >
< div class = "line" >< a name = "l00498" ></ a >< span class = "lineno" > 498</ span >   < span class = "preprocessor" > }\</ span ></ div >
< div class = "line" >< a name = "l00499" ></ a >< span class = "lineno" > 499</ span >   < span class = "preprocessor" > _kernelName (lp, __VA_ARGS__);\</ span ></ div >
< div class = "line" >< a name = "l00500" ></ a >< span class = "lineno" > 500</ span >   < span class = "preprocessor" > } while(0)</ span ></ div >
< div class = "line" >< a name = "l00501" ></ a >< span class = "lineno" > 501</ span >   < span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00502" ></ a >< span class = "lineno" > 502</ span >   < span class = "preprocessor" > #else</ span ></ div >
< div class = "line" >< a name = "l00503" ></ a >< span class = "lineno" > 503</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #warning(DISABLE_GRID_LAUNCH set)</ span ></ div >
< div class = "line" >< a name = "l00504" ></ a >< span class = "lineno" > 504</ span >   < span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00505" ></ a >< span class = "lineno" > 505</ span >   < span class = "preprocessor" > #define hipLaunchKernel(_kernelName, _numBlocks3D, _blockDim3D, _groupMemBytes, _stream, ...) \</ span ></ div >
< div class = "line" >< a name = "l00506" ></ a >< span class = "lineno" > 506</ span >   < span class = "preprocessor" > do {\</ span ></ div >
< div class = "line" >< a name = "l00507" ></ a >< span class = "lineno" > 507</ span >   < span class = "preprocessor" > grid_launch_parm lp;\</ span ></ div >
< div class = "line" >< a name = "l00508" ></ a >< span class = "lineno" > 508</ span >   < span class = "preprocessor" > lp.gridDim.x = _numBlocks3D.x * _blockDim3D.x;</ span >< span class = "comment" > /*Convert from #blocks to #threads*/</ span >< span class = "preprocessor" > \</ span ></ div >
< div class = "line" >< a name = "l00509" ></ a >< span class = "lineno" > 509</ span >   < span class = "preprocessor" > lp.gridDim.y = _numBlocks3D.y * _blockDim3D.y;</ span >< span class = "comment" > /*Convert from #blocks to #threads*/</ span >< span class = "preprocessor" > \</ span ></ div >
< div class = "line" >< a name = "l00510" ></ a >< span class = "lineno" > 510</ span >   < span class = "preprocessor" > lp.gridDim.z = _numBlocks3D.z * _blockDim3D.z;</ span >< span class = "comment" > /*Convert from #blocks to #threads*/</ span >< span class = "preprocessor" > \</ span ></ div >
< div class = "line" >< a name = "l00511" ></ a >< span class = "lineno" > 511</ span >   < span class = "preprocessor" > lp.groupDim.x = _blockDim3D.x; \</ span ></ div >
< div class = "line" >< a name = "l00512" ></ a >< span class = "lineno" > 512</ span >   < span class = "preprocessor" > lp.groupDim.y = _blockDim3D.y; \</ span ></ div >
< div class = "line" >< a name = "l00513" ></ a >< span class = "lineno" > 513</ span >   < span class = "preprocessor" > lp.groupDim.z = _blockDim3D.z; \</ span ></ div >
< div class = "line" >< a name = "l00514" ></ a >< span class = "lineno" > 514</ span >   < span class = "preprocessor" > lp.groupMemBytes = _groupMemBytes;\</ span ></ div >
< div class = "line" >< a name = "l00515" ></ a >< span class = "lineno" > 515</ span >   < span class = "preprocessor" > hc::completion_future cf;\</ span ></ div >
< div class = "line" >< a name = "l00516" ></ a >< span class = "lineno" > 516</ span >   < span class = "preprocessor" > lp.cf = & cf; \</ span ></ div >
< div class = "line" >< a name = "l00517" ></ a >< span class = "lineno" > 517</ span >   < span class = "preprocessor" > lp.av = (ihipLaunchKernel(_stream)); \</ span ></ div >
< div class = "line" >< a name = "l00518" ></ a >< span class = "lineno" > 518</ span >   < span class = "preprocessor" > if (HIP_TRACE_API) {\</ span ></ div >
< div class = "line" >< a name = "l00519" ></ a >< span class = "lineno" > 519</ span >   < span class = "preprocessor" > fprintf(stderr, " hiptrace1: launch ' %s' gridDim:[%d.%d.%d] groupDim:[%d.%d.%d] groupMem:+%d stream=%p\n" , \</ span ></ div >
< div class = "line" >< a name = "l00520" ></ a >< span class = "lineno" > 520</ span >   < span class = "preprocessor" > #_kernelName, lp.gridDim.z, lp.gridDim.y, lp.gridDim.x, lp.groupDim.z, lp.groupDim.y, lp.groupDim.x, lp.groupMemBytes, (void*)(_stream));\</ span ></ div >
< div class = "line" >< a name = "l00521" ></ a >< span class = "lineno" > 521</ span >   < span class = "preprocessor" > }\</ span ></ div >
< div class = "line" >< a name = "l00522" ></ a >< span class = "lineno" > 522</ span >   < span class = "preprocessor" > _kernelName (lp, __VA_ARGS__);\</ span ></ div >
< div class = "line" >< a name = "l00523" ></ a >< span class = "lineno" > 523</ span >   < span class = "preprocessor" > } while(0)</ span ></ div >
< div class = "line" >< a name = "l00524" ></ a >< span class = "lineno" > 524</ span >   < span class = "preprocessor" ></ span >< span class = "comment" > /*end hipLaunchKernel */</ span ></ div >
< div class = "line" >< a name = "l00525" ></ a >< span class = "lineno" > 525</ span >   < span class = "preprocessor" > #endif</ span ></ div >
< div class = "line" >< a name = "l00526" ></ a >< span class = "lineno" > 526</ span >   < span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00527" ></ a >< span class = "lineno" > 527</ span >   < span class = "preprocessor" > #elif defined (__HCC_C__)</ span ></ div >
< div class = "line" >< a name = "l00528" ></ a >< span class = "lineno" > 528</ span >   < span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00529" ></ a >< span class = "lineno" > 529</ span >   < span class = "comment" > //TODO - develop C interface.</ span ></ div >
< div class = "line" >< a name = "l00530" ></ a >< span class = "lineno" > 530</ span >   </ div >
< div class = "line" >< a name = "l00531" ></ a >< span class = "lineno" > 531</ span >   < span class = "preprocessor" > #endif</ span ></ div >
2016-01-26 20:14:33 -06:00
< div class = "line" >< a name = "l00532" ></ a >< span class = "lineno" > 532</ span >   < span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00533" ></ a >< span class = "lineno" > 533</ span >   </ div >
2016-02-18 20:43:03 -06:00
< div class = "line" >< a name = "l00534" ></ a >< span class = "lineno" > 534</ span >   < span class = "preprocessor" > #if not defined(DISABLE_GRID_LAUNCH)</ span ></ div >
< div class = "line" >< a name = "l00535" ></ a >< span class = "lineno" > 535</ span >   < span class = "preprocessor" ></ span >< span class = "comment" > // TODO -In GL these are no-ops and can be removed:</ span ></ div >
< div class = "line" >< a name = "l00536" ></ a >< span class = "lineno" > 536</ span >   < span class = "comment" > // Keep them around for a little while as a fallback.</ span ></ div >
< div class = "line" >< a name = "l00537" ></ a >< span class = "lineno" > 537</ span >   < span class = "preprocessor" > #define KERNELBEGIN</ span ></ div >
< div class = "line" >< a name = "l00538" ></ a >< span class = "lineno" > 538</ span >   < span class = "preprocessor" ></ span >< span class = "preprocessor" > #define KERNELEND</ span ></ div >
< div class = "line" >< a name = "l00539" ></ a >< span class = "lineno" > 539</ span >   < span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00540" ></ a >< span class = "lineno" > 540</ span >   < span class = "preprocessor" > #else</ span ></ div >
< div class = "line" >< a name = "l00541" ></ a >< span class = "lineno" > 541</ span >   < span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00542" ></ a >< span class = "lineno" > 542</ span >   < span class = "comment" > // TODO-GL:</ span ></ div >
< div class = "line" >< a name = "l00543" ></ a >< span class = "lineno" > 543</ span >   < span class = "comment" > // These wrap the kernel in a PFE loop with macros.</ span ></ div >
< div class = "line" >< a name = "l00544" ></ a >< span class = "lineno" > 544</ span >   < span class = "comment" > // Not required with GL but exist here as a fallback.</ span ></ div >
< div class = "line" >< a name = "l00545" ></ a >< span class = "lineno" > 545</ span >   < span class = "preprocessor" > #define KERNELBEGIN \</ span ></ div >
< div class = "line" >< a name = "l00546" ></ a >< span class = "lineno" > 546</ span >   < span class = "preprocessor" > hc::extent< 3> ext(lp.gridDim.x, lp.gridDim.y, lp.gridDim.z);\</ span ></ div >
< div class = "line" >< a name = "l00547" ></ a >< span class = "lineno" > 547</ span >   < span class = "preprocessor" > auto __hipExtTile = ext.tile(lp.groupDim.x, lp.groupDim.y, lp.groupDim.z);\</ span ></ div >
< div class = "line" >< a name = "l00548" ></ a >< span class = "lineno" > 548</ span >   < span class = "preprocessor" > __hipExtTile.set_dynamic_group_segment_size(lp.groupMemBytes);\</ span ></ div >
< div class = "line" >< a name = "l00549" ></ a >< span class = "lineno" > 549</ span >   < span class = "preprocessor" > \</ span ></ div >
< div class = "line" >< a name = "l00550" ></ a >< span class = "lineno" > 550</ span >   < span class = "preprocessor" > hc::completion_future cf = \</ span ></ div >
< div class = "line" >< a name = "l00551" ></ a >< span class = "lineno" > 551</ span >   < span class = "preprocessor" > hc::parallel_for_each (\</ span ></ div >
< div class = "line" >< a name = "l00552" ></ a >< span class = "lineno" > 552</ span >   < span class = "preprocessor" > *lp.av,\</ span ></ div >
< div class = "line" >< a name = "l00553" ></ a >< span class = "lineno" > 553</ span >   < span class = "preprocessor" > __hipExtTile,\</ span ></ div >
< div class = "line" >< a name = "l00554" ></ a >< span class = "lineno" > 554</ span >   < span class = "preprocessor" > [=] (hc::tiled_index< 3> __hipIdx) mutable [[hc]] \</ span ></ div >
< div class = "line" >< a name = "l00555" ></ a >< span class = "lineno" > 555</ span >   < span class = "preprocessor" > {</ span ></ div >
< div class = "line" >< a name = "l00556" ></ a >< span class = "lineno" > 556</ span >   < span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00557" ></ a >< span class = "lineno" > 557</ span >   </ div >
< div class = "line" >< a name = "l00558" ></ a >< span class = "lineno" > 558</ span >   < span class = "preprocessor" > #define KERNELEND \</ span ></ div >
< div class = "line" >< a name = "l00559" ></ a >< span class = "lineno" > 559</ span >   < span class = "preprocessor" > }); \</ span ></ div >
< div class = "line" >< a name = "l00560" ></ a >< span class = "lineno" > 560</ span >   < span class = "preprocessor" > if (HIP_LAUNCH_BLOCKING) {\</ span ></ div >
< div class = "line" >< a name = "l00561" ></ a >< span class = "lineno" > 561</ span >   < span class = "preprocessor" > if (HIP_TRACE_API) {\</ span ></ div >
< div class = "line" >< a name = "l00562" ></ a >< span class = "lineno" > 562</ span >   < span class = "preprocessor" > fprintf(stderr, " hiptrace1: HIP_LAUNCH_BLOCKING ...\n" );\</ span ></ div >
< div class = "line" >< a name = "l00563" ></ a >< span class = "lineno" > 563</ span >   < span class = "preprocessor" > }\</ span ></ div >
< div class = "line" >< a name = "l00564" ></ a >< span class = "lineno" > 564</ span >   < span class = "preprocessor" > cf.wait(); \</ span ></ div >
< div class = "line" >< a name = "l00565" ></ a >< span class = "lineno" > 565</ span >   < span class = "preprocessor" > if (HIP_TRACE_API) {\</ span ></ div >
< div class = "line" >< a name = "l00566" ></ a >< span class = "lineno" > 566</ span >   < span class = "preprocessor" > fprintf(stderr, " hiptrace1: ...completed.\n" );\</ span ></ div >
< div class = "line" >< a name = "l00567" ></ a >< span class = "lineno" > 567</ span >   < span class = "preprocessor" > }\</ span ></ div >
< div class = "line" >< a name = "l00568" ></ a >< span class = "lineno" > 568</ span >   < span class = "preprocessor" > }</ span ></ div >
< div class = "line" >< a name = "l00569" ></ a >< span class = "lineno" > 569</ span >   < span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00570" ></ a >< span class = "lineno" > 570</ span >   < span class = "preprocessor" > #endif </ span >< span class = "comment" > /*DISABLE_GRID_LAUNCH*/</ span >< span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00571" ></ a >< span class = "lineno" > 571</ span >   < span class = "preprocessor" ></ span ></ 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 = "preprocessor" > #endif // __HCC__</ span ></ div >
< div class = "line" >< a name = "l00574" ></ a >< span class = "lineno" > 574</ span >   < span class = "preprocessor" ></ span ></ div >
< div class = "line" >< a name = "l00575" ></ a >< span class = "lineno" > 575</ span >   </ div >
< div class = "line" >< a name = "l00580" ></ a >< span class = "lineno" > 580</ span >   < span class = "keyword" > extern</ span > < span class = "keywordtype" > int</ span > < a class = "code" href = "group__HIP-ENV.html#ga1e1c85dbb250f1acfb484c1be1f3b28a" > HIP_PRINT_ENV</ a > ; </ div >
< div class = "line" >< a name = "l00581" ></ a >< span class = "lineno" > 581</ span >   < span class = "keyword" > extern</ span > < span class = "keywordtype" > int</ span > < a class = "code" href = "group__HIP-ENV.html#gaae9c541f3e25b8f002762337a03fec28" > HIP_TRACE_API</ a > ; </ div >
< div class = "line" >< a name = "l00582" ></ a >< span class = "lineno" > 582</ span >   < span class = "keyword" > extern</ span > < span class = "keywordtype" > int</ span > < a class = "code" href = "group__HIP-ENV.html#ga8049b329f2663b4572d81e7a9aa8a155" > HIP_LAUNCH_BLOCKING</ a > ; </ div >
< div class = "line" >< a name = "l00583" ></ a >< span class = "lineno" > 583</ span >   </ div >
< div class = "line" >< a name = "l00589" ></ a >< span class = "lineno" > 589</ span >   < span class = "comment" > // End doxygen API:</ span ></ div >
2016-01-26 20:14:33 -06:00
< div class = "ttc" id = "group__HIP-ENV_html_gaae9c541f3e25b8f002762337a03fec28" >< div class = "ttname" >< a href = "group__HIP-ENV.html#gaae9c541f3e25b8f002762337a03fec28" > HIP_TRACE_API</ a ></ div >< div class = "ttdeci" > int HIP_TRACE_API</ div >< div class = "ttdoc" > Trace HIP APIs. </ div >< div class = "ttdef" >< b > Definition:</ b > hip_hcc.cpp:57</ div ></ div >
2016-02-18 20:43:03 -06:00
< div class = "ttc" id = "host__defines_8h_html" >< div class = "ttname" >< a href = "host__defines_8h.html" > host_defines.h</ a ></ div >< div class = "ttdoc" > TODO-doc. </ div ></ div >
2016-01-26 20:14:33 -06:00
< div class = "ttc" id = "structihipStream__t_html" >< div class = "ttname" >< a href = "structihipStream__t.html" > ihipStream_t</ a ></ div >< div class = "ttdef" >< b > Definition:</ b > hip_hcc.cpp:82</ div ></ div >
2016-02-18 20:43:03 -06:00
< div class = "ttc" id = "hip__texture_8h_html" >< div class = "ttname" >< a href = "hip__texture_8h.html" > hip_texture.h</ a ></ div >< div class = "ttdoc" > HIP C++ Texture API for hcc compiler. </ div ></ div >
2016-01-26 20:14:33 -06:00
< div class = "ttc" id = "group__HIP-ENV_html_ga1e1c85dbb250f1acfb484c1be1f3b28a" >< div class = "ttname" >< a href = "group__HIP-ENV.html#ga1e1c85dbb250f1acfb484c1be1f3b28a" > HIP_PRINT_ENV</ a ></ div >< div class = "ttdeci" > int HIP_PRINT_ENV</ div >< div class = "ttdoc" > Print all HIP-related environment variables. </ div >< div class = "ttdef" >< b > Definition:</ b > hip_hcc.cpp:56</ div ></ div >
2016-02-18 20:43:03 -06:00
< div class = "ttc" id = "hcc__detail_2hip__runtime__api_8h_html" >< div class = "ttname" >< a href = "hcc__detail_2hip__runtime__api_8h.html" > hip_runtime_api.h</ a ></ div >< div class = "ttdoc" > Contains C function APIs for HIP runtime. This file does not use any HCC builtin or special language ...</ div ></ div >
2016-01-26 20:14:33 -06:00
< div class = "ttc" id = "group__HIP-ENV_html_ga8049b329f2663b4572d81e7a9aa8a155" >< div class = "ttname" >< a href = "group__HIP-ENV.html#ga8049b329f2663b4572d81e7a9aa8a155" > HIP_LAUNCH_BLOCKING</ a ></ div >< div class = "ttdeci" > int HIP_LAUNCH_BLOCKING</ div >< div class = "ttdoc" > Make all HIP APIs host-synchronous. </ div >< div class = "ttdef" >< b > Definition:</ b > hip_hcc.cpp:58</ div ></ div >
</ div > <!-- fragment --> </ div > <!-- contents -->
<!-- start footer part -->
< hr class = "footer" />< address class = "footer" >< small >
2016-02-18 20:43:03 -06:00
Generated on Thu Feb 18 2016 20:42:16 for HIP: Heterogenous-computing Interface for Portability by   < a href = "http://www.doxygen.org/index.html" >
2016-01-26 20:14:33 -06:00
< img class = "footer" src = "doxygen.png" alt = "doxygen" />
</ a > 1.8.6
</ small ></ address >
</ body >
</ html >
