CUDA

CUDA представляет из себя модификацию языка C, поэтому для исходных файлов принято использовать специальное расширение - 'cu' вместо 'c'.

Основные термины:

host - компьютер 'в обычном понимании', управляемый CPU.
device (устройство) - карта с GPU.
kernel (ядро) - функция, которая будет запущена в нескольких экземплярах, каждый из которых будет работать на своём ядре устройства. Для указания, что функция будет ядром, при её описании используется спецификатор
```
___global___
```
Выполнение ядра с распараллеливанием на N потоков описывается следующим образом:
```
MyKernel<<<1, N>>>(параметры);
```

Используемое ПО

CUDA Toolkit

CUDA Toolkit предоставляет полную среду разработки на C и C++ с использованием вычислений на GPU NVIDIA. Включает компилятор для GPU, инструменты для отладки и оптимизации, математические библиотеки и документацию.
Разные версии CUDA Toolkit установлены в поддиректории внутри директории '/opt/shared/nvidia'. Например, версия 8.0 установлена в '/opt/shared/nvidia/cuda-8.0'. ~~При этом символьная ссылка '/opt/shared/nvidia/cuda' указывает на директорию с последней установленной версией.~~
Для использования необходимо настроить переменные окружения следующим образом:
```
export PATH=$PATH:/opt/shared/nvidia/cuda/bin
export LD_LIBRARY_PATH=/opt/shared/nvidia/cuda/lib64:$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 (в зависимости от версии).

Пример

Ниже приведён пример запуска программы, складывающей средствами 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:$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

MPI + CUDA

Для использования 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

MPI_NP=$(wc -l $PBS_NODEFILE | awk '{ print $1 }')

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:$LD_LIBRARY_PATH
mpirun -x LD_LIBRARY_PATH -hostfile $PBS_NODEFILE -np $MPI_NP ./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