Моделирование газодинамических течений с использованием технологии cuda

ЖАЛНИН Р. В., ЛЕЩАНКИНА Т. М.

МОДЕЛИРОВАНИЕ ГАЗОДИНАМИЧЕСКИХ ТЕЧЕНИЙ С ИСПОЛЬЗОВАНИЕМ ТЕХНОЛОГИИ CUDA Аннотация. В статье описан численный алгоритм решения уравнений газовой динамики с использованием технологии CUDA. Для численного эксперимента были смоделированы условия развития неустойчивости при многократном прохождении ударной волны через контактный разрыв с начальными данными Погги. Результаты численного исследования демонстрируют возможности разработанной параллельной версии программы.

Ключевые слова: уравнения газовой динамики, параллельный алгоритм, технология CUDA.

ZHALNIN R. V., LESCHANKINA T. M.

MODELING OF GAS DYNAMIC FLOWS USING CUDA TECHNOLOGY Abstract. This article describes a numerical algorithm for solving equations of gas dynamics using CUDA technology. The Poggie's conditions of instability by repeated passage of a shock wave through a contact discontinuity were simulated for the numerical experiment. The experiment results demonstrate the capabilities of the developed parallel version of the program. Keywords: gas dynamics equations, parallel algorithm, CUDA technology.

Введение. В настоящее время большой практический интерес представляют решения задач газовой динамики. Данного вида задачи используются при проектировании самолетов и автомобилей, предсказании погоды, моделировании климата и т. п. Из-за нелинейности уравнений газовой динамики, для их решения необходимо использовать численные методы, которые на данный момент являются универсальными методами решения этих уравнений. В последние годы было проведено множество исследований в сфере разработки эффективных разностных схем высокого разрешения, которые позволяли бы осуществить математическое моделирование вышеуказанных процессов. Реализация данного рода задач на подобных сетках подразумевает использование больших вычислительных ресурсов. И появление новых технологий, ориентированных на использование в своей разработке графических ускорителей в дополнение к уже существующим позволяет использовать новые возможности для проведения интенсивных вычислений. Увеличение производительности видеокарт привело к тому, что количество вычислительных ядер в них достигло нескольких сотен. На сегодняшний день ведущей технологией разработки для работы на GPU (Graphical Processor Unit) является технология CUDA.

CUDA (Compute Unified Device Architecture) - это программная модель, включающая описание вычислительного параллелизма и иерархичной структуры памяти непосредственно в языке программирования [1]. Эта технология позволяет работать с языками программирования C и FORTRAN. С точки зрения программного обеспечения, реализация CUDA представляет собой кроссплатформенную систему компиляции и исполнения программ, части которых работают на CPU и GPU. CUDA предназначена для разработки GPGPU-приложений без привязки к графическим API и поддерживается всеми GPU NVIDIA.

1. Математическая модель и численная схема. Систему уравнений газовой динамики, записанную в переменных Эйлера, можно представить в виде:

др д(ри) d(pv) dt дх ду

д

dt

дх

ду

d

dt

дх

ду

д(ре) д((ре + р)и) д((ре + р)р) . д1 дх ду

За уравнение состояния принимается уравнение состояния идеального газа с

показателем адиабаты у: р =

р = р(х,1) - плотность газа, V = (и, р) - вектор скорости частиц, р = р(х,1) - давление, £ -

удельная внутренняя энергия на единицу массы, е = е +

u2+v2

полная энергия, у =

2 С„

показатель адиабаты, ср, су - теплоёмкость при постоянном давлении и постоянном объеме соответственно.

В векторной форме система уравнений, принимает вид:

ди дР

dt

+

дх

+

ду

= 0,

где

U =

1Р ри

pv

\ре.

/

,F(U) =

ри ри2 + р puv \(ре + р)и/

/

,G(U) =

pv puv

pv2 + p \(pe + p)v/

Для аппроксимации системы уравнений газовой динамики (2) использовалась нелинейная консервативная квазимонотонная дифференциально-разностная схема [3-5]:

Uij - UH + Fi+i/2j - Pj-!/2j | Gjj+j/2 - Gjj-i/2 _ о

т

где Fi+1/2j, Gij+1/2 - дискретные потоки.

с

р

Дискретные потоки на гранях между ячейками вычислялись как функции двух переменных по интерполированным значениям газодинамических параметров на гранях между ячейками [4; 5]:

F. i = F( Ui,Ur 1

l+2j V i+27 i+27,

G.. i = Y( Ul, 1,Ur. 1

ij+2

ij+2 ij+2,

Здесь ^¿у+1/2, ^¿у+1/2 - "левые" и " правые" значения вектора и на грани между i и i + 1 ячейками, для которой вычисляется поток ^¿+1/2у, ^¿у+1/2, (я = г, I) "левые" и " правые" значения вектора и на грани между ] и ] + 1 ячейками соответственно. Для того чтобы вычислить значения вектора и, на указанных гранях между ячейками введем новый вектор переменных Y = Y(U). Проведем его интерполяцию на грань между ячейками и, затем, пересчитаем искомое значение вектора и = на данной грани. В расчетах полагалось Y = (р, р, и, V), потоки вычислялись по схеме Лакса-Фридрихса-Русанова:

