从其他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）。