libcudf C++ 开发者指南

本文档作为 libcudf C++ 代码贡献者的指南。开发者还应参考以下附加文件以获取更多关于 libcudf 最佳实践的文档。

• 文档指南，用于指导 libcudf 代码的文档编写。
• 测试指南，用于指导编写单元测试。
• 基准测试指南，用于指导编写单元基准测试。

概述

libcudf 是一个 C++ 库，提供 GPU 加速的数据并行算法，用于处理面向列的表格数据。libcudf 提供了包括切片、过滤、排序、各种类型的聚合以及分组和连接等数据库类型操作的算法。libcudf 通过包括 Python 和 Java 在内的多种语言接口为许多客户端提供服务。用户也可以直接从 C++ 代码中使用 libcudf。

术语表

本节定义了 libcudf 中使用的术语。

列 (Column)

列是单一类型的数据数组。与表 (Table) 一样，列是 libcudf 中使用的基本数据结构。大多数 libcudf 算法都在列上操作。列可以有一个有效性掩码，表示每个元素是有效还是空值（无效）。支持嵌套类型的列，这意味着一个列可以有子列。列相当于 cuDF Python 的 Series。

元素 (Element)

列中的单个数据项。也称为行。

标量 (Scalar)

表示数据类型单个元素的类型。

表 (Table)

表是具有相同数量元素的列的集合。表相当于 cuDF Python 的 DataFrame。

视图 (View)

视图是一个非拥有对象，提供对其他对象拥有数据的零拷贝访问（可能带有切片或偏移）。例如列视图和表视图。

目录结构和文件命名

外部/公共 libcudf API 根据功能分组到 cudf/cpp/include/cudf/ 目录下具有适当标题的头文件中。例如，cudf/cpp/include/cudf/copying.hpp 包含与从一个列复制到另一个列相关的函数的 API。注意使用 .hpp 文件扩展名表示 C++ 头文件。

外部/公共 libcudf C++ API 头文件需要使用 CUDF_EXPORT 标记其中的所有符号。这通过将宏放在 namespace cudf 上来完成，如下所示。命名空间上的标记要求它们不能嵌套，因此 cudf 命名空间必须单独存在。

#pragma once
 
namespace CUDF_EXPORT cudf {
namespace lists {
 
...
 
 
} // namespace lists
} // namespace CUDF_EXPORT cudf

外部 API 头文件的命名应与其包含 API 实现源代码文件的文件夹名称保持一致。例如，在 cudf/cpp/include/cudf/copying.hpp 中找到的 API 的实现位于 cudf/src/copying 中。同样，这些 API 的单元测试位于 cudf/tests/copying/ 中。

包含 detail 命名空间定义的内部 API 头文件，如果在 libcudf 内部的多个翻译单元中使用，应放在 include/cudf/detail 中。就像公共 C++ API 头文件一样，任何内部 C++ API 头文件都需要在 cudf 命名空间上标记 CUDF_EXPORT，以便可以对函数进行测试。

cudf 中的所有头文件都应使用 #pragma once 作为 include guard。

文件扩展名

• .hpp : C++ 头文件
• .cpp : C++ 源文件
• .cu : CUDA C++ 源文件
• .cuh : 包含 CUDA 设备代码的头文件

仅在必要时使用 .cu 和 .cuh。一个好的指标是包含 __device__ 以及其他仅由 nvcc 识别的符号。另一个指标是带有设备执行策略的 Thrust 算法 API（在 libcudf 中始终是 rmm::exec_policy）。

代码和文档风格及格式化

libcudf 代码除了少数情况外，所有名称都使用 snake_case 风格：模板参数、单元测试和测试用例名称可以使用 Pascal case，也称为 UpperCamelCase。我们不使用 Hungarian notation，除了有时命名设备数据变量及其对应的主机副本时。私有成员变量通常以 _ 前缀开头。

template <typename IteratorType>
void algorithm_function(int x, rmm::cuda_stream_view s, rmm::device_async_resource_ref mr)
{
  ...
}
 
class utility_class
{
  ...
private
  int _rating{};
std::unique_ptr<cudf::column> _column{};
}
 
TYPED_TEST_SUITE(RepeatTypedTestFixture, cudf::test::FixedWidthTypes);
 
TYPED_TEST(RepeatTypedTestFixture, RepeatScalarCount)
{
  ...
}

C++ 格式化通过 clang-format 强制执行。您应该在您的机器上配置 clang-format，使用 cudf/cpp/.clang-format 配置文件，并在提交之前对所有更改的代码运行 clang-format。最简单的方法是配置您的编辑器实现“保存时格式化”。

本文档中未讨论且无法自动强制执行的代码风格方面通常在代码评审期间发现，或者不强制执行。

C++ 指南

通常，我们建议遵循 C++ Core Guidelines。我们也建议观看 Sean Parent 的 C++ Seasoning 演讲，我们尽量遵循他的规则：“没有裸循环。没有裸指针。没有裸同步原语。”

• 优先使用 STL 和 Thrust 中的算法，而不是裸循环。
• 优先使用 libcudf 和 RMM 的 拥有数据结构和视图，而不是裸指针和裸内存分配。
• libcudf 没有很多 CPU 线程并发，但有一些。目前 libcudf 确实使用了裸同步原语。所以我们应该重新审视 Parent 的第三条规则并在此改进。

libcudf 代码的其他风格指南

• 优先使用“东侧 const”，将 const 放在类型之后。这不会被 clang-format 自动强制执行，因为选项 QualifierAlignment: Right 已被观察到会产生假阴性和假阳性。
• NL.11: Make Literals Readable：十进制值应每千位使用整数分隔符，如 1'234'567。十六进制值应每 4 个字符使用分隔符，如 0x0123'ABCD。

文档在文档指南中讨论。

包含 (Includes)

以下指南适用于组织 #include 行。

• 按库（例如 cuDF、RMM、Thrust、STL）对 includes 进行分组。clang-format 会尊重分组并在组内按字典顺序对单个 includes 进行排序。
• 用空行分隔组。
• 按“最近”到“最远”的顺序排列组。换句话说，先是本地 includes，然后是其他 RAPIDS 库的 includes，然后是相关库的 includes，如 <thrust/...>，然后是随 cuDF 安装的依赖项的 includes，最后是标准头文件（例如 <string>, <iostream>）。
• 我们使用 clang-format 自动分组和排序头文件。详见 cudf/cpp/.clang-format 文件。
• 对所有 includes 使用 <>，除了不在 include 目录中的内部头文件。换句话说，如果它是 cuDF 内部头文件（例如在 src 或 test 目录中），路径不会以 cudf 开头（例如 #include <cudf/some_header.hpp>），因此应该使用引号。示例：#include "io/utilities/hostdevice_vector.hpp"。
• cudf_test 和 nvtext 是 libcudf 仓库内的独立库。因此，它们在 include 中有公共头文件，应该使用 <> 包含。
• 像 clangd 这样的工具通常在可能时自动插入 includes，但它们通常会弄错分组和括号。请纠正引号或括号的使用，然后运行 clang-format 来纠正分组。
• 始终检查 includes 是否仅对包含它们的文件是必需的。尽量避免过度包含，尤其是在头文件中。移除代码时请仔细检查这一点。
• 尽量避免使用 .. 的相对路径。当包含来自与包含文件不在同一目录下的源路径中的（内部）头文件时，使用 .. 的路径是必需的，因为源路径不会随 -I 传递。
• 避免从非内部文件包含库内部头文件。例如，尽量不要在测试或 libcudf 公共头文件中包含 libcudf src 目录中的头文件。如果您发现自己这样做，请发起讨论，考虑将包含的内部头文件（部分）移动到公共头文件中。

