支持動態(tài)并行的CUDA擴(kuò)展功能和最佳應(yīng)用實(shí)踐

D.1. Introduction

D.1.1. Overview

Dynamic Parallelism是 CUDA 編程模型的擴(kuò)展，使 CUDA 內(nèi)核能夠直接在 GPU 上創(chuàng)建新工作并與新工作同步。在程序中需要的任何位置動態(tài)創(chuàng)建并行性提供了令人興奮的新功能。

直接從 GPU 創(chuàng)建工作的能力可以減少在主機(jī)和設(shè)備之間傳輸執(zhí)行控制和數(shù)據(jù)的需要，因?yàn)楝F(xiàn)在可以通過在設(shè)備上執(zhí)行的線程在運(yùn)行時做出啟動配置決策。此外，可以在運(yùn)行時在內(nèi)核內(nèi)內(nèi)聯(lián)生成依賴于數(shù)據(jù)的并行工作，動態(tài)利用 GPU 的硬件調(diào)度程序和負(fù)載平衡器，并根據(jù)數(shù)據(jù)驅(qū)動的決策或工作負(fù)載進(jìn)行調(diào)整。以前需要修改以消除遞歸、不規(guī)則循環(huán)結(jié)構(gòu)或其他不適合平面、單級并行性的構(gòu)造的算法和編程模式可以更透明地表達(dá)。

本文檔描述了支持動態(tài)并行的 CUDA 的擴(kuò)展功能，包括為利用這些功能而對 CUDA 編程模型進(jìn)行必要的修改和添加，以及利用此附加功能的指南和最佳實(shí)踐。

警告：與父塊的子內(nèi)核顯式同步（即在設(shè)備代碼中使用 cudaDeviceSynchronize（））在 CUDA 11.6 中已棄用，并計(jì)劃在未來的 CUDA 版本中刪除。

只有計(jì)算能力為 3.5 或更高的設(shè)備支持動態(tài)并行。

D.1.2. Glossary

本指南中使用的術(shù)語的定義。

Grid：網(wǎng)格是線程的集合。網(wǎng)格中的線程執(zhí)行內(nèi)核函數(shù)并被劃分為線程。

Thread Block：線程塊是在同一多處理器 （SM） 上執(zhí)行的一組線程。線程塊中的線程可以訪問共享內(nèi)存并且可以顯式同步。

Kernel Function：內(nèi)核函數(shù)是一個隱式并行子程序，它在 CUDA 執(zhí)行和內(nèi)存模型下為網(wǎng)格中的每個線程執(zhí)行。

Host:Host 指的是最初調(diào)用 CUDA 的執(zhí)行環(huán)境。通常是在系統(tǒng)的 CPU 處理器上運(yùn)行的線程。

Parent：父線程、線程塊或網(wǎng)格是已啟動新網(wǎng)格、子網(wǎng)格的一種。直到所有啟動的子網(wǎng)格也完成后，父節(jié)點(diǎn)才被視為完成。

Child：子線程、塊或網(wǎng)格是由父網(wǎng)格啟動的線程、塊或網(wǎng)格。子網(wǎng)格必須在父線程、線程塊或網(wǎng)格被認(rèn)為完成之前完成。

Thread Block Scope：具有線程塊作用域的對象具有單個線程塊的生命周期。它們僅在由創(chuàng)建對象的線程塊中的線程操作時具有定義的行為，并在創(chuàng)建它們的線程塊完成時被銷毀。

Device Runtime：設(shè)備運(yùn)行時是指可用于使內(nèi)核函數(shù)使用動態(tài)并行的運(yùn)行時系統(tǒng)和 API。

D.2. Execution Environment and Memory Model

D.2.1. Execution Environment

CUDA 執(zhí)行模型基于線程、線程塊和網(wǎng)格的原語，內(nèi)核函數(shù)定義了線程塊和網(wǎng)格內(nèi)的各個線程執(zhí)行的程序。 當(dāng)調(diào)用內(nèi)核函數(shù)時，網(wǎng)格的屬性由執(zhí)行配置描述，該配置在 CUDA 中具有特殊的語法。 CUDA 中對動態(tài)并行性的支持?jǐn)U展了在新網(wǎng)格上配置、啟動和同步到設(shè)備上運(yùn)行的線程的能力。

警告：與父塊的子內(nèi)核顯式同步（即在設(shè)備代碼中使用 cudaDeviceSynchronize（） 塊）在 CUDA 11.6 中已棄用，并計(jì)劃在未來的 CUDA 版本中刪除。

D.2.1.1. Parent and Child Grids

配置并啟動新網(wǎng)格的設(shè)備線程屬于父網(wǎng)格，調(diào)用創(chuàng)建的網(wǎng)格是子網(wǎng)格。

子網(wǎng)格的調(diào)用和完成是正確嵌套的，這意味著在其線程創(chuàng)建的所有子網(wǎng)格都完成之前，父網(wǎng)格不會被認(rèn)為是完整的。 即使調(diào)用線程沒有在啟動的子網(wǎng)格上顯式同步，運(yùn)行時也會保證父子網(wǎng)格之間的隱式同步。

