2012-12-06 93 views
4

我有一些問題。關於cudaMemcpyAsync功能

最近我正在通過使用CUDA製作程序。

在我的節目,有上的std ::地圖(字符串,向量(INT))編程主機一個大的數據。

通過使用這些DATAS一些矢量(INT)被複制到GPU的全局存儲器和GPU

處理之後,一些結果在GPU上產生,並且這些結果被複制到CPU處理。

這些都是我的節目時間表。

  1. cudaMemcpy(...,cudaMemcpyHostToDevice)
  2. 內核函數
  3. cudaMemcpy(內核函數只可在必要時數據被複制到GPU的全局存儲器完成)(...,cudaMemcpyDeviceToHost)
  4. 重複1〜3次1000次(另一個數據(矢量))

但我想減少處理時間。

所以我決定用cudaMemcpyAsync功能在我的計劃。

搜索一些文件和網頁後,我認識到,使用具有被複制到GPU的全局內存數據必須被分配爲固定內存cudaMemcpyAsync功能主機內存。

但我的程序正在使用的std ::地圖,所以我不能做這樣的std ::地圖數據固定內存。

因此而不是使用這個的,我做了一個緩衝器陣列類型的固定內存和該緩衝區可以隨時處理複製向量的所有情況。

最後,我的程序像這樣工作。

  1. 的memcpy(從複製的std ::地圖數據使用循環緩衝,直到整個數據被複制到緩衝器)只有
  2. cudaMemcpyAsync(...,cudaMemcpyHostToDevice)
  3. 內核(內核函數可被執行時整個數據被複制到GPU全局存儲器)
  4. cudaMemcpyAsync(...,cudaMemcpyDeviceToHost)
  5. 重複1〜1000次4steps(爲另一數據(矢量))

而且我的程序變得比以前的案例快得多。

但問題(我的好奇心)是在這一點上。

我試圖以類似的方式製作另一個程序。

  1. 的memcpy(從標準::地圖拷貝數據僅緩衝用於一個向量)
  2. cudaMemcpyAsync(...,cudaMemcpyHostToDevice)
  3. 循環1〜2,直到整個數據被複制到GPU全局存儲器
  4. 內核(核函數僅可以在必要時數據被複制到GPU全局存儲器)
  5. cudaMemcpyAsync(執行...,cudaMemcpyDeviceToHost )
  6. 重複1〜5個步驟1000倍(對於另一個數據(矢量))

此方法出來爲比上面討論的方法快約10%。

但我不知道爲什麼。

我認爲cudaMemcpyAsync只能與內核函數重疊。

但我的情況我認爲不是。而不是它看起來可以重疊在cudaMemcpyAsync函數之間。

對不起,我很長的問題,但我真的想知道爲什麼。

有人可以教或向我解釋什麼是確切的設施「cudaMemcpyAsync」和什麼功能可以與「cudaMemcpyAsync」重疊?

+0

首先讓我們知道你在哪個SDK工作以及什麼卡(Tesla,Fermi,Kepler),因爲不同的架構有不同的asyncCopy/kernel overlaping –

回答

9

cudaMemcpyAsync(以及內核活動)的複製活動可以與任何主機代碼重疊。此外,從設備拷貝數據(通過cudaMemcpyAsync)可以與內核活動重疊。所有3種活動:主機活動,數據複製活動和內核活動可以彼此異步完成,並且可以相互重疊。正如你所看到和演示的那樣,主機活動和數據拷貝或內核活動可以以相對直接的方式相互重疊:內核啓動會立即返回主機,就像cudaMemcpyAsync一樣。但是,爲了在數據複製和內核活動之間獲得最佳重疊機會,有必要使用一些額外的概念。爲獲得最佳重疊機會,我們需要:

  1. 主機內存緩衝區被固定,例如,通過cudaHostAlloc()CUDA的
  2. 使用流分離不同類型的活動(數據複製和內核計算)cudaMemcpyAsync(而不是cudaMemcpy)的
  3. 使用

當然你的工作也需要被打破了以可分離的方式。這通常意味着如果你的內核正在執行一個特定的功能,你可能需要對這個內核進行多次調用,這樣每次調用都可以在一個單獨的數據上進行。例如,這允許我們在數據塊A的第一個內核調用工作時將數據塊B複製到設備。這樣做,我們有機會與數據塊A的內核處理重疊數據塊B的副本

與cudaMemcpyAsync的主要區別(相對於cudaMemcpy)是:

  1. 它可以(需要流參數)
  2. 通常,它將控制權立即返回給主機(就像內核調用一樣),而不是等待數據複製完成。

項目1是必要的功能,因此數據複製可以與內核計算重疊。第2項是必要的功能,以便數據複製可以與主機活動重疊。

雖然拷貝/計算重疊的概念是非常簡單的,在實踐中實現需要一些工作。有關其他參考資料,請參閱:

  1. Overlap copy/compute section CUDA最佳實踐指南。
  2. 示例代碼顯示basic implementation of copy/compute overlap
  3. 示例代碼顯示完整multi/concurrent kernel copy/compute overlap scenario

請注意,以上討論的一部分是基於具有計算能力2.0或更高的設備(例如併發內核)。此外,不同的設備可以具有一個或2個拷貝引擎,這意味着同時複製到設備並從設備複製只可能在某些設備上。

+0

感謝Robert。你的答案是我想要的。在我的程序中,我也在使用流。就我而言,我的代碼幾乎重複了1000次這些步驟,所以我使用了1000個流。但我不知道在我的情況下使用1000流是最好的方式或不是。使用越來越多的流是更好嗎? – Umbrella

+0

如果你還沒有遇到麻煩的是,那麼我想你沒有打[流限制(http://stackoverflow.com/questions/3565793/max-number-of-streams-in-cuda)。但通常只需要創建3個或最多4個流,以安排複製和計算重疊。您只需重複使用相同的流ID,因爲您正在處理數據塊。如果您正在使用流ID然後將其退出,那麼您的方法沒有問題。但是,你只會從大約16個同時活動或更少,而不是數千個併發。 –

+0

我可以得到一個網站鏈接或上面說明您的解釋的教科書的名稱?因爲我正在寫一篇文章,所以我需要它作爲參考。無論如何,你的建議對我的工作非常有幫助。非常非常感謝你。 – Umbrella