本文档为 libcuspatial C++ 代码贡献者提供指南。开发者还应参考以下附加文件，以进一步了解 libcuspatial 的最佳实践文档。

文档指南，包含 libcuspatial 代码文档化的指导原则。
测试指南，包含编写单元测试的指导原则。
基准测试指南，包含编写单元基准测试的指导原则。
仅头文件 API 指南，包含仅头文件 API 和基于列的 API 的指导原则。

概述

libcuspatial 是一个 C++ 库，提供 GPU 加速的数据并行算法，用于处理地理空间和时空数据。libcuspatial 提供了各种空间关系算法，包括距离计算、包含性（例如，点在多边形内测试）、边界框计算和空间索引。

libcuspatial 目前有两种接口。第一种是基于 libcudf 数据类型（CUDA Dataframe 库的 C++ API）的 C++ API。在本文档中，我们将其称为“基于列的 API”。基于列的 API 将空间数据表示为类型擦除列的表格。

第二种 API 是 cuSpatial 仅头文件 C++ API，它独立于 libcudf，并将数据表示为结构数组（例如，二维点）。仅头文件 API 使用迭代器进行输入和输出，风格类似于 C++ 标准模板库 (STL) 和 Thrust。

术语表

本节定义了 libcuspatial 中使用的术语。对于 libcudf 特定的术语，例如 Column、Table 等，请参阅 libcudf 开发者指南。

待办事项：添加术语

目录结构和文件命名

外部/公共 libcuspatial API 根据功能分组，存放在 cuspatial/cpp/include/cuspatial/ 目录下名称恰当的头文件中。例如，cuspatial/cpp/include/cuspatial/distance.hpp 包含与距离计算相关的公共 API 函数的声明。请注意，.hpp 文件扩展名用于指示可从 .cpp 源文件包含的 C++ 头文件。

头文件应使用 #pragma once 包含守卫。

包含实现 API 的源文件的文件夹应与 API 头文件的名称保持一致。例如，cuspatial/cpp/include/cuspatial/trajectory.hpp 中的 API 实现位于 cuspatial/cpp/src/trajectory 中。这条规则显然不适用于仅头文件 API，因为头文件就是源文件。

同样，单元测试和基准测试也位于与 API 头文件名称对应的文件夹中，例如 distance.hpp 的测试位于 cuspatial/cpp/tests/distance/ 中，基准测试位于 cuspatial/cpp/benchmarks/distance/ 中。

包含 detail 命名空间定义的内部 API 头文件，这些定义在 libcuspatial 内部的翻译单元之间使用，应放在 include/cuspatial/detail 中。

仅头文件 API 文件和基于列的 API 头文件一起存储在 include/cuspatial 中。前者使用 .cuh 扩展名，因为它们几乎普遍需要 CUDA 编译。后者使用 .hpp 扩展名，因为它们可以使用标准 C++ 编译器编译。

文件扩展名

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

仅在必要时使用 .cu 和 .cuh。一个好的指标是包含了只有 nvcc 才能识别的 __device__ 和其他符号。另一个指标是使用了设备执行策略 (在 libcuspatial 中始终是 rmm::exec_policy) 的 Thrust 算法 API。

代码和文档风格及格式

libcuspatial 代码对所有名称使用 snake_case (下划线分隔式)，少数情况除外：模板参数、单元测试和测试用例名称可以使用 Pascal case（帕斯卡命名法），也称为 UpperCamelCase (大驼峰命名法)。我们不使用 Hungarian notation (匈牙利命名法)，除非有时命名设备数据变量及其对应的主机副本（例如 d_data 和 h_data）。私有成员变量通常以前导下划线开头。

示例

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<rmm::device_uvector> _data{};
}
 
using TestTypes = ::testing::Types<float, double>;
 
TYPED_TEST_SUITE(RepeatTypedTestFixture, TestTypes);
 
TYPED_TEST(RepeatTypedTestFixture, RepeatScalarCount)
{
  ...
}

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

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