libcudf 数据结构

libcudf 中的应用程序数据包含在列 (Columns) 和表 (Tables) 中，但在开发 libcudf 代码时，您还会用到各种其他数据结构。

视图和所有权

资源所有权是 libcudf 中的一个重要概念。简而言之，“拥有”对象拥有资源（例如设备内存），它在构造期间获取资源，并在析构期间释放资源（RAII）。“非拥有”对象不拥有资源。libcudf 中任何带有 *_view 后缀的类都是非拥有的。更多详情请参阅 libcudf 演示文稿。

libcudf 函数通常接受视图作为输入（column_view 或 table_view），并生成拥有对象的 unique_ptr 作为输出。例如，

std::unique_ptr<table> sort(table_view const& input);

内存资源 (Memory Resources)

libcudf 通过 RMM 内存资源 (MR) 或 CUDA MR 分配所有设备内存。这两种类型都可以通过 rmm::device_async_resource_ref 参数传递给 libcudf 函数。详见 RMM 文档。

当前设备内存资源

RMM 为每个设备提供一个“默认”内存资源以及访问和设置它的函数。libcudf 在 cpp/include/cudf/utilities/memory_resource.hpp 中提供了这些函数的包装器。所有内存资源参数都应默认使用 cudf::get_current_device_resource_ref() 的返回值。

资源引用 (Resource Refs)

内存资源通过资源引用参数传递。资源引用是内存资源包装器，使消费者能够指定他们期望的资源属性。这些定义在 libcu++ 的 cuda::mr 命名空间中，但 RMM 在 rmm/resource_ref.hpp 中提供了一些便利别名。

• rmm::device_resource_ref 接受一个提供设备可访问内存同步分配的内存资源。
• rmm::device_async_resource_ref 接受一个提供设备可访问内存流有序分配的内存资源。
• rmm::host_resource_ref 接受一个提供主机可访问内存同步分配的内存资源。
• rmm::host_async_resource_ref 接受一个提供主机可访问内存流有序分配的内存资源。
• rmm::host_device_resource_ref 接受一个提供主机和设备可访问内存同步分配的内存资源。
• rmm::host_async_resource_ref 接受一个提供主机和设备可访问内存流有序分配的内存资源。

更多信息请参阅 libcu++ 关于 resource_ref 的文档。

cudf::column

cudf::column 是 libcudf 中核心的拥有数据结构。大多数 libcudf 公共 API 都生成 cudf::column 或 cudf::table 作为输出。一个 column 包含 device_buffers，它们拥有列元素的设备内存以及可选的 null 指示位掩码。

可隐式转换为 column_view 和 mutable_column_view。

可移动和可复制。复制会执行列内容的深拷贝，而移动会将内容从一个列移动到另一个列。

示例

cudf::column col{...};
 
cudf::column copy{col}; // 复制 `col` 的内容
cudf::column const moved_to{std::move(col)}; // 从 `col` 移动内容
 
column_view v = moved_to; // 隐式转换为非拥有 column_view
// mutable_column_view m = moved_to; // 不能为 const column 创建 mutable view

一个 column 可以有嵌套（子）列，取决于列的数据类型。例如，LIST、STRUCT 和 STRING 类型的列。

cudf::column_view

cudf::column_view 是 libcudf 中核心的非拥有数据结构。它是一个不可变的、非拥有的设备内存作为列的视图。大多数 libcudf 公共 API 都接受视图作为输入。

一个 column_view 可以是一个列的“切片”视图。例如，它可以查看具有 1000 行的列的第 75-150 行。这个 column_view 的 size() 将是 75，访问该视图的索引 0 将返回拥有 column 中索引 75 处的元素。在内部，这是通过在视图中存储一个指针、一个偏移量和一个大小来实现的。column_view::data<T>() 返回一个指向 column_view::head<T>() + offset 的指针迭代器。

cudf::mutable_column_view

一个*可变的*、非拥有的设备内存作为列的视图。用于 detail API 和（罕见的）原地修改列的公共 API。

cudf::column_device_view

一个不可变的、非拥有的设备数据作为元素列的视图，可平凡复制并在 CUDA 设备代码中使用。用于将 column_view 数据作为输入传递给 CUDA 内核和设备函数（包括 Thrust 算法）。

cudf::mutable_column_device_view

一个可变的、非拥有的设备数据作为元素列的视图，可平凡复制并在 CUDA 设备代码中使用。用于将 column_view 数据传递给 CUDA 内核和设备函数（包括 Thrust 算法）在设备上进行修改。

cudf::table

拥有类，用于一组具有相同数量元素的 cudf::column。这相当于 C++ 中的数据框。

可隐式转换为 cudf::table_view 和 cudf::mutable_table_view

可移动和可复制。复制会执行所有列的深拷贝，而移动会将所有列从一个表移动到另一个表。

cudf::table_view

一个*不可变的*、非拥有的表视图。

cudf::mutable_table_view

一个*可变的*、非拥有的表视图。

cudf::size_type

cudf::size_type 是用于表示列中元素数量、列内元素的偏移量、访问特定元素的索引、列元素子集的段等类型。它等同于有符号的 32 位整数类型，因此最大值为 2147483647。某些 API 也接受负索引值，这些函数支持最小值为 -2147483648。这个基本类型不仅影响列大小限制的输出值，还影响元素计数的输出值。

Span (连续区间视图)

libcudf 提供了模仿 C++20 std::span 的 span 类，后者是连续对象序列的轻量级视图。libcudf 提供了两个类，host_span 和 device_span，它们可以从多种容器类型构造，也可以从指针（分别为主机或设备）和大小构造，或从迭代器构造。span 类型对于定义适用于多种输入容器类型的通用（内部）接口很有用。device_span 可以从 thrust::device_vector、rmm::device_vector 或 rmm::device_uvector 构造。host_span 可以从 thrust::host_vector、std::vector 或 std::basic_string 构造。

如果您正在定义操作向量的内部（detail）函数，请使用 spans 作为输入向量参数，而不是特定的向量类型，以使您的函数更具通用性。

当 span 指的是不可变元素时，请使用 span<T const>，而不是 span<T> const。由于 span 是轻量级视图，它不会传播 const 属性。因此，const 应该应用于模板类型参数，而不是 span 本身。此外，由于 span 是轻量级视图，它应该按值传递。libcudf 中接受 spans 作为输入的 API 将类似于以下将设备数据复制到主机 std::vector 的函数。

template <typename T>
std::vector<T> make_std_vector_async(device_span<T const> v, rmm::cuda_stream_view stream)

cudf::scalar

cudf::scalar 是一个对象，可以表示 cudf 当前支持的任何类型的一个单独的可空值。每种类型的值都由一个单独的 scalar 类类型表示，所有这些类都派生自 cudf::scalar。例如，numeric_scalar 持有一个数值，string_scalar 持有一个字符串。存储值的数据驻留在设备内存中。

