Мифическая скорость GPU. Параллелизация расчетов GPU+DSP на A7

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

Вообще с этими гистограммами все уже немного переборщили, если честно. Считают их и считают… Я сам себе надоел. Но стоит заметить — задача эта про классические трудности параллельных вычислений на графических процессорах.

И надо, при этом, сказать, в области вычислений на GPGPU в последнее время сложилось много нездоровых легенд. Например утверждается, что чуть ли не вообще все вычисления вскорости заменятся вычислениями на графических ускорителях. Поэтому к теме анализа их производительности возвращаются снова и снова. Вычисление гистограммы в этом смысле некоторая показательная штука. Да и вообще, её можно было бы включить в набор стандартных тестов — типа насколько камни и  GPGPU на них продвинулись в своем развитии. Но и нетрудно показать, что любой современный проц с хорошим набором векторных инструкций посчитает гистограмму большого массива данных если не быстрее, то за сопоставимое время. Вот сегодня этим и займемся — замеряем чего-нибудь. И не говорите вот, что не любите все эти тесты! Инструмент этот безусловно важный и часто используемый, не только в Image Processing.

Совсем немного теории

Кажется, что может быть проще расчета гистограммы? Так и есть, до тех пор пока мы знаем, что имеем дело с линейным алгоритмом. Но вот какая интересная штука во всем этом, очень часто, как только мы пытаемся разложить однопоточный линейный алгоритм на параллельные процессы, натыкаемся на ряд ограничений, которые сводят все наши усилия к неинтересному с точки зрения усилий результату.

Если вы занимались параллельными вычислениями то, скорее всего, сталкивались с некоторыми общепринятыми практиками по реализации различных стратегий организации вычислений и их ускорению. По сути к строгой инженерной дисциплине их отнести нельзя, как нельзя отнести к науке, скажем электротехнику. Но штука эта оказывается крайне практичным методом решения насущных задач. Некоторые фокусы и трюки в организации вычислений на параллельных архитектурах таки несколько улучшают ситуацию. Для меня самое интересное в этой истории как обойти ограничения доступного железа для решения конкретной задачи. Вполне может случиться и так, что реализация алгоритма с помощью одного такого фокуса не будет работать на более свежем устройстве. И вот гистограммный  пост по сути был про это — рассчитать то мы можем, но чем, как, и что это будет стоить — сильно зависит от того на чем производится расчет. Забегая вперед, скажу, что предложенный вариант не будет адекватно работать на процах A7, т.е. еще на казалось бы вполне свежих iPhone 5s.

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

К основным классическим алгоритмам параллелизации относят:

  • Разделение (map)
  • Свертку (reduce)
  • Сборку (gather)
  • Раскидывание (scatter)
  • Сканирование (scan)
  • Поиск (search)
  • Сортировку (sort)

Map/reduce — наверное самая известная технология параллелизации вычислений над большими массивами данных благодаря усилиями большой поисковой компании в приложении этой техники вычислений на больших кластерных системах.

Gather — наверное самый большой класс операций параллельных вычислений. К ним отсносятся почти все spatial-фильтры. К этому же классу можно отнести реализации фильтров с  КИХ или БИХ. И их прикладными имплементациями эмулирующими, например, фильтры размытия Гаусса (IIR Blur). Такие реализации, очень популярны в игровых приложениях. Или вот, например, новый UIBlurEffect — визуализация размытия фона в последних версиях iOS  в реальном времени.

Scan — по сути суммирование последовательности. Search и sort — ну, тоже вроде все ясно.

Scatter — вот эта фигня нам и нужна для параллелизации расчета всяких щняг вроде гистограммы: когда нам нужно кучу разобранных данных раскидать в небольшой массив специфически рассчитанных величин. И сегодня я покажу: каким бы крутым не был GPU Apple — пока  их же собственная реализация проигрывает их же CPU-версии.

Да, но почему тогда столько шума!? Зачем нам GPU-версия расчета гистограммы?

Отвечу на этот не простой вопрос сразу. Прежде чем приступим к разбору кода.

