NIVIDIA DRIVE Orin 系列作為一個萬用 SOC 芯片，可以用于各種不同的感知和通用計算任務(wù)，其優(yōu)質(zhì)的大算力、運(yùn)行性能、完備的兼容性，以及豐富的 I/O 接口，可以減少系統(tǒng)開發(fā)的復(fù)雜度。這些特性使得 Orin 系列的芯片特別適合應(yīng)用在自動駕駛系統(tǒng)。

整體上看，Orin系列芯片頂層SOC架構(gòu)的模塊主要由三部分處理單元組成：即 CPU、GPU 和硬件加速器組成。以當(dāng)前較火的Orin-x作為典型說明英偉達(dá)芯片在其軟件模塊開發(fā)中是如何進(jìn)行調(diào)用的。

1、CPU：

Orin-x中CPU包括 12 個 Cortex-A78，可以提供通用的目標(biāo)高速計算兼容性。同時，Arm Cortex R52 基于功能安全設(shè)計（FSI），可以提供獨立的片上計算資源，這樣就可以不用增加額外的 CPU（ASIL D）芯片用來提供功能安全等級。

CPU 族群所支持的特性包括 Debug 調(diào)試，電源管理，Arm CoreLink 中斷控制器，錯誤檢測與報告。

CPU 需要對芯片進(jìn)行整體性能監(jiān)控，每個核中的性能監(jiān)控單元提供了六個計算單元，每 個單元可以計算處理器中的任何事件。基于 PMUv3 架構(gòu)上，在每個 Runtime 期間這些計算 單元會收集不同的統(tǒng)計值并運(yùn)行在處理器和存儲系統(tǒng)上。

2、GPU：

NVIDIA Ampere GPU 可以提供先進(jìn)的并行處理計算架構(gòu)。開發(fā)者可以使用 CUDA 語言進(jìn)行開發(fā)（后續(xù)將對CUDA架構(gòu)進(jìn)行詳細(xì)說明），并支持 NVIDIA 中各種不同的工具鏈（如開發(fā) Tensor Core 和 RT Core 的應(yīng)用程序接口）。一個深度學(xué)習(xí)接口優(yōu)化器和實時運(yùn)行系統(tǒng)可以傳遞低延遲和高效輸出。Ampere GPU 同時可以提供如下一些的特性來實現(xiàn)對高分辨率、高復(fù)雜度的圖像處理能力（如實時光流追蹤）。

稀疏化：

細(xì)粒度結(jié)構(gòu)化稀疏性使吞吐量翻倍，減少對內(nèi)存消耗。浮點處理能力：每個時鐘周期內(nèi)可實現(xiàn) 2 倍 CUDA 浮點性能。

緩存：

流處理器架構(gòu)可以增加 L1 高速緩存帶寬和共享內(nèi)存，減少緩存未命中延遲。提升異步計算能力，后 L2 緩存壓縮。

3、Domain-Specific：

特定域硬件加速器（DSAs、DLA、PVA）是一組特殊目的硬件引擎，實現(xiàn)計算引擎多任務(wù)、高效、低功率等特性。計算機(jī)視覺和深度學(xué)習(xí)簇包括兩個主要的引擎：可編程視覺加速器 PVA 和深度學(xué)習(xí)加速器 DLA（而在最新的中級算力 Orin n 芯片則取消了 DLA 處理器）。

PVA 是第二代 NVIDIA 視覺DSP架構(gòu)，它是一種特殊應(yīng)用指令矢量處理器，這種處理器是專門針對計算機(jī)視覺、ADAS、ADS、虛擬現(xiàn)實系統(tǒng)。PVA 有一些關(guān)鍵的要素可以很好的適配預(yù)測算法領(lǐng)域，且功耗和延遲性都很低。Orin-x需要通過內(nèi)部的R核（Cortex-R5）子系統(tǒng)可以用于 PVA 控制和任務(wù)監(jiān)控。一個 PVA 簇可以完成如下任務(wù)：雙向量處理單元（VPU）帶有向量核，指令緩存和 3 矢量數(shù)據(jù)存儲單元。每個單元有 7 個可見的插槽，包含可標(biāo)量和向量指令。此外，每個 VPU 還含有 384 KBytes的3端口存儲容量。

DLA 是一個固定的函數(shù)引擎，可用于加速卷積神經(jīng)網(wǎng)絡(luò)中的推理操作。Orin-x 單獨設(shè)置了 DLA 用于實現(xiàn)第二代 NVIDIA 的 DLA架構(gòu)。DLA支持加速 CNN 層的卷積、去卷積、激活、池化、局部歸一化、全連接層。最終支持優(yōu)化結(jié)構(gòu)化稀疏、深度卷積、一個專用的硬件調(diào)度器，以最大限度地提高效率。

