NVBIO
 All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros Groups Pages
proposal.h
Go to the documentation of this file.
1 /*
2  * Copyright (c) 2012-13, NVIDIA CORPORATION. All rights reserved.
3  *
4  * NVIDIA CORPORATION and its licensors retain all intellectual property
5  * and proprietary rights in and to this software, related documentation
6  * and any modifications thereto. Any use, reproduction, disclosure or
7  * distribution of this software and related documentation without an express
8  * license agreement from NVIDIA CORPORATION is strictly prohibited.
9  */
10 
11 #pragma once
12 
14 #include <nvbio/basic/cuda/arch.h>
16 #include <algorithm>
17 
18 namespace cufmi {
19 namespace bowtie2 {
20 namespace cuda {
21 
22 // forward declaration
23 template <typename HitQueuesType> struct HitReference;
24 
25 struct HitQueuesDeviceView;
27 
33 {
34  typedef thrust::device_vector<uint32> links_storage_type;
35 
36  enum Mode {
39  };
40 
41  // alloc enough storage for max_size hits
42  //
43  void resize(const uint32 max_size);
44 
45  // setup the number of input reads
46  //
47  void setup(const uint32 n_hits_per_read, const uint32 in_reads);
48 
49  // return the queue mode
50  //
51  bool mode() const { return m_mode; }
52 
53  // return a view of the data structure
54  //
56 
57 private:
58  links_storage_type m_links;
59  uint32 m_stride;
60  Mode m_mode;
61 };
62 
68 {
69  typedef typename device_view_subtype<ReadHitsIndex::links_storage_type>::type links_storage_type;
70 
74  struct HitArray
75  {
78  CUFMI_FORCEINLINE CUFMI_HOST_DEVICE
79  HitArray(const ReadHitsIndexDeviceView& index, const uint32 read_index) : m_index(index), m_read_index(read_index) {}
80 
83  CUFMI_FORCEINLINE CUFMI_HOST_DEVICE
84  uint32 size() const { return m_index.hit_count( m_read_index ); }
85 
88  CUFMI_FORCEINLINE CUFMI_HOST_DEVICE
89  uint32 operator[] (const uint32 slot) const { return m_index( m_read_index, slot ); }
90 
91  private:
92  const ReadHitsIndexDeviceView& m_index;
93  uint32 m_read_index;
94  };
95 
98  CUFMI_FORCEINLINE CUFMI_HOST_DEVICE
100  uint32* links = NULL,
101  const uint32 links_stride = 0u);
102 
105  CUFMI_FORCEINLINE CUFMI_HOST_DEVICE
106  uint32 hit_count(const uint32 read_idx) const { return m_links ? m_links[ read_index ] : 1u; }
107 
110  CUFMI_FORCEINLINE CUFMI_HOST_DEVICE
111  void set_hit_count(const uint32 read_index, const uint32 count) { m_links[ read_index ] = count; }
112 
115  CUFMI_FORCEINLINE CUFMI_HOST_DEVICE
116  uint32& operator() (const uint32 read_index, const uint32 slot) { return m_links[ read_index + (1u + slot) * m_links_stride ]; }
117 
120  CUFMI_FORCEINLINE CUFMI_HOST_DEVICE
121  uint32 operator() (const uint32 read_index, const uint32 slot) const { return m_links ? m_links[ read_index + (1u + slot) * m_links_stride ] : read_index; }
122 
125  CUFMI_FORCEINLINE CUFMI_HOST_DEVICE
126  HitArray operator[] (const uint32 read_index) const { return HitArray( *this, read_index ); }
127 
128 private:
129  links_storage_type m_links;
130  uint32 m_links_stride;
131 
132  friend struct HitArray;
133 };
134 
138 struct HitQueues
139 {
140  typedef thrust::device_vector<uint32> parent_storage_type;
141  typedef thrust::device_vector<packed_seed> seed_storage_type;
142  typedef thrust::device_vector<uint32> ssa_storage_type;
143  typedef thrust::device_vector<uint32> loc_storage_type;
144  typedef thrust::device_vector<int32> score_storage_type;
145  typedef thrust::device_vector<uint32> sink_storage_type;
146 
149 
159 };
164 {
165  typedef typename device_view_subtype<ScoringQueues::parent_storage_type>::type parent_storage_type;
166  typedef typename device_view_subtype<ScoringQueues::seed_storage_type>::type seed_storage_type;
167  typedef typename device_view_subtype<ScoringQueues::ssa_storage_type>::type ssa_storage_type;
168  typedef typename device_view_subtype<ScoringQueues::loc_storage_type>::type loc_storage_type;
169  typedef typename device_view_subtype<ScoringQueues::score_storage_type>::type score_storage_type;
170  typedef typename device_view_subtype<ScoringQueues::sink_storage_type>::type sink_storage_type;
171 
174 
184 };
185 
190 {
191  typedef PingPongQueues<packed_read> active_reads_storage_type;
195 
199 };
204 {
205  typedef typename device_view_subtype<ScoringQueues::active_reads_storage_type>::type active_reads_storage_type;
206  typedef typename device_view_subtype<ScoringQueues::hits_storage_type>::type hits_storage_type;
207  typedef typename device_view_subtype<ScoringQueues::read_hits_index_type> read_hits_index_type;
208 
212 };
213 
219 template <typename HitQueuesType>
220 struct HitReference
221 {
222  typedef typename reference_subtype<typename HitQueuesType::ssa_storage_type>::type ssa_type;
223  typedef typename reference_subtype<typename HitQueuesType::loc_storage_type>::type loc_type;
224  typedef typename reference_subtype<typename HitQueuesType::sink_storage_type>::type score_type;
225  typedef typename reference_subtype<typename HitQueuesType::score_storage_type>::type sink_type;
226 
231  HitReference(HitQueuesType& hits, const uint32 hit_index);
232 
241 };
242 
246 template <typename ReadHitsIndexType, typename HitQueuesType>
248 {
254  ReadHitsReference(ReadHitsIndexType& index, HitQueuesType& hits, const uint32 read_index);
255 
258  uint32 size() const;
259 
262  CUFMI_FORCEINLINE CUFMI_HOST_DEVICE
264 
265 private:
266  ReadHitsIndexType& m_index;
267  HitQueuesType& m_hits;
268  uint32 m_read_index;
269 };
270 
274 template <typename ScoringQueuesType>
276 {
277  typedef typename ScoringQueuesType::hits_storage_type hits_storage_type;
278  typedef typename ScoringQueuesType::read_hits_index_type read_hits_index_type;
279 
280  typedef typename reference_subtype<typename ScoringQueuesType::active_reads_storage_type>::type info_type;
282 
287  ReadReference(ScoringQueuesType scoring_queues, const uint32 read_index);
288 
291 };
292 
293 } // namespace cuda
294 } // namespace bowtie2
295 } // namespace cufmi