diff --git a/include/cuco/detail/bloom_filter/bloom_filter_impl.cuh b/include/cuco/detail/bloom_filter/bloom_filter_impl.cuh index 2669dd41e..15b70b118 100644 --- a/include/cuco/detail/bloom_filter/bloom_filter_impl.cuh +++ b/include/cuco/detail/bloom_filter/bloom_filter_impl.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2024, NVIDIA CORPORATION. + * Copyright (c) 2024-2025, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -29,10 +29,10 @@ #include // TODO #include once available #include #include +#include #include #include #include -#include #include #include @@ -182,7 +182,7 @@ class bloom_filter_impl { stream.get())); } else { auto const always_true = thrust::constant_iterator{true}; - this->add_if_async(first, last, always_true, thrust::identity{}, stream); + this->add_if_async(first, last, always_true, cuda::std::identity{}, stream); } } @@ -283,7 +283,7 @@ class bloom_filter_impl { cuda::stream_ref stream) const noexcept { auto const always_true = thrust::constant_iterator{true}; - this->contains_if_async(first, last, always_true, thrust::identity{}, output_begin, stream); + this->contains_if_async(first, last, always_true, cuda::std::identity{}, output_begin, stream); } template diff --git a/include/cuco/detail/hash_functions/identity_hash.cuh b/include/cuco/detail/hash_functions/identity_hash.cuh index 995bded1d..cae4504a0 100644 --- a/include/cuco/detail/hash_functions/identity_hash.cuh +++ b/include/cuco/detail/hash_functions/identity_hash.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2024, NVIDIA CORPORATION. + * Copyright (c) 2024-2025, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -16,8 +16,8 @@ #pragma once +#include #include -#include namespace cuco::detail { @@ -33,7 +33,7 @@ namespace cuco::detail { * @tparam Key The type of the values to hash */ template -struct identity_hash : private thrust::identity { +struct identity_hash : private cuda::std::identity { using argument_type = Key; ///< The type of the values taken as argument /// The type of the hash values produced using result_type = cuda::std::conditional_t; @@ -49,7 +49,7 @@ struct identity_hash : private thrust::identity { */ __host__ __device__ result_type operator()(Key const& x) const { - return static_cast(thrust::identity::operator()(x)); + return static_cast(cuda::std::identity::operator()(x)); } }; // identity_hash diff --git a/include/cuco/detail/open_addressing/open_addressing_impl.cuh b/include/cuco/detail/open_addressing/open_addressing_impl.cuh index aece06a12..5dd1b94e9 100644 --- a/include/cuco/detail/open_addressing/open_addressing_impl.cuh +++ b/include/cuco/detail/open_addressing/open_addressing_impl.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2023-2024, NVIDIA CORPORATION. + * Copyright (c) 2023-2025, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -31,6 +31,7 @@ #include #include #include +#include #include #include #include @@ -273,7 +274,7 @@ class open_addressing_impl { size_type insert(InputIt first, InputIt last, Ref container_ref, cuda::stream_ref stream) { auto const always_true = thrust::constant_iterator{true}; - return this->insert_if(first, last, always_true, thrust::identity{}, container_ref, stream); + return this->insert_if(first, last, always_true, cuda::std::identity{}, container_ref, stream); } /** @@ -296,7 +297,7 @@ class open_addressing_impl { cuda::stream_ref stream) noexcept { auto const always_true = thrust::constant_iterator{true}; - this->insert_if_async(first, last, always_true, thrust::identity{}, container_ref, stream); + this->insert_if_async(first, last, always_true, cuda::std::identity{}, container_ref, stream); } /** @@ -494,7 +495,7 @@ class open_addressing_impl { { auto const always_true = thrust::constant_iterator{true}; this->contains_if_async( - first, last, always_true, thrust::identity{}, output_begin, container_ref, stream); + first, last, always_true, cuda::std::identity{}, output_begin, container_ref, stream); } /** @@ -569,7 +570,7 @@ class open_addressing_impl { auto const always_true = thrust::constant_iterator{true}; this->find_if_async( - first, last, always_true, thrust::identity{}, output_begin, container_ref, stream); + first, last, always_true, cuda::std::identity{}, output_begin, container_ref, stream); } /** diff --git a/include/cuco/detail/open_addressing/open_addressing_ref_impl.cuh b/include/cuco/detail/open_addressing/open_addressing_ref_impl.cuh index 1676fb816..b8540a779 100644 --- a/include/cuco/detail/open_addressing/open_addressing_ref_impl.cuh +++ b/include/cuco/detail/open_addressing/open_addressing_ref_impl.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2023-2024, NVIDIA CORPORATION. + * Copyright (c) 2023-2025, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -24,6 +24,7 @@ #include #include +#include #include #include #include @@ -1246,7 +1247,7 @@ class open_addressing_ref_impl { // Fill the buffer if any matching keys are found auto const lane_id = probing_tile.thread_rank(); - if (thrust::any_of(thrust::seq, exists, exists + bucket_size, thrust::identity{})) { + if (thrust::any_of(thrust::seq, exists, exists + bucket_size, cuda::std::identity{})) { if constexpr (IsOuter) { found_match = true; } int32_t num_matches[bucket_size];