Чтение глобального флага не работает для обмена данными CPU>GPU в CUDA

Фундаментальная проблема, которую я вижу в вашем подходе, заключается в том, что потоки cuda не позволят ему работать.

Потоки CUDA имеют два основных принципа:

1. Элементы, выпущенные в один и тот же поток, не будут перекрываться; они будут сериализованы.
2. Элементы, выдаваемые в отдельные созданные потоки, могут перекрываться; нет определенного порядка этих операций, предоставляемых CUDA.

Даже если вы явно не используете потоки, вы работаете в «потоке по умолчанию», и применяется та же семантика потока.

Я не рассказываю все, что нужно знать о потоках, в этом кратком обзоре. Вы можете узнать больше о потоках CUDA в разделе 7 этой серии онлайн-обучения

Из-за потоков CUDA это:

 computeKernel<<<pBlocksPerGPU, aThreadsPerBlock>>>(pCancelRequested);

и это:

 cudaMemcpy(const_cast<bool*>(pCancelRequested), &aTrue, sizeof(bool), cudaMemcpyHostToDevice);

не могут перекрываться (они выпускаются в один и тот же поток CUDA «по умолчанию», поэтому правило 1 выше говорит, что они не могут перекрываться). Но перекрытие необходимо, если вы хотите «сигнализировать» работающему ядру. Мы должны разрешить выполнение операции cudaMemcpy одновременно с запуском ядра.

Мы можем исправить это с помощью прямого применения потоков CUDA (принимая во внимание правило 2 выше), чтобы поместить операцию копирования и операцию вычисления (ядра) в отдельные созданные потоки, чтобы они могли перекрываться. Когда мы это делаем, все работает так, как хотелось бы:

$ cat t2184.cu
#include <iostream>
#include <unistd.h>

__global__ void k(volatile int *flag){

  while (*flag != 0);
}

int main(){

  int *flag, *h_flag = new int;
  cudaStream_t s[2];
  cudaStreamCreate(s+0);
  cudaStreamCreate(s+1);
  cudaMalloc(&flag, sizeof(h_flag[0]));
  *h_flag = 1;
  cudaMemcpy(flag, h_flag, sizeof(h_flag[0]), cudaMemcpyHostToDevice);
  k<<<32, 256, 0, s[0]>>>(flag);
  sleep(5);
  *h_flag = 0;
  cudaMemcpyAsync(flag, h_flag, sizeof(h_flag[0]), cudaMemcpyHostToDevice, s[1]);
  cudaDeviceSynchronize();
}

$ nvcc -o t2184 t2184.cu
$ compute-sanitizer ./t2184
========= COMPUTE-SANITIZER
========= ERROR SUMMARY: 0 errors
$

ПРИМЕЧАНИЯ:

• Хотя это не очевидно из распечатки статического текста, программа тратит около 5 секунд до выхода. Если вы закомментируете такую ​​строку, как *h_flag = 0;, то программа зависнет, указывая на то, что метод сигнала флага работает правильно.
• Обратите внимание на использование volatile. Это необходимо указать компилятору, что любой доступ к этим данным должен быть фактическим доступом, компилятору не разрешено вносить изменения, которые препятствуют чтению или записи памяти в ожидаемом месте.

Такое поведение сигнала хост->устройство также может быть реализовано без явного использования потоков, но с закрепленной хостом памятью в качестве местоположения сигнализации, поскольку она «видима» как хосту, так и коду устройства «одновременно». Вот пример:

#include <iostream>
#include <unistd.h>

__global__ void k(volatile int *flag){

  while (*flag != 0);
}

int main(){

  int *flag;
  cudaHostAlloc(&flag, sizeof(flag[0]), cudaHostAllocDefault);
  *flag = 1;
  k<<<32, 256>>>(flag);
  sleep(5);
  *flag = 0;
  cudaDeviceSynchronize();
}

Для других примеров передачи сигналов, таких как от устройства к хосту, других читателей может заинтересовать это.