Лекция 8. Multi-GPU (cuda_ openmp_ mpi) (Лекции)

Лихогруд Николайn.lihogrud@gmail.comЧасть восьмаяИспользование нескольких GPUCUDA+openmpCUDA+MPIP2P обмены между GPUCUDA Context Аналог процесса CPU Выделения памяти, выполнение операций происходит врамках некоторого контекста (=процесса) Отдельное адресное пространство Выделенная память неявноосвобождается приудалении контекста Операции из разныхконтекстов не могутвыполнятся параллельноCUDA context• Адресноепространство• Ресурсы• ОперацииCUDA Context Контексты устройств неявно создаются при инициализацииCUDA-runtime На каждом устройстве создается по одному контексту –«primary-контекст» Все нити программы совместно их используют Инициализация CUDA-runtime происходит неявно, припервом вызове любой функции, не относящейся к Device /Version Management (см.

Toolkit Reference Manual)CUDA Context В каждой нити может быть только один активный контекст вкаждый момент времениcudaSetDevice(n)- переключение между устройствами(=между контекстами)cudaDeviceReset() - уничтожает primary-контекст,активный в данный момент При этом будет освобождена вся память, выделенная вконтексте При необходимости, новый контекст будет неявно созданв дальнейшемCuda Context & cudaStream/Event cudaStream{Event}Create создает соответствующийресурс в активном контексте Если активный контекст отличен от того, в котором созданпоток/событие: Отправление команды в поток вызовет ошибку cudaEventRecord() для события вызовет ошибку cudaEventElapsedTime() вызовет ошибку, еслисобытия созданы в разных контекстахПримерcudaSetDevice(0);cudaStream_t s0;cudaStreamCreate(&s0); // создать поток на device 0cudaSetDevice(1);// переключить контекст на device 1cudaStream_t s1;cudaStreamCreate(&s1); // создать поток на device 1MyKernel<<<100, 64, 0, s1>>>();MyKernel<<<100, 64, 0, s0>>>(); // ошибкаMulti-GPU & single CPU threadCUDAcontext 0CUDAcontext 1CUDAcontext 2setDevice(…)Поток CPU переключается между контекстамиМодельная задачаfloat *devPtr = NULL, *hostPtr = NULL;int n;loadInputData(&n, &hostPtr);cudaHostRegister(hostPtr, n*sizeof(float),cudaHostRegisterDefault);cudaMalloc(&devPtr, n * sizeof(float));cudaMemcpyAsync(devPtr, hostPtr, n*sizeof(float),cudaMemcpyHostToDevice, 0);kernel<<<(n – 1) / 512 + 1, 512>>>(devPtr, n);cudaMemcpyAsync(hostPtr, devPtr, n*sizeof(float),cudaMemcpyDeviceToHost, 0);cudaDeviceSynchronize();Переписываем на multiGPUfloat *hostPtr = NULL;int n, deviceCount;loadInputData(&n, &hostPtr);cudaGetDeviceCount(&deviceCount);float **devPtr = (float **)malloc(deviceCount *sizeof(float *));• Получили число устройств• Выделили массивуказателей на GPU-памятьВыделение памяти Выделение памяти через cudaMalloc* происходит на устройстве,к которому относится активный контекст При определенных условиях память может быть доступна изядер, работающих на других устройствах (peer-to-peer) cudaHostRegister[Alloc](…) лочит(выделяет) память в рамкахактивного контекста Преимущества доступны другим контекстам только если pinnedпамять является portable:cudaHostRegister(ptr, n, cudaHostRegisterMapped |cudaHostRegisterPortable);Выделение памятиint elemsPerDevice = (n – 1) / deviceCount + 1;for(int device = 0; device < deviceCount; device++) {cudaSetDevice(device);cudaMalloc(devPtr + device, elemsPerDevice *sizeof(float));cudaHostRegister(hostPtr + device * elemsPerDevice,elemsPerDevice * sizeof(float),…);}Выделение памятиблокирует хост-нить!• Рассчитали размер подзадач• Выделили / залочили нужныеобъемы в каждом контекстеОтправка командfor(int device = 0; device < deviceCount; device++) {int offset = device * elemsPerDevice;int elemCount = min(n – offset, elemsPerDevice);cudaSetDevice(device);cudaMemcpyAsync(devPtr[device], hostPtr + offset),elemCount * sizeof(float),…, 0);kernel<<<(elemCount – 1)/512 + 1, 512>>>(devPtr[device], elemCount);cudaMemcpyAsync(hostPtr + offset, devPtr[device],elemCount * sizeof(float),…, 0);}• Асинхронно отправляем командына устройства• Кажое GPU работает со своейпорцией данныхСинхронизацияfor(int device = 0; device < deviceCount; device++) {cudaSetDevice(device);cudaDeviceSynchronize();}• Ожидаем завершения всехкоманд на устройствахПочему синхронизацию нужно делать вотдельном цикле?Результат (1.487 с)Комментарий Неблокирующие запуски команд правильнее будетпоместить в цикл с выделением памяти Команды на первом GPU начнут выполняться пока наостальных выделяется памятьfor(int device = 0; device < deviceCount; device++) {cudaSetDevice(device);cudaMalloc(…);cudaHostRegister(…);cudaMemcpyAsync(…); // pinned CPU <-> GPUkernel<<<…>>>(…);cudaMemcpyAsync(); // pinned CPU <-> GPU}Результат (1.255 с)Multi-GPU & multiple CPU threadsCUDAcontext 0CUDAcontext 1Thread 0setDevice(…)Thread 1Thread 2CUDAcontext 2Компиляция Поддержка OpenMP встроена в популярные компиляторы Intel icc/ifort, gcc/gfortran, MS cl, IBM xlc Обычный компилятор компилирует OpenMP директивы ифункции при указании специального флага компиляции(для распознавания директив) и линковки (для линковкиomp-функций) icc –openmp gcc –fopenmp cl -/openmp xlc -qsmpКомпиляция с NVCC$nvcc –Xcompiler flag –arch=sm_20 main.cu Передает компилятору на стадию компиляции и настадию линковки слово командной строки flag,следующее за –Xcompiler$nvcc –Xcompiler –fopenmp –arch=sm_20 main.cu Компиляция CUDA+OpenMP на Linux (gcc)Компиляция с NVCC Можно раздельно:$nvcc –arch=sm_20 kernel.cu$gcc –fopenmp –I/opt/cuda/include main.c$gcc –fopenmp –L/opt/cuda/lib –lcudartmain.o kernel.oПереписываем под OpenMPint deviceCount;cudaGetDeviceCount(&deviceCount);#pragma omp parallel num_threads(deviceCount){int device = omp_get_thread_num();cudaSetDevice(device);cudaMalloc(devPtr + device, elemsPerDevice *sizeof(float));cudaHostRegister(…);… // прочие командыЗапускаем параллельнуюcudaDeviceSynchronize();секцию на нужном числе}нитейРезультат (1.207 с)Multi-GPU & multiple CPU processesCUDAcontext 0CUDAcontext 1Process 0setDevice(…)Process 1Process 2CUDAcontext 2Компиляция mpicc – обертка над хостовым компилятором Задача mpicc – подставить пути к инклюдам ислинковать объектные файлы с MPI-библиотекамиБольше ничего mpicc не делает! Наша цель: скомпилировать MPI хост код с теми же флагами, скоторыми это делает mpicc Скомпилировать device-код при помощи nvcc Cлинковать все с нужными библиотеками MPI/CUDAКак узнать флаги mpicc? При использовании OpenMPI: Вывести флаги компиляции: $mpicc -showme:compile-I/usr/lib/openmpi/include I/usr/lib/openmpi/include/openmpi –pthread Вывести флаги линковки: $mpicc -showme:link-pthread -L/usr/lib/openmpi/lib -lmpi -lopen-rte lopen-pal -ldl -Wl,--export-dynamic -lnsl -lutil -lm ldlКак узнать флаги mpicc? При использовании OpenMPI: Вывести полную строку, вместе с именемиспользуемого компилятора: $mpicc –showmegcc -I/usr/lib/openmpi/include I/usr/lib/openmpi/include/openmpi -pthread L/usr/lib/openmpi/lib -lmpi -lopen-rte -lopen-pal-ldl -Wl,--export-dynamic -lnsl -lutil -lm -ldlЗамена gcc в mpicc на nvcc При использовании OpenMPI компилятор можно задатьчерез переменную окружения OMPI_CC, OMPI_F77,OMPI_CXX, OMPI_FC $OMP_CC=nvcc mpicc –-showmenvcc -I/usr/lib/openmpi/include I/usr/lib/openmpi/include/openmpi -pthread L/usr/lib/openmpi/lib -lmpi -lopen-rte -lopen-pal-ldl -Wl,--export-dynamic -lnsl -lutil -lm –ldlЗамена gcc в mpicc на nvcc Проблема: nvcc сам парсит флагиПоддерживает только простые –L, -l, -c,-g и своисобственные -v, -arch и ничего не знает о –Wl, pthread Специфические флаги компилятора нужно передаватьчерез –Xcompiler <флаг>,… Специфические флаги линковщика ld нужно передаватьчерез –Xlinker <флаг>,…Раздельная компиляция Подставляем флаги компиляции mpicxx в nvcc Линкуем хостовым компилятором, явно подставляяфлаги из mpicxx и nvccMPI_COMPILE_FLAGS = $(shell mpicxx --showme:compile)MPI_LINK_FLAGS = $(shell mpicxx --showme:link)NVCC_LINK_FLAGS = -L/opt/cuda/lib64 -lcudartall: mainnvcc -Xcompiler "\"$(MPI_COMPILE_FLAGS)\"" main.cu –o main.og++ main.o -o main $(MPI_LINK_FLAGS) $(NVCC_LINK_FLAGS)Замена gcc в nvcc на mpicxx Опция –ccbin позволяет задать используемый компилятор$nvcc –ccbin /usr/bin/mpicxx main.cu –o main cudafe разделит код на host-код и device-код Host-код будет скомпилирован в дальнейшем при помощи/usr/bin/mpicxx Объектники будут слинкованы через /usr/bin/mpicxx => всенужные MPI флаги подставятсяПример main.cu#include <mpi.h>#include <iostream>#include <stdio.h>__global__ void kernel(int procnum, int device ) {printf("Hello from DEVICE %d process %d\n",device,procnum);}int main (int argc, char* argv[]){int rank, size;int numDevices = -1;cudaGetDeviceCount(&numDevices);…Пример main.cu (продолжение)…cudaGetDeviceCount(&numDevices);MPI_Init (&argc, &argv);MPI_Comm_rank (MPI_COMM_WORLD, &rank);MPI_Comm_size (MPI_COMM_WORLD, &size);std::cout<<"Hello from HOST #"<< rank << " see "<<numDevices << " devices" << std::endl;cudaSetDevice(rank % numDevices);kernel<<<1,1>>>(rank % numDevices, rank);cudaDeviceSynchronize();MPI_Finalize();return 0;}Компиляция & запуск$nvcc –arch=sm_20 –ccbin mpicxx main.cu$mpirun –n 6 ./a.outHelloHelloHelloHelloHelloHelloHelloHelloHelloHelloHelloHellofromfromfromfromfromfromfromfromfromfromfromfromHOST #3 see 4 devicesHOST #4 see 4 devicesHOST #2 see 4 devicesHOST #5 see 4 devicesHOST #1 see 4 devicesHOST #0 see 4 devicesDEVICE 1 process 1DEVICE 3 process 3DEVICE 5 process 1DEVICE 2 process 2DEVICE 4 process 0DEVICE 0 process 0Практика Переписать свертку/корреляцию с multi-GPUUVA & peer-to-peer При UVA peer-to-peer обмены между памятью разных GPUделаются неявно при использовании обычных функцийcudaMemcpy* dst и src указывают на память на разных устройствах Если UVA не поддерживается или нужно явно указать, чтоэто peer-to-peer копирование, используются функцииcudaMemcpyPeer*peer-to-peer & non-UVA Нужно явно указать номера устройств, междукоторыми происходит обменcudaError_t cudaMemcpyPeer (void* dst,int dstDevice, const void* src, int srcDevice,size_t count )cudaError_t cudaMemcpyPeerAsync (void* dst,int dstDevice, const void* src, int srcDevice,size_t count, cudaStream_t stream=0)Peer vs PeerAsync Обе функции не блокируют хост cudaMemcpyPeer начнется только когда завершатсявсе команды на обоих устройствах (и на активном),отправленные до неё Параллельно с cudaMemcpyPeer не могут выполнятсядругие команды на обоих устройствах (и на активном) cudaMemcpyPeerAsync лишена этих ограниченийP2P примерcudaSetDevice(0); // Переключились на device 0float* p0, *p1;size_t size = 1024 * sizeof(float);cudaMalloc(&p0, size); // Выделили на device 0cudaSetDevice(1); // Переключились на device 1cudaMalloc(&p1, size); // Выделили на device 1cudaSetDevice(0); // Переключились на device 0MyKernel<<<1000, 128>>>(p0); // Запуск на device 0cudaSetDevice(1); // Переключились на device 1cudaMemcpyPeer(p1, 1, p0, 0, size); // Копировать p0 to p1MyKernel<<<1000, 128>>>(p1); // Запуск на device 0Прямые peer-to-peer обменыШина PCIeПрямые peer-to-peer обменыПрямой P2P доступcudaError_t cudaDeviceCanAccessPeer (int* canAccessPeer, int device,int peerDevice ) Если в canAccessPeer записалось «1», то peer-to-peer копирования межу peerDevice иdevice могут выполняться без буферизации на хосте Память выделенная на peerDevice может бытьдоступна напрямую из ядер, работающих на deviceПрямой P2P доступcudaError_t cudaDeviceCanAccessPeer (int* canAccessPeer, int device,int peerDevice ) Tesla series UVA Compute Capability > 2.0Прямой доступ нужно явно включить!cudaError_t cudaDeviceEnablePeerAccess (int peerDevice, unsigned int flags ) Теперь память, выделяемая на peerDevice доступнанапрямую из ядер, запускаемых на активном, акопирования выполняются без участия памяти хоста Вызов включает доступ только в одну сторону.