dylanyunlong/project_6_89d52222
1
0
Fork 0
Code Issues Pull Requests Actions Projects Releases Wiki Activity
Files
617e11044f51221781f62653bcc3bd7a3e17155f
project_6_89d52222/cccl_upstream/thrust/thrust/detail/merge.inl
EngineX CI 56fd68e7dd [INFRA] Import NVIDIA/CCCL upstream as optimization reference library 
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
2026-07-30 09:35:51 +00:00

276 lines
9.0 KiB
C++
Raw Blame History

// SPDX-FileCopyrightText: Copyright (c) 2008-2013, NVIDIA Corporation. All rights reserved.
// SPDX-License-Identifier: Apache-2.0
#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/iterator/iterator_traits.h>
#include <thrust/merge.h>
#include <thrust/system/detail/generic/select_system.h>
// Include all active backend system implementations (generic, sequential, host and device)
#include <thrust/system/detail/generic/merge.h>
#include <thrust/system/detail/sequential/merge.h>
#include __THRUST_HOST_SYSTEM_ALGORITH_DETAIL_HEADER_INCLUDE(merge.h)
#include __THRUST_DEVICE_SYSTEM_ALGORITH_DETAIL_HEADER_INCLUDE(merge.h)
// Some build systems need a hint to know which files we could include
#if 0
# include <thrust/system/cpp/detail/merge.h>
# include <thrust/system/cuda/detail/merge.h>
# include <thrust/system/omp/detail/merge.h>
# include <thrust/system/tbb/detail/merge.h>
#endif
THRUST_NAMESPACE_BEGIN
_CCCL_EXEC_CHECK_DISABLE
template <typename DerivedPolicy, typename InputIterator1, typename InputIterator2, typename OutputIterator>
_CCCL_HOST_DEVICE OutputIterator merge(
const thrust::detail::execution_policy_base<DerivedPolicy>& exec,
InputIterator1 first1,
InputIterator1 last1,
InputIterator2 first2,
InputIterator2 last2,
OutputIterator result)
{
_CCCL_NVTX_RANGE_SCOPE("thrust::merge");
using thrust::system::detail::generic::merge;
return merge(thrust::detail::derived_cast(thrust::detail::strip_const(exec)), first1, last1, first2, last2, result);
} // end merge()
_CCCL_EXEC_CHECK_DISABLE
template <typename DerivedPolicy,
typename InputIterator1,
typename InputIterator2,
typename OutputIterator,
typename StrictWeakCompare>
_CCCL_HOST_DEVICE OutputIterator merge(
const thrust::detail::execution_policy_base<DerivedPolicy>& exec,
InputIterator1 first1,
InputIterator1 last1,
InputIterator2 first2,
InputIterator2 last2,
OutputIterator result,
StrictWeakCompare comp)
{
_CCCL_NVTX_RANGE_SCOPE("thrust::merge");
using thrust::system::detail::generic::merge;
return merge(
thrust::detail::derived_cast(thrust::detail::strip_const(exec)), first1, last1, first2, last2, result, comp);
} // end merge()
_CCCL_EXEC_CHECK_DISABLE
template <typename DerivedPolicy,
typename InputIterator1,
typename InputIterator2,
typename InputIterator3,
typename InputIterator4,
typename OutputIterator1,
typename OutputIterator2>
_CCCL_HOST_DEVICE ::cuda::std::pair<OutputIterator1, OutputIterator2> merge_by_key(
const thrust::detail::execution_policy_base<DerivedPolicy>& exec,
InputIterator1 keys_first1,
InputIterator1 keys_last1,
InputIterator2 keys_first2,
InputIterator2 keys_last2,
InputIterator3 values_first1,
InputIterator4 values_first2,
OutputIterator1 keys_result,
OutputIterator2 values_result)
{
_CCCL_NVTX_RANGE_SCOPE("thrust::merge_by_key");
using thrust::system::detail::generic::merge_by_key;
return merge_by_key(
thrust::detail::derived_cast(thrust::detail::strip_const(exec)),
keys_first1,
keys_last1,
keys_first2,
keys_last2,
values_first1,
values_first2,
keys_result,
values_result);
} // end merge_by_key()
_CCCL_EXEC_CHECK_DISABLE
template <typename DerivedPolicy,
typename InputIterator1,
typename InputIterator2,
typename InputIterator3,
typename InputIterator4,
typename OutputIterator1,
typename OutputIterator2,
typename Compare>
_CCCL_HOST_DEVICE ::cuda::std::pair<OutputIterator1, OutputIterator2> merge_by_key(
const thrust::detail::execution_policy_base<DerivedPolicy>& exec,
InputIterator1 keys_first1,
InputIterator1 keys_last1,
InputIterator2 keys_first2,
InputIterator2 keys_last2,
InputIterator3 values_first1,
InputIterator4 values_first2,
OutputIterator1 keys_result,
OutputIterator2 values_result,
Compare comp)
{
_CCCL_NVTX_RANGE_SCOPE("thrust::merge_by_key");
using thrust::system::detail::generic::merge_by_key;
return merge_by_key(
thrust::detail::derived_cast(thrust::detail::strip_const(exec)),
keys_first1,
keys_last1,
keys_first2,
keys_last2,
values_first1,
values_first2,
keys_result,
values_result,
comp);
} // end merge_by_key()
template <typename InputIterator1, typename InputIterator2, typename OutputIterator, typename StrictWeakOrdering>
OutputIterator
merge(InputIterator1 first1,
InputIterator1 last1,
InputIterator2 first2,
InputIterator2 last2,
OutputIterator result,
StrictWeakOrdering comp)
{
_CCCL_NVTX_RANGE_SCOPE("thrust::merge");
using thrust::system::detail::generic::select_system;
using System1 = typename thrust::iterator_system<InputIterator1>::type;
using System2 = typename thrust::iterator_system<InputIterator2>::type;
using System3 = typename thrust::iterator_system<OutputIterator>::type;
System1 system1;
System2 system2;
System3 system3;
return thrust::merge(select_system(system1, system2, system3), first1, last1, first2, last2, result, comp);
} // end merge()
template <typename InputIterator1, typename InputIterator2, typename OutputIterator>
OutputIterator
merge(InputIterator1 first1, InputIterator1 last1, InputIterator2 first2, InputIterator2 last2, OutputIterator result)
{
_CCCL_NVTX_RANGE_SCOPE("thrust::merge");
using thrust::system::detail::generic::select_system;
using System1 = typename thrust::iterator_system<InputIterator1>::type;
using System2 = typename thrust::iterator_system<InputIterator2>::type;
using System3 = typename thrust::iterator_system<OutputIterator>::type;
System1 system1;
System2 system2;
System3 system3;
return thrust::merge(select_system(system1, system2, system3), first1, last1, first2, last2, result);
} // end merge()
template <typename InputIterator1,
typename InputIterator2,
typename InputIterator3,
typename InputIterator4,
typename OutputIterator1,
typename OutputIterator2,
typename StrictWeakOrdering>
::cuda::std::pair<OutputIterator1, OutputIterator2> merge_by_key(
InputIterator1 keys_first1,
InputIterator1 keys_last1,
InputIterator2 keys_first2,
InputIterator2 keys_last2,
InputIterator3 values_first1,
InputIterator4 values_first2,
OutputIterator1 keys_result,
OutputIterator2 values_result,
StrictWeakOrdering comp)
{
_CCCL_NVTX_RANGE_SCOPE("thrust::merge_by_key");
using thrust::system::detail::generic::select_system;
using System1 = typename thrust::iterator_system<InputIterator1>::type;
using System2 = typename thrust::iterator_system<InputIterator2>::type;
using System3 = typename thrust::iterator_system<InputIterator3>::type;
using System4 = typename thrust::iterator_system<InputIterator4>::type;
using System5 = typename thrust::iterator_system<OutputIterator1>::type;
using System6 = typename thrust::iterator_system<OutputIterator2>::type;
System1 system1;
System2 system2;
System3 system3;
System4 system4;
System5 system5;
System6 system6;
return thrust::merge_by_key(
select_system(system1, system2, system3, system4, system5, system6),
keys_first1,
keys_last1,
keys_first2,
keys_last2,
values_first1,
values_first2,
keys_result,
values_result,
comp);
} // end merge_by_key()
template <typename InputIterator1,
typename InputIterator2,
typename InputIterator3,
typename InputIterator4,
typename OutputIterator1,
typename OutputIterator2>
::cuda::std::pair<OutputIterator1, OutputIterator2> merge_by_key(
InputIterator1 keys_first1,
InputIterator1 keys_last1,
InputIterator2 keys_first2,
InputIterator2 keys_last2,
InputIterator3 values_first1,
InputIterator4 values_first2,
OutputIterator1 keys_result,
OutputIterator2 values_result)
{
_CCCL_NVTX_RANGE_SCOPE("thrust::merge_by_key");
using thrust::system::detail::generic::select_system;
using System1 = typename thrust::iterator_system<InputIterator1>::type;
using System2 = typename thrust::iterator_system<InputIterator2>::type;
using System3 = typename thrust::iterator_system<InputIterator3>::type;
using System4 = typename thrust::iterator_system<InputIterator4>::type;
using System5 = typename thrust::iterator_system<OutputIterator1>::type;
using System6 = typename thrust::iterator_system<OutputIterator2>::type;
System1 system1;
System2 system2;
System3 system3;
System4 system4;
System5 system5;
System6 system6;
return thrust::merge_by_key(
select_system(system1, system2, system3, system4, system5, system6),
keys_first1,
keys_last1,
keys_first2,
keys_last2,
values_first1,
values_first2,
keys_result,
values_result);
} // end merge_by_key()
THRUST_NAMESPACE_END
Reference in New Issue View Git Blame Copy Permalink