- NVIDIA#1805: Add default constructors to
transform_output_iteratorandtransform_input_output_iterator. Thanks to Mark Harris (@harrism) for this contribution. - NVIDIA#1836: Enable constructions of vectors from
std::initializer_list.
- NVIDIA#1768: Fix type conversion warning in the
thrust::complexutilities. Thanks to Zishi Wu (@zishiwu123) for this contribution. - NVIDIA#1809: Fix some warnings about usage of
__host__functions in__device__code. - NVIDIA#1825: Fix Thrust's CMake install rules. Thanks to Robert Maynard (@robertmaynard) for this contribution.
- NVIDIA#1827: Fix
thrust::reduce_by_keywhen using non-default-initializable iterators. - NVIDIA#1832: Fix bug in device-side CDP
thrust::reducewhen using a large number of inputs.
- NVIDIA#1815: Update Thrust's libcu++ git submodule to version 1.8.1.
- NVIDIA#1841: Fix invalid code in execution policy documentation example. Thanks to Raphaël Frantz (@Eren121) for this contribution.
- NVIDIA#1848: Improve error messages when attempting to launch a kernel on a device that is not supported by compiled PTX versions. Thanks to Zahra Khatami (@zkhatami) for this contribution.
- NVIDIA#1855: Remove usage of deprecated CUDA error codes.
- Disable CDP parallelization of device-side invocations of Thrust algorithms on SM90+. The removal of device-side synchronization support in recent architectures makes Thrust's fork-join model unimplementable on device, so a serial implementation will be used instead. Host-side invocations of Thrust algorithms are not affected.
The Thrust 2.0.0 major release adds a dependency on libcu++ and contains several
breaking changes. These include new diagnostics when inspecting device-only
lambdas from the host, removal of the cub symlink in the Thrust repository
root, and removal of the deprecated THRUST_*_BACKEND macros. It also includes
several minor bugfixes and cleanups.
- NVIDIA#1605: Add libcu++ dependency.
- A suitable version of libcu++ is provided through
the
${THRUST_ROOT}/dependencies/libcudacxx/submodule. - Non-cmake users may need to add the libcu++ include path to their
builds (
-I ${THRUST_ROOT}/dependencies/libcudacxx/include/). - The Thrust CMake packages have been updated to add this include path.
- A suitable version of libcu++ is provided through
the
- NVIDIA#1605: The following macros are no longer defined by default.
They can be re-enabled by defining
THRUST_PROVIDE_LEGACY_ARCH_MACROS. These will be removed completely in a future release.THRUST_IS_HOST_CODE: Replace withNV_IF_TARGET.THRUST_IS_DEVICE_CODE: Replace withNV_IF_TARGET.THRUST_INCLUDE_HOST_CODE: Replace withNV_IF_TARGET.THRUST_INCLUDE_DEVICE_CODE: Replace withNV_IF_TARGET.THRUST_DEVICE_CODE: Replace withNV_IF_TARGET.
- NVIDIA#1661: Thrust's CUDA Runtime support macros have been updated to
support
NV_IF_TARGET. They are now defined consistently across all host/device compilation passes. This should not affect most usages of these macros, but may require changes for some edge cases.THRUST_RUNTIME_FUNCTION: Execution space annotations for functions that invoke CUDA Runtime APIs.- Old behavior:
- RDC enabled: Defined to
__host__ __device__ - RDC not enabled:
- NVCC host pass: Defined to
__host__ __device__ - NVCC device pass: Defined to
__host__
- NVCC host pass: Defined to
- RDC enabled: Defined to
- New behavior:
- RDC enabled: Defined to
__host__ __device__ - RDC not enabled: Defined to
__host__
- RDC enabled: Defined to
- Old behavior:
__THRUST_HAS_CUDART__: No change in behavior, but no longer used in Thrust. Provided for legacy support only. Legacy behavior:- RDC enabled: Defined to 1.
- RDC not enabled:
- NVCC host pass: Defined to 1.
- NVCC device pass: Defined to 0.
THRUST_RDC_ENABLED: New macro, may be combined withNV_IF_TARGETto replace most usages of__THRUST_HAS_CUDART__. Behavior:- RDC enabled: Macro is defined.
- RDC not enabled: Macro is not defined.
- NVIDIA#1701: Remove the
cubsymlink from the root of the Thrust repository.- This symlink caused issues in certain build environments (e.g. NVIDIA#1328).
- Builds that relied on this symlink will need to add the full CUB include
path (
-I ${THRUST_ROOT}/dependencies/cub). - CMake builds that use the Thrust packages via CPM,
add_subdirectory, orfind_packageare not affected.
- NVIDIA#1760: A compile-time error is now emitted when a
__device__-only lambda's return type is queried from host code (requires libcu++ ≥ 1.9.0).- Due to limitations in the CUDA programming model, the result of this query is unreliable, and will silently return an incorrect result. This leads to difficult to debug errors.
- When using libcu++ 1.9.0, an error will be emitted with information about
work-arounds:
- Use a named function object with a
__device__-only implementation ofoperator(). - Use a
__host__ __device__lambda. - Use
cuda::proclaim_return_type(Added in libcu++ 1.9.0)
- Use a named function object with a
- NVIDIA#1761: Removed support for deprecated
THRUST_DEVICE_BACKENDandTHRUST_HOST_BACKENDmacros. TheTHRUST_DEVICE_SYSTEMandTHRUST_HOST_SYSTEMmacros should be used instead.
- NVIDIA#1605: Fix some execution space warnings in the allocator library.
- NVIDIA#1683: Fix bug in
iterator_category_to_traversalmetafunctions. - NVIDIA#1715: Add missing
__thrust_exec_check_disable__annotation tothrust::make_zip_function. Thanks to @mfbalin for this contribution. - NVIDIA#1722: Remove CUDA-specific error handler from code that may be executed on non-CUDA backends. Thanks to @dkolsen-pgi for this contribution.
- NVIDIA#1756: Fix
copy_iffor output iterators that don't support copy assignment. Thanks for @mfbalin for this contribution.
- NVIDIA#1605: Removed special case code for unsupported CUDA architectures.
- NVIDIA#1605: Replace several usages of
__CUDA_ARCH__with<nv/target>to handle host/device code divergence. - NVIDIA#1752: Remove a leftover merge conflict from a documentation file. Thanks to @tabedzki for this contribution.
Thrust 1.17.2 is a minor bugfix release that provides an updated version of CUB.
Thrust 1.17.1 is a minor bugfix release that provides an updated version of CUB.
Thrust 1.17.0 is the final minor release of the 1.X series. This release
provides GDB pretty-printers for device vectors/references, a new unique_count
algorithm, and an easier way to create tagged Thrust iterators. Several
documentation fixes are included, which can be found on the new Thrust
documentation site at https://nvidia.github.io/thrust. We'll be migrating
existing documentation sources to this new location over the next few months.
- NVIDIA#1586: Add new
thrust::make_tagged_iteratorconvenience function. Thanks to @karthikeyann for this contribution. - NVIDIA#1619: Add
unique_countalgorithm. Thanks to @upsj for this contribution. - NVIDIA#1631: Add GDB pretty-printers for device vectors/references
to
scripts/gdb-pretty-printers.py. Thanks to @upsj for this contribution.
- NVIDIA#1671: Fixed
reduce_by_keywhen called with 2^31 elements.
- NVIDIA#1512: Use CUB to implement
adjacent_difference. - NVIDIA#1555: Use CUB to implement
scan_by_key. - NVIDIA#1611: Add new doxybook-based Thrust documentation at https://nvidia.github.io/thrust.
- NVIDIA#1639: Fixed broken link in documentation. Thanks to @jrhemstad for this contribution.
- NVIDIA#1644: Increase contrast of search input text in new doc site. Thanks to @bdice for this contribution.
- NVIDIA#1647: Add
__forceinline__annotations to a functor wrapper. Thanks to @mkuron for this contribution. - NVIDIA#1660: Fixed typo in documentation example for
permutation_iterator. - NVIDIA#1669: Add a new
explicit_cuda_stream.cuexample that shows how to use explicit CUDA streams andpar/par_nosyncexecution policies.
Thrust 1.16.0 provides a new “nosync” hint for the CUDA backend, as well as numerous bugfixes and stability improvements.
Most of Thrust's parallel algorithms are fully synchronous and will block the calling CPU thread until all work is completed. This design avoids many pitfalls associated with asynchronous GPU programming, resulting in simpler and less-error prone usage for new CUDA developers. Unfortunately, this improvement in user experience comes at a performance cost that often frustrates more experienced CUDA programmers.
Prior to this release, the only synchronous-to-asynchronous migration path for
existing Thrust codebases involved significant refactoring, replacing calls
to thrust algorithms with a limited set of future-based thrust::async
algorithms or lower-level CUB kernels. The new thrust::cuda::par_nosync
execution policy provides a new, less-invasive entry point for asynchronous
computation.
par_nosync is a hint to the Thrust execution engine that any non-essential
internal synchronizations should be skipped and that an explicit synchronization
will be performed by the caller before accessing results.
While some Thrust algorithms require internal synchronization to safely compute
their results, many do not. For example, multiple thrust::for_each invocations
can be launched without waiting for earlier calls to complete:
// Queue three `for_each` kernels:
thrust::for_each(thrust::cuda::par_nosync, vec1.begin(), vec1.end(), Op{});
thrust::for_each(thrust::cuda::par_nosync, vec2.begin(), vec2.end(), Op{});
thrust::for_each(thrust::cuda::par_nosync, vec3.begin(), vec3.end(), Op{});
// Do other work while kernels execute:
do_something();
// Must explictly synchronize before accessing `for_each` results:
cudaDeviceSynchronize();Thanks to @fkallen for this contribution.
A future version of Thrust will remove support for CUDA Dynamic Parallelism (CDP).
This will only affect calls to Thrust algorithms made from CUDA device-side code that currently launches a kernel; such calls will instead execute sequentially on the calling GPU thread instead of launching a device-wide kernel.
- Thrust 1.14.0 included a change that aliased the
cubnamespace tothrust::cub. This has caused issues with ambiguous namespaces for projects that declareusing namespace thrust;from the global namespace. We recommend against this practice. - NVIDIA#1572: Removed several unnecessary header includes. Downstream projects may need to update their includes if they were relying on this behavior.
- NVIDIA#1568: Add
thrust::cuda::par_nosyncpolicy. Thanks to @fkallen for this contribution.
- NVIDIA#1511: Use CUB's new
DeviceMergeSortAPI and remove Thrust's internal implementation. - NVIDIA#1566: Improved performance of
thrust::shuffle. Thanks to @djns99 for this contribution. - NVIDIA#1584: Support user-defined
CMAKE_INSTALL_INCLUDEDIRvalues in Thrust's CMake install rules. Thanks to @robertmaynard for this contribution.
- NVIDIA#1496: Fix some issues affecting
iccbuilds. - NVIDIA#1552: Fix some collisions with the
min/maxmacros defined inwindows.h. - NVIDIA#1582: Fix issue with function type alias on 32-bit MSVC builds.
- NVIDIA#1591: Workaround issue affecting compilation with
nvc++. - NVIDIA#1597: Fix some collisions with the
smallmacro defined inwindows.h. - NVIDIA#1599, NVIDIA#1603: Fix some issues with version handling in Thrust's CMake packages.
- NVIDIA#1614: Clarify that scan algorithm results are non-deterministic for pseudo-associative operators (e.g. floating-point addition).
Thrust 1.15.0 provides numerous bugfixes, including non-numeric
thrust::sequence support, several MSVC-related compilation fixes, fewer
conversion warnings, counting_iterator initialization, and documentation
updates.
A future version of Thrust will remove support for CUDA Dynamic Parallelism (CDP).
This will only affect calls to Thrust algorithms made from CUDA device-side code that currently launches a kernel; such calls will instead execute sequentially on the calling GPU thread instead of launching a device-wide kernel.
- NVIDIA#1507: Allow
thrust::sequenceto work with non-numeric types. Thanks to Ben Jude (@bjude) for this contribution. - NVIDIA#1509: Avoid macro collision when calling
max()on MSVC. Thanks to Thomas (@tomintheshell) for this contribution. - NVIDIA#1514: Initialize all members in
counting_iterator's default constructor. - NVIDIA#1518: Fix
std::allocator_traitson MSVC + C++17. - NVIDIA#1530: Fix several
-Wconversionwarnings. Thanks to Matt Stack (@matt-stack) for this contribution. - NVIDIA#1539: Fixed typo in
thrust::for_eachdocumentation. Thanks to Salman (@untamedImpala) for this contribution. - NVIDIA#1548: Avoid name collision with
B0macro in termios.h system header. Thanks to Philip Deegan (@PhilipDeegan) for this contribution.
Thrust 1.14.0 is a major release accompanying the NVIDIA HPC SDK 21.9.
This release adds the ability to wrap the thrust:: namespace in an external
namespace, providing a workaround for a variety of shared library linking
issues. Thrust also learned to detect when CUB's symbols are in a wrapped
namespace and properly import them. To enable this feature, use
#define THRUST_CUB_WRAPPED_NAMESPACE foo to wrap both Thrust and CUB in the
foo:: namespace. See thrust/detail/config/namespace.h for details and more
namespace options.
Several bugfixes are also included: The tuple_size and tuple_element helpers
now support cv-qualified types. scan_by_key uses less memory.
thrust::iterator_traits is better integrated with std::iterator_traits.
See below for more details and references.
- Thrust 1.14.0 included a change that aliased the
cubnamespace tothrust::cub. This has caused issues with ambiguous namespaces for projects that declareusing namespace thrust;from the global namespace. We recommend against this practice.
- NVIDIA#1464: Add preprocessor hooks that allow
thrust::to be wrapped in an external namespace, and support cases when CUB is wrapped in an external namespace.
- NVIDIA#1457: Support cv-qualified types in
thrust::tuple_sizeandthrust::tuple_element. Thanks to Jake Hemstad for this contribution. - NVIDIA#1471: Fixed excessive memory allocation in
scan_by_key. Thanks to Lilo Huang for this contribution. - NVIDIA#1476: Removed dead code from the
expandexample. Thanks to Lilo Huang for this contribution. - NVIDIA#1488: Fixed the path to the installed CUB headers in the CMake
find_packageconfiguration files. - NVIDIA#1491: Fallback to
std::iterator_traitswhen nothrust::iterator_traitsspecialization exists for an iterator type. Thanks to Divye Gala for this contribution.
Thrust 1.13.1 is a minor release accompanying the CUDA Toolkit 11.5.
This release provides a new hook for embedding the thrust:: namespace inside a
custom namespace. This is intended to work around various issues related to
linking multiple shared libraries that use Thrust. The existing CUB_NS_PREFIX
and CUB_NS_POSTFIX macros already provided this capability for CUB; this
update provides a simpler mechanism that is extended to and integrated with
Thrust. Simply define THRUST_CUB_WRAPPED_NAMESPACE to a namespace name, and
both thrust:: and cub:: will be placed inside the new namespace. Using
different wrapped namespaces for each shared library will prevent issues like
those reported in NVIDIA#1401.
- NVIDIA#1464: Add
THRUST_CUB_WRAPPED_NAMESPACEhooks.
- NVIDIA#1488: Fix path to installed CUB in Thrust's CMake config files.
Thrust 1.13.0 is the major release accompanying the NVIDIA HPC SDK 21.7 release.
Notable changes include bfloat16 radix sort support (via thrust::sort) and
memory handling fixes in the reserve method of Thrust's vectors.
The CONTRIBUTING.md file has been expanded to include instructions for
building CUB as a component of Thrust, and API documentation now refers to
cppreference instead of SGI's old STL reference.
- NVIDIA#1459: Remove deprecated aliases
thrust::host_space_tagandthrust::device_space_tag. Use the equivalentthrust::host_system_tagandthrust::device_system_taginstead.
- NVIDIA/cub#306: Add radix-sort support for
bfloat16inthrust::sort. Thanks to Xiang Gao (@zasdfgbnm) for this contribution. - NVIDIA#1423:
thrust::transform_iteratornow supports non-copyable types. Thanks to Jake Hemstad (@jrhemstad) for this contribution. - NVIDIA#1459: Introduce a new
THRUST_IGNORE_DEPRECATED_APImacro that disables deprecation warnings on Thrust and CUB APIs.
- NVIDIA/cub#277: Fixed sanitizer warnings when
thrust::sortcalls intocub::DeviceRadixSort. Thanks to Andy Adinets (@canonizer) for this contribution. - NVIDIA#1442: Reduce extraneous comparisons in
thrust::sort's merge sort implementation. - NVIDIA#1447: Fix memory leak and avoid overallocation when
calling
reserveon Thrust's vector containers. Thanks to Kai Germaschewski (@germasch) for this contribution.
- NVIDIA#1405: Update links to standard C++ documentations from sgi to cppreference. Thanks to Muhammad Adeel Hussain (@AdeilH) for this contribution.
- NVIDIA#1432: Updated build instructions in
CONTRIBUTING.mdto include details on building CUB's test suite as part of Thrust.
Thrust 1.12.1 is a trivial patch release that slightly changes the phrasing of a deprecation message.
Thrust 1.12.0 is the major release accompanying the NVIDIA HPC SDK 21.3
and the CUDA Toolkit 11.4.
It includes a new thrust::universal_vector, which holds data that is
accessible from both host and device. This allows users to easily leverage
CUDA's unified memory with Thrust.
New asynchronous thrust::async:exclusive_scan and inclusive_scan algorithms
have been added, and the synchronous versions of these have been updated to
use cub::DeviceScan directly.
CUB radix sort for floating point types is now stable when both +0.0 and -0.0
are present in the input. This affects some usages of thrust::sort and
thrust::stable_sort.
Many compilation warnings and subtle overflow bugs were fixed in the device
algorithms, including a long-standing bug that returned invalid temporary
storage requirements when num_items was close to (but not
exceeding) INT32_MAX.
This release deprecates support for Clang < 7.0 and MSVC < 2019 (aka
19.20/16.0/14.20).
- NVIDIA#1372: Deprecate Clang < 7 and MSVC < 2019.
- NVIDIA#1376: Standardize
thrust::scan_by_keyfunctors / accumulator types. This may change the results fromscan_by_keywhen input, output, and initial value types are not the same type.
- NVIDIA#1251: Add two new
thrust::async::algorithms:inclusive_scanandexclusive_scan. - NVIDIA#1334: Add
thrust::universal_vector,universal_ptr, anduniversal_allocator.
- NVIDIA#1347: Qualify calls to
make_reverse_iterator. - NVIDIA#1359: Enable stricter warning flags. This fixes several
outstanding issues:
- NVIDIA/cub#221: Overflow in
temp_storage_byteswhennum_itemsclose to (but not over)INT32_MAX. - NVIDIA/cub#228: CUB uses non-standard C++ extensions that break strict compilers.
- NVIDIA/cub#257: Warning when compiling
GridEvenSharewith unsigned offsets. - NVIDIA#974: Conversion warnings in
thrust::transform_reduce. - NVIDIA#1091: Conversion warnings in
thrust::counting_iterator.
- NVIDIA/cub#221: Overflow in
- NVIDIA#1373: Fix compilation error when a standard library type is
wrapped in
thrust::optional. Thanks to Vukasin Milovanovic for this contribution. - NVIDIA#1388: Fix
signbit(double)implementation on MSVC. - NVIDIA#1389: Support building Thrust tests without CUDA enabled.
- NVIDIA#1304: Use
cub::DeviceScanto implementthrust::exclusive_scanandthrust::inclusive_scan. - NVIDIA#1362, NVIDIA#1370: Update smoke test naming.
- NVIDIA#1380: Fix typos in
set_operationdocumentation. Thanks to Hongyu Cai for this contribution. - NVIDIA#1383: Include FreeBSD license in LICENSE.md for
thrust::compleximplementation. - NVIDIA#1384: Add missing precondition to
thrust::gatherdocumentation.
Thrust 1.11.0 is a major release providing bugfixes and performance
enhancements.
It includes a new sort algorithm that provides up to 2x more performance
from thrust::sort when used with certain key types and hardware.
The new thrust::shuffle algorithm has been tweaked to improve the randomness
of the output.
Our CMake package and build system continue to see improvements with
better add_subdirectory support, installation rules, status messages, and
other features that make Thrust easier to use from CMake projects.
The release includes several other bugfixes and modernizations, and received
updates from 12 contributors.
- NVIDIA/cub#204: New implementation for
thrust::sorton CUDA when using 32/64-bit numeric keys on Pascal and up (SM60+). This improved radix sort algorithm provides up to 2x more performance. Thanks for Andy Adinets for this contribution. - NVIDIA#1310, NVIDIA#1312: Various tuple-related APIs have been updated to use variadic templates. Thanks for Andrew Corrigan for these contributions.
- NVIDIA#1297: Optionally add install rules when included with
CMake's
add_subdirectory. Thanks to Kai Germaschewski for this contribution.
- NVIDIA#1309: Fix
thrust::shuffleto produce better quality random distributions. Thanks to Rory Mitchell and Daniel Stokes for this contribution. - NVIDIA#1337: Fix compile-time regression in
transform_inclusive_scanandtransform_exclusive_scan. - NVIDIA#1306: Fix binary search
middlecalculation to avoid overflows. Thanks to Richard Barnes for this contribution. - NVIDIA#1314: Use
size_tfor the index type parameter inthrust::tuple_element. Thanks to Andrew Corrigan for this contribution. - NVIDIA#1329: Fix runtime error when copying an empty
thrust::device_vectorin MSVC Debug builds. Thanks to Ben Jude for this contribution. - NVIDIA#1323: Fix and add test for cmake package install rules. Thanks for Keith Kraus and Kai Germaschewski for testing and discussion.
- NVIDIA#1338: Fix GCC version checks in
thrust::detail::is_podimplementation. Thanks to Anatoliy Tomilov for this contribution. - NVIDIA#1289: Partial fixes for Clang 10 as host compiler. Filed an NVCC bug that will be fixed in a future version of the CUDA Toolkit (NVBug 3136307).
- NVIDIA#1272: Fix ambiguous
iter_swapcall when usingthrust::partitionwith STL containers. Thanks to Isaac Deutsch for this contribution. - NVIDIA#1281: Update our bundled
FindTBB.cmakemodule to support latest MSVC. - NVIDIA#1298: Use semantic versioning rules for our CMake package's compatibility checks. Thanks to Kai Germaschewski for this contribution.
- NVIDIA#1300: Use
FindPackageHandleStandardArgsto print standard status messages when our CMake package is found. Thanks to Kai Germaschewski for this contribution. - NVIDIA#1320: Use feature-testing instead of a language dialect check
for
thrust::remove_cvref. Thanks to Andrew Corrigan for this contribution. - NVIDIA#1319: Suppress GPU deprecation warnings.
- NVIDIA/cub#213: Removed some tuning policies for unsupported hardware (<SM35).
- References to the old Github repository and branch names were updated.
- Github's
thrust/cubrepository is nowNVIDIA/cub. - Development has moved from the
masterbranch to themainbranch.
- Github's
Thrust 1.10.0 is the major release accompanying the NVIDIA HPC SDK 20.9 release and the CUDA Toolkit 11.2 release. It drops support for C++03, GCC < 5, Clang < 6, and MSVC < 2017. It also overhauls CMake support. Finally, we now have a Code of Conduct for contributors: https://github.com/NVIDIA/thrust/blob/main/CODE_OF_CONDUCT.md
- C++03 is no longer supported.
- GCC < 5, Clang < 6, and MSVC < 2017 are no longer supported.
- C++11 is deprecated.
Using this dialect will generate a compile-time warning.
These warnings can be suppressed by defining
THRUST_IGNORE_DEPRECATED_CPP_DIALECTorTHRUST_IGNORE_DEPRECATED_CPP_11. Suppression is only a short term solution. We will be dropping support for C++11 in the near future. - Asynchronous algorithms now require C++14.
- CMake < 3.15 is no longer supported.
- The default branch on GitHub is now called
main. - Allocator and vector classes have been replaced with alias templates.
- NVIDIA#1159: CMake multi-config support, which allows multiple combinations of host and device systems to be built and tested at once. More details can be found here: https://github.com/NVIDIA/thrust/blob/main/CONTRIBUTING.md#multi-config-cmake-options
- CMake refactoring:
- Added install targets to CMake builds.
- Added support for CUB tests and examples.
- Thrust can be added to another CMake project by calling
add_subdirectorywith the Thrust source root (see NVIDIA#976). An example can be found here: https://github.com/NVIDIA/thrust/blob/main/examples/cmake/add_subdir/CMakeLists.txt - CMake < 3.15 is no longer supported.
- Dialects are now configured through target properties.
A new
THRUST_CPP_DIALECToption has been added for single config mode. Logic that modifiedCMAKE_CXX_STANDARDandCMAKE_CUDA_STANDARDhas been eliminated. - Testing related CMake code has been moved to
testing/CMakeLists.txt - Example related CMake code has been moved to
examples/CMakeLists.txt - Header testing related CMake code has been moved to
cmake/ThrustHeaderTesting.cmake - CUDA configuration CMake code has been moved to to
cmake/ThrustCUDAConfig.cmake. - Now we explicitly
include(cmake/*.cmake)files rather than searchingCMAKE_MODULE_PATH- we only want to use the ones in the repo.
thrust::transform_input_output_iterator, a variant of transform iterator adapter that works as both an input iterator and an output iterator. The given input function is applied after reading from the wrapped iterator while the output function is applied before writing to the wrapped iterator. Thanks to Trevor Smith for this contribution.
- Contributor documentation: https://github.com/NVIDIA/thrust/blob/main/CONTRIBUTING.md
- Code of Conduct: https://github.com/NVIDIA/thrust/blob/main/CODE_OF_CONDUCT.md. Thanks to Conor Hoekstra for this contribution.
- Support for all combinations of host and device systems.
- C++17 support.
- NVIDIA#1221: Allocator and vector classes have been replaced with alias templates. Thanks to Michael Francis for this contribution.
- NVIDIA#1186: Use placeholder expressions to simplify the definitions of a number of algorithms. Thanks to Michael Francis for this contribution.
- NVIDIA#1170: More conforming semantics for scan algorithms:
- Follow P0571's guidance regarding intermediate types.
- https://wg21.link/P0571
- The accumulator's type is now:
- The type of the user-supplied initial value (if provided), or
- The input iterator's value type if no initial value.
- Follow C++ standard guidance for default binary operator type.
- https://eel.is/c++draft/exclusive.scan#1
- Thrust binary/unary functors now specialize a default void template parameter. Types are deduced and forwarded transparently.
- Updated the scan's default binary operator to the new
thrust::plus<>specialization.
- The
thrust::intermediate_type_from_function_and_iteratorshelper is no longer needed and has been removed.
- Follow P0571's guidance regarding intermediate types.
- NVIDIA#1255: Always use
cudaStreamSynchronizeinstead ofcudaDeviceSynchronizeif the execution policy has a stream attached to it. Thanks to Rong Ou for this contribution. - NVIDIA#1201: Tests for correct handling of legacy and per-thread default streams. Thanks to Rong Ou for this contribution.
- NVIDIA#1260: Fix
thrust::transform_inclusive_scanwith heterogeneous types. Thanks to Rong Ou for this contribution. - NVIDIA#1258, NVC++ FS #28463: Ensure the CUDA radix sort backend synchronizes before returning; otherwise, copies from temporary storage will race with destruction of said temporary storage.
- NVIDIA#1264: Evaluate
CUDA_CUB_RET_IF_FAILmacro argument only once. Thanks to Jason Lowe for this contribution. - NVIDIA#1262: Add missing
<stdexcept>header. - NVIDIA#1250: Restore some
THRUST_DECLTYPE_RETURNSmacros in async test implementations. - NVIDIA#1249: Use
std::iotainCUDATestDriver::target_devices. Thanks to Michael Francis for this contribution. - NVIDIA#1244: Check for macro collisions with system headers during header testing.
- NVIDIA#1224: Remove unnecessary SFINAE contexts from asynchronous algorithms.
- NVIDIA#1190: Make
out_of_memory_recoverytest trigger faster. - NVIDIA#1187: Elminate superfluous iterators specific to the CUDA backend.
- NVIDIA#1181: Various fixes for GoUDA. Thanks to Andrei Tchouprakov for this contribution.
- NVIDIA#1178, NVIDIA#1229: Use transparent functionals in
placeholder expressions, fixing issues with
thrust::device_referenceand placeholder expressions andthrust::findwith asymmetric equality operators. - NVIDIA/thrust#1153: Switch to placement new instead of assignment to construct items in uninitialized memory. Thanks to Hugh Winkler for this contribution.
- NVIDIA#1050: Fix compilation of asynchronous algorithms when RDC is enabled.
- NVIDIA#1042: Correct return type of
thrust::detail::predicate_to_integralfrombooltoIntegralType. Thanks to Andreas Hehn for this contribution. - NVIDIA#1009: Avoid returning uninitialized allocators. Thanks to Zhihao Yuan for this contribution.
- NVIDIA#990: Add missing
<thrust/system/cuda/memory.h>include to<thrust/system/cuda/detail/malloc_and_free.h>. Thanks to Robert Maynard for this contribution. - NVIDIA#966: Fix spurious MSVC conversion with loss of data warning in sort algorithms. Thanks to Zhihao Yuan for this contribution.
- Add more metadata to mock specializations for testing iterator in
testing/copy.cu. - Add missing include to shuffle unit test.
- Specialize
thrust::wrapped_functionforvoidreturn types because MSVC is not a fan of the patternreturn static_cast<void>(expr);. - Replace deprecated
tbb/tbb_thread.hwith<thread>. - Fix overcounting of initial value in TBB scans.
- Use
thrust::advanceinstead of+=for generic iterators. - Wrap the OMP flags in
-Xcompilerfor NVCC - Extend
ASSERT_STATIC_ASSERTskip for the OMP backend. - Add missing header caught by
tbb.cudaconfigs. - Fix "unsafe API" warnings in examples on MSVC:
s/fopen/fstream/ - Various C++17 fixes.
Thrust 1.9.10-1 is the minor release accompanying the NVIDIA HPC SDK 20.7 release and the CUDA Toolkit 11.1 release.
- #1214, NVBug 200619442: Stop using
std::allocatorAPIs deprecated in C++17. - #1216, NVBug 200540293: Make
thrust::optionalwork with Clang when used with older libstdc++. - #1207, NVBug 200618218: Don't force C++14 with older compilers that don't support it.
- #1218: Wrap includes of
<memory>and<algorithm>to avoid circular inclusion with NVC++.
Thrust 1.9.10 is the release accompanying the NVIDIA HPC SDK 20.5 release.
It adds CMake support for compilation with NVC++ and a number of minor bug fixes
for NVC++.
It also adds CMake find_package support, which replaces the broken 3rd-party
legacy FindThrust.cmake script.
C++03, C++11, GCC < 5, Clang < 6, and MSVC < 2017 are now deprecated.
Starting with the upcoming 1.10.0 release, C++03 support will be dropped
entirely.
- #1082: Thrust now checks that it is compatible with the version of CUB found in your include path, generating an error if it is not. If you are using your own version of CUB, it may be too old. It is recommended to simply delete your own version of CUB and use the version of CUB that comes with Thrust.
- #1089: C++03 and C++11 are deprecated.
Using these dialects will generate a compile-time warning.
These warnings can be suppressed by defining
THRUST_IGNORE_DEPRECATED_CPP_DIALECT(to suppress C++03 and C++11 deprecation warnings) orTHRUST_IGNORE_DEPRECATED_CPP11(to suppress C++11 deprecation warnings). Suppression is only a short term solution. We will be dropping support for C++03 in the 1.10.0 release and C++11 in the near future. - #1089: GCC < 5, Clang < 6, and MSVC < 2017 are deprecated.
Using these compilers will generate a compile-time warning.
These warnings can be suppressed by defining
THRUST_IGNORE_DEPRECATED_COMPILER. Suppression is only a short term solution. We will be dropping support for these compilers in the near future.
- #1130: CMake
find_packagesupport. This is significant because there is a legacyFindThrust.cmakescript authored by a third party in widespread use in the community which has a bug in how it parses Thrust version numbers which will cause it to incorrectly parse 1.9.10. This script only handles the first digit of each part of the Thrust version number correctly: for example, Thrust 17.17.17 would be interpreted as Thrust 1.1.1701717. You can find directions for using the new CMakefind_packagesupport and migrating away from the legacyFindThrust.cmakehere - #1129: Added
thrust::detail::single_device_tls_caching_allocator, a convenient way to get an MR caching allocator for device memory, which is used by NVC++.
- #1129: Refactored RDC handling in CMake to be a global option and not create two targets for each example and test.
- #1129: Fix the legacy
thrust::return_temporary_bufferAPI to support passing a size. This was necessary to enable usage of Thrust caching MR allocators with synchronous Thrust algorithms. This change has allowed NVC++'s C++17 Parallel Algorithms implementation to switch to use Thrust caching MR allocators for device temporary storage, which gives a 2x speedup on large multi-GPU systems such as V100 and A100 DGX wherecudaMallocis very slow. - #1128: Respect
CUDA_API_PER_THREAD_DEFAULT_STREAM. Thanks to Rong Ou for this contribution. - #1131: Fix the one-policy overload of
thrust::async::copyto not copy the policy, resolving use-afer-move issues. - #1145: When cleaning up type names in
unittest::base_class_name, only callstd::string::replaceif we found the substring we are looking to replace. - #1139: Don't use
cxx::__demanglein NVC++. - #1102: Don't use
thrust::detail::normal_distribution_nvccfor Feta because it useserfcinv, a non-standard function that Feta doesn't have.
Thrust 1.9.9 adds support for NVC++, which uses Thrust to implement
GPU-accelerated C++17 Parallel Algorithms.
thrust::zip_function and thrust::shuffle were also added.
C++03, C++11, GCC < 5, Clang < 6, and MSVC < 2017 are now deprecated.
Starting with the upcoming 1.10.0 release, C++03 support will be dropped
entirely.
All other deprecated platforms will be dropped in the near future.
- #1082: Thrust now checks that it is compatible with the version of CUB found in your include path, generating an error if it is not. If you are using your own version of CUB, it may be too old. It is recommended to simply delete your own version of CUB and use the version of CUB that comes with Thrust.
- #1089: C++03 and C++11 are deprecated.
Using these dialects will generate a compile-time warning.
These warnings can be suppressed by defining
THRUST_IGNORE_DEPRECATED_CPP_DIALECT(to suppress C++03 and C++11 deprecation warnings) orTHRUST_IGNORE_DEPRECATED_CPP_11(to suppress C++11 deprecation warnings). Suppression is only a short term solution. We will be dropping support for C++03 in the 1.10.0 release and C++11 in the near future. - #1089: GCC < 5, Clang < 6, and MSVC < 2017 are deprecated.
Using these compilers will generate a compile-time warning.
These warnings can be suppressed by defining
THRUST_IGNORE_DEPRECATED_COMPILER. Suppression is only a short term solution. We will be dropping support for these compilers in the near future.
- #1086: Support for NVC++ aka "Feta".
The most significant change is in how we use
__CUDA_ARCH__. Now, there are four macros that must be used:THRUST_IS_DEVICE_CODE, which should be used in anifstatement around device-only code.THRUST_INCLUDE_DEVICE_CODE, which should be used in an#ifpreprocessor directive inside of theifstatement mentioned in the prior bullet.THRUST_IS_HOST_CODE, which should be used in anifstatement around host-only code.THRUST_INCLUDE_HOST_CODE, which should be used in an#ifpreprocessor directive inside of theifstatement mentioned in the prior bullet.
- #1085:
thrust::shuffle. Thanks to Rory Mitchell for this contribution. - #1029:
thrust::zip_function, a facility for zipping functions that take N parameters instead of a tuple of N parameters asthrust::zip_iteratordoes. Thanks to Ben Jude for this contribution. - #1068:
thrust::system::cuda::managed_memory_pointer, a universal memory strongly typed pointer compatible with the ISO C++ Standard Library.
- #1029: Thrust is now built and tested with NVCC warnings treated as errors.
- #1029: MSVC C++11 support.
- #1029:
THRUST_DEPRECATEDabstraction for generating compile-time deprecation warning messages. - #1029:
thrust::pointer<T>::pointer_to(reference). - #1070: Unit test for
thrust::inclusive_scanwith a user defined types. Thanks to Conor Hoekstra for this contribution.
- #1088: Allow
thrust::replaceto take functions that have non-constoperator(). - #1094: Add missing
constexprtopar_tconstructors. Thanks to Patrick Stotko for this contribution. - #1077: Remove
__device__from CUDA MR-based device allocators to fix obscure "host function called from host device function" warning that occurs when you use the new Thrust MR-based allocators. - #1029: Remove inconsistently-used
THRUST_BEGIN/END_NSmacros. - #1029: Fix C++ dialect detection on newer MSVC.
- #1029 Use
_Pragma/__pragmainstead of#pragmain macros. - #1029: Replace raw
__cpluspluschecks with the appropriate Thrust macros. - #1105: Add a missing
<math.h>include. - #1103: Fix regression of
thrust::detail::temporary_allocatorwith non-CUDA back ends. - #1111: Use Thrust's random number engine instead of
std::s in device code. - #1108: Get rid of a GCC 9 warning about deprecated generation of copy ctors.
Thrust 1.9.8-1 is a variant of 1.9.8 accompanying the NVIDIA HPC SDK 20.3 release. It contains modifications necessary to serve as the implementation of NVC++'s GPU-accelerated C++17 Parallel Algorithms when using the CUDA Toolkit 11.0 release.
Thrust 1.9.8, which is included in the CUDA Toolkit 11.0 release, removes
Thrust's internal derivative of CUB, upstreams all relevant changes too CUB,
and adds CUB as a Git submodule.
It will now be necessary to do git clone --recursive when checking out
Thrust, and to update the CUB submodule when pulling in new Thrust changes.
Additionally, CUB is now included as a first class citizen in the CUDA toolkit.
Thrust 1.9.8 also fixes bugs preventing most Thrust algorithms from working
with more than 2^31-1 elements.
Now, thrust::reduce, thrust::*_scan, and related algorithms (aka most of
Thrust) work with large element counts.
- Thrust will now use the version of CUB in your include path instead of its own internal copy. If you are using your own version of CUB, it may be older and incompatible with Thrust. It is recommended to simply delete your own version of CUB and use the version of CUB that comes with Thrust.
- Refactor Thrust and CUB to support 64-bit indices in most algorithms.
In most cases, Thrust now selects between kernels that use 32-bit indices and
64-bit indices at runtime depending on the size of the input.
This means large element counts work, but small element counts do not have to
pay for the register usage of 64-bit indices if they are not needed.
Now,
thrust::reduce,thrust::*_scan, and related algorithms (aka most of Thrust) work with more than2^31-1elements. Notably,thrust::sortis still limited to less than2^31-1elements. - CUB is now a submodule and the internal copy of CUB has been removed.
- #1051: Stop specifying the
__launch_bounds__minimum blocks parameter because it messes up register allocation and increases register pressure, and we don't actually know at compile time how many blocks we will use (aside from single tile kernels).
- #1020: After making a CUDA API call, always clear the global CUDA error state
by calling
cudaGetLastError. - #1021: Avoid calling destroy in the destructor of a Thrust vector if the vector is empty.
- #1046: Actually throw
thrust::bad_allocwhenthrust::system::cuda::mallocfails instead of just constructing a temporary and doing nothing with it. - Add missing copy constructor or copy assignment operator to all classes that
GCC 9's
-Wdeprecated-copycomplains about - Add missing move operations to
thrust::system::cuda::vector. - #1015: Check that the backend is CUDA before using CUDA-specifics in
thrust::detail::temporary_allocator. Thanks to Hugh Winkler for this contribution. - #1055: More correctly detect the presence of aligned/sized
new/delete. - #1043: Fix ill-formed specialization of
thrust::system::is_error_code_enumforthrust::event_errc. Thanks to Toru Niina for this contribution. - #1027: Add tests for
thrust::tuple_for_eachandthrust::tuple_subset. Thanks to Ben Jude for this contribution. - #1027: Use correct macro in
thrust::tuple_for_each. Thanks to Ben Jude for this contribution. - #1026: Use correct MSVC version formatting in CMake. Thanks to Ben Jude for this contribution.
- Workaround an NVCC issue with type aliases with template template arguments containing a parameter pack.
- Remove unused functions from the CUDA backend which call slow CUDA attribute query APIs.
- Replace
CUB_RUNTIME_FUNCTIONwithTHRUST_RUNTIME_FUNCTION. - Correct typo in
thrust::transformdocumentation. Thanks to Eden Yefet for this contribution.
thrust::sortremains limited to2^31-1elements for now.
Thrust 1.9.7-1 is a minor release accompanying the CUDA Toolkit 10.2 release for Tegra. It is nearly identical to 1.9.7.
- Remove support for GCC's broken nodiscard-like attribute.
Thrust 1.9.7 is a minor release accompanying the CUDA Toolkit 10.2 release.
Unfortunately, although the version and patch numbers are identical, one bug
fix present in Thrust 1.9.7 (NVBug 2646034: Fix incorrect dependency handling
for stream acquisition in thrust::future) was not included in the CUDA
Toolkit 10.2 preview release for AArch64 SBSA.
The tag cuda-10.2aarch64sbsa contains the exact version of Thrust present
in the CUDA Toolkit 10.2 preview release for AArch64 SBSA.
- #967, NVBug 2448170: Fix the CUDA backend
thrust::for_eachso that it supports large input sizes with 64-bit indices. - NVBug 2646034: Fix incorrect dependency handling for stream acquisition in
thrust::future.- Not present in the CUDA Toolkit 10.2 preview release for AArch64 SBSA.
- #968, NVBug 2612102: Fix the
thrust::mr::polymorphic_adaptorto actually use its template parameter.
Thrust 1.9.6-1 is a variant of 1.9.6 accompanying the NVIDIA HPC SDK 20.3 release. It contains modifications necessary to serve as the implementation of NVC++'s GPU-accelerated C++17 Parallel Algorithms when using the CUDA Toolkit 10.1 Update 2 release.
Thrust 1.9.6 is a minor release accompanying the CUDA Toolkit 10.1 Update 2 release.
- NVBug 2509847: Inconsistent alignment of
thrust::complex - NVBug 2586774: Compilation failure with Clang + older libstdc++ that doesn't
have
std::is_trivially_copyable - NVBug 200488234: CUDA header files contain Unicode characters which leads compiling errors on Windows
- #949, #973, NVBug 2422333, NVBug 2522259, NVBug 2528822:
thrust::detail::aligned_reinterpret_castmust be annotated with__host__ __device__. - NVBug 2599629: Missing include in the OpenMP sort implementation
- NVBug 200513211: Truncation warning in test code under VC142
Thrust 1.9.5 is a minor release accompanying the CUDA Toolkit 10.1 Update 1 release.
- NVBug 2502854: Fixed assignment of
thrust::device_vector<thrust::complex<T>>between host and device.
Thrust 1.9.4 adds asynchronous interfaces for parallel algorithms, a new
allocator system including caching allocators and unified memory support, as
well as a variety of other enhancements, mostly related to
C++11/C++14/C++17/C++20 support.
The new asynchronous algorithms in the thrust::async namespace return
thrust::event or thrust::future objects, which can be waited upon to
synchronize with the completion of the parallel operation.
Synchronous Thrust algorithms now block until all of their operations have completed. Use the new asynchronous Thrust algorithms for non-blocking behavior.
thrust::eventandthrust::future<T>, uniquely-owned asynchronous handles consisting of a state (ready or not ready), content (some value; forthrust::futureonly), and an optional set of objects that should be destroyed only when the future's value is ready and has been consumed.- The design is loosely based on C++11's
std::future. - They can be
.wait'd on, and the value of a future can be waited on and retrieved with.getor.extract. - Multiple
thrust::events andthrust::futures can be combined withthrust::when_all. thrust::futures can be converted tothrust::events.- Currently, these primitives are only implemented for the CUDA backend and are C++11 only.
- The design is loosely based on C++11's
- New asynchronous algorithms that return
thrust::event/thrust::futures, implemented as C++20 range style customization points:thrust::async::reduce.thrust::async::reduce_into, which takes a target location to store the reduction result into.thrust::async::copy, including a two-policy overload that allows explicit cross system copies which execution policy properties can be attached to.thrust::async::transform.thrust::async::for_each.thrust::async::stable_sort.thrust::async::sort.- By default the asynchronous algorithms use the new caching allocators.
Deallocation of temporary storage is deferred until the destruction of
the returned
thrust::future. The content ofthrust::futures is stored in either device or universal memory and transferred to the host only upon request to prevent unnecessary data migration. - Asynchronous algorithms are currently only implemented for the CUDA system and are C++11 only.
exec.after(f, g, ...), a new execution policy method that takes a set ofthrust::event/thrust::futures and returns an execution policy that operations on that execution policy should depend upon.- New logic and mindset for the type requirements for cross-system sequence
copies (currently only used by
thrust::async::copy), based on:thrust::is_contiguous_iteratorandTHRUST_PROCLAIM_CONTIGUOUS_ITERATORfor detecting/indicating that an iterator points to contiguous storage.thrust::is_trivially_relocatableandTHRUST_PROCLAIM_TRIVIALLY_RELOCATABLEfor detecting/indicating that a type ismemcpyable (based on principles from P1144).- The new approach reduces buffering, increases performance, and increases correctness.
- The fast path is now enabled when copying CUDA
__halfand vector types withthrust::async::copy.
- All Thrust synchronous algorithms for the CUDA backend now actually
synchronize. Previously, any algorithm that did not allocate temporary
storage (counterexample:
thrust::sort) and did not have a computation-dependent result (counterexample:thrust::reduce) would actually be launched asynchronously. Additionally, synchronous algorithms that allocated temporary storage would become asynchronous if a custom allocator was supplied that did not synchronize on allocation/deallocation, unlikecudaMalloc/cudaFree. So, nowthrust::for_each,thrust::transform,thrust::sort, etc are truly synchronous. In some cases this may be a performance regression; if you need asynchrony, use the new asynchronous algorithms. - Thrust's allocator framework has been rewritten. It now uses a memory
resource system, similar to C++17's
std::pmrbut supporting static polymorphism. Memory resources are objects that allocate untyped storage and allocators are cheap handles to memory resources in this new model. The new facilities live in<thrust/mr/*>.thrust::mr::memory_resource<Pointer>, the memory resource base class, which takes a (possibly tagged) pointer tovoidtype as a parameter.thrust::mr::allocator<T, MemoryResource>, an allocator backed by a memory resource object.thrust::mr::polymorphic_adaptor_resource<Pointer>, a type-erased memory resource adaptor.thrust::mr::polymorphic_allocator<T>, a C++17-style polymorphic allocator backed by a type-erased memory resource object.- New tunable C++17-style caching memory resources,
thrust::mr::(disjoint_)?(un)?synchronized_pool_resource, designed to cache both small object allocations and large repetitive temporary allocations. The disjoint variants use separate storage for management of the pool, which is necessary if the memory being allocated cannot be accessed on the host (e.g. device memory). - System-specific allocators were rewritten to use the new memory resource framework.
- New
thrust::device_memory_resourcefor allocating device memory. - New
thrust::universal_memory_resourcefor allocating memory that can be accessed from both the host and device (e.g.cudaMallocManaged). - New
thrust::universal_host_pinned_memory_resourcefor allocating memory that can be accessed from the host and the device but always resides in host memory (e.g.cudaMallocHost). thrust::get_per_device_resourceandthrust::per_device_allocator, which lazily create and retrieve a per-device singleton memory resource.- Rebinding mechanisms (
rebind_traitsandrebind_alloc) forthrust::allocator_traits. thrust::device_make_unique, a factory function for creating astd::unique_ptrto a newly allocated object in device memory.<thrust/detail/memory_algorithms>, a C++11 implementation of the C++17 uninitialized memory algorithms.thrust::allocate_uniqueand friends, based on the proposed C++23std::allocate_unique.
- New type traits and metaprogramming facilities. Type traits are slowly being
migrated out of
thrust::detail::and<thrust/detail/*>; their new home will bethrust::and<thrust/type_traits/*>.thrust::is_execution_policy.thrust::is_operator_less_or_greater_function_object, which detectsthrust::less,thrust::greater,std::less, andstd::greater.thrust::is_operator_plus_function_object``, which detectsthrust::plusandstd::plus`.thrust::remove_cvref(_t)?, a C++11 implementation of C++20'sthrust::remove_cvref(_t)?.thrust::void_t, and various other new type traits.thrust::integer_sequenceand friends, a C++11 implementation of C++20'sstd::integer_sequencethrust::conjunction,thrust::disjunction, andthrust::disjunction, a C++11 implementation of C++17's logical metafunctions.- Some Thrust type traits (such as
thrust::is_constructible) have been redefined in terms of C++11's type traits when they are available.
<thrust/detail/tuple_algorithms.h>, newstd::tuplealgorithms:thrust::tuple_transform.thrust::tuple_for_each.thrust::tuple_subset.
- Miscellaneous new
std::-like facilities:thrust::optional, a C++11 implementation of C++17'sstd::optional.thrust::addressof, an implementation of C++11'sstd::addressof.thrust::nextandthrust::prev, an implementation of C++11'sstd::nextandstd::prev.thrust::square, a<functional>style unary function object that multiplies its argument by itself.<thrust/limits.h>andthrust::numeric_limits, a customized version of<limits>andstd::numeric_limits.
<thrust/detail/preprocessor.h>, new general purpose preprocessor facilities:THRUST_PP_CAT[2-5], concatenates two to five tokens.THRUST_PP_EXPAND(_ARGS)?, performs double expansion.THRUST_PP_ARITYandTHRUST_PP_DISPATCH, tools for macro overloading.THRUST_PP_BOOL, boolean conversion.THRUST_PP_INCandTHRUST_PP_DEC, increment/decrement.THRUST_PP_HEAD, a variadic macro that expands to the first argument.THRUST_PP_TAIL, a variadic macro that expands to all its arguments after the first.THRUST_PP_IIF, bitwise conditional.THRUST_PP_COMMA_IF, andTHRUST_PP_HAS_COMMA, facilities for adding and detecting comma tokens.THRUST_PP_IS_VARIADIC_NULLARY, returns true if called with a nullary__VA_ARGS__.THRUST_CURRENT_FUNCTION, expands to the name of the current function.
- New C++11 compatibility macros:
THRUST_NODISCARD, expands to[[nodiscard]]when available and the best equivalent otherwise.THRUST_CONSTEXPR, expands toconstexprwhen available and the best equivalent otherwise.THRUST_OVERRIDE, expands tooverridewhen available and the best equivalent otherwise.THRUST_DEFAULT, expands to= default;when available and the best equivalent otherwise.THRUST_NOEXCEPT, expands tonoexceptwhen available and the best equivalent otherwise.THRUST_FINAL, expands tofinalwhen available and the best equivalent otherwise.THRUST_INLINE_CONSTANT, expands toinline constexprwhen available and the best equivalent otherwise.
<thrust/detail/type_deduction.h>, new C++11-only type deduction helpers:THRUST_DECLTYPE_RETURNS*, expand to function definitions with suitable conditionalnoexceptqualifiers and trailing return types.THRUST_FWD(x), expands to::std::forward<decltype(x)>(x).THRUST_MVCAP, expands to a lambda move capture.THRUST_RETOF, expands to a decltype computing the return type of an invocable.
- New CMake build system.
mr_basicdemonstrates how to use the new memory resource allocator system.
- Tagged pointer enhancements:
- New
thrust::pointer_traitsspecialization forvoid const*. nullptrsupport to Thrust tagged pointers.- New
explicit operator boolfor Thrust tagged pointers when using C++11 forstd::unique_ptrinteroperability. - Added
thrust::reinterpret_pointer_castandthrust::static_pointer_castfor casting Thrust tagged pointers.
- New
- Iterator enhancements:
thrust::iterator_systemis now SFINAE friendly.- Removed cv qualifiers from iterator types when using
thrust::iterator_system.
- Static assert enhancements:
- New
THRUST_STATIC_ASSERT_MSG, takes an optional string constant to be used as the error message when possible. - Update
THRUST_STATIC_ASSERT(_MSG)to use C++11'sstatic_assertwhen it's available. - Introduce a way to test for static assertions.
- New
- Testing enhancements:
- Additional scalar and sequence types, including non-builtin types and vectors with unified memory allocators, have been added to the list of types used by generic unit tests.
- The generation of random input data has been improved to increase the range of values used and catch more corner cases.
- New
unittest::truncate_to_max_representableutility for avoiding the generation of ranges that cannot be represented by the underlying element type in generic unit test code. - The test driver now synchronizes with CUDA devices and check for errors after each test, when switching devices, and after each raw kernel launch.
- The
warningtesteruber header is now compiled with NVCC to avoid needing to disable CUDA-specific code with the preprocessor. - Fixed the unit test framework's
ASSERT_*to printchars asints. - New
DECLARE_INTEGRAL_VARIABLE_UNITTESTtest declaration macro. - New
DECLARE_VARIABLE_UNITTEST_WITH_TYPES_AND_NAMEtest declaration macro. thrust::system_errorin the CUDA backend now print out itscudaError_tenumerator in addition to the diagnostic message.- Stopped using conditionally signed types like
char.
- #897, NVBug 2062242: Fix compilation error when using
__device__lambdas withthrust::reduceon MSVC. - #908, NVBug 2089386: Static assert that
thrust::generate/thrust::fillisn't operating on const iterators. - #919 Fix compilation failure with
thrust::zip_iteratorandthrust::complex. - #924, NVBug 2096679, NVBug 2315990: Fix dispatch for the CUDA backend's
thrust::reduceto use two functions (one with the pragma for disabling exec checks, one withTHRUST_RUNTIME_FUNCTION) instead of one. This fixes a regression with device compilation that started in CUDA Toolkit 9.2. - #928, NVBug 2341455: Add missing
__host__ __device__annotations to athrust::complex::operator=to satisfy GoUDA. - NVBug 2094642: Make
thrust::vector_base::clearnot depend on the element type being default constructible. - NVBug 2289115: Remove flaky
simple_cuda_streamsexample. - NVBug 2328572: Add missing
thrust::device_vectorconstructor that takes an allocator parameter. - NVBug 2455740: Update the
range_viewexample to not use device-side launch. - NVBug 2455943: Ensure that sized unit tests that use
thrust::counting_iteratorperform proper truncation. - NVBug 2455952: Refactor questionable
thrust::copy_ifunit tests.
Thrust 1.9.3 unifies and integrates CUDA Thrust and GitHub Thrust.
- #725, #850, #855, #859, #860: Unify the
thrust::iter_swapinterface and fixthrust::device_referenceswapping. - NVBug 2004663: Add a
datamethod tothrust::detail::temporary_arrayand refactor temporary memory allocation in the CUDA backend to be exception and leak safe. - #886, #894, #914: Various documentation typo fixes.
- #724: Provide
NVVMIR_LIBRARY_DIRenvironment variable to NVCC. - #878: Optimize
thrust::min/max_elementto only usethrust::detail::get_iterator_valuefor non-numeric types. - #899: Make
thrust::cuda::experimental::pinned_allocator's comparison operatorsconst. - NVBug 2092152: Remove all includes of
<cuda.h>. - #911: Fix default comparator element type for
thrust::merge_by_key.
- Thanks to Andrew Corrigan for contributing fixes for swapping interfaces.
- Thanks to Francisco Facioni for contributing optimizations for
thrust::min/max_element.
Thrust 1.9.2 brings a variety of performance enhancements, bug fixes and test
improvements.
CUB 1.7.5 was integrated, enhancing the performance of thrust::sort on
small data types and thrust::reduce.
Changes were applied to complex to optimize memory access.
Thrust now compiles with compiler warnings enabled and treated as errors.
Additionally, the unit test suite and framework was enhanced to increase
coverage.
- The
fallback_allocatorexample was removed, as it was buggy and difficult to support.
<thrust/detail/alignment.h>, utilities for memory alignment:thrust::aligned_reinterpret_cast.thrust::aligned_storage_size, which computes the amount of storage needed for an object of a particular size and alignment.thrust::alignment_of, a C++03 implementation of C++11'sstd::alignment_of.thrust::aligned_storage, a C++03 implementation of C++11'sstd::aligned_storage.thrust::max_align_t, a C++03 implementation of C++11'sstd::max_align_t.
- NVBug 200385527, NVBug 200385119, NVBug 200385113, NVBug 200349350, NVBug 2058778: Various compiler warning issues.
- NVBug 200355591:
thrust::reduceperformance issues. - NVBug 2053727: Fixed an ADL bug that caused user-supplied
allocateto be overlooked butdeallocateto be called with GCC <= 4.3. - NVBug 1777043: Fixed
thrust::complexto work withthrust::sequence.
Thrust 1.9.1-2 integrates version 1.7.4 of CUB and introduces a new CUDA backend
for thrust::reduce based on CUB.
- NVBug 1965743: Remove unnecessary static qualifiers.
- NVBug 1940974: Fix regression causing a compilation error when using
thrust::merge_by_keywiththrust::constant_iterators. - NVBug 1904217: Allow callables that take non-const refs to be used with
thrust::reduceandthrust::*_scan.
Thrust 1.9.0-5 replaces the original CUDA backend (bulk) with a new one written using CUB, a high performance CUDA collectives library. This brings a substantial performance improvement to the CUDA backend across the board.
- Any code depending on CUDA backend implementation details will likely be broken.
- New CUDA backend based on CUB which delivers substantially higher performance.
thrust::transform_output_iterator, a fancy iterator that applies a function to the output before storing the result.
transform_output_iteratordemonstrates use of the new fancy iteratorthrust::transform_output_iterator.
- When C++11 is enabled, functors do not have to inherit from
thrust::(unary|binary)_functionanymore to be used withthrust::transform_iterator. - Added C++11 only move constructors and move assignment operators for
thrust::detail::vector_base-based classes, e.g.thrust::host_vector,thrust::device_vector, and friends.
sin(thrust::complex<double>)no longer has precision loss to float.
- Thanks to Manuel Schiller for contributing a C++11 based enhancement
regarding the deduction of functor return types, improving the performance
of
thrust::uniqueand implementingthrust::transform_output_iterator. - Thanks to Thibault Notargiacomo for the implementation of move semantics for
the
thrust::vector_base-based classes. - Thanks to Duane Merrill for developing CUB and helping to integrate it into Thrust's backend.
Thrust 1.8.3 is a small bug fix release.
range_viewdemonstrates the use of a view (a non-owning wrapper for an iterator range with a container-like interface).
thrust::(min|max|minmax)_elementcan now accept raw device pointers when an explicit device execution policy is used.thrust::clearoperations on vector types no longer requires the element type to have a default constructor.
Thrust 1.8.2 is a small bug fix release.
- Avoid warnings and errors concerning user functions called from
__host__ __device__functions. - #632: Fix an error in
thrust::set_intersection_by_keywith the CUDA backend. - #651:
thrust::copybetween host and device now accepts execution policies with streams attached, i.e.thrust::::cuda::par.on(stream). - #664:
thrust::for_eachand algorithms based on it no longer ignore streams attached to execution policys.
- #628:
thrust::reduce_by_keyfor the CUDA backend fails for Compute Capability 5.0 devices.
Thrust 1.8.1 is a small bug fix release.
- #615, #620: Fixed
thrust::for_eachandthrust::reduceto no longer fail on large inputs.
- #628:
thrust::reduce_by_keyfor the CUDA backend fails for Compute Capability 5.0 devices.
Thrust 1.8.0 introduces support for algorithm invocation from CUDA device
code, support for CUDA streams, and algorithm performance improvements.
Users may now invoke Thrust algorithms from CUDA device code, providing a
parallel algorithms library to CUDA programmers authoring custom kernels, as
well as allowing Thrust programmers to nest their algorithm calls within
functors.
The thrust::seq execution policy allows users to require sequential algorithm
execution in the calling thread and makes a sequential algorithms library
available to individual CUDA threads.
The .on(stream) syntax allows users to request a CUDA stream for kernels
launched during algorithm execution.
Finally, new CUDA algorithm implementations provide substantial performance
improvements.
- Algorithms in CUDA Device Code:
- Thrust algorithms may now be invoked from CUDA
__device__and__host__device functions. Algorithms invoked in this manner must be invoked with an execution policy as the first parameter. The following execution policies are supported in CUDA device code:thrust::seqthrust::cuda::parthrust::device, when THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA.
- Device-side algorithm execution may not be parallelized unless CUDA Dynamic Parallelism is available.
- Thrust algorithms may now be invoked from CUDA
- Execution Policies:
- CUDA Streams
- The
thrust::cuda::par.on(stream)syntax allows users to request that CUDA kernels launched during algorithm execution should occur on a given stream. - Algorithms executed with a CUDA stream in this manner may still synchronize with other streams when allocating temporary storage or returning results to the CPU.
- The
thrust::seq, which allows users to require that an algorithm execute sequentially in the calling thread.
- CUDA Streams
thrust::complex, a complex number data type.
- simple_cuda_streams demonstrates how to request a CUDA stream during algorithm execution.
- async_reduce demonstrates ways to achieve algorithm invocations which are asynchronous with the calling thread.
- CUDA sort performance for user-defined types is 300% faster on Tesla K20c for large problem sizes.
- CUDA merge performance is 200% faster on Tesla K20c for large problem sizes.
- CUDA sort performance for primitive types is 50% faster on Tesla K20c for large problem sizes.
- CUDA reduce_by_key performance is 25% faster on Tesla K20c for large problem sizes.
- CUDA scan performance is 15% faster on Tesla K20c for large problem sizes.
- fallback_allocator example is simpler.
- #364: Iterators with unrelated system tags may be used with algorithms invoked with an execution policy
- #371: Do not redefine
__CUDA_ARCH__. - #379: Fix crash when dereferencing transform_iterator on the host.
- #391: Avoid use of uppercase variable names.
- #392: Fix
thrust::copybetweencusp::complexandstd::complex. - #396: Program compiled with gcc < 4.3 hangs during comparison sort.
- #406:
fallback_allocator.cuexample checks device for unified addressing support. - #417: Avoid using
std::less<T>in binary search algorithms. - #418: Avoid various warnings.
- #443: Including version.h no longer configures default systems.
- #578: NVCC produces warnings when sequential algorithms are used with CPU systems.
- When invoked with primitive data types, thrust::sort, thrust::sort_by_key, thrust::stable_sort, & thrust::stable_sort_by_key may
- Sometimes linking fails when compiling with
-rdc=truewith NVCC. - The CUDA implementation of thrust::reduce_by_key incorrectly outputs the last element in a segment of equivalent keys instead of the first.
- Thanks to Sean Baxter for contributing faster CUDA reduce, merge, and scan implementations.
- Thanks to Duane Merrill for contributing a faster CUDA radix sort implementation.
- Thanks to Filipe Maia for contributing the implementation of thrust::complex.
Thrust 1.7.2 is a minor bug fix release.
- Avoid use of
std::minin generic find implementation.
Thrust 1.7.1 is a minor bug fix release.
- Eliminate identifiers in
set_operations.cuexample with leading underscore. - Eliminate unused variable warning in CUDA
reduce_by_keyimplementation. - Avoid deriving function objects from
std::unary_functionandstd::binary_function.
Thrust 1.7.0 introduces a new interface for controlling algorithm execution as well as several new algorithms and performance improvements. With this new interface, users may directly control how algorithms execute as well as details such as the allocation of temporary storage. Key/value versions of thrust::merge and the set operation algorithms have been added, as well stencil versions of partitioning algorithms. thrust::tabulate has been introduced to tabulate the values of functions taking integers. For 32b types, new CUDA merge and set operations provide 2-15x faster performance while a new CUDA comparison sort provides 1.3-4x faster performance. Finally, a new TBB reduce_by_key implementation provides 80% faster performance.
- Dispatch:
- Custom user backend systems' tag types must now inherit from the corresponding system's execution_policy template (e.g. thrust::cuda::execution_policy) instead of the tag struct (e.g. thrust::cuda::tag). Otherwise, algorithm specializations will silently go unfound during dispatch. See examples/minimal_custom_backend.cu and examples/cuda/fallback_allocator.cu for usage examples.
- thrust::advance and thrust::distance are no longer dispatched based on iterator system type and thus may no longer be customized.
- Iterators:
- iterator_facade and iterator_adaptor's Pointer template parameters have been eliminated.
- iterator_adaptor has been moved into the thrust namespace (previously thrust::experimental::iterator_adaptor).
- iterator_facade has been moved into the thrust namespace (previously thrust::experimental::iterator_facade).
- iterator_core_access has been moved into the thrust namespace (previously thrust::experimental::iterator_core_access).
- All iterators' nested pointer typedef (the type of the result of operator->) is now void instead of a pointer type to indicate that such expressions are currently impossible.
- Floating point counting_iterators' nested difference_type typedef is now a signed integral type instead of a floating point type.
- Other:
- normal_distribution has been moved into the thrust::random namespace (previously thrust::random::experimental::normal_distribution).
- Placeholder expressions may no longer include the comma operator.
- Execution Policies:
- Users may directly control the dispatch of algorithm invocations with optional execution policy arguments. For example, instead of wrapping raw pointers allocated by cudaMalloc with thrust::device_ptr, the thrust::device execution_policy may be passed as an argument to an algorithm invocation to enable CUDA execution.
- The following execution policies are supported in this version:
thrust::hostthrust::devicethrust::cpp::parthrust::cuda::parthrust::omp::parthrust::tbb::par
- Algorithms:
thrust::merge_by_keythrust::partitionwith stencilthrust::partition_copywith stencilthrust::set_difference_by_keythrust::set_intersection_by_keythrust::set_symmetric_difference_by_keythrust::set_union_by_keythrust::stable_partition with stencilthrust::stable_partition_copy with stencilthrust::tabulate
- Memory Allocation:
thrust::mallocthrust::freethrust::get_temporary_bufferthrust::return_temporary_buffer
- uninitialized_vector demonstrates how to use a custom allocator to avoid the automatic initialization of elements in thrust::device_vector.
- Authors of custom backend systems may manipulate arbitrary state during algorithm dispatch by incorporating it into their execution_policy parameter.
- Users may control the allocation of temporary storage during algorithm execution by passing standard allocators as parameters via execution policies such as thrust::device.
- THRUST_DEVICE_SYSTEM_CPP has been added as a compile-time target for the device backend.
- CUDA merge performance is 2-15x faster.
- CUDA comparison sort performance is 1.3-4x faster.
- CUDA set operation performance is 1.5-15x faster.
- TBB reduce_by_key performance is 80% faster.
- Several algorithms have been parallelized with TBB.
- Support for user allocators in vectors has been improved.
- The sparse_vector example is now implemented with merge_by_key instead of sort_by_key.
- Warnings have been eliminated in various contexts.
- Warnings about host or device-only functions called from host device functions have been eliminated in various contexts.
- Documentation about algorithm requirements have been improved.
- Simplified the minimal_custom_backend example.
- Simplified the cuda/custom_temporary_allocation example.
- Simplified the cuda/fallback_allocator example.
- #248: Fix broken
thrust::counting_iterator<float>behavior with OpenMP. - #231, #209: Fix set operation failures with CUDA.
- #187: Fix incorrect occupancy calculation with CUDA.
- #153: Fix broken multi GPU behavior with CUDA.
- #142: Eliminate warning produced by
thrust::random::taus88and MSVC 2010. - #208: Correctly initialize elements in temporary storage when necessary.
- #16: Fix compilation error when sorting bool with CUDA.
- #10: Fix ambiguous overloads of
thrust::reinterpret_tag.
- GCC 4.3 and lower may fail to dispatch thrust::get_temporary_buffer correctly causing infinite recursion in examples such as cuda/custom_temporary_allocation.
- Thanks to Sean Baxter, Bryan Catanzaro, and Manjunath Kudlur for contributing a faster merge implementation for CUDA.
- Thanks to Sean Baxter for contributing a faster set operation implementation for CUDA.
- Thanks to Cliff Woolley for contributing a correct occupancy calculation algorithm.
Thrust 1.6.0 provides an interface for customization and extension and a new backend system based on the Threading Building Blocks library. With this new interface, programmers may customize the behavior of specific algorithms as well as control the allocation of temporary storage or invent entirely new backends. These enhancements also allow multiple different backend systems such as CUDA and OpenMP to coexist within a single program. Support for TBB allows Thrust programs to integrate more naturally into applications which may already employ the TBB task scheduler.
- The header <thrust/experimental/cuda/pinned_allocator.h> has been moved to <thrust/system/cuda/experimental/pinned_allocator.h>
- thrust::experimental::cuda::pinned_allocator has been moved to thrust::cuda::experimental::pinned_allocator
- The macro THRUST_DEVICE_BACKEND has been renamed THRUST_DEVICE_SYSTEM
- The macro THRUST_DEVICE_BACKEND_CUDA has been renamed THRUST_DEVICE_SYSTEM_CUDA
- The macro THRUST_DEVICE_BACKEND_OMP has been renamed THRUST_DEVICE_SYSTEM_OMP
- thrust::host_space_tag has been renamed thrust::host_system_tag
- thrust::device_space_tag has been renamed thrust::device_system_tag
- thrust::any_space_tag has been renamed thrust::any_system_tag
- thrust::iterator_space has been renamed thrust::iterator_system
- Backend Systems
- Threading Building Blocks (TBB) is now supported
- Algorithms
thrust::for_each_nthrust::raw_reference_cast
- Types
thrust::pointerthrust::reference
cuda/custom_temporary_allocationcuda/fallback_allocatordevice_ptrexpandminimal_custom_backendraw_reference_castset_operations
thrust::for_eachnow returns the end of the input range similar to most other algorithms.thrust::pairandthrust::tuplehave swap functionality.- All CUDA algorithms now support large data types.
- Iterators may be dereferenced in user
__device__or__global__functions. - The safe use of different backend systems is now possible within a single binary
- #469
min_elementandmax_elementalgorithms no longer require a const comparison operator
- NVCC may crash when parsing TBB headers on Windows.
Thrust 1.5.3 is a minor bug fix release.
- Avoid warnings about potential race due to
__shared__non-POD variable
Thrust 1.5.2 is a minor bug fix release.
- Fixed warning about C-style initialization of structures
Thrust 1.5.1 is a minor bug fix release.
- Sorting data referenced by permutation_iterators on CUDA produces invalid results
Thrust 1.5.0 provides introduces new programmer productivity and performance
enhancements.
New functionality for creating anonymous "lambda" functions has been added.
A faster host sort provides 2-10x faster performance for sorting arithmetic
types on (single-threaded) CPUs.
A new OpenMP sort provides 2.5x-3.0x speedup over the host sort using a
quad-core CPU.
When sorting arithmetic types with the OpenMP backend the combined performance
improvement is 5.9x for 32-bit integers and ranges from 3.0x (64-bit types) to
14.2x (8-bit types).
A new CUDA reduce_by_key implementation provides 2-3x faster
performance.
- device_ptr no longer unsafely converts to device_ptr without an explicit cast. Use the expression device_pointer_cast(static_cast<int*>(void_ptr.get())) to convert, for example, device_ptr to device_ptr.
- Algorithms:
- Stencil-less
thrust::transform_if.
- Stencil-less
- Lambda placeholders
- lambda
- Host sort is 2-10x faster for arithmetic types
- OMP sort provides speedup over host sort
reduce_by_keyis 2-3x fasterreduce_by_keyno longer requires O(N) temporary storage- CUDA scan algorithms are 10-40% faster
host_vectoranddevice_vectorare now documented- out-of-memory exceptions now provide detailed information from CUDART
- improved histogram example
device_referencenow has a specialized swapreduce_by_keyand scan algorithms are compatible withdiscard_iterator
- #44: Allow
thrust::host_vectorto compile whenvalue_typeuses__align__. - #198: Allow
thrust::adjacent_differenceto permit safe in-situ operation. - #303: Make thrust thread-safe.
- #313: Avoid race conditions in
thrust::device_vector::insert. - #314: Avoid unintended ADL invocation when dispatching copy.
- #365: Fix merge and set operation failures.
- None
- Thanks to Manjunath Kudlur for contributing his Carbon library, from which the lambda functionality is derived.
- Thanks to Jean-Francois Bastien for suggesting a fix for #303.
Thrust 1.4.0 is the first release of Thrust to be included in the CUDA Toolkit. Additionally, it brings many feature and performance improvements. New set theoretic algorithms operating on sorted sequences have been added. Additionally, a new fancy iterator allows discarding redundant or otherwise unnecessary output from algorithms, conserving memory storage and bandwidth.
- Eliminations
thrust/is_sorted.hthrust/utility.hthrust/set_intersection.hthrust/experimental/cuda/ogl_interop_allocator.hand the functionality thereinthrust::deprecated::copy_whenthrust::deprecated::absolute_valuethrust::deprecated::copy_whenthrust::deprecated::absolute_valuethrust::deprecated::copy_whenthrust::deprecated::absolute_valuethrust::gatherandthrust::scatterfrom host to device and vice versa are no longer supported.- Operations which modify the elements of a thrust::device_vector are no longer available from source code compiled without nvcc when the device backend is CUDA. Instead, use the idiom from the cpp_interop example.
-
Algorithms:
thrust::copy_nthrust::mergethrust::set_differencethrust::set_symmetric_differencethrust::set_union
-
Types
thrust::discard_iterator
-
Device Support:
- Compute Capability 2.1 GPUs.
- run_length_decoding
- Compilation warnings are substantially reduced in various contexts.
- The compilation time of thrust::sort, thrust::stable_sort, thrust::sort_by_key, and thrust::stable_sort_by_key are substantially reduced.
- A fast sort implementation is used when sorting primitive types with thrust::greater.
- The performance of thrust::set_intersection is improved.
- The performance of thrust::fill is improved on SM 1.x devices.
- A code example is now provided in each algorithm's documentation.
- thrust::reverse now operates in-place
- #212:
thrust::set_intersectionworks correctly for large input sizes. - #275:
thrust::counting_iteratorandthrust::constant_iteratorwork correctly with OpenMP as the backend when compiling with optimization. - #256:
minandmaxcorrectly return their first argument as a tie-breaker - #248:
NDEBUGis interpreted incorrectly
- NVCC may generate code containing warnings when compiling some Thrust algorithms.
- When compiling with
-arch=sm_1x, some Thrust algorithms may cause NVCC to issue benign pointer advisories. - When compiling with
-arch=sm_1xand -G, some Thrust algorithms may fail to execute correctly. thrust::inclusive_scan,thrust::exclusive_scan,thrust::inclusive_scan_by_key, andthrust::exclusive_scan_by_keyare currently incompatible withthrust::discard_iterator.
- Thanks to David Tarjan for improving the performance of set_intersection.
- Thanks to Duane Merrill for continued help with sort.
- Thanks to Nathan Whitehead for help with CUDA Toolkit integration.
Thrust 1.3.0 provides support for CUDA Toolkit 3.2 in addition to many feature and performance enhancements. Performance of the sort and sort_by_key algorithms is improved by as much as 3x in certain situations. The performance of stream compaction algorithms, such as copy_if, is improved by as much as 2x. CUDA errors are now converted to runtime exceptions using the system_error interface. Combined with a debug mode, also new in 1.3, runtime errors can be located with greater precision. Lastly, a few header files have been consolidated or renamed for clarity. See the deprecations section below for additional details.
- Promotions
- thrust::experimental::inclusive_segmented_scan has been renamed thrust::inclusive_scan_by_key and exposes a different interface
- thrust::experimental::exclusive_segmented_scan has been renamed thrust::exclusive_scan_by_key and exposes a different interface
- thrust::experimental::partition_copy has been renamed thrust::partition_copy and exposes a different interface
- thrust::next::gather has been renamed thrust::gather
- thrust::next::gather_if has been renamed thrust::gather_if
- thrust::unique_copy_by_key has been renamed thrust::unique_by_key_copy
- Deprecations
- thrust::copy_when has been renamed thrust::deprecated::copy_when
- thrust::absolute_value has been renamed thrust::deprecated::absolute_value
- The header thrust/set_intersection.h is now deprecated; use thrust/set_operations.h instead
- The header thrust/utility.h is now deprecated; use thrust/swap.h instead
- The header thrust/swap_ranges.h is now deprecated; use thrust/swap.h instead
- Eliminations
- thrust::deprecated::gather
- thrust::deprecated::gather_if
- thrust/experimental/arch.h and the functions therein
- thrust/sorting/merge_sort.h
- thrust/sorting/radix_sort.h
- NVCC 2.3 is no longer supported
-
Algorithms:
thrust::exclusive_scan_by_keythrust::findthrust::find_ifthrust::find_if_notthrust::inclusive_scan_by_keythrust::is_partitionedthrust::is_sorted_untilthrust::mismatchthrust::partition_pointthrust::reversethrust::reverse_copythrust::stable_partition_copy
-
Types:
thrust::system_errorand related types.thrust::experimental::cuda::ogl_interop_allocator.thrust::bit_and,thrust::bit_or, andthrust::bit_xor.
-
Device Support:
- GF104-based GPUs.
- opengl_interop.cu
- repeated_range.cu
- simple_moving_average.cu
- sparse_vector.cu
- strided_range.cu
- Performance of thrust::sort and thrust::sort_by_key is substantially improved for primitive key types
- Performance of thrust::copy_if is substantially improved
- Performance of thrust::reduce and related reductions is improved
- THRUST_DEBUG mode added
- Callers of Thrust functions may detect error conditions by catching thrust::system_error, which derives from std::runtime_error
- The number of compiler warnings generated by Thrust has been substantially reduced
- Comparison sort now works correctly for input sizes > 32M
- min & max usage no longer collides with <windows.h> definitions
- Compiling against the OpenMP backend no longer requires nvcc
- Performance of device_vector initialized in .cpp files is substantially improved in common cases
- Performance of thrust::sort_by_key on the host is substantially improved
- Debug device code now compiles correctly
- thrust::uninitialized_copy and thrust::uninitialized_fill now dispatch constructors on the device rather than the host
- #212 set_intersection is known to fail for large input sizes
- partition_point is known to fail for 64b types with nvcc 3.2
Acknowledgments
- Thanks to Duane Merrill for contributing a fast CUDA radix sort implementation
- Thanks to Erich Elsen for contributing an implementation of find_if
- Thanks to Andrew Corrigan for contributing changes which allow the OpenMP backend to compile in the absence of nvcc
- Thanks to Andrew Corrigan, Cliff Wooley, David Coeurjolly, Janick Martinez Esturo, John Bowers, Maxim Naumov, Michael Garland, and Ryuta Suzuki for bug reports
- Thanks to Cliff Woolley for help with testing
Thrust 1.2.1 is a small bug fix release that is compatible with the CUDA Toolkit 3.1 release.
thrust::inclusive_scanandthrust::exclusive_scanmay fail with very large types.- MSVC may fail to compile code using both sort and binary search algorithms.
thrust::uninitialized_fillandthrust::uninitialized_copydispatch constructors on the host rather than the device.- #109: Some algorithms may exhibit poor performance with the OpenMP backend with large numbers (>= 6) of CPU threads.
thrust::default_random_engine::discardis not accelerated with NVCC 2.3- NVCC 3.1 may fail to compile code using types derived from
thrust::subtract_with_carry_engine, such asthrust::ranlux24andthrust::ranlux48.
Thrust 1.2.0 introduces support for compilation to multicore CPUs and the Ocelot virtual machine, and several new facilities for pseudo-random number generation. New algorithms such as set intersection and segmented reduction have also been added. Lastly, improvements to the robustness of the CUDA backend ensure correctness across a broad set of (uncommon) use cases.
thrust::gather's interface was incorrect and has been removed. The old interface is deprecated but will be preserved for Thrust version 1.2 atthrust::deprecated::gatherandthrust::deprecated::gather_if. The new interface is provided atthrust::next::gatherandthrust::next::gather_if. The new interface will be promoted tothrust::in Thrust version 1.3. For more details, please refer to this thread.- The
thrust::sortingnamespace has been deprecated in favor of the top-level sorting functions, such asthrust::sortandthrust::sort_by_key. - Removed support for
thrust::equalbetween host & device sequences. - Removed support for
thrust::scatterbetween host & device sequences.
- Algorithms:
thrust::reduce_by_keythrust::set_intersectionthrust::unique_copythrust::unique_by_keythrust::unique_copy_by_key
- Types
- Random Number Generation:
thrust::discard_block_enginethrust::default_random_enginethrust::linear_congruential_enginethrust::linear_feedback_shift_enginethrust::subtract_with_carry_enginethrust::xor_combine_enginethrust::minstd_randthrust::minstd_rand0thrust::ranlux24thrust::ranlux48thrust::ranlux24_basethrust::ranlux48_basethrust::taus88thrust::uniform_int_distributionthrust::uniform_real_distributionthrust::normal_distribution(experimental)
- Function Objects:
thrust::project1stthrust::project2nd
thrust::tie- Fancy Iterators:
thrust::permutation_iteratorthrust::reverse_iterator
- Vector Functions:
operator!=rbegincrbeginrendcrenddatashrink_to_fit
- Device Support:
- Multicore CPUs via OpenMP.
- Fermi-class GPUs.
- Ocelot virtual machines.
- Support for NVCC 3.0.
cpp_integrationhistogrammodemonte_carlomonte_carlo_disjoint_sequencespadded_grid_reductionpermutation_iteratorrow_sumrun_length_encodingsegmented_scanstream_compactionsummary_statisticstransform_iteratorword_count
- Integer sorting performance is improved when max is large but (max - min) is small and when min is negative
- Performance of
thrust::inclusive_scanandthrust::exclusive_scanis improved by 20-25% for primitive types.
- #8 cause a compiler error if the required compiler is not found rather than a mysterious error at link time
- #42 device_ptr & device_reference are classes rather than structs, eliminating warnings on certain platforms
- #46 gather & scatter handle any space iterators correctly
- #51 thrust::experimental::arch functions gracefully handle unrecognized GPUs
- #52 avoid collisions with common user macros such as BLOCK_SIZE
- #62 provide better documentation for device_reference
- #68 allow built-in CUDA vector types to work with device_vector in pure C++ mode
- #102 eliminated a race condition in device_vector::erase
- various compilation warnings eliminated
- inclusive_scan & exclusive_scan may fail with very large types
- MSVC may fail to compile code using both sort and binary search algorithms
- uninitialized_fill & uninitialized_copy dispatch constructors on the host rather than the device
- #109 some algorithms may exhibit poor performance with the OpenMP backend with large numbers (>= 6) of CPU threads
- default_random_engine::discard is not accelerated with nvcc 2.3
- Thanks to Gregory Diamos for contributing a CUDA implementation of set_intersection
- Thanks to Ryuta Suzuki & Gregory Diamos for rigorously testing Thrust's unit tests and examples against Ocelot
- Thanks to Tom Bradley for contributing an implementation of normal_distribution
- Thanks to Joseph Rhoads for contributing the example summary_statistics
Thrust 1.1.1 is a small bug fix release that is compatible with the CUDA Toolkit 2.3a release and Mac OSX Snow Leopard.
Thrust 1.1.0 introduces fancy iterators, binary search functions, and several specialized reduction functions. Experimental support for segmented scans has also been added.
thrust::counting_iteratorhas been moved into thethrustnamespace (previouslythrust::experimental).
- Algorithms:
thrust::copy_ifthrust::lower_boundthrust::upper_boundthrust::vectorized lower_boundthrust::vectorized upper_boundthrust::equal_rangethrust::binary_searchthrust::vectorized binary_searchthrust::all_ofthrust::any_ofthrust::none_ofthrust::minmax_elementthrust::advancethrust::inclusive_segmented_scan(experimental)thrust::exclusive_segmented_scan(experimental)
- Types:
thrust::pairthrust::tuplethrust::device_malloc_allocator
- Fancy Iterators:
thrust::constant_iteratorthrust::counting_iteratorthrust::transform_iteratorthrust::zip_iterator
- Computing the maximum absolute difference between vectors.
- Computing the bounding box of a two-dimensional point set.
- Sorting multiple arrays together (lexicographical sorting).
- Constructing a summed area table.
- Using
thrust::zip_iteratorto mimic an array of structs. - Using
thrust::constant_iteratorto increment array values.
- Added pinned memory allocator (experimental).
- Added more methods to host_vector & device_vector (issue #4).
- Added variant of remove_if with a stencil argument (issue #29).
- Scan and reduce use cudaFuncGetAttributes to determine grid size.
- Exceptions are reported when temporary device arrays cannot be allocated.
- #5: Make vector work for larger data types
- #9: stable_partition_copy doesn't respect OutputIterator concept semantics
- #10: scans should return OutputIterator
- #16: make algorithms work for larger data types
- #27: Dispatch radix_sort even when comp=less is explicitly provided
- Using functors with Thrust entry points may not compile on Mac OSX with gcc 4.0.1.
thrust::uninitialized_copyandthrust::uninitialized_filldispatch constructors on the host rather than the device.thrust::inclusive_scan,thrust::inclusive_scan_by_key,thrust::exclusive_scan, andthrust::exclusive_scan_by_keymay fail when used with large types with the CUDA Toolkit 3.1.
First production release of Thrust.
- Rename top level namespace
komradetothrust. - Move
thrust::partition_copy&thrust::stable_partition_copyintothrust::experimentalnamespace until we can easily provide the standard interface. - Rename
thrust::rangetothrust::sequenceto avoid collision with Boost.Range. - Rename
thrust::copy_iftothrust::copy_whendue to semantic differences with C++0xstd::copy_if.
- Add C++0x style
cbegin&cendmethods tothrust::host_vectorandthrust::device_vector. - Add
thrust::transform_iffunction. - Add stencil versions of
thrust::replace_if&thrust::replace_copy_if. - Allow
counting_iteratorto work withthrust::for_each. - Allow types with constructors in comparison
thrust::sortandthrust::reduce.
thrust::merge_sortandthrust::stable_merge_sortare now 2x to 5x faster when executed on the parallel device.
- Komrade 6: Workaround an issue where an incremented iterator causes NVCC to crash.
- Komrade 7: Fix an issue where
const_iterators could not be passed tothrust::transform.