Чтение обновленной памяти из другого потока 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 но теперь что-то еще сломалось. Если я добавлю что-нибудь между двумя вызовами ядра, система вернется к старому поведению, а именно к зависанию и печати всего сразу. Это поведение показано добавлением sleep(2); между set_flag а также read_flag, Кроме того, при установке в другую программу это приводит к блокировке графического процессора. Что я делаю не так сейчас?

Еще раз спасибо.