C++ 指南

通常，我们建议遵循 C++ Core Guidelines (C++ 核心指南)。我们还建议观看 Sean Parent 的 C++ Seasoning 演讲，并努力遵循他的规则：“没有裸循环 (No raw loops)。没有裸指针 (No raw pointers)。没有裸同步原语 (No raw synchronization primitives)。”在可能的情况下，我们还添加了第四条规则：“没有裸内核 (No raw kernels)。”

优先使用 STL 和 Thrust 中的算法，而非裸循环。
优先使用 Thrust 算法，而非裸内核。
对于设备存储，优先使用 libcudf 和 RMM 拥有的数据结构和视图，而非裸指针和裸内存分配。使用指针时，优先使用智能指针（例如 std::shared_ptr 和 std::unique_ptr），而非裸指针。
优先将内核分派到流中，而非显式同步。

文档在文档指南中讨论。

循环和网格步进循环

如上所述，尽可能优先使用算法而非裸循环。然而，并非总是能够避免裸循环。C++ 基于范围的 for 循环可以使裸循环更清晰，cuSpatial 为此目的使用了 Ranger。Ranger 提供带有可传递给基于范围 for 循环的迭代器的范围助手。特别重要的是 ranger::grid_stride_range()，它可用于使用 CUDA 网格的所有线程并行迭代范围。

编写自定义内核时，网格步进范围有助于确保内核适应各种网格形状，特别是在总线程数少于数据项数时。与此相反，

__global__ void foo(int n, int* data) {
  auto const idx = threadIdx.x + blockIdx.x * blockDim.x;
  if (idx < n) return;
 
  // process data
}

网格步进循环确保即使线程数少于 n，也能处理所有数据

__global__ void foo(int n, int* data) {
  for (auto const idx = threadIdx.x + blockIdx.x * blockDim.x;
idx < n;
idx += blockDim.x * gridDim.x) {
    // process data
  }
}

使用 Ranger，代码更加清晰且不易出错

#include <ranger/ranger.hpp>
 
__global__ void foo(int n, int* data) {
  for (auto const idx = ranger::grid_stride_range(n)) {
    // process data
  }
}

包含

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

按库对包含进行分组（例如 cuSpatial、RMM、Thrust、STL）。clang-format 将尊重分组并在组内按字典顺序排序每个包含项。
组之间用空行分隔。
按“最近”到“最远”的顺序排列组。换句话说，本地包含，然后是其他 RAPIDS 库的包含，然后是相关库的包含，如 <thrust/...>，然后是随 cuSpatial 安装的依赖项的包含，最后是标准库头文件（例如 <string>, <iostream>）。
除非头文件与源文件位于同一目录中，否则使用 <> 代替 ""。
像 clangd 这样的工具通常会在可能的情况下自动插入包含项，但它们通常会搞错分组和括号。
始终检查包含项是否只对其所在的包含文件是必需的。尽量避免过度包含，尤其是在头文件中。删除代码时请仔细检查这一点。
使用引号 " 来包含来自同一相对源目录的本地头文件。这只应发生在源文件和非公共头文件中。否则，在包含的头文件名前后使用尖括号 <>。
尽可能避免使用 .. 的相对路径。当从与包含文件不在同一目录的源路径中包含（内部）头文件时，带有 .. 的路径是必需的，因为源路径不随 -I 传递。
避免从非内部文件包含库内部头文件。例如，尽量不要在测试或 libcuspatial 公共头文件中包含 libcuspatial src 目录中的头文件。如果您发现自己在这样做，请讨论将包含的内部头文件（部分）移至公共头文件。

libcuspatial 数据结构

仅头文件 libcuspatial API 不关心应用程序用于存储其数据的容器类型，因为仅头文件 API 基于迭代器（参见迭代器要求）。另一方面，基于 cuDF 的 cuSpatial API 使用 cuDF Column 和 Table 来存储和访问应用程序数据。

有关 cuDF 数据结构（包括视图）的更多信息，请参阅 libcudf 开发者指南。

