Skip to content

Commit

Permalink
rework offsetalator/indexalator dispatch logic
Browse files Browse the repository at this point in the history
  • Loading branch information
@davidwendt
davidwendt committed Nov 1, 2023
1 parent 09afe63 commit 704c853
Show file tree
Hide file tree
Showing 6 changed files with 380 additions and 160 deletions.
128 changes: 126 additions & 2 deletions cpp/include/cudf/detail/indexalator.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -56,7 +56,64 @@ namespace detail {
* auto result = thrust::find(thrust::device, begin, end, size_type{12} );
* @endcode
*/
using input_indexalator = input_normalator<cudf::size_type>;
struct input_indexalator : base_normalator<input_indexalator, cudf::size_type> {
friend struct base_normalator<input_indexalator, cudf::size_type>; // 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 <typename T, std::enable_if_t<cudf::is_index_type<T>()>* = nullptr>
__device__ cudf::size_type operator()(void const* tp)
{
return static_cast<cudf::size_type>(*static_cast<T const*>(tp));
}
template <typename T, std::enable_if_t<not cudf::is_index_type<T>()>* = nullptr>
__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 data_type Type of data in data
*/
CUDF_HOST_DEVICE input_indexalator(void const* data, data_type dtype, cudf::size_type offset = 0)
: base_normalator<input_indexalator, cudf::size_type>(dtype), p_{static_cast<char const*>(data)}
{
p_ += offset * this->width_;
}

char const* p_; /// pointer to the integer data in device memory
};

/**
* @brief The index normalizing output iterator.
Expand All @@ -82,7 +139,74 @@ using input_indexalator = input_normalator<cudf::size_type>;
* thrust::less<Element>());
* @endcode
*/
using output_indexalator = output_normalator<cudf::size_type>;
struct output_indexalator : base_normalator<output_indexalator, cudf::size_type> {
friend struct base_normalator<output_indexalator, cudf::size_type>; // 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 output_indexalator 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_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 <typename T, std::enable_if_t<cudf::is_index_type<T>()>* = nullptr>
__device__ void operator()(void* tp, cudf::size_type const value)
{
(*static_cast<T*>(tp)) = static_cast<T>(value);
}
template <typename T, std::enable_if_t<not cudf::is_index_type<T>()>* = nullptr>
__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 output_indexalator const& 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 data_type Type of data in data
*/
CUDF_HOST_DEVICE output_indexalator(void* data, data_type dtype)
: base_normalator<output_indexalator, cudf::size_type>(dtype), p_{static_cast<char*>(data)}
{
}

char* p_; /// pointer to the integer data in device memory
};

/**
* @brief Use this class to create an indexalator instance.
Expand Down
156 changes: 2 additions & 154 deletions cpp/include/cudf/detail/normalizing_iterator.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -204,7 +204,7 @@ struct base_normalator {

private:
struct integer_sizeof_fn {
template <typename T, std::enable_if_t<not cudf::is_index_type<T>()>* = nullptr>
template <typename T, std::enable_if_t<not cudf::is_fixed_width<T>()>* = nullptr>
CUDF_HOST_DEVICE constexpr std::size_t operator()() const
{
#ifndef __CUDA_ARCH__
Expand All @@ -213,7 +213,7 @@ struct base_normalator {
CUDF_UNREACHABLE("only integral types are supported");
#endif
}
template <typename T, std::enable_if_t<cudf::is_index_type<T>()>* = nullptr>
template <typename T, std::enable_if_t<cudf::is_fixed_width<T>()>* = nullptr>
CUDF_HOST_DEVICE constexpr std::size_t operator()() const noexcept
{
return sizeof(T);
Expand All @@ -233,157 +233,5 @@ struct base_normalator {
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 <typename Integer>
struct input_normalator : base_normalator<input_normalator<Integer>, Integer> {
friend struct base_normalator<input_normalator<Integer>, 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 <typename T, std::enable_if_t<cudf::is_index_type<T>()>* = nullptr>
__device__ Integer operator()(void const* tp)
{
return static_cast<Integer>(*static_cast<T const*>(tp));
}
template <typename T, std::enable_if_t<not cudf::is_index_type<T>()>* = 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<input_normalator<Integer>, Integer>(dtype), p_{static_cast<char const*>(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 <typename Integer>
struct output_normalator : base_normalator<output_normalator<Integer>, Integer> {
friend struct base_normalator<output_normalator<Integer>, 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 <typename T, std::enable_if_t<cudf::is_index_type<T>()>* = nullptr>
__device__ void operator()(void* tp, Integer const value)
{
(*static_cast<T*>(tp)) = static_cast<T>(value);
}
template <typename T, std::enable_if_t<not cudf::is_index_type<T>()>* = 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
*/
CUDF_HOST_DEVICE output_normalator(void* data, data_type dtype)
: base_normalator<output_normalator<Integer>, Integer>(dtype), p_{static_cast<char*>(data)}
{
}

char* p_; /// pointer to the integer data in device memory
};

} // namespace detail
} // namespace cudf
Loading

0 comments on commit 704c853

Please sign in to comment.