如何使用 Metal 将纹理缓冲区数据传递给 Shader？答案

【问题标题】：How to pass texture buffer data to Shader with Metal?如何使用 Metal 将纹理缓冲区数据传递给 Shader？
【发布时间】：2017-09-27 19:44:16
【问题描述】：

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

我正在加载纹理：

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：iOS 的高性能图形和数据并行编程”。附带问题，谁能推荐更多关于金属的书？

【问题讨论】：

使用缓冲区而不是纹理作为纹理的一个问题是，您需要知道图像数据的数据格式。 MTKTextureLoader 从图像数据中推断出纹理像素格式，这意味着它是不可预测的。 Apple 的建议是不要尝试解释任意图像，而是绘制到您选择的格式的位图，然后对其进行解释。
另一个问题：纹理的buffer 属性仅对最初从缓冲区创建的纹理有用。你不能在这里依赖它。很可能是nil。要将纹理数据获取到缓冲区，请使用MTLBlitCommandEncoder 及其copy(from:sourceSlice:sourceLevel:sourceOrigin:sourceSize:to:destinationOffset:destinationBytesPerRow:destinationBytesPerImage:) 方法。

标签： ios swift shader metal compute-shader

【解决方案1】：

您是否应该将数据作为缓冲区或纹理传递，这在一定程度上取决于您想在内核函数中对它做什么。如果您使用缓冲区，您将无法获得纹理的几个好处：定义越界采样时的行为、插值以及将组件从源像素格式自动转换为着色器中请求的组件类型。

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

为了便于讨论，我假设我们希望我们的数据等效于 .rgba8unorm 格式，其中每个组件都是一个字节。

仅仅为了进行这种转换而创建纹理是一种浪费（正如 Ken 在 cmets 中指出的那样，默认情况下纹理没有缓冲区支持，这使我们获取数据的方式变得复杂），所以让我们设置 @987654322 @放在一边，自己做。

假设我们的包中有一个图像，我们有一个 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;
}

注意：

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

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

【讨论】：

感谢百万提供如此有用的答案。我赞成你的回答，一旦我有时间消化它并快速实施它，我几乎肯定会接受它。至于我想要做的基本上是破坏图像，也就是故障艺术，所以实际上我不会从这个应用程序中获得纹理的好处。
我只有一个问题；用makeRenderCommandEncoder(descriptor:) 初始化的MTLRenderCommandEncoder 表明它没有成员setBuffer。我是否使用不同的命令编码器？
您首先要根据您的内核函数创建一个MTLComputePipelineState，然后使用MTLComputeCommandEncoder 来调度计算工作。
谢谢！我得到了一些意想不到的结果。我认为这可能是因为我使用的是MPSUnaryImageKernel，根据标题：A MPSUnaryImageKernel 消耗一个 MTLTexture 并产生一个 MTLTexture。那么我是否改为将MPSKernel 子类化？我希望着色器消耗 1D 图像缓冲区并生成 2D 纹理。
听起来您根本不想使用任何 MPS 内核，因为您正在编写自己的内核函数。