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

Просмтор этого файла доступен только зарегистрированным пользователям. Но у нас супер быстрая регистрация: достаточно только электронной почты!

Текст из файла

Лихогруд Николай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 доступнанапрямую из ядер, запускаемых на активном, акопирования выполняются без участия памяти хоста Вызов включает доступ только в одну сторону.

Характеристики

Тип файла PDF

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

Например, если Вам нужно распечатать чертёж из автокада, Вы сохраните чертёж на флешку, но будет ли автокад в пункте печати? А если будет, то нужная версия с нужными библиотеками? Именно для этого и нужен формат PDF - в нём точно будет показано верно вне зависимости от того, в какой программе создали PDF-файл и есть ли нужная программа для его просмотра.

Список файлов лекций