29 uint32_t getGridDim(uint32_t n, uint16_t d);
30 uint32_t getWorkgroupDim(uint32_t group_id, uint32_t grid_size,
32 uint32_t getNumHardwareThreadsInBlock();
36 uint32_t getThreadIdInWarp();
37 uint32_t getThreadIdInBlock();
38 uint32_t getKernelSize();
39 uint32_t getBlockId();
40 uint32_t getNumberOfBlocks();
41 uint32_t getWarpId();
42 uint32_t getNumberOfWarpsInBlock();
53 uint32_t getGridDim(uint32_t n, uint16_t d) {  in getGridDim()
54   uint32_t q = n / d;  in getGridDim()
58 uint32_t getWorkgroupDim(uint32_t group_id, uint32_t grid_size,  in getWorkgroupDim()
60   uint32_t r = grid_size - group_id * group_size;  in getWorkgroupDim()
64 uint32_t getNumHardwareThreadsInBlock() {  in getNumHardwareThreadsInBlock()
73   uint32_t Lane = mapping::getThreadIdInWarp();  in lanemaskLT()
80   uint32_t Lane = mapping::getThreadIdInWarp();  in lanemaskGT()
88 uint32_t getThreadIdInWarp() {  in getThreadIdInWarp()
92 uint32_t getThreadIdInBlock() { return __builtin_amdgcn_workitem_id_x(); }  in getThreadIdInBlock()
94 uint32_t getKernelSize() { return __builtin_amdgcn_grid_size_x(); }  in getKernelSize()
96 uint32_t getBlockId() { return __builtin_amdgcn_workgroup_id_x(); }  in getBlockId()
98 uint32_t getNumberOfBlocks() {  in getNumberOfBlocks()
103 uint32_t getWarpId() {  in getWarpId()
107 uint32_t getNumberOfWarpsInBlock() {  in getNumberOfWarpsInBlock()
120 uint32_t getNumHardwareThreadsInBlock() {  in getNumHardwareThreadsInBlock()
144 uint32_t getThreadIdInBlock() { return __nvvm_read_ptx_sreg_tid_x(); }  in getThreadIdInBlock()
146 uint32_t getThreadIdInWarp() {  in getThreadIdInWarp()
150 uint32_t getKernelSize() {  in getKernelSize()
155 uint32_t getBlockId() { return __nvvm_read_ptx_sreg_ctaid_x(); }  in getBlockId()
157 uint32_t getNumberOfBlocks() { return __nvvm_read_ptx_sreg_nctaid_x(); }  in getNumberOfBlocks()
159 uint32_t getWarpId() {  in getWarpId()
163 uint32_t getNumberOfWarpsInBlock() {  in getNumberOfWarpsInBlock()
171 uint32_t getWarpSize() { return getGridValue().GV_Warp_Size; }  in getWarpSize()
181   uint32_t MainTId = (mapping::getNumberOfProcessorElements() - 1) &  in isInLastWarp()
216 uint32_t mapping::getThreadIdInWarp() {  in getThreadIdInWarp()
217   uint32_t ThreadIdInWarp = impl::getThreadIdInWarp();  in getThreadIdInWarp()
222 uint32_t mapping::getThreadIdInBlock() {  in getThreadIdInBlock()
223   uint32_t ThreadIdInBlock = impl::getThreadIdInBlock();  in getThreadIdInBlock()
228 uint32_t mapping::getWarpSize() { return impl::getWarpSize(); }  in getWarpSize()
230 uint32_t mapping::getBlockSize(bool IsSPMD) {  in getBlockSize()
231   uint32_t BlockSize = mapping::getNumberOfProcessorElements() -  in getBlockSize()
235 uint32_t mapping::getBlockSize() {  in getBlockSize()
239 uint32_t mapping::getKernelSize() { return impl::getKernelSize(); }  in getKernelSize()
241 uint32_t mapping::getWarpId() {  in getWarpId()
242   uint32_t WarpID = impl::getWarpId();  in getWarpId()
247 uint32_t mapping::getBlockId() {  in getBlockId()
248   uint32_t BlockId = impl::getBlockId();  in getBlockId()
253 uint32_t mapping::getNumberOfWarpsInBlock() {  in getNumberOfWarpsInBlock()
254   uint32_t NumberOfWarpsInBlocks = impl::getNumberOfWarpsInBlock();  in getNumberOfWarpsInBlock()
259 uint32_t mapping::getNumberOfBlocks() {  in getNumberOfBlocks()
260   uint32_t NumberOfBlocks = impl::getNumberOfBlocks();  in getNumberOfBlocks()
265 uint32_t mapping::getNumberOfProcessorElements() {  in getNumberOfProcessorElements()
266   uint32_t NumberOfProcessorElements = impl::getNumHardwareThreadsInBlock();  in getNumberOfProcessorElements()
292 __attribute__((noinline)) uint32_t __kmpc_get_hardware_thread_id_in_block() {  in __kmpc_get_hardware_thread_id_in_block()
297 __attribute__((noinline)) uint32_t __kmpc_get_hardware_num_threads_in_block() {  in __kmpc_get_hardware_num_threads_in_block()
302 __attribute__((noinline)) uint32_t __kmpc_get_warp_size() {  in __kmpc_get_warp_size()