警告：與父塊的子內(nèi)核顯式同步（即在設(shè)備代碼中使用 cudaDeviceSynchronize（））在 CUDA 11.6 中已棄用，并計(jì)劃在未來的 CUDA 版本中刪除。

D.2.1.2. Scope of CUDA Primitives

在主機(jī)和設(shè)備上，CUDA 運(yùn)行時都提供了一個 API，用于啟動內(nèi)核、等待啟動的工作完成以及通過流和事件跟蹤啟動之間的依賴關(guān)系。 在主機(jī)系統(tǒng)上，啟動狀態(tài)和引用流和事件的 CUDA 原語由進(jìn)程內(nèi)的所有線程共享； 但是進(jìn)程獨(dú)立執(zhí)行，可能不共享 CUDA 對象。

設(shè)備上存在類似的層次結(jié)構(gòu)：啟動的內(nèi)核和 CUDA 對象對線程塊中的所有線程都是可見的，但在線程塊之間是獨(dú)立的。 這意味著例如一個流可以由一個線程創(chuàng)建并由同一線程塊中的任何其他線程使用，但不能與任何其他線程塊中的線程共享。

D.2.1.3. Synchronization

警告：與父塊的子內(nèi)核顯式同步（即在設(shè)備代碼中使用 cudaDeviceSynchronize（））在 CUDA 11.6 中已棄用，并計(jì)劃在未來的 CUDA 版本中刪除。

來自任何線程的 CUDA 運(yùn)行時操作，包括內(nèi)核啟動，在線程塊中都是可見的。 這意味著父網(wǎng)格中的調(diào)用線程可以在由該線程啟動的網(wǎng)格、線程塊中的其他線程或在同一線程塊中創(chuàng)建的流上執(zhí)行同步。 直到塊中所有線程的所有啟動都完成后，才認(rèn)為線程塊的執(zhí)行完成。 如果一個塊中的所有線程在所有子啟動完成之前退出，將自動觸發(fā)同步操作。

D.2.1.4. Streams and Events

CUDA 流和事件允許控制網(wǎng)格啟動之間的依賴關(guān)系：啟動到同一流中的網(wǎng)格按順序執(zhí)行，事件可用于創(chuàng)建流之間的依賴關(guān)系。 在設(shè)備上創(chuàng)建的流和事件服務(wù)于這個完全相同的目的。

在網(wǎng)格中創(chuàng)建的流和事件存在于線程塊范圍內(nèi)，但在創(chuàng)建它們的線程塊之外使用時具有未定義的行為。 如上所述，線程塊啟動的所有工作在塊退出時都會隱式同步； 啟動到流中的工作包含在其中，所有依賴關(guān)系都得到了適當(dāng)?shù)慕鉀Q。 已在線程塊范圍之外修改的流上的操作行為未定義。

在主機(jī)上創(chuàng)建的流和事件在任何內(nèi)核中使用時具有未定義的行為，就像在子網(wǎng)格中使用時由父網(wǎng)格創(chuàng)建的流和事件具有未定義的行為一樣。

D.2.1.5. Ordering and Concurrency

從設(shè)備運(yùn)行時啟動內(nèi)核的順序遵循 CUDA Stream 排序語義。在一個線程塊內(nèi)，所有內(nèi)核啟動到同一個流中都是按順序執(zhí)行的。當(dāng)同一個線程塊中的多個線程啟動到同一個流中時，流內(nèi)的順序取決于塊內(nèi)的線程調(diào)度，這可以通過 __syncthreads（） 等同步原語進(jìn)行控制。

請注意，由于流由線程塊內(nèi)的所有線程共享，因此隱式 NULL 流也被共享。如果線程塊中的多個線程啟動到隱式流中，則這些啟動將按順序執(zhí)行。如果需要并發(fā)，則應(yīng)使用顯式命名流。

動態(tài)并行使并發(fā)在程序中更容易表達(dá)；但是，設(shè)備運(yùn)行時不會在 CUDA 執(zhí)行模型中引入新的并發(fā)保證。無法保證設(shè)備上任意數(shù)量的不同線程塊之間的并發(fā)執(zhí)行。

缺乏并發(fā)保證延伸到父線程塊及其子網(wǎng)格。當(dāng)父線程塊啟動子網(wǎng)格時，在父線程塊到達(dá)顯式同步點(diǎn)（例如 cudaDeviceSynchronize（））之前，不保證子網(wǎng)格開始執(zhí)行。

警告：與父塊的子內(nèi)核顯式同步（即在設(shè)備代碼中使用 cudaDeviceSynchronize（））在 CUDA 11.6 中已棄用，并計(jì)劃在未來的 CUDA 版本中刪除。

