<commit_before>/*
 *  Copyright 2008-2009 NVIDIA Corporation
 *
 *  Licensed under the Apache License, Version 2.0 (the "License");
 *  you may not use this file except in compliance with the License.
 *  You may obtain a copy of the License at
 *
 *      http://www.apache.org/licenses/LICENSE-2.0
 *
 *  Unless required by applicable law or agreed to in writing, software
 *  distributed under the License is distributed on an "AS IS" BASIS,
 *  WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
 *  See the License for the specific language governing permissions and
 *  limitations under the License.
 */


/*! \file fill.inl
 *  \brief Inline file for fill.h.
 */

#ifdef __CUDACC__

/////////////////////
// NVCC definition //
/////////////////////

#include <thrust/detail/util/align.h>
#include <thrust/detail/device/dereference.h>
#include <thrust/detail/device/cuda/vectorize.h>
#include <thrust/iterator/iterator_traits.h>
#include <thrust/detail/raw_buffer.h>

namespace thrust
{

namespace detail
{

namespace device
{

namespace detail
{

template <typename ForwardIterator, typename T>
struct fill_functor
{
    ForwardIterator first;
    T exemplar;

    fill_functor(ForwardIterator _first, T _exemplar) 
        : first(_first), exemplar(_exemplar) {}

    template <typename IntegerType>
        __device__
        void operator()(const IntegerType i)
        { 
            thrust::detail::device::dereference(first, i) = exemplar;
        }
}; // end fill_functor


template<typename OutputType, typename T>
  void wide_fill(OutputType * first,
                 OutputType * last,
                 const T &exemplar)
{
    typedef unsigned long long WideType; // type used to pack the Ts

    size_t ALIGNMENT_BOUNDARY = 128; // begin copying blocks at this byte boundary

    size_t n = last - first;

    WideType wide_exemplar;
    for(int i = 0; i < sizeof(WideType)/sizeof(T); i++)
        reinterpret_cast<T *>(&wide_exemplar)[i] = exemplar;

    OutputType * block_first = std::min(first + n,   thrust::detail::util::align_up(first, ALIGNMENT_BOUNDARY));
    OutputType * block_last  = std::max(block_first, thrust::detail::util::align_down(last, sizeof(WideType)));

    size_t pass1 = block_first - first;
    size_t pass2 = (block_last - block_first) / (sizeof(WideType) / sizeof(T));
    size_t pass3 = last - block_last;

    thrust::detail::device::cuda::vectorize(pass1, detail::fill_functor<OutputType *, T       >(first, exemplar));
    thrust::detail::device::cuda::vectorize(pass2, detail::fill_functor<WideType  *, WideType>((WideType *) block_first, wide_exemplar));
    thrust::detail::device::cuda::vectorize(pass3, detail::fill_functor<OutputType *, T       >(block_last, exemplar));
}

template<typename ForwardIterator, typename T>
  void fill(ForwardIterator first,
            ForwardIterator last,
            const T &exemplar,
            thrust::detail::false_type)
{
    detail::fill_functor<ForwardIterator, T> func(first, exemplar); 
    thrust::detail::device::cuda::vectorize(last - first, func);
}

template<typename ForwardIterator, typename T>
  void fill(ForwardIterator first,
            ForwardIterator last,
            const T &exemplar,
            thrust::detail::true_type)
{
    typedef typename thrust::iterator_traits<ForwardIterator>::value_type OutputType;

    
    if ( thrust::detail::util::is_aligned<OutputType>(thrust::raw_pointer_cast(&*first)) )
    {
        detail::wide_fill(thrust::raw_pointer_cast(&*first), thrust::raw_pointer_cast(&*last), exemplar);
    }
    else
    {
        fill(first, last, exemplar, thrust::detail::false_type());
    }
}

} // end detail

template<typename ForwardIterator, typename T>
  void fill(ForwardIterator first,
            ForwardIterator last,
            const T &exemplar)
{
  typedef typename thrust::iterator_traits<ForwardIterator>::value_type      OutputType;
  typedef typename thrust::iterator_traits<ForwardIterator>::difference_type IndexType;

  // we're compiling with nvcc, launch a kernel
  const bool use_wide_fill = thrust::detail::is_trivial_iterator<ForwardIterator>::value
      && (sizeof(OutputType) == 1 || sizeof(OutputType) == 2 || sizeof(OutputType) == 4);
  detail::fill(first, last, exemplar, thrust::detail::integral_constant<bool, use_wide_fill>());
}

} // end namespace device

} // end namespace detail

} // end namespace thrust

#else

///////////////////////////
// C++ (only) definition //
///////////////////////////

#include <thrust/copy.h>
#include <thrust/device_ptr.h>
#include <thrust/distance.h>

