在 本系列的第 1 部分 中，我們引入了新的 API 函數(shù) cudaMallocAsync 和 cudaFreeAsync ，它們使內(nèi)存分配和釋放成為流順序操作。在這篇文章中，我們通過(guò)分享一些大數(shù)據(jù)基準(zhǔn)測(cè)試結(jié)果來(lái)強(qiáng)調(diào)這一新功能的好處，并為修改現(xiàn)有應(yīng)用程序提供代碼 MIG 定量指南。我們還介紹了在多 GPU 訪問(wèn)和 IPC 使用環(huán)境中利用流順序內(nèi)存分配的高級(jí)主題。這一切都有助于提高現(xiàn)有應(yīng)用程序的性能。

GPU 大數(shù)據(jù)基準(zhǔn)

為了衡量新的流式有序分配器在實(shí)際應(yīng)用程序中的性能影響，以下是來(lái)自 RAPIDS GPU 大數(shù)據(jù)基準(zhǔn) （ GPU -bdb］的結(jié)果。 GPU -bdb 是 30 個(gè)查詢的基準(zhǔn)，這些查詢以各種比例因子表示現(xiàn)實(shí)世界的數(shù)據(jù)科學(xué)和機(jī)器學(xué)習(xí)工作流： SF1000 是 1 TB 的數(shù)據(jù)， SF10000 是 10 TB 的數(shù)據(jù)。事實(shí)上，每個(gè)查詢都是一個(gè)模型工作流，可以包括 SQL 、用戶定義函數(shù)、仔細(xì)的子集和聚合以及機(jī)器學(xué)習(xí)。

圖 1 顯示了在 SF1000 上在 NVIDIA DGX-2 上跨 16 個(gè) V100 GPU 執(zhí)行的 gpu-bdb 查詢子集的 cudaMallocAsync 與 cudaMalloc 的性能比較。如您所見(jiàn)，由于內(nèi)存重用和消除無(wú)關(guān)同步，使用 cudaMallocAsync 時(shí)端到端性能提高了 2-5 倍。

圖 1 加速 cudaMallocAsync 結(jié)束 cudaMalloc 對(duì)于 RAPIDS GPU 大數(shù)據(jù)基準(zhǔn)的各種查詢 。

與 CUDA Malloc 和 CUDA Free 的互操作性

應(yīng)用程序可以使用 cudaFreeAsync 釋放 cudaMalloc 分配的指針。在下一次同步傳遞到 cudaFreeAsync 的流之前，不會(huì)釋放基礎(chǔ)內(nèi)存。

cudaMalloc(&ptr, size);
kernel<<<..., stream>>>(ptr);
cudaFreeAsync(ptr, stream);
cudaStreamSynchronize(stream); // The memory for ptr is freed at this point 

類似地，應(yīng)用程序可以使用 cudaFree 釋放使用 cudaMallocAsync 分配的內(nèi)存。但是，在這種情況下， cudaFree 不會(huì)隱式同步，因此應(yīng)用程序必須插入適當(dāng)?shù)耐剑源_保對(duì)要釋放的內(nèi)存的所有訪問(wèn)都已完成。任何有意或無(wú)意依賴 cudaFree 的隱式同步行為的應(yīng)用程序代碼都必須更新。

cudaMallocAsync(&ptr, size, stream);
kernel<<<..., stream>>>(ptr);
cudaStreamSynchronize(stream); // Must synchronize first
cudaFree(ptr);

多 – GPU 訪問(wèn)

默認(rèn)情況下，可以從與指定流關(guān)聯(lián)的設(shè)備訪問(wèn)使用 cudaMallocAsync 分配的內(nèi)存。從任何其他設(shè)備訪問(wèn)內(nèi)存需要啟用從該其他設(shè)備訪問(wèn)整個(gè)池。正如 cudaDeviceCanAccessPeer 所報(bào)告的，它還要求這兩個(gè)設(shè)備具有對(duì)等功能。與 cudaMalloc 分配不同， cudaDeviceEnablePeerAccess 和 cudaDeviceDisablePeerAccess 對(duì)從內(nèi)存池分配的內(nèi)存沒(méi)有影響。

例如，考慮啟用設(shè)備 4Access 到設(shè)備 3 的內(nèi)存池：

