从其他 CUDA 流读取更新的内存

我正在尝试在一个内核函数中设置一个标志并在另一个内核函数中读取它。 基本上，我正在尝试执行以下操作。

#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 锁定。 我现在做错了什么？

再次感谢。