Потоки CUDA блокируются, несмотря на асинхронность

Я работаю над видеопотоком в реальном времени, который пытаюсь обработать с помощью GeForce GTX 960M. (Windows 10, VS 2013, CUDA 8.0)

Каждый кадр должен быть захвачен, слегка размыт, и всякий раз, когда я могу, мне нужно проделать тяжелую работу по 10 последним кадрам. Поэтому мне нужно захватить ВСЕ кадры со скоростью 30 кадров в секунду, и я ожидаю получить результат тяжелой работы со скоростью 5 кадров в секунду.

Моя проблема заключается в том, что я не могу поддерживать захват в нужном темпе: кажется, что тяжелые вычисления замедляют захват кадров либо на уровне процессора, либо на уровне графического процессора. Я пропускаю некоторые кадры...

Я пробовал много решений. Ни один не работал:

1. I tried to set-up jobs on 2 streams (image below):
   • the host gets a frame
   • Первый поток (называемый Stream2): cudaMemcpyAsync копирует кадр на Устройство. Затем первое ядро ​​выполняет базовые вычисления размытия. (На приложенном изображении размытие появляется в виде короткого интервала на 3,07 с и 3,085 с. А затем ничего ... пока не закончится большая часть)
   • хост проверяет, доступен ли второй поток благодаря CudaEvent, и запускает его, если это возможно. Практически поток доступен на 1/2 попытки.
   • Второй поток (называемый Stream4): запускает сложные вычисления в ядре (kernelCalcul_W2), выводит результат и записывает событие.

захват NSight

Практически я написал:

cudaStream_t  sHigh, sLow;
cudaStreamCreateWithPriority(&sHigh, cudaStreamNonBlocking, priority_high);
cudaStreamCreateWithPriority(&sLow, cudaStreamNonBlocking, priority_low);

cudaEvent_t event_1;
cudaEventCreate(&event_1);

if (frame has arrived)
{
    cudaMemcpyAsync(..., sHigh);        // HtoD, to upload images in the GPU
    blur_Image <<<... , sHigh>>> (...)
    if (cudaEventQuery(event_1)==cudaSuccess)) hard_work(sLow);
    else printf("Event 2 not ready\n");
}

void hard_work( cudaStream_t sLow_)
{
    kernelCalcul_W2<<<... , sLow_>>> (...);
    cudaMemcpyAsync(... the result..., sLow_); //DtoH
    cudaEventRecord(event_1, sLow_);    
}

2. I tried to use only one stream. It's the same code as above, but change 1 parameter while launching hard_work.
   • host gets a frame
   • Поток: cudaMemcpyAsync копирует кадр на Устройство. Затем ядро ​​выполняет базовые вычисления размытия. Затем, если CudaEvent Event_1 в порядке, я начинаю тяжелую работу и добавляю Event_1, чтобы получить статус в следующем раунде. Практически поток ВСЕГДА доступен: я никогда не попадаю в «остальную» часть.

Таким образом, пока выполняется тяжелая работа, я ожидал «буферизировать» все кадры для копирования и не потерять ни одного. Но я кое-что теряю: получается, что каждый раз, когда я получаю кадр и копирую его, Event_1 кажется в порядке, поэтому я запускаю тяжелую работу и получаю следующий кадр очень поздно.

3. Я попытался поместить два потока в два разных потока (в C). Не лучше (даже хуже).

Так вот вопрос: как сделать так, чтобы первый поток захватывал ВСЕ кадры? У меня действительно такое ощущение, что разные потоки блокируют ЦП.

Я показываю изображения с OpenGL. Будет ли это мешать?

Любая идея способов улучшить это? Большое спасибо!

EDIT: Как и просили, я разместил здесь MCVE.