cudaMemPool_t mempool;
cudaDeviceGetDefaultMemPool(&mempool, 3);
cudaMemAccessDesc desc = {};
desc.location.type = cudaMemLocationTypeDevice;
desc.location.id = 4;
desc.flags = cudaMemAccessFlagsProtReadWrite;
cudaMemPoolSetAccess(mempool, &desc, 1 /* numDescs */); 

調(diào)用 cudaMemPoolSetAccess 時(shí)，可以使用 cudaMemAccessFlagsProtNone 撤銷對(duì)內(nèi)存池所在設(shè)備以外的設(shè)備的訪問(wèn)。無(wú)法撤消對(duì)內(nèi)存池自身設(shè)備的訪問(wèn)。

進(jìn)程間通信支持

使用與設(shè)備關(guān)聯(lián)的默認(rèn)內(nèi)存池分配的內(nèi)存不能與其他進(jìn)程共享。應(yīng)用程序必須顯式創(chuàng)建自己的內(nèi)存池，以便與其他進(jìn)程共享使用 cudaMallocAsync 分配的內(nèi)存。以下代碼示例顯示如何創(chuàng)建具有進(jìn)程間通信（ IPC ）功能的顯式內(nèi)存池：

cudaMemPool_t exportPool;
cudaMemPoolProps poolProps = {};
poolProps.allocType = cudaMemAllocationTypePinned;
poolProps.handleTypes = cudaMemHandleTypePosixFileDescriptor;
poolProps.location.type = cudaMemLocationTypeDevice;
poolProps.location.id = deviceId;
cudaMemPoolCreate(&exportPool, &poolProps); 

位置類型設(shè)備和位置 ID deviceId 指示必須在特定 GPU 上分配池內(nèi)存。分配類型 pinted 表示內(nèi)存應(yīng)該是 non-migratable ，也稱為不可分頁(yè)。句柄類型 PosixFileDescriptor 表示用戶打算查詢池的文件描述符，以便與其他進(jìn)程共享。

通過(guò) IPC 共享此池中的內(nèi)存的第一步是查詢表示該池的文件描述符：

int fd;
cudaMemAllocationHandleType handleType = cudaMemHandleTypePosixFileDescriptor;
cudaMemPoolExportToShareableHandle(&fd, exportPool, handleType, 0); 

然后，應(yīng)用程序可以與另一個(gè)進(jìn)程共享文件描述符，例如通過(guò) UNIX 域套接字。然后，另一個(gè)進(jìn)程可以導(dǎo)入文件描述符并獲得進(jìn)程本地池句柄：

cudaMemPool_t importPool;
cudaMemAllocationHandleType handleType = cudaMemHandleTypePosixFileDescriptor;
cudaMemPoolImportFromShareableHandle(&importPool, &fd, handleType, 0); 

下一步是導(dǎo)出過(guò)程從池中分配內(nèi)存：

cudaMallocFromPoolAsync(&ptr, size, exportPool, stream); 

cudaMallocAsync還有一個(gè)重載版本，它采用與cudaMallocFromPoolAsync相同的參數(shù)：

cudaMallocAsync(&ptr, size, exportPool, stream); 

通過(guò)這兩個(gè) API 中的任何一個(gè)從該池分配內(nèi)存后，指針就可以與導(dǎo)入進(jìn)程共享。首先，導(dǎo)出過(guò)程獲得一個(gè)表示內(nèi)存分配的不透明句柄：

cudaMemPoolPtrExportData data;
cudaMemPoolExportPointer(&data, ptr); 

然后，可以通過(guò)任何標(biāo)準(zhǔn) IPC 機(jī)制（例如通過(guò)共享內(nèi)存、管道等）與導(dǎo)入進(jìn)程共享此不透明數(shù)據(jù)。導(dǎo)入進(jìn)程然后將不透明數(shù)據(jù)轉(zhuǎn)換為進(jìn)程本地指針：

cudaMemPoolImportPointer(&ptr, importPool, &data); 

現(xiàn)在，兩個(gè)進(jìn)程共享對(duì)相同內(nèi)存分配的訪問(wèn)。在導(dǎo)出過(guò)程中釋放內(nèi)存之前，必須先在導(dǎo)入過(guò)程中釋放內(nèi)存。這是為了確保在導(dǎo)出過(guò)程中，當(dāng)導(dǎo)入過(guò)程仍在訪問(wèn)以前的共享內(nèi)存分配時(shí)，內(nèi)存不會(huì)重新用于另一個(gè) cudaMallocAsync 請(qǐng)求，從而可能導(dǎo)致未定義的行為。