Компромис. Это вообще ключевое слово в практике программиста. Алгоритмисты и математики придумывают прекрасные  и красивые алгоритмы и теории, железячники создают прекрасные и быстрые камни. А нам, простым сантехникам,  приходится вот с этим всем жить. То, что красиво работает в теории, на практике оказывается либо неработающей моделью (см. выше про ограничения), либо потребляет какой-то важный и плохо-восполнимый или не расширяемый ресурс, с которым до сих пор не справились железячники. В мобильных устройствах (а говорим мы за мобильники в основном), надо понимать, это батарейка и оперативная память. В этом посте, чуть позже, я приведу инструментальные замеры энергопотребления различных вариантов реализации расчета гистограммы. И покажу как не только ускориться на том что есть, но и существенно снизить энергопотребление и выделение дополнительной памяти под расчеты именно за счет использования GPU.

Почему опиум  scatter для народа GPU?

Хоть эта штука и очевидная вещь, напишу пояснения просто из уважения к любознательной половине человечества. Главная проблема все та же — когда много данных и одновременно много операций над ними, а свести результаты операций нужно в ограниченный объем ячеек памяти (выходного массива), то получается хочешь не хочешь, а придется как-то фигачить все в итоге в разделяемую память. А потом еще и результаты сборки объединять. А кто нам гарантирует, что совместный доступ к памяти GPU из разных ядер будет эффективным, кроме разработчиков железа? Да и они то врут, в основном. Стратегии скатеринга на GPU для решения общих не графических задач, как ни странно решены давно и используют иногда, что самое странное, графические шейдеры.

Первый известный и самый эффективный алгоритм вычисления гистограммы на GPU основывается на том, что картинка разбивается на вершины с позициями равными значениям интенсивностей пикселов, затем поток полигонов (по сути равных одному пикселу) смешивается в режиме добавления с даунсемплингом результирующей текстуры до размерности гистограммы, где позиция пиксела и содержит значение бина гистограммы. Т.е. на входе у нас NxM картинка, на выходе 256×1 текстура каждый пиксел которой содержит количество бинов соответствующей интенсивности. За один проход расчитывается один канал. Я не буду подробно рассматривать такой подход — он достаточно хорошо реализован в GPUImage, к примеру. Это хоть и не Metal API, с которым мы в основном имеем дело, но можно оценить приблизительные затраты на вычисления таким способом. Основным недостатком такого подхода является фактическое дублирование данных текстуры в вершинные полигоны перед каждым сеансом расчета. Если такой фокус и прокатывает на больших машинах, то маленькие девайсы, обладают вышеупомянутым недостатком — ограничением памяти доступной приложению. Одним из следствий такой реализации может быть отстрел приложения вышедшего за пределы лимита выделения памяти операционной системой. Второй очевидный недостаток — мы никоим образом не сможем решить задачу расчета гистограммы произвольного формата данных. То с чем легко справляются методологии из мира CUDA, Metal, или OpenCL. Поэтому подробно на этом варианте останавливаться не будем.

Второй способ — распространенная практика CUDA-стайл. Когда мы пилим алгоритм вычисления на две фазы: 1-я фаза вычисление частичных гистограмм в каждой группе потоков. Каждое запущенное ядро считает гистограммы в своем изолированном пространстве и пишет расчеты в свои изолированные куски памяти. В этом случае нам не нужно заботится о коллизиях записи в одну ячейку памяти разных из ядер в один вычислительный квант времени. 2-я фаза собирает частичные гистограммы в одну.

Летающий цирк и другие сверкающие шарики камешки в устройстве

Вот с этим вторым способом связана одна красивая легенда. Мол он работает всегда и железно. Так вот нифига не всегда, и нифига не железно. Первая проблема связана с тем, что мы обязаны выбрать вариант группировки и область памяти для сборки одной группы потоков.  Это может быть как локальная разделяемая память устройства, в Metal она определяется модификатором threadgorup:

kernel void my_func(threadgroup float *a [[ threadgroup(0) ]], ...)
{
//
// все ядра будут иметь доступ к этой переменной и
// смогут ее модифицировать
//
threadgroup float b[10]; ...
}

Так и глобальная память устройства, в терминах Metal определяется модификатором device. Глобальная память по сути память разделяемая между CPU и GPU. Предполагается, что она медленнее чем локальная разделяемая. При этом локальная разделяемая как правило сильно меньше чем нам надо, а именно 16Кб. В любом случае эта стратегия подразумевает финальную сборку значений в глобальную память.

