primitives.h

/*
 * CUGAR : Cuda Graphics Accelerator
 *
 * Copyright (c) 2011-2018, 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 <cugar/basic/types.h>
#include <cugar/basic/numbers.h>
#include <cugar/basic/functors.h>
#include <cugar/basic/vector.h>
#include <cugar/basic/algorithms.h>
#include <cugar/basic/cuda/arch.h>
#include <cugar/basic/cuda/timer.h>
#include <thrust/reduce.h>
#include <thrust/scan.h>
#include <thrust/copy.h>
#include <thrust/sort.h>
#include <thrust/binary_search.h>
#include <thrust/merge.h>
#include <thrust/iterator/constant_iterator.h>

#if defined(__CUDACC__)
#include <cugar/basic/cuda/primitives.h>
#include <cugar/basic/cuda/sort.h>
#endif

#if defined (_OPENMP)
#include <omp.h>
#endif


namespace cugar {



template <typename system_tag, typename PredicateIterator>
bool any(
    const uint32            n,
    const PredicateIterator pred);

template <typename system_tag, typename PredicateIterator>
bool all(
    const uint32            n,
    const PredicateIterator pred);

template <typename system_tag, typename Iterator>
bool is_sorted(
    const uint32            n,
    const Iterator          values);

template <typename system_tag, typename Iterator, typename Headflags>
bool is_segment_sorted(
    const uint32            n,
    const Iterator          values,
    const Headflags         flags);


template <typename system_tag, typename Iterator, typename Functor>
void for_each(
    const uint64            n,
    const Iterator          in,
          Functor           functor);

template <typename system_tag, typename Iterator, typename Output, typename Functor>
void transform(
    const uint32            n,
    const Iterator          in,
    const Output            out,
    const Functor           functor);

template <typename system_tag, typename Iterator1, typename Iterator2, typename Output, typename Functor>
void transform(
    const uint32            n,
    const Iterator1         in1,
    const Iterator2         in2,
    const Output            out,
    const Functor           functor);

template <typename system_tag, typename InputIterator, typename BinaryOp>
typename std::iterator_traits<InputIterator>::value_type reduce(
    const uint32                        n,
    InputIterator                       in,
    BinaryOp                            op,
    cugar::vector<system_tag,uint8>&    temp_storage);

template <typename system_tag, typename InputIterator, typename OutputIterator, typename BinaryOp>
void inclusive_scan(
    const uint32                        n,
    InputIterator                       in,
    OutputIterator                      out,
    BinaryOp                            op,
    cugar::vector<system_tag,uint8>&    temp_storage);

template <typename system_tag, typename InputIterator, typename OutputIterator, typename BinaryOp, typename Identity>
void exclusive_scan(
    const uint32                        n,
    InputIterator                       in,
    OutputIterator                      out,
    BinaryOp                            op,
    Identity                            identity,
    cugar::vector<system_tag,uint8>&    temp_storage);

template <typename system_tag, typename InputIterator, typename FlagsIterator, typename OutputIterator>
uint32 copy_flagged(
    const uint32                        n,
    InputIterator                       in,
    FlagsIterator                       flags,
    OutputIterator                      out,
    cugar::vector<system_tag,uint8>&    temp_storage);

template <typename system_tag, typename InputIterator, typename OutputIterator, typename Predicate>
uint32 copy_if(
    const uint32                        n,
    InputIterator                       in,
    OutputIterator                      out,
    const Predicate                     pred,
    cugar::vector<system_tag,uint8>&    temp_storage);

template <typename system_tag, typename InputIterator, typename OutputIterator, typename CountIterator>
uint32 runlength_encode(
    const uint32                        n,
    InputIterator                       in,
    OutputIterator                      out,
    CountIterator                       counts,
    cugar::vector<system_tag,uint8>&    temp_storage);

template <typename system_tag, typename KeyIterator, typename ValueIterator, typename OutputKeyIterator, typename OutputValueIterator, typename ReductionOp>
uint32 reduce_by_key(
    const uint32                        n,
    KeyIterator                         keys_in,
    ValueIterator                       values_in,
    OutputKeyIterator                   keys_out,
    OutputValueIterator                 values_out,
    ReductionOp                         reduction_op,
    cugar::vector<system_tag,uint8>&    temp_storage);

template <typename system_tag, typename KeyIterator, typename ValueIterator, typename OutputIterator>
void lower_bound(
    const uint32                        n,
    ValueIterator                       values,
    const uint32                        n_keys,
    KeyIterator                         keys,
    OutputIterator                      indices);

template <typename system_tag, typename KeyIterator, typename ValueIterator, typename OutputIterator>
void upper_bound(
    const uint32                        n,
    ValueIterator                       values,
    const uint32                        n_keys,
    KeyIterator                         keys,
    OutputIterator                      indices);

template <typename system_tag, typename KeyIterator>
void radix_sort(
    const uint32                        n,
    KeyIterator                         keys,
    cugar::vector<system_tag,uint8>&    temp_storage);

template <typename system_tag, typename KeyIterator, typename ValueIterator>
void radix_sort(
    const uint32                        n,
    KeyIterator                         keys,
    ValueIterator                       values,
    cugar::vector<system_tag,uint8>&    temp_storage);

template <
    typename system_tag,
    typename key_iterator1,
    typename key_iterator2,
    typename value_iterator1,
    typename value_iterator2,
    typename key_output,
    typename value_output>
void merge_by_key(
    const uint32                        A_len,
    const uint32                        B_len,
    const key_iterator1                 A_keys,
    const key_iterator2                 B_keys,
    const value_iterator1               A_values,
    const value_iterator2               B_values,
          key_output                    C_keys,
          value_output                  C_values,
    cugar::vector<system_tag,uint8>&    temp_storage);

template <typename system_tag>
struct for_each_enactor
{
    template <typename Iterator, typename Functor>
    void operator () (
        const uint64    n,
        const Iterator  in,
              Functor   functor)
    {
        for_each<system_tag>( n, in, functor );
    }

    template <typename Functor>
    void operator () (
        const uint64    n,
              Functor   functor)
    {
        for_each<system_tag>( n, thrust::make_counting_iterator<uint64>(0), functor );
    }
};

template <>
struct for_each_enactor<device_tag>
{
    for_each_enactor() :
        m_blocks_lo( 0 ),
        m_blocks_hi( 0 ),
        m_speed_lo( 0.0f ),
        m_speed_hi( 0.0f ) {}

    template <typename Iterator, typename Functor>
    void operator () (
        const uint64            n,
        const Iterator          in,
              Functor           functor);

    template <typename Functor>
    void operator () (
        const uint64    n,
              Functor   functor)
    {
        this->operator()( n, thrust::make_counting_iterator<uint64>(0), functor );
    }

  private:

    template <typename KernelFunction>
    uint32 suggested_blocks(KernelFunction kernel, const uint32 cta_size) const;

    void update(const uint32 n_blocks, const float speed);

    uint32  m_blocks_lo;
    uint32  m_blocks_hi;
    float   m_speed_lo;
    float   m_speed_hi;
};


} // namespace cugar

#include <cugar/basic/primitives_inl.h>