1. Tegra平台的CUDA

本应用说明概述了NVIDIA® Tegra®内存架构，并探讨了将代码从x86系统连接的独立GPU(dGPU)移植到Tegra®集成GPU(iGPU)时的注意事项。文中还讨论了EGL互操作性。

2. 概述

本文档概述了NVIDIA® Tegra®内存架构，并提供了将代码从x86系统连接的独立GPU(dGPU)移植到Tegra®集成GPU(iGPU)的注意事项。同时还讨论了EGL互操作性。

本指南面向已经熟悉CUDA®和C/C++编程，并希望为Tegra® SoC开发应用程序的开发者。

《CUDA C++编程指南》和《CUDA C++最佳实践指南》中提供的性能指南、最佳实践、术语和通用信息适用于所有支持CUDA的GPU架构，包括Tegra®设备。

CUDA C++编程指南和CUDA C最佳实践指南可通过以下网站获取：

CUDA C++ 编程指南：

https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html

CUDA C++ 最佳实践指南：

https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html

3. 内存管理

在Tegra®设备中，CPU（主机）和集成GPU共享SoC DRAM内存。具有独立DRAM内存的独立GPU可以通过PCIe或NVLink连接到Tegra设备。目前该功能仅在NVIDIA DRIVE平台上得到支持。

dGPU连接的Tegra®内存系统概览如图1所示。

图1 连接dGPU的Tegra内存系统

在Tegra系统中，设备内存、主机内存和统一内存都分配在同一个物理SoC DRAM上。而在独立GPU(dGPU)上，设备内存则分配在dGPU的DRAM中。Tegra系统的缓存行为与配备独立GPU的x86系统不同。表1展示了Tegra系统中不同类型内存的缓存和访问行为。

表 1 Tegra系统中不同类型内存的特性

内存类型	CPU	集成显卡	Tegra连接独立显卡
设备内存	无法直接访问	已缓存	已缓存
可分页主机内存	已缓存	无法直接访问	无法直接访问
固定主机内存	计算能力低于7.2时无缓存。 计算能力大于等于7.2时有缓存。	无缓存	无缓存
统一内存	已缓存	已缓存	不支持

在Tegra平台上，由于设备内存、主机内存和统一内存都分配在同一物理SoC DRAM上，因此可以避免重复的内存分配和数据传输。

3.1. I/O一致性

I/O一致性（也称为单向一致性）是一项功能，允许像GPU这样的I/O设备读取CPU缓存中的最新更新。当CPU和GPU共享同一物理内存时，该功能消除了执行CPU缓存管理操作的需求。但由于这种一致性是单向的，仍需要执行GPU缓存管理操作。请注意，当使用托管内存或互操作内存时，CUDA驱动程序会在内部执行GPU缓存管理操作。

从Xavier SOC开始的Tegra设备支持I/O一致性。应用程序无需修改代码即可从这一硬件特性中获益（参见下文第2点）。

以下功能依赖于I/O一致性支持：

1. cudaHostRegister()/cuMemHostRegister() 仅在支持I/O一致性的平台上可用。可以使用设备属性 cudaDevAttrHostRegisterSupported / CU_DEVICE_ATTRIBUTE_HOST_REGISTER_SUPPORTED 来查询主机注册支持情况。

2. 使用cudaMallocHost()/cuMemHostAlloc()/cuMemAllocHost()分配的固定内存的CPU缓存仅在I/O一致的平台上启用。

3.2. 估算集成GPU设备上的总可分配设备内存

cudaMemGetInfo() API 返回可用于GPU分配的内存空闲总量和总量的快照。如果有其他客户端分配内存，空闲内存可能会发生变化。

独立GPU拥有专用的显存(VIDMEM)，与CPU内存分离。cudaMemGetInfo API返回的是独立GPU中可用内存的快照。

在Tegra SoC上，集成GPU与CPU及其他Tegra引擎共享DRAM内存。CPU可以通过将DMAR内容移动到SWAP区域或反向操作来控制DRAM内容并释放DRAM内存。当前cudaMemGetInfo API未计入SWAP内存区域。由于CPU可能通过将页面移动到SWAP区域来释放部分DRAM空间，cudaMemGetInfo API返回的数值可能小于实际可分配内存。为估算可分配设备内存量，CUDA应用开发者应考虑以下因素：

在Linux和Android平台上： 设备可分配内存主要取决于交换空间和主内存的总大小及剩余空间。以下要点可帮助用户估算不同情况下的设备可分配内存总量：

• 主机分配内存 = 总使用物理内存 - 设备分配内存

• 如果（主机分配的内存 < 可用交换空间），则设备可分配内存 = 总物理内存 - 已分配设备内存

• 如果（主机分配内存 > 空闲交换空间），则设备可分配内存 = 总物理内存 - (主机分配内存 - 空闲交换空间)

这里，

• 设备已分配内存是指已在设备上分配的内存。可以从/proc/meminfo中的NvMapMemUsed字段获取，或者从/sys/kernel/debug/nvmap/iovmm/clients的total字段获取。

• 可使用free -m命令获取已使用的总物理内存。Mem行中的used字段即表示该信息。

• 总物理内存是从/proc/meminfo文件中的MemTotal字段获取的。

• 可以使用free -m命令查找空闲交换空间。Swap行中的free字段即表示该信息。

• 如果free命令不可用，可以通过/proc/meminfo获取相同的信息：

  • 已使用的物理内存总量 = MemTotal – MemFree

  • 空闲交换空间 = SwapFree

在QNX平台上： QNX不使用交换空间，因此cudaMemGetInfo.free将是对可分配设备内存的合理估计，因为没有交换空间可以将内存页面移动到交换区域。

4. 移植注意事项

最初为x86系统上的独立GPU(dGPU)开发的CUDA应用程序，可能需要进行修改才能在Tegra系统上高效运行。本节介绍了将此类应用程序移植到Tegra系统时的注意事项，例如选择合适的内存缓冲区类型(固定内存、统一内存等)以及在集成GPU(iGPU)和独立GPU(dGPU)之间进行选择，以实现应用程序的高效性能。

4.1. 内存选择

CUDA应用程序可以使用多种内存缓冲区，例如设备内存、可分页主机内存、固定内存和统一内存。尽管这些内存缓冲区类型分配在同一物理设备上，但每种类型具有不同的访问和缓存行为，如表1所示。选择最适合的内存缓冲区类型对应用程序的高效执行至关重要。

设备内存

对于访问权限仅限于集成GPU(iGPU)的缓冲区，应使用设备内存。例如，在包含多个内核的应用程序中，可能存在仅被应用程序中间内核用作输入或输出的缓冲区。这些缓冲区仅由iGPU访问。此类缓冲区应分配设备内存。

可分页主机内存

对于只能在CPU上访问的缓冲区，使用可分页主机内存。

固定内存

具有不同计算能力的Tegra系统在I/O一致性方面表现出不同的行为。例如，计算能力大于或等于7.2的Tegra系统具有I/O一致性，而其他系统则不具备。在具有I/O一致性的Tegra系统上，固定内存的CPU访问时间与可分页主机内存相当，因为它被缓存在CPU上。然而，在不具备I/O一致性的Tegra系统上，固定内存的CPU访问时间较高，因为它没有被缓存在CPU上。

对于小型缓冲区，建议使用固定内存，因为对此类缓冲区而言缓存效果可以忽略不计，而且与统一内存不同，固定内存不会产生任何额外开销。如果没有额外开销，在iGPU上访问模式不缓存友好的情况下，固定内存也适用于大型缓冲区。对于大型缓冲区，当缓冲区在iGPU上仅以合并方式访问一次时，其在iGPU上的性能可以与统一内存相当。

统一内存

统一内存会被缓存在iGPU和CPU上。在Tegra平台上，应用程序中使用统一内存需要在内核启动、同步和预取提示调用期间执行额外的连贯性和缓存维护操作。对于计算能力低于7.2的Tegra系统，由于缺乏I/O一致性，这种一致性维护开销会略高一些。

在具有I/O一致性（计算能力为7.2或更高）的Tegra设备上，统一内存被缓存在CPU和iGPU上，对于iGPU和CPU频繁访问且iGPU上的访问具有重复性的大型缓冲区，统一内存更为可取，因为重复访问可以抵消缓存维护成本。在不具有I/O一致性（计算能力低于7.2）的Tegra设备上，对于CPU和iGPU频繁访问且iGPU上的访问不具有重复性的大型缓冲区，统一内存仍然比固定内存更可取，因为固定内存不会同时缓存在CPU和iGPU上。这样，应用程序可以利用CPU上的统一内存缓存。

可以使用固定内存或统一内存来减少CPU与集成GPU之间的数据传输开销，因为这两种内存都可以直接从CPU和集成GPU访问。在应用程序中，必须在主机和集成GPU上均可访问的输入和输出缓冲区，可以使用统一内存或固定内存进行分配。

注意

统一内存模型需要驱动程序和系统软件来管理当前Tegra SOC上的数据一致性。软件管理的一致性本质上具有不确定性，在安全关键场景中不推荐使用。这类应用更适合采用零拷贝内存（固定内存）。

评估统一内存开销、固定内存缓存未命中和设备内存数据传输在应用中的影响，以确定正确的内存选择。

4.2. 固定内存

本节提供了将使用固定内存分配的x86系统与dGPU应用移植到Tegra平台的指南。为x86系统外接dGPU开发的CUDA应用程序会使用固定内存来减少数据传输时间，并使数据传输与内核执行时间重叠。关于此主题的具体信息，请参阅以下网站中的"主机与设备间数据传输"和"异步传输与计算重叠"章节。

“主机与设备之间的数据传输”：

https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html#data-transfer-between-host-and-device

"异步计算与重叠传输"：

https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html#asynchronous-transfers-and-overlapping-transfers-with-computation

在没有I/O一致性的Tegra系统上，重复访问固定内存会降低应用程序性能，因为在此类系统中固定内存不会被CPU缓存。

下面展示了一个示例应用，其中对图像应用了一组滤波器和操作（k1、k2和k3）。在配备dGPU的x86系统上分配了固定内存以减少数据传输时间，从而提升整体应用速度。然而，在Tegra设备上运行相同代码会导致readImage()函数的执行时间急剧增加，因为它反复访问未缓存的缓冲区。这会延长整体应用时间。如果readImage()耗时明显高于内核执行时间，建议使用统一内存来缩短readImage()时间。否则，可通过移除不必要的数据传输调用来评估固定内存与统一内存方案，以确定最适合的内存配置。

// Sample code for an x86 system with a discrete GPU
int main()
{
    int *h_a,*d_a,*d_b,*d_c,*d_d,*h_d;
    int height = 1024;
    int width = 1024;
    size_t sizeOfImage = width * height * sizeof(int); // 4MB image

    //Pinned memory allocated to reduce data transfer time
    cudaMallocHost(h_a, sizeOfImage);
    cudaMallocHost(h_d, sizeOfImage);

    //Allocate buffers on GPU
    cudaMalloc(&d_a, sizeOfImage);
    cudaMalloc(&d_b, sizeOfImage);
    cudaMalloc(&d_c, sizeOfImage);
    cudaMalloc(&d_d, sizeOfImage);

    //CPU reads Image;
    readImage(h_a); // Intialize the h_a buffer

    // Transfer image to GPU
    cudaMemcpy(d_a, h_a, sizeOfImage, cudaMemcpyHostToDevice);

    // Data transfer is fast as we used pinned memory
    // ----- CUDA Application pipeline start ----
    k1<<<..>>>(d_a,d_b) // Apply filter 1
    k2<<<..>>>(d_b,d_c)// Apply filter 2
    k3<<<..>>>(d_c,d_d)// Some operation on image data
    // ----- CUDA Application pipeline end ----

    // Transfer processed image to CPU
    cudaMemcpy(h_d, d_d, sizeOfImage, cudaMemcpyDeviceToHost);
    // Data transfer is fast as we used pinned memory

    // Use processed Image i.e h_d in later computations on CPU.
    UseImageonCPU(h_d);
}

