/* * nvbio * Copyright (c) 2011-2014, NVIDIA CORPORATION. All rights reserved. * * Redistribution and use in source and binary forms, with or without * modification, are permitted provided that the following conditions are met: * * Redistributions of source code must retain the above copyright * notice, this list of conditions and the following disclaimer. * * Redistributions in binary form must reproduce the above copyright * notice, this list of conditions and the following disclaimer in the * documentation and/or other materials provided with the distribution. * * Neither the name of the NVIDIA CORPORATION nor the * names of its contributors may be used to endorse or promote products * derived from this software without specific prior written permission. * * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. */ /// ///\file select.h /// #pragma once #include #include #include #include #include #include #include namespace nvbio { namespace bowtie2 { namespace cuda { template struct BaseScoringPipelineState; template struct BestApproxScoringPipelineState; ///@addtogroup nvBowtie ///@{ /// \defgroup Select /// /// The functions in this module implement a pipeline stage in which, for all the active reads in the /// \ref ScoringQueues, a group of seed hits is selected for a new round of \ref Scoring. /// In terms of inputs and outputs, this stage takes the set of input ScoringQueues::active_reads, and /// fills the ScoringQueues::hits with a new set of hits, writing the HitQueues::read_id, HitQueues::seed, /// and HitQueues::loc fields. The hits' location specified by the HitQueues::loc field is expressed in /// Suffix Array coordinates. /// Finally, this stage is also responsible for producing a new set of ScoringQueues::active_reads. /// /// \b inputs: /// - SeedHitDequeArray /// - ScoringQueues::active_reads /// /// \b outputs: /// - SeedHitDequeArray /// - ScoringQueues::active_reads /// - HitQueues::read_id /// - HitQueues::seed /// - HitQueues::loc /// ///@addtogroup Select ///@{ /// /// Initialize the hit-selection pipeline /// void select_init( const uint32 count, const char* read_names, const uint32* read_names_idx, const SeedHitDequeArrayDeviceView hits, uint32* trys, uint32* rseeds, const ParamsPOD params); /// /// Initialize the hit-selection pipeline /// void select_init(BestApproxScoringPipelineState& pipeline, const ParamsPOD& params); /// /// Initialize the hit-selection pipeline /// void select_init(BestApproxScoringPipelineState >& pipeline, const ParamsPOD& params); /// /// a context class for the select_kernel /// struct SelectBestApproxContext { // constructor // // \param trys the per-read vector of extension trys // SelectBestApproxContext(uint32* trys) : m_trys( trys ) {} // stopping function // // \return true iff we can stop the hit selection process // NVBIO_FORCEINLINE NVBIO_DEVICE bool stop(const uint32 read_id) const { return m_trys[ read_id ] == 0; } private: uint32* m_trys; }; /// /// select next hit extensions from the top seed ranges /// __global__ void select_n_from_top_range_kernel( const uint32 begin, const uint32 count, const uint32 n_reads, const SeedHit* hit_data, const uint32* hit_range_scan, uint32* loc_queue, uint32* seed_queue, uint32* read_info); /// /// Prepare for a round of seed extension by selecting the next SA row from each /// of the seed-hit deque arrays (SeedHitDequeArray) bound to the active-reads in /// the scoring queues (ScoringQueues::active_reads). /// void select( const SelectBestApproxContext context, const BestApproxScoringPipelineState& pipeline, const ParamsPOD params); /// /// Prepare for a round of seed extension by selecting the next SA row from each /// of the seed-hit deque arrays (SeedHitDequeArray) bound to the active-reads in /// the scoring queues (ScoringQueues::active_reads). /// void select( const SelectBestApproxContext context, const BestApproxScoringPipelineState >& pipeline, const ParamsPOD params); /// /// Select next hit extensions for all-mapping /// void select_all( const uint64 begin, const uint32 count, const uint32 n_reads, const uint32 n_hit_ranges, const uint64 n_hits, const SeedHitDequeArrayDeviceView hits, const uint32* hit_count_scan, const uint64* hit_range_scan, HitQueuesDeviceView scoring_queues); ///@} // group Select ///@} // group nvBowtie } // namespace cuda } // namespace bowtie2 } // namespace nvbio