NVBIO
 All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros Groups Pages
Functions | Variables
nvbio::cuda::wq Namespace Reference

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()