视图和所有权

资源所有权是 libcudf 中的一个重要概念，因此也是基于 cuDF 的 libcuspatial API 中的重要概念。简而言之，“拥有的”对象拥有资源（例如设备内存）。它在构造期间获取资源，并在销毁时释放资源 (RAII)。“非拥有的”对象不拥有资源。libcudf 中任何带有 *_view 后缀的类都是非拥有的。更多详细信息请参见 libcudf++ 演示文稿。

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

std::unique_ptr<cudf::table> points_in_spatial_window(
  ...,
cudf::column_view const& x,
cudf::column_view const& y);

内存资源 (`rmm::device_memory_resource`)

libcuspatial 通过 RMM 内存资源 (MR) 或 CUDA MR 分配所有设备内存。这两种类型都可以通过 rmm::device_async_resource_ref 参数传递给 libcuspatial 函数。有关详细信息，请参阅 RMM 文档。

当前设备内存资源

RMM 为每个设备提供一个“默认”内存资源，可以通过 rmm::mr::get_current_device_resource() 和 rmm::mr::set_current_device_resource(...) 函数分别访问和更新。所有内存资源参数应默认使用 rmm::mr::get_current_device_resource() 的返回值。

资源引用 (Resource Refs)

内存资源通过 resource ref 参数传递。Resource ref 是内存资源包装器，使消费者能够指定他们期望的资源属性。这些定义在 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_device_async_resource_ref 接受提供流有序分配主机和设备可访问内存的内存资源。

有关详细信息，请参阅 libcu++ 关于 resource_ref 的文档。

libcuspatial API 和实现

本节详细介绍 cuSpatial API 函数的结构和实现。

基于列的 cuSpatial API

libcuspatial 的基于列的 API 旨在与其他 RAPIDS 库（特别是 cuDF）无缝集成。为此，此 API 使用 cudf::column 和 cudf::table 数据结构作为输入和输出。这使得 cuSpatial 能够提供 Python 和其他语言 API（例如 Java），与 cuDF 和 cuML 等其他 RAPIDS 库的 API 无缝集成。这允许用户将空间数据查询和转换集成到端到端 GPU 加速的数据分析和机器学习工作流中。

输入/输出风格

基于列的 API 函数传递输入和返回输出的首选风格如下：

输入参数
- 列
  - column_view const&
- 表
  - table_view const&
- 标量
  - scalar const&
- 其他所有类型
  - 平凡类型或复制开销低的类型
    - 按值传递
  - 非平凡类型或复制开销高的类型
    - 按 const& 传递
输入/输出参数
- 列
  - mutable_column_view&
- 表
  - mutable_table_view&
- 其他所有类型
  - 通过裸指针传递
输出
- 输出应被返回，即没有输出参数
- 列
  - std::unique_ptr<column>
- 表
  - std::unique_ptr<table>
- 标量
  - std::unique_ptr<scalar>

以下是一个基于列的 API 函数示例。

std::unique_ptr<cudf::column> haversine_distance(
cudf::column_view const& a_lon,
cudf::column_view const& a_lat,
cudf::column_view const& b_lon,
cudf::column_view const& b_lat,
  double const radius = EARTH_RADIUS_KM,
rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());

要点

所有输入数据都是 cudf::column_view。这是一个类型擦除的容器，因此必须在运行时确定数据类型。
所有输入都是标量数组。经度和纬度是分开的。
输出是返回的 unique_ptr<cudf::column>。
输出在函数内部使用传递的内存资源分配。
公共 API 不接受流参数。有一个接受流参数的 detail 版本 API。这遵循了 libcudf 的做法，未来可能会改变。

多个返回值

有时函数需要有多个输出。C++ 中有几种方法可以做到这一点（包括为输出创建一个 struct）。一种方便的方法是使用 std::tie 和 std::pair。请注意，传递给 std::pair 的对象会调用对象的复制构造函数或移动构造函数，最好是移动非平凡可复制的对象（对于删除了复制构造函数的类型，如 std::unique_ptr，则是必需的）。

