ssa.h

Go to the documentation of this file.

/*
 * 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.
 */

#pragma once

#include <nvbio/basic/types.h>
#include <nvbio/basic/popcount.h>
#include <nvbio/basic/cuda/arch.h>
#include <nvbio/basic/cuda/ldg.h>
#include <vector_types.h>
#include <vector_functions.h>
#include <cuda_runtime.h>
#include <vector>
#include <stdexcept>
#include <thrust/device_vector.h>

namespace nvbio {




struct SSA_value_multiple_device;

template <uint32 K, typename index_type = uint32>
struct SSA_index_multiple_device;

template <typename SSAIterator, typename BitmaskIterator, typename BlockIterator>
struct SSA_value_multiple_context
{
    NVBIO_FORCEINLINE NVBIO_HOST_DEVICE SSA_value_multiple_context() {}

    NVBIO_FORCEINLINE NVBIO_HOST_DEVICE SSA_value_multiple_context(
        SSAIterator     ssa,
        BitmaskIterator bitmask,
        BlockIterator   blocks) : m_ssa( ssa ), m_bitmask( bitmask ), m_blocks( blocks ) {}

    NVBIO_FORCEINLINE NVBIO_HOST_DEVICE bool fetch(const uint32 i, uint32& r) const;

    NVBIO_FORCEINLINE NVBIO_HOST_DEVICE bool has(const uint32 i) const;

    SSAIterator     m_ssa;
    BitmaskIterator m_bitmask;
    BlockIterator   m_blocks;
};

struct SSA_value_multiple
{
    typedef SSA_value_multiple_context<
        const uint32*,
        const uint32*,
        const uint32*>                  context_type;
    typedef SSA_value_multiple_device   device_type;
    typedef context_type                device_view_type;
    typedef context_type                plain_view_type;

    SSA_value_multiple() : m_n(0), m_stored(0) {}

    SSA_value_multiple(
        const uint32  n,
        const int32*  sa,
        const uint32  K);

    template <typename FMIndexType>
    SSA_value_multiple(
        const FMIndexType& fmi,
        const uint32  K);

    context_type get_context() const
    {
        return context_type( &m_ssa[0], &m_bitmask[0], &m_blocks[0] );
    }

    uint32  m_n;
    uint32  m_stored;
    std::vector<uint32> m_ssa;
    std::vector<uint32> m_bitmask;
    std::vector<uint32> m_blocks;

private:
    uint32 index(const uint32 i) const;
};

struct SSA_value_multiple_device
{
    typedef SSA_value_multiple_context<
        cuda::ldg_pointer<uint32>,
        cuda::ldg_pointer<uint32>,
        cuda::ldg_pointer<uint32> >     context_type;
    typedef context_type                device_view_type;
    typedef context_type                plain_view_type;

    SSA_value_multiple_device(const SSA_value_multiple& ssa);

    ~SSA_value_multiple_device();

    context_type get_context() const
    {
        return context_type( cuda::ldg_pointer<uint32>(m_ssa), cuda::ldg_pointer<uint32>(m_bitmask), cuda::ldg_pointer<uint32>(m_blocks) );
    }

    uint32  m_n;
    uint32  m_stored;
    uint32* m_ssa;
    uint32* m_bitmask;
    uint32* m_blocks;
};

template <uint32 K, typename Iterator = const uint32*>
struct SSA_index_multiple_context
{
    typedef typename std::iterator_traits<Iterator>::value_type index_type;
    typedef index_type                                          value_type;

    NVBIO_FORCEINLINE NVBIO_HOST_DEVICE SSA_index_multiple_context() {}

    NVBIO_FORCEINLINE NVBIO_HOST_DEVICE SSA_index_multiple_context(
        const Iterator ssa) : m_ssa( ssa ) {}

    NVBIO_FORCEINLINE NVBIO_HOST_DEVICE bool fetch(const index_type i, index_type& r) const;

    NVBIO_FORCEINLINE NVBIO_HOST_DEVICE bool has(const index_type i) const;

    Iterator m_ssa;
};

template <uint32 K, typename index_type = uint32>
struct SSA_index_multiple
{
    typedef index_type                                          value_type;
    typedef SSA_index_multiple_context<K, const index_type*>    context_type;
    typedef SSA_index_multiple_device<K,index_type>             device_type;
    typedef context_type                                        device_view_type;
    typedef context_type                                        plain_view_type;

    SSA_index_multiple() : m_n(0) {}

    SSA_index_multiple(
        const index_type  n,
        const index_type* sa);

    template <typename FMIndexType>
    SSA_index_multiple(
        const FMIndexType& fmi);

    SSA_index_multiple(
        const SSA_index_multiple_device<K,index_type>& ssa);

    SSA_index_multiple& operator=(
        const SSA_index_multiple_device<K,index_type>& ssa);

    context_type get_context() const { return context_type( &m_ssa[0] ); }

    index_type              m_n;
    std::vector<index_type> m_ssa;
};

template <uint32 K, typename index_type>
struct SSA_index_multiple_device
{
    typedef index_type                                          value_type;
    typedef SSA_index_multiple_context<K, const index_type*>    context_type;
    typedef context_type                                        device_view_type;
    typedef context_type                                        plain_view_type;

    SSA_index_multiple_device() : m_n(0) {}

    SSA_index_multiple_device(const SSA_index_multiple<K, index_type>& ssa);

#ifdef __CUDACC__
    template <typename FMIndexType>
    SSA_index_multiple_device(
        const FMIndexType& fmi);
    #endif

    ~SSA_index_multiple_device() {}

    #ifdef __CUDACC__
    template <typename FMIndexType>
    void init(const FMIndexType& fmi);
    #endif

    context_type get_context() const { return context_type( thrust::raw_pointer_cast(&m_ssa[0]) ); }

    index_type                        m_n;
    thrust::device_vector<index_type> m_ssa;
};

inline
SSA_value_multiple::plain_view_type plain_view(const SSA_value_multiple& vec) { return vec.get_context(); }

inline
SSA_value_multiple_device::plain_view_type plain_view(const SSA_value_multiple_device& vec) { return vec.get_context(); }

template <uint32 K, typename index_type>
typename SSA_index_multiple<K,index_type>::plain_view_type plain_view(const SSA_index_multiple<K,index_type>& vec) { return vec.get_context(); }



} // namespace nvbio

#include <nvbio/fmindex/ssa_inl.h>