低層級(jí)GPU虛擬內(nèi)存管理引論

Introducing Low-Level GPU Virtual Memory Management

CUDA應(yīng)用程序越來越需要盡可能快速高效地管理內(nèi)存。在CUDA 10.2之前，開發(fā)人員可用的選項(xiàng)數(shù)量僅限于CUDA提供的類似malloc的抽象。

CUDA10.2為虛擬內(nèi)存管理引入了一組新的API函數(shù)，使您能夠構(gòu)建更高效的動(dòng)態(tài)數(shù)據(jù)結(jié)構(gòu)，并更好地控制應(yīng)用程序中的GPU內(nèi)存使用。在這篇文章中，我們將解釋如何使用新的API函數(shù)并瀏覽一些實(shí)際的應(yīng)用程序用例。

在很多應(yīng)用程序中，很難猜測初始分配應(yīng)該有多大。您需要一個(gè)更大的分配，但是您不能承擔(dān)從GPU通過一個(gè)專門的動(dòng)態(tài)數(shù)據(jù)結(jié)構(gòu)來跟蹤指針的性能和開發(fā)成本。您真正想要的是在需要更多內(nèi)存時(shí)增加分配，同時(shí)保持您一直擁有的連續(xù)地址范圍。如果你曾經(jīng)使用過LIB的RealCoc函數(shù)，或者C++的STD:：vector，你可能自己碰到這個(gè)問題。

Growing allocations

看看下面的簡單C++類，它描述了一個(gè)可以擴(kuò)展的向量：

class Vector {

private:

void *d_p;

size_t alloc_sz, reserve_sz;

public:

Vector() : d_p(NULL), alloc_sz(0), reserve_sz(0) {}

// Reserves some extra space in
order to speed up grow()

CUresult reserve(size_t new_sz);

// Actually commits num bytes
of additional memory

CUresult grow(size_t new_sz);

// Frees up all the associated
resources.

~Vector();

};

在CUDA 10.2之前，在CUDA中實(shí)現(xiàn)這個(gè)概念的唯一方法是使用cudamaloc、cudaFree和cudaMemcpy，或者使用cudamalocmanaged和cudaPrefetchAsync來提交需要的內(nèi)存。

CUresult Vector::reserve(size_t new_sz)
{
if (new_sz > reserve_sz)
{
void *new_ptr = nullptr;#ifndef USE_MANAGED_MEMORY
cudaMalloc(&new_ptr, new_sz);
#else
cudaMallocManaged(&new_ptr, new_sz);
#endif
cudaMemcpy(new_ptr, d_p, alloc_sz);
cudaFree(d_p);
d_p = new_ptr;
reserve_sz = new_sz;
}
}
CUresult Vector::grow(size_t new_sz)
{
Vector::reserve(alloc_sz + new_sz);
#ifdef
USE_MANAGED_MEMORY
cudaPrefetchAsync(d_p + alloc_sz, num, dev);
#endif
alloc_sz += new_sz;
}
Vector::~Vector()
{
if (d_p)
cudaFree(d_p);
}

雖然實(shí)現(xiàn)相當(dāng)簡單，但有許多性能影響。

cudaMalloc函數(shù)分配的資源超過了增加分配所需的資源。要增長，您需要保留舊的分配，并分配一個(gè)新的分配，為舊的分配留出足夠的空間和額外的空間，這將大大減少您的增長量。如果設(shè)備只有2 GiB的內(nèi)存，并且您已經(jīng)有1 GiB的向量，則不能將其增大，因?yàn)槟枰? GiB加上您需要的增長量。有效地，你不能增長一個(gè)向量大于一半的GPU內(nèi)存。

每個(gè)分配必須映射到所有對(duì)等上下文，即使它從未在這些對(duì)等上下文中使用過。

cudammcpy調(diào)用為不斷增長的請(qǐng)求增加了延遲，并使用寶貴的內(nèi)存帶寬來復(fù)制數(shù)據(jù)。這樣的帶寬可以更好地用在其他地方。

cudaFree調(diào)用在繼續(xù)之前等待當(dāng)前上下文上的所有掛起工作（以及所有對(duì)等上下文）。

使用托管內(nèi)存解決了其中一些問題，您將在本文后面看到。不幸的是，使用托管內(nèi)存會(huì)增加一些兼容性問題，這些問題可能不適合所有應(yīng)用程序。