功能相关的多个列输出（例如 x 和 y 坐标）应合并到一个 table 中。

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

注意：如果不是因为 Cython 不支持 std::tuple，则可以使用 std::tuple。因此，libcuspatial 公共基于列的 API 必须使用 std::pair，并且仅限于返回两种不同类型的对象。同种类型的多个对象可以通过 std::vector<T> 返回。

或者，可以使用 C++17 结构化绑定来分解多个返回值

auto [out0, out1] = return_two_outputs();

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

auto [out0, out1] = 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.
  // ...
};

仅头文件 cuSpatial API

对于不使用 libcudf 或其他 RAPIDS API 的 C++ 用户和开发者来说，依赖 libcudf 可能会成为采用 libcuspatial 的障碍。libcudf 是一个非常大的库，构建它需要很长时间。

因此，libcuspatial 提供了一个不依赖于 libcudf 的仅头文件 C++ API。仅头文件 API 是一个基于迭代器的接口。这有许多优点。

通过仅头文件 API，用户可以只包含和构建他们使用的部分。
模板 API 可以灵活支持各种基本数据类型，例如用于位置数据的 float 和 double，以及用于索引的不同整数大小。
与 STL 一样，迭代器使得通用算法可以应用于任意容器。
迭代器通过使用“花式 (fancy)”迭代器，使得 cuSpatial 算法可以与输入数据的转换融合。示例包括转换迭代器和计数迭代器。
迭代器使得仅头文件 cuSpatial API 能够使用结构化坐标数据（例如 x,y 坐标对），同时保持与基于列的 API 所需的单独 x 和 y 坐标数组的兼容性。这通过 zip 迭代器实现。在内部，结构化数据（带有算术运算符）使得代码更清晰，更具算术性。
内存资源只需作为分配临时中间存储的 API 的一部分。输出存储在 API 外部分配，并将输出迭代器作为参数传递。

这类 API 的主要缺点是：

仅头文件 API 会增加依赖于它们的代码的编译时间。
一些用户（特别是 cuSpatial Python API）可能更喜欢基于 cuDF 的 API。

基于列的 C++ API 是位于仅头文件 API 之上的简单层。这种方法保护基于列的 API 用户免受缺点的影响，同时保持了仅头文件 API 用户的优点。

输入/输出风格

所有数组输入和输出都是迭代器类型模板，以实现 API 的通用应用。一个示例函数很有帮助。

template <class LonLatItA,
          class LonLatItB,
          class OutputIt,
          class T = typename cuspatial::iterator_vec_base_type<LonLatItA>>
OutputIt haversine_distance(LonLatItA a_lonlat_first,
LonLatItA a_lonlat_last,
LonLatItB b_lonlat_first,
OutputIt distance_first,
T const radius = EARTH_RADIUS_KM,
rmm::cuda_stream_view stream = rmm::cuda_stream_default);

需要注意几个要点。

此 API 与 STL 算法（例如 std::transform）非常相似。
所有数组输入和输出都是迭代器类型模板。
经度/纬度数据作为结构数组传递，使用 cuspatial::vec_2d<T> 类型 (include/cuspatial/vec_2d.hpp)。这通过函数体中的 static_assert 强制执行。（// 通用的二维向量类型。）
浮点类型是一个模板（T），默认等于 LonLatItA 迭代的类型的基本 value_type。libcuspatial 为此提供了 iterator_vec_base_type trait 助手。
两个输入范围（A 和 B）的迭代器类型是不同的模板。这对于组合 A 和 B 可能类型不同的花式迭代器至关重要。
示例 API 中输入和输出范围的大小相等，因此只提供了 A 范围的开始和结束（a_lonlat_first 和 a_lonlat_last）。这类似于 STL API。
此 API 返回一个指向写入输出的最后一个元素之后的元素的迭代器。这受到了 std::transform 的启发，尽管与 transform 一样，许多 cuSpatial API 的使用不需要使用此返回的迭代器。
所有运行 CUDA 设备代码（包括 Thrust 算法）或分配内存的 API 都接受一个 CUDA 流，用于执行设备代码和分配内存。
任何分配并返回设备数据（此处未显示）的 API 也应接受 rmm::device_async_resource_ref，用于输出内存分配。