一个 list_scalar 包含单个列表的底层数据。这意味着底层数据可以是 cudf 支持的任何类型。例如，表示整数列表的 list_scalar 存储一个类型为 INT32 的 cudf::column。表示整数列表的列表的 list_scalar 存储一个类型为 LIST 的 cudf::column，后者又存储一个类型为 INT32 的 column。

值类型	Scalar 类	备注
固定宽度	fixed_width_scalar<T>	T 可以是任何固定宽度类型
数值	numeric_scalar<T>	T 可以是 int8_t, int16_t, int32_t, int64_t, float 或 double
定点	fixed_point_scalar<T>	T 可以是 numeric::decimal32 或 numeric::decimal64
时间戳	timestamp_scalar<T>	T 可以是 timestamp_D, timestamp_s 等
时长	duration_scalar<T>	T 可以是 duration_D, duration_s 等
字符串	string_scalar	此类别对象是不可变的
列表	list_scalar	底层数据可以是 cudf 支持的任何类型

构造 (Construction)

scalars 可以使用其各自的构造函数创建，或使用工厂函数创建，例如 make_numeric_scalar()、make_timestamp_scalar() 或 make_string_scalar()。

转换 (Casting)

所有工厂方法都返回一个 unique_ptr<scalar>，需要先将其静态向下转型为其相应的 scalar 类类型，然后才能访问其值。无需转型即可访问其有效性（空值）。通常，需要从了解值类型的函数（例如从 type_dispatcher 分派的 functor）访问该值。要将给定值类型转型为必需的 scalar 类类型，请使用 type_dispatcher.hpp 中提供的映射实用工具 scalar_type_t。

//unique_ptr<scalar> s = make_numeric_scalar(...);
 
using ScalarType = cudf::scalar_type_t<T>;
// ScalarType 现在是 numeric_scalar<T>
auto s1 = static_cast<ScalarType *>(s.get());

传递到设备 (Passing to device)

每个 scalar 类型，除了 list_scalar，都有一个对应的非拥有设备视图类，允许从设备访问值及其有效性。这可以通过函数 get_scalar_device_view(ScalarType s) 获取。请注意，不为基础 scalar 对象提供设备视图，仅为派生类型 scalar 类对象提供。

list_scalar 的底层数据可以通过 view() 方法访问。对于非嵌套数据，设备视图可以通过函数 column_device_view::create(column_view) 获取。对于嵌套数据，可以通过 lists_column_device_view(column_device_view) 构造一个专门用于列表列的设备视图。

libcudf 策略和设计原则

libcudf 旨在提供线程安全、单 GPU 加速的算法原语，用于解决数据科学中出现的各种问题。API 被编写为在默认 GPU 上执行，调用者可以通过标准的 CUDA 设备 API 或环境变量（如 CUDA_VISIBLE_DEVICES）来控制。我们的目标是使 Spark 或 Pandas 等不同的用例能够从 GPU 的性能中受益，libcudf 依赖于 Spark 或 Dask 等更高级别的层来协调多 GPU 任务。

为了最好地满足这些用例，libcudf 优先考虑性能和灵活性，这有时可能会牺牲便利性。虽然我们欢迎用户直接使用 libcudf，但我们在设计时预计大多数用户将通过 Spark 或 cuDF Python 等更高级的层来使用 libcudf，这些层处理了 libcudf 直接用户必须自己处理的一些细节。我们在此记录这些策略及其背后的原因。

libcudf 不内省数据

libcudf API 通常不执行深度的内省和输入数据验证。原因有很多

1. 它违反了单一职责原则：验证与执行是分开的。
2. 由于 libcudf 数据结构将数据存储在 GPU 上，任何验证都将产生*至少*一次内核启动的开销，并且通常可能非常昂贵。
3. 围绕数据内省的 API 承诺通常会显著增加实现的复杂性。

因此，用户有责任将有效数据传递给此类 API。*请注意，此策略并不意味着 libcudf 根本不执行任何验证*。libcudf API 仍应执行任何不需要内省的验证。为了大致了解哪些应该或不应该验证，这里有一些（非详尽的）示例列表。

libcudf 应该验证的事项:

• 输入列/表的大小或数据类型

libcudf 不应该验证的事项:

• 整数溢出
• 确保给定输入集产生的输出不会超过2GB 大小限制

libcudf 要求嵌套类型具有已清理的 null 掩码

接受嵌套数据类型列（如 LIST 或 STRUCT）的各种 libcudf API 可能会假定这些列已经过清理。在此上下文中，清理指确保具有嵌套 dtype 的列中的 null 元素与嵌套列的元素兼容。具体来说

• 列表列的 null 元素也应该为空。null 元素的起始偏移量应等于结束偏移量。
• 结构体列的 null 元素也应该是底层结构体中的 null 元素。
• 对于复合列，null 只应存在于父列级别。子列不应包含 null。
• 对嵌套列的切片操作不会将偏移量传播到子列。

libcudf API *应该*承诺永远不会返回“脏”列，即包含未清理数据的列。因此，唯一的问题是用户构造了未正确清理的输入列，然后将其传递给 libcudf API。

将 libcudf API 视为异步 API

在主机上调用的 libcudf API 不保证在返回前流已同步。libcudf 中的工作发生在 cudf::get_default_stream().value 上，默认是 CUDA 默认流（流 0）。请注意，如果通过 CUDF_USE_PER_THREAD_DEFAULT_STREAM 启用每线程默认流，流 0 的行为会有所不同。提供给 libcudf 或由 libcudf 返回的、使用单独的非阻塞流的任何数据都需要与默认 libcudf 流同步，以确保流安全。

libcudf 通常不保证排序

libcudf 中的 merge 或 groupby 等函数不保证输出中条目的顺序。通常，保证确定性排序不利于快速并行算法。如果需要排序输出，调用代码负责在事后执行排序。

libcudf 不承诺特定的异常消息

libcudf 文档记录了 API 对不同类型的无效输入将抛出的异常。这些异常的类型（例如 cudf::logic_error）是公共 API 的一部分。然而，这些异常的 what 方法返回的解释性字符串不属于 API，并且可能会更改。调用代码不应依赖 libcudf 错误消息的内容来确定错误的性质。有关 libcudf 在不同情况下抛出的异常类型的信息，请参阅错误处理部分。

libcudf API 和实现

流 (Streams)

libcudf 正在添加使用 CUDA 流支持异步执行。为了方便使用流，所有分配设备内存或执行内核的新 libcudf API 都应在最后接受一个 rmm::cuda_stream_view 参数，其默认值为 cudf::get_default_stream()。这条规则有一个例外：如果 API 还接受内存资源参数，则流参数应放在内存资源*之前*。然后此 API 应将调用转发给具有相同签名的相应 detail API，但 detail API 不应有流的默认参数（detail API 应始终避免默认参数）。实现应完全包含在 detail API 定义中，并且只使用带有流参数的异步版本的 CUDA API。

为了使 detail API 可以从其他 libcudf 函数调用，它应该在放置在 cudf/cpp/include/detail/ 目录中的头文件中暴露。如果其他 libcudf 函数不调用 detail 函数，则不需要声明。

例如

// cpp/include/cudf/header.hpp
void external_function(...,
  rmm::cuda_stream_view 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(..., rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr)
} // namespace detail
 