雖然并發(fā)通常很容易實(shí)現(xiàn)，但它可能會因設(shè)備配置、應(yīng)用程序工作負(fù)載和運(yùn)行時調(diào)度而異。因此，依賴不同線程塊之間的任何并發(fā)性是不安全的。

D.2.1.6. Device Management

設(shè)備運(yùn)行時不支持多 GPU； 設(shè)備運(yùn)行時只能在其當(dāng)前執(zhí)行的設(shè)備上運(yùn)行。 但是，允許查詢系統(tǒng)中任何支持 CUDA 的設(shè)備的屬性。

D.2.2. Memory Model

父網(wǎng)格和子網(wǎng)格共享相同的全局和常量內(nèi)存存儲，但具有不同的本地和共享內(nèi)存。

D.2.2.1. Coherence and Consistency

D.2.2.1.1. Global Memory

父子網(wǎng)格可以連貫地訪問全局內(nèi)存，但子網(wǎng)格和父網(wǎng)格之間的一致性保證很弱。當(dāng)子網(wǎng)格的內(nèi)存視圖與父線程完全一致時，子網(wǎng)格的執(zhí)行有兩點(diǎn)：當(dāng)子網(wǎng)格被父線程調(diào)用時，以及當(dāng)子網(wǎng)格線程完成時（由父線程中的同步 API 調(diào)用發(fā)出信號）。

警告：與父塊的子內(nèi)核顯式同步（即在設(shè)備代碼中使用 cudaDeviceSynchronize（））在 CUDA 11.6 中已棄用，并計(jì)劃在未來的 CUDA 版本中刪除。

在子網(wǎng)格調(diào)用之前，父線程中的所有全局內(nèi)存操作對子網(wǎng)格都是可見的。在父網(wǎng)格完成同步后，子網(wǎng)格的所有內(nèi)存操作對父網(wǎng)格都是可見的。

在下面的示例中，執(zhí)行 child_launch 的子網(wǎng)格只能保證看到在子網(wǎng)格啟動之前對數(shù)據(jù)所做的修改。由于父線程 0 正在執(zhí)行啟動，子線程將與父線程 0 看到的內(nèi)存保持一致。由于第一次 __syncthreads（） 調(diào)用，孩子將看到 data［0］=0， data［1］=1， 。..， data［255］=255（沒有 __syncthreads（） 調(diào)用，只有 data［0］將保證被孩子看到）。當(dāng)子網(wǎng)格返回時，線程 0 保證可以看到其子網(wǎng)格中的線程所做的修改。只有在第二次 __syncthreads（） 調(diào)用之后，這些修改才可用于父網(wǎng)格的其他線程：

__global__ void child_launch(int *data) {
   data[threadIdx.x] = data[threadIdx.x]+1;
}

__global__ void parent_launch(int *data) {
   data[threadIdx.x] = threadIdx.x;

   __syncthreads();

   if (threadIdx.x == 0) {
       child_launch<<< 1, 256 >>>(data);
       cudaDeviceSynchronize();
   }

   __syncthreads();
}

void host_launch(int *data) {
    parent_launch<<< 1, 256 >>>(data);
}

D.2.2.1.2. Zero Copy Memory

零拷貝系統(tǒng)內(nèi)存與全局內(nèi)存具有相同的一致性和一致性保證，并遵循上面詳述的語義。 內(nèi)核可能不會分配或釋放零拷貝內(nèi)存，但可能會使用從主機(jī)程序傳入的指向零拷貝的指針。

D.2.2.1.3. Constant Memory

常量是不可變的，不能從設(shè)備修改，即使在父子啟動之間也是如此。 也就是說，所有 __constant__ 變量的值必須在啟動之前從主機(jī)設(shè)置。 所有子內(nèi)核都從各自的父內(nèi)核自動繼承常量內(nèi)存。

從內(nèi)核線程中獲取常量內(nèi)存對象的地址與所有 CUDA 程序具有相同的語義，并且自然支持將該指針從父級傳遞給子級或從子級傳遞給父級。

D.2.2.1.4. Shared and Local Memory

共享內(nèi)存和本地內(nèi)存分別是線程塊或線程私有的，并且在父子之間不可見或不連貫。 當(dāng)這些位置之一中的對象在其所屬范圍之外被引用時，行為未定義，并且可能導(dǎo)致錯誤。

如果 NVIDIA 編譯器可以檢測到指向本地或共享內(nèi)存的指針作為參數(shù)傳遞給內(nèi)核啟動，它將嘗試發(fā)出警告。 在運(yùn)行時，程序員可以使用 __isGlobal（） 內(nèi)部函數(shù)來確定指針是否引用全局內(nèi)存，因此可以安全地傳遞給子啟動。

請注意，對 cudaMemcpy*Async（） 或 cudaMemset*Async（） 的調(diào)用可能會調(diào)用設(shè)備上的新子內(nèi)核以保留流語義。 因此，將共享或本地內(nèi)存指針傳遞給這些 API 是非法的，并且會返回錯誤。

