天天看點

CUDA學習(九十三)

簡化GPU程式設計:

存儲空間的統一意味着主機和裝置之間不再需要顯式存儲器傳輸。 在托管記憶體空間中建立的任何配置設定都會自動遷移到需要的位置。

程式以兩種方式之一配置設定托管記憶體:通過cudaMallocManaged()例程,它在語義上類似于cudaMalloc(); 或者通過定義一個全局的<code>__managed__</code>變量,這個變量在語義上類似于<code>__device__</code>變量。 這些檔案的精确定義見本文後面。

在具有計算能力6.x的裝置的支援平台上,Unified Memory将使應用程式能夠使用預設系統配置設定器配置設定和共享資料。 這允許GPU在不使用特殊配置設定器的情況下通路整個系統虛拟記憶體。

以下代碼示例說明了如何使用托管記憶體可以更改寫入主機代碼的方式。 首先,一個沒有統一記憶體利益的簡單程式:

第一個示例将GPU上的兩個數字與每個線程ID結合在一起,并将數值傳回到數組中。 如果沒有托管記憶體,則需要用于傳回值的主機和裝置端存儲(在本例中為host_ret和ret),因為兩者之間使用cudaMemcpy()進行顯式拷貝。

将此與程式的統一記憶體版本進行比較,該版本允許從主機直接通路GPU資料。 注意cudaMallocManaged()例程,它傳回一個有效來自主機和裝置代碼的指針。 這允許在沒有單獨的host_ret副本的情況下使用ret,極大地簡化和減小了程式的大小。

最後,語言內建允許直接引用GPU聲明的__managed__變量,并在使用全局變量時進一步簡化程式。

請注意,缺少顯式的cudaMemcpy()指令以及傳回數組ret在CPU和GPU上都可見的事實。

對主機和裝置之間的同步值得評論。 請注意,在非托管示例中,同步cudaMemcpy()例程用于同步核心(即等待它完成運作)并将資料傳輸到主機。 統一記憶體示例不會調用cudaMemcpy(),是以在主機程式可以安全地使用GPU輸出之前需要明确的cudaDeviceSynchronize()。

這裡的另一種方法是設定環境變量CUDA_LAUNCH_BLOCKING = 1,確定所有核心的啟動都是同步完成的。 這通過消除所有顯式同步來簡化代碼,但顯然對整個執行行為具有更廣泛的影響。

CUDA學習(九十三)