現(xiàn)有函數(shù) cudaIpcGetMemHandle 僅適用于通過(guò) cudaMalloc 分配的內(nèi)存，不能用于通過(guò) cudaMallocAsync 分配的任何內(nèi)存，無(wú)論該內(nèi)存是否從顯式池分配。

更改設(shè)備池

如果應(yīng)用程序期望大部分時(shí)間使用顯式內(nèi)存池，則可以考慮通過(guò) cudaDeviceSetMemPool 將其設(shè)置為設(shè)備的當(dāng)前池。這使應(yīng)用程序可以避免每次必須從池中分配內(nèi)存時(shí)都必須指定池參數(shù)。

cudaDeviceSetMemPool(device, pool);
cudaMallocAsync(&ptr, size, stream); // This now allocates from the earlier pool set instead of the device’s default pool. 

這樣做的好處是，使用 cudaMallocAsync 分配的任何其他函數(shù)現(xiàn)在都會(huì)自動(dòng)使用新池作為默認(rèn)池。可以使用 cudaDeviceGetMemPool 查詢與設(shè)備關(guān)聯(lián)的當(dāng)前池。

庫(kù)可組合性

通常，庫(kù)不應(yīng)該更改設(shè)備的池，因?yàn)檫@樣做會(huì)影響整個(gè)頂級(jí)應(yīng)用程序。如果庫(kù)必須分配具有不同于默認(rèn)設(shè)備池屬性的內(nèi)存，它可以創(chuàng)建自己的池，然后使用 cudaMallocFromPoolAsync 從該池進(jìn)行分配。該庫(kù)還可以使用 cudaMallocAsync 的重載版本，該版本將池作為參數(shù)。

為了使應(yīng)用程序的互操作更容易，庫(kù)應(yīng)該考慮為頂級(jí)應(yīng)用程序提供 API 以協(xié)調(diào)所使用的池。例如，庫(kù)可以提供 set 或 get API ，使應(yīng)用程序能夠以更明確的方式控制池。庫(kù)還可以將池作為單個(gè) API 的參數(shù)。

代碼遷移指南

當(dāng)將使用 cudaMalloc 或 cudaFree 的現(xiàn)有應(yīng)用程序移植到新的 cudaMallocAsync 或 cudaFreeAsync API 時(shí)，考慮以下準(zhǔn)則。

確定適當(dāng)人才庫(kù)的指南：

初始默認(rèn)池適用于許多應(yīng)用程序。

今天，顯式構(gòu)造的池只需要在與 CUDA IPC 的進(jìn)程之間共享池內(nèi)存。這可能會(huì)隨著將來(lái)的功能而改變。

為了方便起見(jiàn)，考慮將顯式創(chuàng)建池設(shè)置為設(shè)備的當(dāng)前池，以確保進(jìn)程內(nèi)的所有 cudaMallocAsync 調(diào)用都使用該池。這必須由頂級(jí)應(yīng)用程序而不是庫(kù)來(lái)完成，以避免與頂級(jí)應(yīng)用程序的目標(biāo)沖突。

為所有內(nèi)存池設(shè)置釋放閾值的準(zhǔn)則：

設(shè)備的共享和釋放方式取決于：

對(duì)單個(gè)進(jìn)程是獨(dú)占的 ：使用最大釋放閾值。

在合作進(jìn)程之間共享 ：通過(guò) IPC 協(xié)調(diào)使用相同的池，或?qū)⒚總€(gè)進(jìn)程池設(shè)置為適當(dāng)?shù)闹担员苊馊魏我粋€(gè)進(jìn)程獨(dú)占所有設(shè)備內(nèi)存。

在未知進(jìn)程之間共享： 如果已知，請(qǐng)將閾值設(shè)置為應(yīng)用程序的工作集大小。否則，在使用非零值之前，請(qǐng)將其保留為零，并使用探查器確定分配性能是否是瓶頸。

用 cudaMallocAsync 替換 cudaMalloc 的指南：

確保所有內(nèi)存訪問(wèn)都是在流順序分配之后排序的。

如果需要對(duì)等訪問(wèn)，請(qǐng)使用 cudaMemPoolSetAccess ，因?yàn)?cudaEnablePeerAccess 和 cudaDisablePeerAccesss 對(duì)池內(nèi)存沒(méi)有影響。

與 cudaMalloc 分配不同， cudaDeviceReset 不會(huì)隱式釋放池內(nèi)存，因此必須顯式釋放。