D.2.2.1.5. Local Memory

本地內(nèi)存是執(zhí)行線程的私有存儲，在該線程之外不可見。 啟動子內(nèi)核時將指向本地內(nèi)存的指針作為啟動參數(shù)傳遞是非法的。 從子級取消引用此類本地內(nèi)存指針的結(jié)果將是未定義的。

例如，如果 child_launch 訪問 x_array，則以下內(nèi)容是非法的，具有未定義的行為：

int x_array[10];       // Creates x_array in parent's local memory 
child_launch<<< 1, 1 >>>(x_array);

程序員有時很難知道編譯器何時將變量放入本地內(nèi)存。 作為一般規(guī)則，傳遞給子內(nèi)核的所有存儲都應(yīng)該從全局內(nèi)存堆中顯式分配，或者使用cudaMalloc()、new()或通過在全局范圍內(nèi)聲明__device__存儲。 例如：

// Correct - "value" is global storage
__device__ int value; 
__device__ void x() { 
    value = 5; 
    child<<< 1, 1 >>>(&value); 
}

// Invalid - "value" is local storage
__device__ void y() { 
    int value = 5; 
    child<<< 1, 1 >>>(&value); 
}

D.2.2.1.6. Texture Memory

對紋理映射的全局內(nèi)存區(qū)域的寫入相對于紋理訪問是不連貫的。 紋理內(nèi)存的一致性在子網(wǎng)格的調(diào)用和子網(wǎng)格完成時強(qiáng)制執(zhí)行。 這意味著在子內(nèi)核啟動之前寫入內(nèi)存會反映在子內(nèi)核的紋理內(nèi)存訪問中。 類似地，子進(jìn)程對內(nèi)存的寫入將反映在父進(jìn)程對紋理內(nèi)存的訪問中，但只有在父進(jìn)程同步子進(jìn)程完成之后。 父子并發(fā)訪問可能會導(dǎo)致數(shù)據(jù)不一致。

警告：與父塊的子內(nèi)核顯式同步（即在設(shè)備代碼中使用 cudaDeviceSynchronize（））在 CUDA 11.6 中已棄用，并計(jì)劃在未來的 CUDA 版本中刪除。

D.3. Programming Interface

D.3.1. CUDA C++ Reference

內(nèi)核可以使用標(biāo)準(zhǔn) CUDA 《《《 》》》 語法從設(shè)備啟動：

kernel_name<<< Dg, Db, Ns, S >>>([kernel arguments]);

__global__ void permute(int n, int *data) {
   extern __shared__ int smem[];
   if (n <= 1)
       return;

   smem[threadIdx.x] = data[threadIdx.x];
   __syncthreads();

   permute_data(smem, n);
   __syncthreads();

   // Write back to GMEM since we can't pass SMEM to children.
   data[threadIdx.x] = smem[threadIdx.x];
   __syncthreads();

   if (threadIdx.x == 0) {
       permute<<< 1, 256, n/2*sizeof(int) >>>(n/2, data);
       permute<<< 1, 256, n/2*sizeof(int) >>>(n/2, data+n/2);
   }
}

void host_launch(int *data) {
    permute<<< 1, 256, 256*sizeof(int) >>>(256, data);
}

D.3.1.6.4. Symbol Addresses

設(shè)備端符號（即標(biāo)記為 __device__ 的符號）可以簡單地通過 & 運(yùn)算符從內(nèi)核中引用，因?yàn)樗腥址秶脑O(shè)備變量都在內(nèi)核的可見地址空間中。 這也適用于 __constant__ 符號，盡管在這種情況下指針將引用只讀數(shù)據(jù)。

鑒于可以直接引用設(shè)備端符號，那些引用符號的 CUDA 運(yùn)行時 API（例如 cudaMemcpyToSymbol（） 或 cudaGetSymbolAddress（））是多余的，因此設(shè)備運(yùn)行時不支持。 請注意，這意味著常量數(shù)據(jù)不能在正在運(yùn)行的內(nèi)核中更改，即使在子內(nèi)核啟動之前也是如此，因?yàn)閷?__constant__ 空間的引用是只讀的。

D.3.1.7. API Errors and Launch Failures

與 CUDA 運(yùn)行時一樣，任何函數(shù)都可能返回錯誤代碼。 最后返回的錯誤代碼被記錄下來，并且可以通過 cudaGetLastError（） 調(diào)用來檢索。 每個線程都會記錄錯誤，以便每個線程都可以識別它最近生成的錯誤。 錯誤代碼的類型為 cudaError_t。

與主機(jī)端啟動類似，設(shè)備端啟動可能由于多種原因（無效參數(shù)等）而失敗。 用戶必須調(diào)用 cudaGetLastError（） 來確定啟動是否產(chǎn)生錯誤，但是啟動后沒有錯誤并不意味著子內(nèi)核成功完成。

對于設(shè)備端異常，例如，訪問無效地址，子網(wǎng)格中的錯誤將返回給主機(jī)，而不是由父調(diào)用 cudaDeviceSynchronize（） 返回。