按需頁面遷移并非在所有平臺(tái)上都可用（尤其是在Windows和Tegra移動(dòng)平臺(tái)上）。在這些平臺(tái)上，使用cudamalocmanaged保留一個(gè)VA，然后根據(jù)需要提交它不是一個(gè)選項(xiàng)。
cudamalocmanaged內(nèi)存不能與CUDA進(jìn)程間通信（cudaIpc*）函數(shù)一起使用。要與其他進(jìn)程通信，必須將數(shù)據(jù)復(fù)制到可共享的cudamaloc內(nèi)存中，有效地復(fù)制數(shù)據(jù)以繞過此限制。
cudamalocmanaged內(nèi)存不能與圖形互操作函數(shù)一起使用。在圖形API（如DirectX、OpenGL或Vulkan）中使用此數(shù)據(jù)之前，必須將數(shù)據(jù)復(fù)制到已注冊(cè)的圖形資源。

新的CUDA虛擬內(nèi)存管理功能是低級(jí)的驅(qū)動(dòng)程序功能，允許您實(shí)現(xiàn)不同的分配用例，而不會(huì)出現(xiàn)前面提到的許多缺點(diǎn)。

支持各種用例的需要使得低級(jí)虛擬內(nèi)存分配與像cudamaloc這樣的高級(jí)函數(shù)有很大的不同。與單個(gè)函數(shù)不同，您將使用四個(gè)主要函數(shù)，我們將在后面的章節(jié)中更詳細(xì)地介紹這些函數(shù)：

cuMemCreate創(chuàng)建物理內(nèi)存句柄。

cuMemAddressReserve保留一個(gè)虛擬地址范圍。

cumemap將物理內(nèi)存句柄映射到虛擬地址范圍。

cuMemSetAccess將每個(gè)設(shè)備的內(nèi)存訪問權(quán)限設(shè)置為分配。

這些函數(shù)可以與cudaMalloc和cudamalocmanaged等運(yùn)行時(shí)函數(shù)同時(shí)使用，但它們需要直接從驅(qū)動(dòng)程序加載這些入口點(diǎn)。有關(guān)如何與此類驅(qū)動(dòng)程序函數(shù)交互的更多信息，請(qǐng)參閱本文中包含的示例或隨CUDA工具包分發(fā)的各種示例。下面是這些新的虛擬內(nèi)存管理功能的工作原理。

Allocating physical memory

首先，需要對(duì)物理內(nèi)存進(jìn)行操作，為此需要使用新函數(shù)cuMemCreate。此函數(shù)采用句柄cumemgenericalallocationhandle，它描述要分配的內(nèi)存的屬性，比如該內(nèi)存物理位置在哪里，或者應(yīng)該提供什么類型的可共享句柄。目前，唯一受支持的內(nèi)存類型是當(dāng)前設(shè)備上的固定設(shè)備內(nèi)存，但在將來的CUDA版本中，還會(huì)有更多的屬性。

接下來，你需要尺寸。與cuMemAlloc不同，cuMemCreate只接受與句柄所描述的內(nèi)存的粒度相匹配的大小。使用cuMemGetAllocationGranularity獲取此粒度并使用它填充請(qǐng)求的大小。現(xiàn)在，您擁有創(chuàng)建物理分配所需的所有信息，如下代碼示例所示：

size_t granularity = 0;
CUmemGenericAllocationHandle allocHandle;
CUmemAllocationProp prop = {};
prop.type = CU_MEM_ALLOCATION_TYPE_PINNED;
prop.location.type = CU_MEM_LOCATION_TYPE_DEVICE;
prop.location.id = currentDev;
cuMemGetAllocationGranularity(&granularity, &prop,CU_MEM_ALLOC_GRANULARITY_MINIMUM);
padded_size = ROUND_UP(size, granularity);cuMemCreate(&allocHandle, padded_size, &prop, 0);

您可以使用分配句柄映射分配的內(nèi)存，以便CUDA的其余部分可以訪問它，如下一節(jié)所述。您還可以將此分配句柄導(dǎo)出到可用于進(jìn)程間通信甚至圖形互操作的對(duì)象。我們將在后面的章節(jié)中回到這些用例。

Mapping memory

