А. Куперштох

1. А. Куперштох ПАРАЛЛЕЛЬНЫЕ ВЫЧИСЛЕНИЯ НА МНОГОПРОЦЕССОРНЫХ ВИДЕОКАРТАХ ДЛЯ МОДЕЛИРОВАНИЯ МЕХАНИКИ ЖИДКОСТИ С ФАЗОВЫМИ ПЕРЕХОДАМИ МЕТОДОМ РЕШЕТОЧНЫХ УРАВНЕНИЙ БОЛЬЦМАНА Институт гидродинамикиим. М.А.Лаврентьева Сибирского отделения РАН, Новосибирск, РОССИЯ http://ancient.hydro.nsc.ru/sk/ Международная конференция ИПМСС, Коблево –2011

2. Содержание • Уравненияметода LBE с учетом уравнений состояния • с возможными фазовыми переходами. • 2. Алгоритм расщепления по физическим процессам. • 3. Реализация алгоритма дляпараллельных вычислений на графических ускорителях (GPU). • 4.Архитектура GPU.Технология параллельного программированияCUDA. • 5.Примерымоделирования двухфазных систем • жидкость-пар. Международная конференция ИПМСС, Коблево –2011

3. Уравнение Больцмана с дискретными скоростями Кинетическое уравнение Больцмана На гидродинамическом этапе можно использовать дискретный и конечный набор скоростей частицck (Broadwell, 1964), (Годунов, Султангазин) Характеристики – прямые линии. Обычно используются заселенности Nk Международная конференция ИПМСС, Коблево –2011

4. Регулярные решетки для метода LBE Метод характеристик для LBE - шаг по времени такой, что Трехмерный вариант Двумерный вариант D3Q19 D2Q9

5. Метод решеточных уравнений Больцмана Переносфункций распределениявдоль характеристик Гидродинамические переменные Вычисление равновесныхфункций распределения • кинетическая температура • LBE псевдочастиц. Изменение функций распределения Оператор столкновений в виде BGK (релаксация)

6. Фазовые переходы для заданного УС Чтобы смоделировать притягивающую ветвь потенциала межмолекулярного взаимодействия -- силы притяжения между частицами в соседних узлах. (Shan – Chen, Phys. Rev. E, 1993) Полная сила,действующая на вещество в узле (Qian, Chen, Int. J. Mod. Phys., 1997) Мы ввели Тогда Изотропная конечно-разностная аппроксимация Изменение функций распределения за счет действия сил (Метод точной разности) – Куперштох, 2003. где (Kupershtokh, Medvedev, Karpov, CAMWA, 2009, V. 58, N 5)

7. Алгоритм LBE Использовалась трехмерная LBE модельD3Q19. Шаг по времени состоит изнескольких подшагов (метод расщепления по физическим процессамН.Н. Яненко) 1. Перенос функций распределениявдоль характеристик. 2. Вычисление новых значений плотностиифункцииФ используя конкретное уравнение состояния. 3. Вычислениесил взаимодействия, действующих на узел (градиент псевдопотенциала). 4. Вычислениезначений скорости доипоследействия силна узел. Вычисление соответствующих равновесных функций распределения. 5. Изменение функций распределенияв узле за счет оператора столкновений. 6. Изменение функций распределенияв узле за счет сил взаимодействия, используяметод точной разности. Международная конференция ИПМСС, Коблево –2011

8. Параллельные вычисления на GPU nVIDIAGTX-580: 512 ядер 3 GBвнутренней памяти Есть также nVIDIATesla-C2070: 6 GBвнутренней памяти В алгоритме LBE вычисления - преимущественнолокально в узле за исключениемподшага переносаивычислений градиента псевдопотенциала. Для программирования на GPU использовалась технологияCUDA 4.0и компиляторgcc подWindows XP-32. Первые 3 подшагареализованы как 3 отдельныеядерные функции(kernels)для GPU: MOVE, DENSITY, FORCE. Последние 3 подшагаобъединены в одну ядерную функцию(kernel): COLLISION. Каждая ядерная функция выполняетотносительно простые вычисления в узле.

9. Архитектура GTX-580 (Fermi) 16 мультипроцессоров по 32 ядра в каждом.

10. Параллельные вычисления, CUDA до1024потока в блоке Технология CUDA (Compute Unified Device Architecture): const int i=(blockIdx.z*blockDim.z+threadIdx.z+1)*ny*nx +(blockIdx.y*blockDim.y+threadIdx.y+1)*nx + blockIdx.x*blockDim.x+threadIdx.x+1;

11. Производительность CPU -- GPU 1) Площадь клеток пропорциональна относительному вкладу функциональных блоков в чипе. 2) Внутренняя память GPUработает на скорости на порядок выше, чем оперативная память компьютера 3) Быстрое переключение между потоками. Международная конференция ИПМСС, Коблево –2011

12. GridиBlocksдля З-х мерного моделирования Двухфазная однокомпонентная модель LBE Решетка 256×256×224 для GTX-580-3Gb dim3 threads(16,16,2); dim3 grid(nx/threads.x, ny/threads.y, nz/threads.z); KERNEL<<< grid,threads >>> (….....); Международная конференция ИПМСС, Коблево –2011