D.3.1.7.1. Launch Setup APIs

內(nèi)核啟動是通過設(shè)備運(yùn)行時庫公開的系統(tǒng)級機(jī)制，因此可通過底層 cudaGetParameterBuffer（） 和 cudaLaunchDevice（） API 直接從 PTX 獲得。 允許 CUDA 應(yīng)用程序自己調(diào)用這些 API，其要求與 PTX 相同。 在這兩種情況下，用戶都負(fù)責(zé)根據(jù)規(guī)范以正確的格式正確填充所有必要的數(shù)據(jù)結(jié)構(gòu)。 這些數(shù)據(jù)結(jié)構(gòu)保證了向后兼容性。

與主機(jī)端啟動一樣，設(shè)備端操作符 《《《》》》 映射到底層內(nèi)核啟動 API。 這樣一來，以 PTX 為目標(biāo)的用戶將能夠啟動加載，并且編譯器前端可以將 《《《》》》 轉(zhuǎn)換為這些調(diào)用。

D.3.1.8. API Reference

此處詳細(xì)介紹了設(shè)備運(yùn)行時支持的 CUDA 運(yùn)行時 API 部分。 主機(jī)和設(shè)備運(yùn)行時 API 具有相同的語法； 語義是相同的，除非另有說明。 下表提供了與主機(jī)可用版本相關(guān)的 API 概覽。

D.3.2. Device-side Launch from PTX

本部分適用于以并行線程執(zhí)行 （PTX） 為目標(biāo)并計(jì)劃在其語言中支持動態(tài)并行的編程語言和編譯器實(shí)現(xiàn)者。 它提供了與在 PTX 級別支持內(nèi)核啟動相關(guān)的底層詳細(xì)信息。

D.3.2.1. Kernel Launch APIs

可以使用可從 PTX 訪問的以下兩個 API 來實(shí)現(xiàn)設(shè)備端內(nèi)核啟動：cudaLaunchDevice（） 和 cudaGetParameterBuffer（）。 cudaLaunchDevice（） 使用通過調(diào)用 cudaGetParameterBuffer（） 獲得的參數(shù)緩沖區(qū)啟動指定的內(nèi)核，并將參數(shù)填充到啟動的內(nèi)核。 參數(shù)緩沖區(qū)可以為 NULL，即，如果啟動的內(nèi)核不帶任何參數(shù)，則無需調(diào)用 cudaGetParameterBuffer（）。

D.3.2.1.1. cudaLaunchDevice

在 PTX 級別，cudaLaunchDevice（） 需要在使用前以如下所示的兩種形式之一聲明。

// PTX-level Declaration of cudaLaunchDevice() when .address_size is 64
.extern .func(.param .b32 func_retval0) cudaLaunchDevice 
( 
  .param .b64 func, 
  .param .b64 parameterBuffer, 
  .param .align 4 .b8 gridDimension[12], 
  .param .align 4 .b8 blockDimension[12], 
  .param .b32 sharedMemSize, 
  .param .b64 stream 
) 
;

// PTX-level Declaration of cudaLaunchDevice() when .address_size is 32
.extern .func(.param .b32 func_retval0) cudaLaunchDevice
(
  .param .b32 func,
  .param .b32 parameterBuffer,
  .param .align 4 .b8 gridDimension[12],
  .param .align 4 .b8 blockDimension[12],
  .param .b32 sharedMemSize,
  .param .b32 stream
)
;

下面的 CUDA 級聲明映射到上述 PTX 級聲明之一，可在系統(tǒng)頭文件cuda_device_runtime_api.h中找到。 該函數(shù)在cudadevrt系統(tǒng)庫中定義，必須與程序鏈接才能使用設(shè)備端內(nèi)核啟動功能。

// CUDA-level declaration of cudaLaunchDevice()
extern "C" __device__ 
cudaError_t cudaLaunchDevice(void *func, void *parameterBuffer, 
                             dim3 gridDimension, dim3 blockDimension, 
                             unsigned int sharedMemSize, 
                             cudaStream_t stream);

第一個參數(shù)是指向要啟動的內(nèi)核的指針，第二個參數(shù)是保存已啟動內(nèi)核的實(shí)際參數(shù)的參數(shù)緩沖區(qū)。 參數(shù)緩沖區(qū)的布局在下面的參數(shù)緩沖區(qū)布局中進(jìn)行了說明。 其他參數(shù)指定啟動配置，即網(wǎng)格維度、塊維度、共享內(nèi)存大小以及啟動關(guān)聯(lián)的流（啟動配置的詳細(xì)說明請參見執(zhí)行配置)。

D.3.2.1.2. cudaGetParameterBuffer

cudaGetParameterBuffer()需要在使用前在 PTX 級別聲明。 PTX 級聲明必須采用以下兩種形式之一，具體取決于地址大?。?/p>