那么，怎樣利用英偉達(dá)自身的GPU和CPU上的算力進(jìn)行有效的計算機(jī)視覺開發(fā)能力探測呢？

GPU的軟件架構(gòu)

自動駕駛領(lǐng)域使用的 AI 算法多為并行結(jié)構(gòu)。AI 領(lǐng)域中用于圖像識別的深度學(xué)習(xí)、用于決策和推理的機(jī)器學(xué)習(xí)以及超級計算都需要大規(guī)模的并行計算，更適合采用 GPU 架構(gòu)。由于神經(jīng)網(wǎng)絡(luò)的分層級數(shù)（通常隱藏層的數(shù)量越多，神經(jīng)網(wǎng)絡(luò)模擬的結(jié)果越精確）會很大程度的影響其在預(yù)測結(jié)果。擅長并行處理的 GPU 可以很好的對神經(jīng)網(wǎng)絡(luò)算法進(jìn)行處理和優(yōu)化。因為，神經(jīng)網(wǎng)絡(luò)中的每個計算都是獨立于其他計算的，這意味著任何計算都不依賴于任何其他計算的結(jié)果，所有這些獨立的計算都可以在 GPU 上并行進(jìn)行。通常 GPU 上進(jìn)行的單個卷積計算要比 CPU 慢，但是對于整個任務(wù)來說，CPU 幾乎是串行處理方式，需要要逐個依次完成，因此，其速度要大大慢于 GPU。因此，卷積運(yùn)算可以通過使用并行編程方法和GPU來加速。

英偉達(dá)通過 CPU+GPU+DPU 形成產(chǎn)品矩陣，全面發(fā)力數(shù)據(jù)中心市場。利用 GPU 在AI 領(lǐng)域的先天優(yōu)勢，英偉達(dá)借此切入數(shù)據(jù)中心市場。針對芯片內(nèi)部帶寬以及系統(tǒng)級互聯(lián)等 諸多問題，英偉達(dá)推出了 Bluefield DPU 和 Grace CPU，提升了整體硬件性能。

對于英偉達(dá)的GPU而言，一個 GPC 中有一個光柵引擎（ROP）和 4 個紋理處理集群（TPC），每個引擎可以訪問所有的存儲。

每個 TPC 包含兩路流式多媒體處理器（SM），每個 SM 有 128 個 CUDA cores，被分為 4 個單獨的處理模塊，每個模塊都有自己的指令緩沖區(qū)、調(diào)度程序和張量核心。每個TPC 也包含一個變形引擎、兩個紋理單元和兩個光線追蹤核心（RT Core）。

GPC是一種專門針對硬件模塊開發(fā)的光柵、陰影、紋理結(jié)構(gòu)和計算的單元。GPU的核心圖像處理函數(shù)就是在 GPC 中實現(xiàn)的。在 GPC 中，SM 的 CUDA 核可以執(zhí)行像素級/矢量級/幾何陰影的計算。紋理結(jié)構(gòu)單元執(zhí)行紋理濾波和加載/存儲獲取、數(shù)據(jù)保存到存儲器。

特殊的函數(shù)單元（SFUs）可以處理先驗的和圖像的內(nèi)插指令。

Tensor Cores 執(zhí)行矩陣乘法來極大加速深度學(xué)習(xí)推理。RTcore 單元通過加速邊界體積層次結(jié)構(gòu)（BVH）的遍歷和光流追蹤過程中的場景幾何交叉可以輔助光流追蹤性能。

最后，多元引擎處理可以處理頂點提取、鑲嵌、視角轉(zhuǎn)換、屬性建立和視頻流輸出。SM幾何級、像素級處理執(zhí)行性能可以確保其高適配性。這就可以很好地用于用戶接口和復(fù)雜的應(yīng)用開發(fā)，Ampere GPU 的功率也得到了優(yōu)化，可以確保其在環(huán)境提供低功耗下還能保持較高性能。

核心通用編程架構(gòu)CUDA

CUDA（Compute Unified Device Architecture，統(tǒng)一計算架構(gòu)） 作為連接 AI 的中心節(jié)點，CUDA+GPU 系統(tǒng)極大推動了 AI 領(lǐng)域的發(fā)展。搭載英偉達(dá) GPU 硬件的工作站（Workstation）、服務(wù)器（Server）和云（Cloud）通過 CUDA軟件系統(tǒng)以及開發(fā)的 CUDA-XAI 庫，為自動駕駛系統(tǒng) AI 計算所需要的機(jī)器學(xué)習(xí)、深度學(xué)習(xí)的訓(xùn)練（Train）和推理（Inference）提供了對應(yīng)的軟件工具鏈，來服務(wù)眾多的框架、云服務(wù)等等，是整個英偉達(dá)系列芯片軟件開發(fā)中必不可少的一環(huán)。

