CCCL (CUDA C++ Core Libraries) provides: - CUB: device/block/warp-level GPU primitives (reduce, scan, sort, topk) - Thrust: high-level parallel algorithms (transform_reduce, sort, scan) - libcudacxx: CUDA C++ standard library (atomics, barriers, memory) - cudax: experimental features (memory resources, allocators) - Tuning policies: per-SM hardware-specific algorithm parameters Competition optimization vectors mapped to CCCL: - Output TPS (83% weight): warp_reduce, block_reduce, device_topk - Input TPS (14% weight): device_scan, block_load, prefetch - Cache TPS (3% weight): prefix caching strategy patterns - Memory (0.9 util): pooled/cached/buddy allocators Source: https://github.com/NVIDIA/cccl (shallow clone, HEAD only) License: Apache-2.0
165 lines
4.8 KiB
C++
165 lines
4.8 KiB
C++
/*! \file thrust/zip_function.h
|
|
* \brief Adaptor type that turns an N-ary function object into one that takes
|
|
* a tuple of size N so it can easily be used with algorithms taking zip
|
|
* iterators
|
|
*/
|
|
|
|
#pragma once
|
|
|
|
#include <thrust/detail/config.h>
|
|
|
|
#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)
|
|
# pragma GCC system_header
|
|
#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG)
|
|
# pragma clang system_header
|
|
#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
|
|
# pragma system_header
|
|
#endif // no system header
|
|
|
|
#include <thrust/detail/type_deduction.h>
|
|
|
|
#include <cuda/__functional/address_stability.h>
|
|
#include <cuda/std/__type_traits/decay.h>
|
|
#include <cuda/std/__utility/forward.h>
|
|
#include <cuda/std/__utility/move.h>
|
|
#include <cuda/std/tuple>
|
|
|
|
THRUST_NAMESPACE_BEGIN
|
|
|
|
/*! \addtogroup function_objects Function Objects
|
|
* \{
|
|
*/
|
|
|
|
/*! \addtogroup function_object_adaptors Function Object Adaptors
|
|
* \ingroup function_objects
|
|
* \{
|
|
*/
|
|
|
|
/*! \p zip_function is a function object that allows the easy use of N-ary
|
|
* function objects with \p zip_iterators without redefining them to take a
|
|
* \p tuple instead of N arguments.
|
|
*
|
|
* This means that if a functor that takes 2 arguments which could be used with
|
|
* the \p transform function and \p device_iterators can be extended to take 3
|
|
* arguments and \p zip_iterators without rewriting the functor in terms of
|
|
* \p tuple.
|
|
*
|
|
* The \p make_zip_function convenience function is provided to avoid having
|
|
* to explicitly define the type of the functor when creating a \p zip_function,
|
|
* whic is especially helpful when using lambdas as the functor.
|
|
*
|
|
* \code
|
|
* #include <thrust/iterator/zip_iterator.h>
|
|
* #include <thrust/device_vector.h>
|
|
* #include <thrust/transform.h>
|
|
* #include <thrust/zip_function.h>
|
|
*
|
|
* struct SumTuple {
|
|
* float operator()(auto tup) const {
|
|
* return cuda::std::get<0>(tup) + cuda::std::get<1>(tup) + ::cuda::std::get<2>(tup);
|
|
* }
|
|
* };
|
|
* struct SumArgs {
|
|
* float operator()(float a, float b, float c) const {
|
|
* return a + b + c;
|
|
* }
|
|
* };
|
|
*
|
|
* int main() {
|
|
* thrust::device_vector<float> A{0.f, 1.f, 2.f};
|
|
* thrust::device_vector<float> B{1.f, 2.f, 3.f};
|
|
* thrust::device_vector<float> C{2.f, 3.f, 4.f};
|
|
* thrust::device_vector<float> D(3);
|
|
*
|
|
* auto begin = thrust::make_zip_iterator(A.begin(), B.begin(), C.begin());
|
|
* auto end = thrust::make_zip_iterator(A.end(), B.end(), C.end());
|
|
*
|
|
* // The following four invocations of transform are equivalent:
|
|
* // Transform with 3-tuple
|
|
* thrust::transform(begin, end, D.begin(), SumTuple{});
|
|
*
|
|
* // Transform with 3 parameters
|
|
* thrust::zip_function<SumArgs> adapted{};
|
|
* thrust::transform(begin, end, D.begin(), adapted);
|
|
*
|
|
* // Transform with 3 parameters with convenience function
|
|
* thrust::transform(begin, end, D.begin(), thrust::make_zip_function(SumArgs{}));
|
|
*
|
|
* // Transform with 3 parameters with convenience function and lambda
|
|
* thrust::transform(begin, end, D.begin(), thrust::make_zip_function([] (float a, float b, float c) {
|
|
* return a + b + c;
|
|
* }));
|
|
* return 0;
|
|
* }
|
|
* \endcode
|
|
*
|
|
* \see make_zip_function
|
|
* \see zip_iterator
|
|
*
|
|
* \verbatim embed:rst:leading-asterisk
|
|
* .. versionadded:: 2.2.0
|
|
* \endverbatim
|
|
*/
|
|
template <typename Function>
|
|
class zip_function
|
|
{
|
|
public:
|
|
//! Default constructs the contained function object.
|
|
zip_function() = default;
|
|
|
|
_CCCL_HOST_DEVICE zip_function(Function func)
|
|
: func(::cuda::std::move(func))
|
|
{}
|
|
|
|
template <typename Tuple>
|
|
_CCCL_HOST_DEVICE decltype(auto) operator()(Tuple&& args) const
|
|
{
|
|
return ::cuda::std::apply(func, ::cuda::std::forward<Tuple>(args));
|
|
}
|
|
|
|
//! Returns a reference to the underlying function.
|
|
_CCCL_HOST_DEVICE Function& underlying_function() const
|
|
{
|
|
return func;
|
|
}
|
|
|
|
private:
|
|
mutable Function func;
|
|
};
|
|
|
|
/*! \p make_zip_function creates a \p zip_function from a function object.
|
|
*
|
|
* \param fun The N-ary function object.
|
|
* \return A \p zip_function that takes a N-tuple.
|
|
*
|
|
* \see zip_function
|
|
*
|
|
* \verbatim embed:rst:leading-asterisk
|
|
* .. versionadded:: 2.2.0
|
|
* \endverbatim
|
|
*/
|
|
template <typename Function>
|
|
_CCCL_HOST_DEVICE zip_function<::cuda::std::decay_t<Function>> make_zip_function(Function&& fun)
|
|
{
|
|
using func_t = ::cuda::std::decay_t<Function>;
|
|
return zip_function<func_t>(THRUST_FWD(fun));
|
|
}
|
|
|
|
/*! \} // end function_object_adaptors
|
|
*/
|
|
|
|
/*! \} // end function_objects
|
|
*
|
|
* \verbatim embed:rst:leading-asterisk
|
|
* .. versionadded:: 2.2.0
|
|
* \endverbatim
|
|
*/
|
|
|
|
THRUST_NAMESPACE_END
|
|
|
|
_CCCL_BEGIN_NAMESPACE_CUDA
|
|
template <typename F>
|
|
struct proclaims_copyable_arguments<THRUST_NS_QUALIFIER::zip_function<F>> : proclaims_copyable_arguments<F>
|
|
{};
|
|
_CCCL_END_NAMESPACE_CUDA
|