// PTX-level Declaration of cudaGetParameterBuffer() when .address_size is 64
// When .address_size is 64
.extern .func(.param .b64 func_retval0) cudaGetParameterBuffer
(
  .param .b64 alignment,
  .param .b64 size
)
;

 // PTX-level Declaration of cudaGetParameterBuffer() when .address_size is 32
.extern .func(.param .b32 func_retval0) cudaGetParameterBuffer
(
  .param .b32 alignment,
  .param .b32 size
)
;

cudaGetParameterBuffer()的以下 CUDA 級聲明映射到上述 PTX 級聲明：

// CUDA-level Declaration of cudaGetParameterBuffer()
extern "C" __device__
void *cudaGetParameterBuffer(size_t alignment, size_t size);

第一個參數(shù)指定參數(shù)緩沖區(qū)的對齊要求，第二個參數(shù)以字節(jié)為單位的大小要求。 在當(dāng)前實(shí)現(xiàn)中，cudaGetParameterBuffer（） 返回的參數(shù)緩沖區(qū)始終保證為 64 字節(jié)對齊，忽略對齊要求參數(shù)。 但是，建議將正確的對齊要求值（即要放置在參數(shù)緩沖區(qū)中的任何參數(shù)的最大對齊）傳遞給 cudaGetParameterBuffer（） 以確保將來的可移植性。

D.3.2.2. Parameter Buffer Layout

禁止參數(shù)緩沖區(qū)中的參數(shù)重新排序，并且要求放置在參數(shù)緩沖區(qū)中的每個單獨(dú)的參數(shù)對齊。 也就是說，每個參數(shù)必須放在參數(shù)緩沖區(qū)中的第 n 個字節(jié)，其中 n 是參數(shù)大小的最小倍數(shù)，它大于前一個參數(shù)占用的最后一個字節(jié)的偏移量。 參數(shù)緩沖區(qū)的最大大小為 4KB。

有關(guān) CUDA 編譯器生成的 PTX 代碼的更詳細(xì)說明，請參閱 PTX-3.5 規(guī)范。

D.3.3. Toolkit Support for Dynamic Parallelism

D.3.3.1. Including Device Runtime API in CUDA Code

與主機(jī)端運(yùn)行時 API 類似，CUDA 設(shè)備運(yùn)行時 API 的原型會在程序編譯期間自動包含在內(nèi)。 無需明確包含 cuda_device_runtime_api.h。

D.3.3.2. Compiling and Linking

當(dāng)使用帶有 nvcc 的動態(tài)并行編譯和鏈接 CUDA 程序時，程序?qū)⒆詣渔溄拥届o態(tài)設(shè)備運(yùn)行時庫 libcudadevrt。

設(shè)備運(yùn)行時作為靜態(tài)庫（Windows 上的 cudadevrt.lib，Linux 下的 libcudadevrt.a）提供，必須鏈接使用設(shè)備運(yùn)行時的 GPU 應(yīng)用程序。設(shè)備庫的鏈接可以通過 nvcc 或 nvlink 完成。下面顯示了兩個簡單的示例。

如果可以從命令行指定所有必需的源文件，則可以在一個步驟中編譯和鏈接設(shè)備運(yùn)行時程序：

$ nvcc -arch=sm_35 -rdc=true hello_world.cu -o hello -lcudadevrt

也可以先將 CUDA .cu 源文件編譯為目標(biāo)文件，然后在兩個階段的過程中將它們鏈接在一起：

$ nvcc -arch=sm_35 -dc hello_world.cu -o hello_world.o

$ nvcc -arch=sm_35 -rdc=true hello_world.o -o hello -lcudadevrt

有關(guān)詳細(xì)信息，請參閱 The CUDA Driver Compiler NVCC的使用單獨(dú)編譯部分。

D.4. Programming Guidelines

D.4.1. Basics

設(shè)備運(yùn)行時是主機(jī)運(yùn)行時的功能子集。 API 級別的設(shè)備管理、內(nèi)核啟動、設(shè)備 memcpy、流管理和事件管理從設(shè)備運(yùn)行時公開。

已經(jīng)有 CUDA 經(jīng)驗(yàn)的人應(yīng)該熟悉設(shè)備運(yùn)行時的編程。 設(shè)備運(yùn)行時語法和語義與主機(jī) API 基本相同，但本文檔前面詳細(xì)介紹了任何例外情況。

警告：與父塊的子內(nèi)核顯式同步（即在設(shè)備代碼中使用 cudaDeviceSynchronize（））在 CUDA 11.6 中已棄用，并計(jì)劃在未來的 CUDA 版本中刪除。

以下示例顯示了一個包含動態(tài)并行性的簡單 Hello World 程序：

#include  

__global__ void childKernel() 
{ 
    printf("Hello "); 
} 

__global__ void parentKernel() 
{ 
    // launch child 
    childKernel<<<1,1>>>(); 
    if (cudaSuccess != cudaGetLastError()) { 
        return; 
    }

    // wait for child to complete 
    if (cudaSuccess != cudaDeviceSynchronize()) { 
        return; 
    } 

    printf("World!\n"); 
} 

