MappingDetail

Detailed Description

Classes
struct	nvbio::bowtie2::cuda::detail::seed_mapper< ALG >
	abstract seed-mapper interface More...
struct	nvbio::bowtie2::cuda::detail::seed_mapper< EXACT_MAPPING >
struct	nvbio::bowtie2::cuda::detail::seed_mapper< APPROX_MAPPING >
struct	nvbio::bowtie2::cuda::detail::seed_mapper< CASE_PRUNING_MAPPING >

Enumerations
enum	{ nvbio::bowtie2::cuda::detail::CHECK_EXACT =1u, nvbio::bowtie2::cuda::detail::IGNORE_EXACT =0u }
enum	nvbio::bowtie2::cuda::detail::MappingAlgorithm { nvbio::bowtie2::cuda::detail::EXACT_MAPPING = 0u, nvbio::bowtie2::cuda::detail::APPROX_MAPPING = 1u, nvbio::bowtie2::cuda::detail::CASE_PRUNING_MAPPING = 2u }

Functions
template<typename ReadStream >
NVBIO_DEVICE NVBIO_FORCEINLINE bool	nvbio::bowtie2::cuda::detail::check_N (const ReadStream &seed, const uint32 exact_len, const uint32 seed_len)
template<typename FMType , typename StreamType >
NVBIO_DEVICE NVBIO_FORCEINLINE uint2	nvbio::bowtie2::cuda::detail::match_range (uint2 range, const FMType index, StreamType query, uint32 begin, uint32 end)
NVBIO_DEVICE NVBIO_FORCEINLINE void	nvbio::bowtie2::cuda::detail::store_deque (SeedHitDequeArrayDeviceView hit_deques, const uint32 read_id, const uint32 n_hits, const SeedHit *hitstorage)
template<bool find_exact, typename Stream , typename FMType , typename HitType >
NVBIO_DEVICE NVBIO_FORCEINLINE void	nvbio::bowtie2::cuda::detail::map (const Stream query, uint32 len1, const uint32 len2, const FMType index, const SeedHit::Flags hit_flags, HitType &hitheap, uint32 max_hits, uint32 &range_sum, uint32 &range_count)
template<typename BatchType , typename FMType , typename rFMType >
__global__ void	nvbio::bowtie2::cuda::detail::map_whole_read_kernel (const BatchType read_batch, const FMType fmi, const rFMType rfmi, const nvbio::cuda::PingPongQueuesView< uint32 > queues, uint8 *reseed, SeedHitDequeArrayDeviceView hits, const ParamsPOD params, const bool fw, const bool rc)
template<MappingAlgorithm ALGO, typename BatchType , typename FMType , typename rFMType >
__global__ void	nvbio::bowtie2::cuda::detail::map_queues_kernel (const BatchType read_batch, const FMType fmi, const rFMType rfmi, const uint32 retry, const nvbio::cuda::PingPongQueuesView< uint32 > queues, uint8 *reseed, SeedHitDequeArrayDeviceView hits, const ParamsPOD params, const bool fw, const bool rc)
template<MappingAlgorithm ALGO, typename BatchType , typename FMType , typename rFMType >
__global__ void	nvbio::bowtie2::cuda::detail::map_kernel (const BatchType read_batch, const FMType fmi, const rFMType rfmi, SeedHitDequeArrayDeviceView hits, const uint2 seed_range, const ParamsPOD params, const bool fw, const bool rc)

Enumeration Type Documentation

anonymous enum

Enumerator
CHECK_EXACT
IGNORE_EXACT

Definition at line 117 of file mapping_inl.h.

enum nvbio::bowtie2::cuda::detail::MappingAlgorithm

Enumerator
EXACT_MAPPING
APPROX_MAPPING
CASE_PRUNING_MAPPING

Definition at line 118 of file mapping_inl.h.

Function Documentation

template<typename ReadStream >

NVBIO_DEVICE NVBIO_FORCEINLINE bool nvbio::bowtie2::cuda::detail::check_N	(	const ReadStream &	seed,
		const uint32	exact_len,
		const uint32	seed_len
	)

Definition at line 66 of file mapping_inl.h.

template<bool find_exact, typename Stream , typename FMType , typename HitType >

NVBIO_DEVICE NVBIO_FORCEINLINE void nvbio::bowtie2::cuda::detail::map	(	const Stream	query,
		uint32	len1,
		const uint32	len2,
		const FMType	index,
		const SeedHit::Flags	hit_flags,
		HitType &	hitheap,
		uint32	max_hits,
		uint32 &	range_sum,
		uint32 &	range_count
	)

This is the guts of our mapper. We look for hits that match exactly in the range [0, len1) and have one mismatch in [len1, len2). If find_exact is set, we also check for a perfect match.

Definition at line 132 of file mapping_inl.h.

template<MappingAlgorithm ALGO, typename BatchType , typename FMType , typename rFMType >

__global__ void nvbio::bowtie2::cuda::detail::map_kernel	(	const BatchType	read_batch,
		const FMType	fmi,
		const rFMType	rfmi,
		SeedHitDequeArrayDeviceView	hits,
		const uint2	seed_range,
		const ParamsPOD	params,
		const bool	fw,
		const bool	rc
	)

Seed filtering kernel. Maps the set of seeds for each read in the input queue, filling the corresponding hit priority-deque. Internally performs re-seeding to find a good seeding scheme.

Definition at line 600 of file mapping_inl.h.

template<MappingAlgorithm ALGO, typename BatchType , typename FMType , typename rFMType >

__global__ void nvbio::bowtie2::cuda::detail::map_queues_kernel	(	const BatchType	read_batch,
		const FMType	fmi,
		const rFMType	rfmi,
		const uint32	retry,
		const nvbio::cuda::PingPongQueuesView< uint32 >	queues,
		uint8 *	reseed,
		SeedHitDequeArrayDeviceView	hits,
		const ParamsPOD	params,
		const bool	fw,
		const bool	rc
	)

Seed filtering kernel. Maps the set of seeds for each read in the input queue, filling the corresponding hit priority-deque, and fills the output queue with the list of reads that need reseeding.

Definition at line 512 of file mapping_inl.h.

template<typename BatchType , typename FMType , typename rFMType >

__global__ void nvbio::bowtie2::cuda::detail::map_whole_read_kernel	(	const BatchType	read_batch,
		const FMType	fmi,
		const rFMType	rfmi,
		const nvbio::cuda::PingPongQueuesView< uint32 >	queues,
		uint8 *	reseed,
		SeedHitDequeArrayDeviceView	hits,
		const ParamsPOD	params,
		const bool	fw,
		const bool	rc
	)

Exact whole read filtering kernel. Maps the entire read exactly

Definition at line 435 of file mapping_inl.h.

template<typename FMType , typename StreamType >

NVBIO_DEVICE NVBIO_FORCEINLINE uint2 nvbio::bowtie2::cuda::detail::match_range	(	uint2	range,
		const FMType	index,
		StreamType	query,
		uint32	begin,
		uint32	end
	)

Definition at line 84 of file mapping_inl.h.

NVBIO_DEVICE NVBIO_FORCEINLINE void nvbio::bowtie2::cuda::detail::store_deque	(	SeedHitDequeArrayDeviceView	hit_deques,
		const uint32	read_id,
		const uint32	n_hits,
		const SeedHit *	hitstorage
	)

Definition at line 102 of file mapping_inl.h.