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