namespace thrust<commit_msg>Implement device fill with device generate.<commit_after>/*
 *  Copyright 2008-2009 NVIDIA Corporation
 *
 *  Licensed under the Apache License, Version 2.0 (the "License");
 *  you may not use this file except in compliance with the License.
 *  You may obtain a copy of the License at
 *
 *      http://www.apache.org/licenses/LICENSE-2.0
 *
 *  Unless required by applicable law or agreed to in writing, software
 *  distributed under the License is distributed on an "AS IS" BASIS,
 *  WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
 *  See the License for the specific language governing permissions and
 *  limitations under the License.
 */


/*! \file fill.inl
 *  \brief Inline file for fill.h.
 */

#ifdef __CUDACC__

/////////////////////
// NVCC definition //
/////////////////////

#include <thrust/detail/util/align.h>
#include <thrust/detail/device/generate.h>
#include <thrust/iterator/iterator_traits.h>
#include <thrust/detail/type_traits.h>
#include <thrust/extrema.h>

namespace thrust
{

namespace detail
{

namespace device
{

namespace detail
{

template <typename T>
struct fill_functor
{
  T exemplar;

  fill_functor(T _exemplar) 
    : exemplar(_exemplar) {}

  __device__
  T operator()(void)
  { 
    return exemplar;
  }
}; // end fill_functor


template<typename OutputType, typename T>
  void wide_fill(OutputType * first,
                 OutputType * last,
                 const T &exemplar)
{
  typedef unsigned long long WideType; // type used to pack the Ts

  size_t ALIGNMENT_BOUNDARY = 128; // begin copying blocks at this byte boundary

  size_t n = last - first;

  WideType wide_exemplar;
  for(int i = 0; i < sizeof(WideType)/sizeof(T); i++)
      reinterpret_cast<T *>(&wide_exemplar)[i] = exemplar;

  OutputType * block_first = thrust::min(first + n,   thrust::detail::util::align_up(first, ALIGNMENT_BOUNDARY));
  OutputType * block_last  = thrust::max(block_first, thrust::detail::util::align_down(last, sizeof(WideType)));

  thrust::detail::device::generate(first, block_first, fill_functor<OutputType>(exemplar));
  thrust::detail::device::generate(reinterpret_cast<WideType*>(block_first),
                                   reinterpret_cast<WideType*>(block_last),
                                   fill_functor<WideType>(wide_exemplar));
  thrust::detail::device::generate(block_last, last, fill_functor<OutputType>(exemplar));
}

template<typename ForwardIterator, typename T>
  void fill(ForwardIterator first,
            ForwardIterator last,
            const T &exemplar,
            thrust::detail::false_type)
{
  detail::fill_functor<T> func(exemplar); 
  thrust::detail::device::generate(first, last, func);
}

template<typename ForwardIterator, typename T>
  void fill(ForwardIterator first,
            ForwardIterator last,
            const T &exemplar,
            thrust::detail::true_type)
{
    typedef typename thrust::iterator_traits<ForwardIterator>::value_type OutputType;

    
    if ( thrust::detail::util::is_aligned<OutputType>(thrust::raw_pointer_cast(&*first)) )
    {
        detail::wide_fill(thrust::raw_pointer_cast(&*first), thrust::raw_pointer_cast(&*last), exemplar);
    }
    else
    {
        fill(first, last, exemplar, thrust::detail::false_type());
    }
}

} // end detail

template<typename ForwardIterator, typename T>
  void fill(ForwardIterator first,
            ForwardIterator last,
            const T &exemplar)
{
  typedef typename thrust::iterator_traits<ForwardIterator>::value_type      OutputType;
  typedef typename thrust::iterator_traits<ForwardIterator>::difference_type IndexType;

  // we're compiling with nvcc, launch a kernel
  const bool use_wide_fill = thrust::detail::is_trivial_iterator<ForwardIterator>::value
      && thrust::detail::has_trivial_copy<OutputType>::value
      && (sizeof(OutputType) == 1 || sizeof(OutputType) == 2 || sizeof(OutputType) == 4);
  detail::fill(first, last, exemplar, thrust::detail::integral_constant<bool, use_wide_fill>());
}

} // end namespace device

} // end namespace detail

} // end namespace thrust

#else

///////////////////////////
// C++ (only) definition //
///////////////////////////

#include <thrust/copy.h>
#include <thrust/device_ptr.h>
#include <thrust/distance.h>

namespace thrust
{'fraction_non_alphanumeric': 0.13559322033898305, 'fraction_numerical': 0.0049388259063868, 'mean_word_length': 3.4483275087368948, 'pattern_counts': {'":': 0, '<': 57, '<?xml version=': 0, '>': 55, 'https://': 0, 'lorem ipsum': 0, 'www.': 2, 'xml': 0}, 'pii_count': 0, 'substrings_counts': 0, 'word_list_counts': {'cursed_substrings.json': 2, 'profanity_word_list.json': 0, 'sexual_word_list.json': 0, 'zh_pornsignals.json': 0}}
{'dir': 'git-commits-cleaned', 'id': '4014442', 'n_tokens_mistral': 2716, 'n_tokens_neox': 2618, 'n_words': 654}