CUDA 是一個基于英偉達(dá) GPU 平臺上面定制的特殊計算體系/算法，一般只能在英偉達(dá)的 GPU 系統(tǒng)上使用。這里從開發(fā)者角度我們講講在英偉達(dá) Orin 系列芯片中如何在 CUDA架構(gòu)上進(jìn)行不同軟件級別開發(fā)。

1、CUDA 架構(gòu)及數(shù)據(jù)處理解析

如下圖表示了 CUDA 架構(gòu)示意圖，表示了CPU，GPU，應(yīng)用程序，CUDA 開發(fā)庫，運(yùn)行環(huán)境，驅(qū)動之間的關(guān)系如上圖所示。

那么，如何在涵蓋 CPU 和 GPU 模塊中利用 CUDA 本身的編程語言進(jìn)行開發(fā)調(diào)用呢？

從CUDA 體系結(jié)構(gòu)的組成來說，它包含了三個部分：開發(fā)庫、運(yùn)行期環(huán)境和驅(qū)動。

“Developer Lib 開發(fā)庫”是基于 CUDA 技術(shù)所提供的應(yīng)用開發(fā)庫。例如高度優(yōu)化的通用數(shù)學(xué)庫，即cuBLAS、cuSolver 和 cuFFT。核心庫，例如 Thrust 和 libcu++；通信庫， 例如 NCCL 和 NVSHMEM，以及其他可以在其上構(gòu)建應(yīng)用程序的包和框架。

“Runtime 運(yùn)行期環(huán)境”提供了應(yīng)用開發(fā)接口和運(yùn)行期組件，包括基本數(shù)據(jù)類型的定義和各類計算、類型轉(zhuǎn)換、內(nèi)存管理、設(shè)備訪問和執(zhí)行調(diào)度等函數(shù)。

“Driver 驅(qū)動部分”是 CUDA 使能 GPU 的設(shè)備抽象層，提供硬件設(shè)備的抽象訪問接口。CUDA 提供運(yùn)行期環(huán)境也是通過這一層來實現(xiàn)各種功能的。

目前于 CUDA 開發(fā)的應(yīng)用必須有 NVIDIA CUDA-enable 的硬件支持。對于 CUDA 執(zhí)行過程而言，CPU 擔(dān)任的工作是為控制 GPU 執(zhí)行，調(diào)度分配任務(wù)，并能做一些簡單的計算，而大量需要并行計算的工作都交給 GPU 實現(xiàn)。

在 CUDA 架構(gòu)下，一個程序分為兩個部份：host 端和 device 端。Host 端是指在 CPU 上執(zhí)行的部份，而 device 端則是在顯示芯片（GPU）上執(zhí)行的部份。Device 端的程序又稱為 "kernel"。通常 host 端程序會將數(shù)據(jù)準(zhǔn)備好后，復(fù)制到顯卡的內(nèi)存中，再由顯示芯片執(zhí)行 device 端程序，完成后再由 host 端程序?qū)⒔Y(jié)果從顯卡的內(nèi)存中取回。這里需要注意的是，由于 CPU 存取顯存時只能透過 PCI Express 接口，因此速度較慢 （PCI Express x16 的理論帶寬是雙向各 4GB/s），因此不能經(jīng)常進(jìn)行，以免降低效率。

基于以上分析可知，針對大量并行化問題，采用 CUDA 來進(jìn)行問題處理，可以有效隱藏內(nèi)存的延遲性 latency，且可以有效利用顯示芯片上的大量執(zhí)行單元，同時處理上千個線程 thread 。因此，如果不能處理大量并行化的問題，使用 CUDA 就沒辦法達(dá)到最好的效率了。

對于這一應(yīng)用瓶頸來說，英偉達(dá)也在數(shù)據(jù)存取上做出了較大的努力提升。一方面，優(yōu)化的CUDA 改進(jìn)了 DRAM 的讀寫靈活性，使得GPU與CPU的機(jī)制相吻合。另一方面，CUDA提供了片上（on-chip）共享內(nèi)存，使得線程之間可以共享數(shù)據(jù)。應(yīng)用程序可以利用共享內(nèi)存來減少 DRAM 的數(shù)據(jù)傳送，更少的依賴 DRAM 的內(nèi)存帶寬。

