NVBIO
 All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros Groups Pages
seed_hit_deque_array.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 <nvBowtie/bowtie2/cuda/seed_hit.h>
36 #include <nvbio/basic/cuda/pingpong_queues.h>
37 #include <nvbio/basic/cuda/arch.h>
38 #include <nvbio/basic/thrust_view.h>
39 #include <nvbio/basic/priority_deque.h>
40 #include <nvbio/basic/sum_tree.h>
41 #include <nvbio/basic/vector_view.h>
42 #include <nvbio/basic/vector.h>
43 #include <algorithm>
44 
45 namespace nvbio {
46 namespace bowtie2 {
47 namespace cuda {
48 
51 
54 
55 struct SeedHitDequeArrayDeviceView;
56 
57 template <typename SeedHitDequeArrayType> struct SeedHitDequeReference;
58 
62 struct SeedHitDequeArrayDeviceStorage
63 {
64  typedef thrust::device_vector<SeedHit> hits_storage_type;
65  typedef thrust::device_vector<uint32> index_storage_type;
66  typedef thrust::device_vector<float> prob_storage_type;
67 
68  hits_storage_type m_hits;
69  index_storage_type m_counts;
70  prob_storage_type m_probs;
71  index_storage_type m_index;
72  index_storage_type m_probs_index;
73  index_storage_type m_pool;
74  index_storage_type m_probs_pool;
75 };
76 
80 struct SeedHitDequeArrayHostStorage
81 {
82  typedef thrust::host_vector<SeedHit> hits_storage_type;
83  typedef thrust::host_vector<uint32> index_storage_type;
84  typedef thrust::host_vector<float> prob_storage_type;
85 
88  SeedHitDequeArrayHostStorage(const SeedHitDequeArrayDeviceStorage& other) { (*this)= other; }
89 
92  SeedHitDequeArrayHostStorage& operator=(const SeedHitDequeArrayDeviceStorage& other)
93  {
94  nvbio::cuda::thrust_copy_vector( m_hits, other.m_hits );
95  nvbio::cuda::thrust_copy_vector( m_counts, other.m_counts );
96  nvbio::cuda::thrust_copy_vector( m_probs, other.m_probs );
97  nvbio::cuda::thrust_copy_vector( m_index, other.m_index );
98  nvbio::cuda::thrust_copy_vector( m_probs_index, other.m_probs_index );
99  nvbio::cuda::thrust_copy_vector( m_pool, other.m_pool );
100  nvbio::cuda::thrust_copy_vector( m_probs_pool, other.m_probs_pool );
101  return *this;
102  }
103 
104  hits_storage_type m_hits;
105  index_storage_type m_counts;
106  prob_storage_type m_probs;
107  index_storage_type m_index;
108  index_storage_type m_probs_index;
109  index_storage_type m_pool;
110  index_storage_type m_probs_pool;
111 };
112 
116 struct SeedHitDequeArray : public SeedHitDequeArrayDeviceStorage
117 {
118  typedef SeedHitDequeArrayDeviceStorage::hits_storage_type hits_storage_type;
119  typedef SeedHitDequeArrayDeviceStorage::index_storage_type index_storage_type;
120  typedef SeedHitDequeArrayDeviceStorage::prob_storage_type prob_storage_type;
121 
122  typedef SeedHitDequeArrayDeviceView device_view_type;
123 
127  uint64 resize(const uint32 n_reads, const uint32 max_hits, const bool do_alloc = true);
128 
131  void clear_deques();
132 
135  SeedHitDequeArrayDeviceView device_view();
136 
139  hits_storage_type& hits() { return m_hits; }
140 
143  index_storage_type& counts() { return m_counts; }
144 
147  index_storage_type& index() { return m_index; }
148 
151  prob_storage_type& probs() { return m_probs; }
152 };
153 
157 struct SeedHitDequeArrayDeviceView
158 {
159  typedef typename device_view_subtype< thrust::device_vector<SeedHit> >::type hits_storage_type;
160  typedef typename device_view_subtype< thrust::device_vector<uint32> >::type index_storage_type;
161  typedef typename device_view_subtype< thrust::device_vector<float> >::type prob_storage_type;
162 
163  typedef vector_view<SeedHit*> hit_vector_type;
164  typedef priority_deque<SeedHit, hit_vector_type, hit_compare> hit_deque_type;
165 
166  typedef SeedHitDequeReference<SeedHitDequeArrayDeviceView> reference;
167 
170  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
171  SeedHitDequeArrayDeviceView(
172  index_storage_type counts = index_storage_type(),
173  index_storage_type index = index_storage_type(),
174  hits_storage_type hits = hits_storage_type(),
175  index_storage_type probs_index = index_storage_type(),
176  prob_storage_type probs = prob_storage_type(),
177  index_storage_type pool = index_storage_type(),
178  index_storage_type probs_pool = index_storage_type());
179 
182  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
183  reference operator[] (const uint32 read_id);
184 
187  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
188  SeedHit* alloc_deque(const uint32 read_id, const uint32 size);
189 
193  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
194  void resize_deque(const uint32 read_id, const uint32 size) { m_counts[read_id] = size; }
195 
198  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
199  uint32 get_size(const uint32 read_id) const { return m_counts[read_id]; }
200 
203  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
204  SeedHit* get_data(const uint32 read_id) const { return m_hits.base() + m_index[read_id]; }
205 
208  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
209  float* get_probs(const uint32 read_id) const { return m_probs.base() + m_probs_index[read_id]; }
210 
214  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
215  void erase(const uint32 read_id) { m_counts[read_id] = 0; }
216 
219  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
220  hit_deque_type get_deque(const uint32 read_id, bool build_heap = false) const;
221 
222 //private:
223  hits_storage_type m_hits;
224  index_storage_type m_counts;
225  prob_storage_type m_probs;
226  index_storage_type m_index;
227  index_storage_type m_probs_index;
228  index_storage_type m_pool;
229  index_storage_type m_probs_pool;
230 };
231 
237 template <typename SeedHitDequeArrayType>
238 struct SeedHitDequeReference
239 {
240  typedef vector_view<SeedHit*> hit_vector_type;
241  typedef priority_deque<SeedHit, hit_vector_type, hit_compare> hit_deque_type;
242 
245  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
246  SeedHitDequeReference(SeedHitDequeArrayType& deques, const uint32 read_id) :
247  m_deques( deques ),
248  m_read_id( read_id ),
249  m_deque( deques.get_deque( read_id ) ) {}
250 
253  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
254  ~SeedHitDequeReference()
255  {
256  // update the size of the deque if dirty
257  m_deques.resize_deque( m_read_id, m_deque.size() );
258  }
259 
262  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
263  SeedHit* alloc(const uint32 size)
264  {
265  SeedHit* storage = m_deques.alloc_deque( m_read_id, size );
266  m_deque = m_deques.get_deque( m_read_id ); // re-assign the deque
267  return storage;
268  }
269 
273  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
274  void resize(const uint32 size)
275  {
276  m_deques.resize_deque( m_read_id, size );
277  m_deque = m_deques.get_deque( m_read_id ); // re-assign the deque
278  }
279 
282  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
283  uint32 size() const { return m_deque.size(); }
284 
287  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
288  SeedHit* get_data() const { return m_deques.get_data( m_read_id ); }
289 
292  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
293  float* get_probs() const { return m_deques.get_probs( m_read_id ); }
294 
298  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
299  void erase() { m_deques.erase( m_read_id ); m_deque = get_deque( m_read_id ); }
300 
303  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
304  hit_deque_type get_deque(bool build_heap = false) const { return m_deques.get_deque( m_read_id, build_heap ); }
305 
308  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
309  void push(const SeedHit hit) { m_deque.push( hit ); }
310 
313  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
314  void pop_top() { m_deque.pop_top(); }
315 
318  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
319  void pop_bottom() { m_deque.pop_bottom(); }
320 
323  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
324  void pop() { m_deque.pop(); }
325 
328  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
329  const SeedHit& top() const { return m_deque.top(); }
330 
333  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
334  const SeedHit& bottom() const { return m_deque.bottom(); }
335 
338  NVBIO_FORCEINLINE NVBIO_HOST_DEVICE
339  SeedHit& operator[] (const uint32 i) { return get_data()[i]; }
340 
341 private:
342  SeedHitDequeArrayType& m_deques;
343  uint32 m_read_id;
344  hit_deque_type m_deque;
345 };
346 
349 
350 } // namespace cuda
351 } // namespace bowtie2
352 } // namespace nvbio
353 
354 #include <nvBowtie/bowtie2/cuda/seed_hit_deque_array_inl.h>
Generated on Wed Feb 25 2015 08:33:06 for NVBIO by   doxygen 1.8.4