CUDA學習（六十）

排序和併發：

裝置執行時的核心啟動順序遵循cuda stream排序語義。在乙個執行緒塊內，所有核心啟動到同乙個流中都會按順序執行。由於同一執行緒塊中的多個執行緒啟動到同乙個流中，流內的排序取決於該塊內的執行緒排程，這可以通過同步原語來控制，例如__syncthreads()

請注意，因為流由執行緒塊內的所有執行緒共享，所以隱式null流也被共享。如果執行緒塊中的多個執行緒啟動到隱式流中，那麼這些啟動將按順序執行。如果需要併發性，則應使用顯式命名流。

動態並行可以使併發在程式中更容易表達; 但是，裝置執行時不會在cuda執行模型中引入新的併發保證。不能保證裝置上任何數量的不同執行緒塊之間的併發執行。

缺乏併發性保證擴充套件到父執行緒塊及其子網格。當乙個父執行緒塊啟動乙個子網格時，子節點不能保證開始執行，直到父節點執行緒塊達到顯式同步點（例如cudadevicesynchronize（））。

雖然併發性通常很容易實現，但它可能會隨著裝置配置，應用程式工作負載和執行時間排程而變化。因此依賴不同執行緒塊之間的任何併發是不安全的。

裝置管理:

裝置執行時沒有多gpu支援; 裝置執行時間只能在當前正在執行的裝置上執行。但是，允許查詢系統中任何支援cuda的裝置的屬性。

記憶體模型：

父網格和子網格共享相同的全域性和常量記憶體儲存，但具有不同的本地和共享記憶體。

一致性：

全域性記憶體：

父子網格對全球記憶具有連貫性，並且父子之間的一致性保證較弱。當其記憶體檢視與父執行緒完全一致時，子網格的執行有兩點：當父網格呼叫子網格時，以及父網格中的同步api呼叫發出訊號時子網格完成時線。

在子網格呼叫之前，父執行緒中的所有全域性記憶體操作對子網格都是可見的。子網格的所有記憶體操作在父網格已完成子網格同步後對父級可見。

在以下示例中，執行child_launch的子網格只能保證在子網格啟動之前檢視對資料的修改。由於父級的執行緒0正在執行啟動，所以子級將與父級的執行緒0所見的記憶體一致。由於第乙個__syncthreads（）呼叫，子網格將看到資料[0] = 0，資料[1] = 1，...，資料[255] = 255（不帶__syncthreads（）呼叫，只有資料[0] 將保證被子網格看到）。當子網格返回時，執行緒0將保證看到它的子網格中的執行緒所做的修改。這些修改僅在第二個__syncthreads（）呼叫後才可用於父網格的其他執行緒：

__global__ void child_launch(int *data) 
__global__ void parent_launch(int *data) 
__syncthreads();
}void host_launch(int *data)

零拷貝記憶體:

零拷貝系統記憶體對全域性記憶體具有相同的一致性和一致性保證，並遵循上面詳述的語義。核心可能不會分配或釋放零拷貝記憶體，但可能使用指向從主機程式傳入的零拷貝的指標。

常量記憶體：

常量是不可改變的，並且可能不會從裝置中修改，即使在父級和子級啟動之間也是如此。也就是說，必須在啟動之前從主機設定所有__constant__變數的值。所有子核心都從它們各自的父級中自動繼承常量記憶體。

從核心執行緒中獲取常量記憶體物件的位址與所有cuda程式具有相同的語義，並且自然支援將該指標從父物件傳遞給子物件或從子物件傳遞給父物件。

共享和本地記憶體：

共享和本地記憶體分別對於執行緒塊或執行緒是私有的，並且在父級和子級之間不可見或不一致。如果這些位置中的某個位置中的物件被引用到其所屬的範圍之外並且可能導致錯誤，則行為未定義。

nvidia編譯器會嘗試警告它是否可以檢測到指向本地或共享記憶體的指標作為核心啟動引數傳遞。在執行時，程式設計師可以使用__isglobal（）內部函式確定指標是否引用全域性記憶體，因此可以安全地傳遞給子啟動。

請注意，對cudamemcpy async（）或cudamemset async（）的呼叫可能會在裝置上呼叫新的子核心以保留流語義。因此，將共享或本地記憶體指標傳遞給這些api是非法的，並且會返回錯誤。

本地記憶體:

本地記憶體是執行執行緒的私有儲存空間，並且在該執行緒之外不可見。啟動子核心時，將指標傳遞給本地記憶體作為啟動引數是非法的。取消引用來自子級的這種本地儲存器指標的結果將是未定義的。

例如，以下是非法的，如果x_array被child_launch訪問，則具有未定義的行為：

int x_array[10]; // creates x_array in parent's local memory
child_launch<<< 1, 1 >>>(x_array);

程式設計師有時很難意識到編譯器將變數放入本地記憶體的時間。作為一般規則，傳遞給子核心的所有儲存應該從全域性記憶體堆中明確分配，或者使用cudamalloc（），new（）或者在全域性範圍內宣告__device__儲存。例如：

// correct - "value" is global storage
__device__ int value;
__device__ void x() 
// invalid - "value" is local storage
__device__ void y()

紋理記憶體：

寫入紋理對映的全域性記憶體區域與紋理訪問不相關。在呼叫子網格和子網格完成時，強制執行紋理記憶體的一致性。這意味著在子核心啟動之前寫入記憶體會反映在子級的紋理記憶體訪問中。同樣，由孩子寫入記憶體將反映在父級的紋理記憶體訪問中，但僅在父級和子級完成後進行同步之後。父級和子級同時訪問可能會導致資料不一致。

CUDA學習（六十）

CUDA學習（六十五）

CUDA學習（六十二）

CUDA學習（六十七）

CUDA學習（六十）

CUDA學習（六十五）

CUDA學習（六十二）

CUDA學習（六十七）

相關推薦