此外，CUDA 還可以在程序開始時將數(shù)據(jù)復(fù)制進(jìn) GPU 顯存，然后在 GPU 內(nèi)進(jìn)行計算，直到獲得需要的數(shù)據(jù)，再將其復(fù)制到系統(tǒng)內(nèi)存中。為了讓研發(fā)人員方便使用 GPU 的算力，英偉達(dá)不斷優(yōu)化 CUDA 的開發(fā)庫及驅(qū)動系統(tǒng)。操作系統(tǒng)的多任務(wù)機(jī)制可以同時管理 CUDA 訪問 GPU 和圖形程序的運(yùn)行庫，其計算特性支持利用 CUDA 直觀地編寫 GPU 核心程序。

2、CUDA 編程開發(fā)

英偉達(dá)系列芯片的應(yīng)用基礎(chǔ)是提供一套豐富且成熟的 SDK 和庫，可以基于它們構(gòu)建應(yīng)用程序。CUDA 這種計算架構(gòu)通過優(yōu)化的各類調(diào)度算法和軟件框架，實際為開發(fā)者提供了快速可調(diào)用的底層編程代碼，可以確保開發(fā)者在編程過程中能夠最快速最有效的直接調(diào)用GPU 并行計算資源，從而最大化提升 GPU 的計算效率，助力英偉達(dá) GPU 方便且高效地發(fā)揮其并行計算能力。

CUDA 是一種類 C 語言，本身也兼容 C 語言，所以適合普通開發(fā)者在其上編寫和移植自己的代碼和算法?？傮w來講，英偉達(dá)在編程方式上分成了如下三種不同的模式。如下圖表示了各種編程語言之間的關(guān)系：

? 針對文本主題通常采用標(biāo)準(zhǔn)語言并行開發(fā)；針對性能優(yōu)化的平臺專用語言，如 C++、Fortran、OpenCL 等；其中，OpenCL 也可以實現(xiàn)對 GPU 計算能力的調(diào)用，但由于其通用性較強(qiáng)，整體優(yōu)化效果不如 CUDA， 在大規(guī)模計算中劣勢很大。CUDA 架構(gòu)是 OpenCL 的運(yùn)行平臺之一，OpenCL僅僅是為 CUDA 架構(gòu)提供了一個可編程的 API 而已。也就是相當(dāng)于 API 與執(zhí)行架構(gòu)之間的關(guān)系。

針對如上標(biāo)準(zhǔn)語言和專用語言所定制的編譯器指令，可以通過啟用增量性能優(yōu)化來彌合這兩種方式之間的差距。從而在性能、生產(chǎn)率和代碼可移植性方面進(jìn)行權(quán)衡。

如下圖表示了一種典型的 CUDA 架構(gòu)優(yōu)化后相對于 CPU、GPU 在軟件編程上的性能提升。

為什么CUDA能夠起到比CPU和GPU更好的運(yùn)算效果呢？

實際上是對于同一個計算輸入采用的不同計算方法。舉個例子，考慮1至100的累加，如果是CPU，最終運(yùn)算量計算結(jié)果則是需要調(diào)用執(zhí)行同一個累加函數(shù)100次；而GPU如果并行計算如果有4核，則是4個主線同時運(yùn)算這個算式，則計算量則是前者的1/4；而CUDA計算模式則可以看成優(yōu)化的GPU,即在這種常規(guī)的累加計算上，也會采用一些“討巧”的方式進(jìn)行。比如考慮累加的首尾兩側(cè)加起來都是100（比如1+99,2+98,3+97...）,因此，只需要考慮這種計算的累加次數(shù)就可以直接很快計算出結(jié)果了。從這個例子上看，CUDA的計算方式至少可以比單純地GPU減少一半的計算量。

3、CUDA的典型數(shù)學(xué)庫

類似于CPU編程庫，CUDA庫作為一系列接口的集合，只需要編寫host代碼，調(diào)用相應(yīng)API即可實現(xiàn)各種函數(shù)功能，這樣可以節(jié)約很多開發(fā)時間。并且CUDA的庫都是經(jīng)過大量的編程能力強(qiáng)的大拿經(jīng)過不斷優(yōu)化形成的，因此我們完全可以信任這些庫能夠達(dá)到很好的性能。當(dāng)然，完全依賴于這些庫而對CUDA性能優(yōu)化一無所知也是不行的，我們依然需要手動做一些改進(jìn)來挖掘出更好的性能。

CUDA上常用的庫包括如下幾種：

cuSPARSE線性代數(shù)庫，主要針對稀疏矩陣之類的