И вот тут наступает один интересный момент, на котором зациклились практически все разработчики GPGPU-CUDA-стайл версий скаттеринга: непреодолимое желание произвести сборку на самом же GPU. На Metal Shading Language это будет выглядеть приблизительно так:

...
//
// запускаем ровно kIMP_HistogramChannels - ядер
// выделяем память под частичную гистограмму
//
threadgroup uint temp[kIMP_HistogramChannels];

//
// обнуляем в каждом ядре свой кусок
// tindex - индекс ядра в вычислительной сетке
//
temp[tindex] = 0;

//
// синхронизуем выполение ядер по доступу к локальной
// разделяемой памяти
//
threadgroup_barrier(mem_flags::mem_threadgroup);
temp[intensity]++;
threadgroup_barrier(mem_flags::mem_threadgroup);
...
//
// собираем в финальную гистограмму свой кусок
//
histogram[tindex] += temp[tindex];
...

Тут все бы ничего, но запись в медленную глобальную память напрягает. Поэтому, мне больше нравится стратегия когда можно исключить финальную сборку из GPU и провести её на CPU, используя векторные инструкции  или другом доступном устройстве, например DSP, используя операции сложения векторов. Тут мы воспользуемся простым свойством выбора стратегии: много данных много ядер — выгоднее на GPU, мало данных много инструкций  — на GPU не выгодно. Т.е. очевидно: на GPU не выгодно собирать частичные гистограммы. И да, не забывайте изначально iPhone — это iPod со звонилкой, поэтому внутри нас ждет приятный бонус в виде неплохого камня с DSP и отличной оберткой в виде Accelerate Framework и vDSP движка из этого framework. Поэтому в нашей версии будем использовать комбинированный подход рачета на GPU/DSP и покажем, что иногда даже он не работает. Но иногда все же работает.

Самая свежая пицца

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

///
/// Функция счета частичной гистограммы.
///
kernel void kernel_impPartialRGBYHistogram(
                                           texture2d<float, access::sample>  inTexture   [[ texture(0) ]],
                                           device IMPHistogramPartialBuffer  *outArray   [[ buffer(0) ]],
                                           constant DPCropRegionIn          &regionIn    [[ buffer(1) ]],
                                           constant float                   &scale       [[ buffer(2) ]],
                                           //
                                           // Позиция группы в сетке групп.
                                           // Выходная текстура будет заполняться в соответствие с номером этой группы
                                           //
                                           uint  tindexi [[thread_index_in_threadgroup]],
                                           uint2 gid     [[thread_position_in_grid]]
                                           )
{
    constexpr float3 Im(kIMP_HistogramSize - 1);
 
    //
    // Просто рассчитываем цвет пиксела, если он не попал 
    // в нужный нам прямоугольник делаем ему альфу == 0
    //    
    float4 inColor = histogram::histogramSampledColor(inTexture,regionIn,scale,gid);
    
    //
    // Яркостный канал
    //
    uint   Y    = uint(dot(inColor.rgb,Ym) * inColor.a * Im.x);

    // 
    // Индексы бинов гистограммы по каждому из каналов
    //
    uint4  rgby(uint3(inColor.rgb * Im), Y);
    
    threadgroup_barrier(mem_flags::mem_device);
    if (inColor.a>0){
        for (uint i=0; i<kIMP_HistogramChannels; i++){
            // собираем для каждого из каналов частичную гистограмму
            outArray[tindexi].channel[i][rgby[i]]++;
        }
    }
    threadgroup_barrier(mem_flags::mem_device);
}

Теперь мы эту пиццу съедим, но частями

///
/// Анализатор гистограммы на основе раздельного расчета частичных гистограмм и сборке полной на DSP.
///
class IMPHistogramDSPReduceAnalyzer: IMPHistogramAnalyzer {
    
    ///
    /// Конструктор анализатора с произвольным счетчиком, который
    /// задаем kernel-функцией. Главное условие совместимость с типом IMPHistogramBuffer
    /// как контейнером данных гистограммы.
    ///
    ///
    required init(function: String, context aContext: DPContext!) {
        super.init(function: function, context: aContext)
        //
        // создаем память в устройстве под контейнер счета частями
        //
        histogramPartialUniformBuffer = histogramPartialUniformBuffer ?? self.context.device.newBufferWithLength(sizeof(IMPHistogramPartialBuffers), options: MTLResourceOptions.CPUCacheModeDefaultCache)
    }
    
