#include <thrust/detail/config.h>
#include <thrust/detail/seq.h>
#include <thrust/detail/temporary_array.h>
#include <thrust/detail/type_traits.h>
#include <thrust/execution_policy.h>
#include <thrust/reduce.h>
#include <thrust/system/cuda/detail/bulk.h>
#include <thrust/system/cuda/detail/decomposition.h>
#include <thrust/system/cuda/detail/execute_on_stream.h>
#include <thrust/system/cuda/detail/execution_policy.h>
#include <goofit/detail/ThrustOverrideConfig.h>

Classes
struct	goofit_policy

Macros
#define	OR_GROUPSIZE 128

#define	OR_GRAINSIZE 7

Functions
template<typename InputIterator , typename UnaryFunction , typename OutputType , typename BinaryFunction >
__host__ __device__ OutputType	transform_reduce (goofit_policy &exec, InputIterator first, InputIterator last, UnaryFunction unary_op, OutputType init, BinaryFunction binary_op)

template<typename InputIterator , typename OutputType , typename BinaryFunction >
__host__ __device__ OutputType	reduce (goofit_policy &exec, InputIterator first, InputIterator last, OutputType init, BinaryFunction binary_op)

Macro Definition Documentation

◆ OR_GRAINSIZE

#define OR_GRAINSIZE 7

Definition at line 25 of file ThrustOverride.h.

Referenced by reduce().

◆ OR_GROUPSIZE

#define OR_GROUPSIZE 128

Definition at line 21 of file ThrustOverride.h.

Referenced by reduce().

Function Documentation

◆ reduce()

template<typename InputIterator , typename OutputType , typename BinaryFunction >

__host__ __device__ OutputType reduce	(	goofit_policy &	exec,
		InputIterator	first,
		InputIterator	last,
		OutputType	init,
		BinaryFunction	binary_op
	)

Definition at line 47 of file ThrustOverride.h.

References OR_GRAINSIZE, and OR_GROUPSIZE.

Referenced by getData(), main(), and transform_reduce().

                                                                                                                 {
     typedef typename thrust::iterator_difference<InputIterator>::type size_type;
 
     const size_type n = last - first;
 
     if(n <= 0)
         return init;
 
     // We are no longer expecting a default stream to be passed
     cudaStream_t s = stream(thrust::detail::derived_cast(exec));
 
     // OR_GROUPSIZE=128
     // OR_GRAINSIZE=7
     // printf ("groupsize:%i grainsize:%i\n", OR_GROUPSIZE, OR_GRAINSIZE);
     const size_type groupsize = OR_GROUPSIZE;
     // const size_type groupsize = 128;
     // const size_type grainsize = 7;
     const size_type grainsize    = OR_GRAINSIZE;
     const size_type tile_size    = groupsize * grainsize;
     const size_type num_tiles    = (n + tile_size - 1) / tile_size;
     const size_type subscription = 10;
 
     thrust::system::cuda::detail::bulk_::concurrent_group<thrust::system::cuda::detail::bulk_::agent<grainsize>,
                                                           groupsize>
         g;
 
     const size_type num_groups = thrust::min<size_type>(subscription * g.hardware_concurrency(), num_tiles);
 
     thrust::system::cuda::detail::aligned_decomposition<size_type> decomp(n, num_groups, tile_size);
 
     thrust::detail::temporary_array<OutputType, goofit_policy> partial_sums(exec, decomp.size());
 
     // reduce into partial sums
     thrust::system::cuda::detail::bulk_::async(thrust::system::cuda::detail::bulk_::par(s, g, decomp.size()),
                                                thrust::system::cuda::detail::reduce_detail::reduce_partitions(),
                                                thrust::system::cuda::detail::bulk_::root.this_exec,
                                                first,
                                                decomp,
                                                partial_sums.begin(),
                                                init,
                                                binary_op)
         .wait();
 
     if(partial_sums.size() > 1) {
         // reduce the partial sums
         thrust::system::cuda::detail::bulk_::async(thrust::system::cuda::detail::bulk_::par(s, g, 1),
                                                    thrust::system::cuda::detail::reduce_detail::reduce_partitions(),
                                                    thrust::system::cuda::detail::bulk_::root.this_exec,
                                                    partial_sums.begin(),
                                                    partial_sums.end(),
                                                    partial_sums.begin(),
                                                    binary_op);
     } // end if
 
     return get_value(exec, &partial_sums[0]);
 } // end goofit_reduce()

◆ transform_reduce()

template<typename InputIterator , typename UnaryFunction , typename OutputType , typename BinaryFunction >

__host__ __device__ OutputType transform_reduce	(	goofit_policy &	exec,
		InputIterator	first,
		InputIterator	last,
		UnaryFunction	unary_op,
		OutputType	init,
		BinaryFunction	binary_op
	)

Definition at line 32 of file ThrustOverride.h.

References reduce().

                                                                           {
     // printf ("goofit transform_reduce\n");
     thrust::transform_iterator<UnaryFunction, InputIterator, OutputType> xfrm_first(first, unary_op);
     thrust::transform_iterator<UnaryFunction, InputIterator, OutputType> xfrm_last(last, unary_op);
 
     return thrust::reduce(exec, xfrm_first, xfrm_last, init, binary_op);
 } // end transform_reduce()

Classes