（多个）返回值

在仅头文件 API 中，只要可能，输出数据应写入引用由 API 调用者分配的数据的输出迭代器。在这种情况下，多个“返回值”简单地写入多个输出迭代器。通常，此类 API 返回一个位于主要输出迭代器末尾之后一个位置的迭代器（采用 std::transform() 的风格）。

在输出大小取决于数据的函数中，API 可以分配输出数据并将其作为 rmm::device_uvector 或包含 device_uvector 的其他数据结构返回。

迭代器要求

所有输入和输出迭代器都必须是设备可访问且具有随机访问能力的。它们必须满足 C++ LegacyRandomAccessIterator (旧版随机访问迭代器) 的要求。输出迭代器必须是可变的 (mutable)。

流 (Streams)

CUDA 流尚未在公共基于列的 libcuspatial API 中暴露。执行 GPU 工作或分配 GPU 内存的仅头文件 libcuspatial API 应接受一个流参数。

为了方便将来在公共基于列的 API 中使用流的过渡，所有分配设备内存或执行 GPU 工作（包括内核、Thrust 算法或任何可以接受流的）的 libcuspatial API 都应使用默认流（例如，流 0）上的异步 API 实现。

推荐的模式是让外部 API 的定义调用 detail 命名空间中的内部 API。内部 detail API 具有与公共 API 相同的参数，外加一个没有默认值的 rmm::cuda_stream_view 参数放在末尾。如果 detail API 也接受内存资源参数，流参数理想情况下应放在内存资源参数之前。公共 API 将调用 detail API 并提供 rmm::cuda_stream_default。实现应完全包含在 detail API 定义中，并且只使用带有流参数的异步版本的 CUDA API。

为了使 detail API 可以从其他 libcuspatial 函数调用，它可以暴露在放置在 cuspatial/cpp/include/detail/ 目录中的头文件中。

例如

// cpp/include/cuspatial/header.hpp
void external_function(...);
 
// cpp/include/cuspatial/detail/header.hpp
namespace detail{
void external_function(..., rmm::cuda_stream_view stream)
} // 命名空间 detail
 
// cuspatial/src/implementation.cpp
namespace detail{
    // 在 detail 实现中使用流参数。
    void external_function(..., rmm::cuda_stream_view stream){
        // 实现使用带有异步 API 的流。
rmm::device_buffer buff(...,stream);
        CUSPATIAL_CUDA_TRY(cudaMemcpyAsync(...,stream.value()));
kernel<<<..., stream>>>(...);
thrust::algorithm(rmm::exec_policy(stream), ...);
    }
} // 命名空间 detail
 
void external_function(...){
CUSPATIAL_FUNC_RANGE(); // 为此函数的生命周期生成一个 NVTX range。
detail::external_function(..., rmm::cuda_stream_default);
}

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

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

NVTX 范围

为了帮助性能优化和调试，所有计算密集型 libcuspatial 函数都应有相应的 NVTX 范围。在 libcuspatial 中，我们有一个便利宏 CUSPATIAL_FUNC_RANGE()，它会自动标注包含函数的生命周期，并使用函数名作为 NVTX 范围的名称。有关 NVTX 的更多信息，请参见此处。

流创建

（注意：cuSpatial 尚未需要内部流创建。）以下指导来自 libcudf 的文档。在实现 libcuspatial 功能时，有时使用流内部实现算法重叠可能是有利的。然而，动态创建流可能会很昂贵。RMM 有一个流池类，可以帮助避免动态流创建。但是，这尚未在 libcuspatial 中暴露，因此目前 libcuspatial 功能应避免创建流（即使效率略低）。最好留下一个 // TODO: 注释，表明在哪里使用流会更有益。

