CUDA 6.0以後，出現了：統一記憶體，unified memory，其具有以下特性：
1 其宣告有以下兩種方式：
a 使用關鍵字managed，可選的和device,一起使用，比如：device managed int ret[1000];
b 使用函式cudaMallocManaged():
int *ret;
cudaMallocManaged(&ret,1000*sizeof(int));
cudaFree(ret);

2 使用上面兩種方式宣告的記憶體，能夠被cpu和gpu同時訪問，不需要顯式的在cpu和gpu之間使用cudaMemcpy()在cpu和gpu之間進行記憶體傳輸
3 cuda的managed記憶體和其zero-copy memory之間的區別在於：
a zero-copy分配的是CPU中的固定頁記憶體(pinned memory in cpu system )，這種方式下，根據從哪裡索引這個記憶體，其速度將更快或者更慢，零複製記憶體
實際上是一種特殊形式的記憶體對映，它允許你將主機記憶體直接對映到GPU記憶體空間上。
b 統一記憶體，將記憶體與執行空間分開，這樣所有的資料訪問都很快

cuda中鎖頁記憶體和零複製記憶體
鎖頁記憶體允許GPU上的DMA控制器請求主機傳輸，而不需要CPU主機處理器的參與
CPU仍然可以訪問上述鎖頁記憶體，但是此記憶體是不能移動或換頁到磁碟上的
在GPU上分配的記憶體預設都是鎖頁記憶體，這只是因為GPU不支援將記憶體交換到磁碟上
在主機上分配鎖頁記憶體有以下兩種方式：
a 使用特殊的cudaHostAlloc函式，對用的釋放記憶體使用cudaFreeHost函式進行記憶體釋放
b 使用常規的malloc函式，然後將其註冊為(cudaHostRegister)鎖頁記憶體,註冊為鎖頁記憶體只是設定一些內部標誌位以確保記憶體不被換出，並告訴CUDA驅動程式，該記憶體為鎖頁記憶體，可以直接使用而不需要使用臨時緩衝區

使用鎖頁記憶體需要注意以下幾點：
a 不能分配太多，太多的話會降低系統整體效能
b 鎖頁記憶體和視訊記憶體之間的拷貝速度是6G/s，普通的記憶體和視訊記憶體之間的拷貝速度是3G/s（視訊記憶體之間的拷貝速度是30G/s,CPU之間的速度是10G/s）
c 使用cudaHostAlloc函式分配記憶體，其內的內容需要從普通記憶體拷貝到鎖頁記憶體中，因此會存在：這種拷貝會帶來額外的CPU記憶體拷貝時間開銷，CPU需要把資料從可分頁記憶體拷貝到鎖頁，但是採用cudaHostRegister把普通記憶體改為鎖頁記憶體，則不會帶來額外的cpu記憶體拷貝時間開銷，因為cudaHostAlloc的做法是先分配鎖頁記憶體，這時裡面是沒有資料的，那麼需要將一般的記憶體拷貝過來，而對於cudaHostRegister記憶體，他是之間就使用malloc分配好的，cudaHostRegister只是設定一些內部標誌位以確保其不被換出，相當於只是更改了一些標誌位，就不存在前面說的資料拷貝
d 在某些裝置上，裝置儲存器和主機鎖頁儲存器之間的資料拷貝和核心函式可以併發執行
e 在某些裝置上，可以將主機的鎖頁記憶體對映到裝置地址空間，減少主機和裝置之間的資料拷貝，要訪問資料的時候不是像上面那那樣將資料拷貝過來，而是直接通過主機匯流排到主機上訪問 ，使用cudaHostAlloc分配時傳入cudaHostAllocMapped，或者使用cudaHostRegister時傳入cudaHostRegisterMapped標籤
f 預設情況下，鎖頁記憶體是可以快取的。在使用cudaHostAlloc分配時傳入cudaHostAllocWriteCombined標籤，將其標定為寫結合，這意味著該記憶體沒有一級二級快取，這樣有利用主機寫該記憶體，而如果主機讀取的話，速度將會極其慢，所以這種情況下的記憶體應當只用於那些主機只寫的儲存器

鎖頁記憶體分配的記憶體，也有對應的三種形式：

cudaMemcpyToSymbol可以將資料從host拷貝到global，cudaMemcpy也是從host到>global，這種情況下二個函式有什麼區別嗎？
和各位大佬討論一下後，和大家分享一下~

