Развитие высокопроизводительных систем сопровождается не только увеличением их быстродействия, но и расширением многообразия — на кристалле процессора для одних задач выгоднее разместить много вычислительных ядер и мало памяти, а для других, наоборот, больше памяти и меньше ядер [1, 2]. В результате перенос программы, от которой требуется высокая производительность, с одной архитектуры на другую обычно означает переписывание кода, что долго и дорого. Разработка параллельной программы на переносимом языке OpenCL (Open Computing Language), поддерживаемом производителями большого семейства вычислительных архитектур, не решает проблему, поскольку для многих задач производительность зависит не только от параллельности, но и от использования структуры памяти.

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

Пиковая производительность вычислительной системы — это максимальная производительность выполнения некоторой «идеальной» программы, вычисляемая как количество одновременно работающих вычислительных устройств системы, умноженное на частоту системы. Например, для процессора с частотой 3,5 ГГц с шестью ядрами, каждое из которых имеет SIMD-ускоритель, обрабатывающий одновременно восемь чисел в формате с плавающей запятой, пиковая производительность равна 2*3,5*109 *6*8 FLOPS = 336 GFLOPS. Если программа на таком процессоре достигает производительности 100 GFLOPS, то ее эффективность — 0,3 (100/336). Данный показатель позволяет оценить, в какой степени алгоритм соответствует конкретной архитектуре и насколько удачно на нее отображена программа.

Чтобы программа была эффективна на различных архитектурах, она должна иметь возможность на них настраиваться (переноситься), а мерой переносимости программы для заданного множества архитектур можно считать минимальное и максимальное отношения эффективностей программ для данных архитектур.

Для иллюстрации оценки переносимости рассмотрим тесты CompuBench [compubench.com], написанные на языке OpenCL и выполняемые на четырех разных архитектурах: Intel Core i7-5930K, AMD Opteron Processor 6276, NVIDIA GeForce GTX 980 Ti и AMD Radeon R9 Fury X с пиковой производительностью 672, 665, 5300 и 8600 GFLOPS соответственно. Производительность тестов принято измерять в обрабатываемых мегапикселях в секунду, кадрах в секунду и т. п. Д ля каждого теста приняты свои коэффициенты приведения данных единиц к FLOPS (wi, i = 1, 2,. ..7). Например, для теста Face detection производительность на выбранных архитектурах равна: 0,04w1, 0,024w1, 0,038w1 и 0,018w1. Таким образом, мера переносимости теста Face detection на данном разноплановом наборе архитектур составляет 0,45.

Кроме качества переносимости той или иной программы на данном множестве архитектур можно определять и среднюю относительную эффективность архитектуры на данном множестве тестов. Для этого следует вычислить отношение эффективности конкретной архитектуры на определенном тесте к максимальной эффективности для данного теста и усреднить по всем тестам. Полученная величина и есть средняя относительная эффективность архитектуры.

Если проанализировать эффективность и переносимость на всех тестах CompuBench, то наиболее эффективными программы их реализации будут на графическом ускорителе GeForce GTX 980 Ti. Однако при переносимости программ обработки изображений семейства CompuBench может теряться от 40 до 80% производительности.

Для решения проблемы переносимости программ в 2008 году был разработан стандарт OpenCL, регламентирующий создание приложений для многоядерных процессоров и параллельных ускорителей. Это упростило написание программ, предназначенных для различных вычислителей, однако привело к появлению некоторых заблуждений — например, возникла уверенность в том, что для обеспечения переносимости программы ее достаточно лишь переписать на OpenCL, что не соответствует действительности.

Рассмотрим программу подсчета количества вхождений в заданный массив X элементов массива Y.

