proposal.h

Go to the documentation of this file.

/*
 * Copyright (c) 2012-13, NVIDIA CORPORATION.  All rights reserved.
 *
 * NVIDIA CORPORATION and its licensors retain all intellectual property
 * and proprietary rights in and to this software, related documentation
 * and any modifications thereto.  Any use, reproduction, disclosure or
 * distribution of this software and related documentation without an express
 * license agreement from NVIDIA CORPORATION is strictly prohibited.
 */

#pragma once

#include <nvBowtie/bowtie2/cuda/defs.h>
#include <nvbio/basic/cuda/arch.h>
#include <nvbio/basic/thrust_view.h>
#include <algorithm>

namespace cufmi {
namespace bowtie2 {
namespace cuda {

// forward declaration
template <typename HitQueuesType> struct HitReference;

struct HitQueuesDeviceView;
struct ScoringQueuesDeviceView;

struct ReadHitsIndex
{
    typedef thrust::device_vector<uint32>          links_storage_type;

    enum Mode {
        SingleHitPerRead        = 0u,
        MultipleHitsPerRead     = 1u,
    };

    // alloc enough storage for max_size hits
    //
    void resize(const uint32 max_size);

    // setup the number of input reads
    //
    void setup(const uint32 n_hits_per_read, const uint32 in_reads);

    // return the queue mode
    //
    bool mode() const { return m_mode; }

    // return a view of the data structure
    //
    ReadHitsIndexDeviceView device_view();

private:
    links_storage_type      m_links;
    uint32                  m_stride;
    Mode                    m_mode;
};

struct ReadHitsIndexDeviceView
{
    typedef typename device_view_subtype<ReadHitsIndex::links_storage_type>::type   links_storage_type;

    struct HitArray
    {
        CUFMI_FORCEINLINE CUFMI_HOST_DEVICE
        HitArray(const ReadHitsIndexDeviceView& index, const uint32 read_index) : m_index(index), m_read_index(read_index) {}

        CUFMI_FORCEINLINE CUFMI_HOST_DEVICE
        uint32 size() const { return m_index.hit_count( m_read_index ); }

        CUFMI_FORCEINLINE CUFMI_HOST_DEVICE
        uint32 operator[] (const uint32 slot) const { return m_index( m_read_index, slot ); }

    private:
        const ReadHitsIndexDeviceView& m_index;
        uint32                         m_read_index;
    };

    CUFMI_FORCEINLINE CUFMI_HOST_DEVICE
    ReadHitsIndexDeviceView(
        uint32*         links           = NULL,
        const uint32    links_stride    = 0u);

    CUFMI_FORCEINLINE CUFMI_HOST_DEVICE
    uint32 hit_count(const uint32 read_idx) const { return m_links ? m_links[ read_index ] : 1u; }

    CUFMI_FORCEINLINE CUFMI_HOST_DEVICE
    void set_hit_count(const uint32 read_index, const uint32 count) { m_links[ read_index ] = count; }

    CUFMI_FORCEINLINE CUFMI_HOST_DEVICE
    uint32& operator() (const uint32 read_index, const uint32 slot) { return m_links[ read_index + (1u + slot) * m_links_stride ]; }

    CUFMI_FORCEINLINE CUFMI_HOST_DEVICE
    uint32 operator() (const uint32 read_index, const uint32 slot) const { return m_links ? m_links[ read_index + (1u + slot) * m_links_stride ] : read_index; }

    CUFMI_FORCEINLINE CUFMI_HOST_DEVICE
    HitArray operator[] (const uint32 read_index) const { return HitArray( *this, read_index ); }

private:
    links_storage_type  m_links;
    uint32              m_links_stride;

    friend struct HitArray;
};

struct HitQueues
{
    typedef thrust::device_vector<uint32>          parent_storage_type;
    typedef thrust::device_vector<packed_seed>     seed_storage_type;
    typedef thrust::device_vector<uint32>          ssa_storage_type;
    typedef thrust::device_vector<uint32>          loc_storage_type;
    typedef thrust::device_vector<int32>           score_storage_type;
    typedef thrust::device_vector<uint32>          sink_storage_type;

    typedef HitQueuesDeviceView                    device_view_type;       
    typedef HitReference<HitQueues>                reference;              