13. Шаг переноса Для алгоритма соптимизацией памяти GPU (один буферный массив): const int i=(blockIdx.z*blockDim.z+threadIdx.z+1)*ny*nx +(blockIdx.y*blockDim.y+threadIdx.y+1)*nx + blockIdx.x*blockDim.x+threadIdx.x+1; buff[i]=N3[i+1]; N3[i]=buff [i]; buff[i]=N2d[i-nx+1]; buff[i]=Nz3d[i+nx*ny+1]; buff[i]=Nu4d[i-nx*ny+nx];

14. double *d_A, *d_B; __global__ void LAPLAS(double *A, double *B){ const int i=(blockIdx.y*blockDim.y+threadIdx.y+1)*nx + blockIdx.x*blockDim.x+threadIdx.x+1; B[i] = 0.25*(A[i+1]+A[i-1]+A[i+nx]+A[i-nx]); } void INIT(double *A, double *B){ cudaMalloc((void **)&d_A, DATA_SZ); cudaMalloc((void **)&d_B, DATA_SZ); cudaMemcpy(d_A, A, DATA_SZ, cudaMemcpyHostToDevice) ; cudaMemcpy(d_B, A, DATA_SZ, cudaMemcpyHostToDevice) ; } void SUBR(double *A){ dim3 threads(16, 16); dim3 grid(nx0/threads.x, ny0/threads.y); for (int k=0; k<nk; k++){ LAPLAS<<< grid, threads >>>(d_A, d_B); cudaThreadSynchronize(); LAPLAS<<< grid, threads >>>(d_B, d_A); cudaThreadSynchronize(); } cudaMemcpy(A, d_A, DATA_SZ, cudaMemcpyDeviceToHost); }

15. Компиляторыgccи nvccподWindows XP

16. Время счетаядра “COLLISION” за один шаг по времениалгоритма LBE Международная конференция ИПМСС, Коблево –2011

17. Время расчетадля 3D-LBE Двухфазная однокомпонентная модель D3Q19 Решетка256x256x224=14.7миллионов узлов. Вариантс оптимизацией по памяти! GTX-580-3Gb: Profiler: 10нс/узел Производительность3-х мерноймодели LBE: ~100 MNUPS (MillionsNode Updates PerSecond) 100 шагов по времени за~15секунд. Метод простыхитераций для уравнения Лапласа: 0.2 нс/узел;5000 MNUPS

18. Двухмерный расчет столкновения капель, формирование брызг Кумулятивные струи R1=400, R2=200, u=0.1, s=2.23. t = 40, 1600, 7200, 10800. Решетка 2000 x 4000. Международная конференция ИПМСС, Коблево –2011

19. Эффект Марангони (термокапиллярность) Распределение температуры для , где Решетка до 256×256×224 Международная конференция ИПМСС, Коблево –2011

20. Спинодальная декомпозиция Пар прозрачен

21. Разрушение 3D тонкостенного жидкого пузыря LBE с двумя компонентами. Решетка 144×144×144

22. Заключение Метод LBE достаточнопрост иможет быль легко использовандля моделированияширокого кругатечений жидкостей с возможными фазовыми переходами жидкость-пар. Алгоритм LBE может быть успешно реализован для параллельных расчетовна многоядерных графических ускорителях (GPU),используя технологиюCUDA. Для варианта методаLBEс фазовыми переходами, параллельная реализациянашего алгоритманаGTX-580 при вычислениях с двойной точностью в70-90 разбыстреечем при вычисленияхна одном ядре процессора Intel Core 2 Duo (3.3 ГГц). Производительность GTX-580 для метода LBE с одним компонентом и фазовыми переходами жидкость-пар ~100 MNUPS(Million Node Updates PerSecond). ПроизводительностьGTX-580 для метода LBE сдвумя компонентами и фазовыми переходами жидкость-парв электрическом поле~ 1-10 MNUPS.

23. Будущие устройства Keplerи Maxwell

24. Спасибо за внимание! skn@hydro.nsc.ru http://ancient.hydro.nsc.ru/sk/ Работа выполнена при частичной финансовой поддержкеРФФИ № 10-08-00805, ОЭММПУ №2.14.3, Президиума РАН № 2.12 и интеграционного проекта СО РАН № 58-2009 Международная конференция ИПМСС, Коблево –2011

25. Method of correct incorporation of EOS into theLBE model includes: ● Isotropic finite difference approximation of potential gradient on a lattice. ● The method of exact difference for the implementation of the body force term. ● Special set of reduced variables that allows to control the hydrodynamic Courant number easily to ensure the numerical stability of LBE simulations. Our method allows one to obtain the values of the liquid and vapor densities at the interface corresponding to the saturation curve with high accuracy. Международная конференция ИПМСС, Коблево –2011

26. Распад гомогенной бинарной смесив сильном электрическом поле A = 60 Диэлектрическая жидкость Газ Уравнение Пуассона: A = 27 Ez Возникновение в жидкости паро-газовых каналов вдоль электрического поля. Решетка 128×128×208.

27. GTX-580 (Fermi) 16 мультипроцессоров по 32 ядра в каждом.