如果使用 cudaFree 釋放，請(qǐng)確保在釋放之前通過(guò)適當(dāng)?shù)耐酵瓿伤性L問(wèn)，因?yàn)樵谶@種情況下沒(méi)有隱式同步。依賴隱式同步的任何后續(xù)代碼也可能需要更新。

如果內(nèi)存通過(guò) IPC 與另一個(gè)進(jìn)程共享，請(qǐng)從顯式創(chuàng)建的支持 IPC 的池中進(jìn)行分配，并刪除該指針對(duì) cudaIpcGetMemHandle 、 cudaIpcOpenMemHandle 和 cudaIpcCloseMemHandle 的所有引用。

如果該內(nèi)存必須與 GPU 直接 RDMA 一起使用，請(qǐng)暫時(shí)繼續(xù)使用 cudaMalloc ，因?yàn)橥ㄟ^(guò) cudaMallocAsync 分配的內(nèi)存目前不支持它。 CUDA 打算在將來(lái)支持它。

與使用 cudaMalloc 分配的內(nèi)存不同，使用 cudaMallocAsync 分配的內(nèi)存與 CUDA 上下文不關(guān)聯(lián)。這有以下影響：

使用屬性 CU_POINTER_ATTRIBUTE_CONTEXT 調(diào)用 cuPointerGetAttribute 會(huì)為上下文返回 null 。

當(dāng)使用至少一個(gè)使用 cudaMallocAsync 分配的源或目標(biāo)指針調(diào)用 cudaMemcpy 時(shí)，必須可以從調(diào)用線程的當(dāng)前上下文/設(shè)備訪問(wèn)該內(nèi)存。如果無(wú)法從該上下文或設(shè)備訪問(wèn)，請(qǐng)改用 cudaMemcpyPeer 。

將 cudaFree 替換為 cudaFree 的指南

確保所有內(nèi)存訪問(wèn)都是在按流排序的釋放之前排序的。

在下一次同步操作之前，可能無(wú)法將內(nèi)存釋放回系統(tǒng)。如果釋放閾值設(shè)置為非零值，則在顯式修剪相應(yīng)的池之前，可能無(wú)法將內(nèi)存釋放回系統(tǒng)。

與 cudaFree 不同， cudaFreeAsync 不會(huì)隱式同步設(shè)備。任何依賴此隱式同步的代碼都必須更新為顯式同步。

結(jié)論

CUDA 11 。 2 中添加的流式有序分配器以及 cudaMallocAsync 和 cudaFreeAsync API 函數(shù)通過(guò)將內(nèi)存分配和釋放作為流式有序操作引入 CUDA 流編程模型，擴(kuò)展了 CUDA 流編程模型。這使得分配的范圍能夠限定到內(nèi)核，內(nèi)核使用它們，同時(shí)避免了傳統(tǒng) cudaMalloc/cudaFree 可能發(fā)生的昂貴的設(shè)備范圍同步。

此外，這些 API 函數(shù)在 CUDA 中添加了內(nèi)存池的概念，從而實(shí)現(xiàn)了內(nèi)存的重用，從而避免了代價(jià)高昂的系統(tǒng)調(diào)用并提高了性能。使用指南 MIG 評(píng)估您現(xiàn)有的代碼，并查看您的應(yīng)用程序性能有多大改進(jìn)！

關(guān)于作者

Vivek Kini 是 NVIDIA 的高級(jí)系統(tǒng)軟件工程師。他致力于 CUDA 驅(qū)動(dòng)程序，特別關(guān)注內(nèi)存管理功能。他旨在簡(jiǎn)化 CUDA 應(yīng)用程序的內(nèi)存管理，而不犧牲它們所需的性能。

Jake Hemstad 是一個(gè)高級(jí)開(kāi)發(fā)工程師 NVIDIA ，他在開(kāi)發(fā)高性能 CUDA C ++軟件加速數(shù)據(jù)分析。他同樣關(guān)心開(kāi)發(fā)高質(zhì)量的軟件，正如他實(shí)現(xiàn)最佳的 GPU 性能一樣，也是現(xiàn)代 C ++設(shè)計(jì)的倡導(dǎo)者。在 NVIDIA 之前，他參加了明尼蘇達(dá)大學(xué)的研究生院，在那里他與桑迪亞國(guó)家實(shí)驗(yàn)室在任務(wù)并行 HPC 運(yùn)行時(shí)間和稀疏線性求解器上工作。

審核編輯：郭婷