顯式同步和邏輯GPU活動:
請注意,即使核心在上述示例中快速運作并在CPU觸及y之前完成,也需要顯式同步。 Unified Memory使用邏輯活動來确定GPU是否空閑。 這與CUDA程式設計模型保持一緻,該模型指定核心可以在啟動後随時運作,并且不能保證在主機發出同步調用之前完成。
邏輯上保證GPU完成其工作的任何函數調用都是有效的。 這包括cudaDeviceSynchronize(); cudaStreamSynchronize()和cudaStreamQuery()(隻要它傳回cudaSuccess而不是cudaErrorNotReady)指定的流是仍然在GPU上執行的唯一流; cudaEventSynchronize()和cudaEventQuery()在指定事件沒有被任何裝置工作跟蹤的情況下; 以及被記錄為與主機完全同步的cudaMemcpy()和cudaMemset()的用法。
在流之間建立的依賴關系将通過同步流或事件來推斷其他流的完成。 可以通過cudaStreamWaitEvent()建立依賴關系,或者在使用預設(NULL)流時隐式建立依賴關系。
如果沒有其他可能正在通路托管資料的流在GPU上處于活動狀态,則CPU從流回調中通路托管資料是合法的。 另外,任何裝置工作都沒有遵循的回調可以用于同步:例如,通過在回調中發信号通知條件變量; 否則,CPU通路僅在回調期間有效。
有幾點值得注意:
在GPU處于活動狀态時,CPU始終允許通路非管理的零拷貝資料。
GPU在運作任何核心時都被認為是活動的,即使該核心不使用托管資料。 如果核心可能使用資料,則禁止通路,除非裝置屬性concurrentManagedAccess為1。
除了那些适用于非管理記憶體的多GPU通路的記憶體之外,托管記憶體的并發GPU間通路沒有限制。
并發GPU核心通路托管資料沒有限制。
注意最後一點如何允許GPU核心之間的競争(race),就像目前非管理GPU記憶體的情況一樣。 如前所述,從GPU的角度來看,托管記憶體的功能與非托管記憶體相同。 以下代碼示例說明了這些要點:
使用流管理資料可視性和并行CPU + GPU通路:
到目前為止,假設對于6.x之前的SM體系結構:1)任何活動核心都可以使用任何托管記憶體,以及2)在核心處于活動狀态時使用來自CPU的托管記憶體無效。 在這裡,我們提出了一個更好的粒度控制管理記憶體的系統,該系統設計用于支援托管記憶體的所有裝置,包括concurrentManagedAccess等于0的早期體系結構。
CUDA程式設計模型提供流作為程式機制來訓示核心啟動之間的依賴性和獨立性。 啟動到同一個流中的核心保證連續執行,而啟動到不同流中的核心可以同時執行。 流描述了工作項目之間的獨立性,是以可以通過并發來提高效率。
統一記憶體通過允許CUDA程式将托管配置設定與CUDA流明确關聯,建立在流獨立模型上。 通過這種方式,程式員根據核心是否被啟動到指定的流中來訓示核心使用資料。 這使基于特定程式資料通路模式的并發機會成為可能。 控制這種行為的功能是:
cudaStreamAttachMemAsync()函數将從ptr開始的記憶體的長度位元組與指定的流相關聯。 (目前,長度必須始終為0,表示應該連接配接整個區域)。由于這種關聯,隻要流中的所有操作都已完成,統一記憶體系統就允許CPU通路此記憶體區域,無論其他流 活躍。 實際上,這限制了活動GPU對托管記憶體區域的獨占所有權,以限制每個流的活動而不是整個GPU的活動。
最重要的是,如果配置設定不與特定的流相關聯,則對所有正在運作的核心都可見,而不管它們的流如何。 這是cudaMallocManaged()配置設定或<code>__managed__</code>變量的預設可見性; 是以,在任何核心運作時,CPU可能不會觸及資料的簡單規則。
通過将配置設定與特定的流相關聯,該程式保證隻有啟動到該流中的核心才會觸及該資料。 統一記憶體系統不執行錯誤檢查:程式員有責任確定擔保得到履行。
除了允許更高的并發性之外,使用cudaStreamAttachMemAsync()可以(并且通常會)在統一記憶體系統内啟用資料傳輸優化,這可能會影響延遲和其他開銷。