56fd68e7dd
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
146 lines
4.0 KiB
C++
146 lines
4.0 KiB
C++
// SPDX-FileCopyrightText: Copyright (c) 2008-2018, NVIDIA Corporation. All rights reserved.
|
|
// SPDX-License-Identifier: Apache-2.0
|
|
|
|
/*! \file
|
|
* \brief An allocator which creates new elements in memory accessible by
|
|
* devices.
|
|
*/
|
|
|
|
#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/device_ptr.h>
|
|
#include <thrust/mr/allocator.h>
|
|
#include <thrust/mr/device_memory_resource.h>
|
|
|
|
#include <cuda/std/limits>
|
|
|
|
THRUST_NAMESPACE_BEGIN
|
|
|
|
/** \addtogroup allocators Allocators
|
|
* \ingroup memory_management
|
|
* \{
|
|
*/
|
|
|
|
/*! Memory resource adaptor that turns any memory resource that returns a fancy
|
|
* with the same tag as \p device_ptr, and adapts it to a resource that returns
|
|
* a \p device_ptr.
|
|
*
|
|
* \verbatim embed:rst:leading-asterisk
|
|
* .. versionadded:: 2.2.0
|
|
* \endverbatim
|
|
*/
|
|
template <typename Upstream>
|
|
class device_ptr_memory_resource final : public thrust::mr::memory_resource<device_ptr<void>>
|
|
{
|
|
using upstream_ptr = typename Upstream::pointer;
|
|
|
|
public:
|
|
/*! Initialize the adaptor with the global instance of the upstream resource. Obtains
|
|
* the global instance by calling \p get_global_resource.
|
|
*
|
|
* \verbatim embed:rst:leading-asterisk
|
|
* .. versionadded:: 2.2.0
|
|
* \endverbatim
|
|
*/
|
|
_CCCL_HOST device_ptr_memory_resource()
|
|
: m_upstream(mr::get_global_resource<Upstream>())
|
|
{}
|
|
|
|
/*! Initialize the adaptor with an upstream resource.
|
|
*
|
|
* \param upstream the upstream memory resource to adapt.
|
|
*
|
|
* \verbatim embed:rst:leading-asterisk
|
|
* .. versionadded:: 2.2.0
|
|
* \endverbatim
|
|
*/
|
|
_CCCL_HOST device_ptr_memory_resource(Upstream* upstream)
|
|
: m_upstream(upstream)
|
|
{}
|
|
|
|
[[nodiscard]] _CCCL_HOST pointer
|
|
do_allocate(std::size_t bytes, std::size_t alignment = THRUST_MR_DEFAULT_ALIGNMENT) override
|
|
{
|
|
return pointer(m_upstream->do_allocate(bytes, alignment).get());
|
|
}
|
|
|
|
_CCCL_HOST void do_deallocate(pointer p, std::size_t bytes, std::size_t alignment) override
|
|
{
|
|
m_upstream->do_deallocate(upstream_ptr(p.get()), bytes, alignment);
|
|
}
|
|
|
|
private:
|
|
Upstream* m_upstream;
|
|
};
|
|
|
|
/*! \brief An allocator which creates new elements in memory accessible by
|
|
* devices.
|
|
*
|
|
* \see https://en.cppreference.com/w/cpp/named_req/Allocator
|
|
*
|
|
* \verbatim embed:rst:leading-asterisk
|
|
* .. versionadded:: 2.2.0
|
|
* \endverbatim
|
|
*/
|
|
template <typename T>
|
|
class device_allocator
|
|
: public thrust::mr::stateless_resource_allocator<T, device_ptr_memory_resource<device_memory_resource>>
|
|
{
|
|
using base = thrust::mr::stateless_resource_allocator<T, device_ptr_memory_resource<device_memory_resource>>;
|
|
|
|
public:
|
|
/*! The \p rebind metafunction provides the type of a \p device_allocator
|
|
* instantiated with another type.
|
|
*
|
|
* \tparam U the other type to use for instantiation.
|
|
*
|
|
* \verbatim embed:rst:leading-asterisk
|
|
* .. versionadded:: 2.2.0
|
|
* \endverbatim
|
|
*/
|
|
template <typename U>
|
|
struct rebind
|
|
{
|
|
/*! The alias \p other gives the type of the rebound \p device_allocator.
|
|
*
|
|
* \verbatim embed:rst:leading-asterisk
|
|
* .. versionadded:: 2.2.0
|
|
* \endverbatim
|
|
*/
|
|
using other = device_allocator<U>;
|
|
};
|
|
|
|
/*! Default constructor has no effect. */
|
|
device_allocator() = default;
|
|
|
|
/*! Copy constructor has no effect. */
|
|
_CCCL_HOST_DEVICE device_allocator(const device_allocator& other)
|
|
: base(other)
|
|
{}
|
|
|
|
/*! Constructor from other \p device_allocator has no effect. */
|
|
template <typename U>
|
|
_CCCL_HOST_DEVICE device_allocator(const device_allocator<U>& other)
|
|
: base(other)
|
|
{}
|
|
|
|
device_allocator& operator=(const device_allocator&) = default;
|
|
|
|
/*! Destructor has no effect. */
|
|
~device_allocator() = default;
|
|
};
|
|
|
|
/*! \} // allocators
|
|
*/
|
|
|
|
THRUST_NAMESPACE_END
|