NVBIO
 All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros Groups Pages
scoring_queues.h
Go to the documentation of this file.
1 /*
2  * nvbio
3  * Copyright (c) 2011-2014, NVIDIA CORPORATION. All rights reserved.
4  *
5  * Redistribution and use in source and binary forms, with or without
6  * modification, are permitted provided that the following conditions are met:
7  * * Redistributions of source code must retain the above copyright
8  * notice, this list of conditions and the following disclaimer.
9  * * Redistributions in binary form must reproduce the above copyright
10  * notice, this list of conditions and the following disclaimer in the
11  * documentation and/or other materials provided with the distribution.
12  * * Neither the name of the NVIDIA CORPORATION nor the
13  * names of its contributors may be used to endorse or promote products
14  * derived from this software without specific prior written permission.
15  *
16  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
17  * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
18  * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
19  * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
20  * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
21  * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
22  * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
23  * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
24  * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
25  * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
26  */
27 
31 
32 #pragma once
33 
34 #include <nvBowtie/bowtie2/cuda/defs.h>
35 #include <nvbio/basic/cuda/pingpong_queues.h>
36 #include <nvbio/basic/cuda/arch.h>
37 #include <nvbio/basic/thrust_view.h>
38 #include <algorithm>
39 
40 namespace nvbio {
41 namespace bowtie2 {
42 namespace cuda {
43 
46 
60 
63 
64 // forward declaration
65 template <typename HitQueuesType> struct HitReference;
66 template <typename ScoringQueuesType> struct ReadHitReference;
67 template <typename ScoringQueuesType> struct ReadHitBinder;
68 
69 struct HitQueuesDeviceView;
70 struct ScoringQueuesDeviceView;
71 struct ReadHitsIndexDeviceView;
72 
99 struct ReadHitsIndex
100 {
101  typedef thrust::device_vector<uint32> links_storage_type;
102  typedef ReadHitsIndexDeviceView device_view_type;
103 
104  enum Mode {
105  SingleHitPerRead = 0u,
106  MultipleHitsPerRead = 1u,
107  };
108 
111  uint64 resize(const uint32 max_size, const bool do_alloc);
112 
115  void setup(const uint32 n_hits_per_read, const uint32 in_reads);
116 
119  Mode mode() const { return m_mode; }
120 
123  ReadHitsIndexDeviceView device_view();
124 
125 public:
126  links_storage_type m_links;
127  uint32 m_stride;
128  Mode m_mode;
129 };
130 
135 struct ReadHitsIndexDeviceView
136 {
137  typedef typename device_view_subtype<ReadHitsIndex::links_storage_type>::type links_storage_type;
138 
142  struct HitArray
143  {
146  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
147  HitArray(ReadHitsIndexDeviceView& index, const uint32 read_index) : m_index(index), m_read_index(read_index) {}
148 
151  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
152  void resize(const uint32 size) const { m_index.set_hit_count( m_read_index, size ); }
153 
156  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
157  uint32 size() const { return m_index.hit_count( m_read_index ); }
158 
161  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
162  uint32& operator[] (const uint32 slot) { return m_index( m_read_index, slot ); }
163 
166  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
167  uint32 operator[] (const uint32 slot) const { return m_index( m_read_index, slot ); }
168 
169  private:
170  ReadHitsIndexDeviceView& m_index;
171  uint32 m_read_index;
172  };
173 
174  // define the reference type
175  typedef HitArray reference;
176 
177 
180  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
181  ReadHitsIndexDeviceView(
182  links_storage_type links = links_storage_type(),
183  const uint32 stride = 0u) : m_links( links ), m_stride( stride ) {}
184 
187  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
188  uint32 hit_count(const uint32 read_index) const { return m_links[ read_index ]; }
189 
192  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
193  void set_hit_count(const uint32 read_index, const uint32 count) { m_links[ read_index ] = count; }
194 
197  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
198  uint32& operator() (const uint32 read_index, const uint32 slot) { return m_links[ read_index + (1u + slot) * m_stride ]; }
199 
202  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
203  uint32 operator() (const uint32 read_index, const uint32 slot) const { return m_links[ read_index + (1u + slot) * m_stride ]; }
204 
207  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
208  HitArray operator[] (const uint32 read_index) { return HitArray( *this, read_index ); }
209 
210 public:
211  links_storage_type m_links;
212  uint32 m_stride;
213 
214  friend struct HitArray;
215 };
216 
220 struct HitQueues
221 {
222  typedef thrust::device_vector<uint32> index_storage_type;
223  typedef thrust::device_vector<packed_seed> seed_storage_type;
224  typedef thrust::device_vector<uint32> ssa_storage_type;
225  typedef thrust::device_vector<uint32> loc_storage_type;
226  typedef thrust::device_vector<int32> score_storage_type;
227  typedef thrust::device_vector<uint32> sink_storage_type;
228 
229  typedef HitQueuesDeviceView device_view_type;
230  typedef HitReference<HitQueues> reference;
231 
234  uint64 resize(const uint32 size, const bool do_alloc);
235 
238  HitQueuesDeviceView device_view();
239 
240  index_storage_type read_id;
241  seed_storage_type seed;
242  ssa_storage_type ssa;
243  loc_storage_type loc;
244  score_storage_type score;
245  sink_storage_type sink;
246  loc_storage_type opposite_loc;
247  score_storage_type opposite_score;
248  sink_storage_type opposite_sink;
249  score_storage_type opposite_score2;
250  sink_storage_type opposite_sink2;
251 };
255 struct HitQueuesDeviceView
256 {
257  typedef typename device_view_subtype<HitQueues::index_storage_type>::type index_storage_type;
258  typedef typename device_view_subtype<HitQueues::seed_storage_type>::type seed_storage_type;
259  typedef typename device_view_subtype<HitQueues::ssa_storage_type>::type ssa_storage_type;
260  typedef typename device_view_subtype<HitQueues::loc_storage_type>::type loc_storage_type;
261  typedef typename device_view_subtype<HitQueues::score_storage_type>::type score_storage_type;
262  typedef typename device_view_subtype<HitQueues::sink_storage_type>::type sink_storage_type;
263 
264  typedef HitQueuesDeviceView device_view_type;
265  typedef HitReference<HitQueuesDeviceView> reference;
266 
269  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
270  HitQueuesDeviceView(
271  index_storage_type _read_id = index_storage_type(), // hit -> read mapping
272  seed_storage_type _seed = seed_storage_type(), // hit info
273  ssa_storage_type _ssa = ssa_storage_type(), // hit ssa info
274  loc_storage_type _loc = loc_storage_type(), // hit locations
275  score_storage_type _score = score_storage_type(), // hit scores
276  sink_storage_type _sink = sink_storage_type(), // hit sinks
277  loc_storage_type _opposite_loc = loc_storage_type(), // hit locations, opposite mate
278  score_storage_type _opposite_score = score_storage_type(), // hit scores, opposite mate
279  sink_storage_type _opposite_sink = sink_storage_type(), // hit sinks, opposite mate
280  score_storage_type _opposite_score2 = score_storage_type(), // hit scores, opposite mate
281  sink_storage_type _opposite_sink2 = sink_storage_type()); // hit sinks, opposite mate
282 
285  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
286  HitReference<HitQueuesDeviceView> operator[] (const uint32 i);
287 
290  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
291  HitReference<HitQueuesDeviceView> operator[] (const uint32 i) const;
292 
293  index_storage_type read_id;
294  seed_storage_type seed;
295  ssa_storage_type ssa;
296  loc_storage_type loc;
297  score_storage_type score;
298  sink_storage_type sink;
299  loc_storage_type opposite_loc;
300  score_storage_type opposite_score;
301  sink_storage_type opposite_sink;
302  score_storage_type opposite_score2;
303  sink_storage_type opposite_sink2;
304 };
305 
335 struct ScoringQueues
336 {
337  typedef nvbio::cuda::PingPongQueues<packed_read> active_reads_storage_type;
338  typedef HitQueues hits_storage_type;
339  typedef ReadHitsIndex read_hits_index_type;
340  typedef thrust::device_vector<uint32> pool_type;
341 
342  typedef ScoringQueuesDeviceView device_view_type;
343 
346  ScoringQueues() : hits_pool(1) {}
347 
350  uint64 resize(const uint32 n_reads, const uint32 n_hits, const bool do_alloc);
351 
354  void clear_output() { hits_pool[0] = 0u; active_reads.clear_output(); }
355 
358  uint32 hits_count() { return hits_pool[0]; }
359 
362  ScoringQueuesDeviceView device_view();
363 
364  active_reads_storage_type active_reads;
365  hits_storage_type hits;
366  read_hits_index_type hits_index;
367  pool_type hits_pool;
368 };
372 struct ScoringQueuesDeviceView
373 {
374  typedef typename device_view_subtype<ScoringQueues::active_reads_storage_type>::type active_reads_storage_type;
375  typedef typename device_view_subtype<ScoringQueues::hits_storage_type>::type hits_storage_type;
376  typedef typename device_view_subtype<ScoringQueues::read_hits_index_type>::type read_hits_index_type;
377  typedef typename device_view_subtype<ScoringQueues::pool_type>::type pool_type;
378 
381  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
382  ScoringQueuesDeviceView() {}
383 
386  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
387  ScoringQueuesDeviceView(
388  active_reads_storage_type _active_reads, // set of active reads
389  hits_storage_type _hits, // nested sequence of read hits
390  read_hits_index_type _hits_index, // read -> hits mapping
391  pool_type _hits_pool); // pool counter
392 
395  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
396  uint32 active_read_count() const { return active_reads.in_size; }
397 
400  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
401  packed_read active_read(const uint32 read_index) const { return active_reads.in_queue[read_index]; }
402 
403  active_reads_storage_type active_reads;
404  hits_storage_type hits;
405  read_hits_index_type hits_index;
406  pool_type hits_pool;
407 };
408 
414 template <typename HitQueuesType>
415 struct HitReference
416 {
417  typedef typename reference_subtype<typename HitQueuesType::index_storage_type>::type index_type;
418  typedef typename reference_subtype<typename HitQueuesType::seed_storage_type>::type seed_type;
419  typedef typename reference_subtype<typename HitQueuesType::ssa_storage_type>::type ssa_type;
420  typedef typename reference_subtype<typename HitQueuesType::loc_storage_type>::type loc_type;
421  typedef typename reference_subtype<typename HitQueuesType::score_storage_type>::type score_type;
422  typedef typename reference_subtype<typename HitQueuesType::sink_storage_type>::type sink_type;
423 
428  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
429  HitReference(HitQueuesType& hits, const uint32 hit_index);
430 
431  index_type read_id;
432  seed_type seed;
433  ssa_type ssa;
434  loc_type loc;
435  score_type score;
436  sink_type sink;
437  loc_type opposite_loc;
438  score_type opposite_score;
439  sink_type opposite_sink;
440  score_type opposite_score2;
441  sink_type opposite_sink2;
442 };
443 
447 template <typename ScoringQueuesType>
448 struct ReadHitsReference
449 {
450  typedef typename ScoringQueuesType::read_hits_index_type read_hits_index_type;
451  typedef typename ScoringQueuesType::hits_storage_type hits_storage_type;
452  typedef HitReference<hits_storage_type> reference;
453 
459  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
460  ReadHitsReference(ScoringQueuesType& queues, const uint32 read_index = uint32(-1));
461 
465  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
466  void bind(const uint32 read_index);
467 
470  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
471  uint32 size() const;
472 
476  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
477  HitReference<hits_storage_type> operator[] (const uint32 i) const;
478 
481  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
482  packed_read read_info() const;
483 
486  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
487  uint32 slot(const uint32 i) const;
488 
489 private:
490  ScoringQueuesType& m_queues;
491  uint32 m_read_index;
492 };
493 
494 
498 template <typename ScoringQueuesType>
499 struct ReadHitsBinder
500 {
501  typedef typename ScoringQueuesType::read_hits_index_type read_hits_index_type;
502  typedef typename ScoringQueuesType::hits_storage_type hits_storage_type;
503  typedef HitReference<hits_storage_type> reference;
504 
510  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
511  ReadHitsBinder(ScoringQueuesType& queues, const uint32 read_index = uint32(-1));
512 
516  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
517  void bind(const uint32 read_index);
518 
521  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
522  void resize(const uint32 size);
523 
526  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
527  uint32 size() const;
528 
532  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
533  HitReference<hits_storage_type> operator[] (const uint32 i) const;
534 
537  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
538  packed_read read_info() const;
539 
543  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
544  void set_read_info(const packed_read info);
545 
550  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
551  void bind_hit(const uint32 i, const uint32 slot);
552 
553 private:
554  ScoringQueuesType& m_queues;
555  uint32 m_read_index;
556 };
557 
560 
561 } // namespace cuda
562 } // namespace bowtie2
563 
564 // return a view of the queues
565 //
566 inline
567 bowtie2::cuda::ScoringQueuesDeviceView device_view(bowtie2::cuda::ScoringQueues& queues) { return queues.device_view(); }
568 
569 // return a view of the queues
570 //
571 inline
572 bowtie2::cuda::HitQueuesDeviceView device_view(bowtie2::cuda::HitQueues& queues) { return queues.device_view(); }
573 
574 // return a view of the index
575 //
576 inline
577 bowtie2::cuda::ReadHitsIndexDeviceView device_view(bowtie2::cuda::ReadHitsIndex& index) { return index.device_view(); }
578 
579 } // namespace nvbio
580 
581 #include <nvBowtie/bowtie2/cuda/scoring_queues_inl.h>
Generated on Wed Feb 25 2015 08:33:06 for NVBIO by   doxygen 1.8.4