// cudf/src/implementation.cpp
namespace detail{
// 在 detail 实现中使用 stream 参数。
void external_function(..., rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr){
  // 实现使用 stream 和异步 API。
  rmm::device_buffer buff(..., stream, mr);
  CUDF_CUDA_TRY(cudaMemcpyAsync(...,stream.value()));
kernel<<<..., stream>>>(...);
thrust::algorithm(rmm::exec_policy(stream), ...);
}
} // namespace detail
 
void external_function(..., rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr)
{
CUDF_FUNC_RANGE(); // 为此函数的生命周期生成一个 NVTX 范围。
detail::external_function(..., stream, mr);
}

注意： 仅在*必要时*才同步流非常重要。例如，当 API 返回一个非指针值，该值是异步设备到主机复制的结果时，用于复制的流应在返回前同步。然而，当返回一个 column 时，流不应同步，因为这样做会破坏异步性。

注意： 永远*不*应使用 cudaDeviceSynchronize()。这会限制使用 libcudf API 进行多流/多线程工作的能力。

流创建 (Stream Creation)

在实现 libcudf 功能时，有时在*内部*使用流会很有优势，例如，在实现算法时实现重叠。然而，动态创建流可能会很昂贵。RMM 有一个流池类来帮助避免动态流创建。但是，这尚未在 libcudf 中暴露，因此目前，libcudf 功能应避免创建流（即使效率略低）。留下一个 // TODO: 注释，指出在哪里使用流会更有益，是一个好主意。

内存分配 (Memory Allocation)

设备内存资源在 libcudf 中用于抽象和控制设备内存的分配方式。

输出内存

任何分配内存并*返回*给用户的 libcudf API 都必须接受一个 rmm::device_async_resource_ref 作为最后一个参数。在 API 内部，必须使用此内存资源为返回的对象分配任何内存。因此，应该将其传递给将返回其输出的函数。示例

// 返回的 `column` 包含新分配的内存，
// 因此 API 必须接受一个内存资源指针
std::unique_ptr<column> returns_output_memory(
..., rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref());
 
// 此 API 不分配任何新的 *输出* 内存，因此
// 不需要内存资源
void does_not_allocate_output_memory(...);

此规则自动适用于所有分配内存的 detail API。任何 detail API 都可以由任何公共 API 调用，因此可能会分配并返回给用户的内存。为了支持此类用例，所有分配内存资源的 detail API 都应接受一个 mr 参数。调用者负责根据需要传递提供的 mr 或 cudf::get_current_device_resource_ref()。

临时内存

并非所有在 libcudf API 内部分配的内存都返回给调用者。通常算法必须为中间结果分配临时、scratch 内存。对于临时内存分配，始终使用从 cudf::get_current_device_resource_ref() 获取的默认资源。示例

rmm::device_buffer some_function(
..., rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()) {
    rmm::device_buffer returned_buffer(..., mr); // 返回的 buffer 使用传入的 MR
    ...
    rmm::device_buffer temporary_buffer(...); // 临时 buffer 使用默认 MR
    ...
    return returned_buffer;
}

内存管理 (Memory Management)

libcudf 代码通常避免使用裸指针和直接内存分配。使用基于内存资源的 RMM 类进行设备内存分配，并实现自动生命周期管理。

rmm::device_buffer

使用内存资源分配指定数量的字节，作为无类型、未初始化的设备内存。如果未显式提供 rmm::device_async_resource_ref，则使用 cudf::get_current_device_resource_ref()。

rmm::device_buffer 是可移动的，并且可以在流上复制。复制在指定的流上对 device_buffer 的设备内存进行深拷贝，而移动将设备内存的所有权从一个 device_buffer 移动到另一个。

// 分配至少 100 字节的未初始化设备内存
// 使用指定的资源和流
rmm::device_buffer buff(100, stream, mr);
void * raw_data = buff.data(); // 指向底层设备内存的裸指针
 
// 在 `stream` 上将 `buff` 深拷贝到 `copy`
rmm::device_buffer copy(buff, stream);
 
// 将 `buff` 的内容移动到 `moved_to`
rmm::device_buffer moved_to(std::move(buff));
 
custom_memory_resource *mr...;
// 从 custom_memory_resource 分配 100 字节
rmm::device_buffer custom_buff(100, mr, stream);

rmm::device_scalar<T>

分配指定类型的单个元素，并初始化为指定值。将其用于设备内核的标量输入/输出，例如归约结果、空值计数等。这实际上是长度为 1 的 rmm::device_vector<T> 的便利包装器。

// 使用指定的资源和流为单个 int 分配设备内存
// 并将值初始化为 42
rmm::device_scalar<int> int_scalar{42, stream, mr};
 
// scalar.data() 返回设备内存中值的指针
kernel<<<...>>>(int_scalar.data(),...);
 
// scalar.value() 同步 scalar 的流并复制
// 值从设备到主机并返回值
int host_value = int_scalar.value();

rmm::device_vector<T>

分配指定数量的指定类型元素。如果未提供初始化值，所有元素将默认初始化（这会触发一次内核启动）。

注意：我们已从 libcudf 中移除了所有 rmm::device_vector 和 thrust::device_vector 的用法，在 libcudf 的新代码中，未经仔细考虑不应使用它们。相反，请使用 rmm::device_uvector 以及 device_factories.hpp 中的实用工厂函数。这些实用工具可以从主机端向量创建 uvectors，或者创建零初始化的 uvectors，使其使用起来与 device_vector 一样方便。避免使用 device_vector 有许多好处，详见关于 rmm::device_uvector 的下一节。

rmm::device_uvector<T>

类似于 device_vector，在设备内存中分配一组连续元素，但存在关键区别

• 作为一项优化，元素未初始化，且构造时不发生同步。这限制了类型 T 只能是可平凡复制的类型。
• 所有操作都是流有序的（即，它们接受一个 cuda_stream_view 指定执行操作的流）。这提高了使用非默认流时的安全性。
• device_uvector.hpp 不包含任何 __device__ 代码，与 thrust/device_vector.hpp 不同，这意味着 device_uvectors 可以在 .cpp 文件中使用，而不仅仅在 .cu 文件中使用。

cuda_stream s;
// 在流 `s` 上分配 100 个 `int32_t` 元素的未初始化存储空间，使用
// 默认资源
rmm::device_uvector<int32_t> v(100, s);
// 将元素初始化为 0
thrust::uninitialized_fill(thrust::cuda::par.on(s.value()), v.begin(), v.end(), int32_t{0});
 
auto mr = new my_custom_resource{...};
// 在流 `s` 上使用资源 `mr` 分配 100 个 `int32_t` 元素的未初始化存储空间
rmm::device_uvector<int32_t> v2{100, s, mr};

默认参数 (Default Parameters)

虽然 libcudf 公共 API 可以自由包含默认函数参数，但 detail 函数不应该这样做。默认内存资源参数使得开发者很容易意外地使用错误的资源分配内存。避免默认内存资源强制开发者仔细考虑每一次内存分配。

虽然流目前尚未在 libcudf 的 API 中暴露，但我们计划最终这样做。因此，关于内存资源的相同原因也适用于流。公共 API 默认使用 cudf::get_default_stream()。但是，在 detail API 中包含相同的默认值可能会导致开发者忘记传递用户提供的流（如果公共 API 接收到）。强制每个 detail API 调用显式传递流是为了防止此类错误。

