如何将纹理缓冲区数据传递给带有金属的着色器？ [英] How to pass texture buffer data to Shader with Metal?

问题描述

我希望将纹理数据作为计算着色器中的一维数组进行处理。我读到，最好的方法是将其作为缓冲区传递，而不是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"

如何将缓冲区传递给计算着色器？我是将其作为一种论据来传递，还是作为一种制服来传递？缓冲区的数据类型是什么？

抱歉，如果这些是愚蠢的问题，我才刚刚开始使用金属，我找不到很多可以阅读的东西。我购买并阅读了《Metals by Example：iOS的高性能图形和数据并行编程》。顺便问一句，有没有人推荐更多关于金属的书？

推荐答案

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

但既然您询问了缓冲区，我们就来讨论如何创建包含图像数据的缓冲区以及如何将其传递给内核。

为了便于讨论，我假设我们希望数据的格式相当于.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. 我们将缓冲区视为uchar4，因为每个4字节序列代表一个像素。
2. 我们使用thread_position_in_grid属性的参数来索引缓冲区，该参数表示使用计算命令编码器调度的网格的全局索引。因为我们的图像是一维的，所以这个位置也是一维的。
3. 一般情况下，在GPU上进行整数运算是非常昂贵的。在此函数中执行整数->浮点数转换所花费的时间可能控制了在包含浮点数的缓冲区上操作的额外带宽，至少在某些处理器上是这样。

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

这篇关于如何将纹理缓冲区数据传递给带有金属的着色器？的文章就介绍到这了，希望我们推荐的答案对大家有所帮助，也希望大家多多支持IT屋！