Это история о том, как мы мягко_параллельно замедлился ускоренные расчеты на самом мощном суперкомпьютере в мире .
В апреле этого года наша команда приняла участие в финале Азиатский суперкомпьютерный вызов 2017 , одной из задач которой было ускорение программы моделирования океанских волн Masnum-Wave на китайском суперкомпьютере Sunway TaihuLight. Все началось с отборочного тура в феврале: мы получили доступ к суперкомпьютеру и познакомились с новым другом на ближайшие пару месяцев.
Вычислительные узлы расположены в виде двух овалов, между которыми находится сетевое оборудование.
В узлах используются процессоры Shenwei 26010. Каждый из них состоит из 4 разнородных групп процессоров, включающих одно управляющее и 64 вычислительных ядра с тактовой частотой 1,45 ГГц и размером локального кэша 64 КБ.
Все это охлаждается с помощью системы водоснабжения, расположенной в отдельном здании.
Что касается программного обеспечения, то в нашем распоряжении были компиляторы FORTRAN и C с «поддержкой» OpenACC 2 и потоков (аналог POSIX Threads) и планировщик задач, одновременно сочетающий в себе возможности обычного планировщика и mpirun. Доступ к кластеру осуществлялся через специальный VPN, плагины для работы с которым были доступны только для Windows и Mac OS. Все это придавало работе особый шарм.
Задача заключалась в ускорении Masnum-Wave на этом суперкомпьютере.
Нам предоставили исходный код Masnum-Wave, несколько крошечных файлов readme, описывающих основы работы на кластере, и данные для измерения ускорения.
Masnum-Wave — программа для моделирования движения волн по земному шару.
Он написан на Фортране с использованием MPI. Если вкратце, то он итеративно делает следующее: считывает данные функций внешнего воздействия, синхронизирует граничные области между процессами MPI, рассчитывает ход волн и сохраняет результаты.
Нам дали рабочую нагрузку на 8 модельных месяцев с шагом в 7,5 минут. В первый же день мы нашли в Интернете статью: «Суперкомпьютер Sunway TaihuLight: система и приложения, Haohuan Fu», описывающее ускорение Masnum-Wave на архитектуре Sunway TaihuLight с использованием конвейерной обработки.
Авторы статьи использовали всю мощность кластера (10 649 600 ядер), при этом в нашем распоряжении было 64 вычислительные группы (4 160 ядер).
Masnum-Wave состоит из нескольких модулей: 
Мы впервые столкнулись с таким большим количеством кода, созданного наукой.
Поскольку код представляет собой смесь двух версий Фортрана 90 и 77, это иногда нарушало работу собственного компилятора.
Здесь много замечательных конструкций goto, фрагментов закомментированного кода и, конечно же, комментариев на китайском языке.
Сокращенный пример кода для ясности:
Прежде всего, мы выявили узкие места в коде и кандидатов на оптимизацию, используя вывод времени выполнения каждой функции.do 1201 j=1,jl js=j j11=jp1(mr,j) j12=jp2(mr,j) j21=jm1(mr,j) j22=jm2(mr,j) !**************************************************************** ! do 1202 ia=ix1,ix2 ! do 1202 ic=iy1,iy2 eij=e(ks,js,ia,ic) !if (eij.lt.1.e-20) goto 1201 if (eij.lt.1.e-20) cycle ea1=e(kp ,j11,ia,ic) ea2=e(kp ,j12,ia,ic) ! .eij2=eij**2 zua=2.*eij/al31 ead1=sap/al11+sam/al21 ead2=-2.*sap*sam/al31 ! fcen=fcnss(k,ia,ic)*enh(ia,ic) fcen=fconst0(k,ia,ic)*enh(ia,ic) ad=cwks17*(eij2*ead1+ead2*eij)*fcen adp=ad/al13 adm=ad/al23 delad =cwks17*(eij*2.*ead1+ead2) *fcen deladp=cwks17*(eij2/al11-zua*sam)*fcen/al13 deladm=cwks17*(eij2/al21-zua*sap)*fcen/al23 !* nonlinear transfer se(ks ,js )= se(ks ,js )-2.0*ad se(kp2,j11)= se(kp2,j11)+adp*wp11 se(kp2,j12)= se(kp2,j12)+adp*wp12 se(kp3,j11)= se(kp3,j11)+adp*wp21 se(kp3,j12)= se(kp3,j12)+adp*wp22 se(im ,j21)= se(im ,j21)+adm*wm11 se(im ,j22)= se(im ,j22)+adm*wm12 se(im1,j21)= se(im1,j21)+adm*wm21 se(im1,j22)= se(im1,j22)+adm*wm22 !.
! 1202 continue ! 1212 continue 1201 continue
Больше всего нас заинтересовал модуль warcor, отвечающий за численное решение уравнения модели и функцию записи контрольных точек.
Изучив скудную документацию китайских компиляторов, мы решили использовать OpenACC, так как это стандарт от Nvidia с примерами и спецификациями.
Кроме того, код из ридми к тредам показался нам неоправданно сложным и просто не компилировался.
Как мы ошибались.
Одна из первых идей, которая приходит в голову при оптимизации кода на ускорителях, — использование локальной памяти.
В OpenACC это можно сделать с помощью нескольких директив, но результат всегда должен быть одинаковым: перед началом вычислений данные необходимо скопировать в локальную память.
Для проверки и выбора необходимой директивы мы написали несколько тестовых программ на Фортране, в которых убедились, что они работают и что таким образом можно получить ускорение.
Далее мы расположили эти директивы в Masnum-Wave, дав им имена наиболее используемых переменных.
После компиляции они стали появляться в логах, сопутствующие надписи на китайском не выделялись красным, и мы предполагали, что всё скопировано и работает. 
Но оказалось, что не все так просто.
Компилятор OpenACC не копировал массивы в Masnum-Wave, но корректно работал с тестовыми программами.
Проведя пару дней с Google Translate, мы поняли, что он не копирует объекты, определенные в файлах, включаемых через директивы препроцессора (include)! Следующую неделю мы потратили на миграцию кода Masnum-Wave из включаемых файлов (которых более 30) в исходные файлы, проверяя, чтобы все было идентифицировано и связано в правильном порядке.
Но, поскольку ни у кого из нас не было опыта работы с Фортраном, и все сводилось к тому, «давай смиримся и посмотрим, что получится», без замены некоторых базовых функций на костыльные версии не обошлось.
И вот, когда все модули были перелопачены, и мы в надежде получить свое мизерное ускорение запустили свежескомпилированный код, мы получили очередную порцию разочарования! Директивы, написанные по всем канонам стандарта OpenACC 2.0, выдают ошибки во время выполнения.
В этот момент в голову начала закрадываться мысль, что этот замечательный суперкомпьютер поддерживает какой-то особый стандарт. Тогда мы запросили у организаторов конкурса документацию, и с третьей попытки они нам ее предоставили.
Пара часов работы с Google Translate подтвердили наши опасения: поддерживаемый ими стандарт называется OpenACC 0.5, и он радикально отличается от OpenACC 2.0, который поставляется с компилятором pgi. Например, нашей основной идеей было повторное использование данных на ускорителе.
Чтобы сделать это в стандарте 2.0, нужно обернуть параллельный код в блок данных.
Вот как это делается на примерах от Nvidia: !$acc data copy(A, Anew)
do while ( error .
gt. tol .
and. iter .
lt. iter_max ) error=0.0_fp_kind !$omp parallel do shared(m, n, Anew, A) reduction( max:error ) !$acc kernels do j=1,m-2 do i=1,n-2 Anew(i,j) = 0.25_fp_kind * ( A(i+1,j ) + A(i-1,j ) + & A(i ,j-1) + A(i ,j+1) ) error = max( error, abs(Anew(i,j)-A(i,j)) ) end do end do !$acc end kernels !$omp end parallel do if(mod(iter,100).
eq.0 ) write(*,'(i5,f10.6)'), iter, error
iter = iter +1
!$omp parallel do shared(m, n, Anew, A)
!$acc kernels
do j=1,m-2
do i=1,n-2
A(i,j) = Anew(i,j)
end do
end do
!$acc end kernels
!$omp end parallel do
end do
!$acc end data
Но этот код не будет компилироваться на кластере, поскольку в их стандарте эта операция выполняется путем указания индекса для каждого блока данных: #include <stdio.h>
#include <stdlib.h>
#define NN 128
int A[NN],B[NN],C[NN];
int main()
{
int i;
for (i=0;i<NN;i++)
{
A[i]=1;
B[i]=2;
}
#pragma acc data index(1)
{
#pragma acc parallel loop packin(A,B, at data 1) copyout(C)
for (i=0;i<NN;i++)
{
C[i]=A[i]+B[i];
}
}
for(i=0;i<NN;i++)
{
if(C[i]!=3)
{
printf("Test Error! C[%d] = %d\n", i, C[i]);
exit(-1);
}
}
printf("Test OL!\n");
return 0;
}
Коря себя за выбор OpenACC, мы все же продолжили работу, так как оставалось всего пару дней.
В последний день отборочного тура нам наконец удалось запустить наш «ускоренный» код. Мы получили замедление в 3,5 раза.
Нам ничего не оставалось, как написать в отчете о задании все, что мы думали об их реализации OpenACC в цензурированном виде.
Несмотря на это, мы получили массу положительных эмоций.
Когда еще вам придется удаленно отлаживать код на самом мощном компьютере в мире? P.S.: В итоге мы все-таки дошли до финальной части соревнований и поехали в Китай.
Последним заданием финала стала презентация с описанием решений.
Наилучшего результата добилась местная команда, написавшая свою библиотеку на C с использованием atthread, поскольку OpenACC, по их мнению, не работает. Теги: #HPC #fortran #суперкомпьютеры #вызов #C++ #fortran #C++ #Параллельное программирование
-
Общественный Транспорт В 8 Новых Городах
19 Oct, 24 -
«Шепот» От Соседа Через Канал [Freepbx]
19 Oct, 24
