From 87fd06f60855823446bd7cda52154025f919f9b4 Mon Sep 17 00:00:00 2001 From: Karthikeyan Natarajan Date: Sat, 4 Dec 2021 00:59:38 +0530 Subject: [PATCH 1/4] add output_writer_iterator --- thrust/iterator/output_writer_iterator.cuh | 175 +++++++++++++++++++++ 1 file changed, 175 insertions(+) create mode 100644 thrust/iterator/output_writer_iterator.cuh diff --git a/thrust/iterator/output_writer_iterator.cuh b/thrust/iterator/output_writer_iterator.cuh new file mode 100644 index 000000000..d2757a570 --- /dev/null +++ b/thrust/iterator/output_writer_iterator.cuh @@ -0,0 +1,175 @@ +/* + * Copyright (c) 2021, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +// Output writer iterator +#pragma once + +#include + +namespace thrust { +namespace detail { + +// Proxy reference that calls BinaryFunction with Iterator value and the rhs of assignment operator +template +class output_writer_iterator_proxy { + public: + __host__ __device__ output_writer_iterator_proxy(const Iterator& index_iter, BinaryFunction fun) + : index_iter(index_iter), fun(fun) + { + } + template + __host__ __device__ output_writer_iterator_proxy operator=(const T& x) + { + fun(*index_iter, x); + return *this; + } + + private: + Iterator index_iter; + BinaryFunction fun; +}; + +// Register output_writer_iterator_proxy with 'is_proxy_reference' from +// type_traits to enable its use with algorithms. +template +struct is_proxy_reference> + : public thrust::detail::true_type { +}; + +} // namespace detail +} // namespace thrust + +namespace cudf { +/** + * @brief Transform output iterator with custom writer binary function which takes index and value. + * + * @code {.cpp} + * #include + * #include + * #include + * #include + * + * struct set_bits_field { + * int* bitfield; + * __device__ inline void set_bit(size_t bit_index) + * { + * atomicOr(&bitfield[bit_index/32], (int{1} << (bit_index % 32))); + * } + * __device__ inline void clear_bit(size_t bit_index) + * { + * atomicAnd(&bitfield[bit_index / 32], ~(int{1} << (bit_index % 32))); + * } + * // Index, value + * __device__ void operator()(size_t i, bool x) + * { + * if (x) + * set_bit(i); + * else + * clear_bit(i); + * } + * }; + * + * thrust::device_vector v(1, 0x0000); + * auto result_begin = cudf::make_output_writer_iterator(thrust::make_counting_iterator(0), + * set_bits_field{v.data().get()}); + * auto value = thrust::make_transform_iterator(thrust::make_counting_iterator(0), [] __device__ + * (int x) { return x%2; + * }); + * thrust::copy(thrust::device, value, value+32, result_begin); + * + * #include + * #include + * #include + * #include + * + * struct set_bits_field { + * int* bitfield; + * __device__ inline void set_bit(size_t bit_index) + * { + * atomicOr(&bitfield[bit_index/32], (int{1} << (bit_index % 32))); + * } + * __device__ inline void clear_bit(size_t bit_index) + * { + * atomicAnd(&bitfield[bit_index / 32], ~(int{1} << (bit_index % 32))); + * } + * // Index, value + * __device__ void operator()(size_t i, bool x) + * { + * if (x) + * set_bit(i); + * else + * clear_bit(i); + * } + * }; + * + * thrust::device_vector v(1, 0x0000); + * auto result_begin = cudf::make_output_writer_iterator(thrust::make_counting_iterator(0), + * set_bits_field{v.data().get()}); + * auto value = thrust::make_transform_iterator(thrust::make_counting_iterator(0), + * [] __device__ (int x) { return x%2; }); + * thrust::copy(thrust::device, value, value+32, result_begin); + * int(v[0]); // returns 0xaaaaaaaa; + * @endcode + * + * + * @tparam BinaryFunction Binary function to be called with the Iterator value and the rhs of + * assignment operator. + * @tparam Iterator iterator type that acts as index of the output. + */ +template +class output_writer_iterator + : public thrust::iterator_adaptor< + output_writer_iterator, + Iterator, + thrust::use_default, + thrust::use_default, + thrust::use_default, + thrust::detail::output_writer_iterator_proxy> { + public: + // parent class. + typedef thrust::iterator_adaptor< + output_writer_iterator, + Iterator, + thrust::use_default, + thrust::use_default, + thrust::use_default, + thrust::detail::output_writer_iterator_proxy> + super_t; + // friend thrust::iterator_core_access to allow it access to the private interface dereference() + friend class thrust::iterator_core_access; + __host__ __device__ output_writer_iterator(Iterator const& x, BinaryFunction fun) + : super_t(x), fun(fun) + { + } + + private: + BinaryFunction fun; + + // thrust::iterator_core_access accesses this function + __host__ __device__ typename super_t::reference dereference() const + { + return thrust::detail::output_writer_iterator_proxy( + this->base_reference(), fun); + } +}; + +template +output_writer_iterator __host__ __device__ +make_output_writer_iterator(Iterator out, BinaryFunction fun) +{ + return output_writer_iterator(out, fun); +} // end make_output_writer_iterator +} // namespace cudf From 443ce4e81ea3ace439f0035f7ccfd4924f026d82 Mon Sep 17 00:00:00 2001 From: Karthikeyan Natarajan Date: Sat, 4 Dec 2021 01:14:19 +0530 Subject: [PATCH 2/4] update example code, rename namespace --- thrust/iterator/output_writer_iterator.cuh | 41 ++-------------------- 1 file changed, 3 insertions(+), 38 deletions(-) diff --git a/thrust/iterator/output_writer_iterator.cuh b/thrust/iterator/output_writer_iterator.cuh index d2757a570..67636e992 100644 --- a/thrust/iterator/output_writer_iterator.cuh +++ b/thrust/iterator/output_writer_iterator.cuh @@ -50,47 +50,12 @@ struct is_proxy_reference }; } // namespace detail -} // namespace thrust -namespace cudf { /** * @brief Transform output iterator with custom writer binary function which takes index and value. * * @code {.cpp} - * #include - * #include - * #include - * #include - * - * struct set_bits_field { - * int* bitfield; - * __device__ inline void set_bit(size_t bit_index) - * { - * atomicOr(&bitfield[bit_index/32], (int{1} << (bit_index % 32))); - * } - * __device__ inline void clear_bit(size_t bit_index) - * { - * atomicAnd(&bitfield[bit_index / 32], ~(int{1} << (bit_index % 32))); - * } - * // Index, value - * __device__ void operator()(size_t i, bool x) - * { - * if (x) - * set_bit(i); - * else - * clear_bit(i); - * } - * }; - * - * thrust::device_vector v(1, 0x0000); - * auto result_begin = cudf::make_output_writer_iterator(thrust::make_counting_iterator(0), - * set_bits_field{v.data().get()}); - * auto value = thrust::make_transform_iterator(thrust::make_counting_iterator(0), [] __device__ - * (int x) { return x%2; - * }); - * thrust::copy(thrust::device, value, value+32, result_begin); - * - * #include + * #include * #include * #include * #include @@ -116,7 +81,7 @@ namespace cudf { * }; * * thrust::device_vector v(1, 0x0000); - * auto result_begin = cudf::make_output_writer_iterator(thrust::make_counting_iterator(0), + * auto result_begin = thrust::make_output_writer_iterator(thrust::make_counting_iterator(0), * set_bits_field{v.data().get()}); * auto value = thrust::make_transform_iterator(thrust::make_counting_iterator(0), * [] __device__ (int x) { return x%2; }); @@ -172,4 +137,4 @@ make_output_writer_iterator(Iterator out, BinaryFunction fun) { return output_writer_iterator(out, fun); } // end make_output_writer_iterator -} // namespace cudf +} // namespace thrust From ae217b4cd7b524c612b248934b3a1e074f2a02b0 Mon Sep 17 00:00:00 2001 From: Karthikeyan <6488848+karthikeyann@users.noreply.github.com> Date: Tue, 25 Jan 2022 20:30:10 +0530 Subject: [PATCH 3/4] Apply suggestions from code review (allisonvacanti) Co-authored-by: Allison Vacanti --- thrust/iterator/output_writer_iterator.cuh | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/thrust/iterator/output_writer_iterator.cuh b/thrust/iterator/output_writer_iterator.cuh index 67636e992..7d5d4ce42 100644 --- a/thrust/iterator/output_writer_iterator.cuh +++ b/thrust/iterator/output_writer_iterator.cuh @@ -19,7 +19,7 @@ #include -namespace thrust { +THRUST_NAMESPACE_BEGIN namespace detail { // Proxy reference that calls BinaryFunction with Iterator value and the rhs of assignment operator @@ -80,13 +80,13 @@ struct is_proxy_reference * } * }; * - * thrust::device_vector v(1, 0x0000); + * thrust::device_vector v(1, 0x00000000); * auto result_begin = thrust::make_output_writer_iterator(thrust::make_counting_iterator(0), * set_bits_field{v.data().get()}); * auto value = thrust::make_transform_iterator(thrust::make_counting_iterator(0), * [] __device__ (int x) { return x%2; }); * thrust::copy(thrust::device, value, value+32, result_begin); - * int(v[0]); // returns 0xaaaaaaaa; + * assert(v[0] == 0xaaaaaaaa); * @endcode * * @@ -137,4 +137,4 @@ make_output_writer_iterator(Iterator out, BinaryFunction fun) { return output_writer_iterator(out, fun); } // end make_output_writer_iterator -} // namespace thrust +THRUST_NAMESPACE_END From 1a73995ac9e6f66099496ad903e501ac1dfdabc0 Mon Sep 17 00:00:00 2001 From: Karthikeyan Natarajan Date: Thu, 12 May 2022 00:24:00 +0530 Subject: [PATCH 4/4] rename cuh to h --- .../{output_writer_iterator.cuh => output_writer_iterator.h} | 0 1 file changed, 0 insertions(+), 0 deletions(-) rename thrust/iterator/{output_writer_iterator.cuh => output_writer_iterator.h} (100%) diff --git a/thrust/iterator/output_writer_iterator.cuh b/thrust/iterator/output_writer_iterator.h similarity index 100% rename from thrust/iterator/output_writer_iterator.cuh rename to thrust/iterator/output_writer_iterator.h