    ///
    /// По умолчанию гистограмма инициализируется счетчиком интенсивностей в RGB-пространстве,
    /// с дополнительным вычислением канала яркости.
    ///
     convenience required init!(context aContext: DPContext!) {
        self.init(function: &quot;kernel_impPartialRGBYHistogram&quot;, context:aContext)
    }
    
    //
    // Буфер обмена контейнера счета с GPU
    //
    private var histogramPartialUniformBuffer:MTLBuffer?
    
    //
    // Прикладываем вычисления к тектстуре
    //
    override func apply() {
        
        super.apply(
            MTLSizeMake(Int(kIMP_HistogramSize), Int(1), 1),
            buffer: histogramPartialUniformBuffer!)
        
        // Обновляем структуру гистограммы с которой уже будем работать
        histogram.updateWithData(histogramPartialUniformBuffer!.contents(), dataCount: Int(kIMP_HistogramSize))
        
    }
}

Код сборки чатей в один кусок

...
   func updateWithData(dataIn: UnsafeMutablePointer&lt;Void&gt;, dataCount: Int){
        //
        // Обнуляем
        //
        self.clearHistogram()

        //
        // Пробегаемся, в нашем случае 256 раз, т.е. по всем кускам
        // 
        for i in 0..&lt;dataCount{
            let dataIn = UnsafePointer<IMPHistogramBuffer>(dataIn)+i
            let address = UnsafePointer<UInt3>(dataIn)
            for c in 0..<channels.count{
                var data:[Float] = [Float](count: Int(self.size), repeatedValue: 0)

                // заполняем кусок
                self.updateChannel(&data, address: address, index: c)
                // складываем с предыдущим
                self.addFromData(&data, toChannel: &amp;channels[c])
            }
        }
    }

...
    //
    // Обновление данных контейнера гистограммы
    //
    private func updateChannel(inout channel:[Float], address:UnsafePointer<UInt32>, index:Int){
        let p = address+Int(self.size)*Int(index)
        let dim = self.dim<1 ? 1 : self.dim;
        //
        // конвертим из единственно возможного в текущем MSL (atomic_)[uint] во [float]
        //
        vDSP_vfltu32(p, dim, &amp;channel, 1, self.size);
    }


...
    private func addFromData(inout data:[Float], inout toChannel:[Float]){
        //
        // Чтобы не заставлять CPU работать сверх меры
        // используем экскаватор - вычерпываем в один 
        // проход нужное количество данных и складываем 
        // с исходным куском
        //
        vDSP_vadd(&toChannel, 1, &data, 1, &toChannel, 1, self.size)
    }

    //
    // Обнуляем тоже на DSP, экономя тем самым дорогие такты CPU
    // и заодно батарейку устройства.
    //
    private func clearChannel(inout channel:[Float]){
        var zero:Float = 0
        vDSP_vfill(&zero, &channel, 1, vDSP_Length(self.size))
    }
    
    private func clearHistogram(){
        for c in 0..<channels.count{
            self.clearChannel(&amp;channels[c]);
        }
    }

Вернемся к атомарным операциям

Подозреваю, что у некоторых непостоянных читателей этого блога терпение лопнуло и они пошли пить пиво закрыли нафиг этот пост. Те же, кто дотерпел, задаются вопросом, а почему тогда в предыдущем посте про расчет всего этого безобразия мы отделались простыми атомарными операциями? А вот эта история как раз про нашу нелегкую программистскую жизнь. В A8 Apple не просто улучшила работу с разделяемой памятью CPU/GPU, но наконец таки прикрутила железную организацию атомарного доступа к ней, а это как не трудно догадаться привело к драматическому росту скорости доступа к этой памяти и как следствие отказу от компромиссных решений. Но так ли это? Проверим чуть позже. Код приводить не буду, его можно найти как »
Расчет гистограммы изображения на GPU»
так и в полных исходниках гонялки теста: ImageMetalling-06.

Прежде чем гонять пример, примем окончательное решение по «национальному вопросу»