// Porting the code on Tegra
int main()
{
    int *h_a,*d_b,*d_c,*h_d;
    int height = 1024;
    int width = 1024;
    size_t sizeOfImage = width * height * sizeof(int); // 4MB image

    //Unified memory allocated for input and output
    //buffer of application pipeline
    cudaMallocManaged(h_a, sizeOfImage,cudaMemAttachHost);
    cudaMallocManaged(h_d, sizeOfImage);

    //Intermediate buffers not needed on CPU side.
    //So allocate them on device memory
    cudaMalloc(&d_b, sizeOfImage);
    cudaMalloc(&d_c, sizeOfImage);

    //CPU reads Image;
    readImage (h_a); // Intialize the h_a buffer
    // ----- CUDA Application pipeline start ----
    // Prefetch input image data to GPU
    cudaStreamAttachMemAsync(NULL, h_a, 0, cudaMemAttachGlobal);
    k1<<<..>>>(h_a,d_b)
    k2<<<..>>>(d_b,d_c)
    k3<<<..>>>(d_c,h_d)
    // Prefetch output image data to CPU
    cudaStreamAttachMemAsync(NULL, h_d, 0, cudaMemAttachHost);
    cudaStreamSynchronize(NULL);
    // ----- CUDA Application pipeline end ----

    // Use processed Image i.e h_d on CPU side.
    UseImageonCPU(h_d);
}

cudaHostRegister()函数

在计算能力低于7.2的Tegra设备上不支持cudaHostRegister()函数，因为这些设备不具备I/O一致性。如果设备不支持cudaHostRegister()，请使用其他固定内存分配函数，例如cudaMallocHost()和cudaHostAlloc()。

GNU 对固定内存的原子操作

Tegra CPU不支持对非缓存内存执行GNU原子操作。由于计算能力低于7.2的Tegra设备上固定内存不会被缓存，因此固定内存也不支持GNU原子操作。

4.3. Tegra平台上统一内存的高效使用

在应用程序中使用统一内存需要在内核启动、同步和预取提示调用时执行额外的连贯性和缓存维护操作。这些操作会与其他GPU工作同步执行，可能导致应用程序中出现不可预测的延迟。

通过提供数据预取提示，可以提升Tegra平台上统一内存的性能。驱动程序可利用这些预取提示来优化一致性操作。除了《CUDA C编程指南》"一致性与并发"章节（详见下方链接）中描述的技术外，还可使用cudaStreamAttachMemAsync()函数进行数据预取：

https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#um-coherency-hd

预取数据。统一内存的预取行为，由附件标志状态变化触发，如表2所示。

表 2 统一内存预取行为随附件标志状态变化

前一个标志	当前标志	预取行为
cudaMemAttachGlobal/cudaMemAttachSingle	cudaMemAttachHost	触发预取至CPU
cudaMemAttachHost	cudaMemAttachGlobal/ cudaMemAttachSingle	触发预取至GPU
cudaMemAttachGlobal	cudaMemAttachSingle	不预取到GPU
cudaMemAttachSingle	cudaMemAttachGlobal	不预取到GPU

以下示例展示了使用cudaStreamAttachMemAsync()预取数据的用法。

注意

然而，Tegra设备不支持使用cudaMemPrefetchAsync()的数据预取技术，该技术在以下网站的CUDA C++编程指南的"性能调优"章节中有详细描述：

https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#um-performance-tuning

注意

QNX系统软件存在一些限制，导致无法实现所有UVM优化功能。因此，在QNX上使用cudaStreamAttachMemAsync()进行预取提示并不会带来性能提升。

__global__ void matrixMul(int *p, int *q, int*r, int hp, int hq, int wp, int wq)
{
// Matrix multiplication kernel code
}
void MatrixMul(int hp, int hq, int wp, int wq)
{
    int *p,*q,*r;
    int i;
    size_t sizeP = hp*wp*sizeof(int);
    size_t sizeQ = hq*wq*sizeof(int);
    size_t sizeR = hp*wq*sizeof(int);

    //Attach buffers 'p' and 'q' to CPU and buffer 'r' to GPU
    cudaMallocManaged(&p, sizeP, cudaMemAttachHost);
    cudaMallocManaged(&q, sizeQ, cudaMemAttachHost);
    cudaMallocManaged(&r, sizeR);
    //Intialize with random values
    randFill(p,q,hp,wp,hq,wq);

    // Prefetch p,q to GPU as they are needed in computation
    cudaStreamAttachMemAsync(NULL, p, 0, cudaMemAttachGlobal);
    cudaStreamAttachMemAsync(NULL, q, 0, cudaMemAttachGlobal);
    matrixMul<<<....>>>(p,q,r, hp,hq,wp,wq);

    // Prefetch 'r' to CPU as only 'r' is needed
    cudaStreamAttachMemAsync(NULL, r, 0, cudaMemAttachHost);
    cudaStreamSynchronize(NULL);

    // Print buffer 'r' values
    for(i = 0; i < hp*wq; i++)
    printf("%d ", r[i]);
}

注意

可以在matrixMul内核代码后添加一个额外的cudaStreamSynchronize(NULL)调用，以避免回调线程导致cudaStreamAttachMemAsync()调用出现不可预测性。

4.4. GPU选择

在配备独立GPU(dGPU)的Tegra系统中，决定CUDA应用程序运行在集成GPU(iGPU)还是独立GPU上会影响应用程序的性能。做出此类决策时需要考虑的因素包括内核执行时间、数据传输时间、数据局部性和延迟。例如，要在dGPU上运行应用程序，必须在SoC和dGPU之间传输数据。如果应用程序在iGPU上运行，则可以避免这种数据传输。

4.5. 同步机制选择

cudaSetDeviceFlags API用于控制CPU线程的同步行为。在CUDA 10.1之前，默认情况下，iGPU上的同步机制使用cudaDeviceBlockingSync标志，该标志会在等待设备完成工作时阻塞CPU线程。这种cudaDeviceBlockingSync标志适合有功耗限制的平台。但对于需要低延迟的平台，则需要手动设置cudaDeviceScheduleSpin标志。从CUDA 10.1开始，每个平台的默认同步标志会根据该平台的优化需求自动确定。更多关于同步标志的信息可在cudaSetDeviceFlags API文档中查阅。

4.6. Tegra平台不支持的CUDA功能

Tegra平台支持CUDA的所有核心功能。以下列出了例外情况。

• cudaHostRegister() 函数在 QNX 系统中不受支持。这是由于 QNX 操作系统的限制。该功能在计算能力大于或等于 7.2 的 Linux 系统中受支持。

• 在计算能力低于7.2的Tegra设备上不支持系统范围的原子操作。

• 连接到Tegra的独立GPU不支持统一内存。

• cudaMemPrefetchAsync() 函数不受支持，因为iGPU上尚不支持具有并发访问的统一内存。

• Tegra平台不支持NVIDIA管理库(NVML)。不过，作为监控资源使用情况的替代方案，可以使用tegrastats。

• 自CUDA 11.5起，在计算能力7.x及更高的L4T和嵌入式Linux Tegra设备上，仅支持事件共享IPC API。Tegra平台仍不支持内存共享IPC API。可以使用EGLStream、NvSci或cuMemExportToShareableHandle() / cuMemImportFromShareableHandle() API在两个进程的CUDA上下文之间进行通信。

• 远程直接内存访问(RDMA)仅在运行L4T或嵌入式Linux的Tegra设备上受支持。

• JIT编译可能会消耗大量CPU和带宽资源，可能干扰系统中的其他工作负载。因此，对于确定性嵌入式应用不建议使用PTX-JIT和NVRTC JIT等JIT编译方式，可以通过针对特定GPU目标进行编译来完全绕过JIT编译。例如：如果要为SM版本87进行编译，请使用此nvcc标志--generate-code arch=compute_87,code=sm_87为该设备创建CUDA二进制文件。这样可以避免首次运行时的JIT编译并提高运行时性能。 在安全上下文中，Tegra设备不支持JIT编译。

• Tegra 不支持点对点 (P2P) 通信调用。

• 在运行QNX的Tegra系统上不支持cuSOLVER库。

• 不支持 nvGRAPH 库。

• CUB 在Tegra产品上处于实验阶段。

有关这些功能的更多信息，请访问以下网站：

IPC:

https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#interprocess-communication

NVSCI:

https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#nvidia-softwarcommunication-interface-interoperability-nvsci

RDMA:

https://docs.nvidia.com/cuda/gpudirect-rdma/index.html

点对点(P2P):

https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#peer-to-peer-memory-access

5. EGL互操作性

互操作(interop)是两种API之间共享资源的高效机制。要与多个API共享数据，每个API都必须单独实现一套互操作机制。

EGL提供了互操作扩展功能，使其能够作为连接各种API的枢纽，从而消除多重互操作的需求，并封装共享资源。任何API必须实现这些扩展才能通过EGL与其他API进行互操作。CUDA支持的EGL互操作类型包括EGLStream、EGLImage和EGLSync。

EGL互操作扩展允许应用程序在不同API之间切换而无需重写代码。例如，一个基于EGLStream的应用程序，其中NvMedia作为生产者而CUDA作为消费者，可以修改为使用OpenGL作为消费者，而无需更改生产者代码。

注意

在DRIVE OS平台上，NVSCI作为安全关键应用中EGL互操作性的替代方案提供。更多详情请参阅NVSCI。

5.1. EGLStream

EGLStream互操作性功能支持将帧序列高效地从一种API传输到另一种API，从而可以利用CPU、GPU、ISP等多种Tegra®引擎。

考虑一个应用场景：摄像头持续捕获图像，将其共享给CUDA进行处理，随后使用OpenGL渲染这些图像。在该应用中，图像帧在NvMedia、CUDA和OpenGL之间共享。若缺乏EGLStream互操作性，应用将需要包含多个互操作接口及API间的冗余数据传输。EGLStream采用单一生产者与单一消费者模式。

EGLStream提供以下优势：

• 在生产者与消费者之间高效传输帧。

• 隐式同步处理。

• 跨进程支持。

• 支持dGPU和iGPU。

• 支持Linux、QNX和Android操作系统。

5.1.1. EGLStream数据流

EGLStream流程包含以下步骤：

1. 初始化生产者和消费者API

2. 创建一个EGLStream并连接消费者和生产者。

   注意

   EGLStream使用eglCreateStreamKHR()创建，使用eglDestroyStreamKHR()销毁。

   消费者应始终在生产者之前连接到EGLStream。

   更多信息请参阅以下网站的EGLStream规范：https://www.khronos.org/registry/EGL/extensions/KHR/EGL_KHR_stream.txt

3. 分配用于EGL帧的内存。

4. 生产者填充一个EGL帧并将其呈现给EGLStream。

5. 消费者从EGLStream获取帧，并在处理完成后将其释放回EGLStream。

