Thrust 1.13.0 (NVIDIA HPC SDK 21.7)

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 STL reference.

Breaking Changes

• #1459: Remove deprecated aliases thrust::host_space_tag and thrust::device_space_tag. Use the equivalent thrust::host_system_tag and thrust::device_system_tag instead.

New Features

• NVIDIA/cub#306: Add radix-sort support for bfloat16 in thrust::sort. Thanks to Xiang Gao (@zasdfgbnm) for this contribution.
• #1423: thrust::transform_iterator now supports non-copyable types. Thanks to Jake Hemstad (@jrhemstad) for this contribution.
• #1459: Introduce a new THRUST_IGNORE_DEPRECATED_API macro that disables deprecation warnings on Thrust and CUB APIs.

Bug Fixes

• NVIDIA/cub#277: Fixed sanitizer warnings when thrust::sort calls into cub::DeviceRadixSort. Thanks to Andy Adinets (@canonizer) for this contribution.
• #1442: Reduce extraneous comparisons in thrust::sort's merge sort implementation.
• #1447: Fix memory leak and avoid overallocation when calling reserve on Thrust's vector containers. Thanks to Kai Germaschewski (@germasch) for this contribution.

Other Enhancements

• #1405: Update links to standard C++ documentations from sgi to cppreference. Thanks to Muhammad Adeel Hussain (@AdeilH) for this contribution.
• #1432: Updated build instructions in CONTRIBUTING.md to include details on building CUB's test suite as part of Thrust.

Thrust 1.12.1 (CUDA Toolkit 11.4)

Thrust 1.12.1 is a trivial patch release that slightly changes the phrasing of a deprecation message.

Thrust 1.12.0 (NVIDIA HPC SDK 21.3, CUDA Toolkit 11.4)

Summary

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.

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).

Breaking Changes

• #1372: Deprecate Clang < 7 and MSVC < 2019.
• #1376: Standardize thrust::scan_by_key functors / accumulator types. This may change the results from scan_by_key when input, output, and initial value types are not the same type.

New Features

• #1251: Add two new thrust::async:: algorithms: inclusive_scan and exclusive_scan.
• #1334: Add thrust::universal_vector, universal_ptr, and universal_allocator.

Bug Fixes

• #1347: Qualify calls to make_reverse_iterator.
• #1359: Enable stricter warning flags. This fixes several outstanding issues:
  • NVIDIA/cub#221: Overflow in temp_storage_bytes when num_items close 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 GridEvenShare with unsigned offsets.
  • #974: Conversion warnings in thrust::transform_reduce.
  • #1091: Conversion warnings in thrust::counting_iterator.
• #1373: Fix compilation error when a standard library type is wrapped in thrust::optional. Thanks to Vukasin Milovanovic for this contribution.
• #1388: Fix signbit(double) implementation on MSVC.
• #1389: Support building Thrust tests without CUDA enabled.

Other Enhancements

• #1304: Use cub::DeviceScan to implement thrust::exclusive_scan and thrust::inclusive_scan.
• #1362, #1370: Update smoke test naming.
• #1380: Fix typos in set_operation documentation. Thanks to Hongyu Cai for this contribution.
• #1383: Include FreeBSD license in LICENSE.md for thrust::complex implementation.
• #1384: Add missing precondition to thrust::gather documentation.

Thrust 1.11.0 (CUDA Toolkit 11.3)

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.

New Features

• NVIDIA/cub#204: New implementation for thrust::sort on 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.
• #1310, #1312: Various tuple-related APIs have been updated to use variadic templates. Thanks for Andrew Corrigan for these contributions.
• #1297: Optionally add install rules when included with CMake's add_subdirectory. Thanks to Kai Germaschewski for this contribution.

Bug Fixes

