cuda學習筆記五 cuda stream及 unified memory使用問題
cuda通過多個stream可以降低host到Device的資料傳輸延時，這個沒問題。但是通過stream傳輸就需要通過cudaHostAlloc等重新分配記憶體，那麼這時候就有一個問題，就是這個記憶體需要重新賦值，問題就在於很有可能這段重新賦值的時間會超出接受的範圍。如果用cudaMemcpy就可以直接用原始資料往Device上傳輸了。
本人在實際專案中遇到的問題，印象非常深刻，測試環境是NVIDIA TX2嵌入式板子。因為資料量比較大，原始程式測試通過cudaMemcpy拷貝資料需要20~40ms，那麼很自然就想到用stream。但是現在通過stream，非同步拷貝的時間確實少了，但是時間都消耗在瞭如下所示的第三步，也就是Host端記憶體的初始化上了，而且時間比cudaMemcpy還要多。
cudaHostAlloc((void **)&h_a, sizeof(float)*N, cudaHostAllocDefault);
cudaMalloc((void **)&d_a, sizeof(float)*N);
/**step3**/
for(int i=0; i<N; i++)
{
h_a[i] = src_a[i];
}
/**step3**/
cudaMemcpyAsync(d_a, h_a, sizeof(float)*N/nstreams, cudaMemcpyHostToDevice, stream[i]);
可以說為了遮蔽這段cudaMemcpy資料傳輸時間，各種方法都嘗試了，最終都是不盡如人意。後來求助於gpu world和nvidia的開發者論壇。在屠戮人神的提示下，
（1）如果確定是cudaMemcpy本身消耗了20-40ms，而不是CUDA Runtime初始化之類的因素。可以考慮在TX2上使用managed memory (Unified memory). 一定情況下可以規避傳輸。因為你的TX2本身並不存在單獨的視訊記憶體或者記憶體，只有實質的一體的儲存器，所以常規情況下的cudaMemcpy實質上是在做無用功。
（2）可以對第二次後或者以後的cudaMemcpy進行計時以測試到較為更實際的傳輸時間（不是說你之前的資料是告訴我是錯誤的， 只是為了以防萬一）。確定時間正確的話可以無視這條，直接考慮上面（1）條。
（3）TX2上使用managed memory不建議同時進行CPU和GPU訪問。
我繼續做了相關測試。其實我之前想到過使用cudaMallocManaged，但是當時傳的引數是cudaMemAttachGlobal，導致cpu端初始化的時候存在同樣的問題。在cpu端初始化時用cudaMallocManaged加cudaMemAttachHost引數，然後cpu端計算對這段記憶體賦值，在gpu核函式呼叫之前用cudaStreamAttachMemAsync( stream, gPointI+iRation*i, sizeof(float)*iRation, cudaMemAttachGlobal );將這段記憶體轉換為gpu device可以訪問的。效果確實不錯。