cudaMemcpyToSymbol也有將資料從host拷貝到global的功能，以前只用過這個函式拷貝constant memory。拷貝方式的不同是由目的記憶體申請的方式決定的。申請的是device記憶體，cudaMemcpyToSymbol拷貝就是從host拷貝到global。申請的是constant記憶體，cudaMemcpyToSymbol拷貝就是從host拷貝到constant memory。

CUDA中與記憶體相關的一些函式：
cudaMemcpyToSymbol(主要用於將資料從host拷貝到device的constant memory，但也可以用於將資料從host拷貝到device的global區)

cudamalloc分配，對應使用cudaFree釋放，使用cudaMemcpy在裝置和主機之間進行資料傳輸

cudaHostAlloc(核心是，在主機上分配鎖頁記憶體，對應的釋放函式cudaFreeHost，opencv中的資料結構CudaMem是使用的該函式進行記憶體分配，其優勢有：
d 在某些裝置上，裝置儲存器和主機鎖頁儲存器之間的資料拷貝和核心函式可以併發執行
e 在某些裝置上，可以將主機的鎖頁記憶體對映到裝置地址空間，減少主機和裝置之間的資料拷貝，要訪問資料的時候不是像上面那那樣將資料拷貝過來，而是直接通過主機匯流排到主機上訪問 ，使用cudaHostAlloc分配時傳入cudaHostAllocMapped，或者使用cudaHostRegister時傳入cudaHostRegisterMapped標籤
f 預設情況下，鎖頁記憶體是可以快取的。在使用cudaHostAlloc分配時傳入cudaHostAllocWriteCombined標籤，將其標定為寫結合，這意味著該記憶體沒有一級二級快取，這樣有利用主機寫該記憶體，而如果主機讀取的話，速度將會極其慢，所以這種情況下的記憶體應當只用於那些主機只寫的儲存器)

這裡講下個人對cudaMalloc和cudaHostAlloc的理解：
cudaHostAlloc個人認為主要是在主機和裝置之間進行資料互動的時候使用效率會比cudaMalloc高，另外需要注意的是，cudaHostAlloc分配的是主機上的記憶體，裝置通過某種方式進行訪問，cudaMalloc分配的是裝置上的記憶體，其需要把相應的內容使用cudaMemcpy將資料從主機拷貝到裝置，cudaHostAlloc函式的第三個引數有三個可選列舉，分別是：
cudaHostAllocDefault：
cudaHostAllocMapped(對應到cudaHostRegister函式是cudaHostRegisterMapped)：在該標誌位下，將會分配一塊可對映到裝置地址空間的分頁鎖定主機儲存器。這塊儲存器有兩個地址，一個是主機儲存器上的，一個是裝置儲存器上的，主機指標是cudaHostAlloc返回，裝置指標通過cudaHostGetDevicePointer函式檢索到，可以使用這個裝置指標在核心中訪問這塊儲存器，但是也有例外：
在支援統一虛擬地址空間的裝置上，主機通過cudaHostAlloc和cudaMalloc分配的記憶體使用，單個虛擬地址空間，指標指向哪個儲存空間（主機儲存器或任意一個裝置儲存器），可以通過cudaPointerGetAttributes確定：
當在使用統一地址空間的裝置間複製儲存器時，cudaMemcpy*中的cudaMemcpyKind引數沒有作用，可設定成cudaMemcpyDefault
通過cudaHostAlloc分配的儲存器預設在使用統一地址空間的裝置間是可分享的，cudaHostAlloc返回的指標可被在這些裝置上的核心直接使用，而不需要cudaHostGetDevicePointer獲得裝置指標

cudaHostAllocWriteCombined：主要是去掉一級二級快取，這樣主機寫資料的時候會很快，而主機讀取的時候極慢，所以用於就是用處主機只寫不讀

cudaHostRegister：是得到鎖頁記憶體的另一種方式，和cudaHostAlloc的區別是，cudaHostRegister是將現有的記憶體標記成鎖頁記憶體
cudaMallocPitch(opencv中的GpuMat使用的該函式對影象進行記憶體分配)

注意的是，位於同一個Block中的執行緒才能實現通訊，不同Block中的執行緒不能通過共享記憶體、同步進行通訊，而應採用原子操作或主機介入。