内存资源（以及最终的流）是几乎所有公共 API 的最后参数。为了 API 的一致性，在 libcudf 的内部也是如此。因此，不允许默认流或 MRs 的结果是 detail API 中的任何参数都不能有默认值。

NVTX 范围

为了帮助性能优化和调试，所有计算密集型 libcudf 函数都应有相应的 NVTX 范围。在当前作用域声明 NVTX 范围时，请选择使用 CUDF_FUNC_RANGE 或 cudf::scoped_range

• 如果您想使用函数名作为 NVTX 范围的名称，请使用 CUDF_FUNC_RANGE() 宏
• 使用 cudf::scoped_range rng{"custom_name"}; 为当前作用域的 NVTX 范围提供一个自定义名称

有关 NVTX 的更多信息，请参见此处。

输入/输出风格

输入参数的传递方式和输出结果的返回方式的首选风格如下

• 输入
  • 列
    • column_view const&
  • 表
    • table_view const&
  • 标量
    • scalar const&
  • 其他所有类型
    • Trivial 或拷贝开销小的类型
      • 按值传递
    • Non-trivial 或拷贝开销大的类型
      • 按 const& 传递
• 输入/输出
  • 列
    • mutable_column_view&
  • 表
    • mutable_table_view&
  • 其他所有类型
    • 通过原始指针传递
• 输出
  • 输出应被 返回，即没有输出参数
  • 列
    • std::unique_ptr<column>
  • 表
    • std::unique_ptr<table>
  • 标量
    • std::unique_ptr<scalar>

多返回值

有时函数需要有多个输出。在 C++ 中有几种方法可以实现（包括为输出创建一个 struct）。一种方便的方法是使用 std::tie 和 std::pair。注意，传递给 std::pair 的对象会调用对象的拷贝构造函数或移动构造函数，对于 non-trivially copyable 对象（以及对于拷贝构造函数被删除的类型，如 std::unique_ptr，这是必需的），最好是移动它们。

std::pair<table, table> 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();

注意：如果不是因为 Cython 不支持 std::tuple，则 std::tuple *可以*被使用。因此，libcudf API 必须使用 std::pair，因此只能返回两个不同类型的对象。多个相同类型的对象可以通过 std::vector<T> 返回。

此外，使用 C++17（从 cudf v0.20 开始支持），可以使用 结构化绑定来分解多个返回值

auto [out0, out1] = cudf::return_two_outputs();

请注意，编译器可能不支持在 lambda 中捕获结构化绑定中定义的别名。可以通过使用带初始化的捕获来解决此问题

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.
  // ...
};

基于迭代器的接口

libcudf 越来越倾向于使用带迭代器参数而不是显式 column/table/scalar 参数的内部 (detail) API。与 STL 一样，迭代器使得通用算法可以应用于任意容器。一个很好的例子是 cudf::copy_if_else。此函数接受两个输入和一个布尔掩码。它根据该索引处的掩码值是 true 还是 false，从第一个或第二个输入中选择相应的元素进行拷贝。通过在 detail API 中使用迭代器，简化了为所有 column 和 scalar 参数组合实现 copy_if_else 的工作。

template <typename FilterFn, typename LeftIter, typename RightIter>
std::unique_ptr<column> copy_if_else(
  bool nullable,
LeftIter lhs_begin,
LeftIter lhs_end,
RightIter rhs,
FilterFn filter,
  ...);

LeftIter 和 RightIter 只需实现迭代器所需的接口。libcudf 提供了许多迭代器类型和实用工具，这些工具在 libcudf 的基于迭代器的 API 以及 Thrust 算法中非常有用。大多数都在 include/detail/iterator.cuh 中定义。

Pair 迭代器

pair 迭代器用于将可空列的元素作为包含元素值和有效性的 pair 进行访问。cudf::detail::make_pair_iterator 可用于从 column_device_view 或 cudf::scalar 创建一个 pair 迭代器。make_pair_iterator 不适用于 mutable_column_device_view。

Null-replacement 迭代器

此迭代器将每个元素的 null/validity 值替换为指定的常量 (true 或 false)。使用 cudf::detail::make_null_replacement_iterator 创建。

Validity 迭代器

此迭代器返回底层元素的有效性 (true 或 false)。使用 cudf::detail::make_validity_iterator 创建。

索引归一化迭代器

libcudf 支持的数据类型激增可能导致编译时间延长。其中一个编译时间问题出现在用于存储索引的类型中，这些索引可以是任何整数类型。 "indexalator"，即索引归一化迭代器 (include/cudf/detail/indexalator.cuh)，可用于索引类型（整数），而无需特定于类型的实例。它可用于读取 int8、int16、int32、int64、uint8、uint16、uint32 或 uint64 类型的整数数组的任何迭代器接口。读取特定元素总是返回 cudf::size_type 整数。

使用 indexalator_factory 从 column_view 创建适当的输入迭代器。输入迭代器示例用法

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 );

输出迭代器示例用法

auto result_itr = indexalator_factory::create_output_iterator(indices->mutable_view());
thrust::lower_bound(rmm::exec_policy(stream),
input->begin<Element>(),
input->end<Element>(),
values->begin<Element>(),
values->end<Element>(),
result_itr,
thrust::less<Element>());

偏移归一化迭代器

与 indexalator 类似，"offsetalator"，或偏移归一化迭代器 (include/cudf/detail/offsetalator.cuh)，可用于偏移列类型（仅限 INT32 或 INT64）而无需特定于类型的实例。这在读取或构建 字符串列时很有帮助。归一化类型是 int64，这意味着 input_offsetsalator 对于 INT32 和 INT64 偏移列都将返回 int64 类型值。同样，output_offselator 可以接受 int64 类型值，并将其存储到适当创建的 INT32 或 INT64 输出偏移列中。

使用 cudf::detail::offsetalator_factory 从偏移 column_view 创建适当的输入或输出迭代器。输入迭代器示例用法

// 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

输出迭代器示例用法

// 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

命名空间

外部

所有公共 libcudf API 都应放在 cudf 命名空间中。例如

namespace cudf{
   void public_function(...);
} // namespace cudf

顶级 cudf 命名空间对于大多数公共 API 已经足够。但是，为了逻辑地分组一组广泛的函数，可以使用进一步的命名空间。例如，有许多特定于字符串列的函数。这些函数位于 cudf::strings:: 命名空间中。类似地，仅用于单元测试的功能位于 cudf::test:: 命名空间中。

公共函数预计包含一个对 CUDF_FUNC_RANGE() 的调用，后跟一个对与公共函数具有相同名称和参数的 detail 函数的调用。有关此模式的示例，请参阅 Streams 部分。

内部

许多函数不适用于公共使用，因此根据情况将其放在 detail 或*匿名*命名空间中。

detail 命名空间

将在*多个*翻译单元（即源文件）中使用的函数或对象，应在内部头文件中公开并放在 detail 命名空间中。例如

// some_utilities.hpp
namespace cudf{
namespace detail{
void reusable_helper_function(...);
} // namespace detail
} // namespace cudf

匿名命名空间

仅在*单个*翻译单元中使用的函数或对象，应在使用它的源文件中定义在*匿名*命名空间中。例如

// some_file.cpp
namespace{
void isolated_helper_function(...);
} // anonymous namespace

在头文件中*永远不*应使用匿名命名空间。

