Как понять поведение cudaStreamQuery при запросе потока по умолчанию?

Question

Из руководства CUDA , я знаю, cudaStreamQuery работает в асинхронном потоке:

Запрашивает асинхронный поток для получения статуса завершения.

Я пишу простую программу для проверки ее поведения в «потоке по умолчанию»:

#include <cstdlib>
#include <iostream>

#define cudaSafeCall(call)  \
        do {\
            cudaError_t err = call;\
            if (cudaSuccess != err) \
            {\
                std::cerr << "CUDA error in " << __FILE__ << "(" << __LINE__ << "): " \
                    << cudaGetErrorString(err) << '\n';\
                exit(EXIT_FAILURE);\
            }\
        } while(0)

int main(void)
{
  int N = 1<<20;
  float *x, *d_x;
  x = (float*)malloc(N*sizeof(float));

  cudaSafeCall(cudaMalloc(&d_x, N*sizeof(float))); 

  cudaSafeCall(cudaMemcpy(d_x, x, N*sizeof(float), cudaMemcpyDefault));

  cudaSafeCall(cudaStreamQuery(0));

  cudaSafeCall(cudaFree(d_x));

  free(x);

  return 0;
}

Создайте и запустите ее:

$ nvcc sync_test.cu -o sync_test
$ ./sync_test
CUDA error in sync_test.cu(25): device not ready

После добавления cudaStreamSynchronize(0); между cudaMemcpy и cudaStreamQuery, ошибка исчезла:

cudaSafeCall(cudaMemcpy(d_x, x, N*sizeof(float), cudaMemcpyDefault));
cudaSafeCall(cudaStreamSynchronize(0));
cudaSafeCall(cudaStreamQuery(0));

Я немного запутался:
a) cudaMemcpy() синхронно с кодом хоста, поэтому после возврата cudaMemcpy, работа копирования в этом потоке должна быть выполнена, почему cudaStreamQuery(0) возвращает ошибку в первом случае?
b) Если cudaStreamQuery работает только на асинхронном потоке, а не на потоке по умолчанию, почему cudaStreamQuery не возвращает ошибкуво втором случае?

Answer 1

Ваше утверждение а) не совсем точно для конкретного случая в вашем примере кода:

а) cudaMemcpy() синхронно с кодом хоста, поэтому после возврата cudaMemcpy копированиеработа в этом потоке должна быть сделана,

Документация для cudaMemcpy состояний:

Эта функция показывает синхронный поведение для большинства случаев использования.

Если мы нажмем на ссылку синхронный , мы попадем на страницу, которая дополнительно определяет точное поведение.Возможно, вы захотите прочитать всю эту страницу, она не очень длинная.Я приведу два пункта из определения Синхронный :

Для передач из памяти хоста, доступного для просмотра, в память устройства, синхронизация потока выполняется до начала копирования.Функция вернется, как только буфер для постраничного копирования будет скопирован в промежуточную память для передачи DMA в память устройства, , но DMA до конечного пункта назначения может не завершиться .(выделение добавлено)

и:

Для передачи с устройства в страничную или закрепленную память хоста функция возвращается только после завершения копирования.

Ваш код соответствует первому фрагменту (пункт 2 выше).Этот тип передачи (хост-устройство) может возвращаться до того, как копия будет полностью завершена .

Если вы измените порядок указателей в вашем примере cudaMemcpy, изменив тем самым направление копирования, вы попадете во второй отрывок (пункт 4 выше).В этом случае передача гарантированно завершится, и отчет об ошибке действительно исчезнет из вашего кода.

Такое поведение возможно и учитывается в документации.