6. 生产者从EGLStream收集消费者释放的帧。

7. 生产者向EGLStream呈现相同的帧或新帧。

8. 步骤4-7会重复执行，直到任务完成，期间可能使用旧帧或新帧。

9. 消费者和生产者从EGLStream断开连接。

10. 释放用于EGL帧的内存。

11. 取消初始化生产者和消费者API。

EGLStream应用流程如图2所示。

图2 EGLStream数据流

CUDA的生产者与消费者函数列于表3中。

表 3 CUDA生产者与消费者函数

角色	功能	API
生产者	将生产者连接到EGLStream	cuEGLStreamProducerConnect() cudaEGLStreamProducerConnect()
	将帧呈现给EGLStream	cuEGLStreamProducerPresentFrame() cudaEGLStreamProducerPresentFrame()
	获取已释放的帧	cuEGLStreamProducerReturnFrame() cudaEGLStreamProducerReturnFrame()
	断开与EGLStream的连接	cuEGLStreamProducerDisconnect() cudaEGLStreamProducerDisconnect()
消费者	将消费者连接到EGLStream	cuEGLStreamConsumerConnect() cuEGLStreamConsumeConnectWithFlags() cudaEGLStreamConsumerConnect() cudaEGLStreamConsumerConnectWithFlags()
	从EGLStream获取帧	cuEGLStreamConsumerAcquireFrame() cudaEGLStreamConsumerAcquireFrame()
	释放已消耗的帧	cuEGLStreamConsumerReleaseFrame() cudaEGLStreamConsumerReleaseFrame()
	断开与EGLStream的连接	cuEGLStreamConsumerDisconnect() cudaEGLStreamConsumerDisconnect()

5.1.2. CUDA作为生产者

当CUDA作为生产者时，支持的消费者包括CUDA、NvMedia和OpenGL。表3列出了当CUDA作为生产者时需使用的API函数。除了与EGLStream的连接和断开操作外，所有API调用都是非阻塞式的。

以下生产者端的步骤在随后的示例代码中展示：

1. 准备一个框架（第3-19行）。

2. 将生产者连接到EGLStream（第21行）。

3. 填充帧并呈现给EGLStream（第23-25行）。

4. 从EGLStream获取已释放的帧（第27行）。

5. 任务完成后断开消费者连接。(第31行)

void ProducerThread(EGLStreamKHR eglStream) {
 //Prepares frame
 cudaEglFrame* cudaEgl = (cudaEglFrame *)malloc(sizeof(cudaEglFrame));
 cudaEgl->planeDesc[0].width = WIDTH;
 cudaEgl->planeDesc[0].depth = 0;
 cudaEgl->planeDesc[0].height = HEIGHT;
 cudaEgl->planeDesc[0].numChannels = 4;
 cudaEgl->planeDesc[0].pitch = WIDTH * cudaEgl->planeDesc[0].numChannels;
 cudaEgl->frameType = cudaEglFrameTypePitch;
 cudaEgl->planeCount = 1;
 cudaEgl->eglColorFormat = cudaEglColorFormatARGB;
 cudaEgl->planeDesc[0].channelDesc.f=cudaChannelFormatKindUnsigned
 cudaEgl->planeDesc[0].channelDesc.w = 8;
 cudaEgl->planeDesc[0].channelDesc.x = 8;
 cudaEgl->planeDesc[0].channelDesc.y = 8;
 cudaEgl->planeDesc[0].channelDesc.z = 8;
 size_t numElem = cudaEgl->planeDesc[0].pitch * cudaEgl->planeDesc[0].height;
 // Buffer allocated by producer
 cudaMalloc(&(cudaEgl->pPitch[0].ptr), numElem);
 //CUDA producer connects to EGLStream
 cudaEGLStreamProducerConnect(&conn, eglStream, WIDTH, HEIGHT))
 // Sets all elements in the buffer to 1
 K1<<<...>>>(cudaEgl->pPitch[0].ptr, 1, numElem);
 // Present frame to EGLStream
 cudaEGLStreamProducerPresentFrame(&conn, *cudaEgl, NULL);

 cudaEGLStreamProducerReturnFrame(&conn, cudaEgl, eglStream);
 .
 .
 //clean up
 cudaEGLStreamProducerDisconnect(&conn);

 .
}

一个帧由cudaEglFramestructure表示。cudaEglFrame中的frameType参数指示了帧的内存布局。支持的内存布局包括CUDA数组和设备指针。如果帧的宽度和高度值与cudaEGLStreamProducerConnect()中指定的值不匹配，将导致未定义行为。在示例中，CUDA生产者发送单个帧，但它可以通过循环发送多个帧。CUDA最多只能向EGLStream呈现64个活动帧。

cudaEGLStreamProducerReturnFrame()调用会等待直到从消费者处接收到释放的帧。一旦CUDA生产者向EGLstream提交了第一帧，在生产者断开连接之前，至少会有一帧可供消费者获取。这可以防止从EGLStream中移除最后一帧，否则会阻塞cudaEGLStreamProducerReturnFrame()。

使用EGL_NV_stream_reset扩展将EGLStream属性EGL_SUPPORT_REUSE_NV设为false，以允许从EGLStream中移除最后一帧。这样可以从EGLStream中移除或返回最后一帧。

5.1.3. CUDA作为消费者

当CUDA作为消费者时，支持的生成器包括CUDA、OpenGL、NvMedia、Argus和Camera。表3列出了当CUDA作为消费者时需要使用的API函数。除了与EGLStream的连接和断开操作外，所有API调用都是非阻塞的。

以下消费者端步骤在随后的示例代码中展示：

1. 将消费者连接到EGLStream（第5行）。

2. 从EGLStream获取帧(第8-10行)。

3. 在消费者端处理帧数据（第16行）。

4. 将帧释放回EGLStream（第19行）。

5. 任务完成后断开消费者连接（第22行）。

void ConsumerThread(EGLStreamKHR eglStream) {
.
.
//Connect consumer to EGLStream
cudaEGLStreamConsumerConnect(&conn, eglStream);
// consumer acquires a frame
unsigned int timeout = 16000;
cudaEGLStreamConsumerAcquireFrame(& conn, &cudaResource, eglStream, timeout);
//consumer gets a cuda object pointer
cudaGraphicsResourceGetMappedEglFrame(&cudaEgl, cudaResource, 0, 0);
size_t numElem = cudaEgl->planeDesc[0].pitch * cudaEgl->planeDesc[0].height;
.
.
int checkIfOne = 1;
// Checks if each value in the buffer is 1, if any value is not 1, it sets checkIfOne = 0.
K2<<<...>>>(cudaEgl->pPitch[0].ptr, 1, numElem, checkIfOne);
.
.
cudaEGLStreamConsumerReleaseFrame(&conn, cudaResource, &eglStream);
.
.
cudaEGLStreamConsumerDisconnect(&conn);
.
}

在示例代码中，CUDA消费者接收单个帧，但它也可以通过循环接收多个帧。如果CUDA消费者未能在指定时间限制内使用cudaEGLStreamConsumerAcquireFrame()接收到新帧，它将从EGLStream重新获取前一帧。时间限制由timeout参数指定。

应用程序可以使用eglQueryStreamKHR()来查询新帧的可用性。如果消费者使用已释放的帧，将导致未定义行为。消费者行为仅针对读取操作进行了定义。当消费者对帧进行写入操作时，其行为是未定义的。

如果在连接到EGLStream时CUDA上下文被销毁，流将被置于EGL_STREAM_STATE_DISCONNECTED_KHR状态且连接句柄将失效。

5.1.4. 隐式同步

EGLStream在应用程序中提供隐式同步功能。例如，在前面的代码示例中，生产者和消费者线程并行运行，K1和K2内核进程访问同一帧，但可以确保消费者线程中的K2内核仅在生产者线程中的K1内核完成后才会执行。cudaEGLStreamConsumerAcquireFrame()函数会在GPU端等待直到K1完成，从而确保生产者和消费者之间的同步。变量checkIfOne在消费者线程的K2内核中永远不会被设置为0。

类似地，生产者线程中的cudaEGLStreamProducerReturnFrame()确保只有在K2完成且消费者释放帧后才能获取该帧。这些非阻塞调用允许CPU在间隙执行其他计算，因为同步操作由GPU端处理。

EGLStreams_CUDA_Interop CUDA示例代码详细展示了EGLStream的用法。

5.1.5. 生产者与消费者之间的数据传输

当生产者和消费者位于同一设备上时，可以避免它们之间的数据传输。然而，在包含独立GPU（dGPU）的Tegra®平台（例如NVIDIA DRIVE™ PX 2）中，生产者和消费者可能位于不同设备上。在这种情况下，系统内部需要通过额外的内存拷贝将帧数据在Tegra® SoC DRAM和dGPU DRAM之间移动。EGLStream允许生产者和消费者在任何GPU上运行，而无需修改代码。

注意

在Tegra®设备连接到独立GPU(dGPU)的系统上，如果生产者帧使用CUDA数组，则生产者和消费者应位于同一GPU上。但如果生产者帧使用CUDA设备指针，消费者可以存在于任何GPU上。

5.1.6. EGLStream 流水线

应用程序可以在流水线中使用多个EGL流，将帧从一个API传递到另一个API。例如，在某个应用中，NvMedia将帧发送到CUDA进行计算，计算完成后CUDA再将同一帧发送到OpenGL进行渲染。

EGLStream流水线如图3所示。

图3 EGLStream流水线

NvMedia和CUDA分别作为生产者和消费者连接到一个EGLStream。CUDA和OpenGL分别作为生产者和消费者连接到另一个EGLStream。

在流水线模式下使用多个EGLStreams可以灵活地将帧跨多个API传输，无需分配额外内存或进行显式数据传输。通过上述EGLStream流水线发送帧涉及以下步骤。

1. NvMedia将一帧图像发送到CUDA进行处理。

2. CUDA 使用该框架进行计算，并将结果发送至 OpenGL 进行渲染。

3. OpenGL消耗帧并将其释放回CUDA。

4. CUDA将帧释放回NvMedia。

上述步骤可以在循环中执行，以便于在EGLStream管道中传输多个帧。

5.2. EGLImage

EGLImage互操作允许EGL客户端API与其他EGL客户端API共享图像数据。例如，应用程序可以使用EGLImage互操作在OpenGL纹理与CUDA之间共享数据，而无需分配额外内存。单个EGLImage对象可被多个客户端API共享并进行修改。

EGLImage互操作不提供隐式同步。应用程序必须维护同步以避免竞态条件。

注意

EGLImage通过eglCreateImageKHR()创建，并通过eglDestroyImageKHR()销毁。

更多信息请参阅以下网站上的EGLImage规范：

https://www.khronos.org/registry/EGL/extensions/KHR/EGL_KHR_image_base.txt

5.2.1. CUDA与EGLImage的互操作性

CUDA支持与EGLImage的互操作，允许CUDA读取或修改EGLImage的数据。EGLImage可以是单平面或多平面资源。在CUDA中，单平面EGLImage对象表示为CUDA数组或设备指针。类似地，多平面EGLImage对象表示为设备指针或CUDA数组的集合。EGLImage在运行Linux、QNX或Android操作系统的Tegra®设备上受支持。