弃用和删除代码

libcudf 不断发展以提高性能并更好地满足用户的需求。因此，我们偶尔需要修改或完全删除 API，以响应对我们提供的功能的新的和改进的理解。保持这种自由对于使 libcudf 成为能够快速适应用户需求的敏捷库至关重要。因此，我们并非总是在发布破坏性更改之前提供警告或任何提前通知。在尽力而为的基础上，libcudf 团队将通知用户我们预期会产生重大或广泛影响的更改。

如果可能，在删除 API 之前，使用 deprecated 属性指示待删除的 API，并使用 Doxygen 的 deprecated 命令进行文档记录。当弃用 API 有可用的替代 API 时，请在弃用消息和弃用文档中提及替代项。引入弃用的拉取请求应标记为 "deprecation"，以便于在后续版本中发现和删除。

通过将任何破坏或删除现有 API 的拉取请求标记为 "breaking" 来宣传破坏性更改。这确保了发行说明的 "Breaking" 部分包含对与之前版本相比发生破坏的描述。将包含弃用的拉取请求标记为 "non-breaking"。

错误处理

libcudf 遵循约定（并提供实用工具），用于强制执行编译时和运行时条件，以及检测和处理 CUDA 错误。错误通信始终通过 C++ 异常进行。

运行时条件

使用 CUDF_EXPECTS 宏来强制执行正确执行所需的运行时条件。

示例用法

CUDF_EXPECTS(cudf::have_same_types(lhs, rhs), "Type mismatch", cudf::data_type_error);

第一个参数是在正常条件下预计解析为 true 的条件表达式。CUDF_EXPECTS 的第二个参数是发生的错误的简短描述，用作异常的 what() 消息。如果条件表达式评估为 false，则表示发生错误，并抛出第三个参数中指定的异常类（或默认的 cudf::logic_error）的实例。

有时，如果达到某个特定的代码路径，无论如何都应指示错误。例如，switch 语句的 default 情况通常表示无效的备选项。对于此类错误，请使用 CUDF_FAIL 宏。这实际上等同于调用 CUDF_EXPECTS(false, reason)。

示例

CUDF_FAIL("This code path should not be reached.");

CUDA 错误检查

使用 CUDF_CUDA_TRY 宏检查 CUDA 运行时 API 函数是否成功完成。如果 CUDA API 返回值不是 cudaSuccess，此宏将抛出 cudf::cuda_error 异常。抛出的异常在其 what() 消息中包含 CUDA 错误代码的描述。

示例

CUDF_CUDA_TRY( cudaMemcpy(&dst, &src, num_bytes) );

编译时条件

使用 static_assert 强制执行编译时条件。例如：

template <typename T>
void trivial_types_only(T t){
static_assert(std::is_trivial<T>::value, "This function requires a trivial type.");
...
}

日志记录

libcudf 包含日志记录实用工具（构建在 spdlog 库之上），应使用这些工具记录重要事件（例如用户警告）。此实用工具也可用于记录调试信息，只要使用正确的日志级别。有六个宏可用于在不同级别进行日志记录

• CUDF_LOG_TRACE - 详细调试消息（面向开发者）
• CUDF_LOG_DEBUG - 调试消息（面向开发者）
• CUDF_LOG_INFO - 关于正常执行期间发生的罕见事件（例如每次运行一次）的信息
• CUDF_LOG_WARN - 关于潜在意外行为或弃用的用户警告
• CUDF_LOG_ERROR - 可恢复错误
• CUDF_LOG_CRITICAL - 不可恢复错误（例如内存损坏）

默认情况下，TRACE、DEBUG 和 INFO 消息不包含在日志中。此外，在公共构建中，以 TRACE 和 DEBUG 级别记录的代码会被编译掉。这可以防止记录可能出于调试目的而记录的潜在敏感数据。此外，这也允许开发人员在跟踪/调试日志中包含昂贵的计算，因为在公共构建中不会产生开销。启用的最低日志级别是 WARN，可以通过多种方式进行修改

• CMake 配置变量 LIBCUDF_LOGGING_LEVEL - 设置构建中将编译的最低日志级别。可用级别包括 TRACE、DEBUG、INFO、WARN、ERROR、CRITICAL 和 OFF。
• 环境变量 LIBCUDF_LOGGING_LEVEL - 在初始化期间设置最低日志级别。如果此设置高于编译时 CMake 变量，则介于两个设置之间的任何日志级别将从写入的日志中排除。可用级别与 CMake 变量相同。
• 通过 cudf::logger() 公开的全局 logger 对象 - 在运行时设置最低日志级别。例如，调用 cudf::default_logger().set_level(level_enum::err) 将排除所有不是错误或严重错误的消息。此 API 不应在 libcudf 内部用于操作日志记录，其目的是允许上游用户配置 libcudf 日志记录以适应其应用程序。

默认情况下，日志消息输出到 stderr。设置环境变量 LIBCUDF_DEBUG_LOG_FILE 将日志重定向到指定路径的文件（可以是相对于当前目录）。上游用户还可以操作 cudf::default_logger().sinks() 来添加 sinks 或将日志重定向到标准输出。

数据类型

列可以包含多种类型的数据（参见 include/cudf/types.hpp 中的 enum class type_id）

• 数值数据：有符号和无符号整数（8 位、16 位、32 位或 64 位），浮点数（32 位或 64 位），以及布尔值（8 位）。
• 分辨率为天、秒、毫秒、微秒或纳秒的时间戳数据。
• 分辨率为天、秒、毫秒、微秒或纳秒的持续时间数据。
• 十进制定点数据（32 位或 64 位）。
• 字符串
• 字典
• 任意类型的列表
• 任意类型列的结构体

大多数算法必须支持任何数据类型的列。这导致代码复杂性增加，也是 libcudf 开发者面临的主要挑战之一。有时我们通过逐步增加对更多数据类型的支持来开发新算法，以简化这一过程。通常我们从固定宽度的数据类型（如数值类型和时间戳/持续时间）开始，之后再增加对嵌套类型的支持。

针对不同类型启用不同的算法可以使用模板特化或 SFINAE，如 Specializing Type-Dispatched Code Paths 中所讨论的。

比较数据类型

在比较两列或标量的数据类型时，不要直接比较 a.type() == b.type()。如果只比较顶层类型，则无法正确处理嵌套类型，例如整数列表的结构体。请改用 cudf::have_same_types 函数。

类型分发器

libcudf 将数据（用于列和标量）以 void* 设备内存的形式“类型擦除”存储。这种*类型擦除*使得它能够与其他语言和类型系统（如 Python 和 Java）互操作。为了确定类型，libcudf 算法必须使用存储在列 type() 中的运行时信息，通过将 void* 强制转换为适当的 T* 来重建数据类型 T。

这种所谓的*类型分发*在整个 libcudf 中普遍存在。type_dispatcher 是一个核心实用工具，它自动化了将 data_type 中的运行时类型信息映射到具体 C++ 类型的过程。

概括地说，您调用带有 data_type 和一个带有 operator() 模板的函数对象（也称为*仿函数*）的 type_dispatcher。根据 data_type::id() 的值，类型分发器调用相应的 operator() 模板实例化。

这个简化示例展示了 data_type::id() 的值如何决定调用哪个 F::operator() 模板实例化。