int main(int argc, char *argv[]) 
{ 
    // launch parent 
    parentKernel<<<1,1>>>(); 
    if (cudaSuccess != cudaGetLastError()) { 
        return 1; 
    } 

    // wait for parent to complete 
    if (cudaSuccess != cudaDeviceSynchronize()) { 
        return 2; 
    } 

    return 0; 
}

該程序可以從命令行一步構(gòu)建，如下所示：

$ nvcc -arch=sm_35 -rdc=true hello_world.cu -o hello -lcudadevrt

D.4.2. Performance

D.4.2.1. Synchronization

警告：與父塊的子內(nèi)核顯式同步（即在設(shè)備代碼中使用 cudaDeviceSynchronize（））在 CUDA 11.6 中已棄用，并計(jì)劃在未來的 CUDA 版本中刪除。

一個線程的同步可能會影響同一線程塊中其他線程的性能，即使這些其他線程自己不調(diào)用 cudaDeviceSynchronize（） 也是如此。 這種影響將取決于底層實(shí)現(xiàn)。 通常，與顯式調(diào)用 cudaDeviceSynchronize（） 相比，在線程塊結(jié)束時完成子內(nèi)核的隱式同步更有效。 因此，如果需要在線程塊結(jié)束之前與子內(nèi)核同步，建議僅調(diào)用 cudaDeviceSynchronize（）。

D.4.2.2. Dynamic-parallelism-enabled Kernel Overhead

在控制動態(tài)啟動時處于活動狀態(tài)的系統(tǒng)軟件可能會對當(dāng)時正在運(yùn)行的任何內(nèi)核施加開銷，無論它是否調(diào)用自己的內(nèi)核啟動。 這種開銷來自設(shè)備運(yùn)行時的執(zhí)行跟蹤和管理軟件，并且可能導(dǎo)致性能下降，例如，與從主機(jī)端相比，從設(shè)備進(jìn)行庫調(diào)用時。 通常，鏈接到設(shè)備運(yùn)行時庫的應(yīng)用程序會產(chǎn)生這種開銷。

D.4.3. Implementation Restrictions and Limitations

動態(tài)并行保證本文檔中描述的所有語義，但是，某些硬件和軟件資源依賴于實(shí)現(xiàn)，并限制了使用設(shè)備運(yùn)行時的程序的規(guī)模、性能和其他屬性。

D.4.3.1. Runtime

D.4.3.1.1. Memory Footprint

設(shè)備運(yùn)行時系統(tǒng)軟件為各種管理目的預(yù)留內(nèi)存，特別是用于在同步期間保存父網(wǎng)格狀態(tài)的一個預(yù)留，以及用于跟蹤未決網(wǎng)格啟動的第二個預(yù)留。 配置控制可用于減少這些預(yù)留的大小，以換取某些啟動限制。 有關(guān)詳細(xì)信息，請參閱下面的配置選項(xiàng)。

大多數(shù)保留內(nèi)存被分配為父內(nèi)核狀態(tài)的后備存儲，用于在子啟動時進(jìn)行同步。 保守地說，該內(nèi)存必須支持為設(shè)備上可能的最大活動線程數(shù)存儲狀態(tài)。 這意味著可調(diào)用 cudaDeviceSynchronize（） 的每個父代可能需要多達(dá) 860MB 的設(shè)備內(nèi)存，具體取決于設(shè)備配置，即使它沒有全部消耗，也將無法供程序使用。

D.4.3.1.2. Nesting and Synchronization Depth

使用設(shè)備運(yùn)行時，一個內(nèi)核可能會啟動另一個內(nèi)核，而該內(nèi)核可能會啟動另一個內(nèi)核，以此類推。每個從屬啟動都被認(rèn)為是一個新的嵌套層級，層級總數(shù)就是程序的嵌套深度。同步深度定義為程序在子啟動時顯式同步的最深級別。通常這比程序的嵌套深度小一，但如果程序不需要在所有級別調(diào)用 cudaDeviceSynchronize（） ，則同步深度可能與嵌套深度有很大不同。

警告：與父塊的子內(nèi)核顯式同步（即在設(shè)備代碼中使用 cudaDeviceSynchronize（））在 CUDA 11.6 中已棄用，并計(jì)劃在未來的 CUDA 版本中刪除。

總體最大嵌套深度限制為 24，但實(shí)際上，真正的限制將是系統(tǒng)為每個新級別所需的內(nèi)存量（請參閱上面的內(nèi)存占用量）。任何會導(dǎo)致內(nèi)核處于比最大值更深的級別的啟動都將失敗。請注意，這也可能適用于 cudaMemcpyAsync（），它本身可能會生成內(nèi)核啟動。有關(guān)詳細(xì)信息，請參閱配置選項(xiàng)。

