如何使用 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，我找不到太多可读的东西。我买了阅读"Metal by Example: High-performance graphics and data-parallel programming for 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 高效地转换为我们想要的布局和颜色 space。只需更改几行，我们就可以从 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;
}

注：

我们将缓冲区作为 uchar4，因为每个 4 字节序列代表一个像素。
我们使用属性为 thread_position_in_grid 的参数对缓冲区进行索引，该参数指示我们使用计算命令编码器调度的网格的全局索引。由于我们的"image"是一维的，所以这个位置也是一维的。
一般来说，整数算术运算在 GPU 上的开销非常大。有可能在此函数中执行整数-> 浮点数转换所花费的时间支配了在包含浮点数的缓冲区上操作的额外带宽，至少在某些处理器上是这样。

希望对您有所帮助。如果您告诉我们更多关于您尝试做什么的信息，我们可以就如何加载和处理您的图像数据提出更好的建议。

如何使用 Metal 将纹理缓冲区数据传递给着色器？

How to pass texture buffer data to Shader with Metal?

shader

compute-shader

ios

swift

metal