零復(fù)制(Zero Copy)(零拷貝內(nèi)存)
零復(fù)制是一種特殊形式的內(nèi)存映射，它允許你將主機內(nèi)存直接映射到GPU內(nèi)存空間上回怜。因此距辆，當你對GPU上的內(nèi)存解引用時，如果它是基于GPU的蔽氨，那么你就獲得了全局內(nèi)存的高速帶寬（180GB/s）藐唠。如果GPU代碼讀取一個主機映射變量，它會提交一個PCI-E讀取事務(wù)鹉究，很長時間之后宇立，主機會通過PCI-E總線返回數(shù)據(jù)。
如果程序是計算密集型自赔，那么零復(fù)制可能是一項非常有用的技術(shù)妈嘹。它節(jié)省了設(shè)備顯示傳輸?shù)臅r間。事實上绍妨，是將計算和數(shù)據(jù)傳輸操作重疊了润脸，而且無需執(zhí)行顯式的內(nèi)存管理。
實際上他去，使用零復(fù)制內(nèi)存津函，將傳輸和內(nèi)核操作分成更小的塊，然后以流水線的方式執(zhí)行它們孤页。
然而尔苦，采用零復(fù)制內(nèi)存的傳輸實際上是相當小的。PCI-E總線在兩個方向上的帶寬的總是相同的行施。由于基于PCI-E的內(nèi)存讀取存在高延遲允坚，因此實際上大多數(shù)讀操作應(yīng)該在所有寫操作之前被放入讀隊列。我們可能獲得比顯式內(nèi)存復(fù)制版本節(jié)省大量執(zhí)行時間的效果蛾号。
在論壇上的問答：
Q:我在程序中使用了零拷貝內(nèi)存稠项，程序反而變得更慢了，看書上說
“當輸入內(nèi)存和輸出內(nèi)存都只能使用一次時鲜结，那么在獨立GPU上使用零拷貝內(nèi)存將帶來性能提升”展运。
請問 “當輸入內(nèi)存和輸出內(nèi)存都只能使用一次時” 具體是什么意思活逆？我下面的做法有什么問題？
程序情況是這樣的：
在主機申請了3個 零拷貝內(nèi)存 buf1,buf2,buf3拗胜，然后通過cudaHostGetDevicePointer() 獲得這塊內(nèi)存在GPU上的有效指針dev1,dev2,dev3蔗候。其中 buf1,buf2,buf3 是從影像中讀取的數(shù)據(jù)，然后在核函數(shù)中對dev1,dev2,dev3進行處理埂软，值依然存放在dev1,dev2,dev3中锈遥，然后調(diào)用GDAL將buf1,buf2,buf3 寫出。
A:
LZ您好：
zero copy和普通的cudaMemcpy一樣也是要走pci-e總線的勘畔，只不過cudaMemcpy是一次性全部copy過去所灸，而zero copy是用的時候自動在后臺通過pci-e總線傳輸。
zero copy這樣的機制多少可以利用計算來掩蓋一些copy的時間炫七，而如果使用cudaMemcpy要實現(xiàn)類似的計算和傳輸互相掩蓋的話爬立，需要使用異步版本的cudaMemcpy函數(shù)，并使用頁鎖定內(nèi)存以及多個stream万哪。
zero copy的讀入信息是不在device端緩沖的懦尝，也就是說device端使用幾次就需要從host端走較慢的pci-e 總線讀入幾次。所以壤圃，一般建議只使用一次的數(shù)據(jù)以及少量的返回數(shù)據(jù)可以使用zero copy陵霉，其他情況建議copy到顯存使用，顯存DRAM的帶寬要比pci-e的帶寬高出一個量級伍绳。
以上是對zero copy的簡要介紹踊挠。
使用零復(fù)制的三步
1、啟用零復(fù)制
需要在任何CUDA上下文創(chuàng)建之前進行下面的調(diào)用：
//Enable host mapping to device
memoryCUDA_CALL(cudaSetDeviceFlags(cudaDeviceMapHost));
當CUDA上下文被創(chuàng)建時冲杀，驅(qū)動程序會知道它需要支持主機內(nèi)存映射效床，沒有驅(qū)動程序的支持，零復(fù)制將無法工作权谁。如果該支持在CUDA上下文創(chuàng)建之后完成剩檀，內(nèi)存也無法工作。請注意對cudaHostAlloc
這樣的函數(shù)調(diào)用旺芽，盡管在主機內(nèi)存上執(zhí)行沪猴，也仍然創(chuàng)建一個GPU上下文。
雖然大多數(shù)設(shè)備支持零復(fù)制內(nèi)存采章，但是一些早期的設(shè)備卻不支持运嗜。顯式檢查：
struct cudaDeviceProp device_prop
cudaGetDeviceProperties(&device_prop,device_num);
zero_copy_supported=device_prop.canMapHostMemory;
2、分配主機內(nèi)存
分配了主機內(nèi)存悯舟，這樣它就可以被映射到設(shè)備內(nèi)存担租。我們對cudaHostAlloc
函數(shù)使用額外的標志cudaHostAllocMapped
就可以實現(xiàn)。
//Allocate zero copy pinned
cudaHostAlloc((void**)&host_data_to_device,size_in_bytes,cudaHostAllocWriteCombined|cudaHostAllocMapped);
3抵怎、將常規(guī)的主機指針轉(zhuǎn)換成指向設(shè)備內(nèi)存空間的指針
通過cudaHostGetDevicePointer
函數(shù)：
//Conver to a GPU host
cudaHostGetDevicePointer(&dev_host_data_to_device,host_data_to_device,0);
在這個調(diào)用中奋救，我們將之前在主機內(nèi)存空間分配的host_data_to_device
轉(zhuǎn)換成GPU內(nèi)存空間的指針岭参。在GPU內(nèi)核中，只使用轉(zhuǎn)換后的指針尝艘，原始的指針只出現(xiàn)在主機執(zhí)行的代碼中演侯。因此，為了之后釋放內(nèi)存利耍，需要在主機上執(zhí)行一個操作，其他的調(diào)用保持不變：
//Free pinned
cudaFreeHost(host_data_to_device);
代碼參考
簡單的代碼盔粹，進行兩個數(shù)組之間數(shù)據(jù)的拷貝隘梨，這里只是為了說明零復(fù)制的使用，并無實際意義舷嗡。
#include <numeric>
#include <stdio.h>
#include <stdlib.h>
void checkCUDAError(const char *msg) {
cudaError_t err = cudaGetLastError();
if (cudaSuccess != err) {
fprintf(stderr, "Cuda error: %s: %s.\n", msg, cudaGetErrorString(err));
exit(EXIT_FAILURE);
}
}
__global__ void sumNum( int *data) {
int i = threadIdx.x + blockIdx.x * blockDim.x;
if(i<1000000000){
data[i]=10;
}
}
int main(void) {
size_t size = 1*1000000000 * sizeof(int);//4G
//1.啟用零復(fù)制
cudaSetDeviceFlags (cudaDeviceMapHost);
int* data;
//2.分配主機內(nèi)存
cudaHostAlloc((void**) &data, size,
cudaHostAllocWriteCombined | cudaHostAllocMapped);
checkCUDAError("cudaHostAlloc data");
memset(data, 0, 1*1000000000 * sizeof(int));
int *gpudata;
//3.將常規(guī)的主機指針轉(zhuǎn)換成指向設(shè)備內(nèi)存空間的指針
cudaHostGetDevicePointer(&gpudata, data, 0);
checkCUDAError("cudaHostGetDevicePointer");
//sumNum<<<1000000000/1024+1023, 1024>>>(gpudata);
//注意Ｖ崃浴！因為下面要打印出來測試进萄，所以要先同步數(shù)據(jù)捻脖，這個函數(shù)可以保證cpu等待gpu的kernel函數(shù)結(jié)束才往下運行。如果數(shù)據(jù)暫時用不到中鼠，可以在整體結(jié)束以后再加這句話可婶。明顯等待kernel函數(shù)結(jié)束會占用程序進行的時間。
cudaDeviceSynchronize();
for (int i = 99999999; i < 1000000000; i=i+100000000) {
printf("%d \n", data[i]);
}
//記得零拷貝的free是這個函數(shù)
cudaFreeHost(data);
return 0;
}
