Full Text

1. Теория В ноябре 2006 г. NVIDIA выпустила первую в истории ГП GeForce 8800 GTX, построенную на архитектуре CUDA [3]. Для охвата большой аудитории разработчиков NVIDIA взяла стандартный язык C и дополнила его несколькими новыми ключевыми словами, позволяющими задействовать специальные средства, присущие архитектуре CUDA [4]. Важной характеристикой любого вычислительного устройства является его быстродействие, которое определяется количеством выполняемых операций в секунду. Существуют ограничения по размерам транзисторов (на данный момент это 7 нм), и поэтому рост быстродействия устройств идет за счет увеличения параллельно работающих ядер, т.е. через параллелизм. Согласно закону Амдала, который гласит, что в случае, когда задача разделяется на несколько частей, суммарное время ее выполнения на параллельной системе не может быть меньше времени выполнения самого медленного фрагмента; максимальное ускорение, которое можно получить от распараллеливания программ на N процессоров (ядер), определяется по формуле [4] , (1) где P - это часть времени выполнения программы, которая может быть распараллелена на N процессоров. При N →∞ выигрыш стремится к . Во многих приложениях, имеющих возможность параллельной обработки, удалось добиться повышения производительности на несколько порядков. Также «приложения, работающие на графических процессорах NVIDIA, демонстрируют большую производительность в расчете на доллар вложенных средств и на ватт потребленной энергии по сравнению с реализациями, построенными на базе одних лишь центральных процессоров» [3]. Далее рассмотрим системные и программные требования для написания и запуска параллельных вычислений. 2. Данные и методы Для написания кода на языке CUDA C требуются следующие элементы: • графический процессор, поддерживающий архитектуру CUDA; • драйвер устройства NVIDIA; • комплект средств разработки CUDA (CUDA Toolkit SDK); • компилятор языка C. В качестве графического процессора использовалась видеокарта начального уровня 2017 г. выпуска GeForce 1030 c 384 ядрами CUDA. Стоит отметить, что такое же количество ядер было у видеокарты среднего сегмента GeForce 560 Ti 2011 г. выпуска, на которой были проведены первые эксперименты, в том числе с вычислением максимально возможного ускорения относительно вычислений на процессоре AMD Phenom 955 BE. Видеокарты 1030 нет в официальном списке поддерживаемых видеокарт, но тем не менее она работает и поддерживает технологию CUDA. На самой дешевой видеокарте GeForce 210 с использованием ОС Windows 8.0 не удалось скомпилировать проект. Использовался последний на момент написания статьи драйвер устройства 471.41. В качестве средств разработки CUDA использовалась версия 10.2, поддерживаемая Windows 7. Начиная с 11-й версии CUDA SDK требуется операционная система Windows 10. В качестве среды разработки была выбрана бесплатная среда разработки Visual Studio 2019 Community Edition с поддержкой языка C. Архитектуру ГП можно кратко охарактеризовать как «макроархитектуру вычислительного кластера, реализованного в микромасштабе» [5]. В настоящее время ГП может содержать от нескольких сотен - 384 (GeForce 1030) до нескольких тысяч - 10 496 (GeForce RTX 3090) потоковых процессоров, а это значит, что при полной загруженности потоковых процессоров производительность отличается в 27 раз между видеокартами начального и верхнего уровней. Стоит также отметить, что в 2021 г. цены на видеокарты завышены в 1,5-2,5 раза по сравнению с 2017 и 2019 гг. из-за бума майнинга. К особенностям языка CUDA С можно отнести то, что в нем используются два устройства: центральный процессор (host) и графический процессор (device). Программа на CUDA задействует как ЦП, так и ГП. Последовательный код выполняется на центральном процессоре, а для параллельных вычислений соответствующий код выполняется на ГП как набор одновременно выполняющихся нитей. 3. Модель Методы конечных разностей решают дифференциальные уравнения в частных производных путем аппроксимации дифференциальных уравнений по области интегрирования системой алгебраических уравнений. Они являются средством получения численных решений уравнений в частных производных. Наиболее распространенными методами конечных разностей для решения уравнений в частных производных являются: явный метод, неявный метод, методы Кранка - Николсона [6]. Эти методы тесно связаны, но отличаются стабильностью, точностью и скоростью выполнения. Рассмотрим приведенный ниже рис. 1. Рис. 1. Графическое сравнение методов конечных разностей: а - явный метод, б - неявный метод, в - метод Кранка - Николсона, (авторские результаты) Нестационарный перенос тепла теплопроводностью описывается уравнением Фурье-Кирхгофа, записанным в декартовой системе координат: , (2) где c - удельная теплоемкость, Дж/(кг∙К); ρ - плотность, кг/м3; T(x, y, x, ) - температурное поле в трехмерном пространстве; t - время, с; λ - коэффициент теплопроводности, Вт/(м∙К); Q(x, y, x, , T) - мощность внутренних источников тепловыделения, Вт. Уравнение (2) описывает множество вариантов процесса теплопроводности. В качестве примера рассмотрим решение одномерной задачи теплопроводности через изолированный стержень с постоянными коэффициентами. На одной границе стержня поддерживается постоянная температура Tл, на другой границе температура равна температуре окружающей среды. Начальная температура равна Т0, источники тепла внутри пластины отсутствуют. При таких начальных условиях температура будет изменяться вдоль длины стержня. Предположим, что теплофизические характеристики стержня не зависят от текущей температуры, т.е. являются постоянными величинами. Уравнение (2) в данном случае будет выглядеть следующим образом: , (3) где L - длина стержня, м. Приняв , получим уравнение (3) в виде . (4) Начальные и граничные условия запишутся следующим образом: (5) Эту задачу решаем методом конечных разностей на равномерной сетке. Для этого стержень по длине разбиваем на N-1 равных промежутков (рис. 2). Рис. 2. Конечно-разностная сетка: x2, x3,…, xN-1 - координаты внутренних узлов, x1, xN - координаты граничных узлов (авторские результаты) При одностороннем нагреве использовались граничные условия первого рода, когда температура остается постоянной во всем процессе теплообмена. Это может быть осуществлено при искусственном поддержании постоянной температуры. Слева будет изолированная сторона: , (6) с правой стороны будет постоянная поддерживаемая температура: . (7) Заменим дифференциальные операторы в уравнении (4) на их конечно-разностные аналоги. Для неявной схемы получаем (8) В статье [7] предложена модель теплопроводности с использованием клеточных автоматов, которая требует значительных вычислительных ресурсов и может быть реализована на компьютерах, использующих технологию параллельных вычислений. Задача нагрева пластины для явной схемы рассматривалась в статье [8], поэтому код для расчета на центральном процессоре (ЦП) приводиться не будет. Намного интереснее и необычнее выглядит код на графическом процессоре. Общая схема работы представлена на рис. 3. Из статьи [8] был взят код реализации явной схемы на ЦП, который был модифицирован с учетом граничных условий 1-го рода. Реализация расчетов на ЦП и ГП также была взята из статьи [8], поэтому рассмотрим в общем виде функцию ядра, которая рассчитывается на графическом процессоре (рис. 4). На рис. 5 показана иерархия потоков. Всё начинается с нити. Тридцать две нити (thread) объединяются в один пучок (warp) [9]. Блок может включать в себя 1024 нити и обладает разделяемой памятью внутри блока. Блоки объединяются в сетку. Рис. 3. Общая схема работы центрального и графического процессоров (авторские результаты) Рис. 4. Общая схема работы функции ядра на графическом процессоре (авторские результаты) Следует отметить, что функции, вычисляемые на графическом процессоре, начинаются с ключевого слова языка CUDA C _global_ [10]. Ключевое слово _shared_ [11] служит для объявления разделяемой памяти, куда могут обращаться все нити, принадлежащие одному блоку. А ключевое слово _syncthread() позволяет синхронизировать потоки внутри одного блока [4]. Иными словами, прежде чем перейти к новой итерации, расчеты во всех потоках должны быть завершены. Синхронизация необходима в задачах, где данные из одной нити могут обращаться к данным другой нити. В явной, неявной и схеме Кранка - Николсона идет обращение к соседним ячейкам, в данном случае к соседним нитям, поэтому скорость выполнения будет зависеть от самой медленно выполняющейся нити внутри блока. В авторских статьях [12, 13] были проведены исследования, когда нити выполнялись абсолютно параллельно и скорость вычислений на центральном и графическом процессорах была в 60 раз быстрее на графическом ускорителе. Рис. 5. Иерархия потоков и доступ к частям памяти на графическом процессоре (авторские результаты) Из представленных схем самой точной является схема Кранка - Николсона, так как она берет данные как на текущем шаге, так и на предыдущем. Представим листинг ядра: _global_ void implicit(double* ziro, double* res) { double S = 0.025; //Длина пластины double a = 0.0000064; //коэффициент температуропроводности double dt = 0.0003; //шаг по времени double all_time = 20; //общее время double gamma = -0.0000005; //параметр итерационного процесса double eps = 0.000001; //точность вычисления на итерации double dx = S / SIZE; double _f = a / dx / dx; int NT = all_time / dt; //Массив в памяти графической карты _shared_ double arr[SIZE]; //n + 1 _shared_ double arrs[SIZE]; //n _shared_ double arrt[SIZE]; //t _shared_ double _max; //Переменная, которая доступна всем потокам _max = 1; int idx = blockIdx.x * blockDim.x + threadIdx.x; arr[idx] = ziro[idx]; _syncthreads(); arrs[idx] = ziro[idx]; _syncthreads(); arrt[idx] = ziro[idx]; _syncthreads(); //синхронизация потоков*/ //Время for (int j = 0; j < NT; j++) { //Один шаг _max = 1; while (_max > eps) { arr[0] = 0; arrs[0] = 0; arrt[0] = 0; arr[SIZE - 1] = 100; arrs[SIZE - 1] = 100; arrt[SIZE - 1] = 100; //В центре пластины if ((idx > 0) && (idx < SIZE - 1)) { _syncthreads(); arrs[idx] = arr[idx] + gamma * ((arr[idx] - arrt[idx]) / dt - _f * ((arr[idx - 1] - 2 * arr[idx] + arr[idx + 1]) + (arrt[idx - 1] - 2 * arrt[idx] + arrt[idx + 1])) / 2.0); _max = fabsf(arr[1] - arrs[1]); _syncthreads(); if (_max < fabsf(arr[idx] - arrs[idx])) { _max = fabsf(arr[idx] - arrs[idx]); } _syncthreads(); arr[idx] = arrs[idx]; _syncthreads(); } } //Обновление значений массива для следующей итерации arrt[idx] = arr[idx]; _syncthreads(); } res[idx] = arr[idx]; //Копирование ГП->ЦП _syncthreads(); } К отличительным особенностям реализации схемы Кранка - Николсона относятся три массива разделяемой памяти, которые определяются ключевым словом языка CUDA C _shared_. Это массивы для запоминания значений температуры на текущем, следующем и шаге по времени массивах. Также можно выделить переменную max, необходимую для определения сходимости схемы Кранка - Николсона, которая тоже относится к разделяемой памяти. 4. Полученные результаты В табл. 1 представлены результаты сравнения быстродействия алгоритмов на центральном и графическом процессорах. При расчетах на ГП можно указывать такую величину, как количество нитей, используемых в блоках. При этом это количество должно быть кратно 16, согласно литературе [3, 4]. В табл. 1 и на рис. 6 показано время выполнения решения задачи одномерной теплопроводности на центральном процессоре (ЦП) и графическом процессоре (ГП), при этом для графического процессора были рассмотрены варианты с использованием общей памяти и три варианта с количеством выполняемых нитей 1024, 512 и 256 в блоке. Как видно на рис. 6, для всех вычислений наблюдается прямая зависимость между количеством расчетов и временем выполнения как на центральном, так и на графическом процессоре, при этом большее ускорение получается при 512 нитях. Скорость расчетов на графическом процессоре была всего в 4 раза быстрее, чем на центральном процессоре. Таблица 1 Сравнение времени выполнения по явной схеме на центральном и графическом процессорах при 1024 ячейках (авторские результаты) Заданное время расчета, с ЦП, мс ГП (1024), мс ГП (512), мс ГП (256), мс Количество измерений ГП(512)/ЦП 10 124 46 38 36 10 000 3,26 100 1134 415 311 243 100 000 3,65 500 5412 1587 1230 1210 500 000 4,40 1000 11105 3270 2562 2544 1 000 000 4,33 1500 16010 4747 3689 3757 1 500 000 4,34 2000 21235 6330 4978 4835 2 000 000 4,27 2500 26644 8049 6318 6188 2 500 000 4,22 Следующим этапом было увеличение количества расчетных ячеек, в табл. 1 показаны расчеты для 1024 ячеек. Количество ячеек увеличили в 2 раза и получили результаты, представленные в табл. 2. Рис. 6. Время выполнения на центральном и графическом процессорах (авторские результаты) Таблица 2 Сравнение времени выполнения по явной схеме на центральном и графическом процессорах при 2048 ячейках (авторские результаты) Заданное время расчета, с ЦП, мс ГП (512), мс ГП (256), мс Количество измерений ГП(512)/ЦП 10 512 78 72 20 000 6,56 100 4363 673 495 200 000 6,48 500 21305 2666 2594 1 000 000 7,99 1000 42395 5437 4909 2 000 000 7,80 1500 63649 7995 7462 3 000 000 7,96 Как видно из табл. 2, быстрота расчетов увеличилась почти вдвое и максимальное значение составило 7,96. Увеличение количества ячеек позволило использовать потенциал графической карты и в следующей табл. 3 составило 4096. Таблица 3 Сравнение времени выполнения по явной схеме на центральном и графическом процессорах при 4096 ячейках (авторские результаты) Заданное время расчета, с ЦП, мс ГП (512), мс ГП (256), мс Количество измерений ГП(512)/ЦП 10 4893 323 318 100 000 15,39 100 43 148 3343 3189 1 000 000 13,53 200 85 700 6424 6517 2 000 000 13,15 300 128 236 9782 9557 3 000 000 13,42 Как видно из табл. 3, быстрота расчетов увеличилась и составила в среднем 13,4. Следующая попытка при объеме данных 8192 ячейки не дала значительного прироста по скорости и составила в среднем 14,54 раза по сравнению с теми же расчетами на центральном процессоре. Поскольку явная схема не является самой точной, были также рассмотрены неявная схема и схема Кранка - Николсона. Тривиальный пример с расчетом экспоненциальной функции представлен на рис. 7. Отличительной особенностью реализации неявной схемы и схемы Кранка - Николсона является то, что требуется три одномерных массива для запоминания значений температуры на текущем, следующем шагах и массив для схождения схемы. При этом эти массивы должны располагаться в общей памяти, а также одно значение для запоминания текущего максимального значения. В графических процессорах существуют ограничения по объему такой памяти. Для ГП 1030 этот объем составляет 49 152 байта. Если использовать вещественный тип данных с одинарной точностью (float), одно значение которого занимает 4 байта, то можно использовать 12 288 значений. Если используется вещественный тип данных с двойной точностью (double), объем одного значения 8 байт, следовательно, получится в 2 раза меньше значений. Также требуется значение для поиска максимума, поэтому в расчетах не превышает 4096 - 1 значений для float и 2048 - 1 значений для double. Рис. 7. Сравнение точности разных схем (авторские результаты) Для неявной схемы при 4000 значениях получились результаты расчетов, представленные в табл. 4. Таблица 4 Сравнение времени выполнения по неявной схеме на центральном и графическом процессорах при 4000 ячеек (авторские результаты) Время, с ЦП ГП (256) Ускорение 10 66 702 33 287 2,003 85 20 72 844 55 908 1,302 93 30 76 921 78 626 0,978 32 Как видно из табл. 4, при увеличении времени расчета центральный процессор начинает опережать графический процессор. Это показательный пример того, что не во всех случаях графический процессор будет быстрее. Например, с ростом числа расчетных ячеек графический процессор работает быстрее. Это было видно по явной схеме, а также по схеме Кранка - Николсона. Поскольку самой точной является схема Кранка - Николсона [6, 14], ей было уделено больше внимания. При 1024 ГП оказался быстрее от 3 до 4,2 раза. При 2000 - от 7,2 до 8,3, а самые лучшие результаты отражены в табл. 5. Таблица 5 Сравнение времени выполнения по схеме Кранка - Николсона на центральном и графическом процессорах при 4000 ячеек (авторские результаты) Время, с ЦП, мс ГП (400), мс Ускорение 10 260 612 19 389 13,4412 20 263 631 25 723 10,2488 30 284 941 33 277 8,5627 Заключение В статье рассмотрены алгоритмы реализации одномерной задачи теплопроводности на центральном и графическом процессорах. При обработке данных от 1000 до 4000 ячеек на графическом процессоре получается большая производительность от 0,97 до 13,5 раза. Наибольшая производительность показывается при 4000 ячеек, так как на графическом процессоре задача рассчитывается параллельно. При увеличении количества шагов по явной схеме замедления расчетов по сравнению с центральным процессором не наблюдается. При реализации схем неявной и Кранка - Николсона наблюдается замедление в ускорении вычислений. В дальнейшем планируется применить данную методику и алгоритм параллельного программирования для двухмерной и трехмерной задачи нагрева угольного сырья [15].

About the authors

P. A Sechenov

Siberian State Industrial University

I. A Rybenlko

Siberian State Industrial University

References

Statistics

Views

Abstract - 112

PDF (Russian) - 30

Refbacks

• There are currently no refbacks.