Как использовать Nvidia Multi-Process Service (MPS) для запуска нескольких приложений не-MPI CUDA?

Question

Могу ли я одновременно запускать приложения не MPI CUDA на графических процессорах NVIDIA Kepler с MPS? Я хотел бы сделать это, потому что мои приложения не могут полностью использовать графический процессор, поэтому я хочу, чтобы они работали вместе. Есть ли пример кода для этого?

Answer 1

Необходимые инструкции содержатся в документации для службы MPS. Вы заметите, что эти инструкции на самом деле не зависят от MPI и не вызывают его, поэтому в них нет ничего специфичного для MPI.

Вот пошаговое руководство / пример.

1. Прочитайте раздел 2.3 вышеупомянутой документации для различных требований и ограничений. Для этого я рекомендую использовать CUDA 7, 7.5 или более позднюю версию. Были некоторые различия в конфигурации предыдущих версий CUDA MPS, которые я не буду здесь рассматривать. Кроме того, я продемонстрирую только использование одного сервера / одного графического процессора. Машина, которую я использую для тестирования, - это узел CentOS 6.2, использующий графический процессор K40c (cc3.5/Kepler) с CUDA 7.0. В узле есть другие графические процессоры. В моем случае порядок перечисления CUDA размещает мой K40c на устройстве 0, но порядок перечисления nvidia-smi помещает его как id 2 в порядок. Все эти детали имеют значение в системе с несколькими графическими процессорами, влияя на сценарии, приведенные ниже.

2. Я создам несколько вспомогательных скриптов bash, а также тестовое приложение. Для тестового приложения нам бы хотелось что-то с ядром (ями), которое, очевидно, могло бы работать одновременно с ядрами из других экземпляров приложения, и мы также хотели бы что-то, что делает это очевидным, когда эти ядра (из отдельных приложений / процессов) работают одновременно или нет. Чтобы удовлетворить эти потребности в демонстрационных целях, давайте создадим приложение, которое имеет ядро, которое просто запускается в одном потоке на одном SM и просто ждет некоторый период времени (мы будем использовать ~5 секунд) перед выходом и печатью сообщение. Вот тестовое приложение, которое делает это:

   $ cat t1034.cu
   #include <stdio.h>
   #include <stdlib.h>
   
   #define MAX_DELAY 30
   
   #define cudaCheckErrors(msg) \
     do { \
       cudaError_t __err = cudaGetLastError(); \
       if (__err != cudaSuccess) { \
           fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \
               msg, cudaGetErrorString(__err), \
               __FILE__, __LINE__); \
           fprintf(stderr, "*** FAILED - ABORTING\n"); \
           exit(1); \
       } \
     } while (0)
   
   
   #include <time.h>
   #include <sys/time.h>
   #define USECPSEC 1000000ULL
   
   unsigned long long dtime_usec(unsigned long long start){
   
     timeval tv;
     gettimeofday(&tv, 0);
     return ((tv.tv_sec*USECPSEC)+tv.tv_usec)-start;
   }
   
   #define APPRX_CLKS_PER_SEC 1000000000ULL
   __global__ void delay_kernel(unsigned seconds){
   
     unsigned long long dt = clock64();
     while (clock64() < (dt + (seconds*APPRX_CLKS_PER_SEC)));
   }
   
   int main(int argc, char *argv[]){
   
     unsigned delay_t = 5; // seconds, approximately
     unsigned delay_t_r;
     if (argc > 1) delay_t_r = atoi(argv[1]);
     if ((delay_t_r > 0) && (delay_t_r < MAX_DELAY)) delay_t = delay_t_r;
     unsigned long long difft = dtime_usec(0);
     delay_kernel<<<1,1>>>(delay_t);
     cudaDeviceSynchronize();
     cudaCheckErrors("kernel fail");
     difft = dtime_usec(difft);
     printf("kernel duration: %fs\n", difft/(float)USECPSEC);
     return 0;
   }
   
   
   $ nvcc -arch=sm_35 -o t1034 t1034.cu
   $ ./t1034
   kernel duration: 6.528574s
   $
   