使用cudaGraphicsEGLRegisterImage() API将EGLImage对象注册到CUDA。向CUDA注册EGLImage会创建一个图形资源对象。应用程序可以通过cudaGraphicsResourceGetMappedEglFrame()从图形资源对象获取帧。在CUDA中，帧以cudaEglFrame结构体表示。cudaEglFrame中的frameType参数指示该帧是CUDA设备指针还是CUDA数组。对于单平面图形资源，应用程序可以直接使用cudaGraphicsResourceGetMappedPointer()或cudaGraphicsSubResourceGetMappedArray()分别获取设备指针或CUDA数组。CUDA数组可以绑定到纹理或表面引用以便在内核中访问。此外，多维CUDA数组可以通过cudaMemcpy3D()进行读写操作。

注意

无法从CUDA对象创建EGLImage。cudaGraphicsEGLRegisterImage()函数仅在Tegra®设备上受支持。此外，cudaGraphicsEGLRegisterImage()仅接受'0'标志作为参数，其他API标志留作未来使用。

以下示例代码展示了EGLImage的互操作性。在代码中，使用OpenGL纹理创建了一个EGLImage对象eglImage。该eglImage对象在CUDA中被映射为CUDA数组pArray。将pArray数组绑定到表面对象，以允许在changeTexture中修改OpenGL纹理。函数checkBuf()用于检查纹理是否已更新为新值。

int width = 256;
int height = 256;
int main()
{
 .
 .
 unsigned char *hostSurf;
 unsigned char *pSurf;
 CUarray pArray;
 unsigned int bufferSize = WIDTH * HEIGHT * 4;
 pSurf= (unsigned char *)malloc(bufferSize); hostSurf = (unsigned char *)malloc(bufferSize);
 // Initialize the buffer
 for(int y = 0; y < HEIGHT; y++)
 {
    for(int x = 0; x < WIDTH; x++)
    {
    pSurf[(y*WIDTH + x) * 4 ] = 0; pSurf[(y*WIDTH + x) * 4 + 1] = 0;
    pSurf[(y*WIDTH + x) * 4 + 2] = 0; pSurf[(y*WIDTH + x) * 4 + 3] = 0;
    }
 }

 // NOP call to error-check the above glut calls
 GL_SAFE_CALL({});

 //Init texture
 GL_SAFE_CALL(glGenTextures(1, &tex));
 GL_SAFE_CALL(glBindTexture(GL_TEXTURE_2D, tex));
 GL_SAFE_CALL(glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA, WIDTH, HEIGHT, 0, GL_RGBA, GL_UNSIGNED_BYTE, pSurf));

 EGLDisplay eglDisplayHandle = eglGetCurrentDisplay();
 EGLContext eglCtx = eglGetCurrentContext();

 // Create the EGL_Image
 EGLint eglImgAttrs[] = { EGL_IMAGE_PRESERVED_KHR, EGL_FALSE, EGL_NONE, EGL_NONE };
 EGLImageKHR eglImage = eglCreateImageKHR(eglDisplayHandle, eglCtx, EGL_GL_TEXTURE_2D_KHR, (EGLClientBuffer)(intptr_t)tex, eglImgAttrs);
 glFinish();
 glTexSubImage2D(GL_TEXTURE_2D, 0, 0, 0, WIDTH, HEIGHT, GL_RGBA, GL_UNSIGNED_BYTE, pSurf);
 glFinish();

 // Register buffer with CUDA
cuGraphicsEGLRegisterImage(&pResource, eglImage,0);

 //Get CUDA array from graphics resource object
 cuGraphicsSubResourceGetMappedArray( &pArray, pResource, 0, 0);

 cuCtxSynchronize();

 //Create a CUDA surface object from pArray
 CUresult status = CUDA_SUCCESS;
 CUDA_RESOURCE_DESC wdsc;
 memset(&wdsc, 0, sizeof(wdsc));
 wdsc.resType = CU_RESOURCE_TYPE_ARRAY; wdsc.res.array.hArray = pArray;
 CUsurfObject writeSurface;
 cuSurfObjectCreate(&writeSurface, &wdsc);

 dim3 blockSize(32,32);
 dim3 gridSize(width/blockSize.x,height/blockSize.y);
 // Modifies the OpenGL texture using CUDA surface object
 changeTexture<<<gridSize, blockSize>>>(writeSurface, width, height);
 cuCtxSynchronize();

 CUDA_MEMCPY3D cpdesc;
 memset(&cpdesc, 0, sizeof(cpdesc));
 cpdesc.srcXInBytes = cpdesc.srcY = cpdesc.srcZ = cpdesc.srcLOD = 0;
 cpdesc.dstXInBytes = cpdesc.dstY = cpdesc.dstZ = cpdesc.dstLOD = 0;
 cpdesc.srcMemoryType = CU_MEMORYTYPE_ARRAY; cpdesc.dstMemoryType = CU_MEMORYTYPE_HOST;
 cpdesc.srcArray = pArray; cpdesc.dstHost = (void *)hostSurf;
 cpdesc.WidthInBytes = WIDTH * 4; cpdesc.Height = HEIGHT; cpdesc.Depth = 1;

 //Copy CUDA surface object values to hostSurf
 cuMemcpy3D(&cpdesc);

 cuCtxSynchronize();

 unsigned char* temp = (unsigned char*)(malloc(bufferSize * sizeof(unsigned char)));
 // Get the modified texture values as
 GL_SAFE_CALL(glGetTexImage(GL_TEXTURE_2D, 0, GL_RGBA, GL_UNSIGNED_BYTE,(void*)temp));
 glFinish();
 // Check if the OpenGL texture got modified values
 checkbuf(temp,hostSurf);

 // Clean up CUDA
 cuGraphicsUnregisterResource(pResource);
 cuSurfObjectDestroy(writeSurface);
 .
 .
}
__global__ void changeTexture(cudaSurfaceObject_t arr, unsigned int width, unsigned int height){
 unsigned int x = threadIdx.x + blockIdx.x * blockDim.x;
 unsigned int y = threadIdx.y + blockIdx.y * blockDim.y;
 uchar4 data = make_uchar4(1, 2, 3, 4);
 surf2Dwrite(data, arr, x * 4, y);
}
void checkbuf(unsigned char *ref, unsigned char *hostSurf) {
 for(int y = 0; y < height*width*4; y++){
 if (ref[y] != hostSurf[y])
 printf("mis match at %d\n",y);
 }
}

由于EGLImage不提供隐式同步，上述示例应用程序使用了glFinish()和cudaThreadSynchronize()调用来实现同步。这两个调用都会阻塞CPU线程。为避免阻塞CPU线程，可使用EGLSync来提供同步。下一节将展示一个使用EGLImage和EGLSync的示例。

5.3. EGLSync

EGLSync是一种跨API同步原语。它允许EGL客户端API与其他EGL客户端API共享其同步对象。例如，应用程序可以使用EGLSync互操作将OpenGL同步对象与CUDA共享。

注意

EGLSync对象通过eglCreateSyncKHR()创建，并通过eglDestroySyncKHR()销毁。

更多信息请参阅以下网站上的EGLSync规范：

https://www.khronos.org/registry/EGL/extensions/KHR/EGL_KHR_fence_sync.txt

5.3.1. CUDA与EGLSync的互操作性

在成像应用中，当两个客户端在GPU上运行并共享资源时，由于缺乏跨API的GPU同步对象，客户端不得不使用CPU端同步来避免竞态条件。CUDA与EGLSync的互操作性使应用能够直接在CUDA和其他客户端API之间交换同步对象。这消除了对CPU端同步的需求，并允许CPU执行其他任务。在CUDA中，EGLSync对象被映射为CUDA事件。

注意

目前CUDA与EGLSync的互操作仅在Tegra®设备上受支持。

5.3.2. 从CUDA事件创建EGLSync

以下示例代码展示了如何从CUDA事件创建EGLSync对象。请注意，从CUDA事件创建EGLSync对象的操作应在记录CUDA事件后立即执行。

EGLDisplay dpy = eglGetCurrentDisplay();
// Create CUDA event
cudaEvent_t event;
cudaStream_t *stream;
cudaEventCreate(&event);
cudaStreamCreate(&stream);
// Record the event with cuda event
cudaEventRecord(event, stream);
const EGLAttrib attribs[] = {
 EGL_CUDA_EVENT_HANDLE_NV, (EGLAttrib )event,
 EGL_NONE
};
//Create EGLSync from the cuda event
eglsync = eglCreateSync(dpy, EGL_NV_CUDA_EVENT_NV, attribs);
//Wait on the sync
eglWaitSyncKHR(...);

注意

在创建EGLSync对象之前初始化一个CUDA事件，以避免未定义行为。

5.3.3. 从EGLSync创建CUDA事件

以下示例代码展示了如何从EGLSync对象创建CUDA事件。

EGLSync eglsync;
EGLDisplay dpy = eglGetCurrentDisplay();
// Create an eglSync object from openGL fense sync object
eglsync = eglCreateSyncKHR(dpy, EGL_SYNC_FENCE_KHR, NULL);
cudaEvent_t event;
cudaStream_t* stream;
cudaStreamCreate(&stream);
// Create CUDA event from eglSync
cudaEventCreateFromEGLSync(&event, eglSync, cudaEventDefault);
// Wait on the cuda event. It waits on GPU till OpenGL finishes its
// task
cudaStreamWaitEvent(stream, event, 0);

注意

从EGLSync对象创建的事件不支持cudaEventRecord()和cudaEventElapsedTime()函数。

下面重写了EGLImage部分中的相同示例，以说明EGLSync互操作的使用。在示例代码中，诸如glFinish()和cudaThreadSynchronize()之类的CPU阻塞调用被替换为EGLSync互操作调用。