На самом деле вся вот эта возня, конечно не нужна. Не нужна, если вам нужно посчитать гистограмму один раз, или вы не собираетесь выходить за рамки стандартных процессинговых операций, поскольку все есть в движке vImage из все того-же Accelerate Framework. vImage предоставляет обертку к операциям над картинками с использованием векторных инструкций CPU, и в целом, поскольку в контексте исполнения нам не нужно проводить копирование данных в текстуру и обратно, любая операция будет работать предположительно быстрее чем процессинг или тем более расчет гистограммы на GPU. А ще же есть Core Image…

Но вот если нет, или мы хотим работать со своими уникальными пространственными или гомогенными фильтрами, то скорее всего GPU лучшее решение. Лучше оно еще по одной причине, порождение контекста ядра, т.е. распараллеливание вычислений на GPU практически ничего не стоит и как следствие не порождает дополнительную нагрузку на такой быстро исчерпаемый ресурс как аккумулятор мобильного устройства, в отличии от распараллеливаний векторных вычислений на CPU.

И вот теперь наступает момент истины, а если мы используем в итоге GPU для решения нашей задачи и используем vImage для решения ее частей, чем нам это грозит? Не трудно предположить, что самая первая засада вытекает из необходимости в один момент иметь копию текстуры в GPU и копию данных изображения для передачи в vImage. А как мы помним — iOS не любит приложения прожорливые до памяти, она их просто со временем отстреливает

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

В любом случае покажу как скомбинировать vIMage и Metal:

    override func apply() {

        let width  = Int(floor(Float(self.source.texture.width) * self.downScaleFactor))
        let height = Int(floor(Float(self.source.texture.height) * self.downScaleFactor))
        
        let threadgroupCounts = MTLSizeMake(Int(self.functionThreads), Int(self.functionThreads), 1)
        
        //
        // Вычисляем количество групп вычислительных ядер
        //
        let threadgroups = MTLSizeMake(
            (width  + threadgroupCounts.width ) / threadgroupCounts.width ,
            (height + threadgroupCounts.height) / threadgroupCounts.height,
            1)


        let commandBuffer = self.context.beginCommand()
        
        //
        // Первый большой МИНУС такого подхода: мы вынуждены выделять память под дополнительную текстуру 
        // с подоготовленными для анализа данными.
        //
        // Справедливости ради - это было бы не нужно для простой обработки гистограммы, но мы хотим не просто 
        // получать гистограмму в RGB-прсотранстве, а в произвольном пространстве необходимом для нашего анализа.
        //
        if analizeTexture?.width != self.source.texture.width || analizeTexture?.height != self.source.texture.height {
            let textureDescription = MTLTextureDescriptor.texture2DDescriptorWithPixelFormat(
                self.source.texture.pixelFormat,
                width: width,
                height:height, mipmapped: false)
            analizeTexture = self.context.device.newTextureWithDescriptor(textureDescription)
        }
        
        let commandEncoder = commandBuffer.computeCommandEncoder()
        
        commandEncoder.setComputePipelineState(kernel_impHistogram.pipeline);
        commandEncoder.setTexture(self.source.texture, atIndex:0)
        commandEncoder.setTexture(analizeTexture, atIndex:1)
        commandEncoder.setBuffer(regionUniformBuffer,    offset:0, atIndex:0)
        commandEncoder.setBuffer(scaleUniformBuffer,     offset:0, atIndex:1)

        
        //
        // Запускаем вычисления
        //
        commandEncoder.dispatchThreadgroups(threadgroups, threadsPerThreadgroup:threadgroupCounts);
        commandEncoder.endEncoding()

        //
        // Второй МИНУС такого подхода: выделяем еще один кусок памяти, причем памяти разделяемой CPU/GPU
        // сюда будем коипровать прилетевшую тектсуру
        //
        let imageBufferSize = width*height*4
        if imageBuffer?.length != imageBufferSize {
            imageBuffer = self.context.device.newBufferWithLength( imageBufferSize, options: MTLResourceOptions.CPUCacheModeDefaultCache)
        }
        
        //
        // Запускаем копирование текстуры в память доступную для процессинга гистограммы в движке Accelerate
        //
        let blitEncoder = commandBuffer.blitCommandEncoder()
        
        //
        // Быстрая команда копирования текстуры в разделяемую память.
        // Тут мы кроме памяти ощутимо ничего не теряем. 
        // Копирование в контексте памяти устройства имеет еще один существенных недостаток,
        // для наших целей: память картинки должна приведена к размерности кратной 64 байтам,
        // т.е. произвольный даунсемплинг в таком варианте невозможен
        //
        blitEncoder.copyFromTexture(analizeTexture!,
            sourceSlice: 0,
            sourceLevel: 0,
            sourceOrigin: MTLOrigin(x: 0, y: 0, z: 0),
            sourceSize: MTLSize(width: width, height: height, depth: 1),
            toBuffer: imageBuffer!,
            destinationOffset: 0,
            destinationBytesPerRow: width*4,
            destinationBytesPerImage: 0)
        blitEncoder.endEncoding()
        
        //
        // Выполняем контекст
        //
        self.context.commitCommand()
        
        //
        // Подготавливаемся к вычислению гистограммы на движке акселерации
        //
        var vImage = vImage_Buffer(
            data: (imageBuffer?.contents())!,
            height: vImagePixelCount(analizeTexture!.height),
            width: vImagePixelCount(analizeTexture!.width),
            rowBytes: analizeTexture!.width*4)

        let red   = [vImagePixelCount](count: Int(kIMP_HistogramSize), repeatedValue: 0)
        let green = [vImagePixelCount](count: Int(kIMP_HistogramSize), repeatedValue: 0)
        let blue  = [vImagePixelCount](count: Int(kIMP_HistogramSize), repeatedValue: 0)
        let alpha = [vImagePixelCount](count: Int(kIMP_HistogramSize), repeatedValue: 0)
        
        let redPtr   = UnsafeMutablePointer&lt;vImagePixelCount&gt;(red)
        let greenPtr = UnsafeMutablePointer&lt;vImagePixelCount&gt;(green)
        let bluePtr  = UnsafeMutablePointer&lt;vImagePixelCount&gt; (blue)
        let alphaPtr = UnsafeMutablePointer&lt;vImagePixelCount&gt;(alpha)
        
        let rgba = [redPtr, greenPtr, bluePtr, alphaPtr]

        //
        // Быстро (быстрее чем на GPU) вычисляем гистограмму
        //
        let hist = UnsafeMutablePointer<UnsafeMutablePointer<vImagePixelCount>>(rgba)
        vImageHistogramCalculation_ARGB8888(&amp;vImage, hist, 0)
        
        // обновляем структуру гистограммы с которой уже будем работать
        histogram.updateWithRed(red, green: green, blue: blue, alpha: alpha)
    }