要使用新的CUDA虛擬內(nèi)存管理功能映射分配，必須首先從CUDA請(qǐng)求虛擬地址（VA）范圍。這類似于virtualloc或mmap的工作方式。使用CUDA，使用cuMemAddressReserve獲得合適的地址。接下來，將物理句柄映射到使用cumemap檢索的地址。

/* Reserve a virtual address range /
cuMemAddressReserve(&ptr, padded_size, 0, 0, 0);/ Map the virtual address range * to the physical allocation */
cuMemMap(ptr, padded_size, 0, allocHandle, 0);

繼續(xù)使用前面計(jì)算的填充大小。目前，CUDA不支持物理分配的映射部分，因此需要匹配大小。這在未來可能會(huì)改變。

雖然您現(xiàn)在可以嘗試從設(shè)備訪問地址，但它會(huì)生成設(shè)備故障，就像您訪問了無效內(nèi)存一樣。這是因?yàn)樾掠成涞姆峙涫冀K映射為所有設(shè)備的CU_MEM_ACCESS_FLAGS_PROT_NONE，這意味著從任何設(shè)備對(duì)該VA范圍的訪問無效并觸發(fā)錯(cuò)誤。其原因是使該內(nèi)存的映射操作可伸縮。在本文后面的“用例：可伸縮對(duì)等映射”一節(jié)中，我們將回到這一點(diǎn)。

要啟用對(duì)此內(nèi)存映射的訪問，請(qǐng)初始化訪問描述結(jié)構(gòu)并調(diào)用cuMemSetAccess，如下代碼示例所示：

CUmemAccessDesc accessDesc = {};
accessDesc.location.type = CU_MEM_LOCATION_TYPE_DEVICE;accessDesc.location.id = currentDev;
accessDesc.flags = CU_MEM_ACCESS_FLAGS_PROT_READWRITE;
cuMemSetAccess(ptr, size, &accessDesc, 1);

現(xiàn)在，您可以從當(dāng)前設(shè)備訪問[ptr，ptr+size]范圍內(nèi)的任何地址，而不會(huì)出現(xiàn)問題。

Freeing memory

當(dāng)然，到目前為止描述的所有函數(shù)都有相應(yīng)的自由函數(shù)。若要取消映射映射的VA范圍，請(qǐng)對(duì)整個(gè)VA范圍調(diào)用cummunmap，這會(huì)將VA范圍還原回cummaddressreserve之后的狀態(tài)。完成VA范圍后，cuMemAddressFree會(huì)將其返回給CUDA以用于其他用途。

最后，cuMemRelease使句柄無效，如果沒有映射引用，則將內(nèi)存的備份存儲(chǔ)釋放回操作系統(tǒng)。下面的代碼示例顯示了這種情況：

cuMemUnmap(ptr, size);
cuMemRelease(allocHandle);
cuMemAddressFree(ptr, size);

雖然我們?cè)谶@里沒有詳細(xì)介紹這些函數(shù)，但是您可以查看CUDA示例以及本文中引用的示例，了解它們是如何協(xié)同工作的。

Putting it together

本文前面的部分使用CUDA虛擬內(nèi)存管理功能介紹了cudamaloc的另一種實(shí)現(xiàn)。這些函數(shù)要詳細(xì)得多，并且需要更多關(guān)于應(yīng)用程序如何使用分配的預(yù)先知識(shí)。我們將在本文后面向您展示這種額外冗長的好處。

回到向量的例子。使用CUDA虛擬內(nèi)存管理功能，您可以將內(nèi)存提交到虛擬地址空間的區(qū)域，就像使用cudaPrefetchAsync和cudaMallocManaged一樣。另外，如果您的保留空間不足，則不需要發(fā)出cudammcpy調(diào)用，也不需要分配比原始請(qǐng)求更多的內(nèi)存。只需將您已經(jīng)擁有的分配重新映射到它們的新地址。

首先，您需要一個(gè)VA范圍來映射，這在前面已經(jīng)介紹過了。通常你已經(jīng)有了一個(gè)VA，你只想把它附加到VA上來種植它。cuMemAddressReserve函數(shù)接受一個(gè)fixeddr參數(shù)，該參數(shù)允許您提示所需的VA起始地址。如果CUDA由于任何原因不能使用這個(gè)VA，它會(huì)忽略這個(gè)提示，并嘗試以其他方式完成請(qǐng)求。這對(duì)于向量類很有用：