Есть параметр, который вы можете настроить (#define ADJUST), чтобы увидеть, что происходит. По сути, основная процедура отправляет запросы CUDA в асинхронном режиме, но, похоже, блокирует основной поток. Как вы увидите на изображении, у меня есть «доступ к памяти» (т. е. захваченные изображения) каждые 30 мс, за исключением случаев, когда выполняется тяжелая работа (тогда я просто не получаю изображения).

Последняя деталь: я использую CUDA 7.5 для запуска. Я пытался установить 8.0, но, видимо, компилятор все еще 7.5.

#define _USE_MATH_DEFINES 1
#define _CRT_SECURE_NO_WARNINGS 1

#include <stdio.h>
#include <stdlib.h>
#include <time.h>
#include <Windows.h>

#define ADJUST  400
// adjusting this paramter may make the problem occur.
// Too high => probably watchdog will stop the kernel
// too low => probably the kernel will run smothly

unsigned short * images_as_Unsigned_in_Host;
unsigned short * Images_as_Unsigned_in_Device;
unsigned short * camera;
float * images_as_Output_in_Host;
float *  Images_as_Float_in_Device;
float * imageOutput_in_Device;

unsigned short imageWidth, imageHeight, totNbOfImages, imageSlot;
unsigned long imagePixelSize;
unsigned short lastImageFromCamera;


cudaStream_t  s1, s2;
cudaEvent_t event_2;
clock_t timeRef;

// Basically, in the middle of the image, I average the values. I removed the logic behind to make it simpler.
// This kernel runs fast, and that's the point.
__global__ void blurImage(unsigned short * Images_as_Unsigned_in_Device_, float * Images_as_Float_in_Device_, unsigned short imageWidth_, 
    unsigned long  imagePixelSize_, short blur_distance)
{
    // we start from 'blur_distance' from the edge
    // p0 is the point we will calculate. p is a pointer which will move around for average
    unsigned long p0 = (threadIdx.x + blur_distance) + (blockIdx.x + blur_distance) * imageWidth_;
    unsigned long p = p0;
    unsigned short * us;
    if (p >= imagePixelSize_) return;
    unsigned long tot = 0;
    short a, b, n, k;
    k = 0;
    // p starts from the top edge and will move to the right-bottom
    p -= blur_distance + blur_distance * imageWidth_;
    us = Images_as_Unsigned_in_Device_ + p;
    for (a = 2 * blur_distance; a >= 0; a--)
    {
        for (b = 2 * blur_distance; b >= 0; b--)
        {
            n = *us;
            if (n > 0) { tot += n; k++; }
            us++;
        }
        us += imageWidth_ - 2 * blur_distance - 1;
    }
    if (k > 0) Images_as_Float_in_Device_[p0] = (float)tot / (float)k;
    else Images_as_Float_in_Device_[p0] = 128.f;
}


__global__ void kernelCalcul_W2(float *inputImage, float *outputImage, unsigned long  imagePixelSize_, unsigned short imageWidth_, unsigned short slot, unsigned short totImages)
{
    // point the pixel and crunch it
    unsigned long p = threadIdx.x + blockIdx.x * imageWidth_;
    if (p >= imagePixelSize_)   { return; }
    float result;
    long a, b, n, n0;
    float input;
    b = 3;

    // this is not the right algorithm (which is pretty complex). 
    // I know this is not optimal in terms of memory management. Still, I want a "long" calculation here so I don't care...
    for (n = 0; n < 10; n++)
    {
        n0 = slot - n;
        if (n0 < 0) n0 += totImages;
        input = inputImage[p + n0 * imagePixelSize_]; 
        for (a = 0; a < ADJUST ; a++)
                result += pow(input, inputImage[a + n0 * imagePixelSize_]) * cos(input);
    }
    outputImage[p] = result;
}


void hard_work( cudaStream_t s){

    cudaError err;
    // launch the hard work
    printf("Hard work is launched after image %d is captured  ==> ", imageSlot);
    kernelCalcul_W2 << <340, 500, 0, s >> >(Images_as_Float_in_Device, imageOutput_in_Device, imagePixelSize, imageWidth, imageSlot, totNbOfImages);
    err = cudaPeekAtLastError();
    if (err != cudaSuccess) printf( "running error: %s \n", cudaGetErrorString(err));
    else printf("running ok\n");

    // copy the result back to Host
    //printf(" %p  %p  \n", images_as_Output_in_Host, imageOutput_in_Device);
    cudaMemcpyAsync(images_as_Output_in_Host, imageOutput_in_Device, sizeof(float) *  imagePixelSize, cudaMemcpyDeviceToHost, s);
    cudaEventRecord(event_2, s);
}


void createStorageSpace()
{
    imageWidth = 640;
    imageHeight = 480;
    totNbOfImages = 300;
    imageSlot = 0;
    imagePixelSize = 640 * 480;
    lastImageFromCamera = 0;

    camera = (unsigned short *)malloc(imagePixelSize * sizeof(unsigned short));
    for (int i = 0; i < imagePixelSize; i++) camera[i] = rand() % 255;
    // storing the images in the Host memory. I know I could optimize with cudaHostAllocate.
    images_as_Unsigned_in_Host = (unsigned short *) malloc(imagePixelSize * sizeof(unsigned short) * totNbOfImages);
    images_as_Output_in_Host = (float *)malloc(imagePixelSize * sizeof(float));

    cudaMalloc(&Images_as_Unsigned_in_Device, imagePixelSize * sizeof(unsigned short) * totNbOfImages);
    cudaMalloc(&Images_as_Float_in_Device, imagePixelSize * sizeof(float) * totNbOfImages);

    cudaMalloc(&imageOutput_in_Device, imagePixelSize * sizeof(float));



    int priority_high, priority_low;
    cudaDeviceGetStreamPriorityRange(&priority_low, &priority_high);
    cudaStreamCreateWithPriority(&s1, cudaStreamNonBlocking, priority_high);
    cudaStreamCreateWithPriority(&s2, cudaStreamNonBlocking, priority_low);
    cudaEventCreate(&event_2);

}

void releaseMapFile()
{
    cudaFree(Images_as_Unsigned_in_Device);
    cudaFree(Images_as_Float_in_Device);
    cudaFree(imageOutput_in_Device);
    free(images_as_Output_in_Host);
    free(camera);

    cudaStreamDestroy(s1);
    cudaStreamDestroy(s2);
    cudaEventDestroy(event_2);
}

void putImageCUDA(const void * data)
{       
    // We put the image in a round-robin. The slot to put the image is imageSlot
    printf("\nDealing with image %d\n", imageSlot);
    // Copy the image in the Round Robin
    cudaMemcpyAsync(Images_as_Unsigned_in_Device + imageSlot * imagePixelSize, data, sizeof(unsigned short) *  imagePixelSize, cudaMemcpyHostToDevice, s1);

    // We will blur the image. Let's prepare the memory to get the results as floats
    cudaMemsetAsync(Images_as_Float_in_Device + imageSlot * imagePixelSize, 0., sizeof(float) *  imagePixelSize, s1);

    // blur image
    blurImage << <imageHeight - 140, imageWidth - 140, 0, s1 >> > (Images_as_Unsigned_in_Device + imageSlot * imagePixelSize,
                Images_as_Float_in_Device + imageSlot * imagePixelSize,
                imageWidth, imagePixelSize, 3);


    // launches the hard-work
    if (cudaEventQuery(event_2) == cudaSuccess) hard_work(s2);
    else printf("Hard_work still running, so unable to process after image %d\n", imageSlot);

    imageSlot++;
    if (imageSlot >= totNbOfImages) {
        imageSlot = 0;
    }
}

int main()
{
    createStorageSpace();
    printf("The following loop is supposed to push images in the GPU and do calculations in Async mode, and to wait 30 ms before the next image, so we should have the output on the screen in 10 x 30 ms. But it's far slower...\nYou may adjust a #define ADJUST parameter to see what's happening.");

    for (int i = 0; i < 10; i++)
    {
        putImageCUDA(camera);  // Puts an image in the GPU, does the bluring, and tries to do the hard-work
        Sleep(30);  // to simulate Camera
    }
    releaseMapFile();
    getchar();
}