Протестируем то что получилось

До кучи проверим, что нам покажет вариант реализованный разработчиками Apple в составе Metal Performance Shaders Kit.

Пример использования MPSImageHistogram:

       
...

       //
        // Конструируем гистограмму MPS
        //
        var imageHistogramInfo = MPSImageHistogramInfo(
            numberOfHistogramEntries: Int(kIMP_HistogramSize),
            histogramForAlpha: true,
            minPixelValue: vector_float4(x: 0, y: 0, z: 0, w: 0),
            maxPixelValue: vector_float4(x: 1, y: 1, z: 1, w: 1))
        mpsHistogram = MPSImageHistogram(device: self.context.device, histogramInfo: &amp;imageHistogramInfo)
        


...
        let commandBuffer = self.context.beginCommand()
        
        let blitEncoder = commandBuffer.blitCommandEncoder()
        blitEncoder.fillBuffer(histogramMPSUniformBuffer, range: NSMakeRange(0, sizeof(IMPHistogramBuffer)), value: 0)
        blitEncoder.endEncoding()
        
        if downScaleFactor == 1.0 {
            sampledTexure = self.source.texture
        }
        else {
            scaleFilter?.transform.resampleFactor = CGFloat(self.downScaleFactor)
            scaleFilter?.source=self.source
            sampledTexure = scaleFilter?.destination.texture
        }
                
        mpsHistogram?.encodeToCommandBuffer(commandBuffer, sourceTexture: sampledTexure!, histogram: histogramMPSUniformBuffer, histogramOffset: 0)
        mpsHistogram?.clipRectSource = MTLRegion(origin: MTLOrigin(x: 0, y: 0, z: 0), size: MTLSize(width: self.source.texture.width, height: self.source.texture.height, depth: 1))
        
        self.context.commitCommand()
        
        // обновляем структуру гистограммы с которой уже будем работать
        histogram.updateWithConinuesData(UnsafeMutablePointer&lt;UInt32&gt;(histogramMPSUniformBuffer!.contents()))