int width = 256;
int height = 256;
int main()
{
 .
 .
 unsigned char *hostSurf;
 unsigned char *pSurf;
 cudaArray_t pArray;
 unsigned int bufferSize = WIDTH * HEIGHT * 4;
 pSurf= (unsigned char *)malloc(bufferSize); hostSurf = (unsigned char *)malloc(bufferSize);
 // Intialize the buffer
 for(int y = 0; y < bufferSize; y++)
 pSurf[y] = 0;

 //Init texture
 GL_SAFE_CALL(glGenTextures(1, &tex));
 GL_SAFE_CALL(glBindTexture(GL_TEXTURE_2D, tex));
 GL_SAFE_CALL(glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA, WIDTH, HEIGHT, 0, GL_RGBA, GL_UNSIGNED_BYTE, pSurf));
 EGLDisplay eglDisplayHandle = eglGetCurrentDisplay();
 EGLContext eglCtx = eglGetCurrentContext();

 cudaEvent_t cuda_event;
 cudaEventCreateWithFlags(cuda_event, cudaEventDisableTiming);
 EGLAttribKHR eglattrib[] = { EGL_CUDA_EVENT_HANDLE_NV, (EGLAttrib) cuda_event, EGL_NONE};
 cudaStream_t* stream;
 cudaStreamCreateWithFlags(&stream,cudaStreamDefault);

 EGLSyncKHR eglsync1, eglsync2;
 cudaEvent_t egl_event;

 // Create the EGL_Image
 EGLint eglImgAttrs[] = { EGL_IMAGE_PRESERVED_KHR, EGL_FALSE, EGL_NONE, EGL_NONE };
 EGLImageKHR eglImage = eglCreateImageKHR(eglDisplayHandle, eglCtx, EGL_GL_TEXTURE_2D_KHR, (EGLClientBuffer)(intptr_t)tex, eglImgAttrs);

 glTexSubImage2D(GL_TEXTURE_2D, 0, 0, 0, WIDTH, HEIGHT, GL_RGBA, GL_UNSIGNED_BYTE, pSurf);
 //Creates an EGLSync object from GL Sync object to track
 //finishing of copy.
 eglsync1 = eglCreateSyncKHR(eglDisplayHandle, EGL_SYNC_FENCE_KHR, NULL);

 //Create CUDA event object from EGLSync obejct
 cuEventCreateFromEGLSync(&egl_event, eglsync1, cudaEventDefault);

 //Waiting on GPU to finish GL copy
 cuStreamWaitEvent(stream, egl_event, 0);

 // Register buffer with CUDA
 cudaGraphicsEGLRegisterImage(&pResource, eglImage, cudaGraphicsRegisterFlagsNone);
 //Get CUDA array from graphics resource object
 cudaGraphicsSubResourceGetMappedArray( &pArray, pResource, 0, 0);
 .
 .
 //Create a CUDA surface object from pArray
 struct cudaResourceDesc resDesc;
 memset(&resDesc, 0, sizeof(resDesc));
 resDesc.resType = cudaResourceTypeArray; resDesc.res.array.array = pArray;
 cudaSurfaceObject_t inputSurfObj = 0;
 cudaCreateSurfaceObject(&inputSurfObj, &resDesc);

 dim3 blockSize(32,32);
 dim3 gridSize(width/blockSize.x,height/blockSize.y);
 // Modifies the CUDA array using CUDA surface object
 changeTexture<<<gridSize, blockSize>>>(inputSurfObj, width, height);
 cuEventRecord(cuda_event, stream);
 //Create EGLsync object from CUDA event cuda_event
 eglsync2 = eglCreateSync64KHR(dpy, EGL_SYNC_CUDA_EVENT_NV, eglattrib);
 //waits till kernel to finish
 eglWaitSyncKHR(eglDisplayHandle, eglsync2, 0);
 .
 //Copy modified pArray values to hostSurf
 .
 unsigned char* temp = (unsigned char*)(malloc(bufferSize * sizeof(unsigned char)));
 // Get the modified texture values
 GL_SAFE_CALL(glGetTexImage(GL_TEXTURE_2D, 0, GL_RGBA, GL_UNSIGNED_BYTE,(void*)temp));
 .
 .
 // This function check if the OpenGL texture got modified values
 checkbuf(temp,hostSurf);

 // Clean up CUDA
 cudaGraphicsUnregisterResource(pResource);
 cudaDestroySurfaceObject(inputSurfObj);
 eglDestroySyncKHR(eglDisplayHandle, eglsync1);
 eglDestroySyncKHR(eglDisplayHandle, eglsync2);
 cudaEventDestroy(egl_event);
 cudaEventDestroy(cuda_event);
 .
 .
}

6. 适用于Jetson的CUDA可升级软件包

CUDA从JetPack SDK 5.0开始引入了升级路径，提供了将CUDA驱动程序和CUDA工具包更新至最新版本的选项。

6.1. 安装CUDA升级包

6.1.1. 前提条件

Jetson设备必须安装兼容的NVIDIA JetPack版本。更多信息请参考使用正确的升级包。

6.1.2. 从网络仓库或本地安装程序获取

CUDA下载页面提供了逐步指导，说明如何下载和使用本地安装程序或CUDA网络存储库来安装最新的工具包。CUDA升级包会与适用于Linux-aarch64-jetson设备的相应CUDA工具包一起下载并安装。

对于应用程序构建在主机上且仅需在目标设备上独立安装CUDA升级包的使用场景，可以在CUDA Repos中找到相应的Debian包。以11.8版本为例，可通过运行以下命令进行安装：

$ sudo apt-get install -y cuda-compat-11-8

注意

对于磁盘空间（二级存储）有限的设备，这是推荐的CUDA升级路径。

已安装的升级包位于版本化的工具包目录中。例如，对于11.8版本，它位于/usr/local/cuda-11.8/。

升级包包含以下文件：

• libcuda.so.* - CUDA驱动程序

• libnvidia-nvvm.so.* - 即时链接时间优化（仅限CUDA 11.8及更高版本）

• libnvidia-ptxjitcompiler.so.* - 用于PTX文件的JIT(即时)编译器

• nvidia-cuda-mps-control - CUDA MPS控制可执行程序

• nvidia-cuda-mps-server - CUDA MPS服务器可执行文件

这些文件共同实现了CUDA 11.8的驱动程序接口。

注意

本软件包仅提供文件，不包含系统配置。

示例

以下命令展示了如何安装CUDA升级包并用于运行应用程序。

$ sudo apt-get -y install cuda
Reading package lists...
Building dependency tree...
Reading state information...
The following additional packages will be installed:
  cuda-11-8 cuda-cccl-11-8 cuda-command-line-tools-11-8 cuda-compat-11-8
  ...<snip>...
The following NEW packages will be installed:
  cuda cuda-11-8 cuda-cccl-11-8 cuda-command-line-tools-11-8 cuda-compat-11-8
  ...<snip>...
