Подсчет цветных пикселей на графическом процессоре - Теория

Ответ 1

Действительно, GPU общего назначения (например, устройства Apple от A8 on, например) не только способны, но также предназначены для решения таких проблем с параллельной обработкой данных.

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

Следующие действия выполняются как инструмент командной строки Swift в OS X Sierra и были построены с использованием Xcode 9 (да, я знаю, что он бета). Вы можете получить полный проект из моего github repo.

Как main.swift:

import Foundation
import Metal
import CoreGraphics
import AppKit

guard FileManager.default.fileExists(atPath: "./testImage.png") else {
    print("./testImage.png does not exist")
    exit(1)
}

let url = URL(fileURLWithPath: "./testImage.png")
let imageData = try Data(contentsOf: url)

guard let image = NSImage(data: imageData),
    let imageRef = image.cgImage(forProposedRect: nil, context: nil, hints: nil) else {
    print("Failed to load image data")
    exit(1)
}

let bytesPerPixel = 4
let bytesPerRow = bytesPerPixel * imageRef.width

var rawData = [UInt8](repeating: 0, count: Int(bytesPerRow * imageRef.height))

let bitmapInfo = CGBitmapInfo(rawValue: CGImageAlphaInfo.premultipliedFirst.rawValue).union(.byteOrder32Big)
let colorSpace = CGColorSpaceCreateDeviceRGB()

let context = CGContext(data: &rawData,
                        width: imageRef.width,
                        height: imageRef.height,
                        bitsPerComponent: 8,
                        bytesPerRow: bytesPerRow,
                        space: colorSpace,
                        bitmapInfo: bitmapInfo.rawValue)

let fullRect = CGRect(x: 0, y: 0, width: CGFloat(imageRef.width), height: CGFloat(imageRef.height))
context?.draw(imageRef, in: fullRect, byTiling: false)

// Get access to iPhone or iPad GPU
guard let device = MTLCreateSystemDefaultDevice() else {
    exit(1)
}

let textureDescriptor = MTLTextureDescriptor.texture2DDescriptor(
    pixelFormat: .rgba8Unorm,
    width: Int(imageRef.width),
    height: Int(imageRef.height),
    mipmapped: true)

let texture = device.makeTexture(descriptor: textureDescriptor)

let region = MTLRegionMake2D(0, 0, Int(imageRef.width), Int(imageRef.height))
texture.replace(region: region, mipmapLevel: 0, withBytes: &rawData, bytesPerRow: Int(bytesPerRow))

// Queue to handle an ordered list of command buffers
let commandQueue = device.makeCommandQueue()

// Buffer for storing encoded commands that are sent to GPU
let commandBuffer = commandQueue.makeCommandBuffer()

// Access to Metal functions that are stored in Shaders.metal file, e.g. sigmoid()
guard let defaultLibrary = device.makeDefaultLibrary() else {
    print("Failed to create default metal shader library")
    exit(1)
}

// Encoder for GPU commands
let computeCommandEncoder = commandBuffer.makeComputeCommandEncoder()

// hardcoded to 16 for now (recommendation: read about threadExecutionWidth)
var threadsPerGroup = MTLSize(width:16, height:16, depth:1)
var numThreadgroups = MTLSizeMake(texture.width / threadsPerGroup.width,
                                  texture.height / threadsPerGroup.height,
                                  1);

// b. set up a compute pipeline with Sigmoid function and add it to encoder
let countBlackProgram = defaultLibrary.makeFunction(name: "countBlack")
let computePipelineState = try device.makeComputePipelineState(function: countBlackProgram!)
computeCommandEncoder.setComputePipelineState(computePipelineState)


// set the input texture for the countBlack() function, e.g. inArray
// atIndex: 0 here corresponds to texture(0) in the countBlack() function
computeCommandEncoder.setTexture(texture, index: 0)

// create the output vector for the countBlack() function, e.g. counter
// atIndex: 1 here corresponds to buffer(0) in the Sigmoid function
var counterBuffer = device.makeBuffer(length: MemoryLayout<UInt32>.size,
                                        options: .storageModeShared)