...

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

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

import UIKit
import MetalPerformanceShaders

class IMPTestFilter:DPFilter{
    
    var analyzer:IMPHistogramAnalyzer?
    
    init!(context aContext: DPContext!, analyzer:IMPHistogramAnalyzer?) {
        
        super.init(context: aContext)
        
        if (analyzer != nil) {
            self.analyzer = analyzer
            self.analyzer?.downScaleFactor = 1
            
            self.willStartProcessing = { (source) in
                analyzer!.source = source
            }
        }
    }

    required init!(context aContext: DPContext!) {
        fatalError(&quot;init(context:) has not been implemented&quot;)
    }
    
}

class ViewController: UIViewController {

    let context = DPContext.newContext()
    
    //
    // Замеряем время исполнения функции фильтра по одной и тойже картинке много раз
    //
    func filterTest(filter: IMPTestFilter, provider:DPImageProvider) -> (all: Float, time: Float) {
        
        //
        // Источник
        //
        filter.source = provider
        
        var tac:NSTimeInterval = 0
        var tacl = NSDate.timeIntervalSinceReferenceDate()
        
        filter.analyzer?.analyzerDidUpdate = {
            tac += NSDate.timeIntervalSinceReferenceDate() - tacl
            tacl = NSDate.timeIntervalSinceReferenceDate()
        }
        
        //
        // Первая метка
        //
        let t1 = NSDate.timeIntervalSinceReferenceDate()
        
        //
        // Пусть будет N раз
        //
        let times = 10
        
        for _ in 0..&lt;times{
            //
            // Заставляем фильтр не филонить
            //
            filter.dirty = true
            //
            // А запускать функцию каждый раз
            //
            filter.apply()
        }
        
        //
        // Метка в конце
        //
        let t2 = NSDate.timeIntervalSinceReferenceDate()
        
        return (Float(t2-t1)/Float(times), Float(tac)/Float(times))
        
    }
    
