從設備寫入主機並通知主機

Question

使用CUDA 5與VS 2012和功能3.5（Titan和K20）。

在內核執行的特定階段，我想將生成的數據塊發送到主機內存，並通知主機數據已準備就緒，因此主機將對其進行操作。

我不能等到內核執行結束讀取數據從設備背面，因爲：

1. 一旦計算出的數據不再與設備相關的，所以沒有點保持最後。
2. 數據大小太大而無法放在設備內存上，並等到結束。
3. 主機不應該等到內核執行結束纔開始處理數據。

你能指出我我必須採取的路徑和可能的CUDA概念和功能，我必須使用達到我的要求？簡而言之，如何寫入主機並通知主機塊數據已準備好進行主機處理？

N.B.每個線程不會與任何其他線程共享任何生成的數據，它們將獨立運行。所以，據我所知（如果我錯了，請糾正我），塊，線和經紗的概念不會影響問題。換句話說，如果他們幫助答案，我可以自由改變他們的組合。

下面是一個示例代碼，顯示我試圖做的事：

#pragma once 
#include <conio.h> 
#include <cstdio> 
#include <cuda_runtime_api.h> 

__global__ void Kernel(size_t length, float* hResult) 
{ 
    int tid = threadIdx.x + blockIdx.x * blockDim.x; 
    // Processing multiple data chunks 
    for(int i = 0;i < length;i++) 
    { 
     // Once this is assigned, I don't need it on the device anymore. 
     hResult[i + (tid * length)] = i * 100; 
    } 

} 

void main() 
{ 
    size_t length = 10; 
    size_t threads = 2; 
    float* hResult; 
    // An array that will hold all data from all threads 
    cudaMallocHost((void**)&hResult, threads * length * sizeof(float)); 
    Kernel<<<threads,1>>>(length, hResult); 
    // I DO NOT want to wait to the end and block to get the data 
    cudaError_t error = cudaDeviceSynchronize(); 
    if (error != cudaSuccess) { throw error; } 
    for(int i = 0;i < threads * length;i++) 
    { 
     printf("%f\n", hResult[i]);; 
    } 
    cudaFreeHost(hResult); 
    system("pause"); 
} 

Answer 1

在高層次上，在設備上：

• 你需要將數據寫入或者設備全局內存（先前分配爲cudaMalloc）或直接寫入主機內存（以前分配爲cudaHostAlloc）
• 您可能希望從單個線程塊執行寫入該區域的所有數據，以確保在所有的數據在以下步驟之前被寫
• 然後你會想發出threadfence()（如果您使用的設備的全局存儲器）或threadfence_system()通話（如果使用主機內存）以下步驟之前
• 接下來，您將寫入設備全局內存或主機內存中的特殊位置，我們稱之爲郵箱位置，並指定數據已準備就緒的特定值。
• 可選問題的另一個threadfence或threadfence_system叫

在主機：

• 之前啓動內核，主機需要將郵箱位置設置爲默認值。
• 啓動內核後，主機線程將需要「輪詢」郵箱位置，查找指示數據已準備好的特定值
• 一旦看到特定值，表示數據已準備好，主機就可以消費數據
• （可選）如果要重複此過程，主機可以將郵箱位置重置爲默認值。在用新數據更新數據塊之前，設備可以檢查此默認值。

請注意，即使使用上述過程，如果數據是從多個線程塊生成/創建的，仍需要默認的設備範圍的同步。可用的唯一直接設備範圍的同步是內核啓動（或特定內核的完成）。從單個線程塊複製數據只需將設備範圍內同步的需求從此特定序列中移出（到此序列之前的某個位置）。

你給的理由並不真的暗示我的代碼不能被重構來在內核上創建數據 - 在內核啓動的基礎上啓動，這將很好地解決這些問題，並且不需要上述過程也是如此。

編輯：迴應評論中的問題。 沒有一個具體的例子，關於如何重構代碼來爲每個內核調用提供一個數據塊很困難。

讓我們來看一個圖像處理的情況，我有一個30幀的視頻序列存儲在全局內存中。內核將根據某種算法處理每個幀，然後將處理後的數據提供給主機。

在您的建議中，在內核完成處理幀後，它可以向主機發出數據已準備就緒的信號，並繼續處理下一幀。問題是，如果幀是由多個線程塊處理的，那麼沒有簡單的方法可以知道所有線程塊何時完成處理該幀。設備範圍的同步障礙可能是需要的，但除了通過內核調用機制以外，它不是很方便。然而，想必這樣的內核裏，我們可能有這樣一個順序：

• 而（more_frames）
• • 過程框架
• • 信號主機
• • 增量幀指針

在重構的方法，我們將移動循環在內核之外，託管代碼：

• 而（more_frames）
• • 調用內核處理框架
• • 消耗幀
• • 增量幀指針

通過這樣做，內核標誌着知道該幀何時處理完成所需的顯式同步，並且可以對數據進行消耗。