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 
36 #include <nvbio/basic/cuda/arch.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;
72 
100 {
101  typedef thrust::device_vector<uint32> links_storage_type;
103 
104  enum Mode {
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 
124 
125 public:
129 };
130 
136 {
138 
142  struct HitArray
143  {
147  HitArray(ReadHitsIndexDeviceView& index, const uint32 read_index) : m_index(index), m_read_index(read_index) {}
148 
152  void resize(const uint32 size) const { m_index.set_hit_count( m_read_index, size ); }
153 
157  uint32 size() const { return m_index.hit_count( m_read_index ); }
158 
162  uint32& operator[] (const uint32 slot) { return m_index( m_read_index, slot ); }
163 
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
176 
177 
183  const uint32 stride = 0u) : m_links( links ), m_stride( stride ) {}
184 
188  uint32 hit_count(const uint32 read_index) const { return m_links[ read_index ]; }
189 
193  void set_hit_count(const uint32 read_index, const uint32 count) { m_links[ read_index ] = count; }
194 
198  uint32& operator() (const uint32 read_index, const uint32 slot) { return m_links[ read_index + (1u + slot) * m_stride ]; }
199 
203  uint32 operator() (const uint32 read_index, const uint32 slot) const { return m_links[ read_index + (1u + slot) * m_stride ]; }
204 
208  HitArray operator[] (const uint32 read_index) { return HitArray( *this, read_index ); }
209 
210 public:
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 
231 
234  uint64 resize(const uint32 size, const bool do_alloc);
235 
239 
251 };
256 {
263 
266 
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 
287 
292 
304 };
305 
336 {
340  typedef thrust::device_vector<uint32> pool_type;
341 
343 
347 
350  uint64 resize(const uint32 n_reads, const uint32 n_hits, const bool do_alloc);
351 
355 
358  uint32 hits_count() { return hits_pool[0]; }
359 
363 
368 };
373 {
378 
383 
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 
396  uint32 active_read_count() const { return active_reads.in_size; }
397 
401  packed_read active_read(const uint32 read_index) const { return active_reads.in_queue[read_index]; }
402 
407 };
408 
414 template <typename HitQueuesType>
415 struct HitReference
416 {
423 
429  HitReference(HitQueuesType& hits, const uint32 hit_index);
430 
442 };
443 
447 template <typename ScoringQueuesType>
449 {
450  typedef typename ScoringQueuesType::read_hits_index_type read_hits_index_type;
451  typedef typename ScoringQueuesType::hits_storage_type hits_storage_type;
453 
460  ReadHitsReference(ScoringQueuesType& queues, const uint32 read_index = uint32(-1));
461 
466  void bind(const uint32 read_index);
467 
471  uint32 size() const;
472 
478 
482  packed_read read_info() const;
483 
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>
500 {
501  typedef typename ScoringQueuesType::read_hits_index_type read_hits_index_type;
502  typedef typename ScoringQueuesType::hits_storage_type hits_storage_type;
504 
511  ReadHitsBinder(ScoringQueuesType& queues, const uint32 read_index = uint32(-1));
512 
517  void bind(const uint32 read_index);
518 
522  void resize(const uint32 size);
523 
527  uint32 size() const;
528 
534 
538  packed_read read_info() const;
539 
544  void set_read_info(const packed_read info);
545 
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
568 
569 // return a view of the queues
570 //
571 inline
573 
574 // return a view of the index
575 //
576 inline
578 
579 } // namespace nvbio
580