• #1309: Fix thrust::shuffle to produce better quality random distributions. Thanks to Rory Mitchell and Daniel Stokes for this contribution.
• #1337: Fix compile-time regression in transform_inclusive_scan and transform_exclusive_scan.
• #1306: Fix binary search middle calculation to avoid overflows. Thanks to Richard Barnes for this contribution.
• #1314: Use size_t for the index type parameter in thrust::tuple_element. Thanks to Andrew Corrigan for this contribution.
• #1329: Fix runtime error when copying an empty thrust::device_vector in MSVC Debug builds. Thanks to Ben Jude for this contribution.
• #1323: Fix and add test for cmake package install rules. Thanks for Keith Kraus and Kai Germaschewski for testing and discussion.
• #1338: Fix GCC version checks in thrust::detail::is_pod implementation. Thanks to Anatoliy Tomilov for this contribution.
• #1289: Partial fixes for Clang 10 as host/c++ compiler. Exposed an nvcc bug that will be fixed in a future version of the CUDA Toolkit (NVBug 3136307).
• #1272: Fix ambiguous iter_swap call when using thrust::partition with STL containers. Thanks to Isaac Deutsch for this contribution.
• #1281: Update our bundled FindTBB.cmake module to support latest MSVC.
• #1298: Use semantic versioning rules for our CMake package's compatibility checks. Thanks to Kai Germaschewski for this contribution.
• #1300: Use FindPackageHandleStandardArgs to print standard status messages when our CMake package is found. Thanks to Kai Germaschewski for this contribution.
• #1320: Use feature-testing instead of a language dialect check for thrust::remove_cvref. Thanks to Andrew Corrigan for this contribution.
• #1319: Suppress GPU deprecation warnings.

Other Enhancements

• 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/cub repository is now NVIDIA/cub
  • Development has moved from the master branch to the main branch.

Thrust 1.10.0 (NVIDIA HPC SDK 20.9, CUDA Toolkit 11.2)

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/thrust/thrust/blob/main/CODE_OF_CONDUCT.md

Breaking Changes

• 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_DIALECT or THRUST_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.

New Features

• #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/thrust/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_subdirectory with the Thrust source root (see #976).
    An example can be found here: https://github.com/thrust/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_DIALECT option has been added for single config mode. Logic that modified CMAKE_CXX_STANDARD and CMAKE_CUDA_STANDARD has 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 searching CMAKE_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.

Other Enhancements

• Contributor documentation: https://github.com/thrust/thrust/blob/main/CONTRIBUTING.md
• Code of Conduct: https://github.com/thrust/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.
• #1221: Allocator and vector classes have been replaced with alias templates. Thanks to Michael Francis for this contribution.
• #1186: Use placeholder expressions to simplify the definitions of a number of algorithms. Thanks to Michael Francis for this contribution.
• #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_iterators helper is no longer needed and has been removed.
• #1255: Always use cudaStreamSynchronize instead of cudaDeviceSynchronize if the execution policy has a stream attached to it. Thanks to Rong Ou for this contribution.
• #1201: Tests for correct handling of legacy and per-thread default streams. Thanks to Rong Ou for this contribution.

Bug Fixes

• #1260: Fix thrust::transform_inclusive_scan with heterogeneous types. Thanks to Rong Ou for this contribution.
• #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.
• #1264: Evaluate CUDA_CUB_RET_IF_FAIL macro argument only once. Thanks to Jason Lowe for this contribution.
• #1262: Add missing <stdexcept> header.
• #1250: Restore some THRUST_DECLTYPE_RETURNS macros in async test implementations.
• #1249: Use std::iota in CUDATestDriver::target_devices. Thanks to Michael Francis for this contribution.
• #1244: Check for macro collisions with system headers during header testing.
• #1224: Remove unnecessary SFINAE contexts from asynchronous algorithms.
• #1190: Make out_of_memory_recovery test trigger faster.
• #1187: Elminate superfluous iterators specific to the CUDA backend.
• #1181: Various fixes for GoUDA.
  Thanks to Andrei Tchouprakov for this contribution.
• #1178, #1229: Use transparent functionals in placeholder expressions, fixing issues with thrust::device_reference and placeholder expressions and thrust::find with asymmetric equality operators.
• thrust/thrust#1153: Switch to placement new instead of assignment to construct items in uninitialized memory. Thanks to Hugh Winkler for this contribution.
• #1050: Fix compilation of asynchronous algorithms when RDC is enabled.
• #1042: Correct return type of thrust::detail::predicate_to_integral from bool to IntegralType. Thanks to Andreas Hehn for this contribution.
• #1009: Avoid returning uninitialized allocators. Thanks to Zhihao Yuan for this contribution.
• #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.
• #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_function for void return types because MSVC is not a fan of the pattern return static_cast<void>(expr);.
• Replace deprecated tbb/tbb_thread.h with <thread>.
• Fix overcounting of initial value in TBB scans.
• Use thrust::advance instead of += for generic iterators.
• Wrap the OMP flags in -Xcompiler for NVCC
• Extend ASSERT_STATIC_ASSERT skip for the OMP backend.
• Add missing header caught by tbb.cuda configs.
• Fix "unsafe API" warnings in examples on MSVC: s/fopen/fstream/
• Various C++17 fixes.

