mapping.cu 7.7 KB

123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102103104105106107108109110111112113114115116117118119120121122123124125126127128129130131132133134135136137138139140141142143144145146147148149150151152153154155156157158159160161162163164165166167168169170171172173174175176177178179180181182183184185186187188189190191
  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. #include <nvBowtie/bowtie2/cuda/mapping.h>
  28. #include <nvBowtie/bowtie2/cuda/mapping_impl.h>
  29. namespace nvbio {
  30. namespace bowtie2 {
  31. namespace cuda {
  32. //
  33. // For all i in [0, #seed hit ranges[, output the seed hit range size in
  34. // out_ranges[i].
  35. //
  36. __global__
  37. void gather_ranges_kernel(
  38. const uint32 count,
  39. const uint32 n_reads,
  40. const SeedHitDequeArrayDeviceView hits,
  41. const uint32* hit_counts_scan,
  42. uint64* out_ranges)
  43. {
  44. const uint32 thread_id = threadIdx.x + BLOCKDIM*blockIdx.x;
  45. if (thread_id >= count) return;
  46. // do a binary search, looking for thread_id in hit_counts_scan,
  47. // to find the corresponding read id.
  48. const uint32 read_id = upper_bound_index( thread_id, hit_counts_scan, n_reads );
  49. // at this point we can figure out which seed hit / SA range this thread is
  50. // responsible of
  51. const uint32 count_offset = read_id ? hit_counts_scan[read_id-1] : 0u;
  52. const uint32 range_id = thread_id - count_offset;
  53. const SeedHit* hits_data = hits.get_data( read_id );
  54. const uint2 range = hits_data[ range_id ].get_range();
  55. // and we can compute the corresponding range size
  56. out_ranges[ thread_id ] = range.y - range.x;
  57. }
  58. //
  59. // dispatch the call to gather_ranges_kernel
  60. //
  61. void gather_ranges(
  62. const uint32 count,
  63. const uint32 n_reads,
  64. const SeedHitDequeArrayDeviceView hits,
  65. const uint32* hit_counts_scan,
  66. uint64* out_ranges)
  67. {
  68. const int blocks = (count + BLOCKDIM-1) / BLOCKDIM;
  69. gather_ranges_kernel<<<blocks, BLOCKDIM>>>( count, n_reads, hits, hit_counts_scan, out_ranges );
  70. }
  71. //
  72. // perform exact read mapping
  73. //
  74. void map_whole_read(
  75. const ReadsDef::type& read_batch,
  76. const FMIndexDef::type fmi,
  77. const FMIndexDef::type rfmi,
  78. const nvbio::cuda::PingPongQueuesView<uint32> queues,
  79. uint8* reseed,
  80. SeedHitDequeArrayDeviceView hits,
  81. const ParamsPOD params,
  82. const bool fw,
  83. const bool rc)
  84. {
  85. map_whole_read_t( read_batch, fmi, rfmi, queues, reseed, hits, params, fw, rc );
  86. }
  87. //
  88. // perform one run of exact seed mapping for all the reads in the input queue,
  89. // writing reads that need another run in the output queue
  90. //
  91. void map_exact(
  92. const ReadsDef::type& read_batch,
  93. const FMIndexDef::type fmi,
  94. const FMIndexDef::type rfmi,
  95. const uint32 retry,
  96. const nvbio::cuda::PingPongQueuesView<uint32> queues,
  97. uint8* reseed,
  98. SeedHitDequeArrayDeviceView hits,
  99. const ParamsPOD params,
  100. const bool fw,
  101. const bool rc)
  102. {
  103. map_exact_t( read_batch, fmi, rfmi, retry, queues, reseed, hits, params, fw, rc );
  104. }
  105. //
  106. // perform multiple runs of exact seed mapping in one go and keep the best
  107. //
  108. void map_exact(
  109. const ReadsDef::type& read_batch,
  110. const FMIndexDef::type fmi,
  111. const FMIndexDef::type rfmi,
  112. SeedHitDequeArrayDeviceView hits,
  113. const uint2 seed_range,
  114. const ParamsPOD params,
  115. const bool fw,
  116. const bool rc)
  117. {
  118. map_exact_t( read_batch, fmi, rfmi, hits, seed_range, params, fw, rc );
  119. }
  120. //
  121. // perform one run of approximate seed mapping for all the reads in the input queue,
  122. // writing reads that need another run in the output queue
  123. //
  124. void map_approx(
  125. const ReadsDef::type& read_batch,
  126. const FMIndexDef::type fmi,
  127. const FMIndexDef::type rfmi,
  128. const uint32 retry,
  129. const nvbio::cuda::PingPongQueuesView<uint32> queues,
  130. uint8* reseed,
  131. SeedHitDequeArrayDeviceView hits,
  132. const ParamsPOD params,
  133. const bool fw,
  134. const bool rc)
  135. {
  136. map_approx_t( read_batch, fmi, rfmi, retry, queues, reseed, hits, params, fw, rc );
  137. }
  138. //
  139. // perform multiple runs of approximate seed mapping in one go and keep the best
  140. //
  141. void map_approx(
  142. const ReadsDef::type& read_batch,
  143. const FMIndexDef::type fmi,
  144. const FMIndexDef::type rfmi,
  145. SeedHitDequeArrayDeviceView hits,
  146. const uint2 seed_range,
  147. const ParamsPOD params,
  148. const bool fw,
  149. const bool rc)
  150. {
  151. map_approx_t( read_batch, fmi, rfmi, hits, seed_range, params, fw, rc );
  152. }
  153. //
  154. // perform one run of seed mapping
  155. //
  156. void map(
  157. const ReadsDef::type& read_batch,
  158. const FMIndexDef::type fmi,
  159. const FMIndexDef::type rfmi,
  160. const uint32 retry,
  161. const nvbio::cuda::PingPongQueuesView<uint32> queues,
  162. uint8* reseed,
  163. SeedHitDequeArrayDeviceView hits,
  164. const ParamsPOD params,
  165. const bool fw,
  166. const bool rc)
  167. {
  168. map_t( read_batch, fmi, rfmi, retry, queues, reseed, hits, params, fw, rc );
  169. }
  170. } // namespace cuda
  171. } // namespace bowtie2
  172. } // namespace nvbio