NVBIO
 All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros Groups Pages
Namespaces | Functions | Variables
work_queue_persistent_inl.h File Reference
#include <nvbio/basic/types.h>
#include <nvbio/basic/numbers.h>
#include <nvbio/basic/cuda/arch.h>
#include <cub/cub.cuh>
#include <thrust/copy.h>

Go to the source code of this file.

Namespaces

 nvbio
 Define a vector_view POD type and plain_view() for std::vector.
 
 nvbio::cuda
 
 nvbio::cuda::wq
 

Functions

template<uint32 BLOCKDIM, typename WorkUnitT , typename WorkStreamT >
__global__ void nvbio::cuda::wq::persistent_warps_work_queue_kernel (uint32 *pool, const WorkStreamT stream, WorkQueueStats::View stats)
 
template<uint32 BLOCKDIM, typename WorkUnitT , typename WorkStreamT >
__global__ void nvbio::cuda::wq::__launch_bounds__ (BLOCKDIM, 6) persistent_threads_work_queue_kernel(uint32 *pool
 
 nvbio::cuda::wq::while (__any(active))
 

Variables

__global__ void const uint32 nvbio::cuda::wq::max_inactive_lanes
 
__global__ void const uint32
const WorkStreamT 
nvbio::cuda::wq::stream
 
__global__ void const uint32
const WorkStreamT
WorkQueueStats::View 
nvbio::cuda::wq::stats
 
const uint32 nvbio::cuda::wq::thread_id = threadIdx.x + blockIdx.x*BLOCKDIM
 
WorkUnit nvbio::cuda::wq::unit
 
const uint32 nvbio::cuda::wq::NUM_WARPS = BLOCKDIM >> cuda::Arch::LOG_WARP_SIZE
 
__shared__ volatile uint32 nvbio::cuda::wq::sm_broadcast [NUM_WARPS]
 
const uint32 nvbio::cuda::wq::invalid_unit = uint32(-1)
 
uint32 nvbio::cuda::wq::work_id = invalid_unit
 
uint32 nvbio::cuda::wq::work_iter = 0u
 
bool nvbio::cuda::wq::active = true
 
const uint32 nvbio::cuda::wq::stream_size = stream.size()