diff --git a/CONTRIBUTING.md b/CONTRIBUTING.md index 5f4d7877582d..b79371962967 100644 --- a/CONTRIBUTING.md +++ b/CONTRIBUTING.md @@ -32,7 +32,7 @@ conda install cudf -c rapidsai-nightly -c conda-forge ``` 3. Build and view the docs locally following the instructions in the [Building -documentation docs](https://docs.rapids.ai/api/cudf/stable/developer_guide/documentation/#building-documentation) +documentation docs](https://docs.nvidia.com/cudf/latest/cudf/developer_guide/documentation/#building-documentation) 4. Follow steps 7-10 in the section [Your first issue](#your-first-issue) ## Code contributions @@ -327,7 +327,7 @@ This will bring up an interactive prompt to select which spelling fixes to apply The [C++ Developer Guide](cpp/doxygen/developer_guide/DEVELOPER_GUIDE.md) includes details on contributing to libcudf C++ code. -The [Python Developer Guide](https://docs.rapids.ai/api/cudf/stable/cudf/developer_guide/) includes details on contributing to cuDF Python code. +The [Python Developer Guide](https://docs.nvidia.com/cudf/latest/developer_guide/) includes details on contributing to cuDF Python code. ## Attribution diff --git a/README.md b/README.md index 60ef0be89fa9..ca455fc74079 100644 --- a/README.md +++ b/README.md @@ -8,14 +8,14 @@ Accelerated Data Science suite of libraries. cuDF is composed of multiple libraries including: -* [libcudf](https://docs.rapids.ai/api/libcudf/stable/): A CUDA C++ library with [Apache Arrow](https://arrow.apache.org/) compliant +* [libcudf](https://docs.nvidia.com/cudf/latest/libcudf/): A CUDA C++ library with [Apache Arrow](https://arrow.apache.org/) compliant data structures and fundamental algorithms for tabular data. -* [pylibcudf](https://docs.rapids.ai/api/cudf/stable/pylibcudf/): A Python library providing [Cython](https://cython.org/) bindings for libcudf. -* [cudf](https://docs.rapids.ai/api/cudf/stable/cudf/): A Python library providing +* [pylibcudf](https://docs.nvidia.com/cudf/latest/pylibcudf/): A Python library providing [Cython](https://cython.org/) bindings for libcudf. +* [cudf](https://docs.nvidia.com/cudf/latest/cudf/): A Python library providing - A DataFrame library mirroring the [pandas](https://pandas.pydata.org/) API - - A zero-code change accelerator, [cudf.pandas](https://docs.rapids.ai/api/cudf/stable/cudf_pandas/), for existing pandas code. -* [cudf-polars](https://docs.rapids.ai/api/cudf/stable/cudf_polars/): A Python library providing a GPU engine for [Polars](https://pola.rs/) -* [dask-cudf](https://docs.rapids.ai/api/dask-cudf/stable/): A Python library providing a GPU backend for [Dask](https://www.dask.org/) DataFrames + - A zero-code change accelerator, [cudf.pandas](https://docs.nvidia.com/cudf/latest/cudf_pandas/), for existing pandas code. +* [cudf-polars](https://docs.nvidia.com/cudf/latest/cudf_polars/): A Python library providing a GPU engine for [Polars](https://pola.rs/) +* [dask-cudf](https://docs.nvidia.com/dask-cudf/latest/): A Python library providing a GPU backend for [Dask](https://www.dask.org/) DataFrames Notable projects that use cuDF include: diff --git a/cpp/doxygen/developer_guide/BENCHMARKING.md b/cpp/doxygen/developer_guide/BENCHMARKING.md index df1d5536cfd7..9c155c258a21 100644 --- a/cpp/doxygen/developer_guide/BENCHMARKING.md +++ b/cpp/doxygen/developer_guide/BENCHMARKING.md @@ -37,7 +37,7 @@ performance in repeated iterations. ## Data generation -For generating benchmark input data, helper functions are available at [cpp/benchmarks/common/generate_input.hpp](/cpp/benchmarks/common/generate_input.hpp). The input data generation happens on device, in contrast to any `column_wrapper` where data generation happens on the host. +For generating benchmark input data, helper functions are available at [cpp/benchmarks/common/generate_input.hpp](https://github.com/rapidsai/cudf/blob/main/cpp/benchmarks/common/generate_input.hpp). The input data generation happens on device, in contrast to any `column_wrapper` where data generation happens on the host. * `create_sequence_table` can generate sequence columns starting with value 0 in first row and increasing by 1 in subsequent rows. * `create_random_column` can generate a column filled with random data. The random data parameters are configurable. * `create_random_table` can generate a table of columns filled with random data. The random data parameters are configurable. diff --git a/cpp/doxygen/developer_guide/DEVELOPER_GUIDE.md b/cpp/doxygen/developer_guide/DEVELOPER_GUIDE.md index 84a0e37fd5c3..14b93883cf11 100644 --- a/cpp/doxygen/developer_guide/DEVELOPER_GUIDE.md +++ b/cpp/doxygen/developer_guide/DEVELOPER_GUIDE.md @@ -26,7 +26,7 @@ A column is an array of data of a single type. Along with Tables, columns are th structures used in libcudf. Most libcudf algorithms operate on columns. Columns may have a validity mask representing whether each element is valid or null (invalid). Columns of nested types are supported, meaning that a column may have child columns. A column is the C++ equivalent to a cuDF -Python [Series](https://docs.rapids.ai/api/cudf/stable/user_guide/api_docs/api/cudf.series/). +Python [Series](https://docs.nvidia.com/cudf/latest/cudf/api_docs/series/). ### Element @@ -41,7 +41,7 @@ A type representing a single element of a data type. A table is a collection of columns that all have the same number of elements (rows). A table may also have zero columns while still carrying a row count, mirroring an `(N, 0)` DataFrame. A table is the C++ equivalent to a cuDF Python -[DataFrame](https://docs.rapids.ai/api/cudf/stable/user_guide/api_docs/api/cudf.dataframe/). +[DataFrame](https://docs.nvidia.com/cudf/latest/cudf/api_docs/dataframe/). ### View @@ -151,7 +151,7 @@ recommend watching Sean Parent's [C++ Seasoning talk](https://www.youtube.com/wa and we try to follow his rules: "No raw loops. No raw pointers. No raw synchronization primitives." * Prefer algorithms from STL and Thrust to raw loops. - * Prefer libcudf and RMM [owning data structures and views](#libcudf-data-structures) to raw + * Prefer libcudf and RMM @ref libcudf-data-structures "owning data structures and views" to raw pointers and raw memory allocation. * libcudf doesn't have a lot of CPU-thread concurrency, but there is some. And currently libcudf does use raw synchronization primitives. So we should revisit Parent's third rule and improve @@ -220,7 +220,7 @@ The following guidelines apply to organizing `#include` lines. yourself doing this, start a discussion about moving (parts of) the included internal header to a public header. -# libcudf Data Structures +# libcudf Data Structures {#libcudf-data-structures} Application data in libcudf is contained in Columns and Tables, but there are a variety of other data structures you will use when developing libcudf code. @@ -246,7 +246,7 @@ libcudf allocates all device memory via RMM memory resources (MR) or CUDA MRs. E can be passed to libcudf functions via `rmm::device_async_resource_ref` parameters. See the [RMM documentation](https://github.com/rapidsai/rmm/blob/main/README.md) for details. -### Current Device Memory Resource +### Current Device Memory Resource {#rmmdevice_memory_resource} RMM provides a "default" memory resource for each device and functions to access and set it. libcudf provides wrappers for these functions in `cpp/include/cudf/utilities/memory_resource.hpp`. @@ -347,7 +347,7 @@ An *immutable*, non-owning view of a table. A *mutable*, non-owning view of a table. -## cudf::size_type +## cudf::size_type {#cudfsize_type} The `cudf::size_type` is the type used for the number of elements in a column, indices to address specific elements, segments for subsets of column elements, etc. It is equivalent to a signed, @@ -486,7 +486,7 @@ examples. **Things that libcudf should not validate**: - Integer overflow -- Ensuring that outputs will not exceed the [`size_type`](#cudfsize_type) row count limit for a +- Ensuring that outputs will not exceed the @ref cudfsize_type "size_type" row count limit for a given set of inputs This policy describes libcudf's default behavior. Some APIs offer opt-in overflow-aware variants for @@ -551,7 +551,7 @@ inputs. The types of those exceptions (e.g. `cudf::logic_error`) are part of the However, the explanatory string returned by the `what` method of those exceptions is not part of the API and is subject to change. Calling code should not rely on the contents of libcudf error messages to determine the nature of the error. For information on the types of exceptions that -libcudf throws under different circumstances, see the [section on error handling](#errors). +libcudf throws under different circumstances, see the @ref errors "section on error handling". # libcudf API and Implementation @@ -565,8 +565,8 @@ that allocate device memory or execute a kernel should accept an also accepts a memory resource parameter, the stream parameter should be placed just *before* the memory resource. This API should then forward the call to a corresponding `detail` API with an identical signature, except that the -`detail` API should not have a default parameter for the stream ([detail APIs -should always avoid default parameters](#default-parameters)). The +`detail` API should not have a default parameter for the stream (@ref default-parameters +"detail APIs should always avoid default parameters"). The implementation should be wholly contained in the `detail` API definition and use only asynchronous versions of CUDA APIs with the stream parameter. @@ -635,7 +635,7 @@ where synchronization overhead is proportionally larger. Benchmarks have shown s for many operations on small inputs. This policy aligns with libcudf's existing stream semantics: libcudf APIs called on the host -[do not guarantee that the stream is synchronized before returning](#async-apis). +@ref async-apis "do not guarantee that the stream is synchronized before returning". Callers must explicitly synchronize if they need to access results on the host. Notes on `nosync`: @@ -674,7 +674,7 @@ cudf::detail::copy_if( ## Memory Allocation -Device [memory resources](#rmmdevice_memory_resource) are used in libcudf to abstract and control +Device @ref rmmdevice_memory_resource "memory resources" are used in libcudf to abstract and control how device memory is allocated. ### Output Memory @@ -878,7 +878,7 @@ pinned host memory may be device-accessible despite residing on the host. `cudaM CUDA to infer the valid copy direction from the pointers rather than rejecting such copies based on an explicit policy. -## Default Parameters +## Default Parameters {#default-parameters} While public libcudf APIs are free to include default function parameters, detail functions should not. Default memory resource parameters make it easy for developers to accidentally allocate memory @@ -1037,7 +1037,7 @@ This iterator replaces the null/validity value for each element with a specified This iterator returns the validity of the underlying element (`true` or `false`). Created using `cudf::detail::make_validity_iterator`. -### Index-normalizing iterators +### Index-normalizing iterators {#index-normalizing-iterators} The proliferation of data types supported by libcudf can result in long compile times. One area where compile time was a problem is in types used to store indices, which can be any integer type. @@ -1045,7 +1045,7 @@ The "indexalator", or index-normalizing iterator (`include/cudf/detail/indexalat 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 returns a -[`cudf::size_type`](#cudfsize_type) integer. +@ref cudfsize_type "cudf::size_type" integer. Use the `indexalator_factory` to create an appropriate input iterator from a column_view. Example input iterator usage: @@ -1069,12 +1069,12 @@ thrust::lower_bound(rmm::exec_policy_nosync(stream), cuda::std::less()); ``` -### Offset-normalizing iterators +### Offset-normalizing iterators {#offset-normalizing-iterators} -Like the [indexalator](#index-normalizing-iterators), +Like the @ref index-normalizing-iterators "indexalator", the "offsetalator", or offset-normalizing iterator (`include/cudf/detail/offsetalator.cuh`), can be used for offset column types (`INT32` or `INT64` only) without requiring a type-specific instance. -This is helpful when reading or building [strings columns](#strings-columns). +This is helpful when reading or building @ref strings-columns "strings columns". The normalized type is `int64` which means an `input_offsetsalator` will return `int64` type values for both `INT32` and `INT64` offsets columns. Likewise, an `output_offselator` can accept `int64` type values to store into either an @@ -1123,7 +1123,7 @@ namespace. The public function is expected to contain a call to `CUDF_FUNC_RANGE()` followed by a call to a `detail` function with same name and parameters as the public function. -See the [Streams](#streams) section for an example of this pattern. +See the @ref streams "Streams" section for an example of this pattern. ### Internal @@ -1317,7 +1317,7 @@ gradual support for more data types to make this easier. Typically we start with types such as numeric types and timestamps/durations, adding support for nested types later. Enabling an algorithm differently for different types uses either template specialization or SFINAE, -as discussed in [Specializing Type-Dispatched Code Paths](#specializing-type-dispatched-code-paths). +as discussed in @ref specializing-type-dispatched-code-paths "Specializing Type-Dispatched Code Paths". ## Comparing Data Types @@ -1388,7 +1388,7 @@ dispatched, so a second-level type dispatch results in quadratic growth in compi object code size. As a large library with many types and functions, we are constantly working to reduce compilation time and code size. -## Specializing Type-Dispatched Code Paths +## Specializing Type-Dispatched Code Paths {#specializing-type-dispatched-code-paths} It is often necessary to customize the dispatched `operator()` for different types. This can be done in several ways. @@ -1466,7 +1466,7 @@ For list columns, the parent column's type is `LIST` and contains no data, but i the number of lists in the column, and its null mask represents the validity of each list element. The parent has two children. -1. A non-nullable column of [`size_type`](#cudfsize_type) elements that indicates the offset to the +1. A non-nullable column of @ref cudfsize_type "size_type" elements that indicates the offset to the beginning of each list in a dense column of elements. 2. A column containing the actual data and optional null mask for all elements of all the lists packed together. @@ -1507,7 +1507,7 @@ lists_column = { {{{1, 2}, {3, 4}}, NULL}, {{{10, 20}, {30, 40}}, {{50, 60, 70}, This is related to [Arrow's "Variable-Size List" memory layout](https://arrow.apache.org/docs/format/Columnar.html?highlight=nested%20types#physical-memory-layout). -## Strings columns +## Strings columns {#strings-columns} Strings are represented as a column with a data device buffer and a child offsets column. The parent column's type is `STRING` and its data holds all the characters across all the strings packed together @@ -1523,7 +1523,7 @@ The following image shows an example of this compound column representation of s ![strings](strings.png) The type of the offsets column is either `INT32` or `INT64` depending on the number of bytes in the data buffer. -See [`cudf::strings_view`](#cudfstrings_column_view-and-cudfstring_view) for more information on processing individual string rows. +See @ref cudfstrings_column_view-and-cudfstring_view "cudf::strings_view" for more information on processing individual string rows. ## Structs columns @@ -1632,7 +1632,7 @@ with the corresponding strings from either `destination` or `scatter_values`. libcudf provides view types for nested column types as well as for the data elements within them. -### cudf::strings_column_view and cudf::string_view +### cudf::strings_column_view and cudf::string_view {#cudfstrings_column_view-and-cudfstring_view} A `cudf::strings_column_view` wraps a strings column and contains a parent `cudf::column_view` as a view of the strings column and an offsets `cudf::column_view` @@ -1640,14 +1640,14 @@ which is a child of the parent. The parent view contains the offset, size, and validity mask for the strings column. The offsets view is non-nullable with `offset()==0` and its own size. Since the offset column type can be either `INT32` or `INT64` it is useful to use the -offset normalizing iterators [offsetalator](#offset-normalizing-iterators) to access individual offset values. +offset normalizing iterators @ref offset-normalizing-iterators "offsetalator" to access individual offset values. A `cudf::string_view` is a view of a single string and therefore is the data type of a `cudf::column` of type `STRING` just like `int32_t` is the data type for a `cudf::column` of type `INT32`. As its name implies, this is a read-only object instance that points to device memory inside the strings column. Its lifespan is the same (or less) as the column it views. -An individual strings column row and a `cudf::string_view` is limited to [`size_type`](#cudfsize_type) bytes. +An individual strings column row and a `cudf::string_view` is limited to @ref cudfsize_type "size_type" bytes. Use the `column_device_view::element` method to access an individual row element. Like any other column, do not call `element()` on a row that is null. @@ -1667,13 +1667,14 @@ instance of a class object to represent a null string. The `cudf::string_view` contains comparison operators `<,>,==,<=,>=` that can be used in many cudf functions like `sort` without string-specific code. The data for a `cudf::string_view` instance is -required to be [UTF-8](#utf-8) and all operators and methods expect this encoding. Unless documented +required to be @ref utf-8 "UTF-8" and all operators and methods expect this encoding. Unless documented otherwise, position and length parameters are specified in characters and not bytes. The class also includes a `cudf::string_view::const_iterator` which can be used to navigate through individual characters within the string. `cudf::type_dispatcher` dispatches to the `cudf::string_view` data type when invoked on a `STRING` column. +@anchor utf-8 #### UTF-8 The libcudf strings column only supports UTF-8 encoding for strings data. @@ -1721,7 +1722,7 @@ formats commonly used in data analytics, including CSV, Parquet, ORC, Avro, and Here are some tools that can help with debugging libcudf (besides printf of course): 1. `cuda-gdb`\ - Follow the instructions in the [Contributor to cuDF guide](../../../CONTRIBUTING.md#debugging-cudf) to build + Follow the instructions in the [Contributor to cuDF guide](https://github.com/rapidsai/cudf/blob/main/CONTRIBUTING.md#debugging-cudf) to build and run libcudf with debug symbols. 2. `compute-sanitizer`\ The [CUDA Compute Sanitizer](https://docs.nvidia.com/compute-sanitizer/ComputeSanitizer/index.html) @@ -1731,7 +1732,7 @@ Here are some tools that can help with debugging libcudf (besides printf of cour The `racecheck` and `initcheck` have been known to produce false positives. 3. `cudf::test::print()`\ The `print()` utility can be called within a gtest to output the data in a `cudf::column_view`. - More information is available in the [Testing Guide](TESTING.md#printing-and-accessing-column-data) + More information is available in the @ref md_doxygen_developer_guide_TESTING "Testing Guide" 4. GCC Address Sanitizer\ The GCC ASAN can also be used by adding the `-fsanitize=address` compiler flag. There is a compatibility issue with the CUDA runtime that can be worked around by setting diff --git a/cpp/doxygen/developer_guide/DOCUMENTATION.md b/cpp/doxygen/developer_guide/DOCUMENTATION.md index 14ac0fa247bb..717942542dcb 100644 --- a/cpp/doxygen/developer_guide/DOCUMENTATION.md +++ b/cpp/doxygen/developer_guide/DOCUMENTATION.md @@ -29,7 +29,7 @@ Doxygen recognizes and parses block comments and performs specialized output for There are almost 200 commands (also called tags in this document) that doxygen recognizes in comment blocks. This document provides guidance on which commands/tags to use and how to use them in the libcudf C++ source code. -The doxygen process can be customized using options in the [Doxyfile](../Doxyfile). +The doxygen process can be customized using options in the [Doxyfile](https://github.com/rapidsai/cudf/blob/main/cpp/doxygen/Doxyfile). Here are some of the custom options in the Doxyfile for libcudf. | Option | Setting | Description | @@ -53,7 +53,7 @@ Doxygen comment blocks start with `/**` and end with `*/` only, and with nothing Do not add dashes `-----` or extra asterisks `*****` to the first and last lines of a doxygen block. The block must be placed immediately before the source code line to which it refers. The block may be indented to line up vertically with the item it documents as appropriate. -See the [Example](#the-example) section below. +See the @ref the-example "Example" section below. Each line in the comment block between the `/**` and `*/` lines should start with a space followed by an asterisk. Any text on these lines, including tag declarations, should start after a single space after the asterisk. @@ -71,7 +71,7 @@ For example, there are some limitations on readability with '%' character and pi Avoid using direct HTML tags. Although doxygen supports markdown and markdown supports HTML tags, the HTML support for doxygen's markdown is also limited. -## The Example +## The Example {#the-example} The following example covers most of the doxygen block comment and tag styles for documenting C++ code in libcudf. @@ -177,7 +177,7 @@ The comment description should clearly detail how the output(s) are created from Include any performance and any boundary considerations. Also include any limits on parameter values and if any default values are declared. Don't forget to specify how nulls are handled or produced. -Also, try to include a short [example](#inline-examples) if possible. +Also, try to include a short @ref inline-examples "example" if possible. ### @brief @@ -224,12 +224,12 @@ The following tags should appear near the end of function comment block in the o | Command | Description | | ------- | ----------- | -| [\@throw](#throw) | Specify the conditions in which the function may throw an exception | -| [\@tparam](#tparam) | Description for each template parameter | -| [\@param](#param) | Description for each function parameter | -| [\@return](#return) | Short description of object or value returned | +| @ref throw "@throw" | Specify the conditions in which the function may throw an exception | +| @ref tparam "@tparam" | Description for each template parameter | +| @ref param "@param" | Description for each function parameter | +| @ref return "@return" | Short description of object or value returned | -#### \@throw +#### \@throw {#throw} Add an [\@throw](https://www.doxygen.nl/manual/commands.html#cmdthrow) comment line in the doxygen block for each exception that the function may throw. You only need to include exceptions thrown by the function itself. @@ -243,7 +243,7 @@ Include the name of the exception without backtick marks so doxygen can add refe Using \@throws is also acceptable but VS Code and other tools only do syntax highlighting on \@throw. -#### @tparam +#### @tparam {#tparam} Add a [\@tparam](https://www.doxygen.nl/manual/commands.html#cmdtparam) comment line for each template parameter declared by this function. The name of the parameter specified after the doxygen tag must match exactly to the template parameter name. @@ -256,7 +256,7 @@ The name of the parameter specified after the doxygen tag must match exactly to The definition should detail the requirements of the parameter. For example, if the template is for a functor or predicate, then describe the expected input types and output. -#### @param +#### @param {#param} Add a [\@param](https://www.doxygen.nl/manual/commands.html#cmdparam) comment line for each function parameter passed to this function. The name of the parameter specified after the doxygen tag must match the function's parameter name. @@ -271,7 +271,7 @@ Also include append `[in]`, `[out]` or `[in,out]` to the `@param` if it is not c It is also recommended to vertically aligning the 3 columns of text if possible to make it easier to read in a source code editor. Finally, the description is normally like a title and only needs a period if it is a sentence. -#### @return +#### @return {#return} Add a single [\@return](https://www.doxygen.nl/manual/commands.html#cmdreturn) comment line at the end of the comment block if the function returns an object or value. Include a brief description of what is returned. @@ -284,7 +284,7 @@ Include a brief description of what is returned. Do not include the type of the object returned with the `@return` comment. -### Inline Examples +### Inline Examples {#inline-examples} It is usually helpful to include a source code example inside your comment block when documenting a function or other declaration. Use the [\@code](https://www.doxygen.nl/manual/commands.html#cmdcode) and [\@endcode](https://www.doxygen.nl/manual/commands.html#cmdendcode) pair to include inline examples. @@ -369,12 +369,12 @@ The doxygen output includes a _Modules_ page that organizes items into groups sp These commands can group common functions across header files, source files, and even namespaces. Groups can also be nested by defining new groups within existing groups. -For libcudf, all the group hierarchy is defined in the [doxygen_groups.h](../../include/doxygen_groups.h) header file. -The [doxygen_groups.h](../../include/doxygen_groups.h) file does not need to be included in any other source file, because the definitions in this file are used only by the doxygen tool to generate groups in the _Modules_ page. +For libcudf, all the group hierarchy is defined in the [doxygen_groups.h](https://github.com/rapidsai/cudf/blob/main/cpp/include/doxygen_groups.h) header file. +The [doxygen_groups.h](https://github.com/rapidsai/cudf/blob/main/cpp/include/doxygen_groups.h) file does not need to be included in any other source file, because the definitions in this file are used only by the doxygen tool to generate groups in the _Modules_ page. Modify this file only to add or update groups. The existing groups have been carefully structured and named, so new groups should be added thoughtfully. -When creating a new API, specify its group using the [\@ingroup](https://www.doxygen.nl/manual/commands.html#cmdingroup) tag and the group reference id from the [doxygen_groups.h](../../include/doxygen_groups.h) file. +When creating a new API, specify its group using the [\@ingroup](https://www.doxygen.nl/manual/commands.html#cmdingroup) tag and the group reference id from the [doxygen_groups.h](https://github.com/rapidsai/cudf/blob/main/cpp/include/doxygen_groups.h) file. namespace CUDF_EXPORT cudf { @@ -417,7 +417,7 @@ So include the `@addtogroup` and `@{ ... @}` between the namespace declaration b Summary of groups tags | Tag/Command | Where to use | | ----------- | ------------ | -| `@defgroup` | For use only in [doxygen_groups.h](../../include/doxygen_groups.h) and should include the group's title. | +| `@defgroup` | For use only in [doxygen_groups.h](https://github.com/rapidsai/cudf/blob/main/cpp/include/doxygen_groups.h) and should include the group's title. | | `@ingroup` | Use inside individual doxygen block comments for declaration statements in a header file. | | `@addtogroup` | Use instead of `@ingroup` for multiple declarations in the same file within a namespace declaration. Do not specify a group title. | | `@{ ... @}` | Use only with `@addtogroup`. | diff --git a/cpp/doxygen/developer_guide/TESTING.md b/cpp/doxygen/developer_guide/TESTING.md index 87d4108ec6da..0463fc144f45 100644 --- a/cpp/doxygen/developer_guide/TESTING.md +++ b/cpp/doxygen/developer_guide/TESTING.md @@ -453,6 +453,7 @@ Verifies the bitwise equality of two device memory buffers. Column comparison functions in the `cudf::test::detail` namespace should **NOT** be used directly. +\anchor printing-and-accessing-column-data ### Printing and accessing column data The `` header defines various functions and overloads for printing diff --git a/docs/cudf/source/_static/RAPIDS-logo-purple.png b/docs/cudf/source/_static/RAPIDS-logo-purple.png deleted file mode 100644 index d884e01374dc..000000000000 Binary files a/docs/cudf/source/_static/RAPIDS-logo-purple.png and /dev/null differ diff --git a/docs/cudf/source/conf.py b/docs/cudf/source/conf.py index bf89f727e793..06281d428202 100644 --- a/docs/cudf/source/conf.py +++ b/docs/cudf/source/conf.py @@ -166,6 +166,9 @@ def clean_all_xml_files(path): copybutton_prompt_text = ">>> " autosummary_generate = True +toc_object_entries_show_parents = "hide" +maximum_signature_line_length = 70 + # Enable automatic generation of systematic, namespaced labels for sections myst_heading_anchors = 2 @@ -176,7 +179,7 @@ def clean_all_xml_files(path): # You can specify multiple suffix as a list of string: # # source_suffix = ['.rst', '.md'] -source_suffix = {".rst": "restructuredtext"} +source_suffix = {".rst": "restructuredtext", ".md": "myst-nb"} # The master toctree document. master_doc = "index" @@ -324,22 +327,33 @@ def clean_all_xml_files(path): ) ] +with open("../../../RAPIDS_BRANCH", "r") as f: + branch = f.read().strip() +intersphinx_version = "latest" if branch == "main" else version -# Example configuration for intersphinx: refer to the Python standard library. intersphinx_mapping = { "cupy": ("https://docs.cupy.dev/en/stable/", None), + "dask-cuda": ( + f"https://docs.nvidia.com/dask-cuda/{intersphinx_version}/", + None, + ), + "dask-cudf": ( + f"https://docs.nvidia.com/dask-cudf/{intersphinx_version}/", + None, + ), "dlpack": ("https://dmlc.github.io/dlpack/latest/", None), + "kvikio": (f"https://docs.nvidia.com/kvikio/{intersphinx_version}/", None), "nanoarrow": ("https://arrow.apache.org/nanoarrow/latest/", None), "numpy": ("https://numpy.org/doc/stable/", None), - # Temporarily disable nitpick warnings for pandas: https://github.com/pandas-dev/pandas/issues/64584 - # "pandas": ( - # "https://pandas.pydata.org/pandas-docs/stable/", - # None, - # ), + "pandas": ("https://pandas.pydata.org/pandas-docs/stable/", None), "polars": ("https://docs.pola.rs/api/python/stable/", None), "pyarrow": ("https://arrow.apache.org/docs/", None), "python": ("https://docs.python.org/3/", None), - "rmm": ("https://docs.rapids.ai/api/rmm/nightly/", None), + "rmm": (f"https://docs.nvidia.com/rmm/{intersphinx_version}/", None), + "rapidsmpf": ( + f"https://docs.nvidia.com/rapidsmpf/{intersphinx_version}/", + None, + ), "typing_extensions": ( "https://typing-extensions.readthedocs.io/en/stable/", None, @@ -456,6 +470,8 @@ def _generate_namespaces(namespaces): "type_to_scalar_type_impl", "type_to_scalar_type_impl", "detail", + # Test-only helper types are intentionally not published as API pages. + "classcudf_1_1test_1_1", # kafka objects "python_callable_type", "kafka_oauth_callback_wrapper_type", @@ -479,7 +495,7 @@ def _generate_namespaces(namespaces): _intersphinx_extra_prefixes = ("rmm", "rmm::mr", "mr") _external_intersphinx_aliases = { - # "pandas": "pd", + "pandas": "pd", "pyarrow": "pa", "numpy": "np", "cupy": "cp", @@ -653,6 +669,12 @@ def on_missing_reference(app, env, node, contnode): nitpick_ignore = [ ("py:class", "Dtype"), ("py:class", "pandas.core.indexes.frozen.FrozenList"), + # pandas does not publish these implementation types in its inventory. + ("py:class", "pandas.api.typing.FrozenList"), + ( + "py:class", + "pandas.core.arrays.arrow.extension_types.ArrowIntervalType", + ), ("py:class", "ScalarLike"), ("py:class", "StringColumn"), ("py:class", "ColumnLike"), @@ -694,11 +716,8 @@ def on_missing_reference(app, env, node, contnode): ("py:class", "SupportsCudaArrayInterface"), ("py:class", "T"), ] -# Temporarily disable nitpick warnings for pandas: https://github.com/pandas-dev/pandas/issues/64584 + nitpick_ignore_regex = [ - ("py:.*", "pandas.*"), - ("py:.*", "pd.*"), - ("ref.*", ".*pandas.*"), # External libs without configured intersphinx inventories. ("py:.*", r"rapidsmpf(\..*)?"), ("py:.*", r"kvikio(\..*)?"), diff --git a/docs/cudf/source/cudf/10min.ipynb b/docs/cudf/source/cudf/10min.ipynb index f43b657869b8..05115d0fa716 100644 --- a/docs/cudf/source/cudf/10min.ipynb +++ b/docs/cudf/source/cudf/10min.ipynb @@ -15,7 +15,7 @@ "\n", "[Dask](https://www.dask.org/) is a flexible library for parallel computing in Python that makes scaling out your workflow smooth and simple. On the CPU, Dask uses Pandas to execute operations in parallel on DataFrame partitions.\n", "\n", - "[Dask cuDF](https://github.com/NVIDIA/cudf/tree/main/python/dask_cudf) extends Dask where necessary to allow its DataFrame partitions to be processed using cuDF GPU DataFrames instead of Pandas DataFrames. For instance, when you call `dask_cudf.read_csv(...)`, your cluster's GPUs do the work of parsing the CSV file(s) by calling [`cudf.read_csv()`](https://docs.rapids.ai/api/cudf/stable/cudf/api_docs/api/cudf.read_csv/).\n", + "[Dask cuDF](https://docs.nvidia.com/dask-cudf/) extends Dask where necessary to allow its DataFrame partitions to be processed using cuDF GPU DataFrames instead of Pandas DataFrames. For instance, when you call `dask_cudf.read_csv(...)`, your cluster's GPUs do the work of parsing the CSV file(s) by calling [`cudf.read_csv()`](https://docs.nvidia.com/cudf/latest/).\n", "\n", "\n", "
\n", @@ -2570,7 +2570,7 @@ "id": "fd3fc4f3", "metadata": {}, "source": [ - "Like pandas, cuDF provides string processing methods in the `str` attribute of `Series`. Full documentation of string methods is a work in progress. Please see the [cuDF API documentation](https://docs.rapids.ai/api/cudf/stable/cudf/api_docs/series/#string-handling) for more information." + "Like pandas, cuDF provides string processing methods in the `str` attribute of `Series`. Full documentation of string methods is a work in progress. Please see the [cuDF API documentation](https://docs.nvidia.com/cudf/latest/cudf/api_docs/series/#string-handling) for more information." ] }, { @@ -2635,7 +2635,7 @@ "id": "44fe1243", "metadata": {}, "source": [ - "As well as simple manipulation, We can also match strings using [regular expressions](https://docs.rapids.ai/api/cudf/stable/user_guide/api_docs/api/cudf.core.accessors.string.StringMethods.match.html)." + "As well as simple manipulation, We can also match strings using [regular expressions](https://docs.nvidia.com/cudf/latest/cudf/api_docs/api/cudf.core.accessors.string.StringMethods.match/)." ] }, { diff --git a/docs/cudf/source/cudf/cupy-interop.ipynb b/docs/cudf/source/cudf/cupy-interop.ipynb index 4c09d23c5c7e..ca9cd228de53 100644 --- a/docs/cudf/source/cudf/cupy-interop.ipynb +++ b/docs/cudf/source/cudf/cupy-interop.ipynb @@ -1399,7 +1399,7 @@ "source": [ "From here, we could continue our workflow with a CuPy sparse matrix.\n", "\n", - "For a full list of the functionality built into these libraries, we encourage you to check out the API docs for [cuDF](https://docs.rapids.ai/api/cudf/nightly/) and [CuPy](https://docs.cupy.dev/en/stable/index.html)." + "For a full list of the functionality built into these libraries, we encourage you to check out the API docs for [cuDF](https://docs.nvidia.com/cudf/) and [CuPy](https://docs.cupy.dev/en/stable/index.html)." ] } ], diff --git a/docs/cudf/source/cudf/guide-to-udfs.ipynb b/docs/cudf/source/cudf/guide-to-udfs.ipynb index 589c76529052..9e2483c011cb 100644 --- a/docs/cudf/source/cudf/guide-to-udfs.ipynb +++ b/docs/cudf/source/cudf/guide-to-udfs.ipynb @@ -1700,7 +1700,7 @@ "\n", "For time-series data, we may need to operate on a small \\\"window\\\" of our column at a time, processing each portion independently. We could slide (\\\"roll\\\") this window over the entire column to answer questions like \\\"What is the 3-day moving average of a stock price over the past year?\"\n", "\n", - "We can apply more complex functions to rolling windows to `rolling` Series and DataFrames using `apply`. This example is adapted from cuDF's [API documentation](https://docs.rapids.ai/api/cudf/stable/user_guide/api_docs/api/cudf.dataframe.rolling/). First, we'll create an example Series and then create a `rolling` object from the Series." + "We can apply more complex functions to rolling windows to `rolling` Series and DataFrames using `apply`. This example is adapted from cuDF's [API documentation](https://docs.nvidia.com/cudf/latest/cudf/api_docs/api/cudf.DataFrame.rolling/). First, we'll create an example Series and then create a `rolling` object from the Series." ] }, { @@ -2160,7 +2160,7 @@ "- String UDFs\n", "\n", "\n", - "For more information please see the [cuDF](https://docs.rapids.ai/api/cudf/nightly/), [Numba.cuda](https://numba.readthedocs.io/en/stable/cuda/index.html), and [CuPy](https://docs.cupy.dev/en/stable/) documentation." + "For more information please see the [cuDF](https://docs.nvidia.com/cudf/), [Numba.cuda](https://numba.readthedocs.io/en/stable/cuda/index.html), and [CuPy](https://docs.cupy.dev/en/stable/) documentation." ] } ], diff --git a/docs/cudf/source/cudf/io/io.md b/docs/cudf/source/cudf/io/io.md index ff1f96c8cd49..5d550bb62944 100644 --- a/docs/cudf/source/cudf/io/io.md +++ b/docs/cudf/source/cudf/io/io.md @@ -200,6 +200,6 @@ By default, cuDF's parquet and json readers will try to read the entire file in To better support low memory systems, cuDF provides a "low-memory" reader for parquet and json files. This low memory reader processes data in chunks, leading to lower peak memory usage due to the smaller size of intermediate allocations. -To read a parquet or json file in low memory mode, there are [cuDF options](https://docs.rapids.ai/api/cudf/nightly/cudf/api_docs/options/#api-options) that must be set globally prior to calling the reader. To set those options, call: +To read a parquet or json file in low memory mode, there are {doc}`cuDF options <../api_docs/options>` that must be set globally prior to calling the reader. To set those options, call: - `cudf.set_option("io.parquet.low_memory", True)` for parquet files, or - `cudf.set_option("io.json.low_memory", True)` for json files. diff --git a/docs/cudf/source/cudf/memory-profiling.md b/docs/cudf/source/cudf/memory-profiling.md index f69965c180a6..d984caaf1875 100644 --- a/docs/cudf/source/cudf/memory-profiling.md +++ b/docs/cudf/source/cudf/memory-profiling.md @@ -6,7 +6,7 @@ Peak memory usage is a common concern in GPU programming because GPU memory is t ## Enabling Memory Profiling -First, enable memory profiling in RMM by calling {py:func}`rmm.statistics.enable_statistics()`. This adds a statistics resource adaptor to the current RMM memory resource, which enables cuDF to access memory profiling information. See the [RMM documentation](https://docs.rapids.ai/api/rmm/stable/user_guide/guide/#memory-statistics-and-profiling) for more details. +First, enable memory profiling in RMM by calling {py:func}`rmm.statistics.enable_statistics()`. This adds a statistics resource adaptor to the current RMM memory resource, which enables cuDF to access memory profiling information. See the [RMM documentation](inv:rmm:std:label:#user_guide/guide:memory-statistics-and-profiling) for more details. Second, enable memory profiling in cuDF by setting the `memory_profiling` option to `True`. Use {py:func}`cudf.set_option` or set the environment variable ``CUDF_MEMORY_PROFILING=1`` prior to the launch of the Python interpreter. diff --git a/docs/cudf/source/cudf_pandas/faq.md b/docs/cudf/source/cudf_pandas/faq.md index 8df5f76a84aa..8241ff6507a1 100644 --- a/docs/cudf/source/cudf_pandas/faq.md +++ b/docs/cudf/source/cudf_pandas/faq.md @@ -12,8 +12,7 @@ the cuDF library directly should be considered. from increased performance by using cuDF directly. - cuDF does offer some functions and methods that pandas does not. For - example, cuDF has a [`.list` - accessor](https://docs.rapids.ai/api/cudf/stable/user_guide/api_docs/series/#list-handling) + example, cuDF has a {ref}`.list accessor ` for working with list-like data. If you need access to the additional functionality in cuDF, you will need to use the cuDF package directly. @@ -139,7 +138,7 @@ Both Dask and Apache Spark support accelerated computing through configuration based interfaces. Dask allows you to [configure the dataframe backend](https://docs.dask.org/en/latest/how-to/selecting-the-collection-backend.html) to use cuDF (learn more in [this -blog](https://medium.com/rapids-ai/easy-cpu-gpu-arrays-and-dataframes-run-your-dask-code-where-youd-like-e349d92351d)) and the [RAPIDS Accelerator for Apache Spark](https://nvidia.github.io/spark-rapids/) +blog](https://medium.com/rapids-ai/easy-cpu-gpu-arrays-and-dataframes-run-your-dask-code-where-youd-like-e349d92351d)) and the [RAPIDS Accelerator for Apache Spark](https://docs.nvidia.com/spark-rapids/) provides a similar configuration-based plugin for Spark. ## How do I know if an object is a `cudf.pandas` proxy object? diff --git a/docs/cudf/source/cudf_pandas/index.rst b/docs/cudf/source/cudf_pandas/index.rst index 964dae75ebf0..f8568001d36f 100644 --- a/docs/cudf/source/cudf_pandas/index.rst +++ b/docs/cudf/source/cudf_pandas/index.rst @@ -34,8 +34,10 @@ automatically **falling back to pandas** for other operations. | Nothing changes, not even your `import` statements, when going from CPU to GPU. | Combines the full flexibility of Pandas with blazing fast performance of cuDF | +---------------------------------------------------------------------------------------------+---------------------------------------------------------------------------------------------------------------------+ -``cudf.pandas`` is now Generally Available (GA) as part of the ``cudf`` package. See `RAPIDS -Quick Start `_ to get up-and-running with ``cudf``. +``cudf.pandas`` is available as part of the ``cudf`` package. See the +`installation and deployment guide +_` +to get up-and-running with cuDF. .. toctree:: :maxdepth: 1 diff --git a/docs/cudf/source/cudf_polars/benchmarks.md b/docs/cudf/source/cudf_polars/benchmarks.md index 969be5a0b5a6..9436aa695871 100644 --- a/docs/cudf/source/cudf_polars/benchmarks.md +++ b/docs/cudf/source/cudf_polars/benchmarks.md @@ -9,7 +9,7 @@ The steps below reproduce the PDS-H benchmark results using the Polars GPU engin ### Setup Install `cudf-polars` following the -[RAPIDS installation guide](https://docs.rapids.ai/install/). For nightly wheels, install with +[NVIDIA CUDA-X installation guide](https://developer.nvidia.com/topics/ai/data-science/cuda-x-for-data-science##section-install-and-deploy-in-your-environment). For nightly wheels, install with the `ray` extra (required for multi-GPU benchmarking): ```bash diff --git a/docs/cudf/source/cudf_polars/dask_engine.md b/docs/cudf/source/cudf_polars/dask_engine.md index 1cfa4d9f7ad4..6112ba45bb57 100644 --- a/docs/cudf/source/cudf_polars/dask_engine.md +++ b/docs/cudf/source/cudf_polars/dask_engine.md @@ -194,5 +194,5 @@ created inside an `rrun` cluster. [dask-distributed]: https://distributed.dask.org/en/stable/ [dask-cli]: https://docs.dask.org/en/latest/deploying-cli.html -[dask-cuda]: https://docs.rapids.ai/api/dask-cuda/nightly/ -[dask-cuda-worker]: https://docs.rapids.ai/api/dask-cuda/nightly/quickstart/#dask-cuda-worker +[dask-cuda]: inv:dask-cuda:std:doc:#index +[dask-cuda-worker]: diff --git a/docs/cudf/source/cudf_polars/developer_docs.md b/docs/cudf/source/cudf_polars/developer_docs.md index 069ee07ce00b..12c4ea511c7c 100644 --- a/docs/cudf/source/cudf_polars/developer_docs.md +++ b/docs/cudf/source/cudf_polars/developer_docs.md @@ -2,7 +2,7 @@ You will need: -1. Rust development environment. If you use the rapids [combined +1. Rust development environment. If you use the [combined devcontainer](https://github.com/rapidsai/devcontainers/), add `"./features/src/rust": {"version": "latest", "profile": "default"},` to your preferred configuration. Or else, use @@ -625,25 +625,21 @@ another `nvtx` range (e.g. `Scan.do_evaluate`, `GroupBy.do_evaluate`, etc.). These provide a higher-level grouping over the lower-level libcudf calls (e.g. `read_chunk`, `aggregate`). -Finally, if using [rapidsmpf](https://docs.rapids.ai/api/rapidsmpf/nightly/) -for shuffling, the methods inserting and extracting partitions to shuffle are -annotated with nvtx ranges. - # Query Plans -The module `cudf_polars.experimental.explain` contains functions for dumping +The module `cudf_polars.streaming.explain` contains functions for dumping the query for a given `LazyFrame`. ## Structured Output -`cudf_polars.experimental.explain.serialize_query` can be used to output +`cudf_polars.streaming.explain.serialize_query` can be used to output the query plan in a structured format. ```python >>> import dataclasses >>> import polars as pl ->>> from cudf_polars.experimental.explain import serialize_query +>>> from cudf_polars.streaming.explain import serialize_query >>> q = pl.LazyFrame({"a": ['a', 'b', 'a'], "b": [1, 2, 3]}).group_by("a").agg(pl.len()) >>> dataclasses.asdict(serialize_query(q, engine=pl.GPUEngine())) {'roots': ['526964741'], diff --git a/docs/cudf/source/cudf_polars/index.md b/docs/cudf/source/cudf_polars/index.md index 9f641f33a024..a0b8c2b09690 100644 --- a/docs/cudf/source/cudf_polars/index.md +++ b/docs/cudf/source/cudf_polars/index.md @@ -9,8 +9,10 @@ and runs on the CPU. ## Install -Follow the [RAPIDS installation guide](https://docs.rapids.ai/install/) and pick the -`cudf-polars` package for your CUDA and Python versions. For example, with conda: +Follow the [NVIDIA CUDA-X installation +guide](https://developer.nvidia.com/topics/ai/data-science/cuda-x-for-data-science##section-install-and-deploy-in-your-environment) +and pick the `cudf-polars` package for your CUDA and Python versions. For +example, with conda: ```bash conda install -c rapidsai -c conda-forge -c nvidia cudf-polars diff --git a/docs/cudf/source/cudf_polars/memory_errors.md b/docs/cudf/source/cudf_polars/memory_errors.md index f39a4938d4a9..693a2e5699f8 100644 --- a/docs/cudf/source/cudf_polars/memory_errors.md +++ b/docs/cudf/source/cudf_polars/memory_errors.md @@ -110,4 +110,6 @@ constructing the GPU engine for queries. For the full list of engine configuration options, including `target_partition_size` and `max_concurrent_io_tasks`, see {doc}`options`. For the full list of memory -and spill configuration options see the [RapidsMPF configuration reference](https://docs.rapids.ai/api/rapidsmpf/stable/configuration/#general). +and spill configuration options see the [RapidsMPF configuration reference][rapidsmpf-config]. + +[rapidsmpf-config]: inv:rapidsmpf:std:label:#configuration:general diff --git a/docs/cudf/source/cudf_polars/options.md b/docs/cudf/source/cudf_polars/options.md index d4af20497ff9..321ed9513bc4 100644 --- a/docs/cudf/source/cudf_polars/options.md +++ b/docs/cudf/source/cudf_polars/options.md @@ -136,4 +136,4 @@ These environment variables are intended for library developers and advanced use | `CUDF_POLARS_WARN_UNSTABLE` | Raises a `cudf_polars.UnstableWarning` whenever an unstable cudf-polars feature is used. Set to `1` to enable. | `0` | -[rapidsmpf-config]: https://docs.rapids.ai/api/rapidsmpf/nightly/configuration/ +[rapidsmpf-config]: inv:rapidsmpf:std:doc:#configuration diff --git a/docs/cudf/source/cudf_polars/profiling.md b/docs/cudf/source/cudf_polars/profiling.md index 35e0baa7ffe6..526ab6c473e1 100644 --- a/docs/cudf/source/cudf_polars/profiling.md +++ b/docs/cudf/source/cudf_polars/profiling.md @@ -103,8 +103,8 @@ KvikIO I/O summary ``` Every row is also an attribute, `s.bytes_read`, `s.busy_ns` and so on. See the -[KvikIO reference][kvikio-stats] for the full set, and [busy time and bandwidth][kvikio-busy] -for how the busy figures are measured. +[KvikIO statistics reference][kvikio-stats] for the full set, and [busy time and +bandwidth][kvikio-busy] for how the busy figures are measured. ### What is and is not counted @@ -250,9 +250,9 @@ shape: (2, 3) [nsight]: https://developer.nvidia.com/nsight-systems [nvtx]: https://nvidia.github.io/NVTX/ -[kvikio-stats]: https://docs.rapids.ai/api/kvikio/nightly/statistics/ -[kvikio-busy]: https://docs.rapids.ai/api/kvikio/nightly/statistics/#busy-time-and-bandwidth -[rapidsmpf-stats]: https://docs.rapids.ai/api/rapidsmpf/nightly/statistics/ +[kvikio-stats]: inv:kvikio:std:doc:#statistics +[kvikio-busy]: +[rapidsmpf-stats]: inv:rapidsmpf:std:doc:#statistics [structlog]: https://www.structlog.org/en/stable/ [structlog-configure]: https://www.structlog.org/en/stable/configuration.html [structlog-context]: https://www.structlog.org/en/stable/contextvars.html diff --git a/docs/cudf/source/index.rst b/docs/cudf/source/index.rst index cf06d7670337..74410115846c 100644 --- a/docs/cudf/source/index.rst +++ b/docs/cudf/source/index.rst @@ -1,9 +1,10 @@ NVIDIA cuDF Documentation ========================= -**NVIDIA cuDF** (pronounced "KOO-dee-eff") is a GPU-accelerated library for tabular -data processing. It is part of the `RAPIDS `_ suite of -libraries and is composed of multiple sub-projects: +**NVIDIA cuDF** (pronounced "KOO-dee-eff") is a GPU-accelerated library for +tabular data processing. It is part of `NVIDIA CUDA-X for Data Science +`_ +suite, and is composed of multiple sub-projects: .. list-table:: :header-rows: 1 @@ -15,13 +16,17 @@ libraries and is composed of multiple sub-projects: - A Python library providing a `pandas `_-like DataFrame API and a zero-code change accelerator, `cudf.pandas `_, for existing pandas code. * - `cudf-polars `_ - A Python library providing a GPU engine for `Polars `_. - * - `dask-cudf `_ + * - :doc:`dask-cudf:index` - A Python library providing a GPU backend for `Dask `_ DataFrames. * - `libcudf `_ - A CUDA C++ library with `Apache Arrow `_ compliant data structures and fundamental algorithms for tabular data. * - `pylibcudf `_ - A Python library providing `Cython `_ bindings for libcudf. +See the `installation and deployment guide +_` +to get up-and-running with cuDF. + .. toctree:: :maxdepth: 1 :caption: Libraries diff --git a/docs/cudf/source/libcudf/api_docs/lists_classes.rst b/docs/cudf/source/libcudf/api_docs/lists_classes.rst index 9b89c1647466..9444b202a123 100644 --- a/docs/cudf/source/libcudf/api_docs/lists_classes.rst +++ b/docs/cudf/source/libcudf/api_docs/lists_classes.rst @@ -3,3 +3,6 @@ Lists Classes .. doxygengroup:: lists_classes :members: + +.. doxygenclass:: cudf::list_view + :project: libcudf diff --git a/docs/cudf/source/libcudf/api_docs/structs_classes.rst b/docs/cudf/source/libcudf/api_docs/structs_classes.rst index 2669c2884d63..2f6e2be7c02c 100644 --- a/docs/cudf/source/libcudf/api_docs/structs_classes.rst +++ b/docs/cudf/source/libcudf/api_docs/structs_classes.rst @@ -3,3 +3,6 @@ Structs Classes .. doxygengroup:: structs_classes :members: + +.. doxygenclass:: cudf::struct_view + :project: libcudf diff --git a/docs/cudf/source/libcudf/developer_guide/BENCHMARKING.md b/docs/cudf/source/libcudf/developer_guide/BENCHMARKING.md new file mode 100644 index 000000000000..9c155c258a21 --- /dev/null +++ b/docs/cudf/source/libcudf/developer_guide/BENCHMARKING.md @@ -0,0 +1,99 @@ +# Unit Benchmarking in libcudf + +Unit benchmarks in libcudf are written using [NVBench](https://github.com/NVIDIA/nvbench). +While many existing benchmarks are written using +[Google Benchmark](https://github.com/google/benchmark), new benchmarks should use NVBench. + +The NVBench library is similar to Google Benchmark, but has several quality of life improvements +when doing GPU benchmarking such as displaying the fraction of peak memory bandwidth achieved and +details about the GPU hardware. + +Both NVBench and Google Benchmark provide many options for specifying ranges of parameters to +benchmark, as well as to control the time unit reported, among other options. Refer to existing +benchmarks in `cpp/benchmarks` to understand the options. + +## Directory and File Naming + +The naming of unit benchmark directories and source files should be consistent with the feature +being benchmarked. For example, the benchmarks for APIs in `copying.hpp` should live in +`cpp/benchmarks/copying`. Each feature (or set of related features) should have its own +benchmark source file named `.cu/cpp`. For example, `cpp/src/copying/scatter.cu` has +benchmarks in `cpp/benchmarks/copying/scatter.cu`. + +In the interest of improving compile time, whenever possible, test source files should be `.cpp` +files because `nvcc` is slower than `gcc` in compiling host code. Note that `thrust::device_vector` +includes device code, and so must only be used in `.cu` files. `rmm::device_uvector`, +`rmm::device_buffer` and the various `column_wrapper` types described in [Testing](TESTING.md) +can be used in `.cpp` files, and are therefore preferred in test code over `thrust::device_vector`. + +## CUDA Asynchrony and benchmark accuracy + +CUDA computations and operations like copies are typically asynchronous with respect to host code, +so it is important to carefully synchronize in order to ensure the benchmark timing is not stopped +before the feature you are benchmarking has completed. An RAII helper class `cuda_event_timer` is +provided in `cpp/benchmarks/synchronization/synchronization.hpp` to help with this. This class +can also optionally clear the GPU L2 cache in order to ensure cache hits do not artificially inflate +performance in repeated iterations. + +## Data generation + +For generating benchmark input data, helper functions are available at [cpp/benchmarks/common/generate_input.hpp](https://github.com/rapidsai/cudf/blob/main/cpp/benchmarks/common/generate_input.hpp). The input data generation happens on device, in contrast to any `column_wrapper` where data generation happens on the host. +* `create_sequence_table` can generate sequence columns starting with value 0 in first row and increasing by 1 in subsequent rows. +* `create_random_column` can generate a column filled with random data. The random data parameters are configurable. +* `create_random_table` can generate a table of columns filled with random data. The random data parameters are configurable. + +## NVTX Ranges + +In order to aid in performance measurement, add NVTX ranges to the data generation and/or +the function being measured. Use either the `CUDF_BENCHMARK_RANGE` or `cudf::benchmark::scoped_range` +for declaring NVTX ranges in the current scope as appropriate: +- Use the `CUDF_BENCHMARK_RANGE()` macro if you want to use the name of the function as the name of the +NVTX range +- Use `cudf::benchmark::scoped_range rng{"custom_name"};` to provide a custom name for the current scope's +NVTX range + +For more information about NVTX, see [here](https://github.com/NVIDIA/NVTX/tree/dev/c). + +Do not use the libcudf nvtx range classes in benchmark source code to help separate the execution +scope properly in the NSight Compute tools. + +## What should we benchmark? + +In general, we should benchmark all features over a range of data sizes and types, so that we can +catch regressions across libcudf changes. However, running many benchmarks is expensive, so ideally +we should sample the parameter space in such a way to get good coverage without having to test +exhaustively. + +A rule of thumb is that we should benchmark with enough data to reach the point where the algorithm +reaches its saturation bottleneck, whether that bottleneck is bandwidth or computation. Using data +sets larger than this point is generally not helpful, except in specific cases where doing so +exercises different code and can therefore uncover regressions that smaller benchmarks will not +(this should be rare). + +## Running and Comparing NVBench Benchmarks in libcudf + +### Running Benchmarks + +By default, benchmarks are **not** built as part of the libcudf build process. To enable them, pass the `BUILD_BENCHMARKS=ON` flag to CMake such as: + +```bash +cmake -DBUILD_BENCHMARKS=ON cpp/build/latest +``` +This will build the NVBench benchmark executables under the `cpp/build/latest/benchmarks` +directory. Each benchmark is compiled into its own binary with a `_NVBENCH` suffix. + +To list available benchmarks: + +```bash +ls cpp/build/latest/benchmarks/*_NVBENCH +``` + +To view benchmark options: + +```bash +./cpp/build/latest/benchmarks/_NVBENCH --help +``` + +### Comparing Benchmarks + +To compare two benchmark runs, use the [nvbench_compare.py](https://github.com/NVIDIA/nvbench/blob/main/scripts/nvbench_compare.py) script provided by NVBench. diff --git a/docs/cudf/source/libcudf/developer_guide/DEVELOPER_GUIDE.md b/docs/cudf/source/libcudf/developer_guide/DEVELOPER_GUIDE.md new file mode 100644 index 000000000000..b08e0ca4262c --- /dev/null +++ b/docs/cudf/source/libcudf/developer_guide/DEVELOPER_GUIDE.md @@ -0,0 +1,1759 @@ +(md_developer_guide)= +# libcudf C++ Developer Guide + +```{toctree} +:hidden: +:maxdepth: 1 + +DOCUMENTATION +TESTING +BENCHMARKING +PROFILING +``` + +This document serves as a guide for contributors to libcudf C++ code. Developers should also refer +to these additional files for further documentation of libcudf best practices. + +* [Documentation Guide](DOCUMENTATION.md) for guidelines on documenting libcudf code. +* [Testing Guide](TESTING.md) for guidelines on writing unit tests. +* [Benchmarking Guide](BENCHMARKING.md) for guidelines on writing unit benchmarks. +* [Profiling Guide](PROFILING.md) for guidelines on profiling libcudf code. + +# Overview + +libcudf is a C++ library that provides GPU-accelerated data-parallel algorithms for processing +column-oriented tabular data. libcudf provides algorithms including slicing, filtering, sorting, +various types of aggregations, and database-type operations such as grouping and joins. libcudf +serves a number of clients via multiple language interfaces, including Python and Java. Users may +also use libcudf directly from C++ code. + +## Lexicon + +This section defines terminology used within libcudf. + +### Column + +A column is an array of data of a single type. Along with Tables, columns are the fundamental data +structures used in libcudf. Most libcudf algorithms operate on columns. Columns may have a validity +mask representing whether each element is valid or null (invalid). Columns of nested types are +supported, meaning that a column may have child columns. A column is the C++ equivalent to a cuDF +Python [Series](https://docs.nvidia.com/cudf/latest/cudf/api_docs/series/). + +### Element + +An individual data item within a column. Also known as a row. + +### Scalar + +A type representing a single element of a data type. + +### Table + +A table is a collection of columns that all have the same number of elements (rows). A table may +also have zero columns while still carrying a row count, mirroring an `(N, 0)` DataFrame. A table is +the C++ equivalent to a cuDF Python +[DataFrame](https://docs.nvidia.com/cudf/latest/cudf/api_docs/dataframe/). + +### View + +A view is a non-owning object that provides zero-copy access (possibly with slicing or offsets) to +data owned by another object. Examples are column views and table views. + +# Directory Structure and File Naming + +External/public libcudf APIs are grouped based on functionality into an appropriately titled +header file in `cudf/cpp/include/cudf/`. For example, `cudf/cpp/include/cudf/copying.hpp` +contains the APIs for functions related to copying from one column to another. Note the `.hpp` +file extension used to indicate a C++ header file. + +External/public libcudf C++ API header files need to mark all symbols inside of them with `CUDF_EXPORT`. +This is done by placing the macro on the `namespace cudf` as seen below. Markup on namespace +require them not to be nested, so the `cudf` namespace must be kept by itself. + +```c++ + +#pragma once + +namespace CUDF_EXPORT cudf { +namespace lists { + +... + + +} // namespace lists +} // namespace CUDF_EXPORT cudf + +``` + + +The naming of external API headers should be consistent with the name of the folder that contains +the source files that implement the API. For example, the implementation of the APIs found in +`cudf/cpp/include/cudf/copying.hpp` are located in `cudf/src/copying`. Likewise, the unit tests for +the APIs reside in `cudf/tests/copying/`. + +Internal API headers containing `detail` namespace definitions that are used across translation +units inside libcudf should be placed according to their namespace: + +- For APIs in `cudf::detail`, place headers in `include/cudf/detail/`. +- For APIs in a sub-namespace's detail (e.g., `cudf::hashing::detail` or `cudf::strings::detail`), + place headers in `include/cudf//detail/` (e.g., `include/cudf/hashing/detail/`). + +Internal C++ headers may need `CUDF_EXPORT` if that internal functionality is tested directly (as +opposed to tested via only public APIs). + +All headers in cudf should use `#pragma once` for include guards. + +## File extensions + +- `.hpp` : C++ header files +- `.cpp` : C++ source files +- `.cu` : CUDA C++ source files +- `.cuh` : Headers containing CUDA device code + +Only use `.cu` and `.cuh` if necessary. A good indicator is the inclusion of `__device__` and other +symbols that are only recognized by `nvcc`. Another indicator is Thrust algorithm APIs with a device +execution policy (always `rmm::exec_policy_nosync` in libcudf). + +## Code and Documentation Style and Formatting + +libcudf code uses [snake_case](https://en.wikipedia.org/wiki/Snake_case) for all names except in a +few cases: template parameters, unit tests and test case names may use Pascal case, aka +[UpperCamelCase](https://en.wikipedia.org/wiki/Camel_case). We do not use +[Hungarian notation](https://en.wikipedia.org/wiki/Hungarian_notation), except sometimes when naming +device data variables and their corresponding host copies. Private member variables are typically +prefixed with an underscore. + +```c++ +template +void algorithm_function(int x, cuda::stream_ref s, rmm::device_async_resource_ref mr) +{ + ... +} + +class utility_class +{ + ... +private: + int _rating{}; + std::unique_ptr _column{}; +} + +TYPED_TEST_SUITE(RepeatTypedTestFixture, cudf::test::FixedWidthTypes); + +TYPED_TEST(RepeatTypedTestFixture, RepeatScalarCount) +{ + ... +} +``` + +C++ formatting is enforced using `clang-format`. You should configure `clang-format` on your +machine to use the `cudf/cpp/.clang-format` configuration file, and run `clang-format` on all +changed code before committing it. The easiest way to do this is to configure your editor to +"format on save." + +Aspects of code style not discussed in this document and not automatically enforceable are typically +caught during code review, or not enforced. + +### C++ Guidelines + +In general, we recommend following +[C++ Core Guidelines](https://isocpp.github.io/CppCoreGuidelines/CppCoreGuidelines). We also +recommend watching Sean Parent's [C++ Seasoning talk](https://www.youtube.com/watch?v=W2tWOdzgXHA), +and we try to follow his rules: "No raw loops. No raw pointers. No raw synchronization primitives." + + * Prefer algorithms from STL and Thrust to raw loops. + * Prefer libcudf and RMM [owning data structures and views](#libcudf-data-structures) to raw + pointers and raw memory allocation. + * libcudf doesn't have a lot of CPU-thread concurrency, but there is some. And currently libcudf + does use raw synchronization primitives. So we should revisit Parent's third rule and improve + here. + +Additional style guidelines for libcudf code: + + * Prefer "east const", placing `const` after the type. This is not + automatically enforced by `clang-format` because the option + `QualifierAlignment: Right` has been observed to produce false negatives and + false positives. + * [NL.11: Make Literals + Readable](https://isocpp.github.io/CppCoreGuidelines/CppCoreGuidelines#nl11-make-literals-readable): + Decimal values should use integer separators every thousands place, like + `1'234'567`. Hexadecimal values should use separators every 4 characters, + like `0x0123'ABCD`. + +Documentation is discussed in the [Documentation Guide](DOCUMENTATION.md). + +### Device Code and `constexpr` Functions + +libcudf does **not** enable `--expt-relaxed-constexpr`. This means the CUDA compiler will reject +calls to `constexpr` host functions from `__device__` or `__global__` code. Follow these rules: + + * Every `constexpr` function that may be called from device code **must** be explicitly annotated + with `__device__` or `CUDF_HOST_DEVICE`. A bare `constexpr` without an execution-space + annotation is host-only. + * In `__device__` and `CUDF_HOST_DEVICE` functions, use `cuda::std::` + utilities (type traits, algorithms, math functions) instead of `std::`. The `std::` counterparts + are host-only and will cause compilation errors in device code. For example: + - `cuda::std::is_void_v` instead of `std::is_void_v` + - `cuda::std::min` / `cuda::std::max` instead of `std::min` / `std::max` + - `cuda::std::numeric_limits` instead of `std::numeric_limits` + * Prefer `cuda::std::` type traits and constexpr functions over `std::` in templates that may be + instantiated in device code, even if the template itself is not annotated with `__device__`. + +### Includes + +The following guidelines apply to organizing `#include` lines. + + * Group includes by library (e.g. cuDF, RMM, Thrust, STL). `clang-format` will respect the + groupings and sort the individual includes within a group lexicographically. + * Separate groups by a blank line. + * Order the groups from "nearest" to "farthest". In other words, local includes, then includes + from other RAPIDS libraries, then includes from related libraries, like ``, then + includes from dependencies installed with cuDF, and then standard headers (for example + ``, ``). + * We use clang-format for grouping and sorting headers automatically. See the + `cudf/cpp/.clang-format` file for specifics. + * Use `<>` for all includes except for internal headers that are not in the `include` + directory. In other words, if it is a cuDF internal header (e.g. in the `src` or `test` + directory), the path will not start with `cudf` (e.g. `#include `) so it + should use quotes. Example: `#include "io/utilities/hostdevice_vector.hpp"`. + * `cudf_test` and `nvtext` are separate libraries within the `libcudf` repo. As such, they have + public headers in `include` that should be included with `<>`. + * Tools like `clangd` often auto-insert includes when they can, but they usually get the grouping + and brackets wrong. Correct the usage of quotes or brackets and then run clang-format to correct + the grouping. + * Follow the "include what you use" principle. Directly `#include` headers that declare or forward declare as appropriate every symbol a file uses and do not rely on headers being pulled in transitively through another include. + * Do not include headers whose symbols the file does not use and avoid excessive including especially in header files. Double check this when you remove code. + * Avoid relative paths with `..` when possible. Paths with `..` are necessary when including + (internal) headers from source paths not in the same directory as the including file, + because source paths are not passed with `-I`. + * Avoid including library internal headers from non-internal files. For example, try not to include + headers from libcudf `src` directories in tests or in libcudf public headers. If you find + yourself doing this, start a discussion about moving (parts of) the included internal header + to a public header. + +# libcudf Data Structures + +Application data in libcudf is contained in Columns and Tables, but there are a variety of other +data structures you will use when developing libcudf code. + +## Views and Ownership + +Resource ownership is an essential concept in libcudf. In short, an "owning" object owns a +resource (such as device memory). It acquires that resource during construction and releases the +resource in destruction ([RAII](https://en.cppreference.com/cpp/language/raii)). A "non-owning" +object does not own resources. Any class in libcudf with the `*_view` suffix is non-owning. For more +detail see the [`libcudf` presentation.](https://docs.google.com/presentation/d/1zKzAtc1AWFKfMhiUlV5yRZxSiPLwsObxMlWRWz_f5hA/edit?usp=sharing) + +libcudf functions typically take views as input (`column_view` or `table_view`) +and produce `unique_ptr`s to owning objects as output. For example, + +```c++ +std::unique_ptr sort(table_view const& input); +``` + +## Memory Resources + +libcudf allocates all device memory via RMM memory resources (MR) or CUDA MRs. Either type +can be passed to libcudf functions via `rmm::device_async_resource_ref` parameters. See the +[RMM documentation](https://github.com/rapidsai/rmm/blob/main/README.md) for details. + +(rmmdevice_memory_resource)= +### Current Device Memory Resource + +RMM provides a "default" memory resource for each device and functions to access and set it. libcudf +provides wrappers for these functions in `cpp/include/cudf/utilities/memory_resource.hpp`. +All memory resource parameters should be defaulted to use the return value of +`cudf::get_current_device_resource_ref()`. + +### Resource Refs + +Memory resources are passed via resource ref parameters. A resource ref is a memory resource wrapper +that enables consumers to specify properties of resources that they expect. These are defined +in the `cuda::mr` namespace of libcu++, but RMM provides some convenience aliases in +`rmm/resource_ref.hpp`. + - `rmm::device_resource_ref` accepts a memory resource that provides synchronous allocation + of device-accessible memory. + - `rmm::device_async_resource_ref` accepts a memory resource that provides stream-ordered allocation + of device-accessible memory. + - `rmm::host_resource_ref` accepts a memory resource that provides synchronous allocation of host- + accessible memory. + - `rmm::host_async_resource_ref` accepts a memory resource that provides stream-ordered allocation + of host-accessible memory. + - `rmm::host_device_resource_ref` accepts a memory resource that provides synchronous allocation of + host- and device-accessible memory. + - `rmm::host_async_resource_ref` accepts a memory resource that provides stream-ordered allocation + of host- and device-accessible memory. + +See the libcu++ [docs on `resource_ref`](https://nvidia.github.io/cccl/libcudacxx/extended_api/memory_resource/resource_ref.html) +for more information. + +## cudf::column + +`cudf::column` is a core owning data structure in libcudf. Most libcudf public APIs produce either +a `cudf::column` or a `cudf::table` as output. A `column` contains `device_buffer`s which own the +device memory for the elements of a column and an optional null indicator bitmask. + +Implicitly convertible to `column_view` and `mutable_column_view`. + +Movable and copyable. A copy performs a deep copy of the column's contents, whereas a move moves +the contents from one column to another. + +Example: + +```c++ +cudf::column col{...}; + +cudf::column copy{col}; // Copies the contents of `col` +cudf::column const moved_to{std::move(col)}; // Moves contents from `col` + +column_view v = moved_to; // Implicit conversion to non-owning column_view +// mutable_column_view m = moved_to; // Cannot create mutable view to const column +``` + +A `column` may have nested (child) columns, depending on the data type of the column. For example, +`LIST`, `STRUCT`, and `STRING` type columns. + +### cudf::column_view + +`cudf::column_view` is a core non-owning data structure in libcudf. It is an immutable, +non-owning view of device memory as a column. Most libcudf public APIs take views as inputs. + +A `column_view` may be a view of a "slice" of a column. For example, it might view rows 75-150 of a +column with 1000 rows. The `size()` of this `column_view` would be `75`, and accessing index `0` of +the view would return the element at index `75` of the owning `column`. Internally, this is +implemented by storing in the view a pointer, an offset, and a size. `column_view::data()` +returns a pointer iterator to `column_view::head() + offset`. + +### cudf::mutable_column_view + +A *mutable*, non-owning view of device memory as a column. Used for detail APIs and (rare) public +APIs that modify columns in place. + +### cudf::column_device_view + +An immutable, non-owning view of device data as a column of elements that is trivially copyable and +usable in CUDA device code. Used to pass `column_view` data as input to CUDA kernels and device +functions (including Thrust algorithms) + +### cudf::mutable_column_device_view + +A mutable, non-owning view of device data as a column of elements that is trivially copyable and +usable in CUDA device code. Used to pass `column_view` data to be modified on the device by CUDA +kernels and device functions (including Thrust algorithms). + +## cudf::table + +Owning class for a set of `cudf::column`s all with equal number of elements. This is the C++ +equivalent to a data frame. + +Implicitly convertible to `cudf::table_view` and `cudf::mutable_table_view` + +Movable and copyable. A copy performs a deep copy of all columns, whereas a move moves all columns +from one table to another. + +### cudf::table_view + +An *immutable*, non-owning view of a table. + +### cudf::mutable_table_view + +A *mutable*, non-owning view of a table. + +## cudf::size_type + +The `cudf::size_type` is the type used for the number of elements in a column, indices to address +specific elements, segments for subsets of column elements, etc. It is equivalent to a signed, +32-bit integer type and therefore has a maximum value of 2147483647. Some APIs also accept negative +index values and those functions support a minimum value of -2147483648. This fundamental type also +influences output values not just for column size limits but for counting elements as well. + +Offset to elements within a column should be either `int32_t` or `int64_t` appropriately, except +for `LIST` columns that only support `int32_t` offsets. + +## Spans + +libcudf provides `span` classes that mimic C++20 `std::span`, which is a lightweight +view of a contiguous sequence of objects. libcudf provides two classes, `host_span` and +`device_span`, which can be constructed from multiple container types, or from a pointer +(host or device, respectively) and size, or from iterators. `span` types are useful for defining +generic (internal) interfaces which work with multiple input container types. `device_span` can be +constructed from `thrust::device_vector`, `rmm::device_vector`, or `rmm::device_uvector`. +`host_span` can be constructed from `thrust::host_vector`, `std::vector`, or `std::basic_string`. + +If you are defining internal (detail) functions that operate on vectors, use spans for the input +vector parameters rather than a specific vector type, to make your functions more widely applicable. + +When a `span` refers to immutable elements, use `span`, not `span const`. Since a span +is lightweight view, it does not propagate `const`-ness. Therefore, `const` should be applied to +the template type parameter, not to the `span` itself. Also, `span` should be passed by value +because it is a lightweight view. APIS in libcudf that take spans as input will look like the +following function that copies device data to a host `std::vector`. + +```c++ +template +std::vector make_std_vector_async(device_span v, cuda::stream_ref stream) +``` + +### When to use `host_span` vs `std::span` + +For host-side data, prefer `std::span` and reserve `cudf::host_span` for the cases where its +libcudf-specific extensions are actually needed. + +Use `std::span` when the parameter is purely a host buffer view and the function does not need +to know whether the memory is device-accessible. + +Use `cudf::host_span` only when one of the following applies: + +1. The function needs to query `is_device_accessible()` to take a different code path for pinned + or otherwise device-reachable host memory (for example, to enable copy-engine optimizations or + skip an explicit host-to-device copy). +2. The function must accept a libcudf-specific container that `std::span` cannot be constructed + from directly. + +## cudf::scalar + +A `cudf::scalar` is an object that can represent a singular, nullable value of any of the types +currently supported by cudf. Each type of value is represented by a separate type of scalar class +which are all derived from `cudf::scalar`. e.g. A `numeric_scalar` holds a single numerical value, +a `string_scalar` holds a single string. The data for the stored value resides in device memory. + +A `list_scalar` holds the underlying data of a single list. This means the underlying data can be +any type that cudf supports. For example, a `list_scalar` representing a list of integers stores a +`cudf::column` of type `INT32`. A `list_scalar` representing a list of lists of integers stores a +`cudf::column` of type `LIST`, which in turn stores a column of type `INT32`. + +|Value type|Scalar class|Notes| +|-|-|-| +|fixed-width|`fixed_width_scalar`| `T` can be any fixed-width type| +|numeric|`numeric_scalar` | `T` can be `int8_t`, `int16_t`, `int32_t`, `int64_t`, `float` or `double`| +|fixed-point|`fixed_point_scalar` | `T` can be `numeric::decimal32` or `numeric::decimal64`| +|timestamp|`timestamp_scalar` | `T` can be `timestamp_D`, `timestamp_s`, etc.| +|duration|`duration_scalar` | `T` can be `duration_D`, `duration_s`, etc.| +|string|`string_scalar`| This class object is immutable| +|list|`list_scalar`| Underlying data can be any type supported by cudf | + +### Construction +`scalar`s can be created using either their respective constructors or using factory functions like +`make_numeric_scalar()`, `make_timestamp_scalar()` or `make_string_scalar()`. + +### Casting +All the factory methods return a `unique_ptr` which needs to be statically downcasted to +its respective scalar class type before accessing its value. Their validity (nullness) can be +accessed without casting. Generally, the value needs to be accessed from a function that is aware +of the value type e.g. a functor that is dispatched from `type_dispatcher`. To cast to the +requisite scalar class type given the value type, use the mapping utility `scalar_type_t` provided +in `type_dispatcher.hpp` : + +```c++ +//unique_ptr s = make_numeric_scalar(...); + +using ScalarType = cudf::scalar_type_t; +// ScalarType is now numeric_scalar +auto s1 = static_cast(s.get()); +``` + +### Passing to device +Each scalar type, except `list_scalar`, has a corresponding non-owning device view class which +allows access to the value and its validity from the device. This can be obtained using the function +`get_scalar_device_view(ScalarType s)`. Note that a device view is not provided for a base scalar +object, only for the derived typed scalar class objects. + +The underlying data for `list_scalar` can be accessed via `view()` method. For non-nested data, +the device view can be obtained via function `column_device_view::create(column_view)`. For nested +data, a specialized device view for list columns can be constructed via +`lists_column_device_view(column_device_view)`. + +# libcudf Policies and Design Principles + +`libcudf` is designed to provide thread-safe, single-GPU accelerated algorithm primitives for +solving a wide variety of problems that arise in data science. APIs are written to execute on the +default GPU, which can be controlled by the caller through standard CUDA device APIs or environment +variables like `CUDA_VISIBLE_DEVICES`. Our goal is to enable diverse use cases like Spark or Pandas +to benefit from the performance of GPUs, and libcudf relies on these higher-level layers like Spark +or Dask to orchestrate multi-GPU tasks. + +To best satisfy these use-cases, libcudf prioritizes performance and flexibility, which sometimes +may come at the cost of convenience. While we welcome users to use libcudf directly, we design with +the expectation that most users will be consuming libcudf through higher-level layers like Spark or +cuDF Python that handle some of details that direct users of libcudf must handle on their own. We +document these policies and the reasons behind them here. + +## libcudf does not introspect data + +libcudf APIs generally do not perform deep introspection and validation of input data. +There are numerous reasons for this: +1. It violates the single responsibility principle: validation is separate from execution. +2. Since libcudf data structures store data on the GPU, any validation incurs _at minimum_ the + overhead of a kernel launch, and may in general be prohibitively expensive. +3. API promises around data introspection often significantly complicate implementation. + +Users are therefore responsible for passing valid data into such APIs. +_Note that this policy does not mean that libcudf performs no validation whatsoever_. +libcudf APIs should still perform any validation that does not require introspection. +To give some idea of what should or should not be validated, here are (non-exhaustive) lists of +examples. + +**Things that libcudf should validate**: +- Input column/table sizes or data types + +**Things that libcudf should not validate**: +- Integer overflow +- Ensuring that outputs will not exceed the [`size_type`](#cudfsize_type) row count limit for a + given set of inputs + +This policy describes libcudf's default behavior. Some APIs offer opt-in overflow-aware variants for +callers that need strict semantics, such as the `SUM_OVERFLOW` aggregation, which reports +overflow through an output flag, and the overflow-checking AST arithmetic operators +(`ADD_OVERFLOW`, `SUB_OVERFLOW`, and similar) that raise on overflow. +These are explicit, caller-selected behaviors rather than validation performed on every API. + + +## libcudf expects nested types to have sanitized null masks + +Various libcudf APIs accepting columns of nested data types (such as `LIST` or `STRUCT`) may assume +that these columns have been sanitized. In this context, sanitization refers to ensuring that the +null elements in a column with a nested dtype are compatible with the elements of nested columns. +Specifically: +- Null elements of list columns should also be empty. The starting offset of a null element should + be equal to the ending offset. +- Null elements of struct columns should also be null elements in the underlying structs. +- For compound columns, nulls should only be present at the level of the parent column. Child + columns should not contain nulls. +- Slice operations on nested columns do not propagate offsets to child columns. + +libcudf APIs _should_ promise to never return "dirty" columns, i.e. columns containing unsanitized +data. Therefore, the only problem is if users construct input columns that are not correctly +sanitized and then pass those into libcudf APIs. + +## Null values of fixed-width columns are undefined + +For columns of fixed-width types (such as integers, floats, timestamps, and durations), the values +corresponding to null elements (where the validity mask bit is set to null) are **undefined** and +may contain arbitrary data. libcudf makes no guarantees about the initialization or content of these +values. + +Code should not assume that null rows in fixed-width columns contain any particular value, including +zero. Algorithms must rely solely on the validity mask to determine nullness and should not inspect +the underlying data values for null elements. + +This policy applies only to fixed-width types. It does **not** apply to variable-width types +(strings) or nested types (lists, structs), which have their own requirements as described in the +sections above. + +(async-apis)= +## Treat libcudf APIs as if they were asynchronous + +libcudf APIs called on the host do not guarantee that the stream is synchronized before returning. +Work in libcudf occurs on `cudf::get_default_stream().value`, which defaults to the CUDA default +stream (stream 0). Note that the stream 0 behavior differs if [per-thread default stream is +enabled](https://docs.nvidia.com/cuda/cuda-runtime-api/stream-sync-behavior.html) via +`CUDF_USE_PER_THREAD_DEFAULT_STREAM`. Any data provided to or returned by libcudf that uses a +separate non-blocking stream requires synchronization with the default libcudf stream to ensure +stream safety. + +## libcudf generally does not make ordering guarantees + +Functions like merge or groupby in libcudf make no guarantees about the order of entries in the +output. Promising deterministic ordering is not, in general, conducive to fast parallel algorithms. +Calling code is responsible for performing sorts after the fact if sorted outputs are needed. + +## libcudf does not promise specific exception messages + +libcudf documents the exceptions that will be thrown by an API for different kinds of invalid +inputs. The types of those exceptions (e.g. `cudf::logic_error`) are part of the public API. +However, the explanatory string returned by the `what` method of those exceptions is not part of the +API and is subject to change. Calling code should not rely on the contents of libcudf error +messages to determine the nature of the error. For information on the types of exceptions that +libcudf throws under different circumstances, see the [section on error handling](#errors). + +# libcudf API and Implementation + +(streams)= +## Streams + +libcudf is in the process of adding support for asynchronous execution using +CUDA streams. In order to facilitate the usage of streams, all new libcudf APIs +that allocate device memory or execute a kernel should accept an +`cuda::stream_ref` parameter at the end with a default value of +`cudf::get_default_stream()`. There is one exception to this rule: if the API +also accepts a memory resource parameter, the stream parameter should be placed +just *before* the memory resource. This API should then forward the call to a +corresponding `detail` API with an identical signature, except that the +`detail` API should not have a default parameter for the stream ([detail APIs +should always avoid default parameters](#default-parameters)). The +implementation should be wholly contained in the `detail` API definition and +use only asynchronous versions of CUDA APIs with the stream parameter. + +In order to make the `detail` API callable from other libcudf functions, it should be exposed in a +header placed in the `cudf/cpp/include/detail/` directory. +The declaration is not necessary if no other libcudf functions call the `detail` function. + +For example: + +```c++ +// cpp/include/cudf/header.hpp +void external_function(..., + cuda::stream_ref stream = cudf::get_default_stream(), + rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); + +// cpp/include/cudf/detail/header.hpp +namespace detail{ +void external_function(..., cuda::stream_ref stream, rmm::device_async_resource_ref mr) +} // namespace detail + +// cudf/src/implementation.cpp +namespace detail{ +// Use the stream parameter in the detail implementation. +void external_function(..., cuda::stream_ref stream, rmm::device_async_resource_ref mr){ + // Implementation uses the stream with async APIs. + rmm::device_buffer buff(..., stream, mr); + CUDF_CUDA_TRY(cudaMemcpyAsync(...,stream.value())); + kernel<<<..., stream>>>(...); + thrust::algorithm(rmm::exec_policy_nosync(stream), ...); +} +} // namespace detail + +void external_function(..., cuda::stream_ref stream, rmm::device_async_resource_ref mr) +{ + CUDF_FUNC_RANGE(); // Generates an NVTX range for the lifetime of this function. + detail::external_function(..., stream, mr); +} +``` + +**Note:** It is important to synchronize the stream if *and only if* it is necessary. For example, +when a non-pointer value is returned from the API that is the result of an asynchronous +device-to-host copy, the stream used for the copy should be synchronized before returning. However, +when a column is returned, the stream should not be synchronized because doing so will break +asynchrony. + +**Note:** `cudaDeviceSynchronize()` should *never* be used. +This limits the ability to do any multi-stream/multi-threaded work with libcudf APIs. + +### Stream Creation + +There may be times in implementing libcudf features where it would be advantageous to use streams +*internally*, i.e., to accomplish overlap in implementing an algorithm. However, dynamically +creating a stream can be expensive. RMM has a stream pool class to help avoid dynamic stream +creation. However, this is not yet exposed in libcudf, so for the time being, libcudf features +should avoid creating streams (even if it is slightly less efficient). It is a good idea to leave a +`// TODO:` note indicating where using a stream would be beneficial. + +### Thrust Execution Policy + +libcudf uses `rmm::exec_policy_nosync(stream)` for all Thrust algorithm calls. This execution policy +avoids internal stream synchronizations except when required for correctness, such as when an +algorithm returns a value to the host (e.g., `thrust::reduce`). + +Using `nosync` provides significant performance improvements, particularly for small data sizes +where synchronization overhead is proportionally larger. Benchmarks have shown speedups of 10-40% +for many operations on small inputs. + +This policy aligns with libcudf's existing stream semantics: libcudf APIs called on the host +[do not guarantee that the stream is synchronized before returning](#async-apis). +Callers must explicitly synchronize if they need to access results on the host. + +Notes on `nosync`: + +- Algorithms that return values to the host (like `thrust::reduce`) will still synchronize + internally as needed for correctness. +- Ensure that stream-ordered accesses are considered in the context of RAII. Host objects that could + go out of scope (whose lifetime could expire before stream-ordered access) may require an explicit + synchronization before returning or exiting that scope. +- All new code should use `rmm::exec_policy_nosync(stream)` rather than `rmm::exec_policy(stream)`. + If a stream sync is needed, call `stream.synchronize()` explicitly. + +### Device lambdas + +Always declare the return type of an extended `__device__` lambda that is passed to a device +algorithm. A bare `__device__` lambda cannot have its return type queried from host code, and some +host-side APIs (for example CUB device algorithms such as `cub::DeviceSelect::If`, and the transform +iterators) do exactly that, failing to compile with a static assertion in ``. +Not every API requires it, but declaring it uniformly keeps call sites correct as the underlying +APIs evolve. + +Prefer a trailing return type; it is the lightest and most readable. Fall back to +`cuda::proclaim_return_type(lambda)` only when a trailing return type is impractical. + +```c++ +// Fails: return type of a bare __device__ lambda is not visible to host code +cudf::detail::copy_if(begin, end, output, [d] __device__(auto i) { return d[i] > 0; }, stream); + +// Preferred: trailing return type +cudf::detail::copy_if(begin, end, output, [d] __device__(auto i) -> bool { return d[i] > 0; }, stream); + +// Alternative: proclaim_return_type +cudf::detail::copy_if( + begin, end, output, cuda::proclaim_return_type([d] __device__(auto i) { return d[i] > 0; }), stream); +``` + +## Memory Allocation + +Device [memory resources](#rmmdevice_memory_resource) are used in libcudf to abstract and control +how device memory is allocated. + +### Output Memory + +Any libcudf API that allocates memory that is *returned* to a user must accept a +`rmm::device_async_resource_ref` as the last parameter. Inside the API, this memory resource must +be used to allocate any memory for returned objects. It should therefore be passed into functions +whose outputs will be returned. Example: + +```c++ +// Returned `column` contains newly allocated memory, +// therefore the API must accept a memory resource pointer +std::unique_ptr returns_output_memory( + ..., rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); + +// This API does not allocate any new *output* memory, therefore +// a memory resource is unnecessary +void does_not_allocate_output_memory(...); +``` + +This rule automatically applies to all detail APIs that allocate memory. Any detail API may be +called by any public API, and therefore could be allocating memory that is returned to the user. +To support such uses cases, all detail APIs allocating memory resources should accept an `mr` +parameter. Callers are responsible for either passing through a provided `mr` or +`cudf::get_current_device_resource_ref()` as needed. + +### Temporary Memory + +Not all memory allocated within a libcudf API is returned to the caller. Often algorithms must +allocate temporary, scratch memory for intermediate results. Always use the default resource +obtained from `cudf::get_current_device_resource_ref()` for temporary memory allocations. Example: + +```c++ +rmm::device_buffer some_function( + ..., rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()) { + rmm::device_buffer returned_buffer(..., mr); // Returned buffer uses the passed in MR + ... + rmm::device_buffer temporary_buffer(...); // Temporary buffer uses default MR + ... + return returned_buffer; +} +``` + +### Memory Management + +libcudf code generally eschews raw pointers and direct memory allocation. Use RMM classes built to +use memory resources for device memory allocation with automated lifetime management. + +#### rmm::device_buffer +Allocates a specified number of bytes of untyped, uninitialized device memory using a +memory resource. If no `rmm::device_async_resource_ref` is explicitly provided, it uses +`cudf::get_current_device_resource_ref()`. + +`rmm::device_buffer` is movable and copyable on a stream. A copy performs a deep copy of the +`device_buffer`'s device memory on the specified stream, whereas a move moves ownership of the +device memory from one `device_buffer` to another. + +```c++ +// Allocates at least 100 bytes of uninitialized device memory +// using the specified resource and stream +rmm::device_buffer buff(100, stream, mr); +void * raw_data = buff.data(); // Raw pointer to underlying device memory + +// Deep copies `buff` into `copy` on `stream` +rmm::device_buffer copy(buff, stream); + +// Moves contents of `buff` into `moved_to` +rmm::device_buffer moved_to(std::move(buff)); + +custom_memory_resource *mr...; +// Allocates 100 bytes from the custom_memory_resource +rmm::device_buffer custom_buff(100, mr, stream); +``` + +#### cudf::detail::device_scalar +A self-contained device scalar for internal libcudf code that needs a single trivially copyable +value in device memory, such as a reduction result, temporary counter, or kernel status value. + +Use this for internal scalar input/output with device kernels. Public libcudf APIs should use +`cudf::scalar` and derived public scalar classes instead of this detail type. + +It exposes `data()` for kernels and `value()`/`set_value_async()` for stream-ordered host/device +transfers. + +```c++ +// Allocates device memory for a single int using the specified resource and stream +// and initializes the value to 42 +cudf::detail::device_scalar int_scalar{42, stream, mr}; + +// scalar.data() returns pointer to value in device memory +kernel<<<..., stream>>>(int_scalar.data(), ...); + +// value() copies the device value to the host on the specified stream +int host_value = int_scalar.value(stream); +``` + +#### rmm::device_vector + +Allocates a specified number of elements of the specified type. If no initialization value is +provided, all elements are default initialized (this incurs a kernel launch). + +**Note**: We have removed all usage of `rmm::device_vector` and `thrust::device_vector` from +libcudf, and you should not use it in new code in libcudf without careful consideration. Instead, +use `rmm::device_uvector` along with the utility factories in `device_factories.hpp`. These +utilities enable creation of `uvector`s from host-side vectors, or creating zero-initialized +`uvector`s, so that they are as convenient to use as `device_vector`. Avoiding `device_vector` has +a number of benefits, as described in the following section on `rmm::device_uvector`. + +#### rmm::device_uvector + +Similar to a `device_vector`, allocates a contiguous set of elements in device memory but with key +differences: +- As an optimization, elements are uninitialized and no synchronization occurs at construction. +This limits the types `T` to trivially copyable types. +- All operations are stream ordered (i.e., they accept a `cuda::stream_ref` specifying the stream +on which the operation is performed). This improves safety when using non-default streams. +- `device_uvector.hpp` does not include any `__device__` code, unlike `thrust/device_vector.hpp`, + which means `device_uvector`s can be used in `.cpp` files, rather than just in `.cu` files. + +```c++ +cuda_stream s; +// Allocates uninitialized storage for 100 `int32_t` elements on stream `s` using the +// default resource +rmm::device_uvector v(100, s); +// Initializes the elements to 0 +thrust::uninitialized_fill(thrust::cuda::par.on(s.value()), v.begin(), v.end(), int32_t{0}); + +auto mr = new my_custom_resource{...}; +// Allocates uninitialized storage for 100 `int32_t` elements on stream `s` using the resource `mr` +rmm::device_uvector v2{100, s, mr}; +``` + +## Memory Copies + +libcudf code should prefer `cudf::detail::cuda_memcpy_async` and `cudf::detail::memcpy_async` over +direct calls to `cudaMemcpyAsync`. These cudf utilities try to use `cudaMemcpyBatchAsync` on CUDA +13.0+, but not primarily for the "batch" properties: +- `cudaMemcpyBatchAsync` can be lower-overhead, it is actually asynchronous in certain cases where + `cudaMemcpyAsync` cannot be asynchronous +- `cudaMemcpyBatchAsync` may also reduce multi-thread lock contention compared to `cudaMemcpyAsync` + +For host-to-device or device-to-host copies, prefer the typed span-based wrappers: + +```c++ +cudf::detail::cuda_memcpy_async(device_span{dst}, host_span{src}, stream); +cudf::detail::cuda_memcpy_async(host_span{dst}, device_span{src}, stream); +``` + +For device-to-device copies, or when a raw `void*` interface is required, use +`cudf::detail::memcpy_async` (single buffer) or `cudf::detail::memcpy_batch_async` (multiple +buffers) and check errors with `CUDF_CUDA_TRY` at the call site: + +```c++ +// Single buffer copy +CUDF_CUDA_TRY(cudf::detail::memcpy_async(dst, src, size_bytes, stream)); + +// Batch copy of multiple buffers +CUDF_CUDA_TRY(cudf::detail::memcpy_batch_async(dsts, srcs, sizes, count, stream)); +``` + +`memcpy_async` is a thin wrapper around `memcpy_batch_async` with `count = 1`. Prefer +`memcpy_batch_async` directly when copying multiple buffers, as it issues all copies in a single +`cudaMemcpyBatchAsync` call. + +**Important:** Both functions use `cudaMemcpyBatchAsync` with `cudaMemcpySrcAccessOrderStream`, +which defers reading the source buffers until the stream reaches the copies. The **source buffers +must remain valid until the stream has executed the copies**. For device memory this is naturally +satisfied; for host memory the caller must ensure the sources are not freed before the stream is +synchronized. + +When a temporary host buffer is used as the source of an async copy, its destructor must not free +the memory before the stream has consumed it. Buffers allocated with a stream-ordered allocator +(e.g., the pinned memory pool) are safe — deallocation is deferred until the stream catches up. +Buffers from non-stream-ordered allocators (e.g., `new_delete_resource`) require an explicit +`stream.synchronize()` before the buffer is destroyed. Prefer `make_pinned_vector_async` for +temporary host staging buffers to avoid the sync: + +```c++ +// UNSAFE — std::vector uses pageable memory; must synchronize before it goes out of scope +{ + auto staging = std::vector(size); + fill(staging); + cudf::detail::cuda_memcpy_async(d_buf, staging, stream); + stream.synchronize(); // required: pageable allocator is not stream-ordered +} + +// SAFE — stream-ordered pinned buffer, no sync needed +{ + auto staging = cudf::detail::make_pinned_vector_async(size, stream); + fill(staging); + cudf::detail::cuda_memcpy_async(d_buf, staging, stream); +} // deallocation is stream-ordered → no race +``` + +The same stream-safety requirements apply to `memcpy_async` and `memcpy_batch_async`. + +If CUDA memory copy APIs must be called directly, always use `cudaMemcpyDefault` instead of an +explicit host/device copy policy. Copy correctness depends on whether the source and destination +pointers are accessible from the host or device, not where the memory is resident. For example, +pinned host memory may be device-accessible despite residing on the host. `cudaMemcpyDefault` allows +CUDA to infer the valid copy direction from the pointers rather than rejecting such copies based on +an explicit policy. + +## Default Parameters + +While public libcudf APIs are free to include default function parameters, detail functions should +not. Default memory resource parameters make it easy for developers to accidentally allocate memory +using the incorrect resource. Avoiding default memory resources forces developers to consider each +memory allocation carefully. + +While streams are not currently exposed in libcudf's API, we plan to do so eventually. As a result, +the same reasons for memory resources also apply to streams. Public APIs default to using +`cudf::get_default_stream()`. However, including the same default in detail APIs opens the door for +developers to forget to pass in a user-provided stream if one is passed to a public API. Forcing +every detail API call to explicitly pass a stream is intended to prevent such mistakes. + +The memory resources (and eventually, the stream) are the final parameters for essentially all +public APIs. For API consistency, the same is true throughout libcudf's internals. Therefore, a +consequence of not allowing default streams or MRs is that no parameters in detail APIs may have +defaults. + +## NVTX Ranges + +In order to aid in performance optimization and debugging, all compute intensive libcudf functions +should have a corresponding NVTX range. Choose between `CUDF_FUNC_RANGE` or `cudf::scoped_range` +for declaring NVTX ranges in the current scope: +- Use the `CUDF_FUNC_RANGE()` macro if you want to use the name of the function as the name of the +NVTX range +- Use `cudf::scoped_range rng{"custom_name"};` to provide a custom name for the current scope's +NVTX range + +For more information about NVTX, see [here](https://github.com/NVIDIA/NVTX/tree/dev/c). + +## Input/Output Style + +The preferred style for how inputs are passed in and outputs are returned is the following: +- Inputs + - Columns: + - `column_view const&` + - Tables: + - `table_view const&` + - Scalar: + - `scalar const&` + - Everything else: + - Trivial or inexpensively copied types + - Pass by value + - Non-trivial or expensive to copy types + - Pass by `const&` +- In/Outs + - Columns: + - `mutable_column_view&` + - Tables: + - `mutable_table_view&` + - Everything else: + - Pass by via raw pointer +- Outputs + - Outputs should be *returned*, i.e., no output parameters + - Columns: + - `std::unique_ptr` + - Tables: + - `std::unique_ptr
` + - Scalars: + - `std::unique_ptr` + + +### Multiple Return Values + +Sometimes it is necessary for functions to have multiple outputs. There are a few ways this can be +done in C++ (including creating a `struct` for the output). One convenient way to do this is +using `std::tie` and `std::pair`. Note that objects passed to `std::pair` will invoke +either the copy constructor or the move constructor of the object, and it may be preferable to move +non-trivially copyable objects (and required for types with deleted copy constructors, like +`std::unique_ptr`). + +```c++ +std::pair return_two_tables(void){ + cudf::table out0; + cudf::table out1; + ... + // Do stuff with out0, out1 + + // Return a std::pair of the two outputs + return std::pair(std::move(out0), std::move(out1)); +} + +cudf::table out0; +cudf::table out1; +std::tie(out0, out1) = cudf::return_two_outputs(); +``` + +Note: `std::tuple` _could_ be used if not for the fact that Cython does not support +`std::tuple`. Therefore, libcudf APIs must use `std::pair`, and are therefore limited to return +only two objects of different types. Multiple objects of the same type may be returned via a +`std::vector`. + +Alternatively, with C++17 (supported from cudf v0.20), +[structured binding](https://en.cppreference.com/cpp/language/structured_binding) +may be used to disaggregate multiple return values: + +```c++ +auto [out0, out1] = cudf::return_two_outputs(); +``` + +Note that the compiler might not support capturing aliases defined in a structured binding +in a lambda. One may work around this by using a capture with an initializer instead: + +```c++ +auto [out0, out1] = cudf::return_two_outputs(); + +// Direct capture of alias from structured binding might fail with: +// "error: structured binding cannot be captured" +// auto foo = [out0]() {...}; + +// Use an initializing capture: +auto foo = [&out0 = out0] { + // Use out0 to compute something. + // ... +}; +``` + +## Iterator-based interfaces + +Increasingly, libcudf is moving toward internal (`detail`) APIs with iterator parameters rather +than explicit `column`/`table`/`scalar` parameters. As with STL, iterators enable generic +algorithms to be applied to arbitrary containers. A good example of this is `cudf::copy_if_else`. +This function takes two inputs, and a Boolean mask. It copies the corresponding element from the +first or second input depending on whether the mask at that index is `true` or `false`. Implementing +`copy_if_else` for all combinations of `column` and `scalar` parameters is simplified by using +iterators in the `detail` API. + +```c++ +template +std::unique_ptr copy_if_else( + bool nullable, + LeftIter lhs_begin, + LeftIter lhs_end, + RightIter rhs, + FilterFn filter, + ...); +``` + +`LeftIter` and `RightIter` need only implement the necessary interface for an iterator. libcudf +provides a number of iterator types and utilities that are useful with iterator-based APIs from +libcudf as well as Thrust algorithms. Most are defined in `include/detail/iterator.cuh`. + +### Pair iterator + +The pair iterator is used to access elements of nullable columns as a pair containing an element's +value and validity. `cudf::detail::make_pair_iterator` can be used to create a pair iterator from a +`column_device_view` or a `cudf::scalar`. `make_pair_iterator` is not available for +`mutable_column_device_view`. + +### Null-replacement iterator + +This iterator replaces the null/validity value for each element with a specified constant (`true` or +`false`). Created using `cudf::detail::make_null_replacement_iterator`. + +### Validity iterator + +This iterator returns the validity of the underlying element (`true` or `false`). Created using +`cudf::detail::make_validity_iterator`. + +(index-normalizing-iterators)= +### Index-normalizing iterators + +The proliferation of data types supported by libcudf can result in long compile times. One area +where compile time was a problem is in types used to store indices, which can be any integer type. +The "indexalator", or index-normalizing iterator (`include/cudf/detail/indexalator.cuh`), 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 returns a +[`cudf::size_type`](#cudfsize_type) integer. + +Use the `indexalator_factory` to create an appropriate input iterator from a column_view. Example +input iterator usage: + +```c++ +auto begin = indexalator_factory::create_input_iterator(gather_map); +auto end = begin + gather_map.size(); +auto result = detail::gather( source, begin, end, IGNORE, stream, mr ); +``` + +Example output iterator usage: + +```c++ +auto result_itr = indexalator_factory::create_output_iterator(indices->mutable_view()); +thrust::lower_bound(rmm::exec_policy_nosync(stream), + input->begin(), + input->end(), + values->begin(), + values->end(), + result_itr, + cuda::std::less()); +``` + +(offset-normalizing-iterators)= +### Offset-normalizing iterators + +Like the [indexalator](#index-normalizing-iterators), +the "offsetalator", or offset-normalizing iterator (`include/cudf/detail/offsetalator.cuh`), can be +used for offset column types (`INT32` or `INT64` only) without requiring a type-specific instance. +This is helpful when reading or building [strings columns](#strings-columns). +The normalized type is `int64` which means an `input_offsetsalator` will return `int64` type values +for both `INT32` and `INT64` offsets columns. +Likewise, an `output_offselator` can accept `int64` type values to store into either an +`INT32` or `INT64` output offsets column created appropriately. + +Use the `cudf::detail::offsetalator_factory` to create an appropriate input or output iterator from an offsets column_view. +Example input iterator usage: + +```c++ + // convert the sizes to offsets + auto [offsets, char_bytes] = cudf::strings::detail::make_offsets_child_column( + output_sizes.begin(), output_sizes.end(), stream, mr); + auto d_offsets = + cudf::detail::offsetalator_factory::make_input_iterator(offsets->view()); + // use d_offsets to address the output row bytes +``` + +Example output iterator usage: + +```c++ + // create offsets column as either INT32 or INT64 depending on the number of bytes + auto offsets_column = cudf::strings::detail::create_offsets_child_column(total_bytes, + offsets_count, + stream, mr); + auto d_offsets = + cudf::detail::offsetalator_factory::make_output_iterator(offsets_column->mutable_view()); + // write appropriate offset values to d_offsets +``` + +## Namespaces + +### External +All public libcudf APIs should be placed in the `cudf` namespace. Example: + +```c++ +namespace cudf{ + void public_function(...); +} // namespace cudf +``` + +The top-level `cudf` namespace is sufficient for most of the public API. However, to logically +group a broad set of functions, further namespaces may be used. For example, there are numerous +functions that are specific to columns of Strings. These functions reside in the `cudf::strings::` +namespace. Similarly, functionality used exclusively for unit testing is in the `cudf::test::` +namespace. + +The public function is expected to contain a call to `CUDF_FUNC_RANGE()` followed by a call to +a `detail` function with same name and parameters as the public function. +See the [Streams](#streams) section for an example of this pattern. + +### Internal + +Many functions are not meant for public use, so place them in either the `detail` or an *anonymous* +namespace, depending on the situation. + +#### detail namespace + +Functions or objects that will be used across *multiple* translation units (i.e., source files), +should be exposed in an internal header file and placed in the `detail` namespace. Example: + +```c++ +// some_utilities.hpp +namespace cudf{ +namespace detail{ +void reusable_helper_function(...); +} // namespace detail +} // namespace cudf +``` + +#### Anonymous namespace + +Functions or objects that will only be used in a *single* translation unit should be defined in an +*anonymous* namespace in the source file where it is used. Example: + +```c++ +// some_file.cpp +namespace{ +void isolated_helper_function(...); +} // anonymous namespace +``` + +[**Anonymous namespaces should *never* be used in a header file.**](https://cmu-sei.github.io/secure-coding-standards/sei-cert-cpp-coding-standard/rules/declarations-and-initialization-dcl/dcl59-cpp/) + +# Deprecating and Removing Code + +libcudf is constantly evolving to improve performance and better meet our users' needs. As a +result, we occasionally need to break or entirely remove APIs to respond to new and improved +understanding of the functionality we provide. Remaining free to do this is essential to making +libcudf an agile library that can rapidly accommodate our users needs. As a result, we do not +always provide a warning or any lead time prior to releasing breaking changes. On a best effort +basis, the libcudf team will notify users of changes that we expect to have significant or +widespread effects. + +Where possible, indicate pending API removals using the +[deprecated](https://en.cppreference.com/cpp/language/attributes/deprecated) attribute and +document them using Doxygen's +[deprecated](https://www.doxygen.nl/manual/commands.html#cmddeprecated) command prior to removal. +When a replacement API is available for a deprecated API, mention the replacement in both the +deprecation message and the deprecation documentation. Pull requests that introduce deprecations +should be labeled "deprecation" to facilitate discovery and removal in the subsequent release. + +Advertise breaking changes by labeling any pull request that breaks or removes an existing API with +the "breaking" tag. This ensures that the "Breaking" section of the release notes includes a +description of what has broken from the past release. Label pull requests that contain deprecations +with the "non-breaking" tag. + + +(errors)= +# Error Handling + +libcudf follows conventions (and provides utilities) enforcing compile-time and run-time +conditions and detecting and handling CUDA errors. Communication of errors is always via C++ +exceptions. + +## Runtime Conditions + +Use the `CUDF_EXPECTS` macro to enforce runtime conditions necessary for correct execution. + +Example usage: + +```c++ +CUDF_EXPECTS(cudf::have_same_types(lhs, rhs), "Type mismatch", cudf::data_type_error); +``` + +The first argument is the conditional expression expected to resolve to `true` under normal +conditions. The second argument to `CUDF_EXPECTS` is a short description of the error that has +occurred and is used for the exception's `what()` message. If the conditional evaluates to +`false`, then an error has occurred and an instance of the exception class in the third argument +(or the default, `cudf::logic_error`) is thrown. + +The condition (first argument) of `CUDF_EXPECTS` should be a pure predicate that only inspects +state without modifying it. If the result of an operation with a side effect needs to be checked, +capture the result in a variable first: + +```c++ +// WRONG — side effect (initialization) inside the condition: +CUDF_EXPECTS(reader.init(source), "Failed to initialize reader"); + +// RIGHT — capture the result, then check it: +auto const init_ok = reader.init(source); +CUDF_EXPECTS(init_ok, "Failed to initialize reader"); +``` + +There are times where a particular code path, if reached, should indicate an error no matter what. +For example, often the `default` case of a `switch` statement represents an invalid alternative. +Use the `CUDF_FAIL` macro for such errors. This is effectively the same as calling +`CUDF_EXPECTS(false, reason)`. + +Example: + +```c++ +CUDF_FAIL("This code path should not be reached."); +``` + +Prefer `CUDF_EXPECTS` over `if (condition) { CUDF_FAIL(reason); }` when the condition has no side +effects. The `if`/`CUDF_FAIL` pattern is appropriate when cleanup or other actions must be +performed before throwing, or when the condition itself involves side effects that cannot be +separated from the check. + +### CUDA Error Checking + +Use the `CUDF_CUDA_TRY` macro to check for the successful completion of CUDA runtime API functions. +This macro throws a `cudf::cuda_error` exception if the CUDA API return value is not `cudaSuccess`. +The thrown exception includes a description of the CUDA error code in its `what()` message. + +Example: + +```c++ +CUDF_CUDA_TRY( cudaMemcpy(&dst, &src, num_bytes) ); +``` + +## Compile-Time Conditions + +Use `static_assert` to enforce compile-time conditions. For example, + +```c++ +template +void trivial_types_only(T t){ + static_assert(std::is_trivial::value, "This function requires a trivial type."); +... +} +``` + +# Logging + +libcudf includes logging utilities (built on top of [spdlog](https://github.com/gabime/spdlog) +library), which should be used to log important events (e.g. user warnings). This utility can also +be used to log debug information, as long as the correct logging level is used. There are six macros +that should be used for logging at different levels: + +* `CUDF_LOG_TRACE` - verbose debug messages (targeted at developers) +* `CUDF_LOG_DEBUG` - debug messages (targeted at developers) +* `CUDF_LOG_INFO` - information about rare events (e.g. once per run) that occur during normal +execution +* `CUDF_LOG_WARN` - user warnings about potentially unexpected behavior or deprecations +* `CUDF_LOG_ERROR` - recoverable errors +* `CUDF_LOG_CRITICAL` - unrecoverable errors (e.g. memory corruption) + +By default, `TRACE`, `DEBUG` and `INFO` messages are excluded from the log. In addition, in public +builds, the code that logs at `TRACE` and `DEBUG` levels is compiled out. This prevents logging of +potentially sensitive data that might be done for debug purposes. Also, this allows developers to +include expensive computation in the trace/debug logs, as the overhead will not be present in the +public builds. +The minimum enabled logging level is `WARN`, and it can be modified in multiple ways: + +* CMake configuration variable `LIBCUDF_LOGGING_LEVEL` - sets the minimum level of logging that +will be compiled in the build. +Available levels are `TRACE`, `DEBUG`, `INFO`, `WARN`, `ERROR`, `CRITICAL`, and `OFF`. +* Environment variable `LIBCUDF_LOGGING_LEVEL` - sets the minimum logging level during +initialization. If this setting is higher than the compile-time CMake variable, any logging levels +in between the two settings will be excluded from the written log. The available levels are the same +as for the CMake variable. +* Global logger object exposed via `cudf::logger()` - sets the minimum logging level at runtime. +For example, calling `cudf::default_logger().set_level(level_enum::err)`, will exclude any messages that +are not errors or critical errors. This API should not be used within libcudf to manipulate logging, +its purpose is to allow upstream users to configure libcudf logging to fit their application. + +By default, logging messages are output to stderr. +Setting the environment variable `LIBCUDF_DEBUG_LOG_FILE` redirects the log to a file with the +specified path (can be relative to the current directory). +Upstream users can also manipulate `cudf::default_logger().sinks()` to add sinks or divert the log to +standard output. + +# Data Types + +Columns may contain data of a number of types (see `enum class type_id` in `include/cudf/types.hpp`) + + * Numeric data: signed and unsigned integers (8-, 16-, 32-, or 64-bit), floats (32- or 64-bit), and + Booleans (8-bit). + * Timestamp data with resolution of days, seconds, milliseconds, microseconds, or nanoseconds. + * Duration data with resolution of days, seconds, milliseconds, microseconds, or nanoseconds. + * Decimal fixed-point data (32- or 64-bit). + * Strings + * Dictionaries + * Lists of any type + * Structs of columns of any type + +Most algorithms must support columns of any data type. This leads to complexity in the code, and +is one of the primary challenges a libcudf developer faces. Sometimes we develop new algorithms with +gradual support for more data types to make this easier. Typically we start with fixed-width data +types such as numeric types and timestamps/durations, adding support for nested types later. + +Enabling an algorithm differently for different types uses either template specialization or SFINAE, +as discussed in [Specializing Type-Dispatched Code Paths](#specializing-type-dispatched-code-paths). + +## Comparing Data Types + +When comparing the data types of two columns or scalars, do not directly compare +`a.type() == b.type()`. Nested types such as lists of structs of integers will not be handled +properly if only the top level type is compared. Instead, use the `cudf::have_same_types` function. + +# Type Dispatcher + +libcudf stores data (for columns and scalars) "type erased" in `void*` device memory. This +*type-erasure* enables interoperability with other languages and type systems, such as Python and +Java. In order to determine the type, libcudf algorithms must use the run-time information stored +in the column `type()` to reconstruct the data type `T` by casting the `void*` to the appropriate +`T*`. + +This so-called *type dispatch* is pervasive throughout libcudf. The `type_dispatcher` is a +central utility that automates the process of mapping the runtime type information in `data_type` +to a concrete C++ type. + +At a high level, you call the `type_dispatcher` with a `data_type` and a function object (also +known as a *functor*) with an `operator()` template. Based on the value of `data_type::id()`, the +type dispatcher invokes the corresponding instantiation of the `operator()` template. + +This simplified example shows how the value of `data_type::id()` determines which instantiation of +the `F::operator()` template is invoked. + +```c++ +template +void type_dispatcher(data_type t, F f){ + switch(t.id()) + case type_id::INT32: f.template operator()() + case type_id::INT64: f.template operator()() + case type_id::FLOAT: f.template operator()() + ... +} +``` + +The following example shows a function object called `size_of_functor` that returns the size of the +dispatched type. + +```c++ +struct size_of_functor{ + template + int operator()(){ return sizeof(T); } +}; + +cudf::type_dispatcher(data_type{type_id::INT8}, size_of_functor{}); // returns 1 +cudf::type_dispatcher(data_type{type_id::INT32}, size_of_functor{}); // returns 4 +cudf::type_dispatcher(data_type{type_id::FLOAT64}, size_of_functor{}); // returns 8 +``` + +By default, `type_dispatcher` uses `cudf::type_to_id` to provide the mapping of `cudf::type_id` +to dispatched C++ types. However, this mapping may be customized by explicitly specifying a +user-defined trait for the `IdTypeMap`. For example, to always dispatch `int32_t` for all values of +`cudf::type_id`: + +```c++ +template struct always_int{ using type = int32_t; } + +// This will always invoke `operator()` +cudf::type_dispatcher(data_type, f); +``` + +## Avoid Multiple Type Dispatch + +Avoid multiple type-dispatch if possible. The compiler creates a code path for every type +dispatched, so a second-level type dispatch results in quadratic growth in compilation time and +object code size. As a large library with many types and functions, we are constantly working to +reduce compilation time and code size. + +## Specializing Type-Dispatched Code Paths + +It is often necessary to customize the dispatched `operator()` for different types. This can be +done in several ways. + +The first method is to use explicit, full template specialization. This is useful for specializing +behavior for single types. The following example function object prints `"int32_t"` or `"double"` +when invoked with either of those types, or `"unhandled type"` otherwise. + +```c++ +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"; } +``` + +The second method is to use [SFINAE](https://en.cppreference.com/cpp/language/sfinae) with +`std::enable_if_t`. This is useful to partially specialize for a set of types with a common trait. +The following example functor prints `integral` or `floating point` for integral or floating point +types, respectively. + +```c++ +struct integral_or_floating_point { +template ::value and + not std::is_floating_point::value>* = nullptr> +void operator()() { std::cout << "neither integral nor floating point\n"; } + +template ::value>* = nullptr> +void operator()() { std::cout << "integral\n"; } + +template < typename ColumnType, + std::enable_if_t::value>* = nullptr> +void operator()() { std::cout << "floating point\n"; } +}; +``` + +For more info on SFINAE with `std::enable_if`, [see this post](https://eli.thegreenplace.net/2014/sfinae-and-enable_if). + +There are a number of traits defined in `include/cudf/utilities/traits.hpp` that are useful for +partial specialization of dispatched function objects. For example `is_numeric()` can be used to +specialize for any numeric type. + +# Variable-Size and Nested Data Types + +libcudf supports a number of variable-size and nested data types, including strings, lists, and +structs. + + * `string`: Simply a character string, but a column of strings may have a different-length string + in each row. + * `list`: A list of elements of any type, so a column of lists of integers has rows with a list of + integers, possibly of a different length, in each row. + * `struct`: In a column of structs, each row is a structure comprising one or more fields. These + fields are stored in structure-of-arrays format, so that the column of structs has a nested + column for each field of the structure. + +As the heading implies, list and struct columns may be nested arbitrarily. One may create a column +of lists of structs, where the fields of the struct may be of any type, including strings, lists and +structs. Thinking about deeply nested data types can be confusing for column-based data, even with +experience. Therefore it is important to carefully write algorithms, and to test and document them +well. + +## List columns + +In order to represent variable-width elements, libcudf columns contain a vector of child columns. +For list columns, the parent column's type is `LIST` and contains no data, but its size represents +the number of lists in the column, and its null mask represents the validity of each list element. +The parent has two children. + +1. A non-nullable column of [`size_type`](#cudfsize_type) elements that indicates the offset to the + beginning of each list in a dense column of elements. +2. A column containing the actual data and optional null mask for all elements of all the lists + packed together. + +With this representation, `data[offsets[i]]` is the first element of list `i`, and the size of list +`i` is given by `offsets[i+1] - offsets[i]`. + +Note that the data may be of any type, and therefore the data column may itself be a nested column +of any type. Note also that not only is each list nullable (using the null mask of the parent), but +each list element may be nullable. So you may have a lists column with null row 3, and also null +element 2 of row 4. + +The underlying data for a lists column is always bundled into a single leaf column at the very +bottom of the hierarchy (ignoring structs, which conceptually "reset" the root of the hierarchy), +regardless of the level of nesting. So a `List>>>` column has a single `int` +column at the very bottom. The following is a visual representation of this. + +``` +lists_column = { {{{1, 2}, {3, 4}}, NULL}, {{{10, 20}, {30, 40}}, {{50, 60, 70}, {0}}} } + + List>> (2 rows): + Length : 2 + Offsets : 0, 2, 4 + Children : + List>: + Length : 4 + Offsets : 0, 2, 2, 4, 6 + Null count: 1 + 1101 + Children : + List: + Length : 6 + Offsets : 0, 2, 4, 6, 8, 11, 12 + Children : + Column of ints + 1, 2, 3, 4, 10, 20, 30, 40, 50, 60, 70, 0 +``` + +This is related to [Arrow's "Variable-Size List" memory layout](https://arrow.apache.org/docs/format/Columnar.html?highlight=nested%20types#physical-memory-layout). + +## Strings columns + +Strings are represented as a column with a data device buffer and a child offsets column. +The parent column's type is `STRING` and its data holds all the characters across all the strings packed together +but its size represents the number of strings in the column and its null mask represents the +validity of each string. + +The strings column contains a single, non-nullable child column +of offset elements that indicates the byte position offset to the beginning of each +string in the dense data buffer of all characters. With this representation, `data[offsets[i]]` is the +first character of string `i`, and the size of string `i` is given by `offsets[i+1] - offsets[i]`. +The following image shows an example of this compound column representation of strings. + +![strings](strings.png) + +The type of the offsets column is either `INT32` or `INT64` depending on the number of bytes in the data buffer. +See [`cudf::strings_view`](#cudfstrings_column_view-and-cudfstring_view) for more information on processing individual string rows. + +## Structs columns + +A struct is a nested data type with a set of child columns each representing an individual field +of a logical struct. Field names are not represented. + +A structs column with `N` fields has `N` children. Each child is a column storing all the data +of a single field packed column-wise, with an optional null mask. The parent column's type is +`STRUCT` and contains no data, its size represents the number of struct rows in the column, and its +null mask represents the validity of each struct element. + +With this representation, `child[0][10]` is row 10 of the first field of the struct, `child[1][42]` +is row 42 of the second field of the struct. + +Notice that in addition to the struct column's null mask, each struct field column has its own +optional null mask. A struct field's validity can vary independently from the corresponding struct +row. For instance, a non-null struct row might have a null field. However, the fields of a null +struct row are deemed to be null as well. For example, consider a struct column of type +`STRUCT`. If the contents are `[ {1.0, 2}, {4.0, 5}, null, {8.0, null} ]`, the +struct column's layout is as follows. (Note that null masks should be read from right to left.) + +``` +{ + type = STRUCT + null_mask = [1, 1, 0, 1] + null_count = 1 + children = { + { + type = FLOAT32 + data = [1.0, 4.0, X, 8.0] + null_mask = [ 1, 1, 0, 1] + null_count = 1 + }, + { + type = INT32 + data = [2, 5, X, X] + null_mask = [1, 1, 0, 0] + null_count = 2 + } + } +} +``` + +The last struct row (index 3) is not null, but has a null value in the `INT32` field. Also, row 2 of +the struct column is null, making its corresponding fields also null. Therefore, bit 2 is unset in +the null masks of both struct fields. + +## Dictionary columns + +Dictionaries provide an efficient way to represent low-cardinality data by storing a single copy +of each value. A dictionary comprises a column of keys and a column containing an index into +the keys column for each row of the parent column. The keys column may have any fixed-width data_type +or STRING data_type. The indices represent the corresponding positions of each +element's value in the keys. The indices child column can have any signed integer type +(`INT8`, `INT16`, `INT32`, or `INT64`). + +The `cudf::dictionary::encode()` API is non-deterministic. That is, calling this encode twice on the same +input column will produce equivalent dictionary columns but the keys may be in a different order +and therefore the indices will not match as well. Using `cudf::dictionary::decode()` on both dictionary +columns should produce the same result. + +The libcudf APIs also accept dictionary columns with non-unique keys. +However, output dictionary columns will generally contain unique keys in an unspecified order. +The exceptions are `cudf::make_dictionary_column()`, which accepts keys and indices without +changing them, and `cudf::dictionary::set_keys()`, which strictly honors the given keys +(both order and duplicates). + +## Nested column challenges + +The first challenge with nested columns is that it is effectively impossible to do any operation +that modifies the length of any string or list in place. For example, consider trying to append the +character `'a'` to the end of each string. This requires dynamically resizing the characters column +to allow inserting `'a'` at the end of each string, and then modifying the offsets column to +indicate the new size of each element. As a result, every operation that can modify the strings or +lists in a column must be done out-of-place. + +The second challenge is that in an out-of-place operation on a strings column, unlike with fixed- +width elements, the size of the output cannot be known *a priori*. For example, consider scattering +into a column of strings: + + destination: {"this", "is", "a", "column", "of", "strings"} + scatter_map: {1, 3, 5} + scatter_values: {"red", "green", "blue"} + + result: {"this", "red", "a", "green", "of", "blue"} + +In this example, the strings "red", "green", and "blue" will respectively be scattered into +positions `1`, `3`, and `5` of `destination`. Recall from above that this operation cannot be done +in place, therefore `result` will be generated by selectively copying strings from `destination` and +`scatter_values`. Notice that `result`'s child column of characters requires storage for `19` +characters. However, there is no way to know ahead of time that `result` will require `19` +characters. Therefore, most operations that produce a new output column of strings use a two-phase +approach: + +1. Determine the number and size of each string in the result. This amounts to materializing the + output offsets column. +2. Allocate sufficient storage for all of the output characters and materialize each output string. + +In scatter, the first phase consists of using the `scatter_map` to determine whether string `i` in +the output will come from `destination` or from `scatter_values` and use the corresponding size(s) +to materialize the offsets column and determine the size of the output. Then, in the second phase, +sufficient storage is allocated for the output's characters, and then the characters are filled +with the corresponding strings from either `destination` or `scatter_values`. + +## Nested Type Views + +libcudf provides view types for nested column types as well as for the data elements within them. + +(cudfstrings_column_view-and-cudfstring_view)= +### cudf::strings_column_view and cudf::string_view + +A `cudf::strings_column_view` wraps a strings column and contains a parent +`cudf::column_view` as a view of the strings column and an offsets `cudf::column_view` +which is a child of the parent. +The parent view contains the offset, size, and validity mask for the strings column. +The offsets view is non-nullable with `offset()==0` and its own size. +Since the offset column type can be either `INT32` or `INT64` it is useful to use the +offset normalizing iterators [offsetalator](#offset-normalizing-iterators) to access individual offset values. + +A `cudf::string_view` is a view of a single string and therefore +is the data type of a `cudf::column` of type `STRING` just like `int32_t` is the +data type for a `cudf::column` of type `INT32`. As its name implies, this is a +read-only object instance that points to device memory inside the strings column. +Its lifespan is the same (or less) as the column it views. +An individual strings column row and a `cudf::string_view` is limited to [`size_type`](#cudfsize_type) bytes. + +Use the `column_device_view::element` method to access an individual row element. Like any other +column, do not call `element()` on a row that is null. + +```c++ + cudf::strings_column_view scv; + auto d_strings = cudf::column_device_view::create(scv.parent(), stream); + ... + if( d_strings.is_valid(row_index) ) { + string_view d_str = d_strings.element(row_index); + ... + } +``` + +A null string is not the same as an empty string. Use the `cudf::string_scalar` class if you need an +instance of a class object to represent a null string. + +The `cudf::string_view` contains comparison operators `<,>,==,<=,>=` that can be used in many cudf +functions like `sort` without string-specific code. The data for a `cudf::string_view` instance is +required to be [UTF-8](#utf-8) and all operators and methods expect this encoding. Unless documented +otherwise, position and length parameters are specified in characters and not bytes. The class also +includes a `cudf::string_view::const_iterator` which can be used to navigate through individual characters +within the string. + +`cudf::type_dispatcher` dispatches to the `cudf::string_view` data type when invoked on a `STRING` column. + +(utf-8)= +#### UTF-8 + +The libcudf strings column only supports UTF-8 encoding for strings data. +[UTF-8](https://en.wikipedia.org/wiki/UTF-8) is a variable-length character encoding wherein each +character can be 1-4 bytes. This means the length of a string is not the same as its size in bytes. +For this reason, it is recommended to use the `cudf::string_view` class to access these characters for +most operations. + +The `cudf/strings/detail/utf8.hpp` header also includes some utility methods for reading and writing +(`to_char_utf8/from_char_utf8`) individual UTF-8 characters to/from byte arrays. + +### cudf::lists_column_view and cudf::lists_view + +`cudf::lists_column_view` is a view of a lists column. `cudf::list_view` is a view of a single list, +and therefore `cudf::list_view` is the data type of a `cudf::column` of type `LIST`. + +`cudf::type_dispatcher` dispatches to the `list_view` data type when invoked on a `LIST` column. + +### cudf::structs_column_view and cudf::struct_view + +`cudf::structs_column_view` is a view of a structs column. `cudf::struct_view` is a view of a single +struct, and therefore `cudf::struct_view` is the data type of a `cudf::column` of type `STRUCT`. + +`cudf::type_dispatcher` dispatches to the `struct_view` data type when invoked on a `STRUCT` column. + +# Empty Columns + +The libcudf columns support empty, typed content. These columns have no data and no validity mask. +Empty strings or lists columns may or may not contain a child offsets column. +It is undefined behavior (UB) to access the offsets child of an empty strings or lists column. +Nested columns like lists and structs may require other children columns to provide the +nested structure of the empty types. + +Use `cudf::make_empty_column()` to create fixed-width and strings columns. +Use `cudf::empty_like()` to create an empty column from an existing `cudf::column_view`. + +# cuIO: file reading and writing + +cuIO is a component of libcudf that provides GPU-accelerated reading and writing of data file +formats commonly used in data analytics, including CSV, Parquet, ORC, Avro, and JSON_Lines. + +// TODO: add more detail and move to a separate file. + +# Debugging Tips + +Here are some tools that can help with debugging libcudf (besides printf of course): +1. `cuda-gdb`\ + Follow the instructions in the [Contributor to cuDF guide](https://github.com/rapidsai/cudf/blob/main/CONTRIBUTING.md#debugging-cudf) to build + and run libcudf with debug symbols. +2. `compute-sanitizer`\ + The [CUDA Compute Sanitizer](https://docs.nvidia.com/compute-sanitizer/ComputeSanitizer/index.html) + tool can be used to locate many CUDA reported errors by providing a call stack + close to where the error occurs even with a non-debug build. The sanitizer includes various + tools including `memcheck`, `racecheck`, and `initcheck` as well as others. + The `racecheck` and `initcheck` have been known to produce false positives. +3. `cudf::test::print()`\ + The `print()` utility can be called within a gtest to output the data in a `cudf::column_view`. + More information is available in the {ref}`Testing Guide ` +4. GCC Address Sanitizer\ + The GCC ASAN can also be used by adding the `-fsanitize=address` compiler flag. + There is a compatibility issue with the CUDA runtime that can be worked around by setting + environment variable `ASAN_OPTIONS=protect_shadow_gap=0` before running the executable. + Note that the CUDA `compute-sanitizer` can also be used with GCC ASAN by setting the + environment variable `ASAN_OPTIONS=protect_shadow_gap=0,alloc_dealloc_mismatch=0`. diff --git a/docs/cudf/source/libcudf/developer_guide/DOCUMENTATION.md b/docs/cudf/source/libcudf/developer_guide/DOCUMENTATION.md new file mode 100644 index 000000000000..c527c972e62f --- /dev/null +++ b/docs/cudf/source/libcudf/developer_guide/DOCUMENTATION.md @@ -0,0 +1,446 @@ +# libcudf C++ Documentation Guide + +These guidelines apply to documenting all libcudf C++ source files using doxygen style formatting although only public APIs and classes are actually [published](https://docs.rapids.ai/api/libcudf/stable/index.html). + +## Copyright License + +The copyright comment is included here but may also be mentioned in a coding guideline document as well. +The following is the license header comment that should appear at the beginning of every C++ source file. + + /* + * SPDX-FileCopyrightText: Copyright (c) 2021-2022, NVIDIA CORPORATION. + * SPDX-License-Identifier: Apache-2.0 + */ + +The comment should start with `/*` and not `/**` so it is not processed by doxygen. + +Also, here are the rules for the copyright year. + +- A new file should have the year in which it was created +- A modified file should span the year it was created and the year it was modified (e.g. `2019-2021`) + +Changing the copyright year may not be necessary if no content has changed (e.g. reformatting only). + +## Doxygen + +The [doxygen tool](https://www.doxygen.nl/manual/index.html) is used to generate HTML pages from the C++ comments in the source code. +Doxygen recognizes and parses block comments and performs specialized output formatting when it encounters [doxygen commands](https://www.doxygen.nl/manual/commands.html). + +There are almost 200 commands (also called tags in this document) that doxygen recognizes in comment blocks. +This document provides guidance on which commands/tags to use and how to use them in the libcudf C++ source code. + +The doxygen process can be customized using options in the [Doxyfile](https://github.com/rapidsai/cudf/blob/main/cpp/doxygen/Doxyfile). + +Here are some of the custom options in the Doxyfile for libcudf. +| Option | Setting | Description | +| ------ | ------- | ----------- | +| PROJECT_NAME | libcudf | Title used on the main page | +| PROJECT_NUMBER | 22.02.00 | Version number | +| EXTENSION_MAPPING | cu=C++ cuh=C++ | Process `cu` and `cuh` as C++ | +| INPUT | main_page.md regex.md unicode.md ../include | Embedded markdown files and source code directories to process | +| FILE_PATTERNS | *.cpp *.hpp *.h *.c *.cu *.cuh | File extensions to process | + +## Block Comments + +Use the following style for block comments describing functions, classes and other types, groups, and files. + + /** + * description text and + * doxygen tags go here + */ + +Doxygen comment blocks start with `/**` and end with `*/` only, and with nothing else on those lines. +Do not add dashes `-----` or extra asterisks `*****` to the first and last lines of a doxygen block. +The block must be placed immediately before the source code line to which it refers. +The block may be indented to line up vertically with the item it documents as appropriate. +See the [Example](#the-example) section below. + +Each line in the comment block between the `/**` and `*/` lines should start with a space followed by an asterisk. +Any text on these lines, including tag declarations, should start after a single space after the asterisk. + +## Tag/Command names + +Use @ to prefix doxygen commands (e.g. \@brief, \@code, etc.) + +## Markdown + +The doxygen tool supports a limited set of markdown format in the comment block including links, tables, lists, etc. +In some cases a trade-off may be required for readability in the source text file versus the readability in the doxygen formatted web pages. +For example, there are some limitations on readability with '%' character and pipe character '|' within a markdown table. + +Avoid using direct HTML tags. +Although doxygen supports markdown and markdown supports HTML tags, the HTML support for doxygen's markdown is also limited. + +## The Example + +The following example covers most of the doxygen block comment and tag styles +for documenting C++ code in libcudf. + + /** + * @file source_file.cpp + * @brief Description of source file contents + * + * Longer description of the source file contents. + */ + + /** + * @brief One line description of the class + * + * @ingroup optional_predefined_group_id + * + * Longer, more detailed description of the class. + * + * @tparam T Short description of each template parameter + * @tparam U Short description of each template parameter + */ + template + class example_class { + + void get_my_int(); ///< Simple members can be documented like this + void set_my_int( int value ); ///< Try to use descriptive member names + + /** + * @brief Short, one line description of the member function + * + * A more detailed description of what this function does and what + * its logic does. + * + * @code + * example_class inst; + * inst.set_my_int(5); + * int output = inst.complicated_function(1,dptr,fptr); + * @endcode + * + * @param[in] first This parameter is an input parameter to the function + * @param[in,out] second This parameter is used both as an input and output + * @param[out] third This parameter is an output of the function + * + * @return The result of the complex function + */ + T complicated_function(int first, double* second, float* third) + { + // Do not use doxygen-style block comments + // for code logic documentation. + } + + private: + int my_int; ///< An example private member variable + }; + + /** + * @brief Short, one line description of this free function + * + * @ingroup optional_predefined_group_id + * + * A detailed description must start after a blank line. + * + * @code + * template + * struct myfunctor { + * bool operator()(T input) { return input % 2 > 0; } + * }; + * free_function(myfunctor{},12); + * @endcode + * + * @throw cudf::logic_error if `input_argument` is negative or zero + * + * @tparam functor_type The type of the functor + * @tparam input_type The datatype of the input argument + * + * @param[in] functor The functor to be called on the input argument + * @param[in] input_argument The input argument passed into the functor + * @return The result of calling the functor on the input argument + */ + template + bool free_function(functor_type functor, input_type input_argument) + { + CUDF_EXPECTS( input_argument > 0, "input_argument must be positive"); + return functor(input_argument); + } + + /** + * @brief Short, one line description + * + * @ingroup optional_predefined_group_id + * + * Optional, longer description. + */ + enum class example_enum { + first_enum, ///< Description of the first enum + second_enum, ///< Description of the second enum + third_enum ///< Description of the third enum + }; + +## Descriptions + +The comment description should clearly detail how the output(s) are created from any inputs. +Include any performance and any boundary considerations. +Also include any limits on parameter values and if any default values are declared. +Don't forget to specify how nulls are handled or produced. +Also, try to include a short [example](#inline-examples) if possible. + +### @brief + +The [\@brief](https://www.doxygen.nl/manual/commands.html#cmdbrief) text should be a short, one line description. +Doxygen does not provide much space to show this text in the output pages. +Always follow the \@brief line with a blank comment line. Normally this is like a title and not sentence +and therefore does not need a period. Only use a period if it is a sentence. + +The longer description is the rest of the comment text that is not tagged with any doxygen command. + + /** + * @brief Short description or title + * + * Long description. + * + +### \@copydoc + +Documentation for declarations in headers should be clear and complete. +You can use the [\@copydoc](https://www.doxygen.nl/manual/commands.html#cmdcopydoc) tag to avoid duplicating the comment block for a function definition. + + /** + * @copydoc complicated_function(int,double*,float*) + * + * Any extra documentation. + */ + +Also, \@copydoc is useful when documenting a `detail` function that differs only by the `stream` parameter. + + /** + * @copydoc cudf::segmented_count_set_bits(bitmask_type const*,std::vector const&) + * + * @param[in] stream Optional CUDA stream on which to execute kernels + */ + std::vector segmented_count_set_bits(bitmask_type const* bitmask, + std::vector const& indices, + cuda::stream_ref stream = cudf::get_default_stream()); + +Note, you must specify the whole signature of the function, including optional parameters, so that doxygen will be able to locate it. + +### Function parameters + +The following tags should appear near the end of function comment block in the order specified here: + +| Command | Description | +| ------- | ----------- | +| [\@throw](#throw) | Specify the conditions in which the function may throw an exception | +| [\@tparam](#tparam) | Description for each template parameter | +| [\@param](#param) | Description for each function parameter | +| [\@return](#return) | Short description of object or value returned | + +(throw)= +#### \@throw + +Add an [\@throw](https://www.doxygen.nl/manual/commands.html#cmdthrow) comment line in the doxygen block for each exception that the function may throw. +You only need to include exceptions thrown by the function itself. +If the function calls another function that may throw an exception, you do not need to document those exceptions here. + +Include the name of the exception without backtick marks so doxygen can add reference links correctly. + + * + * @throw cudf::logic_error if `input_argument` is negative or zero + * + +Using \@throws is also acceptable but VS Code and other tools only do syntax highlighting on \@throw. + +(tparam)= +#### @tparam + +Add a [\@tparam](https://www.doxygen.nl/manual/commands.html#cmdtparam) comment line for each template parameter declared by this function. +The name of the parameter specified after the doxygen tag must match exactly to the template parameter name. + + * + * @tparam functor_type The type of the functor + * @tparam input_type The datatype of the input argument + * + +The definition should detail the requirements of the parameter. +For example, if the template is for a functor or predicate, then describe the expected input types and output. + +(param)= +#### @param + +Add a [\@param](https://www.doxygen.nl/manual/commands.html#cmdparam) comment line for each function parameter passed to this function. +The name of the parameter specified after the doxygen tag must match the function's parameter name. +Also include append `[in]`, `[out]` or `[in,out]` to the `@param` if it is not clear from the declaration and the parameter name itself. + + * + * @param[in] first This parameter is an input parameter to the function + * @param[in,out] second This parameter is used both as an input and output + * @param[out] third This parameter is an output of the function + * + +It is also recommended to vertically aligning the 3 columns of text if possible to make it easier to read in a source code editor. +Finally, the description is normally like a title and only needs a period if it is a sentence. + +(return)= +#### @return + +Add a single [\@return](https://www.doxygen.nl/manual/commands.html#cmdreturn) comment line at the end of the comment block if the function returns an object or value. +Include a brief description of what is returned. + + /** + * ... + * + * @return A new column of type INT32 and no nulls + */ + +Do not include the type of the object returned with the `@return` comment. + +(inline-examples)= +### Inline Examples + +It is usually helpful to include a source code example inside your comment block when documenting a function or other declaration. +Use the [\@code](https://www.doxygen.nl/manual/commands.html#cmdcode) and [\@endcode](https://www.doxygen.nl/manual/commands.html#cmdendcode) pair to include inline examples. + +Doxygen supports syntax highlighting for C++ and several other programming languages (e.g. Python, Java). +By default, the \@code tag uses syntax highlighting based on the source code in which it is found. + + * + * @code + * auto result = cudf::make_column( ); + * @endcode + * + +You can specify a different language by indicating the file extension in the tag: + + * + * @code{.py} + * import cudf + * s = cudf.Series([1,2,3]) + * @endcode + * + +If you wish to use pseudocode in your example, use the following: + + * + * Sometimes pseudocode is clearer. + * @code{.pseudo} + * s = int column of [ 1, 2, null, 4 ] + * r = fill( s, [1, 2], 0 ) + * r is now [ 1, 0, 0, 4 ] + * @endcode + * + +When writing example snippets, using fully qualified class names allows doxygen to add reference links to the example. + + * + * @code + * auto result1 = make_column( ); // reference link will not be created + * auto result2 = cudf::make_column( ); // reference link will be created + * @endcode + * + +Although using 3 backtick marks \`\`\` for example blocks will work too, they do not stand out as well in VS Code and other source editors. + +Do not use the `@example` tag in the comments for a declaration, or doxygen will interpret the entire source file as example source code. +The source file is then published under a separate _Examples_ page in the output. + +### Deprecations + +Add a single [\@deprecated](https://www.doxygen.nl/manual/commands.html#cmddeprecated) comment line +to comment blocks for APIs that will be removed in future releases. Mention alternative / +replacement APIs in the deprecation comment. + + /** + * ... + * + * @deprecated This function is deprecated. Use another new function instead. + */ + +## Namespaces + +Doxygen output includes a _Namespaces_ page that shows all the namespaces declared with comment blocks in the processed files. +Here is an example of a doxygen description comment for a namespace declaration. + + /** + * @brief cuDF interfaces + * + * This is the top-level namespace which contains all cuDF functions and types. + */ + namespace CUDF_EXPORT cudf { + +A description comment should be included only once for each unique namespace declaration. +Otherwise, if more than one description is found, doxygen aggregates the descriptions in an arbitrary order in the output pages. + +If you introduce a new namespace, provide a description block for only one declaration and not for every occurrence. + +## Groups/Modules + +Grouping declarations into modules helps users to find APIs in the doxygen pages. +Generally, common functions are already grouped logically into header files but doxygen does not automatically group them this way in its output. +The doxygen output includes a _Modules_ page that organizes items into groups specified using the [Grouping doxygen commands](https://www.doxygen.nl/manual/grouping.html). +These commands can group common functions across header files, source files, and even namespaces. +Groups can also be nested by defining new groups within existing groups. + +For libcudf, all the group hierarchy is defined in the [doxygen_groups.h](https://github.com/rapidsai/cudf/blob/main/cpp/include/doxygen_groups.h) header file. +The [doxygen_groups.h](https://github.com/rapidsai/cudf/blob/main/cpp/include/doxygen_groups.h) file does not need to be included in any other source file, because the definitions in this file are used only by the doxygen tool to generate groups in the _Modules_ page. +Modify this file only to add or update groups. +The existing groups have been carefully structured and named, so new groups should be added thoughtfully. + +When creating a new API, specify its group using the [\@ingroup](https://www.doxygen.nl/manual/commands.html#cmdingroup) tag and the group reference id from the [doxygen_groups.h](https://github.com/rapidsai/cudf/blob/main/cpp/include/doxygen_groups.h) file. + + namespace CUDF_EXPORT cudf { + + /** + * @brief ... + * + * @ingroup transformation_fill + * + * @param ... + * @return ... + */ + std::unique_ptr fill(table_view const& input,...); + + } // namespace cudf + +You can also use the \@addtogroup with a `@{ ... @}` pair to automatically include doxygen comment blocks as part of a group. + + namespace CUDF_EXPORT cudf { + /** + * @addtogroup transformation_fill + * @{ + */ + + /** + * @brief ... + * + * @param ... + * @return ... + */ + std::unique_ptr fill(table_view const& input,...); + + /** @} */ + } // namespace cudf + +This just saves adding \@ingroup to individual doxygen comment blocks within a file. +Make sure a blank line is included after the \@addtogroup command block so doxygen knows it does not apply to whatever follows in the source code. +Note that doxygen will not assign groups to items if the \@addtogroup with `@{ ... @}` pair includes a namespace declaration. +So include the `@addtogroup` and `@{ ... @}` between the namespace declaration braces as shown in the example above. + +Summary of groups tags +| Tag/Command | Where to use | +| ----------- | ------------ | +| `@defgroup` | For use only in [doxygen_groups.h](https://github.com/rapidsai/cudf/blob/main/cpp/include/doxygen_groups.h) and should include the group's title. | +| `@ingroup` | Use inside individual doxygen block comments for declaration statements in a header file. | +| `@addtogroup` | Use instead of `@ingroup` for multiple declarations in the same file within a namespace declaration. Do not specify a group title. | +| `@{ ... @}` | Use only with `@addtogroup`. | + +## Build Doxygen Output + +We recommend installing Doxygen using conda (`conda install doxygen`) or a Linux package manager (`sudo apt install doxygen`). +Alternatively you can [build and install doxygen from source](https://www.doxygen.nl/manual/install.html). + +To build the libcudf HTML documentation simply run the `doxygen` command from the `cpp/doxygen` directory containing the `Doxyfile`. +The libcudf documentation can also be built using `cmake --build . --target docs_cudf` from the cmake build directory (e.g. `cpp/build`). +Doxygen reads and processes all appropriate source files under the `cpp/include/` directory. +The output is generated in the `cpp/doxygen/html/` directory. +You can load the local `index.html` file generated there into any web browser to view the result. + +To view docs built on a remote server, you can run a simple HTTP server using Python: `cd html && python -m http.server`. +Then open `:8000` in your local web browser, inserting the IP address of the machine on which you ran the HTTP server. + +The doxygen output is intended for building documentation only for the public APIs and classes. +For example, the output should not include documentation for `detail` or `/src` files, and these directories are excluded in the `Doxyfile` configuration. +When published by the build/CI system, the doxygen output will appear on our external [RAPIDS web site](https://docs.rapids.ai/api/libcudf/stable/index.html). diff --git a/docs/cudf/source/libcudf/developer_guide/PROFILING.md b/docs/cudf/source/libcudf/developer_guide/PROFILING.md new file mode 100644 index 000000000000..589a5829b53b --- /dev/null +++ b/docs/cudf/source/libcudf/developer_guide/PROFILING.md @@ -0,0 +1,46 @@ +# Profiling libcudf + +Profiling is essential for understanding performance characteristics and identifying bottlenecks in libcudf. This guide covers GPU profiling using NVIDIA Nsight Systems. + +## NVIDIA Nsight Systems + +[NVIDIA Nsight Systems](https://developer.nvidia.com/nsight-systems) is a system-wide performance analysis tool that provides detailed timeline views of CPU and GPU activity. +It's the recommended tool for profiling CUDA applications and understanding kernel execution, memory transfers, and API calls. + +### Installation + +Nsight Systems is included with the CUDA Toolkit, or can be downloaded from https://developer.nvidia.com/nsight-systems. The command-line tool is `nsys`. Verify installation: + +```bash +nsys --version +``` + +### Recommended Profile Command + +When profiling cuDF workloads, use the following flags: + +```bash +nsys profile --trace=nvtx,cuda,osrt --cuda-memory-usage=true --gpu-metrics-devices=0 --nvtx-domain-exclude=CCCL python script.py +``` + +**Options explained:** +- `--trace=nvtx,cuda,osrt`: Trace NVTX ranges, CUDA API calls, and OS runtime libraries +- `--cuda-memory-usage=true`: Track CUDA memory allocation and usage +- `--gpu-metrics-devices=0`: Collect GPU metrics from device 0 +- `--nvtx-domain-exclude=CCCL`: Exclude verbose CCCL (CUDA C++ Core Libraries) NVTX ranges + +### Profiling Specific GPUs + +When working with multi-GPU systems, you may want to profile a specific GPU. +To profile GPUs other than device 0, use both `--gpu-metrics-devices=N` and `--env-var CUDA_VISIBLE_DEVICES=N` to ensure the application and profiler target the same device. + +For example, modify the flags like this for profiling GPU 4: + +```bash +nsys profile --trace=nvtx,cuda,osrt --cuda-memory-usage=true --gpu-metrics-devices=4 --env-var CUDA_VISIBLE_DEVICES=4 python script.py +``` + +### Analyzing Results + +After profiling, open the `.nsys-rep` file in the Nsight Systems GUI to analyze CPU and GPU activity over time. +The interface shows individual kernel launches and durations, memory allocations and transfers, and metrics like memory bandwidth utilization. diff --git a/docs/cudf/source/libcudf/developer_guide/TESTING.md b/docs/cudf/source/libcudf/developer_guide/TESTING.md new file mode 100644 index 000000000000..19f6a00c26e7 --- /dev/null +++ b/docs/cudf/source/libcudf/developer_guide/TESTING.md @@ -0,0 +1,536 @@ +# Unit Testing in libcudf + +Unit tests in libcudf are written using +[Google Test](https://github.com/google/googletest/blob/master/docs/primer.md). + +**Important:** Instead of including `gtest/gtest.h` directly, use +`#include `. + +Also, write test code in the global namespace. That is, +do not write test code in the `cudf` or the `cudf::test` namespace or their +sub-namespaces. +Likewise, do not use `using namespace cudf;` or `using namespace cudf::test;` +in the global namespace. + + +## Best Practices: What Should We Test? + +In general we should test to make sure all code paths are covered. This is not always easy or +possible. But generally this means we test all supported combinations of algorithms and data types, +and all operators supported by algorithms that support multiple operators (e.g. reductions, +groupby). Here are some other guidelines. + + * In general empty input is not an error in libcudf. Typically empty input results in empty output. + Tests should verify this. + + * Anything that involves manipulating bitmasks (especially hand-rolled kernels) should have tests + that check varying number of rows, especially around boundaries like the warp size (32). So, test + fewer than 32 rows, more than 32 rows, exactly 32 rows, and greater than 64 rows. + + * Most algorithms should have one or more tests exercising inputs with a large enough number of + rows to require launching multiple thread blocks, especially when values are ultimately + communicated between blocks (e.g. reductions). This is especially important for custom kernels + but also applies to Thrust and CUB algorithm calls with lambdas / functors. + + * For anything involving strings or lists, test exhaustive combinations of empty strings/lists, + null strings/lists and strings/lists with null elements. + + * Strings tests should include a mixture of non-ASCII UTF-8 characters like `é` in test data. + + * Test sliced columns as input (that is, columns that have a nonzero `offset`). This is an easy to + forget case. + + * Tests that verify various forms of "degenerate" column inputs, for example: empty + string columns that have no children (not many paths in cudf can generate these but it + does happen); columns with zero size but that somehow have non-null data pointers; and struct + columns with no children. + + * Decimal types are not included in the `cudf::test::NumericTypes` type list, but are included in + `cudf::test::FixedWidthTypes`, so be careful that tests either include or exclude decimal types as + appropriate. + + +## Directory and File Naming + +The naming of unit test directories and source files should be consistent with the feature being +tested. For example, the tests for APIs in `copying.hpp` should live in `cudf/cpp/tests/copying`. +Each feature (or set of related features) should have its own test source file named +`_tests.cu/cpp`. For example, `cudf/cpp/src/copying/scatter.cu` has tests in +`cudf/cpp/tests/copying/scatter_tests.cu`. + +In the interest of improving compile time, whenever possible, test source files should be `.cpp` +files because `nvcc` is slower than `gcc` in compiling host code. Note that `thrust::device_vector` +includes device code, and so must only be used in `.cu` files. `rmm::device_uvector`, +`rmm::device_buffer` and the various `column_wrapper` types described later can be used in `.cpp` +files, and are therefore preferred in test code over `thrust::device_vector`. + +## Base Fixture + +All libcudf unit tests should make use of a GTest ["Test Fixture"](https://github.com/google/googletest/blob/master/docs/primer.md#test-fixtures-using-the-same-data-configuration-for-multiple-tests-same-data-multiple-tests). +Even if the fixture is empty, it should inherit from the base fixture `cudf::test::BaseFixture` +found in `include/cudf_test/base_fixture.hpp`. This ensures that RMM is properly initialized and +finalized. `cudf::test::BaseFixture` already inherits from `testing::Test` and therefore it is +not necessary for your test fixtures to inherit from it. + +Example: + + class MyTestFixture : public cudf::test::BaseFixture {...}; + +## Typed Tests + +In general, libcudf features must work across all of the supported types (there are exceptions e.g. +not all binary operations are supported for all types). In order to automate the process of running +the same tests across multiple types, we use GTest's +[Typed Tests](https://github.com/google/googletest/blob/master/docs/advanced.md#typed-tests). +Typed tests allow you to write a test once and run it across a list of types. + +For example: + +```c++ +// Fixture must be a template +template +class TypedTestFixture : cudf::test::BaseFixture {...}; +using TestTypes = cudf::test:Types; // Notice custom cudf type list type +TYPED_TEST_SUITE(TypedTestFixture, TestTypes); +TYPED_TEST(TypedTestFixture, FirstTest){ + // Access the current type using `TypeParam` + using T = TypeParam; +} +``` + +To specify the list of types to use, instead of GTest's `testing::Types<...>`, libcudf provides `cudf::test::Types<...>` which is a custom, drop-in replacement for `testing::Types`. +In this example, all tests using the `TypedTestFixture` fixture will run once for each type in the +list defined in `TestTypes` (`int, float, double`). + +### Type Lists + +The list of types that are used in tests should be consistent across all tests. To ensure +consistency, several sets of common type lists are provided in +`include/cudf_test/type_lists.hpp`. For example, `cudf::test::NumericTypes` is a type list of all numeric types, +`FixedWidthTypes` is a list of all fixed-width element types, and `cudf::test::AllTypes` is a list of every +element type that libcudf supports. + +```c++ +#include + +// All tests using TypeTestFixture will be invoked once for each numeric type +TYPED_TEST_SUITE(TypedTestFixture, cudf::test::NumericTypes); +``` + +Whenever possible, use one of the type list provided in `include/utilities/test/type_lists.hpp` +rather than creating new custom lists. + +#### Advanced Type Lists + +Sometimes it is necessary to generate more advanced type lists than the simple lists of single types +in the `TypeList` example above. libcudf provides a set of meta-programming utilities in +`include/cudf_test/type_list_utilities.hpp` for generating and composing more advanced type lists. + +For example, it may be useful to generate a *nested* type list where each element in the list is two +types. In a nested type list, each element in the list is itself another list. In order to access +the `N`th type within the nested list, use `GetType`. + +Imagine testing all possible two-type combinations of ``. This could be done manually: + +```c++ +template +TwoTypesFixture : cudf::test::BaseFixture{...}; +using TwoTypesList = Types< Types, Types, + Types, Types >; +TYPED_TEST_SUITE(TwoTypesFixture, TwoTypesList); +TYPED_TEST(TwoTypesFixture, FirstTest){ + // TypeParam is a list of two types, i.e., a "nested" type list + // Use `cudf::test::GetType` to retrieve the individual types + using FirstType = GetType; + using SecondType = GetType; +} +``` + +The above example manually specifies all pairs composed of `int` and `float`. `CrossProduct` is a +utility in `type_list_utilities.hpp` which materializes this cross product automatically. + +```c++ +using TwoTypesList = Types< Types, Types, + Types, Types >; +using CrossProductTypeList = CrossProduct< Types, Types >; +// TwoTypesList and CrossProductTypeList are identical +``` + +`CrossProduct` can be used with an arbitrary number of type lists to generate nested type lists of +two or more types. **However**, overuse of `CrossProduct` can dramatically inflate compile time. +The cross product of two type lists of size `n` and `m` will result in a new list with +`n*m` nested type lists. This means `n*m` templates will be instantiated; `n` and `m` need not be +large before compile time becomes unreasonable. + +There are a number of other utilities in `type_list_utilities.hpp`. For more details, see the +documentation in that file and their associated tests in +`cudf/cpp/tests/utilities_tests/type_list_tests.cpp`. + +## Utilities + +libcudf provides a number of utilities in `include/cudf_test` to make common testing operations more +convenient. Before creating your own test utilities, look to see if one already exists that does +what you need. If not, consider adding a new utility to do what you need. However, make sure that +the utility is generic enough to be useful for other tests and is not overly tailored to your +specific testing need. + +### Column Wrappers + +In order to make generating input columns easier, libcudf provides the `*_column_wrapper` classes in +`include/cudf_test/column_wrapper.hpp`. These classes wrap a `cudf::column` and provide constructors +for initializing a `cudf::column` object usable with libcudf APIs. Any `*_column_wrapper` class is +implicitly convertible to a `column_view` or `mutable_column_view` and therefore may be +transparently passed to any API expecting a `column_view` or `mutable_column_view` argument. + +#### fixed_width_column_wrapper + +The `cudf::test::fixed_width_column_wrapper` class should be used for constructing and initializing columns of +any fixed-width element type, e.g., numeric types, timestamp types, Boolean, etc. +`cudf::test::fixed_width_column_wrapper` provides constructors that accept an iterator range to generate each +element in the column. For nullable columns, an additional iterator can be provided to indicate the +validity of each element. There are also constructors that accept a `std::initializer_list` for +the column elements and optionally for the validity of each element. + +Example: + +```c++ +// Creates a non-nullable column of INT32 elements with 5 elements: {0, 1, 2, 3, 4} +auto elements = cudf::detail::make_counting_transform_iterator(0, [](auto i){return i;}); +cudf::test::fixed_width_column_wrapper w(elements, elements + 5); + +// Creates a nullable column of INT32 elements with 5 elements: {null, 1, null, 3, null} +auto elements = cudf::detail::make_counting_transform_iterator(0, [](auto i){return i;}); +auto validity = cudf::detail::make_counting_transform_iterator(0, [](auto i){return i % 2;}) +cudf::test::fixed_width_column_wrapper w(elements, elements + 5, validity); + +// Creates a non-nullable INT32 column with 4 elements: {1, 2, 3, 4} +cudf::test::fixed_width_column_wrapper w{{1, 2, 3, 4}}; + +// Creates a nullable INT32 column with 4 elements: {1, NULL, 3, NULL} +cudf::test::fixed_width_column_wrapper w{ {1,2,3,4}, {1, 0, 1, 0}}; +``` + +#### fixed_point_column_wrapper + +The `cudf::test::fixed_point_column_wrapper` class should be used for constructing and initializing columns of +any fixed-point element type (DECIMAL32 or DECIMAL64). `cudf::test::fixed_point_column_wrapper` provides +constructors that accept an iterator range to generate each element in the column. For nullable +columns, an additional iterator can be provided to indicate the validity of each element. +Constructors also take the scale of the fixed-point values to create. + +Example: + +```c++ +// Creates a non-nullable column of 4 DECIMAL32 elements of scale 3: {1000, 2000, 3000, 4000} +auto elements = cudf::detail::make_counting_transform_iterator(0, [](auto i){ return i; }); +cudf::test::fixed_point_column_wrapper w(elements, elements + 4, 3); + +// Creates a nullable column of 5 DECIMAL32 elements of scale 2: {null, 100, null, 300, null} +auto elements = cudf::detail::make_counting_transform_iterator(0, [](auto i){ return i; }); +auto validity = cudf::detail::make_counting_transform_iterator(0, [](auto i){ return i % 2; }); +cudf::test::fixed_point_column_wrapper w(elements, elements + 5, validity, 2); +``` + +#### dictionary_column_wrapper + +The `cudf::test::dictionary_column_wrapper` class should be used to create dictionary columns. +`cudf::test::dictionary_column_wrapper` provides constructors that accept an iterator range to generate each +element in the column. For nullable columns, an additional iterator can be provided to indicate the +validity of each element. There are also constructors that accept a `std::initializer_list` for +the column elements and optionally for the validity of each element. + +Example: + +```c++ +// Creates a non-nullable dictionary column of INT32 elements with 5 elements +// keys = {0, 2, 6}, indices = {0, 1, 1, 2, 2} +std::vector elements{0, 2, 2, 6, 6}; +cudf::test::dictionary_column_wrapper w(element.begin(), elements.end()); + +// Creates a nullable dictionary column with 5 elements and a validity iterator. +std::vector elements{0, 2, 0, 6, 0}; +// Validity iterator here sets even rows to null. +auto validity = cudf::detail::make_counting_transform_iterator(0, [](auto i){return i % 2;}) +// keys = {2, 6}, indices = {NULL, 0, NULL, 1, NULL} +cudf::test::dictionary_column_wrapper w(elements, elements + 5, validity); + +// Creates a non-nullable dictionary column with 4 elements. +// keys = {1, 2, 3}, indices = {0, 1, 2, 0} +cudf::test::dictionary_column_wrapper w{{1, 2, 3, 1}}; + +// Creates a nullable dictionary column with 4 elements and validity initializer. +// keys = {1, 3}, indices = {0, NULL, 1, NULL} +cudf::test::dictionary_column_wrapper w{ {1, 0, 3, 0}, {1, 0, 1, 0}}; + +// Creates a nullable column of dictionary elements with 5 elements and validity initializer. +std::vector elements{0, 2, 2, 6, 6}; +// keys = {2, 6}, indices = {NULL, 0, NULL, 1, NULL} +cudf::test::dictionary_width_column_wrapper w(elements, elements + 5, {0, 1, 0, 1, 0}); + +// Creates a non-nullable dictionary column with 7 string elements +std::vector strings{"", "aaa", "bbb", "aaa", "bbb", "ccc", "bbb"}; +// keys = {"","aaa","bbb","ccc"}, indices = {0, 1, 2, 1, 2, 3, 2} +cudf::test::dictionary_column_wrapper d(strings.begin(), strings.end()); + +// Creates a nullable dictionary column with 7 string elements and a validity iterator. +// Validity iterator here sets even rows to null. +// keys = {"a", "bb"}, indices = {NULL, 1, NULL, 1, NULL, 0, NULL} +auto validity = cudf::detail::make_counting_transform_iterator(0, [](auto i){return i % 2;}); +cudf::test::dictionary_column_wrapper d({"", "bb", "", "bb", "", "a", ""}, validity); +``` + +#### strings_column_wrapper + +The `cudf::test::strings_column_wrapper` class should be used to create columns of strings. It provides +constructors that accept an iterator range to generate each string in the column. For nullable +columns, an additional iterator can be provided to indicate the validity of each string. There are +also constructors that accept a `std::initializer_list` for the column's strings and +optionally for the validity of each element. + +Example: + +```c++ +// Creates a non-nullable STRING column with 7 string elements: +// {"", "this", "is", "a", "column", "of", "strings"} +std::vector strings{"", "this", "is", "a", "column", "of", "strings"}; +cudf::test::strings_column_wrapper s(strings.begin(), strings.end()); + +// Creates a nullable STRING column with 7 string elements: +// {NULL, "this", NULL, "a", NULL, "of", NULL} +std::vector strings{"", "this", "is", "a", "column", "of", "strings"}; +auto validity = cudf::detail::make_counting_transform_iterator(0, [](auto i){return i % 2;}); +cudf::test::strings_column_wrapper s(strings.begin(), strings.end(), validity); + +// Creates a non-nullable STRING column with 7 string elements: +// {"", "this", "is", "a", "column", "of", "strings"} +cudf::test::strings_column_wrapper s({"", "this", "is", "a", "column", "of", "strings"}); + +// Creates a nullable STRING column with 7 string elements: +// {NULL, "this", NULL, "a", NULL, "of", NULL} +auto validity = cudf::detail::make_counting_transform_iterator(0, [](auto i){return i % 2;}); +cudf::test::strings_column_wrapper s({"", "this", "is", "a", "column", "of", "strings"}, validity); +``` + +#### lists_column_wrapper + +The `cudf::test::lists_column_wrapper` class should be used to create columns of lists. It provides +constructors that accept an iterator range to generate each list in the column. For nullable +columns, an additional iterator can be provided to indicate the validity of each list. There are +also constructors that accept a `std::initializer_list` for the column's lists and +optionally for the validity of each element. A number of other constructors are available. + +Example: + +```c++ +// Creates an empty LIST column +// [] +cudf::test::lists_column_wrapper l{}; + +// Creates a LIST column with 1 list composed of 2 total integers +// [{0, 1}] +cudf::test::lists_column_wrapper l{0, 1}; + +// Creates a LIST column with 3 lists +// [{0, 1}, {2, 3}, {4, 5}] +cudf::test::lists_column_wrapper l{ {0, 1}, {2, 3}, {4, 5} }; + +// Creates a LIST of LIST columns with 2 lists on the top level and +// 4 below +// [ {{0, 1}, {2, 3}}, {{4, 5}, {6, 7}} ] +cudf::test::lists_column_wrapper l{ {{0, 1}, {2, 3}}, {{4, 5}, {6, 7}} }; + +// Creates a LIST column with 1 list composed of 5 total integers +// [{0, 1, 2, 3, 4}] +auto elements = cudf::detail::make_counting_transform_iterator(0, [](auto i){return i*2;}); +cudf::test::lists_column_wrapper l(elements, elements+5); + +// Creates a LIST column with 1 lists composed of 2 total integers +// [{0, NULL}] +auto validity = cudf::detail::make_counting_transform_iterator(0, [](auto i){return i % 2;}); +cudf::test::lists_column_wrapper l{{0, 1}, validity}; + +// Creates a LIST column with 1 lists composed of 5 total integers +// [{0, NULL, 2, NULL, 4}] +auto elements = cudf::detail::make_counting_transform_iterator(0, [](auto i){return i*2;}); +auto validity = cudf::detail::make_counting_transform_iterator(0, [](auto i){return i % 2;}); +cudf::test::lists_column_wrapper l(elements, elements+5, validity); + +// Creates a LIST column with 1 list composed of 2 total strings +// [{"abc", "def"}] +cudf::test::lists_column_wrapper l{"abc", "def"}; + +// Creates a LIST of LIST columns with 2 lists on the top level and 4 below +// [ {{0, 1}, NULL}, {{4, 5}, NULL} ] +auto validity = cudf::detail::make_counting_transform_iterator(0, [](auto i){return i % 2;}); +cudf::test::lists_column_wrapper l{ {{{0, 1}, {2, 3}}, validity}, {{{4, 5}, {6, 7}}, validity} }; +``` + +#### structs_column_wrapper + +The `cudf::test::structs_column_wrapper` class should be used to create columns of structs. It provides +constructors that accept a vector or initializer list of pre-constructed columns or column wrappers +for child columns. For nullable columns, an additional iterator can be provided to indicate the +validity of each struct. + +Examples: + +```c++ +// The following constructs a column for struct< int, string >. +auto child_int_col = cudf::test::fixed_width_column_wrapper{ 1, 2, 3, 4, 5 }.release(); +auto child_string_col = cudf::test::string_column_wrapper {"All", "the", "leaves", "are", "brown"}.release(); + +std::vector> child_columns; +child_columns.push_back(std::move(child_int_col)); +child_columns.push_back(std::move(child_string_col)); + +cudf::test::struct_col wrapper wrapper{ + child_cols, + {1,0,1,0,1} // Validity +}; + +auto struct_col {wrapper.release()}; + +// The following constructs a column for struct< int, string >. +cudf::test::fixed_width_column_wrapper child_int_col_wrapper{ 1, 2, 3, 4, 5 }; +cudf::test::string_column_wrapper child_string_col_wrapper {"All", "the", "leaves", "are", "brown"}; + +cudf::test::structs_column_wrapper wrapper{ + {child_int_col_wrapper, child_string_col_wrapper} + {1,0,1,0,1} // Validity +}; + +auto struct_col {wrapper.release()}; + +// The following constructs a column for struct< int, string >. +cudf::test::fixed_width_column_wrapper child_int_col_wrapper{ 1, 2, 3, 4, 5 }; +cudf::test::string_column_wrapper child_string_col_wrapper {"All", "the", "leaves", "are", "brown"}; + +cudf::test::structs_column_wrapper wrapper{ + {child_int_col_wrapper, child_string_col_wrapper} + cudf::detail::make_counting_transform_iterator(0, [](auto i){ return i % 2; }) // Validity +}; + +auto struct_col {wrapper.release()}; +``` + +### Column Comparison Utilities + +A common operation in testing is verifying that two columns are equal, or equivalent, or that they +have the same metadata. + +#### CUDF_TEST_EXPECT_COLUMN_PROPERTIES_EQUAL + +Verifies that two columns have the same type, size, and nullability. For nested types, recursively +verifies the equality of type, size and nullability of all nested children. + +#### CUDF_TEST_EXPECT_COLUMN_PROPERTIES_EQUIVALENT + +Verifies that two columns have equivalent type and equal size, ignoring nullability. For nested +types, recursively verifies the equivalence of type, and equality of size of all nested children, +ignoring nullability. + +Note "equivalent type". Most types are equivalent if and only they are equal. `fixed_point` types +are one exception. They are equivalent if the representation type is equal, even if they have +different scales. Nested type columns can be equivalent in the case where they both have zero size, +but one has children (also empty) and the other does not. For columns with nonzero size, both equals +and equivalent expect equal number of children. + +#### CUDF_TEST_EXPECT_COLUMNS_EQUAL + +Verifies that two columns have equal properties and verifies elementwise equality of the column +data. Null elements are treated as equal. + +#### CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT + +Verifies that two columns have equivalent properties and verifies elementwise equivalence of the +column data. Null elements are treated as equivalent. + +#### CUDF_TEST_EXPECT_EQUAL_BUFFERS + +Verifies the bitwise equality of two device memory buffers. + +#### Caveats + +Column comparison functions in the `cudf::test::detail` namespace should **NOT** be used directly. + +(printing-and-accessing-column-data)= +### Printing and accessing column data + +The `` header defines various functions and overloads for printing +columns (`print`), converting column data to string (`to_string`, `to_strings`), and copying data to +the host (`to_host`). For example, to print a `cudf::column_view` contents or `column_wrapper` instance +to the console use the `cudf::test::print()`: +```cpp + cudf::test::fixed_width_column_wrapper input({1,2,3,4}); + auto splits = cudf::split(input,{2}); + cudf::test::print(input); + cudf::test::print(splits.front()); +``` +Fixed-width and strings columns output as comma-separated entries including null rows. +Nested columns are also supported and output includes the offsets and data children as well as +the null mask bits. + +## Validating Stream Usage + +### Background + +libcudf employs a custom-built [preload +library](https://man7.org/linux/man-pages/man8/ld.so.8.html) to validate its internal stream usage +(the code may be found +[`here`](https://github.com/NVIDIA/cudf/blob/main/cpp/tests/utilities/identify_stream_usage.cpp)). +This library wraps every asynchronous CUDA runtime API call that accepts a stream with a check to +ensure that the passed CUDA stream is a valid one, immediately throwing an exception if an invalid +stream is detected. Running tests with this library loaded immediately triggers errors if any test +accidentally runs code on an invalid stream. + +Stream validity is determined by overloading the definition of libcudf's default stream. Normally, in +libcudf `cudf::get_default_stream` returns one of `rmm`'s default stream values (depending on +whether or not libcudf is compiled with per thread default stream enabled). In the preload library, +this function is redefined to instead return a new user-created stream managed using a +function-local static `rmm::cuda_stream`. An invalid stream in this situation is defined as any of +CUDA's default stream values (cudaStreamLegacy, cudaStreamDefault, or cudaStreamPerThread), since +any kernel that properly uses `cudf::get_default_stream` will now instead be using the custom stream +created by the preload library. + +The preload library supports two different modes, `cudf` mode and `testing` mode. The previous +paragraph describes the behavior of `cudf` mode, where `cudf::get_default_stream` is overloaded. In +`cudf` mode, the preload library ensures that all CUDA runtime APIs are being provided cudf's +default stream. This will detect oversights where, for example, a Thrust call has no stream specified, or +when one of CUDA's default stream values is explicitly specified to a kernel. However, it will not +detect cases where a stream is not correctly forwarded down the call stack, for instance if +some `detail` function that accepts a stream parameter fails to forward it along and instead +erroneously calls `cudf::get_default_stream` instead. + +In `testing` mode, the library instead overloads `cudf::test::get_default_stream`. This function +defined in the `cudf::test` namespace enables a more stringent mode of testing. In `testing` mode, +the preload library instead verifies that all CUDA runtime APIs are instead called using the test +namespace's default stream. This distinction is important because cudf internals never use +`cudf::test::get_default_stream`, so this stream value can only appear internally if it was provided +to a public API and forwarded properly all the way down the call stack. While `testing` mode is more +strict than `cudf` mode, it is also more intrusive. `cudf` mode can operate with no changes to the +library or the tests because the preload library overwrites the relevant APIs in place. `testing` +mode, however, can only be used to validate tests that are correctly passing +`cudf::test::get_default_stream` to public libcudf APIs. + +In addition to the preload library, the test suite also implements a [custom memory +resource](https://github.com/NVIDIA/cudf/blob/main/cpp/include/cudf_test/stream_checking_resource_adaptor.hpp) +that performs analogous stream verification when its `do_allocate` method is called. During testing +this rmm's default memory resource is set to use this adaptor for additional stream validation. + +### Usage + +When writing tests for a libcudf API, a special set of additional tests should be added to validate +the API's stream usage. These tests should be placed in the `cpp/tests/streams` directory in a file +corresponding to the header containing the tested APIs, e.g. `cpp/tests/streams/copying_test.cpp` +for all APIs declared in `cpp/include/cudf/copying.hpp`. These tests should contain a minimal +invocation of the tested API with no additional assertions since they are solely designed to check +stream usage. When adding these tests to `cpp/tests/CMakeLists.txt`, the `ConfigureTest` CMake +function should be provided the arguments `STREAM_MODE testing`. This change is sufficient for +CTest to set up the test to automatically load the preload library compiled in `testing` mode when +running the test. + +The rest of the test suite is configured to run with the preload library in `cudf` mode. As a +result, all test runs with `ctest` will always include stream validation. Since this configuration +is managed via CMake and CTest, direct execution of the test executables will not use the preload +library at all. Tests will still run and pass normally in this situation, however (with the +exception of the test of the preload library itself). diff --git a/docs/cudf/source/libcudf/developer_guide/strings.png b/docs/cudf/source/libcudf/developer_guide/strings.png new file mode 100644 index 000000000000..1d18ea8a407e Binary files /dev/null and b/docs/cudf/source/libcudf/developer_guide/strings.png differ diff --git a/docs/cudf/source/libcudf/index.rst b/docs/cudf/source/libcudf/index.rst index 9f390f647ef0..a3d6673e338c 100644 --- a/docs/cudf/source/libcudf/index.rst +++ b/docs/cudf/source/libcudf/index.rst @@ -8,3 +8,4 @@ libcudf api_docs/index.rst md_regex unicode_limitations + developer_guide/DEVELOPER_GUIDE diff --git a/docs/dask_cudf/source/_static/RAPIDS-logo-purple.png b/docs/dask_cudf/source/_static/RAPIDS-logo-purple.png deleted file mode 100644 index d884e01374dc..000000000000 Binary files a/docs/dask_cudf/source/_static/RAPIDS-logo-purple.png and /dev/null differ diff --git a/docs/dask_cudf/source/best_practices.rst b/docs/dask_cudf/source/best_practices.rst index 675e5fc1c11a..037de0e1c5cb 100644 --- a/docs/dask_cudf/source/best_practices.rst +++ b/docs/dask_cudf/source/best_practices.rst @@ -3,8 +3,8 @@ Dask cuDF Best Practices ======================== -This page outlines several important guidelines for using `Dask cuDF -`__ effectively. +This page outlines several important guidelines for using +:doc:`Dask cuDF ` effectively. .. note:: Since Dask cuDF is a backend extension for @@ -22,24 +22,23 @@ Use Dask-CUDA ~~~~~~~~~~~~~ To execute a Dask workflow on multiple GPUs, a Dask cluster must -be deployed with `Dask-CUDA `__ +be deployed with :doc:`Dask-CUDA ` and `Dask.distributed `__. -When running on a single machine, the `LocalCUDACluster `__ +When running on a single machine, the :class:`~dask_cuda.LocalCUDACluster` convenience function is strongly recommended. No matter how many GPUs are -available on the machine (even one!), using `Dask-CUDA has many advantages -`__ +available on the machine (even one!), using :ref:`Dask-CUDA has many advantages +` over default (threaded) execution. Just to list a few: * Dask-CUDA makes it easy to pin workers to specific devices. * Dask-CUDA makes it easy to configure memory-spilling options. * The distributed scheduler collects useful diagnostic information that can be viewed on a dashboard in real time. -Please see `Dask-CUDA's API `__ -and `Best Practices `__ +Please see :doc:`Dask-CUDA's API ` +and :doc:`Best Practices ` documentation for detailed information. Typical ``LocalCUDACluster`` usage -is also illustrated within the multi-GPU section of `Dask cuDF's -`__ documentation. +is also shown in :ref:`multiple_gpus`. .. note:: When running on cloud infrastructure or HPC systems, it is usually best to @@ -47,7 +46,7 @@ is also illustrated within the multi-GPU section of `Dask cuDF's `__ and `Dask-Jobqueue `__. - Please see `the RAPIDS deployment documentation `__ + Please see `the cloud deployment documentation `__ for further details and examples. @@ -71,23 +70,23 @@ Enable cuDF spilling ~~~~~~~~~~~~~~~~~~~~ When using Dask cuDF for classic ETL workloads, it is usually best -to enable `native spilling support in cuDF -`__. -When using :class:`dask_cuda.LocalCUDACluster`, this is easily accomplished by +to enable :ref:`native spilling support in cuDF +`. +When using :class:`~dask_cuda.LocalCUDACluster`, this is easily accomplished by setting ``enable_cudf_spill=True``. Use RMM ~~~~~~~ Memory allocations in cuDF are significantly faster and more efficient when -the `RAPIDS Memory Manager (RMM) `__ -library is configured appropriately on worker processes. In most cases, the best way to manage +:doc:`NVIDIA RMM ` +is configured appropriately on worker processes. In most cases, the best way to manage memory is by initializing an RMM pool on each worker before executing a -workflow. When using :class:`dask_cuda.LocalCUDACluster`, this is easily accomplished +workflow. When using :class:`~dask_cuda.LocalCUDACluster`, this is easily accomplished by setting ``rmm_pool_size`` to a large fraction (e.g. ``0.9``). -See the `Dask-CUDA memory-management documentation -`__ +See the :ref:`Dask-CUDA memory-management documentation +` for more details. Use the Dask DataFrame API @@ -289,15 +288,15 @@ bottleneck is typically device-to-host memory spilling. Although every workflow is different, the following guidelines are often recommended: -* Use a distributed cluster with `Dask-CUDA `__ workers +* Use a distributed cluster with :doc:`Dask-CUDA ` workers -* Use native cuDF spilling whenever possible (`Dask-CUDA spilling documentation `__) +* Use native cuDF spilling whenever possible (:doc:`Dask-CUDA spilling documentation `) * Avoid shuffling whenever possible * Use ``split_out=1`` for low-cardinality groupby aggregations * Use ``broadcast=True`` for joins when at least one collection comprises a small number of partitions (e.g. ``<=5``) -* `Use UCX `__ if communication is a bottleneck. +* :doc:`Use UCX ` if communication is a bottleneck. .. note:: UCX enables Dask-CUDA workers to communicate using high-performance diff --git a/docs/dask_cudf/source/conf.py b/docs/dask_cudf/source/conf.py index 635ba31f5754..1e31ab671998 100644 --- a/docs/dask_cudf/source/conf.py +++ b/docs/dask_cudf/source/conf.py @@ -59,7 +59,7 @@ htmlhelp_basename = "dask-cudfdoc" html_use_modindex = True -html_static_path = ["_static"] +html_static_path = [] pygments_style = "sphinx" @@ -83,11 +83,11 @@ "cupy": ("https://docs.cupy.dev/en/stable/", None), "numpy": ("https://numpy.org/doc/stable/", None), "pyarrow": ("https://arrow.apache.org/docs/", None), - "cudf": ("https://docs.rapids.ai/api/cudf/stable/", None), + "cudf": (f"https://docs.nvidia.com/cudf/{version}/", None), "dask": ("https://docs.dask.org/en/stable/", None), - # Temporarily disable pandas intersphinx: https://github.com/pandas-dev/pandas/issues/64584 - # "pandas": ("https://pandas.pydata.org/docs/", None), - "dask-cuda": ("https://docs.rapids.ai/api/dask-cuda/stable/", None), + "pandas": ("https://pandas.pydata.org/docs/", None), + "dask-cuda": (f"https://docs.nvidia.com/dask-cuda/{version}/", None), + "rmm": (f"https://docs.nvidia.com/rmm/{version}/", None), } numpydoc_show_inherited_class_members = True diff --git a/docs/dask_cudf/source/index.rst b/docs/dask_cudf/source/index.rst index eee1bc39fc4b..0bf4f5a4f50b 100644 --- a/docs/dask_cudf/source/index.rst +++ b/docs/dask_cudf/source/index.rst @@ -21,12 +21,11 @@ as the ``"cudf"`` dataframe backend for of the GPU and networking hardware. If you are familiar with Dask and `pandas `__ or -`cuDF `__, then Dask cuDF +:doc:`cuDF `, then Dask cuDF should feel familiar to you. If not, we recommend starting with `10 minutes to Dask `__ followed -by `10 minutes to cuDF and Dask cuDF -`__. +by :doc:`10 minutes to cuDF and Dask cuDF `. After reviewing the sections below, please see the :ref:`Best Practices ` page for further guidance on @@ -120,7 +119,7 @@ automatic query planning (see the next section). Query Planning ~~~~~~~~~~~~~~ -Dask cuDF now provides automatic query planning by default (RAPIDS 24.06+). +Since version 24.06, Dask cuDF provides automatic query planning by default. As long as the ``"dataframe.query-planning"`` configuration is set to ``True`` (the default) when ``dask.dataframe`` is first imported, `Dask Expressions `__ will be used under the hood. @@ -149,6 +148,8 @@ Simplified expression graph (``df.simplify().pprint()``):: (via :func:`dask.compute` or :func:`dask.persist`). You do not need to optimize or simplify the graph yourself. +.. _multiple_gpus: + Using Multiple GPUs and Multiple Nodes ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ @@ -193,8 +194,7 @@ to define a client object. For example:: Please see the :doc:`dask-cuda:index` documentation for more information about deploying GPU-aware clusters -(including `best practices -`__). +(including :doc:`best practices `). API Reference @@ -202,7 +202,7 @@ API Reference Generally speaking, Dask cuDF tries to offer exactly the same API as Dask DataFrame. There are, however, some minor differences mostly because -cuDF does not `perfectly mirror `__ +cuDF does not :doc:`perfectly mirror ` the pandas API, or because cuDF provides additional configuration flags (these mostly occur in data reading and writing interfaces). diff --git a/python/cudf/cudf/core/groupby/groupby.py b/python/cudf/cudf/core/groupby/groupby.py index ce3bb1dcd28e..d9cc17506a32 100644 --- a/python/cudf/cudf/core/groupby/groupby.py +++ b/python/cudf/cudf/core/groupby/groupby.py @@ -519,7 +519,7 @@ def _collect_series_key_column_names(obj, by) -> dict[int, Hashable]: class GroupByNthSelector: - """Mirror of :class:`pandas.core.groupby.indexing.GroupByNthSelector`. + """Mirror of ``pandas.core.groupby.indexing.GroupByNthSelector``. ``GroupBy.nth`` supports both the call form ``gb.nth(n, dropna=...)`` and the index form ``gb.nth[n]``. @@ -1505,8 +1505,8 @@ def _reduce( Computed {op} of values within each group. .. pandas-compat:: - :meth:`pandas.core.groupby.DataFrameGroupBy.{op}`, - :meth:`pandas.core.groupby.SeriesGroupBy.{op}` + :meth:`pandas.api.typing.DataFrameGroupBy.{op}`, + :meth:`pandas.api.typing.SeriesGroupBy.{op}` The numeric_only, min_count """ @@ -2640,8 +2640,8 @@ def mult(df): 6 2 6 12 .. pandas-compat:: - :meth:`pandas.core.groupby.DataFrameGroupBy.apply`, - :meth:`pandas.core.groupby.SeriesGroupBy.apply` + :meth:`pandas.api.typing.DataFrameGroupBy.apply`, + :meth:`pandas.api.typing.SeriesGroupBy.apply` cuDF's ``groupby.apply`` is limited compared to pandas. In some situations, Pandas returns the grouped keys as part of @@ -3593,8 +3593,8 @@ def shift( Object shifted within each group. .. pandas-compat:: - :meth:`pandas.core.groupby.DataFrameGroupBy.shift`, - :meth:`pandas.core.groupby.SeriesGroupBy.shift` + :meth:`pandas.api.typing.DataFrameGroupBy.shift`, + :meth:`pandas.api.typing.SeriesGroupBy.shift` Parameter ``freq`` is unsupported. """ diff --git a/python/cudf/cudf/core/indexed_frame.py b/python/cudf/cudf/core/indexed_frame.py index 37b88e55c19a..53ce03856b43 100644 --- a/python/cudf/cudf/core/indexed_frame.py +++ b/python/cudf/cudf/core/indexed_frame.py @@ -434,7 +434,7 @@ def flags(self) -> pd.Flags: The available flags are - * :attr:`pandas.Flags.allows_duplicate_labels` + * ``allows_duplicate_labels`` See Also -------- diff --git a/python/cudf/cudf/core/multiindex.py b/python/cudf/cudf/core/multiindex.py index 5be6335601cd..cbfa76c83055 100644 --- a/python/cudf/cudf/core/multiindex.py +++ b/python/cudf/cudf/core/multiindex.py @@ -594,7 +594,7 @@ def __repr__(self) -> str: @property @_external_only_api("Use ._codes instead") @_performance_tracking - def codes(self) -> pd.core.indexes.frozen.FrozenList: + def codes(self) -> pd.api.typing.FrozenList: """ Returns the codes of the underlying MultiIndex. diff --git a/python/cudf/pyproject.toml b/python/cudf/pyproject.toml index a0d16b62935d..818e8f934bf6 100644 --- a/python/cudf/pyproject.toml +++ b/python/cudf/pyproject.toml @@ -85,7 +85,7 @@ cudf-pandas-tests = [ [project.urls] Homepage = "https://github.com/NVIDIA/cudf" -Documentation = "https://docs.rapids.ai/api/cudf/stable/" +Documentation = "https://docs.nvidia.com/cudf/latest/" [tool.pydistcheck] select = [ diff --git a/python/cudf_polars/docs/cudf-polars-mp.md b/python/cudf_polars/docs/cudf-polars-mp.md index 0aaf27d073d8..e0dc29917dd9 100644 --- a/python/cudf_polars/docs/cudf-polars-mp.md +++ b/python/cudf_polars/docs/cudf-polars-mp.md @@ -779,6 +779,6 @@ argument; user-supplied keys are merged with reserved entries set by `SPMDEngine [spmd-wiki]: https://en.wikipedia.org/wiki/Single_program,_multiple_data [ray-docs]: https://docs.ray.io/en/latest/ [ray-actors]: https://docs.ray.io/en/latest/ray-core/actors.html -[rapidsmpf-communicator]: https://docs.rapids.ai/api/rapidsmpf/stable/glossary/#term-Communicator -[rapidsmpf-context]: https://docs.rapids.ai/api/rapidsmpf/stable/glossary/#term-Context +[rapidsmpf-communicator]: https://docs.nvidia.com/rapidsmpf/latest/glossary/#term-Communicator +[rapidsmpf-context]: https://docs.nvidia.com/rapidsmpf/latest/glossary/#term-Context [polars-gpuengine]: https://docs.pola.rs/api/python/stable/reference/api/polars.GPUEngine.html diff --git a/python/cudf_polars/docs/overview.md b/python/cudf_polars/docs/overview.md index f2a11f91bc71..f715f77f2068 100644 --- a/python/cudf_polars/docs/overview.md +++ b/python/cudf_polars/docs/overview.md @@ -647,10 +647,6 @@ another `nvtx` range (e.g. `Scan.do_evaluate`, `GroupBy.do_evaluate`, etc.). These provide a higher-level grouping over the lower-level libcudf calls (e.g. `read_chunk`, `aggregate`). -Finally, if using [rapidsmpf](https://docs.rapids.ai/api/rapidsmpf/nightly/) -for shuffling, the methods inserting and extracting partitions to shuffle are -annotated with nvtx ranges. - # Query Plans The module `cudf_polars.streaming.explain` contains functions for dumping diff --git a/python/dask_cudf/README.md b/python/dask_cudf/README.md index d0d4eee7db62..273c72d86b8e 100644 --- a/python/dask_cudf/README.md +++ b/python/dask_cudf/README.md @@ -3,11 +3,11 @@ Dask cuDF (a.k.a. dask-cudf or `dask_cudf`) is an extension library for [Dask DataFrame](https://docs.dask.org/en/stable/dataframe.html) that provides a Pandas-like API for parallel and larger-than-memory DataFrame computing on GPUs. When installed, Dask cuDF is automatically registered as the `"cudf"` [dataframe backend](https://docs.dask.org/en/stable/how-to/selecting-the-collection-backend.html) for Dask DataFrame. > [!IMPORTANT] -> Dask cuDF does not provide support for multi-GPU or multi-node execution on its own. You must also deploy a distributed cluster (ideally with [Dask-CUDA](https://docs.rapids.ai/api/dask-cuda/stable/)) to leverage multiple GPUs efficiently. +> Dask cuDF does not provide support for multi-GPU or multi-node execution on its own. You must also deploy a distributed cluster (ideally with [Dask-CUDA](https://docs.nvidia.com/dask-cuda/latest/)) to leverage multiple GPUs efficiently. ## Using Dask cuDF -Please visit [the official documentation page](https://docs.rapids.ai/api/dask-cudf/stable/) for detailed information about using Dask cuDF. +Please visit [the official documentation page](https://docs.nvidia.com/dask-cudf/latest/) for detailed information about using Dask cuDF. ## Installation @@ -15,11 +15,11 @@ See the [RAPIDS install page](https://docs.rapids.ai/install/) for the most up-t ## Resources -- [Dask cuDF documentation](https://docs.rapids.ai/api/dask-cudf/stable/) -- [Best practices](https://docs.rapids.ai/api/dask-cudf/stable/best_practices/) -- [cuDF documentation](https://docs.rapids.ai/api/cudf/stable/) -- [10 Minutes to cuDF and Dask cuDF](https://docs.rapids.ai/api/cudf/latest/user_guide/10min/) -- [Dask-CUDA documentation](https://docs.rapids.ai/api/dask-cuda/stable/) +- [Dask cuDF documentation](https://docs.nvidia.com/dask-cudf/latest/) +- [Best practices](https://docs.nvidia.com/dask-cudf/latest/best_practices/) +- [cuDF documentation](https://docs.nvidia.com/cudf/latest/) +- [10 Minutes to cuDF and Dask cuDF](https://docs.nvidia.com/cudf/latest/cudf/10min/) +- [Dask-CUDA documentation](https://docs.nvidia.com/dask-cuda/latest/) - [Deployment](https://docs.rapids.ai/deployment/stable/) - [RAPIDS Community](https://rapids.ai/learn-more/#get-involved): Get help, contribute, and collaborate. @@ -59,6 +59,6 @@ if __name__ == "__main__": query.head() ``` -If you do not have multiple GPUs available, using `LocalCUDACluster` is optional. However, it is still a good idea to [enable cuDF spilling](https://docs.rapids.ai/api/cudf/stable/cudf/developer_guide/library_design/#spilling-to-host-memory). +If you do not have multiple GPUs available, using `LocalCUDACluster` is optional. However, it is still a good idea to [enable cuDF spilling](https://docs.nvidia.com/cudf/latest/cudf/developer_guide/library_design/#spilling-to-host-memory). If you wish to scale across multiple nodes, you will need to use a different mechanism to deploy your Dask-CUDA workers. Please see [the RAPIDS deployment documentation](https://docs.rapids.ai/deployment/stable/) for more instructions.