@hsa.jit(device=True)
def shuf_device_inclusive_scan(data, temp):
 tid = hsa.get_local_id(0)
 lane = (tid & (_WARPSIZE - 1))
 warpid = (tid >> 6)
 warp_scan_res = shuf_wave_inclusive_scan(data)
 hsa.barrier()
 if (lane == (_WARPSIZE - 1)):
  temp[warpid] = warp_scan_res
 hsa.barrier()
 if (warpid == 0):
  shuf_wave_inclusive_scan(temp[lane])
 hsa.barrier()
 blocksum = 0
 if (warpid > 0):
  blocksum = temp[(warpid - 1)]
 return (warp_scan_res + blocksum)

@ hsa . jit ( device = True ) def shuf_device_inclusive_scan ( data , temp ) : tid = hsa . get_local_id ( 0 ) lane = ( tid & ( _WARPSIZE - 1 ) ) warpid = ( tid >> 6 ) warp_scan_res = shuf_wave_inclusive_scan ( data ) hsa . barrier ( ) if ( lane == ( _WARPSIZE - 1 ) ) : temp [ warpid ] = warp_scan_res hsa . barrier ( ) if ( warpid == 0 ) : shuf_wave_inclusive_scan ( temp [ lane ] ) hsa . barrier ( ) blocksum = 0 if ( warpid > 0 ) : blocksum = temp [ ( warpid - 1 ) ] return ( warp_scan_res + blocksum )
args data: scalar input for tid temp: shared memory for temporary work .