3. Мы будем использовать скрипт bash для запуска сервера MPS:

   $ cat start_as_root.bash
   #!/bin/bash
   # the following must be performed with root privilege
   export CUDA_VISIBLE_DEVICES="0"
   nvidia-smi -i 2 -c EXCLUSIVE_PROCESS
   nvidia-cuda-mps-control -d
   $
   

4. И скрипт bash для запуска 2-х копий нашего тестового приложения "одновременно":

   $ cat mps_run
   #!/bin/bash
   ./t1034 &
   ./t1034
   $
   

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

   $ cat stop_as_root.bash
   #!/bin/bash
   echo quit | nvidia-cuda-mps-control
   nvidia-smi -i 2 -c DEFAULT
   $
   

6. Теперь, когда мы просто запускаем наше тестовое приложение, используя mps_run Сценарий выше, но без фактического включения сервера MPS, мы получаем ожидаемое поведение, что один экземпляр приложения занимает ожидаемые ~5 секунд, тогда как другой экземпляр занимает примерно вдвое больше (~10 секунд), потому что, поскольку он не работает одновременно с приложением из другого процесса оно ждет 5 секунд, пока другое приложение / ядро ​​работает, а затем тратит 5 секунд на запуск своего собственного ядра, в общей сложности ~ 10 секунд:

   $ ./mps_run
   kernel duration: 6.409399s
   kernel duration: 12.078304s
   $
   

7. С другой стороны, если мы сначала запустим сервер MPS и повторим тест:

   $ su
   Password:
   # ./start_as_root.bash
   Set compute mode to EXCLUSIVE_PROCESS for GPU 0000:82:00.0.
   All done.
   # exit
   exit
   $ ./mps_run
   kernel duration: 6.167079s
   kernel duration: 6.263062s
   $
   

   мы видим, что запуск обоих приложений занимает одинаковое количество времени, потому что ядра работают одновременно из-за MPS.

8. Вы можете экспериментировать, как считаете нужным. Если эта последовательность работает правильно для вас, но запуск вашего собственного приложения не дает ожидаемых результатов, одной из возможных причин может быть то, что ваше приложение / ядра не могут работать одновременно с другими экземплярами приложения / ядра из-за к построению ваших ядер, никак не связанных с MPS. Возможно, вы захотите проверить требования к параллельным ядрам и / или изучить пример приложения concurrentKernels.

9. Большая часть информации здесь была взята из теста / работы, проделанной здесь, хотя представление здесь с отдельными приложениями отличается от случая MPI, представленного там.

ОБНОВЛЕНИЕ: Поведение планировщика в случае не-MPS при запуске ядер из нескольких процессов изменилось с Pascal и более новыми GPU. Приведенные выше результаты теста по-прежнему верны для графических процессоров, на которых проводилось тестирование (например, Kepler), но при запуске вышеуказанного тестового примера на Pascal или более новом GPU, в случае без MPS будут получены другие результаты. Планировщик описан как планировщик с разделением по времени в последнем документе MPS, и, по-видимому, происходит то, что вместо того, чтобы ждать завершения ядра одного процесса, планировщик может, в соответствии с некоторыми неопубликованными правилами, выбрать предварительно - освободить работающее ядро, чтобы оно могло переключиться на другое ядро ​​из другого процесса. Это по-прежнему не означает, что ядра из отдельных процессов работают "одновременно" при традиционном использовании этого слова в документации CUDA, но приведенный выше код "обманут" планировщиком с временными интервалами (на Pascal и новее), потому что это зависит на использовании часов SM для установки продолжительности ядра. Комбинация планировщика с временными интервалами плюс использование часов SM делает этот тестовый пример работающим "одновременно". Однако, как описано в документе MPS, код из ядра A не выполняется в том же тактовом цикле (ах), что и код из ядра B, когда A и B происходят из отдельных процессов в случае без MPS.

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