NVBIO
|
Functions | |
template<uint32 BLOCKDIM, typename WorkUnitT , typename WorkStreamT > | |
__global__ void | inplace_work_queue_kernel (const WorkStreamT stream, WorkQueueStats::View stats) |
template<uint32 BLOCKDIM, typename WorkUnitT , typename WorkStreamT , bool DO_LOADS> | |
__global__ void | mk_work_queue_kernel (const uint32 n_tile_grids, typename WorkQueue< MultiPassQueueTag, WorkUnitT, BLOCKDIM >::Context context, const uint32 in_queue_id, const uint32 in_queue_size, const WorkStreamT stream, uint32 stream_begin) |
template<uint32 BLOCKDIM, typename WorkUnitT , typename WorkStreamT > | |
__global__ void | mk_load_kernel (const uint32 n_tile_grids, typename WorkQueue< MultiPassQueueTag, WorkUnitT, BLOCKDIM >::Context context, const uint32 in_queue_id, const uint32 in_queue_size, const WorkStreamT stream, uint32 stream_begin) |
template<uint32 BLOCKDIM, typename WorkUnitT , typename WorkStream , typename WorkMover > | |
__global__ void | mk_move_kernel (const uint32 n_tile_grids, typename WorkQueue< MultiPassQueueTag, WorkUnitT, BLOCKDIM >::Context context, const uint32 in_queue_id, const uint32 in_queue_size, const WorkStream stream, const WorkMover mover) |
template<uint32 BLOCKDIM, typename WorkUnitT , typename WorkStreamT , typename WorkMoverT > | |
__global__ void | work_queue_kernel (const uint32 n_tile_grids, typename WorkQueue< OrderedQueueTag, WorkUnitT, BLOCKDIM >::Context context, const WorkStreamT stream, const WorkMoverT mover) |
template<uint32 BLOCKDIM, typename WorkUnitT , typename WorkStreamT > | |
__global__ void | persistent_warps_work_queue_kernel (uint32 *pool, const WorkStreamT stream, WorkQueueStats::View stats) |
template<uint32 BLOCKDIM, typename WorkUnitT , typename WorkStreamT > | |
__global__ void | __launch_bounds__ (BLOCKDIM, 6) persistent_threads_work_queue_kernel(uint32 *pool |
while (__any(active)) | |
Variables | |
__global__ void const uint32 | max_inactive_lanes |
__global__ void const uint32 const WorkStreamT | stream |
__global__ void const uint32 const WorkStreamT WorkQueueStats::View | stats |
const uint32 | thread_id = threadIdx.x + blockIdx.x*BLOCKDIM |
WorkUnit | unit |
const uint32 | NUM_WARPS = BLOCKDIM >> cuda::Arch::LOG_WARP_SIZE |
__shared__ volatile uint32 | sm_broadcast [NUM_WARPS] |
const uint32 | invalid_unit = uint32(-1) |
uint32 | work_id = invalid_unit |
uint32 | work_iter = 0u |
bool | active = true |
const uint32 | stream_size = stream.size() |