diff --git a/cpp/include/cudf/column/column_device_view.cuh b/cpp/include/cudf/column/column_device_view.cuh index 35851a99822..b1ff0bbaea7 100644 --- a/cpp/include/cudf/column/column_device_view.cuh +++ b/cpp/include/cudf/column/column_device_view.cuh @@ -16,6 +16,7 @@ #pragma once #include +#include #include #include #include @@ -442,10 +443,11 @@ class alignas(16) column_device_view : public detail::column_device_view_base { __device__ T element(size_type element_index) const noexcept { size_type index = element_index + offset(); // account for this view's _offset - auto const* d_offsets = d_children[strings_column_view::offsets_column_index].data(); char const* d_strings = d_children[strings_column_view::chars_column_index].data(); - size_type offset = d_offsets[index]; - return string_view{d_strings + offset, d_offsets[index + 1] - offset}; + auto const offsets = d_children[strings_column_view::offsets_column_index]; + auto const itr = cudf::detail::input_offsetalator(offsets.head(), offsets.type()); + auto const offset = itr[index]; + return string_view{d_strings + offset, static_cast(itr[index + 1] - offset)}; } private: diff --git a/cpp/include/cudf/detail/indexalator.cuh b/cpp/include/cudf/detail/indexalator.cuh index 6532dae3695..4d261c54b29 100644 --- a/cpp/include/cudf/detail/indexalator.cuh +++ b/cpp/include/cudf/detail/indexalator.cuh @@ -56,10 +56,69 @@ namespace detail { * auto result = thrust::find(thrust::device, begin, end, size_type{12} ); * @endcode */ -using input_indexalator = input_normalator; +struct input_indexalator : base_normalator { + friend struct base_normalator; // for CRTP + + using reference = cudf::size_type const; // this keeps STL and thrust happy + + input_indexalator() = default; + input_indexalator(input_indexalator const&) = default; + input_indexalator(input_indexalator&&) = default; + input_indexalator& operator=(input_indexalator const&) = default; + input_indexalator& operator=(input_indexalator&&) = default; + + /** + * @brief Indirection operator returns the value at the current iterator position + */ + __device__ inline cudf::size_type operator*() const { return operator[](0); } + + /** + * @brief Dispatch functor for resolving a Integer value from any integer type + */ + struct normalize_type { + template ())> + __device__ cudf::size_type operator()(void const* tp) + { + return static_cast(*static_cast(tp)); + } + template ())> + __device__ cudf::size_type operator()(void const*) + { + CUDF_UNREACHABLE("only integral types are supported"); + } + }; + + /** + * @brief Array subscript operator returns a value at the input + * `idx` position as a `Integer` value. + */ + __device__ inline cudf::size_type operator[](size_type idx) const + { + void const* tp = p_ + (idx * this->width_); + return type_dispatcher(this->dtype_, normalize_type{}, tp); + } + + /** + * @brief Create an input index normalizing iterator + * + * Use the indexalator_factory to create an iterator instance. + * + * @param data Pointer to an integer array in device memory. + * @param dtype Type of data in data + * @param offset Applied to the data pointer per size of the type + */ + CUDF_HOST_DEVICE input_indexalator(void const* data, data_type dtype, cudf::size_type offset = 0) + : base_normalator(dtype), p_{static_cast(data)} + { + p_ += offset * this->width_; + } + + protected: + char const* p_; /// pointer to the integer data in device memory +}; /** - * @brief The index normalizing output iterator. + * @brief The index normalizing output iterator * * This is an iterator that can be used for index types (integers) without * requiring a type-specific instance. It can be used for any iterator @@ -82,7 +141,75 @@ using input_indexalator = input_normalator; * thrust::less()); * @endcode */ -using output_indexalator = output_normalator; +struct output_indexalator : base_normalator { + friend struct base_normalator; // for CRTP + + using reference = output_indexalator const&; // required for output iterators + + output_indexalator() = default; + output_indexalator(output_indexalator const&) = default; + output_indexalator(output_indexalator&&) = default; + output_indexalator& operator=(output_indexalator const&) = default; + output_indexalator& operator=(output_indexalator&&) = default; + + /** + * @brief Indirection operator returns this iterator instance in order + * to capture the `operator=(Integer)` calls. + */ + __device__ inline reference operator*() const { return *this; } + + /** + * @brief Array subscript operator returns an iterator instance at the specified `idx` position. + * + * This allows capturing the subsequent `operator=(Integer)` call in this class. + */ + __device__ inline output_indexalator const operator[](size_type idx) const + { + output_indexalator tmp{*this}; + tmp.p_ += (idx * this->width_); + return tmp; + } + + /** + * @brief Dispatch functor for setting the index value from a size_type value. + */ + struct normalize_type { + template ())> + __device__ void operator()(void* tp, cudf::size_type const value) + { + (*static_cast(tp)) = static_cast(value); + } + template ())> + __device__ void operator()(void*, cudf::size_type const) + { + CUDF_UNREACHABLE("only index types are supported"); + } + }; + + /** + * @brief Assign an Integer value to the current iterator position + */ + __device__ inline reference operator=(cudf::size_type const value) const + { + void* tp = p_; + type_dispatcher(this->dtype_, normalize_type{}, tp, value); + return *this; + } + + /** + * @brief Create an output normalizing iterator + * + * @param data Pointer to an integer array in device memory. + * @param dtype Type of data in data + */ + CUDF_HOST_DEVICE output_indexalator(void* data, data_type dtype) + : base_normalator(dtype), p_{static_cast(data)} + { + } + + protected: + char* p_; /// pointer to the integer data in device memory +}; /** * @brief Use this class to create an indexalator instance. @@ -92,14 +219,12 @@ struct indexalator_factory { * @brief A type_dispatcher functor to create an input iterator from an indices column. */ struct input_indexalator_fn { - template ()>* = nullptr> + template ())> input_indexalator operator()(column_view const& indices) { return input_indexalator(indices.data(), indices.type()); } - template ()>* = nullptr> + template ())> input_indexalator operator()(Args&&... args) { CUDF_FAIL("indices must be an index type"); @@ -110,16 +235,14 @@ struct indexalator_factory { * @brief Use this class to create an indexalator to a scalar index. */ struct input_indexalator_scalar_fn { - template ()>* = nullptr> + template ())> input_indexalator operator()(scalar const& index) { // note: using static_cast const&>(index) creates a copy auto const scalar_impl = static_cast const*>(&index); return input_indexalator(scalar_impl->data(), index.type()); } - template ()>* = nullptr> + template ())> input_indexalator operator()(Args&&... args) { CUDF_FAIL("scalar must be an index type"); @@ -130,14 +253,12 @@ struct indexalator_factory { * @brief A type_dispatcher functor to create an output iterator from an indices column. */ struct output_indexalator_fn { - template ()>* = nullptr> + template ())> output_indexalator operator()(mutable_column_view const& indices) { return output_indexalator(indices.data(), indices.type()); } - template ()>* = nullptr> + template ())> output_indexalator operator()(Args&&... args) { CUDF_FAIL("indices must be an index type"); diff --git a/cpp/include/cudf/detail/normalizing_iterator.cuh b/cpp/include/cudf/detail/normalizing_iterator.cuh index 35a695d47df..8f90afc3e57 100644 --- a/cpp/include/cudf/detail/normalizing_iterator.cuh +++ b/cpp/include/cudf/detail/normalizing_iterator.cuh @@ -33,7 +33,7 @@ namespace detail { * @tparam Integer The type the iterator normalizes to */ template -struct base_normalator { +struct alignas(16) base_normalator { static_assert(cudf::is_index_type()); using difference_type = std::ptrdiff_t; using value_type = Integer; @@ -204,7 +204,7 @@ struct base_normalator { private: struct integer_sizeof_fn { - template ()>* = nullptr> + template ())> CUDF_HOST_DEVICE constexpr std::size_t operator()() const { #ifndef __CUDA_ARCH__ @@ -213,7 +213,7 @@ struct base_normalator { CUDF_UNREACHABLE("only integral types are supported"); #endif } - template ()>* = nullptr> + template ())> CUDF_HOST_DEVICE constexpr std::size_t operator()() const noexcept { return sizeof(T); @@ -229,160 +229,16 @@ struct base_normalator { width_ = static_cast(type_dispatcher(dtype, integer_sizeof_fn{})); } - int32_t width_; /// integer type width = 1,2,4, or 8 - data_type dtype_; /// for type-dispatcher calls -}; - -/** - * @brief The integer normalizing input iterator - * - * This is an iterator that can be used for index types (integers) without - * requiring a type-specific instance. It can be used for any iterator - * interface for reading an array of integer values of type - * int8, int16, int32, int64, uint8, uint16, uint32, or uint64. - * Reading specific elements always return a type of `Integer` - * - * @tparam Integer Type returned by all read functions - */ -template -struct input_normalator : base_normalator, Integer> { - friend struct base_normalator, Integer>; // for CRTP - - using reference = Integer const; // this keeps STL and thrust happy - - input_normalator() = default; - input_normalator(input_normalator const&) = default; - input_normalator(input_normalator&&) = default; - input_normalator& operator=(input_normalator const&) = default; - input_normalator& operator=(input_normalator&&) = default; - - /** - * @brief Indirection operator returns the value at the current iterator position - */ - __device__ inline Integer operator*() const { return operator[](0); } - - /** - * @brief Dispatch functor for resolving a Integer value from any integer type - */ - struct normalize_type { - template ()>* = nullptr> - __device__ Integer operator()(void const* tp) - { - return static_cast(*static_cast(tp)); - } - template ()>* = nullptr> - __device__ Integer operator()(void const*) - { - CUDF_UNREACHABLE("only integral types are supported"); - } - }; - /** - * @brief Array subscript operator returns a value at the input - * `idx` position as a `Integer` value. - */ - __device__ inline Integer operator[](size_type idx) const - { - void const* tp = p_ + (idx * this->width_); - return type_dispatcher(this->dtype_, normalize_type{}, tp); - } - - /** - * @brief Create an input index normalizing iterator. - * - * Use the indexalator_factory to create an iterator instance. - * - * @param data Pointer to an integer array in device memory. - * @param data_type Type of data in data - */ - CUDF_HOST_DEVICE input_normalator(void const* data, data_type dtype, cudf::size_type offset = 0) - : base_normalator, Integer>(dtype), p_{static_cast(data)} - { - p_ += offset * this->width_; - } - - char const* p_; /// pointer to the integer data in device memory -}; - -/** - * @brief The integer normalizing output iterator - * - * This is an iterator that can be used for index types (integers) without - * requiring a type-specific instance. It can be used for any iterator - * interface for writing an array of integer values of type - * int8, int16, int32, int64, uint8, uint16, uint32, or uint64. - * Setting specific elements always accept the `Integer` type values. - * - * @tparam Integer The type used for all write functions - */ -template -struct output_normalator : base_normalator, Integer> { - friend struct base_normalator, Integer>; // for CRTP - - using reference = output_normalator const&; // required for output iterators - - output_normalator() = default; - output_normalator(output_normalator const&) = default; - output_normalator(output_normalator&&) = default; - output_normalator& operator=(output_normalator const&) = default; - output_normalator& operator=(output_normalator&&) = default; - - /** - * @brief Indirection operator returns this iterator instance in order - * to capture the `operator=(Integer)` calls. - */ - __device__ inline output_normalator const& operator*() const { return *this; } - - /** - * @brief Array subscript operator returns an iterator instance at the specified `idx` position. - * - * This allows capturing the subsequent `operator=(Integer)` call in this class. - */ - __device__ inline output_normalator const operator[](size_type idx) const - { - output_normalator tmp{*this}; - tmp.p_ += (idx * this->width_); - return tmp; - } - - /** - * @brief Dispatch functor for setting the index value from a size_type value. - */ - struct normalize_type { - template ()>* = nullptr> - __device__ void operator()(void* tp, Integer const value) - { - (*static_cast(tp)) = static_cast(value); - } - template ()>* = nullptr> - __device__ void operator()(void*, Integer const) - { - CUDF_UNREACHABLE("only index types are supported"); - } - }; - - /** - * @brief Assign an Integer value to the current iterator position - */ - __device__ inline output_normalator const& operator=(Integer const value) const - { - void* tp = p_; - type_dispatcher(this->dtype_, normalize_type{}, tp, value); - return *this; - } - - /** - * @brief Create an output normalizing iterator - * - * @param data Pointer to an integer array in device memory. - * @param data_type Type of data in data + * @brief Constructor assigns width and type member variables for base class. */ - CUDF_HOST_DEVICE output_normalator(void* data, data_type dtype) - : base_normalator, Integer>(dtype), p_{static_cast(data)} + explicit CUDF_HOST_DEVICE base_normalator(data_type dtype, int32_t width) + : width_(width), dtype_(dtype) { } - char* p_; /// pointer to the integer data in device memory + int32_t width_; /// integer type width = 1,2,4, or 8 + data_type dtype_; /// for type-dispatcher calls }; } // namespace detail diff --git a/cpp/include/cudf/detail/offsets_iterator.cuh b/cpp/include/cudf/detail/offsets_iterator.cuh new file mode 100644 index 00000000000..3eb77b32353 --- /dev/null +++ b/cpp/include/cudf/detail/offsets_iterator.cuh @@ -0,0 +1,165 @@ +/* + * Copyright (c) 2023, 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. + */ + +#pragma once + +#include +#include + +namespace cudf { +namespace detail { + +/** + * @brief The offsets normalizing input iterator + * + * This is an iterator that can be used for offsets where the underlying + * type may be int32_t or int64_t. + * + * Use the offsetalator_factory to create an appropriate input iterator + * from an offsets column_view. + */ +struct input_offsetalator : base_normalator { + friend struct base_normalator; // for CRTP + + using reference = int64_t const; // this keeps STL and thrust happy + + input_offsetalator() = default; + input_offsetalator(input_offsetalator const&) = default; + input_offsetalator(input_offsetalator&&) = default; + input_offsetalator& operator=(input_offsetalator const&) = default; + input_offsetalator& operator=(input_offsetalator&&) = default; + + /** + * @brief Indirection operator returns the value at the current iterator position + */ + __device__ inline int64_t operator*() const { return operator[](0); } + + /** + * @brief Array subscript operator returns a value at the input + * `idx` position as a int64_t value. + */ + __device__ inline int64_t operator[](size_type idx) const + { + void const* tp = p_ + (idx * this->width_); + return this->width_ == sizeof(int32_t) ? static_cast(*static_cast(tp)) + : *static_cast(tp); + } + + /** + * @brief Create an input index normalizing iterator. + * + * Use the indexalator_factory to create an iterator instance. + * + * @param data Pointer to an integer array in device memory. + * @param dtype Type of data in data + */ + CUDF_HOST_DEVICE input_offsetalator(void const* data, data_type dtype) + : base_normalator( + dtype, dtype.id() == type_id::INT32 ? sizeof(int32_t) : sizeof(int64_t)), + p_{static_cast(data)} + { +#ifndef __CUDA_ARCH__ + CUDF_EXPECTS(dtype.id() == type_id::INT32 || dtype.id() == type_id::INT64, + "Unexpected offsets type"); +#else + cudf_assert((dtype.id() == type_id::INT32 || dtype.id() == type_id::INT64) && + "Unexpected offsets type"); +#endif + } + + protected: + char const* p_; /// pointer to the integer data in device memory +}; + +/** + * @brief The offsets normalizing output iterator + * + * This is an iterator that can be used for storing offsets values + * where the underlying type may be either int32_t or int64_t. + * + * Use the offsetalator_factory to create an appropriate output iterator + * from a mutable_column_view. + * + */ +struct output_offsetalator : base_normalator { + friend struct base_normalator; // for CRTP + + using reference = output_offsetalator const&; // required for output iterators + + output_offsetalator() = default; + output_offsetalator(output_offsetalator const&) = default; + output_offsetalator(output_offsetalator&&) = default; + output_offsetalator& operator=(output_offsetalator const&) = default; + output_offsetalator& operator=(output_offsetalator&&) = default; + + /** + * @brief Indirection operator returns this iterator instance in order + * to capture the `operator=(int64)` calls. + */ + __device__ inline output_offsetalator const& operator*() const { return *this; } + + /** + * @brief Array subscript operator returns an iterator instance at the specified `idx` position. + * + * This allows capturing the subsequent `operator=(int64)` call in this class. + */ + __device__ inline output_offsetalator const operator[](size_type idx) const + { + output_offsetalator tmp{*this}; + tmp.p_ += (idx * this->width_); + return tmp; + } + + /** + * @brief Assign an offset value to the current iterator position + */ + __device__ inline output_offsetalator const& operator=(int64_t const value) const + { + void* tp = p_; + if (this->width_ == sizeof(int32_t)) { + (*static_cast(tp)) = static_cast(value); + } else { + (*static_cast(tp)) = value; + } + return *this; + } + + /** + * @brief Create an output offsets iterator + * + * @param data Pointer to an integer array in device memory. + * @param dtype Type of data in data + */ + CUDF_HOST_DEVICE output_offsetalator(void* data, data_type dtype) + : base_normalator( + dtype, dtype.id() == type_id::INT32 ? sizeof(int32_t) : sizeof(int64_t)), + p_{static_cast(data)} + { +#ifndef __CUDA_ARCH__ + CUDF_EXPECTS(dtype.id() == type_id::INT32 || dtype.id() == type_id::INT64, + "Unexpected offsets type"); +#else + cudf_assert((dtype.id() == type_id::INT32 || dtype.id() == type_id::INT64) && + "Unexpected offsets type"); +#endif + } + + protected: + char* p_; /// pointer to the integer data in device memory +}; + +} // namespace detail +} // namespace cudf diff --git a/cpp/include/cudf/detail/offsets_iterator_factory.cuh b/cpp/include/cudf/detail/offsets_iterator_factory.cuh new file mode 100644 index 00000000000..5b4c6b825d2 --- /dev/null +++ b/cpp/include/cudf/detail/offsets_iterator_factory.cuh @@ -0,0 +1,47 @@ +/* + * Copyright (c) 2023, 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. + */ + +#pragma once + +#include +#include + +namespace cudf { +namespace detail { + +/** + * @brief Use this class to create an offsetalator instance. + */ +struct offsetalator_factory { + /** + * @brief Create an input offsetalator instance from an offsets column + */ + static input_offsetalator make_input_iterator(column_view const& offsets) + { + return input_offsetalator(offsets.head(), offsets.type()); + } + + /** + * @brief Create an output offsetalator instance from an offsets column + */ + static output_offsetalator make_output_iterator(mutable_column_view const& offsets) + { + return output_offsetalator(offsets.head(), offsets.type()); + } +}; + +} // namespace detail +} // namespace cudf diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index 196a4f2d038..5b3fdd09cd8 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -393,6 +393,7 @@ set_tests_properties( ConfigureTest( ITERATOR_TEST iterator/indexalator_test.cu + iterator/offsetalator_test.cu iterator/optional_iterator_test_chrono.cu iterator/optional_iterator_test_numeric.cu iterator/pair_iterator_test_chrono.cu diff --git a/cpp/tests/iterator/indexalator_test.cu b/cpp/tests/iterator/indexalator_test.cu index 3e8bcd5cb0d..0c10853ec02 100644 --- a/cpp/tests/iterator/indexalator_test.cu +++ b/cpp/tests/iterator/indexalator_test.cu @@ -157,40 +157,3 @@ TYPED_TEST(IndexalatorTest, output_iterator) expected = cudf::test::fixed_width_column_wrapper({0, 1, 1, 2, 3, 4, 5, 5, 7}); CUDF_TEST_EXPECT_COLUMNS_EQUAL(d_col2, expected); } - -/** - * For testing creating and using the indexalator in device code. - */ -struct device_functor_fn { - cudf::column_device_view const d_col; - __device__ cudf::size_type operator()(cudf::size_type idx) - { - auto itr = cudf::detail::input_indexalator(d_col.head(), d_col.type()); - return itr[idx] * 3; - } -}; - -TYPED_TEST(IndexalatorTest, device_indexalator) -{ - using T = TypeParam; - - auto d_col1 = - cudf::test::fixed_width_column_wrapper({0, 6, 7, 14, 23, 33, 43, 45, 63}); - auto d_col2 = - cudf::test::fixed_width_column_wrapper({0, 0, 0, 0, 0, 0, 0, 0, 0}); - auto input = cudf::column_view(d_col1); - auto output = cudf::mutable_column_view(d_col2); - auto stream = cudf::get_default_stream(); - - auto d_input = cudf::column_device_view::create(input, stream); - - thrust::transform(rmm::exec_policy(stream), - thrust::counting_iterator(0), - thrust::counting_iterator(input.size()), - output.begin(), - device_functor_fn{*d_input}); - - auto expected = - cudf::test::fixed_width_column_wrapper({0, 18, 21, 42, 69, 99, 129, 135, 189}); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(d_col2, expected); -} diff --git a/cpp/tests/iterator/offsetalator_test.cu b/cpp/tests/iterator/offsetalator_test.cu new file mode 100644 index 00000000000..e569e58f42a --- /dev/null +++ b/cpp/tests/iterator/offsetalator_test.cu @@ -0,0 +1,140 @@ +/* + * Copyright (c) 2023, 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. + */ + +#include + +#include +#include + +#include + +#include +#include +#include +#include +#include +#include +#include +#include + +using TestingTypes = cudf::test::Types; + +template +struct OffsetalatorTest : public IteratorTest {}; + +TYPED_TEST_SUITE(OffsetalatorTest, TestingTypes); + +TYPED_TEST(OffsetalatorTest, input_iterator) +{ + using T = TypeParam; + + auto host_values = cudf::test::make_type_param_vector({0, 6, 0, -14, 13, 64, -13, -20, 45}); + + auto d_col = cudf::test::fixed_width_column_wrapper(host_values.begin(), host_values.end()); + + auto expected_values = thrust::host_vector(host_values.size()); + std::transform(host_values.begin(), host_values.end(), expected_values.begin(), [](auto v) { + return static_cast(v); + }); + + auto it_dev = cudf::detail::offsetalator_factory::make_input_iterator(d_col); + this->iterator_test_thrust(expected_values, it_dev, host_values.size()); +} + +TYPED_TEST(OffsetalatorTest, output_iterator) +{ + using T = TypeParam; + + auto d_col1 = cudf::test::fixed_width_column_wrapper({0, 6, 7, 14, 23, 33, 43, 45, 63}); + auto d_col2 = cudf::test::fixed_width_column_wrapper({0, 0, 0, 0, 0, 0, 0, 0, 0}); + auto itr = cudf::detail::offsetalator_factory::make_output_iterator(d_col2); + auto input = cudf::column_view(d_col1); + auto stream = cudf::get_default_stream(); + + auto map = cudf::test::fixed_width_column_wrapper({0, 2, 4, 6, 8, 1, 3, 5, 7}); + auto d_map = cudf::column_view(map); + thrust::gather(rmm::exec_policy_nosync(stream), + d_map.begin(), + d_map.end(), + input.begin(), + itr); + auto expected = cudf::test::fixed_width_column_wrapper({0, 7, 23, 43, 63, 6, 14, 33, 45}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(d_col2, expected); + + thrust::scatter(rmm::exec_policy_nosync(stream), + input.begin(), + input.end(), + d_map.begin(), + itr); + expected = cudf::test::fixed_width_column_wrapper({0, 33, 6, 43, 7, 45, 14, 63, 23}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(d_col2, expected); + + thrust::fill(rmm::exec_policy(stream), itr, itr + input.size(), 77); + expected = cudf::test::fixed_width_column_wrapper({77, 77, 77, 77, 77, 77, 77, 77, 77}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(d_col2, expected); + + thrust::sequence(rmm::exec_policy(stream), itr, itr + input.size()); + expected = cudf::test::fixed_width_column_wrapper({0, 1, 2, 3, 4, 5, 6, 7, 8}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(d_col2, expected); + + auto offsets = + cudf::test::fixed_width_column_wrapper({0, 10, 20, 30, 40, 50, 60, 70, 80}); + auto d_offsets = cudf::column_view(offsets); + thrust::lower_bound(rmm::exec_policy(stream), + d_offsets.begin(), + d_offsets.end(), + input.begin(), + input.end(), + itr); + expected = cudf::test::fixed_width_column_wrapper({0, 1, 1, 2, 3, 4, 5, 5, 7}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(d_col2, expected); +} + +namespace { +/** + * For testing creating and using the offsetalator in device code. + */ +struct device_functor_fn { + cudf::column_device_view const d_col; + __device__ int32_t operator()(int idx) + { + auto const itr = cudf::detail::input_offsetalator(d_col.head(), d_col.type()); + return static_cast(itr[idx] * 3); + } +}; +} // namespace + +TYPED_TEST(OffsetalatorTest, device_offsetalator) +{ + using T = TypeParam; + + auto d_col1 = cudf::test::fixed_width_column_wrapper({0, 6, 7, 14, 23, 33, 43, 45, 63}); + auto d_col2 = cudf::test::fixed_width_column_wrapper({0, 0, 0, 0, 0, 0, 0, 0, 0}); + auto input = cudf::column_view(d_col1); + auto output = cudf::mutable_column_view(d_col2); + auto stream = cudf::get_default_stream(); + + auto d_input = cudf::column_device_view::create(input, stream); + + thrust::transform(rmm::exec_policy(stream), + thrust::counting_iterator(0), + thrust::counting_iterator(input.size()), + output.begin(), + device_functor_fn{*d_input}); + + auto expected = + cudf::test::fixed_width_column_wrapper({0, 18, 21, 42, 69, 99, 129, 135, 189}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(d_col2, expected); +}