Использование узлов с GPU

  • GPGPU (general-purpose graphics processing units) - концепция использования для выполнения расчётов графического процессора вместо центрального процессора. В процессе эволюции видеокарт в определённый момент стало понятно, что GPU не только не уступают CPU в производительности, но на определённых задачах значительно превосходят за счёт наличия существенно большего количества ядер.
  • Важный момент при использовании GPU заключается в том, что они реализуют принцип SIMD (single instruction, multiple data - одиночный поток команд, множественный поток данных) и для решения задач другого типа подходят плохо.
  • Существует несколько реализаций GPGPU:
    1. AMD FireStream (ATI Stream).
    2. CUDA (Compute Unified Device Architecture) - архитектура параллельных вычислений, поддерживаемая NVIDIA.
    3. DirectCompute, входит в состав Microsoft DirectX (начиная с версии 11).
    4. OpenCL (Open Computing Language) - открытый стандарт для использования любых вычислительных ресурсов: CPU, GPU, карт расширения, … . Поддерживается как AMD так и NVIDIA.
  • В ИВЦ НГУ используются вычислительные узлы HP SL390s G7, каждый из которых имеет:
    • Два 6-ядерных CPU Xeon X5670 (2.93GHz)
    • 96 ГБ ОЗУ
    • Три карты NVIDIA Tesla M2090 на архитектуре Fermi (compute capability 2.0), у каждой:
      • 1 GPU c 512 ядрами
      • 6 ГБ памяти GDDR5 при выключенном контроле чётности (при включении ECC некая часть будет тратиться для обеспечения контроля)
      • 665 Гигафлопс пиковой производительности для вычислений с двойной точностью, 1331 Гфлопс для одинарной

Основные реализации GPGPU для имеющегося оборудования - CUDA и OpenCL от NVidia.

CUDA Toolkit

  • CUDA Toolkit предоставляет полную среду разработки на C и C++ с использованием вычислений на GPU NVIDIA. Включает компилятор для GPU, инструменты для отладки и оптимизации, математические библиотеки и документацию.
  • Разные версии CUDA Toolkit установлены в поддиректории внутри директории '/opt/shared/nvidia'. Например, версия 6.0 установлена в '/opt/shared/nvidia/cuda-6.0'. При этом символьная ссылка '/opt/shared/nvidia/cuda' указывает на директорию с последней установленной версией.
  • Для использования необходимо настроить переменные окружения следующим образом:
    export PATH=$PATH:/opt/shared/nvidia/cuda/bin
    export LD_LIBRARY_PATH=/opt/shared/nvidia/cuda/lib64:/opt/shared/nvidia/cuda/lib:$LD_LIBRARY_PATH

    Тоже самое можно сделать с помощью утилиты module:

    module load nvidia/cuda-toolkit
  • Если необходимо использовать какую-то другую установленную версию, пути (или параметр для 'module') необходимо поменять соответствующим образом.
  • Если необходимо использовать версию, не установленную у нас, вы можете самостоятельно установить её в свою домашнюю директорию. Это не требует повышенных прав.

CUDA Code Samples

  • CUDA Code Samples (предыдущее название - GPU Computing SDK) содержит примеры кода и официальные документы, призванные помочь создавать ПО, использующее NVIDIA GPU, с помощью CUDA C/C++, OpenCL или DirectCompute.
  • Установлено в директории вида /opt/shared/nvidia/NVIDIA_CUDA-6.0_Samples (в зависимости от версии).

Прикладное ПО

