从其他 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）中断了您的程序。