    parent_storage_type    parent;             
    seed_storage_type      hit;                
    ssa_storage_type       ssa;                
    loc_storage_type       loc;                
    score_storage_type     score;              
    sink_storage_type      sink;               
    loc_storage_type       opposite_loc;       
    score_storage_type     opposite_score;     
    sink_storage_type      opposite_sink;      
};
struct HitQueuesDeviceView
{
    typedef typename device_view_subtype<ScoringQueues::parent_storage_type>::type           parent_storage_type;
    typedef typename device_view_subtype<ScoringQueues::seed_storage_type>::type             seed_storage_type;
    typedef typename device_view_subtype<ScoringQueues::ssa_storage_type>::type              ssa_storage_type;
    typedef typename device_view_subtype<ScoringQueues::loc_storage_type>::type              loc_storage_type;
    typedef typename device_view_subtype<ScoringQueues::score_storage_type>::type            score_storage_type;
    typedef typename device_view_subtype<ScoringQueues::sink_storage_type>::type             sink_storage_type;

    typedef HitQueuesDeviceView                    device_view_type;       
    typedef HitReference<HitQueuesView>            reference;              

    parent_storage_type    parent;             
    seed_storage_type      hit;                
    ssa_storage_type       ssa;                
    loc_storage_type       loc;                
    score_storage_type     score;              
    sink_storage_type      sink;               
    loc_storage_type       opposite_loc;       
    score_storage_type     opposite_score;     
    sink_storage_type      opposite_sink;      
};

struct ScoringQueues
{
    typedef PingPongQueues<packed_read>            active_reads_storage_type;
    typedef HitQueues                              hits_storage_type;
    typedef ReadHitsIndex                          read_hits_index_type;
    typedef ScoringQueuesDeviceView                device_view_type;

    active_reads_storage_type  active_reads;       
    hits_storage_type          hits;               
    read_hits_index_type       hits_index;         
};
struct ScoringQueuesDeviceView
{
    typedef typename device_view_subtype<ScoringQueues::active_reads_storage_type>::type     active_reads_storage_type;
    typedef typename device_view_subtype<ScoringQueues::hits_storage_type>::type             hits_storage_type;
    typedef typename device_view_subtype<ScoringQueues::read_hits_index_type>                read_hits_index_type;

    active_reads_storage_type  active_reads;       
    hits_storage_type          hits;               
    read_hits_index_type       hits_index;         
};

template <typename HitQueuesType>
struct HitReference
{
    typedef typename reference_subtype<typename HitQueuesType::ssa_storage_type>::type             ssa_type;
    typedef typename reference_subtype<typename HitQueuesType::loc_storage_type>::type             loc_type;
    typedef typename reference_subtype<typename HitQueuesType::sink_storage_type>::type            score_type;
    typedef typename reference_subtype<typename HitQueuesType::score_storage_type>::type           sink_type;

    HitReference(HitQueuesType& hits, const uint32 hit_index);

    uint32     parent;            
    ssa_type   ssa;               
    loc_type   loc;               
    score_type score;             
    sink_type  sink;              
    loc_type   opposite_loc;      
    score_type opposite_score;    
    sink_type  opposite_sink;     
};

template <typename ReadHitsIndexType, typename HitQueuesType>
struct ReadHitsReference
{
    ReadHitsReference(ReadHitsIndexType& index, HitQueuesType& hits, const uint32 read_index);

    uint32 size() const;

    CUFMI_FORCEINLINE CUFMI_HOST_DEVICE
    HitReference<HitQueuesType> operator[] (const uint32 i) const;

private:
    ReadHitsIndexType& m_index;        
    HitQueuesType&     m_hits;         
    uint32             m_read_index;   
};

template <typename ScoringQueuesType>
struct ReadReference
{
    typedef typename ScoringQueuesType::hits_storage_type             hits_storage_type;
    typedef typename ScoringQueuesType::read_hits_index_type          read_hits_index_type;

    typedef typename reference_subtype<typename ScoringQueuesType::active_reads_storage_type>::type   info_type;
    typedef ReadHitsReference<read_hits_index_type, hits_storage_type>                                hits_type;

    ReadReference(ScoringQueuesType scoring_queues, const uint32 read_index);

    info_type  info;              
    hits_type  hits;              
};

} // namespace cuda
} // namespace bowtie2
} // namespace cufmi