Программирование для GPU с использованием NVidia CUDA.

Программирование для GPU с использованием NVidia CUDA. Половинкин А.Н.

Содержание • Постановка задачи • Алгоритм умножения матриц на GPU • Программная реализация

Постановка задачи • C = A*B • A – прямоугольная матрица • число строк – hA • число столбцов – wA • B – прямоугольная матрица • число строк – hB • число столбцов - wB

Алгоритм умножения матриц на GPU (1) Bsub Asub

Алгоритм умножения матриц на GPU (2) • Каждый блок потоков занимается вычислением одной подматрицыCsubматрицы С • Каждый поток внутри блока потоков занимается вычислением одного элемента подматрицы Csub

Алгоритм умножения матриц на GPU (3) • Csub = Asub * Bsub • Asubразмерности (block_size, wA) • Bsubразмерности (wA, block_size) • Матрицы Asub и Bsubв общем случае могут не помещаться в общей памяти устройства, что приведёт к потере производительности • Решение: разбить AsubиBsubна блоки Asub,iиBsub,iразмерности (block_size, block_size), вычислять Csubкак сумму произведений этих блоков:

Алгоритм умножения матриц на GPU (4) • for i = 1 to wA/block_size • загрузить Asub,iи Bsub,iиз глобальной (device) памяти в общую (shared) памятьблока потоков. Каждый поток загружает «свой» элемент! • каждый поток вычисляет «свой» элемент в произведении Asub,i*Bsub,iи сохраняет аккумулированное значение • end • каждый поток загружает «свой» вычисленный элемент в глобальную (device) память

Структура проекта • matrixMul.h– содержит определения (через define) размера блока и размеров матриц • matrixMul_gold.cpp • computeGold • matrixMul.cu • main • randomInit • printDiff • runMultiplication • matrixMul_kernel.cu • matrixMul (kernel)

Реализация функции-ядра (1) • Используется схема хранения матриц по строкам в виде одномерного массива • __global__ voidmatrixMul( float* C, float* A, float* B, intwA, intwB) • вычисляем координаты текущего блока потоков и сохраняем их в переменные bx, by intbx = blockIdx.x int by = blockIdx.y • вычисляем координаты текущего потока в блоке потоков и сохраняем их в переменные tx, ty

Реализация функции-ядра (2) • вычисляем индекс элемента в массиве, хранящем исходную матрицу A, который соответствует первому элементу первой обрабатываемой подматрицы Asub,1 intaBegin = wA * BLOCK_SIZE * by; • вычисляем шаг для перехода к первому элементу следующей обрабатываемой подматрицы intaStep = BLOCK_SIZE; • вычисляем индекс (условие останова перебора подматриц Asub,i) intaEnd = aBegin + wA - 1;

Реализация функции-ядра (3) • вычисляем индекс элемента в массиве, хранящем исходную матрицу B, который соответствует первому элементу первой обрабатываемой подматрицы Bsub,1 intbBegin = … • вычисляем шаг для перехода к первому элементу следующей обрабатываемой подматрицы intbStep = … • объявляем переменную, в которой хранится элемент произведения, вычисляемый текущим потоком float Csub = 0.f;

Реализация функции-ядра (4) • В цикле по всем подматрицам Asub,i, Bsub,iвыполняем следующие действия for (int a = aBegin, b = bBegin;a <= aEnd; a += aStep, b += bStep) {… } • объявляем в общей (shared) памяти подматрицы Asub,iи Bsub,i __shared__ float As[BLOCK_SIZE][BLOCK_SIZE]; __shared__ float As[BLOCK_SIZE][BLOCK_SIZE]; • копируем из глобальной памяти в общую элементы соответствующих подматриц (каждый поток копирует «свой» элемент) • синхронизируем все потоки в блоке потоков __syncthreads();

Реализация функции-ядра (5) • каждый поток вычисляет свой элемент Csubв произведении подматриц Asub,i*Bsub,i • синхронизируем потоки в блоке потоков • конец цикла • загрузить вычисленный элемент Csubв соответствующий элемент матрицы С (в глобальную память) intc = wB * BLOCK_SIZE * by + BLOCK_SIZE * bx; C[c + wB * ty + tx] = Csub;

Реализация функции runMultiplication (1) • void runMultiplication(intargc, char** argv) • инициализируем устройство (device) CUT_DEVICE_INIT(argc, argv); • выделяем память на хосте для хранения матриц A и B unsigned intsize_A = WA * HA; unsigned intmem_size_A = sizeof(float) * size_A; float* h_A = (float*)malloc(mem_size_A); unsigned intsize_B = WB * HB; unsigned intmem_size_B = sizeof(float) * size_B; float* h_B = (float*)malloc(mem_size_B);

Реализация функции runMultiplication (2) • инициализируем матрицы A и B случайными значениями randomInit(h_A, size_A); randomInit(h_B, size_B); • выделяем память под матрицы A и B на устройстве, копируем данные с хоста на устройство float* d_A; CUDA_SAFE_CALL(cudaMalloc((void**)&d_A, mem_size_A)); float* d_B; CUDA_SAFE_CALL(cudaMalloc((void**)&d_B, mem_size_B)); CUDA_SAFE_CALL(cudaMemcpy(d_A, h_A, mem_size_A, cudaMemcpyHostToDevice) ); CUDA_SAFE_CALL(cudaMemcpy(d_B, h_B, mem_size_B, cudaMemcpyHostToDevice) ); • выделяем память под матрицу C на хосте и на устройстве (имена соответствующих переменных h_C, d_C, size_C, mem_size_C)

Реализация функции runMultiplication(3) • создаем и инициализируем таймер unsigned int timer = 0; CUT_SAFE_CALL(cutCreateTimer(&timer)); CUT_SAFE_CALL(cutStartTimer(timer)); • определяем конфигурацию выполнения ядра (размер решетки блоков потоков и блока потоков) dim3 threads(BLOCK_SIZE, BLOCK_SIZE); dim3 grid(WC / threads.x, HC / threads.y); • запускаем ядро • копируем вычисленную матрицу С с устройства на хост

Реализация функции runMultiplication(4) • останавливаем таймер, выводим время вычислений, освобождаем ресурсы таймера CUT_SAFE_CALL(cutStopTimer(timer)); printf("Processing time: %f (ms) \n", cutGetTimerValue(timer)); CUT_SAFE_CALL(cutDeleteTimer(timer)); • вычисляем то же самое произведение на CPU float* reference = (float*) malloc(mem_size_C); computeGold(reference, h_A, h_B, HA, WA, WB);

Реализация функции runMultiplication(5) • сравниваем результат, полученный на GPU, с результатом, полученным на CPU (по евклидовой норме) CUTBooleanres = cutCompareL2fe(reference, h_C, size_C, 1e-6f); printf("Test %s \n", (1 == res) ? "PASSED" : "FAILED"); if (res!=1) printDiff(reference, h_C, WC, HC); • освобождаем память

Литература • Nvidia CUDA Programming Guide • Многочисленные курсы по CUDA: • http://courses.ece.uiuc.edu/ece498/al1/Syllabus.html • http://www.nvidia.ru/object/cuda_state_university_courses_ru.html(на русском языке)

Вопросы ?

Программирование для GPU с использованием NVidia CUDA.