内存分配

libcuspatial 中使用设备内存资源来抽象和控制设备内存的分配方式。

输出内存

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

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

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

临时内存

并非 libcuspatial API 中分配的所有内存都返回给调用者。算法通常必须分配临时、刮擦内存用于中间结果。对于临时内存分配，始终使用从 rmm::mr::get_current_device_resource() 获取的默认资源。示例

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

内存管理

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

rmm::device_buffer

使用内存资源分配指定数量的无类型、未初始化设备内存的字节。如果未明确提供 rmm::device_async_resource_ref，则使用 rmm::mr::get_current_device_resource()。

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>

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

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

rmm::device_vector<T>

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

注意：（待办事项：这在 libcuspatial 中尚不完全正确，但我们应该为此努力。以下内容复制自 libcudf 的开发者指南。）我们已从 libcuspatial 中删除了所有 rmm::device_vector 和 thrust::device_vector 的使用，在新代码中应谨慎使用。相反，请使用 rmm::device_uvector 以及 device_factories.hpp 中的实用工厂函数。这些实用函数能够从主机端向量创建 uvector，或创建零初始化的 uvector，使其使用起来像 device_vector 一样方便。避免使用 device_vector 有许多好处，如下一节关于 rmm::device_uvector 的描述。

rmm::device_uvector<T>

类似于 device_vector，在设备内存中分配连续的元素集，但有一些关键区别：

作为优化，元素是未初始化的，并且构造时不会发生同步。这将类型 T 限制为平凡可复制类型。
所有操作都是流有序的（即，它们接受一个 cuda_stream_view 来指定执行操作的流）。这提高了使用非默认流时的安全性。
与 thrust/device_vector.hpp 不同，device_uvector.hpp 不包含任何 __device__ 代码，这意味着 device_uvector 可以在 .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});
 
rmm::mr::device_memory_resource * mr = new my_custom_resource{...};
// 在流 `s` 上使用资源 `mr` 分配 100 个 `int32_t` 元素的未初始化存储空间
rmm::device_uvector<int32_t> v2{100, s, mr};

命名空间

外部

所有公共 libcuspatial API 都应放在 cuspatial 命名空间中。示例

namespace cuspatial{
   void public_function(...);
} // 命名空间 cuspatial

顶层 cuspatial 命名空间对于大多数公共 API 来说已经足够。但是，为了逻辑上对大量函数进行分组，可以使用进一步的命名空间。

内部

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

detail 命名空间

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

// some_utilities.hpp
namespace cuspatial{
namespace detail{
namespace detail{
} // 命名空间 detail
} // 命名空间 cuspatial

void reusable_helper_function(...);

匿名命名空间

只在单个翻译单元中使用的函数或对象应在使用它的源文件中定义在匿名命名空间中。示例
// some_file.cpp{
namespace{
void isolated_helper_function(...);

} // 匿名命名空间

匿名命名空间绝不应用于头文件中。

弃用和移除代码

libcuspatial 不断发展，以提高性能并更好地满足用户的需求。因此，我们有时需要破坏或完全移除 API，以响应对我们提供功能的新的和改进的理解。保持自由地这样做对于使 libcuspatial 成为一个能够快速适应用户需求的灵活库至关重要。因此，我们并非总是在发布破坏性更改之前提供警告或任何提前通知。在尽最大努力的基础上，libcuspatial 团队将通知用户我们预计会产生重大或广泛影响的更改。

在可能的情况下，在移除之前使用 deprecated (已弃用) 属性指示即将移除的 API，并使用 Doxygen 的 deprecated 命令进行文档记录。如果弃用的 API 有替代 API，请在弃用消息和弃用文档中提及替代方案。引入弃用的 Pull Request 应标记为“deprecation”，以便在后续版本中更容易发现和移除。

通过将任何破坏或移除现有 API 的 Pull Request 标记为“breaking”标签来宣传破坏性更改。这确保了发布说明的“Breaking”部分包含对与过去版本相比有哪些更改的描述。将包含弃用的 Pull Request 标记为“non-breaking”标签。

错误处理

libcuspatial 遵循强制执行编译时和运行时条件以及检测和处理 CUDA 错误的约定（并提供实用程序）。错误总是通过 C++ 异常进行通信。

运行时条件

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

示例用法

CUSPATIAL_EXPECTS(lhs.type() == rhs.type(), "Column type mismatch");

CUSPATIAL_EXPECTS

用于检查（前置）条件的宏，当条件不满足时抛出异常。

定义 error.hpp:76

第一个参数是预期在正常条件下解析为 true 的条件表达式。如果条件评估为 false，则表示发生了错误，并抛出一个 cuspatial::logic_error 实例。CUSPATIAL_EXPECTS 的第二个参数是发生的错误的简短描述，用于异常的 what() 消息。

有时，特定的代码路径如果被执行到，无论如何都应该表示错误。例如，switch 语句的 default 情况通常表示无效的替代方案。对于此类错误，请使用 CUSPATIAL_FAIL 宏。这实际上等同于调用 CUSPATIAL_EXPECTS(false, reason)。

示例

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

CUSPATIAL_FAIL

指示已执行到错误的代码路径。

定义 error.hpp:119

CUDA 错误检查

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

CUSPATIAL_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.");
   数据类型
...
}