Thrust 1.9.10-1 (NVIDIA HPC SDK 20.7, CUDA Toolkit 11.1)

Thrust 1.9.10-1 is the minor release accompanying the NVIDIA HPC SDK 20.7 release and the CUDA Toolkit 11.1 release.

Bug Fixes

• #1214, NVBug 200619442: Stop using std::allocator APIs deprecated in C++17.
• #1216, NVBug 200540293: Make thrust::optional work 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 (NVIDIA HPC SDK 20.5)

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. All other deprecated platforms will be dropped in the near future.

Breaking Changes

• #1130: CMake find_package support. This is significant because there is a legacy FindThrust.cmake script 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 CMake find_package support and migrating away from the legacy FindThrust.cmake here
• #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 verison 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) or THRUST_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.

New Features

• #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++.

Other Enhancements

• #1129: Refactored RDC handling in CMake to be a global option and not create two targets for each example and test.

Bug Fixes

• #1129: Fix the legacy thrust::return_temporary_buffer API 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 where cudaMalloc is 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::copy to not copy the policy, resolving use-afer-move issues.
• #1145: When cleaning up type names in unittest::base_class_name, only call std::string::replace if we found the substring we are looking to replace.
• #1139: Don't use cxx::__demangle in NVC++.
• #1102: Don't use thrust::detail::normal_distribution_nvcc for Feta because it uses erfcinv, a non-standard function that Feta doesn't have.

Thrust 1.9.9 (CUDA Toolkit 11.0)

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. As of this release, C++03, C++11, GCC < 5, Clang < 6, and MSVC < 2017 are 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.

Breaking Changes

• #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 verison 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) or THRUST_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.

New Features

• #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 an if statement around device-only code.
  • THRUST_INCLUDE_DEVICE_CODE, which should be used in an #if preprocessor directive inside of the if statement mentioned in the prior bullet.
  • THRUST_IS_HOST_CODE, which should be used in an if statement around host-only code.
  • THRUST_INCLUDE_HOST_CODE, which should be used in an #if preprocessor directive inside of the if statement 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 as thrust::zip_iterator does. 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.

Other Enhancements

• #1029: Thrust is now built and tested with NVCC warnings treated as errors.
• #1029: MSVC C++11 support.
• #1029: THRUST_DEPRECATED abstraction for generating compile-time deprecation warning messages.
• #1029: thrust::pointer<T>::pointer_to(reference).
• #1070: Unit test for thrust::inclusive_scan with a user defined types. Thanks to Conor Hoekstra for this contribution.

Bug Fixes

• #1088: Allow thrust::replace to take functions that have non-const operator().
• #1094: Add missing constexpr to par_t constructors. 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_NS macros.
• #1029: Fix C++ dialect detection on newer MSVC.
• #1029 Use _Pragma/__pragma instead of #pragma in macros.
• #1029: Replace raw __cplusplus checks with the appropriate Thrust macros.
• #1105: Add a missing <math.h> include.
• #1103: Fix regression of thrust::detail::temporary_allocator with 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 (NVIDIA HPC SDK 20.3)

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.

Thrust 1.9.8 (CUDA Toolkit 11.0 Early Access)

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.

Breaking Changes

• 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.

Other Enhancements

• 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 than 2^31-1 elements. Notably, thrust::sort is still limited to less than 2^31-1 elements.
• 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).

Bug Fixes

• #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_alloc when thrust::system::cuda::malloc fails 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-copy complains 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_enum for thrust::event_errc. Thanks to Toru Niina for this contribution.
• #1027: Add tests for thrust::tuple_for_each and thrust::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_FUNCTION with THRUST_RUNTIME_FUNCTION.
• Correct typo in thrust::transform documentation. Thanks to Eden Yefet for this contribution.

Known Issues

• thrust::sort remains limited to 2^31-1 elements for now.