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

允许编译器进行相当激进的优化。此外，Fermi 设备上的 L1 缓存不能保证是一致的。要解决这些问题，请尝试将 volatile 关键字添加到 flag 变量的函数用法中，如下所示：

__global__ void set_flag(volatile int *flag)       

和

__global__ void read_flag(volatile int *flag)     

一般而言，当用于驻留在全局内存中的变量时，这将导致编译器发出绕过 L1 缓存的加载，并且通常还会阻止将这些变量优化到寄存器中。

我认为你会有更好的结果。

由于这些问题，您发布的代码可能会出现死锁。因此，您看到的观察结果实际上可能是操作系统（例如 Windows TDR）中断了您的程序。