0 upgraded, 48 newly installed, 0 to remove and 38 not upgraded.
Need to get 15.7 MB/1,294 MB of archives.
After this operation, 4,375 MB of additional disk space will be used.
Get:1 https://developer.download.nvidia.com/compute/cuda/repos/ubuntu2004/arm64  cuda-compat-11-8 11.8.31339915-1 [15.8 MB]
Fetched 15.7 MB in 12s (1,338 kB/s)
Selecting previously unselected package cuda-compat-11-8.
(Reading database ...
  ...<snip>...
(Reading database ... 100%
(Reading database ... 148682 files and directories currently installed.)
Preparing to unpack .../00-cuda-compat-11-8_11.8.30682616-1_arm64.deb ...
Unpacking cuda-compat-11-8 (11.8.30682616-1) ...
  ...<snip>...
Unpacking cuda-11-8 (11.8.0-1) ...
Selecting previously unselected package cuda.
Preparing to unpack .../47-cuda_11.8.0-1_arm64.deb ...
Unpacking cuda (11.8.0-1) ...
Setting up cuda-toolkit-config-common (11.8.56-1) ...
Setting up cuda-nvml-dev-11-8 (11.8.56-1) ...
Setting up cuda-compat-11-8 (11.8.30682616-1) ...
  ...<snip>...


$ ls -l /usr/local/cuda-11.8/compat
total 55300
lrwxrwxrwx 1 root root       12 Jan  6 19:14 libcuda.so -> libcuda.so.1
lrwxrwxrwx 1 root root       14 Jan  6 19:14 libcuda.so.1 -> libcuda.so.1.1
-rw-r--r-- 1 root root 21702832 Jan  6 19:14 libcuda.so.1.1
lrwxrwxrwx 1 root root       19 Jan  6 19:14 libnvidia-nvvm.so -> libnvidia-nvvm.so.4
lrwxrwxrwx 1 root root       23 Jan  6 19:14 libnvidia-nvvm.so.4 -> libnvidia-nvvm.so.4.0.0
-rw-r--r-- 1 root root 24255256 Jan  6 19:14 libnvidia-nvvm.so.4.0.0
-rw-r--r-- 1 root root 10665608 Jan  6 19:14 libnvidia-ptxjitcompiler.so
lrwxrwxrwx 1 root root       27 Jan  6 19:14 libnvidia-ptxjitcompiler.so.1 -> libnvidia-ptxjitcompiler.so

$ export PATH=/usr/local/cuda-11.8/bin:$PATH
$ export LD_LIBRARY_PATH=/usr/local/cuda-11.8/lib64:$LD_LIBRARY_PATH

用户可以在运行CUDA 11.8应用程序之前，设置LD_LIBRARY_PATH以包含升级包安装的库：

$ LD_LIBRARY_PATH=/usr/local/cuda-11.8/compat:$LD_LIBRARY_PATH ~/Samples/1_Utilities/deviceQuery

CUDA Device Query (Runtime API) version (CUDART static linking)

Detected 1 CUDA Capable device(s)

Device 0: "Orin"
  CUDA Driver Version / Runtime Version          11.8 / 11.8
  CUDA Capability Major/Minor version number:    8.7
      ...<snip>...
deviceQuery, CUDA Driver = CUDART, CUDA Driver Version = 11.8, CUDA Runtime Version = 11.8, NumDevs = 1
Result = PASS

在给定系统上，任何时候只能安装一个CUDA升级包。安装新的CUDA升级包时，之前安装的升级包版本将被移除并替换为新版本。安装程序会保留默认驱动程序（最初随NVIDIA JetPack安装，属于L4T BSP的一部分）。应用程序可以选择使用默认版本的CUDA（最初随NVIDIA JetPack安装）或升级包安装的版本。可通过LD_LIBRARY_PATH环境变量来选择所需版本。

除了LD_LIBRARY_PATH外，CUDA MPS用户还必须设置PATH环境变量，以便在启动MPS和运行使用MPS的CUDA应用程序之前，能够使用升级包安装的nvidia-cuda-mps-*可执行文件。升级包安装的MPS可执行文件仅与同一升级包安装的CUDA驱动程序兼容，反之亦然，这可以通过版本信息进行检查。

如果升级包与NVIDIA JetPack版本不兼容，安装将会失败。

6.2. CUDA升级包的部署注意事项

6.2.1. 使用正确的升级包

CUDA升级包的命名基于其支持的最高工具包版本。例如，如果您当前使用的是NVIDIA JetPack SDK 5.0（11.4版）驱动程序，但需要11.8版应用支持，则应安装适用于11.8版的CUDA升级包。

每个CUDA版本仅支持特定NVIDIA JetPack版本的升级。下表显示了各CUDA版本所支持的NVIDIA JetPack SDK版本。

JetPack SDK	CUDA 11.4	CUDA 11.8	CUDA 12.0	CUDA 12.1	CUDA 12.2	CUDA 12.3 及更高版本
5.x	默认	C	C	C	C	X

JetPack SDK	CUDA 12.2	CUDA 12.3	CUDA 12.4	CUDA 12.5	CUDA 12.6
6.x	默认	X	C	C	C

下表展示了NVIDIA JetPack 5.x版本中CUDA UMD与CUDA Toolkit的版本兼容性：

CUDA UMD					CUDA 工具包
11.4 (默认 - 包含在NVIDIA JetPack中)	11.8	12.0	12.1	12.2
11.4 (默认 - 包含在NVIDIA JetPack中)	C	C (Minor Version Compatibility)	X	X	X
11.8 (通过升级包)	C (二进制兼容性)	C	X	X	X
12.0 (通过升级包)	C (二进制兼容性)	C (二进制兼容性)	C	C (小版本兼容性)	C (小版本兼容性)
12.1 (通过升级包)	C (二进制兼容性)	C (二进制兼容性)	C (二进制兼容性)	C	C (次要版本兼容性)
12.2 (通过升级包)	C (二进制兼容性)	C (二进制兼容性)	C (二进制兼容性)	C (二进制兼容性)	C

下表展示了NVIDIA JetPack 6.x版本中CUDA UMD与CUDA Toolkit的版本兼容性：

CUDA UMD						CUDA 工具包
12.2 (默认 - 包含在NVIDIA JetPack中)	12.4	12.5	12.6	12.7	12.8
12.2 (默认 - 包含在NVIDIA JetPack中)	C	C (Minor Version Compatibility)	C (Minor Version Compatibility)	C (Minor Version Compatibility)	X	C (Minor Version Compatibility)
12.4 (通过升级包)	C (二进制兼容性)	C	C (小版本兼容性)	C (小版本兼容性)	X	C (小版本兼容性)
12.5 (通过升级包)	C (二进制兼容性)	C (二进制兼容性)	C	C (小版本兼容性)	X	C (小版本兼容性)
12.6 (通过升级包)	C (二进制兼容性)	C (二进制兼容性)	C (二进制兼容性)	C	X	C (小版本兼容性)
12.7	X	X	X	X	X	X
12.8 (通过升级包)	C (二进制兼容性)	C (二进制兼容性)	C (二进制兼容性)	C (二进制兼容性)	X	C

C - 兼容

X - 不兼容

注意

NVIDIA JetPack SDK 5.x上的CUDA升级包从CUDA 11.8版本开始提供。

6.2.2. 功能异常

CUDA升级包仅更新CUDA驱动程序接口，而保持NVIDIA JetPack SDK其余组件不变。如果最新CUDA驱动程序中的新功能需要更新的NVIDIA JetPack SDK组件/接口，则可能无法正常工作并在使用时报错。

6.2.3. 检查兼容性支持

除了CUDA驱动和某些编译器组件外，NVIDIA JetPack中的其他驱动程序仍保持默认版本。CUDA升级路径仅适用于CUDA。

编写良好的应用程序应使用以下错误代码来确定是否支持CUDA升级。系统管理员应了解这些错误代码，以判断部署过程中是否存在错误。

1. CUDA_ERROR_SYSTEM_DRIVER_MISMATCH = 803。该错误表明升级后的CUDA驱动版本与系统上已安装的驱动版本不匹配。

2. CUDA_ERROR_COMPAT_NOT_SUPPORTED_ON_DEVICE = 804. 该错误表明系统已更新为使用CUDA升级包运行，但CUDA检测到的可见硬件不支持此配置。

7. cuDLA

DLA（深度学习加速器）是NVIDIA Tegra SoC上搭载的固定功能加速器，专为推理应用设计。该硬件在性能/功耗比方面表现卓越，能够原生运行现代神经网络中的多数层，因此成为嵌入式AI应用的理想选择。DLA编程通常包含离线与在线两个阶段：离线阶段中，DLA编译器会解析输入网络并将其编译为可加载文件；在线阶段则通过DLA硬件执行该文件以生成推理结果。NVIDIA当前提供的在线执行软件栈包含NvMediaDla和DLA运行时/KMD，这些API共同支持用户向DLA硬件提交推理任务。主要功能路径如下图所示。

图4 DLA软件栈

根据上述模型可知，希望在同一应用中使用GPU和DLA的用户必须借助EGLStreams/NvSci等互操作机制来共享缓冲区，并在GPU与DLA之间实现同步原语。这些互操作机制通常需要为每个共享缓冲区执行多个步骤，且难以精细调节GPU与DLA之间的任务调度。cuDLA作为CUDA编程模型的扩展，将深度学习加速器(DLA)与CUDA深度集成，使得开发者能够使用CUDA流(streams)和图(graphs)等编程结构来提交DLA任务。cuDLA会透明处理共享缓冲区的管理以及GPU与DLA间的任务同步，让程序员能专注于高层级用例开发。

7.1. 开发者指南

本节介绍使用cuDLA API对DLA硬件进行编程的关键原则。cuDLA接口提供了初始化设备、管理内存和提交DLA任务的机制。因此，本节将讨论如何利用cuDLA API实现这些用例。这些API的详细规范请参阅API规范文档，在编写cuDLA应用程序时应作为参考。

由于cuDLA是CUDA的扩展，它被设计为与执行GPU管理、上下文管理等CUDA功能的API协同工作。因此，在评估cuDLA API行为时，应用程序当前状态（包括所选GPU和当前活动上下文及其生命周期）都是重要的考量因素。

7.1.1. 设备模型

要执行任何DLA操作，应用程序首先需要创建一个cuDLA设备句柄。cudlaCreateDevice() API会创建一个cuDLA设备的逻辑实例，其中选定的DLA硬件实例会通过CUDA与当前活动的GPU进行耦合。例如，以下代码片段将创建一个由当前GPU（通过cudaSetDevice()设置）和DLA硬件0组成的逻辑实例。目前，cuDLA仅支持Tegra上的集成GPU，若尝试将当前GPU设置为独立GPU来创建设备句柄，将会在cudlaCreateDevice()过程中导致设备创建错误。

cudlaDevHandle devHandle;
cudlaStatus ret;
ret = cudlaCreateDevice(0, &devHandle, CUDLA_CUDA_DLA);

图5 设备模型

用户可以使用cudlaCreateDevice()创建任意数量的此类逻辑实例，并采用GPU和DLA硬件实例的任意组合（取决于系统资源可用性）：

图6 设备模型 - 多实例

此外，cudlaCreateDevice()在创建设备时支持一个替代标志——CUDLA_STANDALONE。当应用程序希望在独立模式下创建cuDLA设备（即不与GPU设备耦合）时，可以使用此标志。所有设备提交任务也可以使用独立模式下的cuDLA完成，但此模式下不支持CUDA交互。因此，在后续描述特定API或特定用例时，会考虑两种执行模式：混合模式和独立模式。API规范中详细说明了每种模式下支持哪些API。

7.1.2. 加载与查询模块

在提交任何DLA任务之前，cuDLA设备句柄需要关联一个适当的可加载文件。该可加载文件通常使用TensorRT离线创建，其中包含输入输出张量的数量及其各自元数据信息，应用程序可查询获取这些信息。成功初始化cuDLA设备后，典型应用流程如下（穿插部分调试日志）：

DPRINTF("Device created successfully\n");

// Load the loadable from 'loadableData' in which the loadable binary has
// been copied from the location of the loadable - disk or otherwise.
err = cudlaModuleLoadFromMemory(devHandle, loadableData, file_size, &moduleHandle, 0);
if (err != cudlaSuccess)
{
    // handle error
}

// Get tensor attributes.
uint32_t numInputTensors = 0;
uint32_t numOutputTensors = 0;
cudlaModuleAttribute attribute;

err = cudlaModuleGetAttributes(moduleHandle, CUDLA_NUM_INPUT_TENSORS, &attribute);
if (err != cudlaSuccess)
{
    // handle error
}
numInputTensors = attribute.numInputTensors;
DPRINTF("numInputTensors = %d\n", numInputTensors);

err = cudlaModuleGetAttributes(moduleHandle, CUDLA_NUM_OUTPUT_TENSORS, &attribute);
if (err != cudlaSuccess)
{
    // handle error
}
numOutputTensors = attribute.numOutputTensors;
DPRINTF("numOutputTensors = %d\n", numOutputTensors);

cudlaModuleTensorDescriptor* inputTensorDesc =
        (cudlaModuleTensorDescriptor*)malloc(sizeof(cudlaModuleTensorDescriptor)
              *numInputTensors);
cudlaModuleTensorDescriptor* outputTensorDesc =
        (cudlaModuleTensorDescriptor*)malloc(sizeof(cudlaModuleTensorDescriptor)
              *numOutputTensors);

if ((inputTensorDesc == NULL) || (outputTensorDesc == NULL))
{
    // handle error
}

attribute.inputTensorDesc = inputTensorDesc;
err = cudlaModuleGetAttributes(moduleHandle,
                               CUDLA_INPUT_TENSOR_DESCRIPTORS,
                               &attribute);
if (err != cudlaSuccess)
{
    // handle error
}

attribute.outputTensorDesc = outputTensorDesc;
err = cudlaModuleGetAttributes(moduleHandle,
                               CUDLA_OUTPUT_TENSOR_DESCRIPTORS,
                               &attribute);
if (err != cudlaSuccess)
{
    // handle error
}

应用程序可以使用检索到的张量描述符来设置其数据缓冲区的大小和格式。有关张量描述符内容的详细信息，请参阅API规范部分中的cudlaModuleGetAttributes()。

7.1.3. 内存模型

GPU和DLA拥有各自不同的MMU（内存管理单元），在执行各自功能时负责管理虚拟地址(VA)到物理地址(PA)的转换。下图展示了一个示例场景：GMMU负责处理GPU虚拟地址的转换，而SMMU则为来自DLA的虚拟地址执行类似的功能。

图7 虚拟地址到物理地址的转换

在混合模式下，在DLA可以访问CUDA指针之前，必须先将该指针注册到DLA。这个注册步骤会在SMMU中创建一个条目，并返回相应的虚拟地址(VA)用于任务提交。以下代码片段展示了使用CUDLA_CUDA_DLA标志创建的设备句柄的注册示例：

// Allocate memory on GPU.
void* buffer;
uint32_t size = 100;

result = cudaMalloc(&inputBufferGPU, size);
if (result != cudaSuccess)
{
    // handle error
}

// Register the CUDA-allocated buffers.
uint64_t* bufferRegisteredPtr = NULL;

err = cudlaMemRegister(devHandle,
                       (uint64_t* )inputBufferGPU,
                       size,
                       &bufferRegisteredPtr,
                       0);
if (err != cudlaSuccess)
{
    // handle error
}

在独立模式下，cuDLA无需底层CUDA设备即可运行。因此在此模式下，应用程序执行的内存分配（后续需要注册）必须来自CUDA之外。在Tegra系统上，cuDLA支持通过cudlaImportExternalMemory() API注册NvSciBuf分配，如下代码片段所示：

// Allocate the NvSciBuf object.
NvSciBufObj inputBufObj;
sciError = NvSciBufObjAlloc(reconciledInputAttrList, &inputBufObj);
if (sciError != NvSciError_Success)
{
    // handle error
}

uint64_t* inputBufObjRegPtr = NULL;

// importing external memory
cudlaExternalMemoryHandleDesc memDesc = { 0 };
memset(&memDesc, 0, sizeof(memDesc));
memDesc.extBufObject = (void *)inputBufObj;
memDesc.size = size;
err = cudlaImportExternalMemory(devHandle, &memDesc, &inputBufObjRegPtr, 0);
if (err != cudlaSuccess)
{
   // handle error
}

7.1.4. 任务执行与同步模型

7.1.4.1. 任务执行

提交DLA任务执行类似于向GPU提交CUDA内核。cuDLA原生支持CUDA流，并与流语义无缝协作，确保只有当流上前序任务完成执行后，所有预定给DLA的任务才会由DLA硬件执行。这使得应用程序能够使用熟悉的流语义在GPU和DLA之间建立复杂处理工作流，而无需管理GPU与DLA之间的内存一致性和执行依赖关系。下图展示了该执行模型的示意图。在给定流或多个流中，DLA任务可以与GPU任务交替排列，cudlaSubmitTask()会处理所有内存/执行依赖关系。

图8 cuDLA任务执行模型

提交任务API需要以注册到DLA的地址形式提供输入和输出张量（使用cudlaMemRegister()）。应用程序可以预先将所有需要的指针注册到cuDLA，然后在cudlaSubmitTask()期间使用这些已注册的指针。该API会确保在DLA开始执行当前任务之前，与已注册指针对应的底层内存上先前操作的结果对DLA可见。下面代码片段展示了一个典型的包含CUDA和cuDLA操作的应用程序代码：

DPRINTF("ALL MEMORY REGISTERED SUCCESSFULLY\n");

// Copy data from CPU buffers to GPU buffers.
result = cudaMemcpyAsync(inputBufferGPU, inputBuffer, inputTensorDesc[0].size, cudaMemcpyHostToDevice, stream);
if (result != cudaSuccess)
{
    // handle error
}

result = cudaMemsetAsync(outputBufferGPU, 0, outputTensorDesc[0].size, stream);
if (result != cudaSuccess)
{
    // handle error
}

// Enqueue a cuDLA task.
cudlaTask task;
task.moduleHandle = moduleHandle;
task.outputTensor = &outputBufferRegisteredPtr;
task.numOutputTensors = 1;
task.numInputTensors = 1;
task.inputTensor = &inputBufferRegisteredPtr;
task.waitEvents = NULL;
task.signalEvents = NULL;
err = cudlaSubmitTask(devHandle, &task, 1, stream, 0);
if (err != cudlaSuccess)
{
    // handle error
}
DPRINTF("SUBMIT IS DONE !!!\n");

result = cudaMemcpyAsync(outputBuffer, outputBufferGPU, outputTensorDesc[0].size, cudaMemcpyDeviceToHost, stream);
if (result != cudaSuccess)
{
    // handle error
}

在独立模式下，cudlaSubmitTask()中的stream参数必须指定为NULL，因为cuDLA独立于CUDA运行。在这种情况下，提交给DLA的任务将按FIFO顺序执行。

7.1.4.1.1. 多线程用户提交

用户可以在向特定设备句柄提交时指定CUDLA_SUBMIT_SKIP_LOCK_ACQUIRE标志，前提是确保该设备句柄仅在此线程中使用，并且该设备句柄与可能在其他并行线程中使用的任何其他设备句柄之间不存在应用程序级别的共享数据。此标志有助于在提交路径中进行某些优化，从应用程序的角度来看可能会缩短提交时间。

7.1.4.2. 同步

混合模式下的任务同步不需要使用不同的API。由于DLA任务被提交到CUDA流中，只需等待流完成其工作即可确保在该流上提交的所有DLA任务都已完成。在这方面，DLA任务同步与CUDA中可用的任何同步机制（事件、流、设备）兼容，整个CUDA机制可供应用程序用于设置不同的流程和用例。

然而在独立模式下，由于cuDLA独立于CUDA运行，其同步机制有所不同。在此模式下，cudlaTask结构体提供了指定等待事件和信号事件的机制，这些事件将分别作为cudlaSubmitTask()调用的一部分由cuDLA进行等待和触发。每个提交的任务将在开始执行前等待其所有等待事件被触发，并提供一个信号事件（如果在cudlaSubmitTask()调用期间请求了该事件），应用程序（或其他实体）可通过等待该信号事件来确保提交的任务已完成执行。在cuDLA 1.0中，仅支持将NvSciSync栅栏作为等待事件的一部分。此外，只有NvSciSync对象能被注册并作为信号事件触发，且被触发事件对应的栅栏将作为cudlaSubmitTask()调用的返回值返回。

与所有内存操作类似，事件的基础存储支持（在本例中为NvSciSync对象）必须在使用cuDLA提交任务前完成注册。以下代码片段展示了一个示例流程：应用程序创建输入和输出的NvSciSync对象并注册它们，创建对应的围栏，在cudlaSubmitTask()调用中将相应围栏标记为等待/信号状态，随后触发输入围栏并等待输出围栏。

7.1.4.2.1. 注册外部信号量:

sciError = NvSciSyncObjAlloc(nvSciSyncReconciledListObj1, &syncObj1);
if (sciError != NvSciError_Success)
{
    // handle error
}

sciError = NvSciSyncObjAlloc(nvSciSyncReconciledListObj2, &syncObj2);
if (sciError != NvSciError_Success)
{
    // handle error
}

// importing external semaphore
uint64_t* nvSciSyncObjRegPtr1 = NULL;
uint64_t* nvSciSyncObjRegPtr2 = NULL;
cudlaExternalSemaphoreHandleDesc semaMemDesc = { 0 };
memset(&semaMemDesc, 0, sizeof(semaMemDesc));
semaMemDesc.extSyncObject = syncObj1;
err = cudlaImportExternalSemaphore(devHandle,
                                   &semaMemDesc,
                                   &nvSciSyncObjRegPtr1,
                                   0);
if (err != cudlaSuccess)
{
    // handle error
}

memset(&semaMemDesc, 0, sizeof(semaMemDesc));
semaMemDesc.extSyncObject = syncObj2;
err = cudlaImportExternalSemaphore(devHandle,
                                   &semaMemDesc,
                                   &nvSciSyncObjRegPtr2,
                                   0);
if (err != cudlaSuccess)
{
    // handle error
}

DPRINTF("ALL EXTERNAL SEMAPHORES REGISTERED SUCCESSFULLY\n");

7.1.4.2.2. cudlaSubmitTask()的事件设置

// Wait events
NvSciSyncFence preFence = NvSciSyncFenceInitializer;
sciError = NvSciSyncObjGenerateFence(syncObj1, &preFence);
if (sciError != NvSciError_Success)
{
    // handle error
}

cudlaWaitEvents* waitEvents;
waitEvents = (cudlaWaitEvents *)malloc(sizeof(cudlaWaitEvents));
if (waitEvents == NULL)
{
    // handle error
}

waitEvents->numEvents = 1;
CudlaFence* preFences = (CudlaFence *)malloc(waitEvents->numEvents *
                                             sizeof(CudlaFence));
if (preFences == NULL)
{
    // handle error
}

preFences[0].fence = &preFence;
preFences[0].type = CUDLA_NVSCISYNC_FENCE;
waitEvents->preFences = preFences;

// Signal Events
cudlaSignalEvents* signalEvents;
signalEvents = (cudlaSignalEvents *)malloc(sizeof(cudlaSignalEvents));
if (signalEvents == NULL)
{
    // handle error
}

signalEvents->numEvents = 1;
uint64_t** devPtrs = (uint64_t **)malloc(signalEvents->numEvents *
                                         sizeof(uint64_t *));
if (devPtrs == NULL)
{
    // handle error
}

devPtrs[0] = nvSciSyncObjRegPtr2;
signalEvents->devPtrs = devPtrs;

signalEvents->eofFences = (CudlaFence *)malloc(signalEvents->numEvents *
                                               sizeof(CudlaFence));
if (signalEvents->eofFences == NULL)
{
    // handle error
}

NvSciSyncFence eofFence = NvSciSyncFenceInitializer;
signalEvents->eofFences[0].fence = &eofFence;
signalEvents->eofFences[0].type = CUDLA_NVSCISYNC_FENCE;

// Enqueue a cuDLA task.
cudlaTask task;
task.moduleHandle = moduleHandle;
task.outputTensor = &outputBufObjRegPtr;
task.numOutputTensors = 1;
task.numInputTensors = 1;
task.inputTensor = &inputBufObjRegPtr;
task.waitEvents = waitEvents;
task.signalEvents = signalEvents;
err = cudlaSubmitTask(devHandle, &task, 1, NULL, 0);
if (err != cudlaSuccess)
{
    // handle error
}
DPRINTF("SUBMIT IS DONE !!!\n");

7.1.4.2.3. 等待信号事件

// Signal wait events.
// For illustration purposes only. In practice, this signal will be done by another
// entity or driver that provides the data input for this particular submitted task.
NvSciSyncObjSignal(syncObj1);

// Wait for operations to finish.
// For illustration purposes only. In practice, this wait will be done by
// another entity or driver that is waiting for the output of the submitted task.
sciError = NvSciSyncFenceWait(reinterpret_cast<NvSciSyncFence*>(signalEvents->eofFences[0].fence),
                              nvSciCtx, -1);
if (sciError != NvSciError_Success)
{
    // handle error
}

7.1.4.2.4. cuDLA中支持的同步原语

cuDLA支持两种类型的NvSciSync对象原语，分别是同步点和确定性信号量。默认情况下，cuDLA会优先选择同步点原语而非确定性信号量原语，并在应用程序通过cudlaGetNvSciSyncAttributes()请求时，将这些优先级设置在NvSciSync属性列表中。

对于确定性信号量，用于创建NvSciSync对象的NvSciSync属性列表必须将NvSciSyncAttrKey_RequireDeterministicFences键的值设置为true。确定性栅栏允许用户在对应该信号入队之前就对信号量对象进行等待操作入队。对于此类信号量对象，cuDLA保证每个信号操作将使栅栏值递增'1'。用户需要自行跟踪在信号量对象上入队的信号，并相应地插入等待操作。

7.1.4.2.5. 在NvSciSyncAttrList中设置NvSciSyncAttrKey_RequireDeterministicFences键

// Set NvSciSyncAttrKey_RequireDeterministicFences key to true in
// NvScisyncAtrrList that is used to create NvSciSync object with
// Deterministic Semaphore primitive.
NvSciSyncAccessPerm cpuPerm = NvSciSyncAccessPerm_SignalOnly;
    keyValue[0].attrKey = NvSciSyncAttrKey_RequiredPerm;
    keyValue[0].value = (void*) &cpuPerm;
    keyValue[0].len = sizeof(cpuPerm);

bool detFenceReq = true;
    keyValue[1].attrKey = NvSciSyncAttrKey_RequireDeterministicFences;
    keyValue[1].value   = (const void*)&detFenceReq;
    keyValue[1].len     = sizeof(detFenceReq);

    return NvSciSyncAttrListSetAttrs(list, keyValue, 2);

7.1.4.2.6. NvSciFence的时间戳支持

cuDLA 在独立模式下支持 NvSci 的时间戳功能。

时间戳支持功能允许用户获取特定围栏(fence)被触发的时间。该时间值是以微秒为单位的DLA时钟快照。

cuDLA用户可以通过在填充NvSci等待者属性列表时，将NvSciSyncAttrKey_WaiterRequireTimestamps键的值设置为true来请求时间戳支持。

用户可以利用该时间戳，配合SOF（帧起始）围栏和EOF（帧结束）围栏，分别在任务开始前和任务完成后获取DLA时钟的快照。这使得用户能够计算DLA执行所提交任务所花费的时间。

7.1.4.2.7. 为NvSciSync对象请求时间戳支持

sciError fillCpuWaiterAttrList(NvSciSyncAttrList list)
{
    bool cpuWaiter = true;
    NvSciSyncAttrKeyValuePair keyValue[3];
    memset(keyValue, 0, sizeof(keyValue));
    keyValue[0].attrKey = NvSciSyncAttrKey_NeedCpuAccess;
    keyValue[0].value = (void*) &cpuWaiter;
    keyValue[0].len = sizeof(cpuWaiter);

    NvSciSyncAccessPerm cpuPerm = NvSciSyncAccessPerm_WaitOnly;
    keyValue[1].attrKey = NvSciSyncAttrKey_RequiredPerm;
    keyValue[1].value = (void*) &cpuPerm;
    keyValue[1].len = sizeof(cpuPerm);

    bool cpuRequiresTimeStamp = true;
    keyValue[2].attrKey = NvSciSyncAttrKey_WaiterRequireTimestamps;
    keyValue[2].value = (void*) &cpuRequiresTimeStamp;
    keyValue[2].len = sizeof(cpuRequiresTimeStamp);

    return NvSciSyncAttrListSetAttrs(list, keyValue, 3);
 }

NvSciSyncCpuWaitContext nvSciCtx;
NvSciSyncModule syncModule;
NvSciSyncAttrList waiterAttrListObj = nullptr;
NvSciSyncAttrList signalerAttrListObj = nullptr;
NvSciSyncAttrList syncAttrListObj[2];
NvSciSyncAttrList nvSciSyncConflictListObj;
NvSciSyncAttrList nvSciSyncReconciledListObj;

sciError = NvSciSyncModuleOpen(&syncModule);
if (sciError != NvSciError_Success) {
    //handle error
}

sciError = NvSciSyncAttrListCreate(syncModule, &signalerAttrListObj);
if (sciError != NvSciError_Success) {
    //handle error
}

sciError = NvSciSyncAttrListCreate(syncModule, &waiterAttrListObj);
if (sciError != NvSciError_Success) {
 //handle error
}

err = cudlaGetNvSciSyncAttributes(reinterpret_cast<uint64_t*>(signalerAttrListObj),
                                  CUDLA_NVSCISYNC_ATTR_SIGNAL);
if (err != cudlaSuccess) {
    //handle error
}

sciError = fillCpuWaiterAttrList(waiterAttrListObj);
if (sciError != NvSciError_Success) {
    //handle error
}

syncAttrListObj[0] = signalerAttrListObj;
syncAttrListObj[1] = waiterAttrListObj;
sciError = NvSciSyncAttrListReconcile(syncAttrListObj,
                                      2,
                                      &nvSciSyncReconciledListObj,
                                      &nvSciSyncConflictListObj3);
if (sciError != NvSciError_Success) {
    //handle error
}

sciError = NvSciSyncObjAlloc(nvSciSyncReconciledListObj, &syncObj);
if (sciError != NvSciError_Success) {
    //handle error
}

sciError = NvSciSyncCpuWaitContextAlloc(syncModule, &nvSciCtx);
if (sciError != NvSciError_Success) {
    //handle error
}

7.1.4.2.8. 从Fence中提取时间戳值

更多信息请参考以下部分：

• 注册外部信号量：

• 为cudaSubmitTask()设置事件

• 等待信号事件

// To extract Timestamp of the fence
// Timestamp will be valid only after fence is signaled
// hence Fence must be waited up on before extracting timestamp value

uint64_t eofTimestampUS = 0UL;
sciError =  NvSciSyncFenceGetTimestamp(reinterpret_cast<NvSciSyncFence*>(signalEvents->eofFences.fence), &(eofTimestampUS));
if ((sciError != NvSciError_Success) || (eofTimestampUS == 0UL)) {
    //handle error
}

7.1.4.3. 故障诊断

要对DLA硬件进行故障诊断，用户需要指定CUDLA_MODULE_ENABLE_FAULT_DIAGNOSTICS标志来加载模块，并在任务提交时设置CUDLA_SUBMIT_DIAGNOSTICS_TASK。该任务可用于探测DLA硬件的状态。启用此标志后，在独立模式下不允许用户仅提交事件（即张量信息为NULL且任务中仅包含等待/信号或两者的事件），因为该任务始终运行在内部加载的诊断模块上。此诊断模块不需要任何输入张量，因此也不要求输入张量内存。但用户需要查询输出张量的数量，分配输出张量内存，并在提交任务时传递这些信息。

7.1.4.4. 空操作提交

用户可以在调用cudlaSubmitTask()时将某些任务标记为无操作任务。

这是通过在cudlaSubmitTask()的flags参数中传递CUDLA_SUBMIT_NOOP来实现的。空操作提交意味着所有其他提交语义都被保留。具体来说，任务会被提交到DLA，前后会考虑等待/信号事件，并遵循流语义。关键区别在于DLA会跳过该任务的执行。此功能在混合模式和独立模式下均受支持。

7.1.5. 错误报告模型

任务执行的异步特性会导致两种错误可能通过cuDLA API报告：

• 同步错误

• 异步错误

同步错误是指cuDLA API在应用程序中调用时通过其返回代码报告的错误。异步错误则是指相对于程序顺序执行而言较晚检测到的错误。典型场景是每个提交给DLA硬件的任务会在特定时间后执行，因此如果任务执行中出现错误，无法通过任务提交API进行报告。根据错误发生时机，这些错误会在后续cuDLA API调用期间或同步操作后被报告。作为cuDLA API一部分报告的硬件执行错误在应用层处理较为直接。然而，如果当前应用程序中没有正在执行或即将执行的cuDLA API调用，则应用程序需要采取额外步骤来处理异步错误。

在混合模式下，DLA硬件错误可以通过CUDA同步操作报告。如设备模型部分所述，cuDLA在逻辑上将DLA与GPU关联以执行任务。因此，任何DLA硬件错误都会通过CUDA传播给用户。用户需要从CUDA同步操作中检查特定于DLA的错误，然后使用cudlaGetLastError()检查cuDLA设备句柄以获取确切错误。如果应用程序中有多个cuDLA设备句柄，并且每个句柄都以混合模式向cuDLA提交了一些任务，那么必须检查每个设备句柄是否存在错误。这里的基本模型是使用CUDA检测DLA硬件错误，然后在相关句柄上使用cudlaGetLastError()来报告确切错误。以下代码片段展示了一个示例：

result = cudaStreamSynchronize(stream);
if (result != cudaSuccess)
{
    DPRINTF("Error in synchronizing stream = %s\n", cudaGetErrorName(result));

    if (result == cudaErrorExternalDevice)
    {
        cudlaStatus hwStatus = cudlaGetLastError(devHandle);
        if (hwStatus != cudlaSuccess)
        {
            DPRINTF("Asynchronous error in HW = %u\n", hwStatus);
        }
    }
}

该错误报告模型也与CUDA Driver API兼容，因此如果应用程序使用CUDA Driver API进行同步，类似的错误代码和错误处理流程同样适用。

在独立模式下，该模型类似，区别在于没有相应的机制来检测同步操作中的错误。在此模式下，应用程序等待已提交任务的唯一选择是等待最新提交返回的NvSciSync围栏。截至本文撰写时，NvSciSync不支持报告DLA硬件错误，因此应用程序需要等待围栏，然后查询cudlaGetLastError()以获取执行期间的任何错误。

7.2. 从NvMediaDla迁移到cuDLA

NvMediaDla和cuDLA具有不同的编程模型，在各自API提供的功能上存在一定程度的重叠。下表提供了从NvMediaDla API到等效cuDLA API或功能的映射关系。本表旨在作为将NvMediaDla应用迁移至cuDLA应用时的参考指南。

NvMediaDla	cuDLA
NvMediaDlaGetVersion()	cudlaGetVersion()
NvMediaDlaPingById()	不需要，因为ping操作在cudlaCreateDevice内部完成，只有ping成功时设备句柄创建才会成功。
NvMediaDlaCreate()	cudlaCreateDevice()
NvMediaDlaDestroy()	cudlaDestroyDevice()
NvMediaDlaGetUMDVersion()	不可用
NvMediaDlaGetNumEngines()	cudlaDeviceGetCount()
NvMediaDlaGetMaxOutstandingTasks()	不可用
NvMediaDlaInit()	cudlaCreateDevice（但无法指定输入任务的数量）
NvMediaDlaGetInstanceId()	不可用
NvMediaDlaGetNumTasks()	不可用
NvMediaDlaLoadableCreate()	不需要，因为声明一个cudlaModule类型的变量并配合cudlaModuleLoadFromMemory()就足够了。
NvMediaDlaLoadableDestroy()	不需要，因为cuDLA模块被声明为cudlaModule类型的变量。
NvMediaDlaAppendLoadable()	不需要，因为这已在cudlaModuleLoadFromMemory()内部完成。
NvMediaDlaSetCurrentLoadable()	不需要，因为这已在cudlaModuleLoadFromMemory()内部完成。
NvMediaDlaGetNumOfInputTensors()	cudlaModuleGetAttributes()
NvMediaDlaGetInputTensorDescriptor()	cudlaModuleGetAttributes()
NvMediaDlaGetNumOfOutputTensors()	cudlaModuleGetAttributes()
NvMediaDlaGetOutputTensorDescriptor()	cudlaModuleGetAttributes()
NvMediaDlaDataRegister()	cudlaMemRegister()
NvMediaDlaDataUnregister()	cudlaMemUnregister()
NvMediaDlaLoadLoadable()	cudlaModuleLoadFromMemory()
NvMediaDlaRemoveLoadable()	cudlaModuleUnload()
NvMediaDlaSubmit()	cudlaSubmitTask()
NvMediaDlaNvSciSyncGetVersion()	不可用
NvMediaDlaFillNvSciSyncAttrList()	cudlaGetNvSciSyncAttributes()
NvMediaDlaRegisterNvSciSyncObj()	cudlaImportExternalSemaphore()
NvMediaDlaUnregisterNvSciSyncObj()	cudlaMemUnregister()
NvMediaDlaSetNvSciSyncObjforEOF()	由于cudlaTask结构已具备指定此功能的能力，因此无需额外操作。
NvMediaDlaInsertPreNvSciSyncFence()	不需要，因为cudlaTask结构已具备指定此功能的能力。
NvMediaDlaGetEOFNvSciSyncFence()	不需要，因为cudlaTask结构已具备检索此信息的能力。

7.3. 分析cuDLA应用

可以使用NVIDIA Nsight Systems对cuDLA API进行性能分析。以下命令可用于生成cuDLA API的跟踪记录，这些跟踪记录可在Nsight中查看。

$ nsys profile --trace nvtx -e CUDLA_NVTX_LEVEL=1 --output <file> <cudla_App>

7.4. cuDLA 版本说明

cuDLA 1.2.1中的已知问题：

• 在混合模式下，cuDLA内部会使用CUDA的主上下文分配内存。因此，在销毁/重置CUDA主上下文之前，必须销毁所有cuDLA设备的初始化。

• 在销毁一个cuDLA设备句柄之前，必须确保之前提交到该设备的所有任务都已完成。否则可能导致应用程序崩溃，因为内部内存分配仍在使用中。

• 应用程序分配的NvSciBuf缓冲区必须遵守DLA对齐约束。

• 应用程序需确保在提交任务时，等待事件中指定的围栏没有重复。

• 通常，cuDLA API返回的任何同步或异步错误都应被视为不可恢复的错误。在这种情况下，应用程序需要重新启动并再次初始化cuDLA以提交DLA任务。此规则的例外是cudlaErrorMemoryRegistered，当应用程序尝试重复注册特定内存而未先取消注册时，cuDLA会返回此错误。

• cuDLA不支持CUDA和DLA之间的UVM。

• cuDLA不支持CUDA Graph。

• cuDLA不支持每线程默认流。

• cuDLA不支持CNP（DLA函数不能与CNP一起使用）。

• cuDLA不支持块线性内存。

• cuDLA目前不支持CUDA VMM API。

• cuDLA不支持dGPU。

• 在某些情况下，DLA FW可能会在执行特定任务时卡住。这可能导致应用程序在混合模式及独立模式下均出现挂起现象。应用程序应能检测到此类情况并作出相应处理。

• 支持加载多个模块。

• 当加载多个模块时，不支持分层统计功能。

• 当加载单个模块且该模块同时用于任务提交和统计信息转换时，支持分层统计。

8. 公告

8.1. 注意事项

本文档仅供信息参考之用，不应视为对产品功能、状态或质量的保证。NVIDIA公司（“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对客户就本文所述产品的全部及累计责任应受产品销售条款的限制。

8.2. OpenCL

OpenCL是苹果公司的商标，经Khronos Group Inc.授权使用。

8.3. 商标

NVIDIA和NVIDIA标识是美国及其他国家NVIDIA公司的商标或注册商标。其他公司及产品名称可能是其各自关联公司的商标。