如何将纹理缓冲区数据传递给Shader with Metal？

Question

我想在计算着色器中将纹理数据用作一维数组.我读到最好的方法是将其作为缓冲区而不是1D纹理传递.

我正在加载纹理:

let textureLoader = MTKTextureLoader(device: device)

do {
    if let image = UIImage(named: "testImage") {
        let options = [ MTKTextureLoaderOptionSRGB : NSNumber(value: false) ]
        try kernelSourceTexture = textureLoader.newTexture(with: image.cgImage!, options: options)
            kernelDestTexture = device.makeTexture(descriptor: kernelSourceTexture!.matchingDescriptor())
    } else {
        print("Failed to load texture image from main bundle")
    }
}
catch let error {
    print("Failed to create texture from image, error \(error)")
}

我正在创建缓冲区(不确定这是否正确):

var textureBuffer: MTLBuffer! = nil
var currentVertPtr = kernelSourceTexture!.buffer!.contents()
textureBuffer = device.makeBuffer(bytes: &currentVertPtr, length: kernelSourceTexture!.buffer!.length, options: [])
uniformBuffer.label = "textureData"

如何将缓冲区传递给计算着色器？我把它作为论据或制服传递给我吗？缓冲区的数据类型是什么？

对不起,如果这些都是愚蠢的问题,我刚开始使用Metal,我找不到太多的阅读.我买了并阅读了"金属示例:适用于iOS的高性能图形和数据并行编程".附带问题,有人可以推荐更多有关Metal的书籍吗？

Answer 1

是否应该将数据作为缓冲区或纹理传递取决于您在内核函数中要对其执行的操作.如果使用缓冲区,则不会获得纹理的几个好处:在超出范围的采样,插值以及从源像素格式到着色器中请求的组件类型的组件自动转换时定义的行为.

但是既然你问过缓冲区,那么我们来谈谈如何创建一个包含图像数据的缓冲区以及如何将它传递给内核.

为了讨论起见,我假设我们希望我们的数据等效于.rgba8unorm格式,其中每个组件都是单个字节.

为了进行这种转换而创建纹理是浪费的(正如Ken在评论中指出的那样,默认情况下纹理没有缓冲区支持,这使得我们获取数据的方式变得复杂),所以让MTKTextureLoader我们自己做并自己做.

假设我们的包中有一个图像,我们有一个URL.然后我们可以使用如下方法加载它,确保它以所需的格式,并MTLBuffer用最少的副本包装数据:

func bufferWithImageData(at url: URL, resourceOptions: MTLResourceOptions, device: MTLDevice) -> MTLBuffer? {
    guard let imageSource = CGImageSourceCreateWithURL(url as CFURL, nil) else { return nil }
    if CGImageSourceGetCount(imageSource) != 1 { return nil }
    guard let image = CGImageSourceCreateImageAtIndex(imageSource, 0, nil) else { return nil }
    guard let colorspace = CGColorSpace(name: CGColorSpace.genericRGBLinear) else { return nil }

    let bitsPerComponent = UInt32(8)
    let bytesPerComponent = bitsPerComponent / 8
    let componentCount = UInt32(4)
    let bytesPerPixel = bytesPerComponent * componentCount
    let rowBytes = UInt32(image.width) * bytesPerPixel
    let imageSizeBytes = rowBytes * UInt32(image.height)

    let pageSize = UInt32(getpagesize())
    let allocSizeBytes = (imageSizeBytes + pageSize - 1) & (~(pageSize - 1))

    var dataBuffer: UnsafeMutableRawPointer? = nil
    let allocResult = posix_memalign(&dataBuffer, Int(pageSize), Int(allocSizeBytes))
    if allocResult != noErr { return nil }

    var targetFormat = vImage_CGImageFormat()
    targetFormat.bitsPerComponent = bitsPerComponent
    targetFormat.bitsPerPixel = bytesPerPixel * 8
    targetFormat.colorSpace = Unmanaged.passUnretained(colorspace)
    targetFormat.bitmapInfo = CGBitmapInfo(rawValue: CGImageAlphaInfo.premultipliedLast.rawValue)

    var imageBuffer = vImage_Buffer(data: dataBuffer, height: UInt(image.height), width: UInt(image.width), rowBytes: Int(rowBytes))
    let status = vImageBuffer_InitWithCGImage(&imageBuffer, &targetFormat, nil, image, vImage_Flags(kvImageNoAllocate))
    if status != kvImageNoError {
        free(dataBuffer)
        return nil
    }

    return device.makeBuffer(bytesNoCopy: imageBuffer.data, length: Int(allocSizeBytes), options: resourceOptions, deallocator: { (memory, size) in
        free(memory)
    })
}

(请注意,您需要import Accelerate使用vImage函数.)

以下是如何调用此方法的示例:

let resourceOptions: MTLResourceOptions = [ .storageModeShared ]
let imageURL = Bundle.main.url(forResource: "my_image", withExtension: "png")!
let inputBuffer = bufferWithImageData(at: imageURL, resourceOptions: resourceOptions, device: device)

这可能看起来不必要地复杂,但其优点在于,对于各种各样的输入格式,我们可以使用vImage有效地转换为我们想要的布局和色彩空间.通过仅改变几行,我们可以从RGBA8888转到BGRAFFFF,或许多其他格式.

以常规方式创建计算管道状态和您要使用的任何其他资源.您可以通过将其分配给任何缓冲区参数槽来传递刚刚创建的缓冲区:

computeCommandEncoder.setBuffer(inputBuffer, offset: 0, at: 0)

以通常的方式调度您的计算网格.

为了完整起见,这是一个在我们的缓冲区上运行的内核函数.它绝不是计算此结果的最有效方法,但这只是为了说明:

kernel void threshold(constant uchar4 *imageBuffer [[buffer(0)]],
                      device uchar *outputBuffer [[buffer(1)]],
                      uint gid [[thread_position_in_grid]])
{
    float3 p = float3(imageBuffer[gid].rgb);
    float3 k = float3(0.299, 0.587, 0.114);
    float luma = dot(p, k);
    outputBuffer[gid] = (luma > 127) ? 255 : 0;
}

注意:

1. 我们将缓冲区视为a uchar4,因为每个4字节序列代表一个像素.
2. 我们使用一个参数来索引缓冲区,该参数thread_position_in_grid表示我们使用计算命令编码器调度的网格的全局索引.由于我们的"图像"是1D,因此该位置也是一维的.
3. 通常,整数算术运算在GPU上非常昂贵.在此函数中执行integer-> float转换所花费的时间可能会占用在包含浮点数的缓冲区上运行的额外带宽,至少在某些处理器上是这样.

希望有所帮助.如果您告诉我们有关您要执行的操作的更多信息,我们可以就如何加载和处理您的图像数据提出更好的建议.