template <typename F>
void type_dispatcher(data_type t, F f){
    switch(t.id())
       case type_id::INT32: f.template operator()<int32_t>()
       case type_id::INT64: f.template operator()<int64_t>()
       case type_id::FLOAT: f.template operator()<float>()
       ...
}

以下示例显示了一个名为 size_of_functor 的函数对象，它返回分发类型的 size。

struct size_of_functor{
  template <typename T>
  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

默认情况下，type_dispatcher 使用 cudf::type_to_id<t> 提供 cudf::type_id 到分发 C++ 类型的映射。但是，可以通过显式指定用户定义的 trait 作为 IdTypeMap 来定制此映射。例如，对于所有 cudf::type_id 值都分发 int32_t

template<cudf::type_id t> struct always_int{ using type = int32_t; }
 
// This will always invoke `operator()<int32_t>`
cudf::type_dispatcher<always_int>(data_type, f);

避免多次类型分发

如果可能，避免多次类型分发。编译器为每种分发类型创建一个代码路径，因此第二级类型分发会导致编译时间和目标代码大小呈二次增长。作为一个拥有许多类型和函数的大型库，我们一直在努力减少编译时间和代码大小。

特化类型分发的代码路径

通常需要为不同的类型自定义分发的 operator()。这可以通过几种方式完成。

第一种方法是使用显式完全模板特化。这对于特化单个类型的行为很有用。以下示例函数对象在用 int32_t 或 double 调用时打印 "int32_t" 或 "double"，否则打印 "unhandled type"。

struct type_printer {
template <typename ColumnType>
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()<int32_t>() { std::cout << "int32_t\n"; }
 
template <>
void type_printer::operator()<double>() { std::cout << "double\n"; }

第二种方法是使用 SFINAE 和 std::enable_if_t。这对于对一组具有共同特征的类型进行部分特化很有用。以下示例仿函数分别对整数类型和浮点类型打印 integral 或 floating point。

struct integral_or_floating_point {
template <typename ColumnType,
std::enable_if_t<not std::is_integral<ColumnType>::value and
not std::is_floating_point<ColumnType>::value>* = nullptr>
void operator()() { std::cout << "neither integral nor floating point\n"; }
 
template <typename ColumnType,
std::enable_if_t<std::is_integral<ColumnType>::value>* = nullptr>
void operator()() { std::cout << "integral\n"; }
 
template < typename ColumnType,
std::enable_if_t<std::is_floating_point<ColumnType>::value>* = nullptr>
void operator()() { std::cout << "floating point\n"; }
};

有关 SFINAE 与 std::enable_if 的更多信息，请参见此文章。

include/cudf/utilities/traits.hpp 中定义了许多 trait，可用于分发函数对象的部分特化。例如，is_numeric<T>() 可用于特化任何数值类型。

可变大小和嵌套数据类型

libcudf 支持多种可变大小和嵌套数据类型，包括字符串、列表和结构体。

• string：简单地说就是一个字符字符串，但字符串列中的每行可能包含不同长度的字符串。
• list：任意类型元素的列表，因此一个整数列表列的每一行都包含一个整数列表，其长度可能不同。
• struct：在一个结构体列中，每一行都是一个结构体，包含一个或多个字段。这些字段以结构体数组格式存储，因此结构体列为结构体的每个字段都有一个嵌套列。

如标题所示，列表列和结构体列可以任意嵌套。可以创建列表结构体列，其中结构体的字段可以是任意类型，包括字符串、列表和结构体。即使有经验，考虑深度嵌套的数据类型对于基于列的数据也可能令人困惑。因此，仔细编写算法，并对其进行良好的测试和文档记录非常重要。

列表列

为了表示可变宽度的元素，libcudf 列包含一个子列向量。对于列表列，父列的类型为 LIST 且不包含数据，但其大小表示列中的列表数量，其空值掩码表示每个列表元素的有效性。父列有两个子列。

1. 一个非空值列，包含 size_type 元素，指示密集元素列中每个列表的起始偏移量。
2. 一个列，包含所有列表的所有元素的实际数据和可选的空值掩码，它们被打包在一起。

使用这种表示方式，data[offsets[i]] 是列表 i 的第一个元素，而列表 i 的大小由 offsets[i+1] - offsets[i] 给出。

注意，数据可以是任何类型，因此数据列本身可以是任何类型的嵌套列。另请注意，不仅每个列表可以是可空的（使用父列表的空值掩码），而且每个列表元素也可以是可空的。因此，您可能有一个列表列，其中第 3 行是空值，同时第 4 行的第 2 个元素也是空值。

列表列的底层数据总是打包到层次结构的最低层的单个叶子列中（忽略结构体，它在概念上“重置”了层次结构的根），无论嵌套级别如何。因此，一个 List<List<List<List<int>>>> 列在最底层有一个单独的 int 列。以下是对此的可视化表示。

lists_column = { {{{1, 2}, {3, 4}}, NULL}, {{{10, 20}, {30, 40}}, {{50, 60, 70}, {0}}} }
 
List<List<List<int>>> (2 行)
长度 : 2
偏移量 : 0, 2, 4
子列
List<List<int>>
长度 : 4
偏移量 : 0, 2, 2, 4, 6
空值计数: 1
        1101
子列
List<int>
长度 : 6
偏移量 : 0, 2, 4, 6, 8, 11, 12
子列
Int 列
          1, 2, 3, 4, 10, 20, 30, 40, 50, 60, 70, 0

这与 Arrow 的“可变大小列表”内存布局有关。

字符串列

字符串表示为一个列，其中包含一个数据设备缓冲区和一个子偏移列。父列的类型为 STRING，其数据包含所有字符串的所有字符打包在一起，但其大小表示列中的字符串数量，其空值掩码表示每个字符串的有效性。

字符串列包含一个单独的非空值子列，该子列包含偏移量元素，指示所有字符的密集数据缓冲区中每个字符串起始的字节位置偏移量。使用这种表示方式，data[offsets[i]] 是字符串 i 的第一个字符，而字符串 i 的大小由 offsets[i+1] - offsets[i] 给出。下图展示了这种复合列表示字符串的示例。

偏移列的类型取决于数据缓冲区中的字节数，可以是 INT32 或 INT64。有关处理单个字符串行（string row）的更多信息，请参见 cudf::strings_view。

结构体列

结构体是一种嵌套数据类型，包含一组子列，每个子列代表逻辑结构体的一个独立字段。字段名称不予表示。

具有 N 个字段的结构体列有 N 个子列。每个子列都是一个列，存储着单个字段的所有数据，这些数据按列打包，并且可选地带有一个空值掩码。父列的类型是 STRUCT，不包含数据，其大小表示列中的结构体行数，其空值掩码表示每个结构体元素的有效性。

使用这种表示方式，child[0][10] 是结构体第一个字段的第 10 行，child[1][42] 是结构体第二个字段的第 42 行。

请注意，除了结构体列的空值掩码外，每个结构体字段列都有自己的可选空值掩码。结构体字段的有效性可以独立于相应的结构体行。例如，一个非空的结构体行可能有一个空值字段。然而，空结构体行的字段也被视为空值。例如，考虑一个类型为 STRUCT<FLOAT32, INT32> 的结构体列。如果内容是 [ {1.0, 2}, {4.0, 5}, null, {8.0, null} ]，则结构体列的布局如下。（注意，空值掩码应从右向左读取。）

{
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
    }
  }
}

