cudaStreamDestroy() не синхронизирует/блокирует?

Question

Я использую CUDA 4.2 на Quadro NVS 295 на машине Win7 x64. Из Руководства по программированию CUDA C Я прочитал это:

»... Потоки освобождаются вызовом cudaStreamDestroy()

for (int i = 0; i < 2; ++i) 
cudaStreamDestroy(stream[i]); 

cudaStreamDestroy() ждет всех предыдущих команд в данном потоке в. , прежде чем уничтожить поток и вернуть управление хост-потоку ».

Действительно ли это так? Я написал небольшой код, где я делаю более или менее следующее (я буду ставить только псевдокод):

//transfer input buffer to device 
cudaMemcpyToArrayAsync(... , stream[1]); 

//launch kernel 
my_kernel <<<dimGrid, dimBlock, 0, stream[1]>>> (...); 

//transfer from device to host 
cudaMemcpyAsync(.., cudaMemcpyDeviceToHost, stream[1]); 

//Destroy stream. In theory this should block the host until everything on the stream is completed! 
ret = cudaStreamDestroy(stream[1]); 

В этом примере, кажется, что cudaStreamDestroy() вызов немедленно вернуться к хозяину, то есть не дожидаясь для вызова cudaMemcpyAsync() и других инструкций strem. Если я поставлю «cudaStreamSynchronize (поток [1]); звоните перед уничтожением потока, все идет хорошо, но медленнее. Итак, что я делаю неправильно?

Большое спасибо за ваши ответы!

Answer 1

Я не уверен, какую версию документации вы ищете, но это не то же самое, что мое. В моей документации CUDA 4.2 указано следующее:

Уничтожает и очищает асинхронный поток, указанный потоком.

В случае, устройство по-прежнему делает работу в потоке потока при cudaStreamDestroy() вызывается функция будет возвращать сразу и ресурсы, связанные с потоком будет выпущен автоматически, как только устройство завершило все работы в потоке ,

И, по моему опыту, это именно то, что он делает. Драйвер ждет, пока поток не станет пустым и уничтожит его. Но cudaStreamDestroyне блокирует вызывающую нить.

Вы можете подтвердить это, запустив этот пример:

#include <stdio.h> 
#include <assert.h> 
#include <unistd.h> 

__global__ void kernel(int * inout, const int N) 
{ 
    int gid = threadIdx.x + blockIdx.x * blockDim.x; 
    int gstride = gridDim.x * blockDim.x; 

    for (; gid < N; gid+= gstride) inout[gid] *= 2; 
} 

#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); } 
inline void gpuAssert(cudaError_t code, char *file, int line, bool abort=true) 
{ 
    if (code != cudaSuccess) 
    { 
     fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line); 
     if (abort) exit(code); 
    } 
} 

int main(void) 
{ 
    const int N = 2<<20, sz = N * sizeof(int); 

    int * inputs, * outputs, * _inout; 

    gpuErrchk(cudaMallocHost((void **)&inputs, sz)); 
    gpuErrchk(cudaMallocHost((void **)&outputs, sz)); 
    gpuErrchk(cudaMalloc((void **)&_inout, sz)); 

    for(int i=0; i<N; i++) { inputs[i] = i; outputs[i] = 0; } 

    cudaStream_t stream[2]; 
    for (int i = 0; i < 2; i++) 
     gpuErrchk(cudaStreamCreate(&stream[i])); 

    gpuErrchk(cudaMemcpyAsync(_inout, inputs, sz, cudaMemcpyHostToDevice, stream[1])); 

    kernel<<<128, 128, 0, stream[1]>>>(_inout, N); 
    gpuErrchk(cudaPeekAtLastError()); 

    gpuErrchk(cudaMemcpyAsync(outputs, _inout, sz, cudaMemcpyDeviceToHost, stream[1])); 

    for(int i = 0; i < 2; i++) 
     gpuErrchk(cudaStreamDestroy(stream[i])); 

    sleep(1); // remove the sleep and see what happens.... 

    for(int i = 0; i < N; i++) 
     assert((2 * inputs[i]) == outputs[i]); 

    cudaDeviceReset(); 

    return 0; 
} 

без sleep() код потерпит неудачу, потому что ГПУ не закончена, но с ним, assert пройдет.Обратите внимание, что sleep делает что-то тонко отличное от использования явного примитива синхронизации потока до вызовов cudaStreamDestroy, даже если результат тот же. Если поток не был пустым, когда он был уничтожен, проверка результата никогда не пройдет.

Answer 2

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

«cudaStreamDestroy() ждет всех предыдущих команд в данном потоке, чтобы завершить перед разрушением потока и возвращения управления в хост-нить.»

означает, что поток не может быть уничтожен, пока код устройства не будет завершен.