Мониторинг активных деформаций и потоков во время расходящегося запуска CUDA

Question

Я реализовал некоторый код CUDA. Он работает нормально, но алгоритм по своей сути создает сильное расхождение потоков. Это ожидаемо.

Позже я попытаюсь уменьшить расхождение. Но на данный момент я был бы счастлив иметь возможность измерить его.

Есть ли простой способ (предпочтительно с помощью вызова API во время выполнения или инструмента CLI) проверить, сколько из моих первоначально запланированных деформаций и/или потоков все еще активны?

Answer 1

Помимо решений, приведенных в комментариях, вы можете использовать Nsight Computeдля профилирования ваших ядер. Вы можете попробовать его CLI, а затем увидеть результаты в его графическом интерфейсе, например:

      ncu --export output --force-overwrite --target-processes application-only \
  --replay-mode kernel --kernel-regex-base function --launch-skip-before-match 0 \
  --section InstructionStats \
  --section Occupancy \
  --section SchedulerStats \
  --section SourceCounters \
  --section WarpStateStats \
  --sampling-interval auto \
  --sampling-max-passes 5 \
  --profile-from-start 1 --cache-control all --clock-control base \
  --apply-rules yes --import-source no --check-exit-code yes \
  your-appication [arguments]

Затем в его графическом интерфейсе вы можете увидеть некоторую полезную информацию. Например, в разделе счетчики источников вы можете увидеть что-то вроде этого:

Answer 2

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

      __device__ void printConvergentThreadCount(int line) // Pass __LINE__
{
   const int count = __popc(__activemask());
   const int threadId = blockIdx.x * blockDim.x + threadIdx.x;
   if (threadId == 0) // Filter
   {
      printf("Line %i: %i\n", line, count);
   }
}

Тем не менее, это не дает цифр, пока работают ядра.