<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 gather.inl
 *  \brief Inline file for gather.h
 */

#pragma once

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

namespace thrust
{

namespace detail
{

namespace device
{

namespace detail
{

//////////////
// Functors //
//////////////
template <typename ForwardIterator, typename InputIterator, typename RandomAccessIterator>
struct gather_functor
{
    ForwardIterator first;
    InputIterator map;
    RandomAccessIterator input;

    gather_functor(ForwardIterator _first, InputIterator _map, RandomAccessIterator _input)
      : first(_first), map(_map), input(_input) {}
  
    template <typename IntegerType>
        __device__
        void operator()(const IntegerType& i)
        { 
            //output[i] = input[map[i]];
            thrust::detail::device::dereference(first, i) = thrust::detail::device::dereference(input, thrust::detail::device::dereference(map, i)); 
        }
}; // end gather_functor

template <typename ForwardIterator, typename InputIterator1, typename InputIterator2, typename RandomAccessIterator, typename Predicate>
struct gather_if_functor
{
    ForwardIterator first;
    InputIterator1 map;
    InputIterator2 stencil;
    RandomAccessIterator input;
    Predicate pred;

    gather_if_functor(ForwardIterator _first, InputIterator1 _map, InputIterator2 _stencil, RandomAccessIterator _input, Predicate _pred)
      : first(_first), map(_map), stencil(_stencil), input(_input), pred(_pred) {}
  
    template <typename IntegerType>
        __device__
        void operator()(const IntegerType& i)
        { 
            //if(pred(stencil[i])
            //    output[i] = input[map[i]];
            if(pred(thrust::detail::device::dereference(stencil, i)))
                thrust::detail::device::dereference(first, i) = thrust::detail::device::dereference(input, thrust::detail::device::dereference(map, i)); 
        }
}; // end gather_functor

} // end detail


template<typename ForwardIterator,
         typename InputIterator,
         typename RandomAccessIterator>
  void gather(ForwardIterator first,
              ForwardIterator last,
              InputIterator  map,
              RandomAccessIterator input)
{
    detail::gather_functor<ForwardIterator, InputIterator, RandomAccessIterator> func(first, map, input);
    thrust::detail::device::cuda::vectorize(last - first, func);
} // end gather()


template<typename ForwardIterator,
         typename InputIterator1,
         typename InputIterator2,
         typename RandomAccessIterator,
         typename Predicate>
  void gather_if(ForwardIterator first,
                 ForwardIterator last,
                 InputIterator1 map,
                 InputIterator2 stencil,
                 RandomAccessIterator input,
                 Predicate pred)
{
    detail::gather_if_functor<ForwardIterator, InputIterator1, InputIterator2, RandomAccessIterator, Predicate> func(first, map, stencil, input, pred);
    thrust::detail::device::cuda::vectorize(last - first, func);
} // end gather_if()


} // end namespace device

} // end namespace detail

} // end namespace thrust
<commit_msg>Implement device gather with device for_each.<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 gather.inl
 *  \brief Inline file for gather.h
 */

#pragma once

#include <thrust/detail/device/for_each.h>
#include <thrust/detail/device/dereference.h>
#include <thrust/distance.h>

namespace thrust
{

namespace detail
{

namespace device
{

namespace detail
{


template <typename RandomAccessIterator>
struct gather_functor
{
  RandomAccessIterator input;

  gather_functor(RandomAccessIterator _input)
    : input(_input) {}
  
  template <typename Tuple>
  __device__
  void operator()(Tuple t)
  { 
    thrust::get<0>(t) = thrust::detail::device::dereference(input, thrust::get<1>(t)); 
  }
}; // end gather_functor


template <typename RandomAccessIterator, typename Predicate>
struct gather_if_functor
{
  RandomAccessIterator input;
  Predicate pred;

  gather_if_functor(RandomAccessIterator _input, Predicate _pred)
    : input(_input), pred(_pred) {}
  
  template <typename Tuple>
  __device__
  void operator()(Tuple t)
  { 
    if(pred(thrust::get<2>(t)))
      thrust::get<0>(t) = thrust::detail::device::dereference(input, thrust::get<1>(t)); 
  }
}; // end gather_functor


} // end detail


template<typename ForwardIterator,
         typename InputIterator,
         typename RandomAccessIterator>
  void gather(ForwardIterator first,
              ForwardIterator last,
              InputIterator  map,
              RandomAccessIterator input)
{
  detail::gather_functor<RandomAccessIterator> func(input);
  thrust::detail::device::for_each(thrust::make_zip_iterator(thrust::make_tuple(first, map)),
                                   thrust::make_zip_iterator(thrust::make_tuple(last,  map + thrust::distance(first, last))),
                                   func);
} // end gather()


template<typename ForwardIterator,
         typename InputIterator1,
         typename InputIterator2,
         typename RandomAccessIterator,
         typename Predicate>
  void gather_if(ForwardIterator first,
                 ForwardIterator last,
                 InputIterator1 map,
                 InputIterator2 stencil,
                 RandomAccessIterator input,
                 Predicate pred)
{
  detail::gather_if_functor<RandomAccessIterator, Predicate> func(input, pred);
  thrust::detail::device::for_each(thrust::make_zip_iterator(thrust::make_tuple(first, map, stencil)),
                                   thrust::make_zip_iterator(thrust::make_tuple(last,  map + thrust::distance(first, last), stencil + thrust::distance(first, last))),
                                   func);
} // end gather_if()


} // end namespace device

} // end namespace detail

} // end namespace thrust

{'fraction_non_alphanumeric': 0.10092385218365062, 'fraction_numerical': 0.006298992161254199, 'mean_word_length': 3.2203189604252804, 'pattern_counts': {'":': 0, '<': 30, '<?xml version=': 0, '>': 30, 'https://': 0, 'lorem ipsum': 0, 'www.': 2, 'xml': 0}, 'pii_count': 0, 'substrings_counts': 0, 'word_list_counts': {'cursed_substrings.json': 4, 'profanity_word_list.json': 0, 'sexual_word_list.json': 0, 'zh_pornsignals.json': 0}}
{'dir': 'git-commits-cleaned', 'id': '2340427', 'n_tokens_mistral': 1988, 'n_tokens_neox': 1923, 'n_words': 521}