CUresult Vector::reserve(size_t new_sz) {// …// Try to reserve at the end of
old_ptrstatus = cuMemAddressReserve(&new_ptr, (aligned_sz - reserve_sz), 0ULL, old_ptr + reserve_sz, 0ULL);
if ((status != CUDA_SUCCESS) || (new_ptr != (old_ptr + reserve_sz)))
{ // Nope, something went wrong. You couldn’t get the address you wanted, // so fall back to the slow path.
if (new_ptr != 0ULL)
{ // Don’t leak new_ptr if you got one.
(void)cuMemAddressFree(new_ptr, (aligned_sz - reserve_sz));
} // Now reserve the new, bigger VA range. status = cuMemAddressReserve(&new_ptr, aligned_sz,0ULL, 0ULL, 0ULL); // You have a new address range reserved, so remap. // …
}

既然您有了VA范圍，就需要時(shí)間來創(chuàng)建所需的塊，映射它，并提供對(duì)它的訪問權(quán)限。存儲(chǔ)信息以供以后使用，如句柄和分配大小。

CUresult Vector::grow(size_t new_sz)
{// …// Pad the size to the correct granularity
padded_sz = ROUND_UP(new_sz - alloc_sz, chunk_sz);// Create the chunk that you need
cuMemCreate(&handle, padded_sz, &prop, 0);// Map it at the end of ptr
cuMemMap(ptr + alloc_sz, padded_sz, 0ULL, handle, 0ULL);// Set the access
cuMemSetAccess(ptr + alloc_sz, padded_sz, &accessDesc, 1ULL);// Keep track of the metadata (for later)
handles.push_back(handle);
handle_sizes.push_back(padded_sz);
}

在某些情況下，您可能無法在當(dāng)前VA范圍之后立即保留相鄰的VA。可能是另一個(gè)分配。您可以退回到釋放虛擬地址并將其重新映射到新的更大地址范圍的較慢路徑。返回Vector：：reserve并實(shí)現(xiàn)此回退路徑。

因?yàn)榫浔痛笮∈前捶峙漤樞螂[藏的，所以您只需取消映射舊的VA范圍，然后在正確的偏移量將每個(gè)句柄映射到更大的VA范圍。下面的代碼示例顯示了這種情況：

CUresult Vector::reserve(size_t new_sz) {// …// You have a new address range reserved, so remap.
CUdeviceptr ptr = new_ptr;
cuMemUnmap(d_p, alloc_sz); // And remap them to the new VA range, enabling their access
for (size_t i = 0ULL; i < handles.size(); i++) {
const size_t hdl_sz = handle_sizes[i];
cuMemMap(ptr, hdl_sz, 0ULL, handles[i], 0ULL); ptr += hdl_sz;}
cuMemSetAccess(new_ptr, new_sz, &accessDesc, 1ULL);// Free up our previous VA range
for (size_t i = 0ULL; i < va_ranges.size(); i++)
{
cuMemAddressFree(va_ranges[i].start, va_ranges[i].sz);
}

這里有一個(gè)新的CUDA虛擬內(nèi)存管理功能的矢量類的工作實(shí)現(xiàn)。

Performance results

現(xiàn)在您開始看到使用CUDA虛擬內(nèi)存管理功能的好處。雖然帶有保留的標(biāo)準(zhǔn)cumemaloc（cudamaloc）路徑是最快的，但它也是最占用內(nèi)存的路徑：它提交它保留的所有內(nèi)存，即使它不需要它。cuMemAlloc without reservation方法中的內(nèi)存使用峰值是您需要增加的額外分配。尖峰會(huì)隨著你需要增長的數(shù)量呈指數(shù)增長。

另一方面，對(duì)于帶有預(yù)保留的cumemalocmanaged版本，應(yīng)用程序分配它需要保留的1 GiB。然后它調(diào)用cummprefetchasync并在向量需要增長時(shí)進(jìn)行同步。如果沒有保留，應(yīng)用程序會(huì)像在cudaMalloc實(shí)現(xiàn)中那樣分配一個(gè)更大的緩沖區(qū)并執(zhí)行一個(gè)拷貝，但是在接觸到該分配之前，不會(huì)對(duì)其進(jìn)行分頁。

因?yàn)橹挥|及了分配的一部分（要復(fù)制到的部分），所以只需要前一個(gè)分配的大小。然后釋放舊的緩沖區(qū)，并預(yù)取未觸及的部分，確保您永遠(yuǎn)不需要超過以前的緩沖區(qū)大小。也就是說，這個(gè)方法確實(shí)會(huì)釋放一個(gè)臟的分配回操作系統(tǒng)，在預(yù)取數(shù)組的未觸及部分之后，最終會(huì)得到一個(gè)干凈的分配。

CUDA虛擬內(nèi)存管理功能與cumemalocmanaged保持著密切的同步，但是在是否可以附加到VA范圍并因此返回到前面描述的慢路徑上存在一些抖動(dòng)。即便如此，這條緩慢的路徑仍然比其他實(shí)現(xiàn)快得多。

當(dāng)您使用cuMemAddressReserve預(yù)先保留整個(gè)VA范圍，并在增長時(shí)分配新塊并將其映射到中時(shí)，您會(huì)看到您與cumemalocmanaged+reserve非常匹配，甚至在64 MiB大小調(diào)整后擴(kuò)展得更好。

由于在任何時(shí)候都不會(huì)分配比所需更多的內(nèi)存，即使是慢速重新映射，也總是低于分配的預(yù)算，就像cumemalocmanaged一樣。這兩種方法的區(qū)別在于不需要復(fù)制到新緩沖區(qū)，因此將提交內(nèi)存的需要推遲到“預(yù)取”或塊創(chuàng)建時(shí)間。

查看通過自己運(yùn)行vector_example代碼可以獲得什么樣的性能優(yōu)勢。

Application use case: Join operation in OLAP

在數(shù)據(jù)分析中可以找到不斷增長的分配器的一個(gè)重要用例。數(shù)據(jù)庫應(yīng)用程序中計(jì)算最密集的操作是連接操作。

聯(lián)接的輸出大小依賴于數(shù)據(jù)，并且事先不知道。通常，輸出大小估計(jì)器被實(shí)現(xiàn)以向探測內(nèi)核提供輸出緩沖區(qū)。然而，一個(gè)估計(jì)永遠(yuǎn)不是100%準(zhǔn)確的，所以你最終會(huì)分配比需要更多的內(nèi)存。如何將未使用的物理內(nèi)存?zhèn)鬟f回驅(qū)動(dòng)程序？對(duì)于cudaMalloc，這將需要分配一個(gè)新的緩沖區(qū)，從舊的緩沖區(qū)復(fù)制數(shù)據(jù)，并釋放舊的緩沖區(qū)，類似于前面討論的不斷增長的分配示例，如圖6所示。