computeCommandEncoder.setBuffer(counterBuffer, offset: 0, index: 0)

computeCommandEncoder.dispatchThreadgroups(numThreadgroups, threadsPerThreadgroup: threadsPerGroup)

computeCommandEncoder.endEncoding()
commandBuffer.commit()
commandBuffer.waitUntilCompleted()

// a. Get GPU data
// outVectorBuffer.contents() returns UnsafeMutablePointer roughly equivalent to char* in C
var data = NSData(bytesNoCopy: counterBuffer.contents(),
                  length: MemoryLayout<UInt32>.size,
                  freeWhenDone: false)
// b. prepare Swift array large enough to receive data from GPU
var finalResultArray = [UInt32](repeating: 0, count: 1)

// c. get data from GPU into Swift array
data.getBytes(&finalResultArray, length: MemoryLayout<UInt>.size)

print("Found \(finalResultArray[0]) non-white pixels")

// d. YOU'RE ALL SET!

Кроме того, в Shaders.metal:

#include <metal_stdlib>
using namespace metal;

kernel void
countBlack(texture2d<float, access::read> inArray [[texture(0)]],
           volatile device uint *counter [[buffer(0)]],
           uint2 gid [[thread_position_in_grid]]) {

    // Atomic as we need to sync between threadgroups
    device atomic_uint *atomicBuffer = (device atomic_uint *)counter;
    float3 inColor = inArray.read(gid).rgb;
    if(inColor.r != 1.0 || inColor.g != 1.0 || inColor.b != 1.0) {
        atomic_fetch_add_explicit(atomicBuffer, 1, memory_order_relaxed);
    }
}

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

Источники:

http://flexmonkey.blogspot.com.ar/2016/05/histogram-equalisation-with-metal.html

http://metalbyexample.com/introduction-to-compute/

http://memkite.com/blog/2014/12/15/data-parallel-programming-with-metal-and-swift-for-iphoneipad-gpu/

Ответ 2

Здесь можно сделать GPU.

Я не уверен, что вы ищете алгоритм здесь, но я могу указать на широко используемую библиотеку GPU, которая реализует эффективную процедуру подсчета. Взгляните на функцию count в библиотеке thrust: https://thrust.github.io/doc/group__counting.html

Он работает как входной функцией предиката. Он подсчитывает количество вхождений входных данных, которые удовлетворяют предикату.

Следующее подсчитывает количество элементов в data, равных нулю.

template <typename T>
struct zero_pixel{
  __host__ __device__ bool operator()(const T &x) const {return x == 0;}
};
thrust::count_if(data.begin(), data.end(), zero_pixel<T>())

Рабочий пример: https://github.com/thrust/thrust/blob/master/testing/count.cu

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

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

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

Взгляните на этот пример: https://github.com/thrust/thrust/blob/master/examples/histogram.cu

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

Ответ 3

Ваш вопрос: Я просто хотел бы знать, может ли это сделать/может делать GPU?

Ответ: Да, GPU может обрабатывать ваши вычисления. Все номера выглядят очень дружелюбными к GPU:

• размер основы: 32 (16x2)
• Максимальное количество потоков на блок: 1024 (8x128) (8x8x16)
• Максимальное количество потоков на мультипроцессор: 2048... и т.д.

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

Процедура: Как правило, использование GPU означает, что вы копируете данные из памяти CPU в память GPU, затем выполняете вычисления на графическом процессоре и, наконец, копируете результат в CPU для дальнейшего расчеты. Важно подумать, что вся эта передача данных осуществляется через PCI-e-связь между процессором и графическим процессором, что очень медленно по сравнению с обоими.

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

Мой совет: Если после анализа (профилирования) всей вашей программы вы считаете, что эта процедура получения черного количества пикселей является настоящим узким местом, попробуйте:

• рекурсивный алгоритм деления и покорения, или

• распараллеливание ваших вычислений в нескольких ядрах процессора.