На сайте NVIDIA имеется список прикладного ПО сторонних разработчиков, умеющего использовать GPU этой фирмы. Если используемое вами ПО есть в этом списке, может быть полезно сравнить скорость его работы на CPU и на узлах с GPU и, возможно, начать использовать GPU.

  • Узлы с GPU выделены в отдельную очередь, называемую 'teslaq', которую необходимо указывать команде 'qsub'.
  • Ресурсы каждой видеокарты всегда должны использоваться одним пользователем монопольно. Тем не менее, наличие трёх GPU на каждом сервере позволяет одновременно работать 3-м пользователям. Для этого каждый из серверов делится на три виртуальных узла (virtual node, vnode), каждый из которых содержит 4 процессорных ядра, 32 ГБ ОЗУ и 1 GPU. Стоит заметить, что название «виртуальный узел» подразумевает только логическое выделение части ресурсов физического сервера с точки зрения планировщика PBS, но не использование виртуализации.
  • Для запроса необходимого количества GPU используется ресурс ngpus:
    #PBS -l select=1:ngpus=1:ncpus=4:mem=32gb
  • После запуска задачи она должна узнать, какой виртуальный узел (или узлы) ей выделен и по его названию понять порядковый номер GPU, который она должна использовать. Например, виртуальные узлы сервера 'sl002' называются sl002[0], sl002[1] и sl002[2]. Если задаче выделен виртуальный узел sl002[2], значит, она должна использовать именно GPU №2 на узле sl002.
  • Определить предоставленный виртуальный узел можно с помощью значения параметра 'exec_vnode' в выводе команды 'qstat -f $PBS_JOBID'. Данную операцию можно произвести как из скрипта (примеры приведены ниже), так и из самого запускаемого приложения, если имеется возможность модифицировать его код. Следует учитывать, что при запросе нескольких GPU значение 'exec_vnode' может занимать более одной строки.
  • GPU нумеруются начиная с '0', индекс соответствует номеру устройства в CUDA (например, в команде cudaSetDevice) или в выводе команды 'nvidia-smi':
    nvidia-smi -L
    GPU 0: Tesla M2090 (S/N: 0341615011546)
    GPU 1: Tesla M2090 (S/N: 0341615011399)
    GPU 2: Tesla M2090 (S/N: 0341615011254)
  • :!: Автоматического ограничения доступа к остальным GPU узла средствами PBS или драйверов NVIDIA не происходит, пользователь должен сам настраивать свои задачи для использования только предоставленных ему GPU.

Пример 1

  • В случае, если запрашивается ровно один GPU, можно быть уверенным, что значение 'exec_vnode' поместится в одну строку. Кроме того, нет необходимости определять, какой физический сервер нам выделен, т.к. именно на нём задача и запустится.
  • Код, определяющий предоставленный GPU, его SN и UUID:
    #!/bin/sh
    
    #PBS -q teslaq
    #PBS -l walltime=0:01:00
    #PBS -l select=1:ngpus=1:ncpus=4:mem=32gb
    
    cd $PBS_O_WORKDIR
    
    vnodes=$(qstat -f $PBS_JOBID|grep exec_vnode|sed -e 's/ *//')
    echo "$vnodes"
    if [ $(echo $vnodes|grep -c '+') != 0 ] ; then
        echo "Error: several vnodes are provided."
        exit 100
    fi
    gpu=$(echo $vnodes|sed 's/.*\[//'|sed 's/\].*//')
    sn=$(nvidia-smi -q -i $gpu_nbr|grep 'Serial Number'|awk -F\: '{print $2}'|sed -e 's/ *//')
    uuid=$(nvidia-smi -q -i $gpu_nbr|grep 'GPU UUID'|awk -F\: '{print $2}'|sed -e 's/ *//')
    echo "GPU  = $gpu"
    echo "SN   = $sn"
    echo "UUID = $uuid"
  • В результате работы получим в файле стандартного вывода примерно следующее:
    exec_vnode = (sl001[0]:ncpus=4:ngpus=1:mem=33554432kb)
    GPU  = 0
    SN   = 0341615011546
    UUID = GPU-df433275cee7fdcd-53619fe0-b727ce0a-4aa635a9-69819036549a1e42288e1692