Figure 6. Example pseudo-code for the probe phase of a join operation. This includes resizing the join output buffer to free up unused GPU memory.

下面是RAPIDS cuDF 0.13 join實(shí)現(xiàn)中的相應(yīng)代碼：

rmm::device_vector<size_type> left_indices;
rmm::device_vector<size_type> right_indices;…
left_indices.resize(estimated_size);
right_indices.resize(estimated_size); …
probe_hash_table<<<…>>>(…); …
join_size = write_index.value(); …
left_indices.resize(join_size);
right_indices.resize(join_size);

GPU內(nèi)存分配/釋放和內(nèi)存復(fù)制開銷隱藏在rmm：：device_vector類中。當(dāng)前實(shí)現(xiàn)的問題是，必須為輸出緩沖區(qū)提供兩倍的可用GPU內(nèi)存，并且在調(diào)整大小操作期間，可以很容易地耗盡內(nèi)存。這正是前一節(jié)中提出的向量類可以解決的問題。

可以使用前面討論過的CUDA虛擬內(nèi)存管理功能改進(jìn)rmm：：device_vector類，這將允許您支持更大的連接輸出，并通過刪除副本來提高性能。NVIDIA正在考慮將其添加到RAPIDS內(nèi)存管理器庫中。

Use case: Avoiding device synchronization on cudaFree

今天使用cudaFree會(huì)產(chǎn)生應(yīng)用程序所依賴的意外副作用：同步。當(dāng)調(diào)用cudaFree時(shí)，設(shè)備上的任何正在運(yùn)行的工作都將完成，并且調(diào)用該函數(shù)的CPU線程將被阻塞，直到完成所有這些工作。這有一些編程模型的優(yōu)點(diǎn)和缺點(diǎn)，但是直到現(xiàn)在應(yīng)用程序才真正能夠靈活地選擇不使用這種行為。

