Блог

Я выполняю несколько простых параллельных вычислений над массивом с помощью CUDA C . Все работает нормально, ядро выдает правильные результаты (проверено последовательным кодом процессора), но во время выполнения ядра весь мой экран становится черным на все время выполнения ядра. Я новичок в CUDA, поэтому, вероятно, я делаю что-то не так, я просто не могу понять, что.

 #define KERNEL_FOR_ITERS 1e6

__global__ void addKernel(float *c, const float *a, const float *b)
{
    int i = blockIdx.x * blockDim.x   threadIdx.x;
    if (i % 2 == 0)
        for (int j = 0; j < KERNEL_FOR_ITERS; j  )
            c[i]  = sqrt(abs(sin(a[i] * b[i])));
    else
        for (int j = 0; j < KERNEL_FOR_ITERS; j  )
            c[i]  = sqrt(abs(cos(a[i] * b[i])));
}
  

Мой вопрос в том, могу ли я предотвратить потемнение экрана на все время выполнения ядра без ущерба для слишком высокой скорости?

Это несколько полезно, если вы опишете свои настройки, включая ОС и графический процессор, на котором вы работаете, независимо от того, управляет ли графический процессор отображением, и если в ОС Windows графический процессор находится в режиме WDDM или TCC.

Однако мы можем сделать некоторые общие утверждения и без этого.

Как указано в комментариях, в настоящее время графический процессор, на котором запущено ядро CUDA, не будет обслуживать запросы на отображение, если он также поддерживает отображение. Это означает, что дисплей, по-видимому, «зависнет» или, возможно, станет черным во время работы ядра GPU. Конечно, возможно, что это может измениться в будущем, но это текущее и ожидаемое поведение.

Обычные предложения в этом случае — использовать 2-й графический процессор для запуска CUDA, если вы вообще не хотите мешать отображению, и если вы используете Windows, лучше всего, если этот графический процессор способен и переведен в режим TCC.

Чтобы смягчить эффект при использовании только одного графического процессора и, действительно, обеспечить поддержку CUDA для производственных целей в среде отображения с одним графическим процессором, важно, чтобы CUDA-сторона приложения была спроектирована таким образом, чтобы продолжительность работы ядра была ограничена. Для хорошей интерактивности разумной отправной точкой является ограничение продолжительности работы ядра до 0,1 секунды или меньше, поскольку этот уровень потери интерактивности может быть не особенно заметен. Если вы или кто-то другой не согласен с этим утверждением о человеческом факторе, все в порядке; нам не нужно спорить об этом. Сокращение продолжительности работы ядра до любого уровня, который вы решите, приведет к хорошей интерактивности отображения.

Ситуация еще более усложняется в случае Windows (насколько мне известно, не в случае Linux) из-за пакетной обработки команд WDDM. Для повышения производительности команды могут быть пакетными, а пакетное выполнение последовательных вызовов ядра может привести к более длительным периодам потери интерактивности, чем это можно было бы предположить при одном вызове ядра. Насколько мне известно, нет методов, позволяющих формально обойти это. Возможно, вам удастся «очистить» очередь команд WDDM, выполнив ложную (т. Е. не необходимую иным образом) операцию CUDA, такую как cudaStreamQuery() , после каждого вызова ядра. Опять же, я не знаю официально документированных методов для этого, и в некоторой степени это может зависеть от дизайна вашего приложения.

Что касается производительности, запуск ядра CUDA обычно требует затрат на запуск где-то около 100 микросекунд или меньше (мы можем назвать это потраченным впустую временем). Следовательно, если мы разобьем долго работающее ядро на «куски» по 100 миллисекунд, и каждый фрагмент добавит ~ 100 микросекунд накладных расходов, то итоговый эффект на производительность может составить порядка снижения пропускной способности вычислений CUDA на 0,1% (при условии, что задачи отображения тривиальны).

Используя предоставленный вами код в качестве примера, вы могли бы захотеть разбить это ядро на последовательность ядер, сравнив / синхронизировав его на выбранном вами графическом процессоре, чтобы ядро работало не дольше, чем около 100 миллисекунд (или количество по вашему выбору).

 #define KERNEL_FOR_ITERS 1e6


__global__ void addKernel(float *c, const float *a, const float *b,const int iters)
{
    int i = blockIdx.x * blockDim.x   threadIdx.x;
    if (i % 2 == 0)
        for (int j = 0; j < iters; j  )
            c[i]  = sqrt(abs(sin(a[i] * b[i])));
    else
        for (int j = 0; j < iters; j  )
            c[i]  = sqrt(abs(cos(a[i] * b[i])));
}

...
const int loop_iters = 1e4; // chosen by tuning or benchmarking
cudaStream_t str;
cudaStreamCreate(amp;str);
for (int i = 0; i < KERNEL_FOR_ITERS; i = loop_iters){
  addKernel<<<...,0,str>>>(d_c, d_a, d_b, loop_iters);
  cudaStreamQuery(str);//probably unnecessary on linux}
  

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

 __global__ void addKernel(float *c, const float *a, const float *b,const int iters)
{
    int i = blockIdx.x * blockDim.x   threadIdx.x;
    float val = a[i] * b[i];
    if (i % 2 == 0)
       val = sin(val);
    else 
       val = cos(val);
    for (int j = 0; j < iters; j  )
            c[i]  = sqrt(abs(val));
}
  

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