Пример 2

  • Может потребоваться использовать в одной задаче несколько GPU, в том числе принадлежащих разным физическим серверам. Например, для гибридной задачи, использующей как GPGPU, так и MPI. В этом случае анализ предоставленных задаче виртуальных узлов будет немного сложнее.
    #!/bin/sh
    
    #PBS -q teslaq
    #PBS -l walltime=0:01:00
    #PBS -l select=3:mpiprocs=2:ngpus=2:ncpus=8:mem=32gb,place=scatter
    
    cd $PBS_O_WORKDIR
    vnodes=$(qstat -f $PBS_JOBID|tr -d '\n'' ''\t'|sed 's/Hold_Types.*//'|sed 's/.*exec_vnode=//'|tr -d \(\)|tr + '\n'|sed 's/:.*//'|sort)
    echo "My vnodes:"
    for vnode in $vnodes ; do
        node=$(echo $vnode|sed 's/\[.*//')
        gpu=$(echo $vnode|sed 's/.*\[//'|sed 's/\]//')
        echo "$vnode = Node $node, GPU $gpu"
    done
    echo
    echo "PBS_NODEFILE:"
    cat $PBS_NODEFILE
  • В результате работы получим в файле стандартного вывода примерно следующее:
    My vnodes:
    sl001[1] = Node sl001, GPU 1
    sl001[2] = Node sl001, GPU 2
    sl002[0] = Node sl002, GPU 0
    sl002[1] = Node sl002, GPU 1
    sl004[0] = Node sl004, GPU 0
    sl004[1] = Node sl004, GPU 1
    
    PBS_NODEFILE:
    sl001
    sl001
    sl004
    sl004
    sl002
    sl002
  • CUDA представляет из себя модификацию языка C, поэтому для исходных файлов принято использовать специальное расширение - 'cu' вместо 'c'.
  • Основные термины:
    • host - компьютер 'в обычном понимании', управляемый CPU.
    • device (устройство) - карта с GPU.
    • kernel (ядро) - функция, которая будет запущена в нескольких экземплярах, каждый из которых будет работать на своём ядре устройства. Для указания, что функция будет ядром, при её описании используется спецификатор
      ___global___

      Выполнение ядра с распараллеливанием на N потоков описывается следующим образом:

      MyKernel<<<1, N>>>(параметры);


  • Ниже приведён пример запуска программы, складывающей средствами CUDA два вектора 'A' и 'B' (т.е. два массива поэлементно) и сохраняющей сумму в вектор 'С'.
  • Поскольку GPU обрабатывает данные, находящиеся в своей собственной памяти, а не в ОЗУ компьютера, требуются дополнительные действия - выделение памяти на устройстве, копирование туда исходных данных, копирование полученного результата обратно на компьютер.
  • Номер устройства, которое будет использоваться, программа будет получать при запуске в качестве первого параметра.


  • Создать файл 'addvectors.cu' следующего содержания:
    #include <stdio.h>
    #include "cuda.h"
    
    #define N 128
    
    int assigned_device;
    int used_device;
    
    // Data on the host system
    int HostA[N];
    int HostB[N];
    int HostC[N];
    
    // Pointers to data on the device
    int *DeviceA;
    int *DeviceB;
    int *DeviceC;
    
    //----------------------------------------------------------
    __global__ void AddVectors(int* a, int* b, int* c) {
        int i = threadIdx.x;
        c[i] = a[i] + b[i];
    }
    
    //----------------------------------------------------------
    int main(int argc, char** argv) {
    
    // Define the device to use:
        if (argc < 2) {
            printf ("Error: device number is absent\n");
            return 100;
        }
        assigned_device=atoi(argv[1]);
        if ( strlen(argv[1]) > 1 or ( assigned_device == 0 and strcmp(argv[1],"0") != 0 ) ) {
            printf ("Error: device number is incorrect\n");
            return 110;
        }
    
    // Select the used device:
        if ( cudaSetDevice(assigned_device) != cudaSuccess or
             cudaGetDevice( &used_device ) != cudaSuccess or
             used_device != assigned_device
           ) {
            printf ("Error: unable to set device %d\n", assigned_device);
            return 120;
        }
        printf ("Used device: %d\n", used_device);
    
    // Initialize summands:
        for (int i=0; i<N; i++) {
            HostA[i]=i*2;
            HostB[i]=i*3;
        }
    
    // Allocate memory on the device:
        cudaMalloc((void**)&DeviceA, N*sizeof(int));
        cudaMalloc((void**)&DeviceB, N*sizeof(int));
        cudaMalloc((void**)&DeviceC, N*sizeof(int));
    
    // Copy summands from host to device:
        cudaMemcpy(DeviceA, HostA, N*sizeof(int), cudaMemcpyHostToDevice);
        cudaMemcpy(DeviceB, HostB, N*sizeof(int), cudaMemcpyHostToDevice);
    
    // Execute kernel:
        AddVectors<<<1, N>>>(DeviceA, DeviceB, DeviceC);
    
    // Copy result from device to host:
        cudaMemcpy(HostC, DeviceC, N*sizeof(int), cudaMemcpyDeviceToHost);
    
    // Show result:
        for (int i=0; i<N; i++) {
            printf ("%d + %d = %d\n",HostA[i],HostB[i],HostC[i]);
        }
    
        cudaFree(DeviceA);
        cudaFree(DeviceB);
        cudaFree(DeviceC);
    }
  • Для установки переменных окружения выполнить:
    module load nvidia/cuda-toolkit
  • Для компилирования используется утилита nvcc. Необходимо выполнить:
    nvcc -arch=compute_20 addvectors.cu -o addvectors

    Либо можно использовать 'makefile' следующего содержания

    addvectors : addvectors.cu
            nvcc -arch=compute_20 addvectors.cu -o $@
  • В результате должен быть создан исполняемый файл 'addvectors'
  • Для взаимодействия с планировщиком PBS создать файл 'submit.sh' следующего содержания:
    #!/bin/sh
    
    #PBS -q teslaq
    #PBS -l walltime=0:01:00
    #PBS -l select=1:ngpus=1:ncpus=4:mem=32gb
    
    cd $PBS_O_WORKDIR
    
    vnodes=$(qstat -f $PBS_JOBID|grep exec_vnode|sed -e 's/ *//')
    if [ $( echo $vnodes|grep -c '+') != 0 ] ; then
        echo "Error: several vnodes are provided."
        exit 100
    fi
    gpu_nbr=$(echo $vnodes|sed 's/.*\[//'|sed 's/\].*//')
    echo "GPU number from vnodes = $gpu_nbr"
    echo
    
    export PATH=$PATH:/opt/shared/nvidia/cuda/bin
    export LD_LIBRARY_PATH=/opt/shared/nvidia/cuda/lib64:/opt/shared/nvidia/cuda/lib:$LD_LIBRARY_PATH
    
    ./addvectors $gpu_nbr
  • Поставить задачу в очередь:
    qsub submit.sh
  • После завершения в файле стандартного вывода будет получено примерно следующее:
    GPU number from vnodes = 0
    
    Used device: 0
    0 + 0 = 0
    2 + 3 = 5
    4 + 6 = 10
    6 + 9 = 15
    8 + 12 = 20
    ...
    246 + 369 = 615
    248 + 372 = 620
    250 + 375 = 625
    252 + 378 = 630
    254 + 381 = 635
  • Для использования CUDA совместно с MPI необходимо использовать два компилятора - nvcc для кода CUDA и mpicc (mpiCC/mpicxx/…) для остального. Поэтому исходный текст должен быть разбит на несколько файлов таким образом, чтобы разделить CUDA и MPI код. Каждый файл компилируется соответствующим компилятором, но не до исполняемого файла, а до объектного. Затем все полученные объектные файлы объединяются в один исполняемый.
  • Кроме того, каждый процесс MPI должен знать, с каким именно GPU, присутствующем на сервере, он должен работать. Данная информация не предоставляется явно средствами PBS, MPI или CUDA и программа должна выяснять её самостоятельно.
  • Команда mpirun должна получить корректное значение переменной окружения LD_LIBRARY_PATH. Это необходимо, чтобы обеспечить установку этой переменной на всех используемых узлах. Без этого программа не найдёт библиотеки CUDA и завершится аварийно. При использовании OpenMPI это делается при помощи параметра '-x', для Intel MPI - параметром '-genvlist'.
  • Ниже приведён рабочий пример, в котором каждый MPI-процесс работает со своим GPU.
  • Выбрать необходимую реализацию MPI:
    mpi-selector --set openmpi_gcc-1.4.4
  • Переподключить SSH-соединение.
  • Установить переменные окружения для работы компилятора nvcc:
    module load nvidia/cuda-toolkit
  • Создать файл 'cuda_part.cu' следующего содержания:
    #include <stdio.h>
    #include "cuda.h"
    
    #define N 4
    
    int used_device;
    
    // Data on the host system
    int HostA[N];
    int HostB[N];
    int HostC[N];
    
    // Pointers to data on the device
    int *DeviceA;
    int *DeviceB;
    int *DeviceC;
    
    //----------------------------------------------------------
    __global__ void AddVectors(int* a, int* b, int* c) {
        int i = threadIdx.x;
        c[i] = a[i] + b[i];
    }
    
    //----------------------------------------------------------
    extern "C" void exec_cuda(int mpi_rank, int assigned_device) {
    
    // Select the used device:
        if ( cudaSetDevice(assigned_device) != cudaSuccess or
             cudaGetDevice( &used_device ) != cudaSuccess or
             used_device != assigned_device
           ) {
            printf ("Error: unable to set device %d\n", assigned_device);
            return;
        }
    
    // Initialize summands:
        for (int i=0; i<N; i++) {
            HostA[i]=(i+mpi_rank*100)*2;
            HostB[i]=(i+mpi_rank*100)*3;
        }
    
    // Allocate memory on the device:
        cudaMalloc((void**)&DeviceA, N*sizeof(int));
        cudaMalloc((void**)&DeviceB, N*sizeof(int));
        cudaMalloc((void**)&DeviceC, N*sizeof(int));
    
    // Copy summands from host to device:
        cudaMemcpy(DeviceA, HostA, N*sizeof(int), cudaMemcpyHostToDevice);
        cudaMemcpy(DeviceB, HostB, N*sizeof(int), cudaMemcpyHostToDevice);
    
    // Execute kernel:
        AddVectors<<<1, N>>>(DeviceA, DeviceB, DeviceC);
    
    // Copy result from device to host:
        cudaMemcpy(HostC, DeviceC, N*sizeof(int), cudaMemcpyDeviceToHost);
    
    // Show result:
        for (int i=0; i<N; i++) {
            printf ("%d + %d = %d\n",HostA[i],HostB[i],HostC[i]);
        }
    
        cudaFree(DeviceA);
        cudaFree(DeviceB);
        cudaFree(DeviceC);
    }
  • Основное отличие от примера, приведённого выше для CUDA (без MPI) - вместо стандартной функции 'main' используется функция с другим именем и параметрами, описанная как 'extern'.


  • Создать файл 'mpi_part.c':
    #include <mpi.h>
    #include <stdio.h>
    #include <unistd.h>
    #include <stdlib.h>
    
    extern void exec_cuda(int, int);
    
    int *gpu_by_rank;
    int *vnode_is_used;
    
    //------------------------------------------------------------
    int define_gpu (int rank, char *host, int argc, char** argv) {
        int i, gpu = -1;
       
    // 'argv' looks like the following:
    // ./mpi_cuda sl002 0 sl002 1 sl002 2 sl003 0 sl003 1 sl003 2
    
        for (i=1; i<=(argc-1)/2; i++) {
            if ( strcmp(host,argv[i*2-1]) != 0 )
                continue;
            if ( vnode_is_used[i-1] !=0 )
                continue;
            gpu=atoi(argv[i*2]);
            gpu_by_rank[rank] = gpu;
            vnode_is_used[i-1] = 1;
            break;
        }
        return gpu;
    }
    
    //------------------------------------------------------------
    int main(int argc, char** argv)
    {
        int size, rank, i, k, gpu;
        char host[32];
        MPI_Status status;
    
        MPI_Init(&argc,&argv);
        MPI_Comm_size(MPI_COMM_WORLD,&size);
        MPI_Comm_rank(MPI_COMM_WORLD,&rank);
    
    // Validate arguments:
        if ( size != (argc-1)/2 ) {
            MPI_Finalize();
            printf ("Error: amount of GPUs do not match the MPI size!\n");
            return 100;
        }
    
    // Define correspondence between ranks and vnodes gpu:
        if ( rank == 0 ) {
            gpu_by_rank = (int *)malloc(size*sizeof(int));
            vnode_is_used = (int *)malloc(size*sizeof(int));
            for ( i=0; i<size; i++ ) {
                gpu_by_rank[i]=-1;
                vnode_is_used[i]=0;
            }
            for ( i=1; i<size; i++ ) {
                MPI_Recv(host,32,MPI_CHAR,i,0,MPI_COMM_WORLD, &status);
                gpu = define_gpu(i,host, argc, argv);
                MPI_Send(&gpu,1,MPI_INT,i,0,MPI_COMM_WORLD);
            }
            gethostname(host,32);
            gpu=define_gpu(0,host, argc, argv);
            free(gpu_by_rank);
            free(vnode_is_used);
        } else {
            gethostname(host,32);
            MPI_Send(host,32,MPI_CHAR,0,0,MPI_COMM_WORLD);
            MPI_Recv(&gpu,1,MPI_INT,0,0,MPI_COMM_WORLD, &status);
        }
    
        printf("I'm number %d from %d and I run on host %s, gpu %d.\n",rank,size,host,gpu);
        exec_cuda(rank,gpu);
    
        MPI_Finalize();
        return 0;
    }


  • Скомпилировать до объектных файлов (*.o):
    nvcc -arch=compute_20 -c cuda_part.cu
    mpicc -c mpi_part.c

    и затем собрать в исполняемый 'mpi_cuda':

    mpicc mpi_part.o cuda_part.o -lm -lcudart -L/opt/shared/nvidia/cuda/lib64 -I/opt/shared/nvidia/cuda/include -o mpi_cuda
  • Либо можно использовать 'makefile' следующего содержания:
    mpi_cuda : mpi_part.o cuda_part.o
            mpicc mpi_part.o cuda_part.o -lm -lcudart -L/opt/shared/nvidia/cuda/lib64 -I/opt/shared/nvidia/cuda/include -o mpi_cuda
    
    cuda_part.o : cuda_part.cu
            nvcc -arch=compute_20 -c cuda_part.cu
    
    mpi_part.o : mpi_part.c
            mpicc -c mpi_part.c


  • Создать файл 'submit.sh':
    #!/bin/sh
    
    #PBS -q teslaq
    #PBS -l walltime=0:01:00
    #PBS -l select=2:ngpus=3:ncpus=3:mpiprocs=3:mem=32gb,place=scatter
    
    cd $PBS_O_WORKDIR
    
    vnodes=$(qstat -f $PBS_JOBID|tr -d '\n'' ''\t'|sed 's/Hold_Types.*//'|sed 's/.*exec_vnode=//'|tr -d \(\)|tr + '\n'|sed 's/:.*//'|sort)
    echo "My vnodes:"
    for vnode in $vnodes ; do
        node=$(echo $vnode|sed 's/\[.*//')
        gpu=$(echo $vnode|sed 's/.*\[//'|sed 's/\]//')
        echo " $vnode = Node $node, GPU $gpu"
    done
    echo
    
    ## Replace all '[' and ']' by spaces before passing to program:
    vnodes=$(echo $vnodes|sed 's/\[/ /g'|sed 's/\]/ /g')
    
    export LD_LIBRARY_PATH=/opt/shared/nvidia/cuda/lib64:/opt/shared/nvidia/cuda/lib:$LD_LIBRARY_PATH
    mpirun -x LD_LIBRARY_PATH -hostfile $PBS_NODEFILE ./mpi_cuda $vnodes
  • Передать планировщику:
    qsub submit.sh


  • После завершения файл стандартного вывода будет содержать примерно следующее:
    My vnodes:
     sl002[0] = Node sl002, GPU 0
     sl002[1] = Node sl002, GPU 1
     sl002[2] = Node sl002, GPU 2
     sl003[0] = Node sl003, GPU 0
     sl003[1] = Node sl003, GPU 1
     sl003[2] = Node sl003, GPU 2
    
    I'm number 1 from 6 and I run on host sl002, gpu 0.
    I'm number 2 from 6 and I run on host sl002, gpu 1.
    I'm number 3 from 6 and I run on host sl003, gpu 0.
    I'm number 0 from 6 and I run on host sl002, gpu 2.
    I'm number 4 from 6 and I run on host sl003, gpu 1.
    I'm number 5 from 6 and I run on host sl003, gpu 2.
    0 + 0 = 0
    2 + 3 = 5
    4 + 6 = 10
    6 + 9 = 15
    400 + 600 = 1000
    402 + 603 = 1005
    404 + 606 = 1010
    406 + 609 = 1015
    600 + 900 = 1500
    602 + 903 = 1505
    604 + 906 = 1510
    606 + 909 = 1515
    200 + 300 = 500
    202 + 303 = 505
    204 + 306 = 510
    206 + 309 = 515
    800 + 1200 = 2000
    802 + 1203 = 2005
    804 + 1206 = 2010
    806 + 1209 = 2015
    1000 + 1500 = 2500
    1002 + 1503 = 2505
    1004 + 1506 = 2510
    1006 + 1509 = 2515