使用CUDA虛擬內(nèi)存管理功能，您不能假設(shè)在調(diào)用cummunmap或cummsetaccess期間先前的工作會(huì)同步。但是，這些功能可能在某些平臺(tái)配置上同步，例如具有Maxwell或較舊GPU架構(gòu)的系統(tǒng)。

Example

下面的示例顯示了使用cudamaloc和cudaFree進(jìn)行同步的效果。在這里，N個(gè)獨(dú)立的線程都在獨(dú)立的、非阻塞的流上啟動(dòng)工作。在理想的情況下，您應(yīng)該在GPU上觀察N個(gè)并發(fā)的spinKernel啟動(dòng)，并且每個(gè)流中很少有間隙。直觀地說，引入同時(shí)分配和釋放自己內(nèi)存的線程0不應(yīng)該有任何效果：

global void spinKernel(); // thread 1…N
while (keep_going) { spinKernel<<<1,1, stream[i]>>>();} // thread 0
for (size_t i = 0; i < 100; i++)
{ cudaMalloc(&x, 1);
cudaFree(x);
}

Optimizing

在所有CUDA虛擬內(nèi)存管理調(diào)用中，重疊量都在增加。與以前的版本相比，在修改設(shè)備的內(nèi)存布局時(shí)，GPU上沒有任何地方?jīng)]有運(yùn)行任何東西。

當(dāng)在多GPU平臺(tái)中使用cudaEnablePeerAccess啟用點(diǎn)對(duì)點(diǎn)訪問時(shí)，您還可以使用cudaFree看到這種同步效果。在這種情況下，您最終會(huì)同步每個(gè)cudaFree調(diào)用上的所有對(duì)等映射設(shè)備，即使分配僅由單個(gè)設(shè)備使用。有了新的CUDA虛擬內(nèi)存管理功能，這不再是一個(gè)問題。

Use case: Scalable peer mappings

cudaenableeracess函數(shù)用于啟用對(duì)等設(shè)備對(duì)分配的訪問，但在調(diào)用時(shí)，它會(huì)強(qiáng)制所有先前的cudamaloc分配映射到啟用的目標(biāo)對(duì)等設(shè)備。此外，cudaenableeracess還強(qiáng)制將所有未來的cudamaloc分配映射到目標(biāo)對(duì)等設(shè)備以及源設(shè)備。

為了便于開發(fā)，自動(dòng)對(duì)等映射是非常理想的，因?yàn)樗烁櫭總€(gè)設(shè)備的分配映射狀態(tài)的需要，并且避免了調(diào)試可能遇到的無效設(shè)備地址訪問問題。

不幸的是，cudaEnablePeerAccess提供的易用性可能會(huì)帶來性能上的損失，而直接讀取源代碼是不明顯的。典型的cudaMalloc調(diào)用的運(yùn)行時(shí)復(fù)雜性為O（lg（N）），其中N是先前分配的數(shù)量。這主要是由于內(nèi)部記賬。

同時(shí)，cudaenableeracessapi的運(yùn)行時(shí)復(fù)雜性大約為O（Nlg（N）），其中N是在源設(shè)備上進(jìn)行的需要映射到目標(biāo)設(shè)備的分配數(shù)。通常，這是為每個(gè)設(shè)備對(duì)調(diào)用的，以啟用完全雙向?qū)Φ仍L問，即總O（DDNlg（N）），其中D是設(shè)備數(shù)。此外，如前所述，cudamaloc現(xiàn)在必須將其分配映射到啟用對(duì)等訪問的所有設(shè)備。這意味著運(yùn)行時(shí)復(fù)雜性現(xiàn)在可以擴(kuò)展為O（D*lg（N））。

許多應(yīng)用程序通常只需要使用少量分配進(jìn)行通信，這意味著并非所有分配都必須映射到所有設(shè)備。但是，當(dāng)您只需要一些映射時(shí)，就需要支付這些額外映射的成本。

這里是新的CUDA虛擬內(nèi)存管理功能可以幫助的地方。cuMemSetAccess函數(shù)允許您將特定分配目標(biāo)設(shè)置為對(duì)等映射到特定設(shè)備集。雖然這仍然隨著訪問它的設(shè)備的數(shù)量而變化，但是只有一個(gè)設(shè)備的常見情況仍然是O（lg（N））。此外，您不再需要cudaenableeracess，讓cudamaloc調(diào)用速度更快，只在需要時(shí)支付額外映射的費(fèi)用。

