從其他CUDA流中讀取更新的內存

Question

我試圖在一個內核函數中設置一個標誌，並在另一個內核函數中讀取它。基本上，我試圖做到以下幾點。

#include <iostream>                
#include <cuda.h>                 
#include <cuda_runtime.h>               

#define FLAGCLEAR 0                
#define FLAGSET 1                

using namespace std;                

__global__ void set_flag(int *flag)            
{                     
    *flag = FLAGSET;                

    // Wait for flag to reset.             
    while (*flag == FLAGSET);              
}                     

__global__ void read_flag(int *flag)            
{                     
    // wait for the flag to set.             
    while (*flag != FLAGSET);              

    // Clear it for next time.             
    *flag = FLAGCLEAR;               
}                     

int main(void)                 
{                     
    // Setup memory for flag              
    int *flag;                 
    cudaMalloc(&flag, sizeof(int));            

    // Setup streams                
    cudaStream_t stream0, stream1;            
    cudaStreamCreate(&stream0);             
    cudaStreamCreate(&stream1);             

    // Print something to let me know that we started.       
    cout << "Starting the flagging" << endl;          

    // do the flag test               
    set_flag <<<1,1,0,stream0>>>(flag);           
    read_flag <<<1,1,0,stream1>>>(flag);           

    // Wait for the streams              
    cudaDeviceSynchronize();              

    // Getting here is a painful process! 
    cout << "Finished the flagging" << endl;          

    // Clean UP!                 
    cudaStreamDestroy(stream0);             
    cudaStreamDestroy(stream1);             
    cudaFree(flag);                

} 

我最終得到第二打印，但計算機凍結15秒後，才和我在同一時間同時獲得打印輸出。這些流應該並行運行，而不是讓系統陷入停滯狀態。我究竟做錯了什麼？我怎樣才能解決這個問題？

謝謝。

編輯

它好像一個特殊情況是通過添加volitile解決，但現在別的東西壞了。如果我在兩次內核調用之間添加任何內容，系統將恢復到原來的行爲，即一次凍結並打印所有內容。此行爲通過在set_flag和read_flag之間添加sleep(2);來顯示。另外，當放入其他程序時，會導致GPU鎖定。我現在做錯了什麼？

再次感謝。

Answer 1

允許編譯器進行相當積極的優化。此外，費米器件上的L1高速緩存不能保證一致。要解決這些問題，嘗試添加volatile關鍵字到功能flag變量的使用，像這樣：

__global__ void set_flag(volatile int *flag)  

和

__global__ void read_flag(volatile int *flag)  

一般來說，在全局內存變量居民使用時，這會導致編譯器發出繞過L1高速緩存的負載，並且通常也會阻止優化這些變量到寄存器中。

我想你會有更好的結果。

您發佈的代碼有可能因這些問題導致死鎖。因此，您看到的觀察可能實際上是中斷程序的操作系統（例如，Windows TDR）。