列可以包含多种类型的数据。cuDF 支持 cuSpatial 中不使用的多种类型。cuSpatial 函数主要操作数值和时间戳数据。有关 libcudf 数据类型的更多信息，请参阅 libcudf 开发者指南。

类型分派器 (Type Dispatcher)

`cudf::column` 在 `void` 设备内存中存储“类型擦除”的数据（用于列和标量）。这种类型擦除使得能够与其他语言和类型系统（如 Python 和 Java）互操作。为了确定类型，函数必须使用存储在列 `type()` 中的运行时信息，通过将 `void` 转换为相应的 `T*` 来重构数据类型 `T`。

这种所谓的类型分派在 libcudf 和基于列的 libcuspatial API 中普遍存在。cudf::type_dispatcher 是一个核心实用程序，它自动化了将 data_type 中的运行时类型信息映射到具体 C++ 类型的过程。有关详细信息，请参阅 libcudf 开发者指南。

由 1.13.2 生成

概述

术语表

目录结构和文件命名

文件扩展名

代码和文档风格及格式

C++ 指南

循环和网格步进循环

包含

libcuspatial 数据结构

视图和所有权

内存资源 (rmm::device_memory_resource)

当前设备内存资源

资源引用 (Resource Refs)

libcuspatial API 和实现

基于列的 cuSpatial API

输入/输出风格

多个返回值

仅头文件 cuSpatial API

输入/输出风格

（多个）返回值

迭代器要求

流 (Streams)

NVTX 范围

流创建

内存分配

输出内存

临时内存

内存管理

rmm::device_buffer

rmm::device_scalar<T>

rmm::device_vector<T>

rmm::device_uvector<T>

命名空间

外部

内部

detail 命名空间

void reusable_helper_function(...);

匿名命名空间绝不应用于头文件中。

通过将任何破坏或移除现有 API 的 Pull Request 标记为“breaking”标签来宣传破坏性更改。这确保了发布说明的“Breaking”部分包含对与过去版本相比有哪些更改的描述。将包含弃用的 Pull Request 标记为“non-breaking”标签。

libcuspatial 遵循强制执行编译时和运行时条件以及检测和处理 CUDA 错误的约定（并提供实用程序）。错误总是通过 C++ 异常进行通信。

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

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

列可以包含多种类型的数据。cuDF 支持 cuSpatial 中不使用的多种类型。cuSpatial 函数主要操作数值和时间戳数据。有关 libcudf 数据类型的更多信息，请参阅 libcudf 开发者指南。

内存资源 (`rmm::device_memory_resource`)

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

使用 `static_assert` 来强制执行编译时条件。例如，