Я пытался объяснить глобальную память тому, кто плохо знаком с CUDA. Я придумал следующее фиктивное ядро, которое блокирует другие потоки в других деформациях, пока выбранный деформация не установит глобальную переменную в другое значение:
__global__ void with_sync()
{
while (threadIdx.x / 32 != 0)
{
if (is_done != 0)
{
break;
}
}
if (threadIdx.x / 32 == 0)
{
is_done = 1;
printf("I'm done!\n");
}
}
Переменная is_done объявлена вне функции как __device__ __managed__ int (что, исправьте меня, если я ошибаюсь, означает, что переменная будет находиться в глобальном пространстве памяти.
Однако, когда я запускаю это ядро (1024 1D потоков в одном блоке), примерно так:
with_sync<<<1, 1024>>>();
cudaDeviceSynchronize();
I'm done распечатывается как ожидалось. Однако программа CUDA не завершает свою работу (я поместил cudaDeviceSynchronize() в коде хоста, чтобы он ожидал всех потоков). Это заставляет меня задуматься, не получили ли другие деформации изменения в переменную is_done. Однако я понимаю, что глобальная память подразумевает, что значение можно увидеть на уровне устройства (то есть, по крайней мере, все блоки в сетке).
У меня следующий вопрос: существует ли какое-либо кэширование / оптимизация, выполненная CUDA, которая делает его таким, чтобы могло возникнуть это несовместимое представление глобальной памяти? Есть ли способ получить доступ к «последнему» значению из переменной, которая находится в глобальной памяти?