cuBLAS是CUDA標(biāo)準(zhǔn)的線性代數(shù)庫，該庫沒有專門針對稀疏矩陣的操作

cuFFT傅里葉變換，用于內(nèi)核加載時機(jī)和方式

cuRAND隨機(jī)數(shù)，用于隨機(jī)數(shù)的計算。

cuBLASLt

cuBLASLt 使用新的 FP8 數(shù)據(jù)類型公開混合精度乘法運(yùn)算。這些操作還支持 BF16 和 FP16 偏差融合，以及 FP16 偏差與 GELU 激活融合，用于具有 FP8 輸入和輸出數(shù)據(jù)類型的 GEMM。

在性能方面，與 A100 上的 BF16 相比，F(xiàn)P8 GEMM 在 H100 PCIe 和 SXM 上的速度分別提高了 3 倍和 4.5 倍。CUDA Math API 提供 FP8 轉(zhuǎn)換以方便使用新的 FP8 矩陣乘法運(yùn)算。

cuBLAS 12.0 擴(kuò)展了 API 以支持 64 位整數(shù)問題大小、前導(dǎo)維度和向量增量。這些新函數(shù)與其對應(yīng)的 32 位整數(shù)函數(shù)具有相同的 API，只是它們在名稱中具有 _64 后綴并將相應(yīng)的參數(shù)聲明為 int64_t。

cublasStatus_t cublasIsamax(cublasHandle_t handle, int n, const float *x, int incx, int *result);

對應(yīng)的 64 位整數(shù)如下：

cublasStatus_t cublasIsamax_64(cublasHandle_t handle, int64_t n, const float *x, int64_t incx, int64_t *result);

cuFFT

在計劃初始化期間，cuFFT 執(zhí)行一系列步驟（包括試探法）是用來確定使用了哪些內(nèi)核以及內(nèi)核模塊如何加載。從 CUDA 12.0 開始，cuFFT 使用 CUDA 并行線程執(zhí)行 (PTX) 匯編形式而不是二進(jìn)制形式交付了大部分內(nèi)核。

當(dāng)cuFFT 計劃初始化時，cuFFT 內(nèi)核的 PTX 代碼在運(yùn)行時由 CUDA 設(shè)備驅(qū)動程序加載并進(jìn)一步編譯為二進(jìn)制代碼。由于新的實施，第一個可用的改進(jìn)將為 NVIDIA Maxwell、NVIDIA Pascal、NVIDIA Volta 和 NVIDIA Turing 架構(gòu)啟用許多新的加速內(nèi)核。

cuSPARSE

為了減少稀疏矩陣乘法 (SpGEMM) 所需的工作空間，NVIDIA 發(fā)布了兩種內(nèi)存使用率較低的新算法。第一種算法計算中間產(chǎn)物數(shù)量的嚴(yán)格界限，而第二種算法允許將計算分成塊。這些新算法有利于使用較小內(nèi)存存儲設(shè)備的客戶。

最新INT8 已支持添加到 cusparseGather、cusparseScatter 和 cusparseCsr2cscEx2幾種不同的函數(shù)模塊中。

最后需要說明的是，最新版本中的CUDA也加入了延遲加載作為其中一部分。隨后的 CUDA 版本繼續(xù)增強(qiáng)和擴(kuò)展它。從應(yīng)用程序開發(fā)的角度來看，選擇延遲加載不需要任何特定的東西。現(xiàn)有的應(yīng)用程序可以按原樣使用延遲加載。延遲加載是一種延遲加載內(nèi)核和 CPU 端模塊直到應(yīng)用程序需要加載的技術(shù)。默認(rèn)是在第一次初始化庫時搶先加載所有模塊。這不僅可以顯著節(jié)省設(shè)備和主機(jī)內(nèi)存，還可以縮短算法的端到端執(zhí)行時間。

總結(jié)

本文從英偉達(dá)中的幾個核心軟件模塊GPU和CUDA開發(fā)角度分析了其在整個軟件框架構(gòu)建，軟件算法調(diào)度及軟件函數(shù)構(gòu)建和應(yīng)用上進(jìn)行了詳細(xì)的分析。重點是可以很好的掌握整個Orin系列的軟件架構(gòu)，與前序文章的硬件架構(gòu)相呼應(yīng)。當(dāng)然，從整個軟件開發(fā)的完整流程上，本文還只是入門級的說明，對于精細(xì)化的開發(fā)還是遠(yuǎn)遠(yuǎn)不夠的。后續(xù)文章我們將從操作系統(tǒng)Drive OS/驅(qū)動模塊Drive Work等角度進(jìn)行更為詳細(xì)的分析。