1

F1 =п l+2J 2

S+Г1

-ai( Ur i -U i

V l+27 l+27;

- a2 ( Ur i - U i

\ 1 ^+2

где

«1 = max{| u0-| + ciy, | Uj+iy | + q+iy}, a2 = max{| | + ciy, | viy+i | + Qy+i},

2. Реализация алгоритма на языке CUDA C. Графический ускоритель использовался для проведения расчетов численных потоков на границах ячеек и вычисления ограничителей.

Для проведения расчетов на GPU необходимо выполнить следующие действия: 1. Выделить память для данных и потоков на GPU. Пример: выделяется память под газодинамические параметры:

cudaMalloc((void**)&r_dev, (ntri_new_m[0])*sizeof(float)); cudaMalloc((void**)&p_dev, (ntri_new_m[0])*sizeof(float)); cudaMalloc((void**)&u_dev, (ntri_new_m[0])*sizeof(float)); cudaMalloc((void**)&v_dev, (ntri_new_m[0])*sizeof(float)); где первым параметром идет указатель, в который записывается адрес выделенной памяти, а вторым - размер выделяемой памяти в байтах.

2. Скопировать данные на GPU. Пример: газодинамические параметры (r,p,u,v) копируем на GPU по средствам использования функции:

cudaMemcpy(r_dev, r, (ntri_new_m[0])*sizeof(float), cudaMemcpyHostToDevice); cudaMemcpy(p_dev, p, (ntrinew _m[0])*sizeof(float), cudaMemcpyHostToDevice); cudaMemcpy(u_dev, u, (ntrinew _m[0])*sizeof(float), cudaMemcpyHostToDevice); cudaMemcpy(v_dev, v, (ntri_new_m[0])*sizeof(float), cudaMemcpyHostToDevice); где первым параметром идет указатель, содержащий адрес места-назначения копирования; вторым - указатель, содержащий адрес источника копирования; следом - размер копируемого ресурса в байтах; далее - параметр, указывающий направление копирования.

3. Распределить ядра для вычисления.

Расчет газодинамических параметров состоит из последовательности вызываемых процедур, запускаемых на GPU на каждом шаге по времени:

while (t <= tmax) { boundcond<<<blocks, threads>>> (rdev, pdev, udev, vdev, dldev, d2_dev, d3_dev, d4_dev, d5_dev, d6_dev, otrudev, otruldev, otr_v_dev, otr_v1_dev, alphaldev, alpha2_dev, sootvdev, cm); calcU<<<blocks, threads>>>(r_dev, p dev, u dev, v dev,

rodev, rudev, rvdev, endev, ntridev); ENO_mass<<<blocks, threads>>>(ux_m_dev, uymdev, vxmdev, vymdev,

rxmdev, rymdev, pxmdev, pymdev, ntri dev, trineighldev, cmdev, r dev, u dev, v dev, p dev, ntrinewdev); calcFluxes<<<blocks, threads>>>(r_dev, p dev, u dev, v dev, rtmp_dev, utmp_dev, vtmp_dev, etmp_dev, uxmdev, uy m dev, vx m dev, vy m dev, rx m dev, ry m dev, px m dev, py m dev, xldev, tri neighl dev, tri_v1_dev, cm dev, ntri dev, ntri new dev, nxnewdev); gasDinStep<<<blocks, threads>>>(rn_dev, pndev, undev, vndev, ro dev, ru dev, rv dev, en dev, rtmp_dev, utmp_dev, vtmp dev, etmp dev, s dev, ntri dev); t = t+tau; step++;

}

Процедура boundcond<<<blocks, threads>>>() определяет граничные условия.

4

Процедура calcU<<< blocks, threads>>>() вычисляет консервативные переменные в соответствие с формулой (3).

В процедуре ENO_mass<<<blocks, threads>>>() реализован алгоритм реконструкции газодинамических параметров на границе ячеек.

В процедуре calcFluxes<<<blocks, threads>>>() рассчитываем дискретные потоки по схеме Лакса-Фридрихса в соответствии с формулой (8).

Процедура gasDinStep<<<blocks, threads>>>() рассчитывает значения консервативных, а затем примитивных переменных на каждом шаге по времени.

4. Скопировать полученные данные обратно в CPU. Пример: газодинамические параметры (r,p,u,v) копируем на CPU

cudaMemcpy(r, rndev, (ntri_new_m[0])*sizeof(float), cudaMemcpyDeviceToHost); cudaMemcpy(p, pndev, (ntri_new_m[0])*sizeof(float), cudaMemcpyDeviceToHost); cudaMemcpy(u, undev, (ntri_new_m[0])*sizeof(float), cudaMemcpyDeviceToHost); cudaMemcpy(v, vndev, (ntri_new_m[0])*sizeof(float), cudaMemcpyDeviceToHost);

5. Освободить память GPU. cudaFree(r_dev); cudaFree(pdev); cudaFree(udev); cudaFree(vdev);

где единственным параметром функции является указатель освобождаемой памяти.

3. Вычислительный эксперимент. Большой интерес с практической точки зрения представляет задача взаимодействия ударной волны с зоной турбулентного перемешивания. Такие виды задач являются хорошими тестовыми примерами для проверки различных алгоритмов, методов и программ. На практике такие задачи позволяют изучить ход образования вихрей при контактном взаимодействии нескольких газов с разными плотностями и уравнениями состояний [6].

Выполнено численное исследование динамики развития газодинамической неустойчивости в контактной границе элегаз-воздух для условий эксперимента на ударной трубе [7]. Параметры ударной трубы следующие: длина равна 0,4 м, а ширина равна 0,04 м. Труба разделена на три области:

1. область длиной 0,08 м, заполнена элегазом с высоким давлением (SF6_HI);

2. область длиной 0,02 м, заполнена элегазом с низким давлением (SF6_LO);

3. область длиной 0,03 м, заполнена воздухом (Л/ ft).

В качестве расчетных данных была принята следующая физическая постановка задачи: начальная температура 291 К в камере низкого давления давление составляет 10-4

ГПа = 1 бар. Непосредственно за ударной волной в элегазе давление составляет 2,152 бар, плотность - 1,209 * 10-2 г/см3, скорость ударной волны - 195,2 м/с, скорость течения за ударной волной - 97,76 м/с. Начальные плотности элегаза и воздуха в камере низкого давления равны 6,037 * 10 и 1,198 * 10-3 г/смз соответственно. Физические свойства элегаза и воздуха следующие: оба вещества являются невязкими, нетеплопроводными и идеальными газами с показателями адиабаты у = 1,094 (5Р6) и у = 1.4 (воздух), отношение молекулярных масс (5Р6 /воздух) принято равным 5,04.

Необходимо смоделировать развитие неустойчивости при многократном прохождении ударной волны через контактный разрыв.

На рисунке 1 приведен график плотности в начальный момент времени. На рисунках 2 и 3 представлены изменения плотности при первом и повторном прохождении ударной волны через контактную границу.

Рис. 1. Изменение плотности в начальный момент времени.

Рис. 2. Изменение плотности при первом прохождении ударной волны через контактную границу.

Рис. 3. Изменение плотности при втором прохождении ударной волны через контактную границу.

Заключение. Разностная схема и программа расчета были протестированы на задаче с известными экспериментальными данными [7]. Результаты численных расчетов показали соответствие теоретическим предложениям и экспериментальным данным. Таким образом, проведенное тестирование предложенной схемы и программы моделирования

гидродинамических течений показывает, что данную схему можно использовать при решении задач нестационарной газовой динамики.

ЛИТЕРАТУРА

1. NVIDIA CUDA. Programming Guide [Электронный ресурс]. - Режим доступа: http://developer.nvidia.com/cuda-toolkit-40.

2. Kurganov A., Chi-Tien L. On the reduction of numerical dissipation in central-upwind schemes // Commun. Comput. Phys. February. - 2007. - Vol. 2, No. 1. - pp. 141-163.

3. Жалнин Р. В. О построении параллельного вычислительного алгоритма высокого порядка точности для гиперболических систем уравнений // Журнал Средневолжского математического общества. - 2007. - Т. 9, № 1. - С. 145-153.

4. Жалнин Р. В., Змитренко Н. В., Ладонкина М. Е., Тишкин В. Ф. Численное моделирование развития неустойчивости Рихтмайера-Мешкова с использованием схем высокого порядка точности // Математическое моделирование. - 2007. -Т. 19, № 10. - С. 61-66.

5. Ладонкина М. Е. Численное моделирование турбулентного перемешивания с использованием высокопроизводительных систем: дисс. ... канд. физ.-мат. наук. -М., 2005. - 157 с.

6. Lebo I. G., Nikishin V. V., Rozanov V. B., Tishkin V. F. On the development of instability near the contact boundary between two equal-density gases in passage of a shock wave // Bulletin of the Lebedev Physics Institute (Kratkie Soobsheniya po Fizike). - New York: Allerton Press, Inc, 1997.

7. Poggi F., Thorembey M.-H., Rodriguez G. Velocity measurements in turbulent gaseous mixtures induced by Richtmyer-Meshkov instability // Physics of Fluids. - 1998. -Vol. 10, No. 11. - P. 2698-2700.