    override func viewDidAppear(animated: Bool) {
        super.viewDidAppear(animated)
        
        let analyzerAT:IMPHistogramAnalyzer?     = IMPHistogramATAnalyzer.newWithContext(self.context)
        let analyzerMPS:IMPHistogramAnalyzer?    = IMPHistogramMPSAnalyzer.newWithContext(self.context)
        let analyzerDSP:IMPHistogramAnalyzer?    = IMPHistogramDSPReduceAnalyzer.newWithContext(self.context)
        let analyzerVImage:IMPHistogramAnalyzer? = IMPHistogramVImageAnalyzer.newWithContext(self.context)

        let analizers = [
            (analizer: analyzerDSP,   title: "через сборку на DSP", name: "DSP"),
            (analizer: analyzerMPS,   title: "через MPS          ", name: "MPS"),
            (analizer: analyzerVImage,title: "через vImage       ", name: "VImage"),
            (analizer: analyzerAT,    title: "через atomic types ", name: "AT"),
        ]
        
        //
        // Загружаем картинку
        //
        let provider = DPUIImageProvider.newWithImage(UIImage(named: "test.jpg"), context: context)
        
        let size = Float(provider.texture.width * provider.texture.height * 4)
        let mb   = powf(1024, 2)

        var messages = [String]()
        
        for a in analizers{
            if a.analizer?.isHardwareSupported == false{
                messages.append(String(format:" *** %@ не поддерживается...",a.name))
                continue
            }
            
            if let filter = IMPTestFilter(context: context, analyzer: a.analizer){
                let t = self.filterTest(filter, provider: provider)
                let rate = size/(t.time == 0 ? 0 : t.time)/mb
                print(" \(a.name) = \(filter.analyzer!.histogram.channels[0]);")
                let s = String(format:" *** скорость расчета гистограммы %@: %.2fMb/s время счета = %.4fs общее время фильтра = %.4f", a.title, rate, t.time, t.all)
                messages.append(s)
                filter.flush()
            }
            print("\n----\n")
            sleep(1)
        }
        
        NSLog(" *** Модель: %@:%@", UIDevice .currentDevice().model, provider.context.device.name!)
        for m in messages{
            NSLog("%@",m)
        }
    }
    
}
2015-12-10 22:21:43.867 ImageMetalling-06[1246:355394]  *** Модель: iPhone:Apple A8 GPU
2015-12-10 22:21:43.868 ImageMetalling-06[1246:355394]  *** скорость расчета гистограммы через сборку на DSP : 108.56Mb/s время счета = 0.2808s общее время фильтра = 0.2808
2015-12-10 22:21:43.869 ImageMetalling-06[1246:355394]  *** скорость расчета гистограммы через MPS           : 316.80Mb/s время счета = 0.0962s общее время фильтра = 0.0962
2015-12-10 22:21:43.869 ImageMetalling-06[1246:355394]  *** скорость расчета гистограммы через vImage        : 487.80Mb/s время счета = 0.0625s общее время фильтра = 0.0625
2015-12-10 22:21:43.869 ImageMetalling-06[1246:355394]  *** скорость расчета гистограммы через atomic types  : 326.62Mb/s время счета = 0.0933s общее время фильтра = 0.0933
015-12-10 22:20:51.815 ImageMetalling-06[1219:185616]  *** Модель: iPhone:Apple A7 GPU
2015-12-10 22:20:51.816 ImageMetalling-06[1219:185616]  *** скорость расчета гистограммы через сборку на DSP : 43.44Mb/s время счета = 0.7017s общее время фильтра = 0.7018
2015-12-10 22:20:51.816 ImageMetalling-06[1219:185616]  *** MPS не поддерживается...
2015-12-10 22:20:51.817 ImageMetalling-06[1219:185616]  *** скорость расчета гистограммы через vImage        : 74.13Mb/s время счета = 0.4112s общее время фильтра = 0.4112
2015-12-10 22:20:51.817 ImageMetalling-06[1219:185616]  *** скорость расчета гистограммы через atomic types  : 15.21Mb/s время счета = 2.0041s общее время фильтра = 2.0041

Так же замерим энергопотребление каждого из вариантов.

Как видно атомарные операции не дают преимущества перед сборкой частичных гистограмм на DSP
Сравнение энергозатрат на AT/DSP scatter

Но зато вычисления на базе vImage оказались на порядок более прожорливы
Сравнение энергозатрат на вычисление гистограммы на vImage и GPU

Не будем оригинальными, сделаем выводы

Наше дотошное исследование показало:

  1. нет в мире совершенства
  2. не все GPU одинаково хороши, новые явно лучше
  3. Apple не хочет поддерживать MPS для 5s
  4. CPU иногда быстрее чем GPU
  5. vImage работает с гистограммой пока всегда быстрее, но требует в два раза больше памяти и в 10 раз больше батарейки
  6. Атомарные операции на A8, и видимо выше, лучше подходят для скаттеринга, чем раздельное вычисление и раздельная сборка
  7. Атомарные операции на A7 не работают должным образом
  8. Иногда чуть более чем быстрее чем Apple можно что-то посчитать на A7 с помощью GPU

В итоге мы получили, как минимум, один интересный результат: считать на GPU для A7 и A8 приходится по разному — A7 просто не умеет работать с атомарными операциями, а на A8 не имеет смысла их не использовать — они сильно быстрее чем простой доступ к памяти. И вот с этим нам теперь предстоит жить и работать.

Всем инжой!


Еще раз напомню: авторы блога не преследуют задачи быть предельно корректным, но если заметили явную ашипку, если написали явную глупость, если что-то не понятно: комментируйте или пишите на: imagemetalling [*] gmail.com .

Реклама

Добавить комментарий

Заполните поля или щелкните по значку, чтобы оставить свой комментарий:

Логотип WordPress.com

Для комментария используется ваша учётная запись WordPress.com. Выход / Изменить )

Фотография Twitter

Для комментария используется ваша учётная запись Twitter. Выход / Изменить )

Фотография Facebook

Для комментария используется ваша учётная запись Facebook. Выход / Изменить )

Google+ photo

Для комментария используется ваша учётная запись Google+. Выход / Изменить )

Connecting to %s