所以我很清楚如何编写使用
CIFilter
或 CIColorKernel
的 CoreImage CIBlendKernel
。本质上,这些方法是将 1px 输入转换为 1px 输出。
这不是我的问题。使用现成的过滤器也不是我的问题。
是否可以通过典型的
CIColorKernel
函数迭代所有像素,同时访问我作为指针传入的值?
或者,是否可以计算 ROI 的平均值,然后返回宽度/高度 1 的图像?即过滤器调用 Metal 函数一次?
基本上,我想计算具有非零 Alpha 通道的任何像素的中值颜色。那么这个术语是“减少过滤器”吗?
我正在寻找这种金属函数的一般形式,以及它的内核类型。
这可能吗?我想将大量计算转移到 GPU 上,而 Apple 文档有点简单,超出了表面案例。
我理解您为什么考虑 1x1 像素输出纹理,但这不一定是您想要的。
在“hello着色器”示例中,您将看到经过处理后产生输出纹理的输入纹理。一个简单易懂的例子是从 RGB 计算灰度。每个像素 RGB 都会产生输出灰度像素。 GPU 允许同时“粗略”地处理每个像素。重要的一点是输出像素“大致”同时写入所有像素。 如果您将自己限制为 1x1 像素输出,则所有输入像素都会竞争,并且将不确定地覆盖其输出。 您真正追求的是写入输出“像素”(或者更确切地说,您将用来计算最终像素的东西)的同步原子方式。 在我提供的代码中,这是通过
device atomic_int
金属类型和
atomic_fetch_add_explicit
实现的。显然,同时访问越少,性能越好。
注意每个 RGB 通道都有一个单独的 Int32 累加器。这是故意的,因为我们很容易溢出,并且 Int32 大小的device atomic_int
目前是唯一具有原子数学运算的 Metal 支持类型。
除了输入纹理之外,我们还引入了 4 个缓冲区:
2个输入:ROI - 代表您感兴趣区域的 4 个整数数组 Alpha 阈值 - 截止 alpha,低于该值则丢弃像素。
2个输出:
rgbSum - 一个 3 整数数组,其中索引 0,1,2 分别代表 RGB 通道总和
count - 处理后的像素数,以便我们可以计算算术平均值
请注意,我们不会选择“纯”中位数,因为它需要排序,并且会削弱并行处理的优势。
这个想法非常简单 - 我们将 R、G、B 通道的总值相加,并计算它们基于访问像素数的平均值。当 GPU 完成时,您可以在 Swift 完成处理程序中获得结果。
class MetalHelper {
let metalDevice = MTLCreateSystemDefaultDevice()!
var computePipelineState: MTLComputePipelineState?
var textureCache: CVMetalTextureCache!
lazy var commandQueue: MTLCommandQueue? = {
return self.metalDevice.makeCommandQueue()
}()
init(shaderFunctionName: String) {
let defaultLibrary = metalDevice.makeDefaultLibrary()!
let kernelFunction = defaultLibrary.makeFunction(name: shaderFunctionName)
do {
computePipelineState = try metalDevice.makeComputePipelineState(function: kernelFunction!)
} catch {
print("Could not create pipeline state: \(error)")
}
}
func process(pixelBuffer: CVPixelBuffer) {
guard let commandQueue = device.makeCommandQueue() else { return }
guard let commandBuffer = commandQueue.makeCommandBuffer() else { return }
guard let commandEncoder = commandBuffer.makeComputeCommandEncoder() else { return }
guard let inputTexture = makeTextureFromCVPixelBuffer(pixelBuffer: pixelBuffer, textureFormat: .bgra8Unorm) else { return nil }
commandEncoder.setTexture(inputTexture, index: 0)
let intputRoi: [Int32] = [Int32(0), // origin X
Int32(0), // origin Y
Int32(100),// width
Int32(100)]// height
commandEncoder.setBuffer(metalDevice.makeBuffer(bytes: intputRoi, length: intputRoi.count * MemoryLayout<UInt32>.stride), offset: 0, index: 0)
let inputThesholdAlpha: [Float] = [0.5]
let bufferSize = inputThesholdAlpha.count * MemoryLayout<Float>.
let inputBuffer = device.makeBuffer(bytes: inputThesholdAlpha, length: bufferSize)
commandEncoder.setBuffer(metalDevice.makeBuffer(bytes: inputThesholdAlpha, length: 1 * MemoryLayout<Float>.stride), offset: 0, index: 1)
let rgbSumOutput = metalDevice.makeBuffer(length: 3 * MemoryLayout<UInt32>.stride)
commandEncoder.setBuffer(metalDevice.makeBuffer(bytes: rgbOutput, length: 3 * MemoryLayout<UInt32>.stride), offset: 0, index: 2)
let count = metalDevice.makeBuffer(length: 1 * MemoryLayout<UInt32>.stride)
commandEncoder.setBuffer(metalDevice.makeBuffer(bytes: count, length: 1 * MemoryLayout<UInt32>.stride), offset: 0, index: 3)
// Set up the thread groups.
let width = computePipelineState!.threadExecutionWidth
let height = computePipelineState!.maxTotalThreadsPerThreadgroup / width
let threadsPerThreadgroup = MTLSizeMake(width, height, 1)
let threadgroupsPerGrid = MTLSize(width: (inputTexture.width + width - 1) / width,
height: (inputTexture.height + height - 1) / height,
depth: 1)
commandEncoder.dispatchThreadgroups(threadgroupsPerGrid, threadsPerThreadgroup: threadsPerThreadgroup)
commandEncoder.endEncoding()
commandBuffer.addCompletedHandler { [weak self] (mtlCommandBuffer: MTLCommandBuffer) in
let rgbSumOutputPointer = bufferForHough!.contents().bindMemory(to: Int32.self, capacity: houghSpaceCount)
let countPointer = angleScores!.contents().bindMemory(to: Int32.self, capacity: 180)
let red = rgbSumOutputPointer[0]
let blue = rgbSumOutputPointer[1]
let green = rgbSumOutputPointer[2]
let finalPixel = (Float(rgbSumOutputPointer[0]) / Float(countPointer), Float(rgbSumOutputPointer[1]) / Float(countPointer), Float(rgbSumOutputPointer[2]) / Float(countPointer))
// do something with final pixel
}
}
实际的内核着色器函数
#include <metal_stdlib>
#include <metal_atomic>
using namespace metal;
kernel void averageCompute(texture2d<half, access::read> inTexture [[ texture(0) ]],
device int *roi [[ buffer(0) ]],
device float *alphaThreshold [[ buffer(1) ]],
device atomic_int *rgbSum [[ buffer(2) ]],
device atomic_int *count [[ buffer(3) ]]),
uint2 gid [[ thread_position_in_grid ]]) {
// check if within RegionOfInterest
if (gid.x < roi[0] || gid.y < roi[1] || gid.x > roi[2] || gid.y > roi[3]) {
return;
}
half3 pixelAtCoordinates = inTexture.read(gid);
// filter out pixels with too low alpha
if (pixelAtCoordinates.a < alphaThreshold[0]) {
return;
}
int red = (int)(pixelAtCoordinates.r * 255.0);
atomic_fetch_add_explicit(&rgbSum[0],
red,
memory_order_relaxed);
int blue = (int)(pixelAtCoordinates.b * 255.0);
atomic_fetch_add_explicit(&rgbSum[1],
blue,
memory_order_relaxed);
int green = (int)(pixelAtCoordinates.g * 255.0);
atomic_fetch_add_explicit(&rgbSum[2],
green,
memory_order_relaxed);
atomic_fetch_add_explicit(&count[0],
1,
memory_order_relaxed);
}
用途:
MetalHelper("averageCompute")