事实上,通用 GPU(例如从 A8 开始的 Apple 设备中的 GPU)不仅有能力,而且旨在能够解决此类并行数据处理问题。
Apple 在其平台中使用 Metal 引入了数据并行处理,只需一些简单的代码,您就可以使用 GPU 解决像您这样的问题。即使这也可以使用其他框架来完成,我也包含了一些 Metal+Swift 案例的代码作为概念证明。
以下内容在 OS X Sierra 上作为 Swift 命令行工具运行,并使用 Xcode 9 构建(是的,我知道它是测试版)。您可以从我的网站获得完整的项目github 仓库 https://github.com/pevasquez/MetalCounter/tree/master.
As 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);
}
}
我通过这个问题了解了一些有关 Metal 和数据并行计算的知识,因此大部分代码都是从网上文章中用作样板并进行编辑的。请花时间访问下面提到的来源以获取更多示例。另外,代码几乎是针对这个特定问题进行硬编码的,但是您在调整它时应该不会遇到很多麻烦。
Sources:
http://flexmonkey.blogspot.com.ar/2016/05/histogram-equalization-with-metal.html http://flexmonkey.blogspot.com.ar/2016/05/histogram-equalisation-with-metal.html
http://metalbyexample.com/introduction-to-compute/ http://metalbyexample.com/introduction-to-compute/
http://memkite.com/blog/2014/12/15/data-parallel-programming-with-metal-and-swift-for-iphoneipad-gpu/ http://memkite.com/blog/2014/12/15/data-parallel-programming-with-metal-and-swift-for-iphoneipad-gpu/