kernel void calc1(__global int *x, __global int *y, uint N, uint M, __global uint *result ) {

int i,j, id=get_global_id(0), sz = get_global_size(0);

for (int i = id; i


  Если выполнить эту программу для M = 32 и N = 2 000 000 на трех различных архитектурах: двухъядерном процессоре Intel T2310, ускорителе Nvidia Tesla C2075 и видеокарте AMD Radeon HD6970 с пиковыми производительностями 14 GFLOPS, 1030 GFLOPS и 2700 GFLOPS, то ее эффективность составит 9,2, 0,8 и 0,65% соответственно. Следовательно, переносимость будет равна 0,07, что мало по сравнению со специально предназначенными для запуска на различных архитектурах тестами CompuBench. Для данной программы не характерна высокая интенсивность выполнения арифметических операций, поэтому отношения максимальной пропускной способности памяти к пиковой производительности на вычислительных архитектурах различного типа для нее сильно отличаются.



  В приведенном примере размер внутреннего цикла M = 32 удобен для векторизации, а большая длина внешнего цикла удобна для количества ядер всех трех приведенных архитектур. Тем не менее при хорошей распараллеливаемости эффективная переносимость этой программы оказывается низкой.



  Арифметически интенсивная программа расчета суммы ряда с M слагаемыми для заданного набора точек x[i], i=1...N будет более эффективна — 6,5, 7,5 и 4,1% соответственно, а переносимость будет равна 0,55. Для повышения производительности программы при выполнении на системах с векторными регистрами требуется провести векторизацию кода, что повлечет за собой снижение переносимости. Так как вычислители Nvidia не имеют векторных регистров, компилятор генерирует вдвое менее эффективный код — 3,4%, хотя на остальных системах эффективность повышается в несколько раз: 24% на T2310 и 13% на Radeon HD6970. Следовательно, при увеличении параллелизма (использовании векторизации) переносимость составит 0,14, что в четыре раза меньше, чем в невекторизованном случае.



  Таким образом, хорошая распараллеливаемость еще не гарантирует высокой производительности — производительность зависит и от других факторов, которые следует учитывать разработчикам.



  Перемены в компьютерных архитектурах оказали влияние на вычислительные методы и алгоритмы — например, существенно изменилось отношение времени выполнения вычислительных операций и операций обращения к памяти. Ежегодно скорость выполнения арифметических операций растет в среднем на 30%, а скорость обращений к памяти — на 9%. Чтобы максимально загрузить вычислительные устройства (ядра, векторизаторы и т. п.), необходимо на эти устройства подавать данные со скоростью не меньшей, чем происходит их обработка, однако скорость чтения из оперативной памяти пока еще уступает скорости их обработки — тем более, параллельной. Если для выполнения каждой операции считывать аргументы этой операции из оперативной памяти, то вычислительные устройства будут простаивать, поэтому создается промежуточная память (кэш, локальная), более быстрая, чем оперативная. К сожалению, пока эту быструю память сложно сделать такого же размера, как и оперативную, чтобы вместить все необходимые данные, поэтому возникают дополнительные пересылки данных. 



  Различные источники свидетельствуют о повышении эффективности программ за счет перехода к блочным вычислениям (тайлинг), это объясняется тем, что задача разбивается на более мелкие подзадачи, у каждой из которых данные подаются в кэш-память [3]. В этом случае с данными из кэш-памяти выполняется много операций, уменьшается число пересылок данных между оперативной памятью и кэшем, что влечет за собой минимизацию времени выполнения.



  Хорошее быстродействие достигается в том случае, когда программа удачно отображается на архитектуру — например, если у каждой части алгоритма все данные помещаются в кэш-память и количество таких частей мало, то отображение считается удачным. Однако у каждой из вычислительных систем свои характеристики элементов памяти, поэтому программа, оптимизированная для одной вычислительной системы, может «тормозить» на другой.



  Для программ, использующих блочные вычисления на системах с иерархической общей памятью, одним из параметров управления производительностью может быть размер блока. Если данных, используемых блоком вычислений, значительно меньше, чем вычислительных операций, и эти данные помещаются в кэш-памяти, то программа должна быть эффективной.



  Предположим, что программа решает задачу, исходные данные которой помещаются в большой медленной памяти, но не помещаются в быстрой. Идея ускорения программ за счет блочных вычислений состоит в том, чтобы разбить задачу на подзадачи так, чтобы у каждой данные помещались в быструю память и при этом подзадач было как можно меньше. Размер блока вычислений — это и есть параметр, управляющий объемом памяти для подзадачи. Оптимальные значения этого параметра для разных архитектур обычно различаются, поскольку различным может быть объем малой быстрой памяти.



  Следует отметить, что для системы, имеющей более двух уровней памяти, размер блока в блочной программе при одних значениях может разбивать задачу на подзадачи, попадающие в память одного уровня, а при других — другого. Если блоки вычислений программы разбить на блоки второго порядка, то можно подстроить вычисления под два уровня памяти.



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



  ***



  Распараллеливание ускоряет выполнение программ, однако может и замедлять их, если не учитывается иерархия памяти. В практике разработки программ давно используются параметры для указания размеров обрабатываемых данных, количества задействованных узлов и пр. Сегодня целесообразно с помощью параметров варьировать также объемы блоков данных высокопроизводительных программ. Такая параметризация позволяет адаптировать программы к вычислительным архитектурам с различной иерархией памяти для повышения производительности.



  Литература


    
  
  1. 
    Абу-Халил Ж. М., Морылев Р. И., Штейнберг Б. Я. Параллельный алгоритм глобального выравнивания с оптимальным использованием памяти // Современные проблемы науки и образования. — 2013. — № 1.
    2. 
  
  2. 
    Штейнберг Б. Я. Зависимость оптимального распределения площади кристалла процессора между памятью и вычислительными ядрами от алгоритма. PACO’2012/ Труды международной конференции «Параллельные вычисления и задачи управления». М., 24–26 октября 2012 г., ИПУ РАН, С. 99–108.
    3. 
  
  3. 
    Лев Гервич, Борис Штейнберг, Михаил Юрушкин. Программирование экзафлопсных систем // Открытые системы.СУБД. — 2013. — № 8. — С. 26–29. URL: http://www.osp.ru/os/2013/08/13037853 (дата обращения: 18.12.2015).
    4. 




  Жумана Абу-Халил, Сергей Гуда, Борис Штейнберг (borsteinb@mail.ru) — Южный федеральный университет (Ростов-на-Дону). Статья подготовлена на основе материалов доклада, представленного авторами на  VI Московский суперкомпьютерный форум  (МСКФ-2015, грант РФФИ 15-07-20824-г).