Это несколько полезно, если вы описываете свои настройки, включая ОС и графический процессор, на котором вы работаете, независимо от того, управляет ли графический процессор дисплеем, и если в операционной системе Windows графический процессор находится в режиме WDDM или TCC.
Однако мы можем сделать некоторые общие утверждения без этого.
Как указано в комментариях, в настоящее время графический процессор с ядром CUDA не будет обслуживать запросы на отображение, если он также поддерживает отображение. Это означает, что дисплей будет «зависать» или, возможно, станет черным во время работы ядра графического процессора. Вполне возможно, что это может измениться в будущем, но это текущее и ожидаемое поведение.
Обычно в этом случае рекомендуется использовать 2-й графический процессор для запуска CUDA, если вы вообще не хотите мешать отображению, и если вы работаете в Windows, лучше всего, если этот графический процессор способен и помещен в Режим TCC.
Чтобы смягчить эффект при использовании только одного графического процессора и действительно обеспечить поддержку CUDA для производственных целей в среде отображения с одним графическим процессором, важно, чтобы сторона приложения CUDA была разработана таким образом, чтобы продолжительность ядра ограничено. Для хорошей интерактивности разумной отправной точкой является ограничение продолжительности ядра до 0,1 секунды или менее, поскольку этот уровень потери интерактивности может быть не особенно заметным. Если вы или кто-то не согласен с этим утверждением о человеческом факторе, это нормально; нам не нужно спорить об этом. Уменьшите продолжительность ядра до любого уровня, который вы решите, что приведет к хорошей интерактивности дисплея.
Ситуация еще более усложняется в случае Windows (не, насколько мне известно, в случае Linux) из-за пакетирования команд WDDM. Для повышения производительности команды могут быть пакетированы, и пакетные вызовы ядра могут привести к предполагаемым более длительным периодам потери интерактивности, чем может указывать только один вызов ядра. Я не знаю никаких методов, чтобы формально обойти это. Вы можете быть в состоянии «очистить» очередь команд WDDM, выполнив ложную (то есть не обязательную) операцию CUDA, такую как cudaStreamQuery(), после каждого вызова ядра. Опять же, я не знаю формально документированных методов для этого, и в некоторой степени это может зависеть от дизайна вашего приложения.
Что касается производительности, запуски ядра CUDA обычно занимают где-то около 100 микросекунд или меньше издержек запуска (мы можем назвать это потерянным временем). Следовательно, если мы разбиваем долго работающее ядро на 100-миллисекундные «чанки», а каждый чанк добавляет ~ 100 микросекунд накладных расходов, то общее влияние на производительность может составить порядка 0,1% снижения пропускной способности вычислений CUDA (при условии, что Задачи отображения тривиальны).
Используя в качестве примера предоставленный вами код, вы захотите разбить это ядро на последовательность ядер, сравнив / опробовав его на выбранном вами графическом процессоре, чтобы ядро работало не более примерно 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(&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));
}
Компилятор в любом случае может вычислить такого рода сжатие, но я обычно стараюсь дать ему наилучшее из возможных "начальное преимущество".