默認(rèn)情況下，為兩級同步保留足夠的存儲空間。這個最大同步深度（以及因此保留的存儲）可以通過調(diào)用 cudaDeviceSetLimit（） 并指定 cudaLimitDevRuntimeSyncDepth 來控制。必須在主機(jī)啟動頂層內(nèi)核之前配置要支持的層數(shù)，以保證嵌套程序的成功執(zhí)行。在大于指定最大同步深度的深度調(diào)用 cudaDeviceSynchronize（） 將返回錯誤。

在父內(nèi)核從不調(diào)用 cudaDeviceSynchronize（） 的情況下，如果系統(tǒng)檢測到不需要為父狀態(tài)保留空間，則允許進(jìn)行優(yōu)化。在這種情況下，由于永遠(yuǎn)不會發(fā)生顯式父/子同步，因此程序所需的內(nèi)存占用量將遠(yuǎn)小于保守的最大值。這樣的程序可以指定較淺的最大同步深度，以避免過度分配后備存儲。

D.4.3.1.3. Pending Kernel Launches

啟動內(nèi)核時，會跟蹤所有關(guān)聯(lián)的配置和參數(shù)數(shù)據(jù)，直到內(nèi)核完成。 此數(shù)據(jù)存儲在系統(tǒng)管理的啟動池中。

啟動池分為固定大小的池和性能較低的虛擬化池。 設(shè)備運(yùn)行時系統(tǒng)軟件將首先嘗試跟蹤固定大小池中的啟動數(shù)據(jù)。 當(dāng)固定大小的池已滿時，虛擬化池將用于跟蹤新的啟動。

固定大小啟動池的大小可通過從主機(jī)調(diào)用 cudaDeviceSetLimit（） 并指定 cudaLimitDevRuntimePendingLaunchCount 來配置。

D.4.3.1.4. Configuration Options

設(shè)備運(yùn)行時系統(tǒng)軟件的資源分配通過主機(jī)程序的 cudaDeviceSetLimit（） API 進(jìn)行控制。 限制必須在任何內(nèi)核啟動之前設(shè)置，并且在 GPU 正在運(yùn)行程序時不得更改。

警告：與父塊的子內(nèi)核顯式同步（即在設(shè)備代碼中使用 cudaDeviceSynchronize（））在 CUDA 11.6 中已棄用，并計(jì)劃在未來的 CUDA 版本中刪除。

可以設(shè)置以下命名限制：

D.4.3.1.5. Memory Allocation and Lifetime

cudaMalloc（） 和 cudaFree（） 在主機(jī)和設(shè)備環(huán)境之間具有不同的語義。 當(dāng)從主機(jī)調(diào)用時，cudaMalloc（） 從未使用的設(shè)備內(nèi)存中分配一個新區(qū)域。 當(dāng)從設(shè)備運(yùn)行時調(diào)用時，這些函數(shù)映射到設(shè)備端的 malloc（） 和 free（）。 這意味著在設(shè)備環(huán)境中，總可分配內(nèi)存限制為設(shè)備 malloc（） 堆大小，它可能小于可用的未使用設(shè)備內(nèi)存。 此外，在設(shè)備上由 cudaMalloc（） 分配的指針上從主機(jī)程序調(diào)用 cudaFree（） 是錯誤的，反之亦然。

請注意，在 PTX 中，%smid 和 %warpid 被定義為 volatile 值。 設(shè)備運(yùn)行時可以將線程塊重新調(diào)度到不同的 SM 上，以便更有效地管理資源。 因此，依賴 %smid 或 %warpid 在線程或線程塊的生命周期內(nèi)保持不變是不安全的。

D.4.3.1.7. ECC Errors

CUDA 內(nèi)核中的代碼沒有可用的 ECC 錯誤通知。 整個啟動樹完成后，主機(jī)端會報(bào)告 ECC 錯誤。 在嵌套程序執(zhí)行期間出現(xiàn)的任何 ECC 錯誤都將生成異?；蚶^續(xù)執(zhí)行（取決于錯誤和配置）。

關(guān)于作者

Ken He 是 NVIDIA 企業(yè)級開發(fā)者社區(qū)經(jīng)理 & 高級講師，擁有多年的 GPU 和人工智能開發(fā)經(jīng)驗(yàn)。自 2017 年加入 NVIDIA 開發(fā)者社區(qū)以來，完成過上百場培訓(xùn)，幫助上萬個開發(fā)者了解人工智能和 GPU 編程開發(fā)。在計(jì)算機(jī)視覺，高性能計(jì)算領(lǐng)域完成過多個獨(dú)立項(xiàng)目。并且，在機(jī)器人和無人機(jī)領(lǐng)域，有過豐富的研發(fā)經(jīng)驗(yàn)。對于圖像識別，目標(biāo)的檢測與跟蹤完成過多種解決方案。曾經(jīng)參與 GPU 版氣象模式GRAPES，是其主要研發(fā)者。

審核編輯：郭婷