要了解多GPU處理在實(shí)際中的工作方式，請(qǐng)參閱vectorAddDrvMMAP示例。

Other notable use cases

下面是一些需要考慮的其他用例：

操作系統(tǒng)本機(jī)進(jìn)程間通信

導(dǎo)出到圖形

Operating system native interprocess communication

新的CUDA虛擬內(nèi)存管理功能不支持其內(nèi)存中的傳統(tǒng)cuIpc*功能。相反，它們公開了一種新的進(jìn)程間通信機(jī)制，這種機(jī)制在每個(gè)受支持的平臺(tái)上都能更好地工作。這種新機(jī)制是基于操作特定于系統(tǒng)的句柄。在Windows上，它們是HANDLE或D3DKMT_HANDLE類型，而在基于Linux的平臺(tái)上，它們是文件描述符。

為了獲得這些特定于操作系統(tǒng)的句柄之一，引入了新函數(shù)cummexporttoshareablehandle。必須將適當(dāng)?shù)恼?qǐng)求句柄類型傳遞給cuMemCreate。默認(rèn)情況下，內(nèi)存不可導(dǎo)出，因此可共享句柄不能與默認(rèn)屬性一起使用。

將分配導(dǎo)出到特定于操作系統(tǒng)的句柄后，可以按通常的方式將句柄傳輸?shù)搅硪粋€(gè)進(jìn)程：Linux可以使用Unix域套接字，Windows可以使用DuplicateHandle。然后，另一個(gè)進(jìn)程可以使用cummimportfromshareablehandle并返回CUDA虛擬內(nèi)存管理函數(shù)可以使用的cummgenericallocationhandle值。

CUDA示例memMapIpcDrv顯示了這在實(shí)踐中的工作方式。此示例適用于支持CUDA虛擬內(nèi)存管理功能的所有Linux和Windows平臺(tái)。

Export to graphics

有些情況下，您希望CUDA應(yīng)用程序在完全無頭模式下工作，而不涉及任何圖形。其他時(shí)候，就像大型基于物理的模擬一樣，您必須以某種方式可視化結(jié)果。

在CUDA 10.2之前，應(yīng)用程序和庫必須提前知道他們想要為圖形導(dǎo)出內(nèi)存，以及他們需要使用或綁定到什么圖形庫。然后，他們必須實(shí)現(xiàn)該圖形庫的代碼來分配內(nèi)存并將其導(dǎo)入到CUDA中使用。

或者，他們可以要求應(yīng)用程序向臨時(shí)緩沖區(qū)發(fā)出memcpy調(diào)用，該緩沖區(qū)已經(jīng)注冊(cè)到應(yīng)用程序所需的圖形庫中。然而，如前所述，memcpy增加了很多延遲，浪費(fèi)了內(nèi)存帶寬。

遵循用于進(jìn)程間通信的相同代碼路徑，您還可以將操作系統(tǒng)特定的共享句柄用于其他用戶模式驅(qū)動(dòng)程序，如Vulkan或OpenGL。這允許您使用CUDA虛擬內(nèi)存管理功能分配內(nèi)存，并將該內(nèi)存導(dǎo)入所有支持操作系統(tǒng)特定句柄的圖形庫。

雖然我們還沒有公開此特定功能的示例，但您可以查看以下Vulkan和OpenGL擴(kuò)展，并將其與前面的memMapIpcDrv示例組合在一起：

·
VkMemoryAllocateInfo

·
GL_EXT_memory_objects

Conclusion

CUDA 10.2引入了新的CUDA虛擬內(nèi)存管理功能。這些新功能支持許多新的用例和性能優(yōu)化，使用CUDA的應(yīng)用程序可以利用這些新的用例和性能優(yōu)化。我們?cè)谶@篇文章中描述了其中的一些用例，但是我們有興趣了解您可以如何使用這個(gè)新特性。

看看與CUDA 10.2工具包一起發(fā)布的一些CUDA示例，或者查看本文中引用的完整代碼示例。

總結(jié)

以上是生活随笔為你收集整理的低层级GPU虚拟内存管理引论的全部內(nèi)容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網(wǎng)站內(nèi)容還不錯(cuò)，歡迎將生活随笔推薦給好友。