1. NVIDIA GPUDirect Storage 最佳实践指南#
本最佳实践指南由熟悉 NVIDIA® GPUDirect® Storage (GDS) 的专家提供指导。
2. 简介#
本最佳实践指南的目的是提供熟悉 NVIDIA® GPUDirect® Storage (GDS) 的专家的指导。本指南还提供了关于构建和扩展大规模 GPU 加速 I/O 存储基础设施的经验教训。目标受众包括数据中心规划人员、系统构建商、开发人员和存储供应商。
3. 软件设置#
本节介绍 GDS 所需的设置。
为了获得最佳性能，整个系统需要进行多项软件设置，并且某些设置特定于您正在使用的文件系统。
有关更多信息，请参阅GPUDirect Storage 安装和故障排除指南。
3.1. 系统设置#
对于基于 Grace CPU 的 DGX™ (Grace Hopper) 平台上的 GDS p2p 支持，应启用 IOMMU，并禁用直通设置。
以下是我们建议在基于裸机 x86_64 的平台上获得最佳性能的系统设置。
PCIe 访问控制服务 (ACS)。
ACS 强制 P2P PCIe 事务通过 PCIe 根联合体，这不允许 GDS 在包含 PCIe 交换机的系统中绕过网络适配器或 NVMe 与 GPU 之间的 CPU 路径。
为了获得最佳 GDS 性能，请禁用 ACS。
注意
要列出所有已启用 ACS 的 PCI 交换机，请执行
/usr/local/cuda/gds/tools/gdscheck -p。IOMMU
当启用 IOMMU 设置时，PCIe 流量将通过 CPU 根端口路由。对于 GPU 和 NIC 位于同一 PCIe 交换机下的配置，此路由限制了最大可实现吞吐量。在安装 GDS **之前**，您**必须**禁用 IOMMU。有关更多信息，请参阅安装 GPUDirect Storage。
注意
要确定是否启用了 IOMMU 设置，请检查 cat /proc/cmdline 的输出或使用
gdscheck命令。例如，以下输出显示在此系统上启用了 IOMMU
$ cat /proc/cmdline BOOT_IMAGE=/boot/vmlinuz-5.19.0-38-generic root=UUID=fb2a25a8-9d2e-4e1c-9d8a-efabdf165adc ro rootflags=data=ordered amd_iommu=on
同样，如果系统上禁用了 IOMMU，则使用
gdscheck您应该看到以下输出$ /usr/local/cuda/gds/tools/gdscheck -p IOMMU: disabled Platform verification succeeded
NIC 亲和性
为了使 P2P DMA 有效运行，NIC、NVMe 和 GPU 应尽可能位于 PCIe 交换机下。为了使 P2P DMA 在基于 NVIDIA DGX™ 的平台上有效运行，请确保至少一个 NIC 与 GPU 位于同一 CPU 插槽中。
避免 NIC 分配在需要 PCIe 流量跨越 CPU 根端口或跨越使用 QPI 的 CPU 插槽的配置。
NIC 版本
当使用 Mellanox ConnectX-5 或更高版本时，HCA 必须配置为 InfiniBand 或 RoCE v2 模式。
对于 GDS 支持，需要 MLNX_OFED 5.4 或更高版本，或 DOCA 2.9.0 或更高版本。
3.2. 在 GPU 内核和存储 IO 中使用 CUDA 上下文#
在某些情况下，GDS 工作负载数据可以通过称为反弹缓冲区的中间缓冲区发布。因此，涉及到从这些 GPU 反弹缓冲区到应用程序的 GPU 缓冲区的 D2D 复制。cuFile 库在主 CUDA 上下文上创建的流上发布这些 IO。如果一个繁重的计算作业或应用程序内核以 GPU 内核的形式在单独的上下文（非主上下文）中在后台运行，则它可能会干扰 D2D 复制并增加 D2D 复制启动时间。如果计算内核在主上下文中运行，则不会发生此问题，因此建议应用程序在主上下文而不是使用单独的上下文启动 GPU 内核。
注意
如果应用程序使用 CUDA 运行时 API，则内核启动将默认在主上下文中发生。
3.3. cuFile 配置设置#
GDS 中的 cuFile 配置设置存储在 /etc/cufile.json 文件中。
您可以编辑该文件以获得应用程序的最佳性能，如下所示。有关文件中参数的信息，请参阅https://docs.nvda.net.cn/gpudirect-storage/configuration-guide/index.html#gds-parameters。
要显示配置设置，请运行以下命令
$ cat /etc/cufile.json
示例输出的一部分
"properties": {
// max IO size issued by cuFile to nvidia-fs driver (in KB)
"max_direct_io_size_kb" : 16384,
...
}
对于请求的 IO 大小，GDS 基于 max_direct_io_size 参数，以读取/写入块的顺序发出 IO 请求。较大的 max_direct_io_size 值将减少对 IO 堆栈的调用次数，并可能导致更高的吞吐量。
max_direct_io_size_kb 参数可以设置为 64K 的倍数的值。此过程定义了在 cuFileBufRegister 期间用于每个缓冲区的额外系统内存，最大值为 properties:max_direct_io_size_kb 参数的 16MB。此值可以减小到 1MB，以减少每个缓冲区使用的系统内存量。
可以从 nvidia-fs 统计信息中获得使用的总系统内存。
在此示例中，256 个线程中的每个线程都为 GDS 注册了 1MB 缓冲区。
运行以下命令
$ cat /proc/driver/nvidia-fs/stats
查看输出
NVFS statistics(ver:1.0) Active Shadow-Buffer (MB): 256...
cufile.json中有许多可调参数。请参阅GPUDirect Storage 参数。
4. API 用法#
本节介绍使用 GDS API 时要记住的最佳实践。
cuFile API 被设计为线程安全的。
在初始化 cuFile 库后，不应使用 fork 系统调用。在子进程中，fork 系统调用后的 API 行为未定义。
带有 GPU 缓冲区的 API 应在有效的 CUDA 上下文中调用。
下表概述了各种特定于 IO 的用例及其对应的最适合的 cuFile API 的建议。
模式 |
IO 行为 |
用例 |
优点/缺点 |
|---|---|---|---|
|
同步提交 同步完成 |
单线程应用程序，使用标准文件系统调用处理单个大文件和大型缓冲区（>16MB） |
优点
缺点
|
启用 cuFile 线程池
|
同步提交 同步完成 |
单线程应用程序，使用标准文件系统调用处理单个大文件和大型缓冲区 多线程应用程序，使用标准文件系统调用处理多个文件和缓冲区。 应用程序具有用于其 IO 管道的线程池。 |
优点
缺点
|
|
同步提交 异步完成 |
单线程应用程序，使用标准文件系统调用，为多个不连续的文件偏移量、大小和 GPU 缓冲区执行 IO。 每个 IO 请求都很小（< 64KB） 可以异步跟踪 IO 的完成情况，或在同一线程中等待。 |
优点
缺点
|
|
异步提交 异步完成 |
单线程应用程序，使用标准文件系统调用，为多个不连续的文件偏移量、大小和 GPU 缓冲区执行 IO。 IO 大小 - 缓冲区数据取决于先前的 CUDA 工作。 |
优点
缺点
|
4.1. cuFileDriverOpen#
cuFileDriverOpen API 每个进程应仅调用一次，并且必须在调用任何其他 cuFile API **之前**发生。应用程序应调用此例程以避免驱动程序初始化的延迟，否则该延迟将在第一次 IO 调用中产生。
4.2. cuFileHandleRegister#
cuFileHandleRegister API 将文件描述符转换为 cuFileHandle，并检查命名文件在其挂载点上是否可以通过此平台上的 GDS 支持。此例程是调用所有采用 cuFileHandle 参数的 cuFile API 调用所必需的。
注意
每个文件描述符应仅创建一个句柄。
同一个句柄可以由多个线程共享。有关多个线程使用同一句柄的更多信息，请参阅示例程序。
注意
在兼容模式下，可以在文件上打开一个额外的文件描述符，而无需 O_DIRECT 模式。即使在 POSIX 无法处理的情况下，此模式也可以处理未对齐的读取/写入。
4.3. cuFileBufRegister、cuFileRead、cuFileWrite、cuFileBatchIOSubmit、cuFileBatchIOGetStatus、cuFileReadAsync、cuFileWriteAsync 和 cuFileStreamRegister#
GPU 缓冲区需要暴露给第三方设备，以启用这些设备的 DMA。GPU 虚拟地址空间中跨越这些缓冲区的页面集需要映射到基地址寄存器 (BAR) 空间，并且此映射会产生开销。
完成此映射的机制称为注册。使用 cuFileBufRegister API 进行显式 GPU 缓冲区注册是可选的。如果未注册用户缓冲区，则将使用 cuFile 实现拥有的中间预注册 GPU 缓冲区，并且从那里到用户缓冲区会进行额外的复制。下表和 IO 模式描述提供了关于注册是否有利的指导。
注意
IO 模式 1 是一个次优的基线案例，并且未在本表中引用。
用例 |
描述 |
建议 |
|---|---|---|
一个 4KB 对齐的 GPU 缓冲区被重用为中间缓冲区，以使用存储系统的最佳 IO 大小（4KB 的倍数）来读取或写入数据。 |
GPU 缓冲区用作中间缓冲区，以流式传输内容或填充 GPU 内存中的不同数据结构。 您可以为具有 DSG 的 IO 库实现此用例。 |
注册此可重用的中间缓冲区，以避免通过在 cuFile 库中使用 GPU 反弹缓冲区来进行额外的数据内部暂存。 有关建议的用法，请参阅IO 模式 2。 |
为一个用途填充大型 GPU 缓冲区。 |
GPU 缓冲区是数据的最终位置。由于缓冲区不会被重用，因此注册成本将不会被摊销。一个使用示例是读取大型预格式化的检查点二进制数据。 注册大型缓冲区可能会在注册缓冲区时产生延迟影响。 |
这也可能导致 BAR 内存耗尽，因为运行多个线程或应用程序将竞争 BAR 内存。 在不进行缓冲区注册的情况下读取或写入数据。 有关建议的用法，请参阅IO 模式 3。 |
分区 GPU 缓冲区以供多个线程访问。 |
主线程分配一个大型内存缓冲区并创建多个线程。每个线程独立注册内存缓冲区的一部分，并像在IO 模式 2中那样使用它。 您还可以在父线程中注册整个缓冲区，并将此注册缓冲区与大小和 |
在每个线程中独立分配、注册和注销缓冲区，以实现简单的 IO 工作流程。 对于 GPU 内存预先分配的情况，每个线程都可以设置适当的上下文并独立注册缓冲区。 有关建议的用法，请参阅 IO 模式 6。 安装 GDS 软件包后，请参阅 |
GPU 偏移量、文件偏移量和 IO 请求大小未对齐。 |
IO 读取或写入大多未对齐。可能需要中间对齐的缓冲区来处理 GPU 偏移量、文件偏移量和 IO 大小方面的对齐问题。 |
**不要**注册缓冲区。 |
在 BAR 空间与可用 GPU 内存相比很小的 GPU 上工作。 |
在某些 GPU SKU 中，BAR 内存小于设备总内存。 |
为了避免因 BAR 内存耗尽而导致的故障，请勿注册缓冲区。 请参阅IO 模式 3。 |
4.3.1. IO 模式 1#
以下是 IO 模式 1 的代码示例。
1 #define MB(x) ((x)*1024*1024L)
2 #define GB(x) ((x)*1024*1024L*1024L)
3
4
5 void thread_func(CUfileHandle_t cuHandle)
6 {
7 void *devPtr_base;
8 int readSize = MB(100);
9 int devPtr_offset = 0;
10 int file_offset = 0;
11 int ret = 0;
12
13 cudaSetDevice(0);
14 cudaMalloc(&devPtr_base, GB(1));
15
16 for (int i = 0; i < 10; i++) {
17
18 cuFileBufRegister((char *)devPtr_base + devPtr_offset, readSize, 0);
19
20 ret = cuFileRead(cuHandle, (char *)devPtr_base + devPtr_offset,
readSize, file_offset, 0);
21
22
<... launch cuda kernel using contents at devPtr_base + devPtr_offset … >
23 file_offset += readSize;
24 devPtr_offset += readSize;
25
26 cuFileBufDeregister((char *)devPtr_base + devPtr_offset);
27 }
28 }
使用
cudaMalloc分配 1 GB 的 GPU 内存。通过一次从文件中读取 100 MB 来填充 1 GB，如下面的循环所示
在第 18 行，注册了 100 MB 的 GPU 缓冲区。
提交 100MB 的读取（readsize 为 100 MB）。
在第 26 行，注销了 100 MB 的 GPU 缓冲区。
虽然在语义上是正确的，但此循环可能无法提供最佳性能，因为 cuFileBufRegister 和 cuFileBufDeregister 在循环中不断发出。例如，可以按照IO 模式 2所示的方式解决此问题。
4.3.2. IO 模式 2#
以下是 IO 模式 2 的代码示例。
1 #define MB(x) ((x)*1024*1024L)
2 #define GB(x) ((x)*1024*1024L*1024L)
3
4
5 void thread_func(CUfileHandle_t cuHandle)
6 {
7 void *devPtr_base;
8 int readSize = MB(100);
9 int devPtr_offset = 0;
10 int file_offset = 0;
11 int ret = 0;
12
13 cudaSetDevice(0);
14 cudaMalloc(&devPtr_base, GB(1));
15 cuFileBufRegister(devPtr_base, GB(1), 0);
16
17 for (int i = 0; i < 10; i++) {
18
19 ret = cuFileRead(cuHandle, devPtr_base,
readSize, file_offset, devPtr_offset);
20
21 <... launch cuda kernel using contents at devPtr_base + devPtr_offset … >
22
23 file_offset += readSize;
24 devPtr_offset += readSize;
25
26 }
27 cuFileBufDeregister(devPtr_base);
28 }
4.3.3. IO 模式 3#
以下是 IO 模式 3 的代码示例。
1 #define MB(x) ((x)*1024*1024L)
2 #define GB(x) ((x)*1024*1024L*1024L)
3
4
5 void thread_func(CUfileHandle_t cuHandle)
6 {
7 void *devPtr_base;
8 int readSize = MB(100);
9 int devPtr_offset = 0;
10 int file_offset = 0;
11 int ret = 0;
12
13 cudaSetDevice(0);
14 cudaMalloc(&devPtr_base, GB(1));
15
16 for (int i = 0; i < 10; i++) {
17
18 ret = cuFileRead(cuHandle, (char *)devPtr_base,
readSize, file_offset, devPtr_offset);
19
20 <... launch cuda kernel using contents at devPtr_base + devPtr_offset … >
21
22 file_offset += readSize;
23 devPtr_offset += readSize;
24 }
25 }
此示例演示了在不使用 cuFileBufRegister 和 cuFileBufDeRegister API 的情况下使用 cuFileRead/cuFileWrite API。IO 模式 3 代码片段与 IO 模式 1 和 IO 模式 2 代码片段相同，但未使用 cuFileBufRegister API。
分配 1 GB 的 GPU 内存。
通过一次从文件中读取 100 MB 来填充 1 GB 的整个 GPU 内存，如下面的循环所示。
注意
虽然在语义上是正确的，但此循环可能不是最佳的。
在内部，GDS 使用 GPU 反弹缓冲区来执行 IO。反弹缓冲区是 GDS 内部的 GPU 内存分配，这些缓冲区由 GDS 库注册和管理。反弹缓冲区的数量根据 max_device_cache_size （表示反弹缓冲区缓存的总大小）和 per_buffer_cache_size （表示每个缓冲区的大小）设置在 /etc/cufile.json 文件中进行限制。max_device_cache_size 和 per_buffer_cache_size 的默认值分别为 128MB 和 1MB，默认情况下总共为 128 个反弹缓冲区。
4.3.4. IO 模式 4#
以下是 IO 模式 4 的代码示例。这是一个由于文件偏移量未对齐而导致的未对齐 IO。
1 #define MB(x) ((x)*1024*1024L)
2 #define GB(x) ((x)*1024*1024L*1024L)
3
4
5 void thread_func(CUfileHandle_t cuHandle)
6 {
7 void *devPtr_base;
8 int readSize = MB(100);
9 int devPtr_offset = 0;
10 int file_offset = 3; // Start from odd offset
11 int ret = 0;
12
13 cudaSetDevice(0);
14 cudaMalloc(&devPtr_base, GB(1));
15 cuFileBufRegister(devPtr_base, GB(1), 0);
16
17 for (int i = 0; i < 10; i++) {
18 // IO issued at offsets which are not 4K aligned
19 ret = cuFileRead(cuHandle, devPtr_base,
readSize, file_offset, devPtr_offset);
20 assert(ret >= 0);
<... launch cuda kernel using contents at devPtr_base + devPtr_offset … >
21
22 file_offset += readSize;
23 devPtr_offset += readSize;
24
25 }
26 cuFileBufDeRegister(devPtr_base);
27 }
此示例演示了当 IO 未对齐时如何使用 cuFileRead 或 cuFileWrite。
如果以下条件之一为真，则 IO 未对齐
在
cuFileRead或cuFileWrite中发出的file_offset未 4K 对齐。在
cuFileRead或cuFileWrite中发出的大小未 4K 对齐。在
cuFileRead或cuFileWrite中发出的devPtr_base未 4K 对齐。在
cuFileRead或cuFileWrite中发出的devPtr_offset未 4K 对齐。
注意
在上面的示例中，file_offset 的初始化在第 10 行。
在分配 1 GB 的 GPU 内存后，立即为 1 GB 的整个范围调用
cuFileBufRegister，如第 15 行所示。通过一次从文件中读取 100 MB 来填充整个 1 GB GPU 内存，如下面的循环所示
初始 file_offset 为 3，并且每次迭代都以 100MB 的
readSize值在偏移量 3 处提交读取。因此，每次读取期间的
file_offset都未 4K 对齐。由于
file_offset未 4K 对齐，GDS 库将在内部使用 GPU 反弹缓冲区来完成 IO。GPU 反弹缓冲区机制与IO 模式 3相同。
未对齐的 IO 可能不是最佳的，应通过读取以 4KB 的倍数指定的大小值和以 4KB 的倍数指定
file_offsets值来避免。在上面的示例中，使用
cuFileBufRegister注册了整个 1GB 的 GPU 内存。但是，由于 IO 未对齐，GDS 库无法直接对这些注册的缓冲区执行 IO。为了处理未对齐的 IO，库将使用 GPU 反弹缓冲区来执行 IO，并将数据从反弹缓冲区复制到应用程序缓冲区。作为最佳实践，如果应用程序通常执行未对齐的 IO，则应用程序缓冲区不需要使用 GDS 库进行注册。IO 模式 4 中的示例演示了当
file_offset未对齐时会发生什么情况；如果任何未对齐条件为真，则先前提到的点是准确的。
如果应用程序无法发出 4K 对齐的 IO，请使用 cuFileRead 或 cuFileWrite API，如 IO 模式 2 中所述，而不是使用 cuFileBufRegister API。
注意
当写入工作负载未对齐时，GDS 在内部使用 POSIX 模式使用读取-修改-写入。
4.3.5. IO 模式 5#
以下是 IO 模式 5 的代码示例。此 IO 是由于缓冲区指针和偏移量未 4K 对齐而导致的未对齐 IO。
1 #define MB(x) ((x)*1024*1024L)
2 #define GB(x) ((x)*1024*1024L*1024L)
3
4
5 void thread_func(CUfileHandle_t cuHandle)
6 {
7 void *devPtr_base;
8 int readSize = MB(100);
9 int devPtr_offset = 3; // Start from odd offset
10 int file_offset = 0;
11 int ret = 0;
12
13 cudaSetDevice(0);
14 cudaMalloc(&devPtr_base, GB(1));
15 cuFileBufRegister(devPtr_base, GB(1), 0);
16
17 for (int i = 0; i < 10; i++) {
18 // IO issued at gpu buffer offsets which are not 4K aligned
19 ret = cuFileRead(cuHandle, devPtr_base,
readSize, file_offset, devPtr_offset);
20 assert (ret >= 0);
<... launch cuda kernel using contents at devPtr_base + devPtr_offset … >
21
22 file_offset += readSize;
23 devPtr_offset += readSize;
24
25 }
26 cuFileBufDeRegister(devPtr_base);
27 }
此示例演示了当 IO 未对齐时如何使用 cuFileRead/cuFileWrite。devPtr_base + devPtr_offset 发送到 cuFileRead 或 cuFileWrite 的未 4K 对齐。
如果 IO 未对齐，则 cuFile 库将通过其内部 GPU 反弹缓冲区缓存发出 IO。但是，如果内部缓存的分配失败，则 IO 将失败。为了避免在这种情况下发生 IO 失败，您可以在 /etc/cufile.json 文件中将 allow_compat_mode 设置为 true。通过此设置，IO 将回退到在 GDS 中使用 POSIX API 调用。
4.3.6. IO 模式 6#
以下程序代码段演示了 cuFile 批处理 API 的使用。
int main(int argc, char *argv[]) {
int fd[MAX_BATCH_IOS];
void *devPtr[MAX_BATCH_IOS];
CUfileDescr_t cf_descr[MAX_BATCH_IOS];
CUfileHandle_t cf_handle[MAX_BATCH_IOS];
CUfileIOParams_t io_batch_params[MAX_BATCH_IOS];
CUfileIOEvents_t io_batch_events[MAX_BATCH_IOS];
<Get program inputs>
status = cuFileDriverOpen();
if (status.err != CU_FILE_SUCCESS) {
std::cerr << "cufile driver open error: "
<< cuFileGetErrorString(status) << std::endl;
return -1;
}
<Open files and call cuFileHandleRegister for each of the batch entry file handles>
<Allocate cuda memory and register buffers using cuFileBufRegister for each of the
batch entries>
for(i = 0; i < batch_size; i++) {
io_batch_params[i].mode = CUFILE_BATCH;
io_batch_params[i].fh = cf_handle[i];
io_batch_params[i].u.batch.devPtr_base = devPtr[i];
io_batch_params[i].u.batch.file_offset = i * size;
io_batch_params[i].u.batch.devPtr_offset = 0;
io_batch_params[i].u.batch.size = size;
io_batch_params[i].opcode = CUFILE_READ;
}
std::cout << "Setting Up Batch" << std::endl;
errorBatch = cuFileBatchIOSetUp(&batch_id, batch_size);
if(errorBatch.err != 0) {
std::cerr << "Error in setting Up Batch" << std::endl;
goto error;
}
errorBatch = cuFileBatchIOSubmit(batch_id, batch_size, io_batch_params, flags);
if(errorBatch.err != 0) {
std::cerr << "Error in IO Batch Submit" << std::endl;
goto error;
}
// Setting min_nr to batch_size for this example.
min_nr = batch_size;
while(num_completed != min_nr) {
memset(io_batch_events, 0, sizeof(*io_batch_events));
nr = batch_size;
errorBatch = cuFileBatchIOGetStatus(batch_id, batch_size, &nr, io_batch_events, NULL);
if(errorBatch.err != 0) {
std::cerr << "Error in IO Batch Get Status" << std::endl;
goto error;
}
std::cout << "Got events " << nr << std::endl;
num_completed += nr;
<Copy to the user buffer>
}
cuFileBatchIODestroy(batch_id);
< Deregister the device memory using cuFileBufDeregister>
status = cuFileDriverClose();
std::cout << "cuFileDriverClose Done" << std::endl;
if (status.err != CU_FILE_SUCCESS) {
...
}
ret = 0;
return ret;
...
}
此程序演示了一个简单的用例，其中可以使用 cuFile 批处理 API 来执行指定批处理大小的读取。它提供了一系列调用的示例，其中每个条目都使用每个单独文件描述符上的注册缓冲区。
值得一提的是，在上面的示例中传递给 cuFileBatchIOGetStatus() 的 min_nr 设置为 batch_size。可以将 min_nr 设置为小于 batch_size 的值，并且随着 min_nr 个 I/O 完成，可以将许多后续 I/O 提交到 I/O 管道，从而提高 I/O 吞吐量。
4.3.7. IO 模式 7#
以下程序代码段使用基于 cuFile 流的异步 I/O API 来执行数据完整性测试。
typedef struct io_args_s
{
void *devPtr;
size_t max_size;
off_t offset;
off_t buf_off;
ssize_t read_bytes_done;
ssize_t write_bytes_done;
} io_args_t;
int main(int argc, char *argv[]) {
unsigned char iDigest[SHA256_DIGEST_LENGTH],
oDigest[SHA256_DIGEST_LENGTH];
<Get inputs>
<Create a data file using some random data>
// Allocate device Memory and register with cuFile
check_cudaruntimecall(cudaMalloc(&args.devPtr, args.max_size));
// Register buffers. For unregistered buffers, this call is not required.
status = cuFileBufRegister(args.devPtr, args.max_size, 0);
if (status.err != CU_FILE_SUCCESS) {
goto error;
}
< Open the data file just created for read and create a new data file to write the content
read from the datafile>
<Register the filehandles>
// Create stream for I/O.
check_cudaruntimecall(cudaStreamCreateWithFlags(&io_stream,
cudaStreamNonBlocking));
// Register Streams for best performance
// If all the inputs i.e. size, offset and buf_off are known and they are page aligned, then
// use CU_FILE_STREAM_FIXED_AND_ALIGNED flag. If they are not known but will
// always be page aligned then use CU_FILE_STREAM_PAGE_ALIGNED_INPUTS flag
// flag.
check_cudaruntimecall(cuFileStreamRegister(io_stream,
CU_FILE_STREAM_FIXED_AND_ALIGNED));
// special case for holes
check_cudaruntimecall(cudaMemsetAsync(args.devPtr, 0, args.max_size, io_stream));
status = cuFileReadAsync(cf_rhandle, (unsigned char *)args.devPtr,
&args.max_size, &args.offset, &args.buf_off,
&args.read_bytes_done, io_stream);
if (status.err != CU_FILE_SUCCESS) {
std::cerr << "read failed : "
<< cuFileGetErrorString(status) << std::endl;
ret = -1;
goto error;
}
// Write loaded data from GPU memory to a new file
status = cuFileWriteAsync(cf_whandle, (unsigned char *)args.devPtr,
(size_t *)&args.max_size, &args.offset, &args.buf_off,
&args.write_bytes_done, io_stream);
if (status.err != CU_FILE_SUCCESS) {
goto error;
}
std::cout << "writing submit done to file :" << TEST_WRITEFILE << std::endl;
check_cudaruntimecall(cudaStreamSynchronize(io_stream));
if((args.read_bytes_done < (ssize_t)args.max_size) ||
(args.write_bytes_done < args.read_bytes_done))
{
std::cerr << "io error issued size:" << args.max_size <<
" read:" << args.read_bytes_done <<
" write:" << args.write_bytes_done << std::endl;
goto error;
}
// Compare file signatures
ret = SHASUM256(TEST_READWRITEFILE, iDigest, args.max_size);
if(ret < 0) {
...
}
DumpSHASUM(iDigest);
ret = SHASUM256(TEST_WRITEFILE, oDigest, args.max_size);
if(ret < 0) {
...
}
DumpSHASUM(oDigest);
if (memcmp(iDigest, oDigest, SHA256_DIGEST_LENGTH) != 0) {
std::cerr << "SHA SUM Mismatch" << std::endl;
ret = -1;
} else {
std::cout << "SHA SUM Match" << std::endl;
ret = 0;
}
if(io_stream) {
check_cudaruntimecall(cuFileStreamDeregister(io_stream));
check_cudaruntimecall(cudaStreamDestroy(io_stream));
}
<Free up all the resources>
return ret;
error:
...
}
此程序演示了一个简单的用例，其中可以使用 cuFile 流 API 来使用单个流执行数据完整性测试。它首先使用随机内容创建一个数据文件。然后，它通过 I/O 流读取内容，并将该内容写入新文件。最后，它使用 SHA（简单哈希算法）比较新创建的数据文件的内容与原始内容。可能在一开始不知道确切的大小，而稍后才会知道。在这种情况下，可以在调用 cuFileReadAsync 或 cuFileWriteAsync API 之前，在同一流上使用 CUDA 主机回调函数 (cuLaunchHostFunc) 设置实际大小。
4.4. cuFileHandleDeregister#
**先决条件**：在调用此 API 之前，应用程序必须确保该句柄上的 IO 已完成且不再使用。文件描述符应仍然打开。
为了在进程结束前回收资源，请始终调用 cuFileHandleDeregister API。
4.5. cuFileBufDeregister#
**先决条件**：在调用此 API 之前，应用程序必须确保使用该缓冲区的全部 cuFile IO 操作已完成。
对于使用 cuFileBufRegister 注册的每个缓冲区，请使用此 API 通过使用用于注册的同一设备指针来注销它。此过程确保在进程结束前回收所有资源。
4.6. cuFileStreamRegister#
cuFileStreamRegister API 将文件描述符转换为 cuFileHandle，并检查命名文件在其挂载点上是否可以通过此平台上的 GDS 支持。
使用 cuFileStreamRegister API 进行显式流注册是可选的。如果注册了流，则将为后续流 I/O 预先分配一些内部缓冲区和关联的元数据资源，并可能提高 I/O 延迟。此外，这些资源将一直重用，直到使用 cuFileStreamUnregister 注销。如果没有此 API，所有这些资源将在实际 I/O 期间分配。
4.7. cuFileStreamDeregister#
**先决条件**：在调用此 API 之前，应用程序必须确保该流上的 I/O 已完成，并且该流不再使用
对于使用 cuFileStreamRegister 注册的每个流，请使用此 API 通过使用用于注册的同一流来注销它。为了在进程结束前回收资源，请始终调用此 API。
4.8. cuFileDriverClose#
**先决条件**：在调用此 API 之前，应用程序必须确保所有 cuFile IO 操作已完成，并且所有缓冲区和句柄都已注销。
为了减少启用 GDS 的应用程序的拆卸时间（即加速释放固定的 GPU 缓冲区和其他 cuFile 资源），强烈建议在应用程序结束时调用 cuFileDriverClose() API。
5. 声明#
本文档仅供参考，不得视为对产品的特定功能、状况或质量的保证。NVIDIA Corporation（“NVIDIA”）不对本文档中包含的信息的准确性或完整性做出任何明示或暗示的陈述或保证，并且对本文档中包含的任何错误不承担任何责任。NVIDIA 对因使用此类信息或因其使用而可能导致的侵犯第三方专利或其他权利的行为的后果或使用不承担任何责任。本文档不承诺开发、发布或交付任何材料（如下定义）、代码或功能。
NVIDIA 保留随时对此文档进行更正、修改、增强、改进和任何其他更改的权利，恕不另行通知。
客户应在下订单前获取最新的相关信息，并应验证此类信息是否为最新且完整。
NVIDIA 产品根据订单确认时提供的 NVIDIA 标准销售条款和条件进行销售，除非 NVIDIA 和客户的授权代表签署的单独销售协议（“销售条款”）另有约定。NVIDIA 特此明确反对将任何客户一般条款和条件应用于购买本文档中引用的 NVIDIA 产品。本文档既不直接也不间接地形成合同义务。
NVIDIA 产品并非设计、授权或保证适用于医疗、军事、航空、航天或生命维持设备，也不适用于 NVIDIA 产品的故障或故障可能合理预期会导致人身伤害、死亡或财产或环境损害的应用。NVIDIA 对在上述设备或应用中包含和/或使用 NVIDIA 产品不承担任何责任，因此，此类包含和/或使用风险由客户自行承担。
NVIDIA 不保证基于本文档的产品将适用于任何特定用途。NVIDIA 不一定对每个产品的全部参数进行测试。客户全权负责评估和确定本文档中包含的任何信息的适用性，确保产品适合且满足客户计划的应用，并为应用执行必要的测试，以避免应用或产品的默认设置。客户产品设计中的缺陷可能会影响 NVIDIA 产品的质量和可靠性，并可能导致超出本文档中包含的附加或不同的条件和/或要求。对于可能基于或归因于以下原因的任何默认设置、损坏、成本或问题，NVIDIA 不承担任何责任：(i) 以任何与本文档相悖的方式使用 NVIDIA 产品，或 (ii) 客户产品设计。
本文档未授予任何 NVIDIA 专利权、版权或其他 NVIDIA 知识产权下的任何明示或暗示的许可。NVIDIA 发布的关于第三方产品或服务的信息不构成 NVIDIA 授予使用此类产品或服务的许可，也不构成对其的保证或认可。使用此类信息可能需要获得第三方的专利或其他知识产权下的许可，或获得 NVIDIA 的专利或其他 NVIDIA 知识产权下的许可。
仅当事先获得 NVIDIA 书面批准，且在复制时不进行更改并完全遵守所有适用的出口法律和法规，并附带所有相关的条件、限制和声明时，才允许复制本文档中的信息。
本文件以及所有 NVIDIA 设计规范、参考板、文件、图纸、诊断程序、列表和其他文档（统称为“资料”，单独或合称）均“按现状”提供。NVIDIA 对这些资料不作任何明示、暗示、法定或其他方面的保证，并且明确声明不承担所有关于不侵权、适销性和针对特定用途适用性的默示保证。在法律未禁止的范围内，NVIDIA 在任何情况下均不对任何损害承担责任，包括但不限于因使用本文件而引起的任何直接、间接、特殊、附带、惩罚性或后果性损害，无论何种原因造成，也无论基于何种责任理论，即使 NVIDIA 已被告知可能发生此类损害。尽管客户可能因任何原因遭受任何损害，NVIDIA 对客户就本文所述产品承担的总体和累积责任应根据产品的销售条款进行限制。
6. OpenCL#
OpenCL 是 Apple Inc. 的商标，已授权 Khronos Group Inc. 使用。
7. 商标#
NVIDIA、NVIDIA 徽标、CUDA、DGX、DGX-1、DGX-2、DGX-A100、Tesla 和 Quadro 是 NVIDIA Corporation 在美国和其他国家/地区的商标和/或注册商标。其他公司和产品名称可能是与其相关的各自公司的商标。