最后一个结构体行（索引 3）不是空值，但在 INT32 字段中有一个空值。此外，结构体列的第 2 行是空值，使其对应的字段也为空值。因此，在两个结构体字段的空值掩码中，位 2 未设置。

字典列

字典提供了一种有效的方式来表示低基数数据，通过存储每个值的单个副本。字典包含一个已排序的键列和一个索引列，该索引列为父列的每一行包含一个指向键列的索引。键列可以是任何 libcudf 数据类型，例如数值类型或字符串。索引表示每个元素值在键中的相应位置。索引子列可以是任何无符号整数类型（UINT8、UINT16、UINT32 或 UINT64）。

嵌套列挑战

嵌套列的第一个挑战是，实际上不可能原地修改任何字符串或列表的长度。例如，考虑尝试在每个字符串末尾添加字符 'a'。这需要动态调整字符列的大小，以允许在每个字符串末尾插入 'a'，然后修改偏移列以指示每个元素的新大小。因此，对列中的字符串或列表进行修改的每个操作都必须异地完成（out-of-place）。

第二个挑战是，在对字符串列进行异地操作时，与固定宽度元素不同，输出的大小是无法*先验*知道的。例如，考虑分散（scattering）到字符串列中

destination:    {"this", "is", "a", "column", "of", "strings"}
scatter_map:    {1, 3, 5}
scatter_values: {"red", "green", "blue"}

result:         {"this", "red", "a", "green", "of", "blue"}

在此示例中，字符串 "red", "green", 和 "blue" 将分别分散到 destination 的位置 1, 3, 和 5。回想一下上面提到的，此操作不能原地完成，因此 result 将通过选择性地从 destination 和 scatter_values 中拷贝字符串来生成。请注意，result 的字符子列需要 19 个字符的存储空间。然而，无法提前知道 result 需要 19 个字符。因此，大多数生成新的字符串输出列的操作采用两阶段方法

1. 确定结果中每个字符串的数量和大小。这相当于实例化输出偏移列。
2. 为所有输出字符分配足够的存储空间并实例化每个输出字符串。

在分散操作中，第一阶段使用 scatter_map 来确定输出中的字符串 i 将来自 destination 还是 scatter_values，并使用相应的 size 来实例化偏移列并确定输出的大小。然后，在第二阶段，为输出的字符分配足够的存储空间，然后用来自 destination 或 scatter_values 的相应字符串填充这些字符。

嵌套类型视图

libcudf 为嵌套列类型及其中的数据元素提供了视图类型。

cudf::strings_column_view 和 cudf::string_view

cudf::strings_column_view 封装了一个字符串列，包含一个作为字符串列视图的父 cudf::column_view 和一个作为父列子列的偏移 cudf::column_view。父视图包含字符串列的偏移量、大小和有效性掩码。偏移视图是非空值，offset() 为 0，并有自己的大小。由于偏移列类型可以是 INT32 或 INT64，因此使用偏移归一化迭代器 offsetalator 访问单个偏移值非常有用。

cudf::string_view 是单个字符串的视图，因此它是 STRING 类型 cudf::column 的数据类型，就像 int32_t 是 INT32 类型 cudf::column 的数据类型一样。顾名思义，这是一个只读的对象实例，指向字符串列内部的设备内存。其生命周期与它所查看的列相同（或更短）。单个字符串列行和 cudf::string_view 被限制为 size_type 字节。

使用 column_device_view::element 方法访问单个行元素。与其他列一样，不要在空值行上调用 element()。

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<string_view>(row_index);
   ...
}

空字符串与空字符串不同。如果您需要一个类对象实例来表示空字符串，请使用 cudf::string_scalar 类。

cudf::string_view 包含比较运算符 <,>,==,<=,>=，可以在许多 cudf 函数（如 sort）中使用，而无需字符串特定代码。cudf::string_view 实例的数据要求是 UTF-8 编码，所有运算符和方法都期望此编码。除非另有说明，位置和长度参数以字符为单位指定，而不是字节。该类还包含一个 cudf::string_view::const_iterator，可用于在字符串内遍历单个字符。

在对 STRING 列调用时，cudf::type_dispatcher 会分发到 cudf::string_view 数据类型。

UTF-8

libcudf 字符串列仅支持 UTF-8 编码的字符串数据。UTF-8 是一种变长字符编码，其中每个字符可以是 1-4 字节。这意味着字符串的长度与其字节大小不同。因此，对于大多数操作，建议使用 cudf::string_view 类来访问这些字符。

cudf/strings/detail/utf8.hpp 头文件还包含一些用于从字节数组读取和写入（to_char_utf8/from_char_utf8）单个 UTF-8 字符的实用方法。

cudf::lists_column_view 和 cudf::lists_view

cudf::lists_column_view 是列表列的视图。cudf::list_view 是单个列表的视图，因此 cudf::list_view 是 LIST 类型 cudf::column 的数据类型。

在对 LIST 列调用时，cudf::type_dispatcher 会分发到 list_view 数据类型。

cudf::structs_column_view 和 cudf::struct_view

cudf::structs_column_view 是结构体列的视图。cudf::struct_view 是单个结构体的视图，因此 cudf::struct_view 是 STRUCT 类型 cudf::column 的数据类型。

在对 STRUCT 列调用时，cudf::type_dispatcher 会分发到 struct_view 数据类型。

空列

libcudf 列支持空的、类型化的内容。这些列没有数据，也没有有效性掩码。空的字符串或列表列可能包含或不包含子偏移列。访问空的字符串或列表列的偏移子列是未定义行为（UB）。列表和结构体等嵌套列可能需要其他子列来提供空类型的嵌套结构。

使用 cudf::make_empty_column() 创建固定宽度和字符串列。使用 cudf::empty_like() 从现有 cudf::column_view 创建空列。

cuIO: 文件读写

cuIO 是 libcudf 的一个组件，提供 GPU 加速的数据文件格式读写功能，这些格式常用于数据分析，包括 CSV、Parquet、ORC、Avro 和 JSON_Lines。

// TODO: add more detail and move to a separate file.

调试技巧

以下是一些可以帮助调试 libcudf 的工具（当然除了 printf）

1. cuda-gdb\ 按照 cuDF 贡献者指南中的说明，使用调试符号构建和运行 libcudf。
2. compute-sanitizer\ CUDA Compute Sanitizer 工具可用于定位许多 CUDA 报告的错误，即使在非调试构建中也能提供接近错误发生位置的调用堆栈。sanitizer 包括各种工具，如 memcheck、racecheck 和 initcheck 等等。racecheck 和 initcheck 已知会产生误报。
3. cudf::test::print()\ 可以在 gtest 中调用 print() 实用工具来输出 cudf::column_view 中的数据。更多信息可在 测试指南 中找到。
4. GCC Address Sanitizer\ 通过添加 -fsanitize=address 编译器标志也可以使用 GCC ASAN。存在与 CUDA 运行时兼容性问题，可以通过在运行可执行文件之前设置环境变量 ASAN_OPTIONS=protect_shadow_gap=0 来解决。注意，CUDA compute-sanitizer 也可以与 GCC ASAN 一起使用，通过设置环境变量 ASAN_OPTIONS=protect_shadow_gap=0,alloc_dealloc_mismatch=0。