Лабораторная работа Графические системы. Лабораторная графические системы. Параллельное программирование для графического процессора в среде nvidia cuda
 Скачать 163 Kb.
|
|
Министерство Образования и Науки РФ Новосибирский Государственный Технический Университет Кафедра Вычислительной техники Лабораторная работа №1 По дисциплине «Графические Системы» На тему «Параллельное программирование для графического процессора в среде NVidia CUDA» Факультет: АВТФ Группа: АВТ-816 Студенты: Грохотов С. О., Каманин А. А. Преподаватель: Ильиных С. П. Новосибирск 2019 г. Цель работы: изучить архитектуру CUDA и основы разработки параллельных программ для совместного использования CPU и GPU, написать и отладить на языке Си параллельную программу решения поставленной задачи на графическом процессоре: Нахождение определителя квадратной матрицы.Архитектура CUDACUDA – это архитектура параллельных вычислений от NVIDIA, позволяющая существенно увеличить вычислительную производительность благодаря использованию GPU (графических процессоров).CUDA SDK позволяет программистам реализовывать на специальном упрощённом диалекте языка программирования Си алгоритмы, выполнимые на графических процессорах NVIDIA, и включать специальные функции в текст программы на Cи. Архитектура CUDA даёт разработчику возможность по своему усмотрению организовывать доступ к набору инструкций графического ускорителя и управлять его памятью.Технология CUDA – это программно-аппаратная вычислительная архитектура NVIDIA, основанная на расширении языка Си, которая даёт возможность организации доступа к набору инструкций графического ускорителя и управления его памятью при организации параллельных вычислений. CUDA помогает реализовывать алгоритмы, выполнимые на графических процессорах видеоускорителей GeForce восьмого поколения и старше (серии GeForce 8, GeForce 9, GeForce 200), а также Quadro и Tesla. Графический процессор представляется в виде набора независимых мультипроцессоров(multiprocessors). Каждый мультипроцессор состоит из нескольких CUDA-ядер(CUDA cores), нескольких модулей для вычисления математических функций (SFU), конвейера, а также разделяемой памяти (shared memory) и, кэша (для определенных видов памяти). Графический процессор основан на так называемой архитектуре SIMT (Single Instruction, Multiple Thread). Технология CUDA позволяет определять специальные функции – ядра (kernels), которые выполняются параллельно на CPU в виде множества различных потоков (threads). Таким образом, ядро является аналогом потоковой функции. Каждый поток исполняется на одном CUDA-ядре, используя собственный стек инструкций и локальную память. Программы для CUDA (соответствующие файлы обычно имеют расширение .cu) пишутся на «расширенном» С и компилируются при помощи команды nvcc. Вводимые в CUDA расширения языка С состоят из спецификаторов функций, показывающих где будет выполняться функция и откуда она может быть вызвана; спецификаторы переменных, задающие тип памяти, используемый для данной переменных; директива, служащая для запуска ядра, задающая как данные, так и иерархию потоков; встроенные переменные, содержащие информацию о текущем потоке; runtime, включающий в себя дополнительные типы данных Таблица 1. Спецификаторы функций в CUDA
При этом спецификаторы __host__ и __device__ могут быть использованы вместе (это значит, что соответствующая функция может выполняться как на GPU, так и на CPU - соответствующий код для обеих платформ будет автоматически сгенерирован компилятором). Спецификаторы __global__ и __host__ не могут быть использованы вместе. Спецификатор __global__ обозначает ядро и соответствующая функция должна возвращать значение типа void. Метод распараллеливания и алгоритм программы Содержание оперативной памяти (матрица размерности N*N типа float; вектор полей длины N float (заполненный единицами), произведение элементов которых даст определитель матрицы; булевская переменная для определения надобности продолжения алгоритма) копируется в глобальную память GPU. Затем происходит цикл, каждый шаг которого охватывает квадратную подматрицу исходной матрицы, расположенной в нижней-правой части. Изначально подматрица равна матрице, при каждом последующем шаге цикла линейный размер Hstr подматрицы будет уменьшаться на 1, пока не станет равным 1. Все одновременные действия будут выполняться в индивидуальном ядре как одномерное множество потоков. На каждом шаге цикла происходит следующая последовательность действий над подматрицей: Все крайние левые элементы «подстроки» длины Hstr подматрицы последовательно проверяются на равенство 0. Если все элементы равны 0, то выполнение программы прекращается, определитель равен 0; В каждой подстроке, обрабатываемой одновременно с остальными: На крайний левый элемент, если он не равен 0 или 1, последовательно делятся все остальные элементы; Элемент массива определителя исходной матрицы (номер которого равен номеру строки в исходной матрице, в которой находится данная подстрока) помножается на значение крайнего левого элемента; Крайний левый элемент, если он не равен 0 или 1, принимает значение 1. В каждом столбце подматрицы происходит последовательное вычитание крайнего верхнего элемента из остальных нижестоящих элементов; Результат – все крайние левые элементы подматрицы, кроме верхнего со значением 1, равны 0, следовательно, определитель подматрицы равен определителю той же подматрицы без всех крайних верхних и крайних левых элементов. Пройдя весь цикл до конца мы помножаем последний элемент вектора определителя на крайний нижний правый элемент исходной матрицы приступаем к обработке массива элементов определителя исходной матрицы, произведение которых даст нам сам определитель. Для этого проведём циклическую операцию, при которой каждый чётный не обработанный элемент (начиная с 0-го) будет перемножаться с соседним правым нечётным не обработанным элементом, пока не останется всего 1 необработанный элемент. Каждый нечётный элемент, на который был помножен соседний левый чётный член, считается обработанным, поэтому при каждом шаге цикла список чётных и нечётных необработанных элементов уменьшается не более чем в 2 раза. Нулевой оставшийся необработанный элемент будет иметь значение определителя матрицы. Алгоритм CPU Вычисление определителя для CPU организовано с помощью метода Гаусса, при котором необходимо привести матрицу к треугольному виду. Алгоритм заключается в следующем: 1. Разделим элементы каждой строки на первый элемент соответствующей строки. 2. Вычтем из элементов всех строк, начиная со второй, элементы первой строки. 3. Повторим данный алгоритм для внутренних строк. Продолжать будем до тех пор, пока размер конечной матрицы не станет размера 1x1. В ходе тестирования обоих программ выяснились следующие вещи: Как и в CPU, так и в GPU имелись отклонения получаемого ответа от истинного (для GPU в меньшей степени, в то время как в CPU при размере матрицы в 7 элементов отклонение от истины составляло единицы процентов. Возможно, это объясняется более качественной реализацией машинной обработки операций с типом float, чем в центральном процессоре); Также CPU выходило из строя при размере обрабатываемой матрицы более 750, в то время как размер матрицы, без ошибок обрабатываемой в GPU, составлял 3000. В отличие от GPU, для CPU серьёзной проблемой был замер времени работы программы, так как это время при доступных размерностях матрицы возрастало крайне незначительно. В то же время временные промежутки для GPU считывались с высокой точностью в доли миллисекунд в широком диапазоне размерностей матрицы. Результаты работы Таблица 2. Зависимость времени выполнения от размера матрицы  График зависимости времени выполнения от размера матрицы на CPU и GPU Как видно из диаграммы выше, при малых значениях размера матрицы CPU работает быстрее, т.к. в CUDA происходит копирование данных. Но при размере матрицы около 800 работа CPU начинает стремительно замедляться, а работа CUDA значительно не изменяется и при больших значениях размера матрицы работает значительно быстрее, чем CPU. Код программы CUDA #include #include #include #include #include "time.h" #include #define N 20 float HostA[N*N]; float HostB[N]; bool HTrue; float *DeviceA; float *DeviceB; bool *DTrue; int Hstep; int*Dstep; __global__ void Step0(float* a,int H,int NN,bool *B,int*st) { int ii,s=(NN+1)*(NN-H); // s – смещение в столбцах от начала матрицы (Проверяем с нулевой строки) (*st)++; for(ii = 0; (ii < H) && (a[ii*NN+s] == 0); ii++)(*st)++; if(ii==H) { *B=true;(*st)++;// Если все элементы равны нулю, то завершаем } } __global__ void Step1(float* a,float* b,int NN,int*st) { int ii = threadIdx.x;// Номер потока (индекс) if (ii==0)*st++; int H=blockDim.x; // Размерность блока (смещение в столбцах) if (ii==0)(*st)++; int s=(NN+1)*(NN-H); // Смещение в строках if (threadIdx.x==0) (*st)++; float a0 = a[NN*ii+s]; // Крайний левый элемент всех строк(так как ii = threadIdx.x) то здесь и идет распаралеливание if (ii==0)(*st)++; if((a0 != 0) && (a0 != 1)) { for (int i0 = 1; i0 < H; i0++) { a[s+NN*ii+i0]/=a0; // Делим все строки на крайний левый(у каждого свой) if (ii==0)(*st)++; } b[ii] *= a0; // Умножаем всю строку матрицы определителя на левое значение if (ii==0)(*st)++; a[NN*ii+s] = 1; // Все крайние левые элементы - единицы if (ii==0)(*st)++; } } __global__ void Step2(float*a,int NN,int*st) { int ii = threadIdx.x; /*номер столбца*/ if (ii==0)(*st)++; int H=blockDim.x; if (ii==0)*st++; /*высота*/ int s=(NN+1)*(NN-H); if (ii==0)*st++; for (int i0=1; i0 { a[s+ii+NN*i0]-=a[ii+s]; // Вычитаем из всех столбцов верхний элемент if (ii==0)*st++; } } __global__ void Step3(float*a,float*b,int NN,int*st) { b[NN-1] *= a[NN*NN-1]; // Умножаем элемент определителя на элементы главной диагонали if (threadIdx.x==0)(*st)++; } __global__ void ProB (float*b, int sm, int NN,int*st) { int a = sm*threadIdx.x; if (threadIdx.x==0) (*st)++; int a0 = a + sm/2; if (threadIdx.x==0)(*st)++; if (a0 < NN) { b[a]*=b[a0]; if (threadIdx.x==0)(*st)++; } } __global__ void STPP (int*st){(*st)++;} int main() { Hstep=1; cudaEvent_t start, stop; float time; cudaEventCreate(&start); cudaEventCreate(&stop); int Hstr,i,j; for (i=0; i for (j=0; j { HostA[i*N+j]=(float)(((int)(j*i*j*i*j+i*i*i+j+7))%10); } for(i = 0; i < N; i++) HostB[i]=1; HTrue = false; cudaEventRecord(start,0); cudaMalloc((void**)&Dstep, sizeof(int)); cudaMalloc((void**)&DeviceA,N*N*sizeof(float)); cudaMalloc((void**)&DeviceB,N*sizeof(float)); cudaMalloc((void**)&DTrue,sizeof(bool)); cudaMemcpy(Dstep, &Hstep, sizeof(int),cudaMemcpyHostToDevice); cudaMemcpy(DeviceA,HostA, N*N*sizeof(float),cudaMemcpyHostToDevice); cudaMemcpy(DeviceB,&HostB, N*sizeof(float),cudaMemcpyHostToDevice); cudaMemcpy(DTrue, &HTrue, sizeof(bool), cudaMemcpyHostToDevice); for (Hstr=N; Hstr>1; Hstr--) // Hstr - Размерность подматрицы { Step0<<<1,1>>>(DeviceA,Hstr,N,DTrue,Dstep); // Проверка на нулевую продолжимость cudaMemcpy(&HTrue, DTrue, sizeof(bool), cudaMemcpyDeviceToHost); { if(HTrue)break;STPP<<<1,1>>>(Dstep); } Step1<<<1,Hstr>>>(DeviceA,DeviceB,N,Dstep); // Заединичивание строк Step2<<<1,Hstr>>>(DeviceA,N,Dstep); //Вычитание } cudaMemcpy(HostB, DeviceB, sizeof(float)*N, cudaMemcpyDeviceToHost); if(Hstr==1) { Step3<<<1,1>>>(DeviceA,DeviceB,N,Dstep); cudaMemcpy(&HostB, DeviceB, sizeof(float), cudaMemcpyDeviceToHost); } else { HostB[0]=0;STPP<<<1,1>>>(Dstep);} // Последнее число cudaMemcpy(HostB, DeviceB, sizeof(float)*N, cudaMemcpyDeviceToHost); for(int s = 2; (s >> 1) < N; s <<= 1) { int y=(N / s) + (N % s) / ( (s >> 1) + 1); ProB<<<1,y>>>(DeviceB,s,N,Dstep); } cudaMemcpy(HostB, DeviceB, sizeof(float)*N, cudaMemcpyDeviceToHost); cudaMemcpy(&Hstep, Dstep, sizeof(int), cudaMemcpyDeviceToHost); cudaFree(DeviceA); cudaFree(DeviceB); cudaFree(DTrue); cudaFree(Dstep); cudaEventRecord(stop,0); cudaEventSynchronize(stop); cudaEventElapsedTime(&time, start, stop); printf("%f\n",HostB[0]); printf("%f (ms)\n",time); printf("%d (steps)\n",Hstep); getch(); return 0; } Код программы на C++ (CPU) // Det.cpp: определяет точку входа для консольного приложения. // #include "stdafx.h" #include #include #include using namespace std; #pragma comment(lib,"Winmm") float detCalculate(float **&M,int n){ int CurrStroke = 0; int CurrStolbec = 0; float count = 1; float delitel = 0; int normStolb = -6; int normStroke = -6; float yj = 0; //Деление строки for(int z = 0; z < n; z++){ for(int j = CurrStroke; j < n; j++){ delitel = M[j][CurrStolbec]; if(delitel != 0) count *= M[j][CurrStolbec]; if(M[CurrStolbec][CurrStroke] == 0){ for(int k = 0; k < n; k++){ if(M[k][CurrStolbec] !=0){ yj = 7; } } if(yj!=7)return 0; normStolb = -6; yj = 0; for(int g =0; g if(M[CurrStroke][g] != 0){ yj = 1; } if(yj == 1){ normStolb = g; } } if(normStolb < 0){ return 0; } for(int z = 0; z if(M[h][CurrStolbec] == 0){ count = -count; for( int b = 0; b < n; b++){ float temp = M[b][CurrStolbec]; M[b][CurrStolbec] = M[b][normStolb]; M[b][normStolb] = temp; } h = n+1; z = n+1; break; } } } } for(int i = CurrStolbec; i < n; i++){ if(delitel != 0 ) M[j][i] /= delitel; if( j != CurrStroke && delitel != 0 ){ M[j][i] -= M[CurrStroke][i]; } } } CurrStroke++; CurrStolbec++; } return count; } int _tmain(int argc, _TCHAR* argv[]) { //Инициализация int n = 1000; float **Matrix = new float* [n]; for(int i = 0; i < n; i++){ Matrix[i] = new float [n]; } for(int i = 0; i < n; i++) for(int j = 0; j < n; j++){ Matrix[i][j] = rand()%100 +1; } for(int i = 10; i < 1000; i+=100){ n = i; float startTime = timeGetTime(); detCalculate(Matrix,1000); float delta = (timeGetTime() - startTime); cout << delta << " "; } for(int i =0; i < 1000; i++){ delete [] Matrix[i]; } getchar(); getchar(); return 0; } Вывод: в ходе лабораторной работы №1 мы познакомились с параллельным программированием на графическом процессоре NVidia CUDA. По результатам данной работы, можно сделать вывод, что на больших объемах данных GPU работает в разы быстрее, чем CPU, за счет распараллеливания вычислений. Но на малых объемах данных, CPU выигрывает за счет необходимости копирования данных из глобальной памяти GPU в регистровую (но самих данных для обработки немного). |