From 964d63d3feb17f61b3be7d06588b43fa240be2a3 Mon Sep 17 00:00:00 2001 From: Charles Hastings Date: Fri, 15 Nov 2024 09:56:34 -0800 Subject: [PATCH] Add type dispatching, generic (type erased) device vector and edge properties object that utilizes them --- cpp/CMakeLists.txt | 3 + cpp/include/cugraph/device_vector.hpp | 113 +++++++ cpp/include/cugraph/edge_properties.hpp | 172 ++++++++++ .../utilities/cugraph_data_type_id.hpp | 295 ++++++++++++++++++ cpp/src/structure/edge_properties_impl.hpp | 159 ++++++++++ cpp/src/structure/edge_properties_v32.cpp | 145 +++++++++ cpp/src/structure/edge_properties_v64.cpp | 51 +++ cpp/src/utilities/cugraph_data_type_id.cpp | 42 +++ 8 files changed, 980 insertions(+) create mode 100644 cpp/include/cugraph/device_vector.hpp create mode 100644 cpp/include/cugraph/edge_properties.hpp create mode 100644 cpp/include/cugraph/utilities/cugraph_data_type_id.hpp create mode 100644 cpp/src/structure/edge_properties_impl.hpp create mode 100644 cpp/src/structure/edge_properties_v32.cpp create mode 100644 cpp/src/structure/edge_properties_v64.cpp create mode 100644 cpp/src/utilities/cugraph_data_type_id.cpp diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 2cea2e504ab..a7908383170 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -187,6 +187,9 @@ set(CUGRAPH_SOURCES src/community/detail/maximal_independent_moves_mg_v32_e32.cu src/detail/utility_wrappers_32.cu src/detail/utility_wrappers_64.cu + src/structure/edge_properties_v32.cpp + src/structure/edge_properties_v64.cpp + src/utilities/cugraph_data_type_id.cpp src/structure/graph_view_mg_v64_e64.cu src/structure/graph_view_mg_v32_e32.cu src/structure/remove_self_loops_sg_v32_e32.cu diff --git a/cpp/include/cugraph/device_vector.hpp b/cpp/include/cugraph/device_vector.hpp new file mode 100644 index 00000000000..44fb61f451f --- /dev/null +++ b/cpp/include/cugraph/device_vector.hpp @@ -0,0 +1,113 @@ +/* + * Copyright (c) 2024, 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 "cugraph_c/types.h" + +#include +#include +#include + +#include + +#include +#include +#include + +namespace cugraph { + +namespace detail { + +inline rmm::device_buffer allocate_buffer(raft::handle_t const& handle, + cugraph_data_type_id_t t, + size_t size) +{ + return (t == BOOL) ? rmm::device_buffer(cugraph::packed_bool_size(size) * data_type_size(t), + handle.get_stream()) + : rmm::device_buffer(size * data_type_size(t), handle.get_stream()); +} + +} // namespace detail + +/** + * Class wrapping a type-erased device vector. + * */ +class device_vector_t { + public: + /** + * Constructor creating a device vector + * + * @param data_type The data type of the vector + * @param size The number of elements in the vector + */ + device_vector_t(raft::handle_t const& handle, cugraph_data_type_id_t data_type, size_t size) + : data_(detail::allocate_buffer(handle, data_type, size)), type_(data_type), size_(size) + { + } + + /** + * Constructor initializing device vector from an rmm device uvector + * + * @tparam T type for the array/vector + * @param vector Vector to + */ + template + device_vector_t(rmm::device_uvector&& vector) : type_(type_to_id()), size_(vector.size()) + { + data_ = vector.release(); + } + + template + T* begin() + { + return reinterpret_cast(data_.data()); + } + + template + T const* begin() const + { + return reinterpret_cast(data_.data()); + } + + template + T* end() + { + return reinterpret_cast(data_.data()) + size_; + } + + template + T const* end() const + { + return reinterpret_cast(data_.data()) + size_; + } + + cugraph_data_type_id_t type() const { return type_; } + size_t size() const { return size_; } + + void clear(rmm::cuda_stream_view stream_view) + { + data_.resize(0, stream_view); + data_.shrink_to_fit(stream_view); + } + + private: + rmm::device_buffer data_{}; + cugraph_data_type_id_t type_{NTYPES}; + size_t size_{0}; +}; + +} // namespace cugraph diff --git a/cpp/include/cugraph/edge_properties.hpp b/cpp/include/cugraph/edge_properties.hpp new file mode 100644 index 00000000000..f2b8a6997b9 --- /dev/null +++ b/cpp/include/cugraph/edge_properties.hpp @@ -0,0 +1,172 @@ +/* + * Copyright (c) 2024, 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 "cugraph/device_vector.hpp" +#include "cugraph_c/types.h" + +#include +#include +#include + +#include + +namespace cugraph { + +namespace detail { + +class edge_property_impl_t { + public: + edge_property_impl_t(raft::handle_t const& handle, + cugraph_data_type_id_t data_type, + std::vector const& edge_counts); + + edge_property_impl_t(cugraph_data_type_id_t data_type, + std::vector&& vectors); + + template + edge_property_impl_t(std::vector>&& vectors, + std::vector const& edge_counts); + + template + edge_property_view_t view(std::vector const& edge_counts) const; + + template + edge_property_view_t mutable_view(std::vector const& edge_counts); + + cugraph_data_type_id_t data_type() const { return dtype_; } + + private: + cugraph_data_type_id_t dtype_{NTYPES}; + std::vector vectors_{}; +}; + +} // namespace detail + +/** + * Class for containing a collection of edge properties. Semantic interpretation of + * the properties is for the caller to interpret. + * + * Edge properties are labeled as in a vector, from 0 to n-1. It is up to the caller to + * handle proper usage of the properties. + */ +class edge_properties_t { + public: + /** + * Constructor initializing properties from a graph view + * + * @tparam GraphViewType type for the graph view + * @param graph_view Graph view object + */ + template + edge_properties_t(GraphViewType const& graph_view); + + /** + * Adds an empty property of the specified type. + * + * Fails if the property for @p idx is already defined + * + * @param handle Handle for resources + * @param idx Index of which property to add + * @param data_type Data type of the property to add + */ + void add_property(raft::handle_t const& handle, size_t idx, cugraph_data_type_id_t data_type); + + /** + * Adds a property of the specified type initialized with the provided values + * + * Fails if the property for @p idx is already defined + * + * @tparam value_type Type of the property + * @param idx Index of which property to add + * @param buffers Initial value of the property + */ + template + void add_property(size_t idx, std::vector>&& buffers); + + /** + * Adds a property initialized with the provided values + * + * Fails if the property for @p idx is already defined + * + * @param idx Index of which property to add + * @param vectors Type erased device vectors for the initial value + */ + void add_property(size_t idx, std::vector&& vectors); + /** + * Clears the specified property, releasing any allocated memory + * + * @param idx Index of which property to clear + */ + void clear_property(size_t idx); + + /** + * clears all properties, releasing any allocate memory + */ + void clear_all_properties(); + + /** + * Returns true if property @p idx is defined + */ + bool is_defined(size_t idx); + + /** + * Returns data type of property @p idx + */ + cugraph_data_type_id_t data_type(size_t idx); + + /** + * Returns a read-only edge property view of the property using the provided types + * + * Throws exception if idx does not refer to a defined property. + * Throws exception if value_t does not match the property type + * + * @tparam edge_t Typename for the edge + * @tparam value_t Typename for the property + * @param idx Index of the property + * + * @returns a read-only view for accessing the property + */ + template + edge_property_view_t view(size_t idx) const; + + /** + * Returns a read-write edge property view of the property using the provided types + * + * Throws exception if idx does not refer to a defined property. + * Throws exception if value_t does not match the property type + * + * @tparam edge_t Typename for the edge + * @tparam value_t Typename for the property + * @param idx Index of the property + * + * @returns a read-write view for accessing the property + */ + template + edge_property_view_t mutable_view(size_t idx); + + /** + * Return list of defined properties + */ + std::vector defined() const; + + private: + std::vector> properties_{}; + std::vector edge_counts_{}; +}; + +} // namespace cugraph diff --git a/cpp/include/cugraph/utilities/cugraph_data_type_id.hpp b/cpp/include/cugraph/utilities/cugraph_data_type_id.hpp new file mode 100644 index 00000000000..ee87e8f00da --- /dev/null +++ b/cpp/include/cugraph/utilities/cugraph_data_type_id.hpp @@ -0,0 +1,295 @@ +/* + * Copyright (c) 2024, 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 + +#include +#include + +namespace cugraph { + +// Type conversion logic adapted from CUDF + +/** + * @brief Maps a C++ type to its corresponding `cugraph_data_type_id_t` + * + * When explicitly passed a template argument of a given type, returns the + * appropriate `cugraph_data_type_id_t` enum for the specified C++ type. + * + * For example: + * + * ``` + * return cugraph::type_to_id(); // Returns INT32 + * ``` + * + * @tparam T The type to map to a `cugraph_data_type_id_t` + * @return The `cugraph_data_type_id_t` corresponding to the specified type + */ +template +inline constexpr cugraph_data_type_id_t type_to_id() +{ + return cugraph_data_type_id_t::NTYPES; +}; + +/** + * @brief Maps a `cugraph_data_type_id_t` types to its corresponding C++ type name string + * + */ +struct type_to_name_impl { + /** + * @brief Maps a `cugraph_data_type_id_t` types to its corresponding C++ type name string + * + * @return The C++ type name as string + */ + template + inline std::string operator()() + { + return "void"; + } +}; + +template +struct id_to_type_impl { + using type = void; +}; + +/** + * @brief Maps a `cugraph_data_type_id_t` to its corresponding concrete C++ type + * + * Example: + * ``` + * static_assert(std::is_same); + * ``` + * @tparam t The `cugraph_data_type_id_t` to map + */ +template +using id_to_type = typename id_to_type_impl::type; + +/** + * @brief Macro used to define a mapping between a concrete C++ type and a + *`cugraph_data_type_id_t` enum. + + * @param Type The concrete C++ type + * @param Id The `cugraph_data_type_id_t` enum + */ +#ifndef CUGRAPH_TYPE_MAPPING +#define CUGRAPH_TYPE_MAPPING(Type, Id) \ + template <> \ + constexpr inline cugraph_data_type_id_t type_to_id() \ + { \ + return Id; \ + } \ + template <> \ + inline std::string type_to_name_impl::operator()() \ + { \ + return #Type; \ + } \ + template <> \ + struct id_to_type_impl { \ + using type = Type; \ + }; +#endif + +// Defines all of the mappings between C++ types and their corresponding `cugraph_data_type_id_t` +// values. +CUGRAPH_TYPE_MAPPING(int8_t, cugraph_data_type_id_t::INT8) +CUGRAPH_TYPE_MAPPING(int16_t, cugraph_data_type_id_t::INT16) +CUGRAPH_TYPE_MAPPING(int32_t, cugraph_data_type_id_t::INT32) +CUGRAPH_TYPE_MAPPING(int64_t, cugraph_data_type_id_t::INT64) +CUGRAPH_TYPE_MAPPING(uint8_t, cugraph_data_type_id_t::UINT8) +CUGRAPH_TYPE_MAPPING(uint16_t, cugraph_data_type_id_t::UINT16) +CUGRAPH_TYPE_MAPPING(uint32_t, cugraph_data_type_id_t::UINT32) +CUGRAPH_TYPE_MAPPING(uint64_t, cugraph_data_type_id_t::UINT64) +CUGRAPH_TYPE_MAPPING(float, cugraph_data_type_id_t::FLOAT32) +CUGRAPH_TYPE_MAPPING(double, cugraph_data_type_id_t::FLOAT64) + +// These are duplicative of uint8_t and uint64_t +// CUGRAPH_TYPE_MAPPING(bool, cugraph_data_type_id_t::BOOL) +// CUGRAPH_TYPE_MAPPING(size_t, cugraph_data_type_id_t::SIZE_T) + +/** + * @brief Invokes an `operator()` template with the type instantiation based on + * the specified `cudf::data_type`'s `id()`. + * + * Example usage with a functor that returns the size of the dispatched type: + * + * @code + * struct size_of_functor{ + * template + * int operator()(){ + * return sizeof(T); + * } + * }; + * cugraph_data_type_id_t t{INT32}; + * cugraph::type_dispatcher(t, size_of_functor{}); // returns 4 + * @endcode + * + * The `type_dispatcher` uses `cugraph::type_to_id` to provide a default mapping + * of `cugraph_data_type_id_t`s to dispatched C++ types. However, this mapping may be + * customized by explicitly specifying a user-defined trait struct for the + * `IdTypeMap`. For example, to always dispatch `int32_t` + * + * @code + * template struct always_int{ using type = int32_t; } + * + * // This will always invoke operator() + * cugraph::type_dispatcher(data_type, f); + * @endcode + * + * It is sometimes necessary to customize the dispatched functor's + * `operator()` for different types. This can be done in several ways. + * + * The first method is to use explicit template specialization. This is useful + * for specializing behavior for single types. For example, a functor that + * prints `int32_t` or `double` when invoked with either of those types, else it + * prints `unhandled type`: + * + * @code + * struct type_printer { + * template + * void operator()() { std::cout << "unhandled type\n"; } + * }; + * + * // Due to a bug in g++, explicit member function specializations need to be + * // defined outside of the class definition + * template <> + * void type_printer::operator()() { std::cout << "int32_t\n"; } + * + * template <> + * void type_printer::operator()() { std::cout << "double\n"; } + * @endcode + * + * A second method is to use SFINAE with `std::enable_if_t`. This is useful for + * specializing for a set of types that share some property. For example, a + * functor that prints `integral` or `floating point` for integral or floating + * point types: + * + * @code + * struct integral_or_floating_point { + * template and + * not std::is_floating_point_v >* = nullptr> + * void operator()() { + * std::cout << "neither integral nor floating point\n "; } + * + * template >* = nullptr> + * void operator()() { std::cout << "integral\n"; } + * + * template >* = nullptr> + * void operator()() { std::cout << "floating point\n"; } + * }; + * @endcode + * + * For more info on SFINAE and `std::enable_if`, see + * https://eli.thegreenplace.net/2014/sfinae-and-enable_if/ + * + * The return type for all template instantiations of the functor's "operator()" + * lambda must be the same, else there will be a compiler error as you would be + * trying to return different types from the same function. + * + * @tparam id_to_type_impl Maps a `cugraph_data_type_id_t` its dispatched C++ type + * @tparam Functor The callable object's type + * @tparam Ts Variadic parameter pack type + * @param dtype The `cugraph_data_type_id_t` whose `id()` determines which template + * instantiation is invoked + * @param f The callable whose `operator()` template is invoked + * @param args Parameter pack of arguments forwarded to the `operator()` + * invocation + * @return Whatever is returned by the callable's `operator()` + */ +// This pragma disables a compiler warning that complains about the valid usage +// of calling a __host__ functor from this function which is __host__ __device__ +#ifdef __CUDACC__ +#pragma nv_exec_check_disable +#endif +template