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

Question

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

3

我希望在计算着色器中使用纹理数据作为一维数组。我了解到最好的方法是将其作为缓冲区而不是一维纹理进行传递。

我正在使用以下代码加载纹理：

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的书籍吗？

- Jeshua Lacock

使用缓冲区而不是纹理作为纹理的一个问题是，您需要了解图像数据的数据格式。MTKTextureLoader会从图像数据中推断纹理像素格式，这意味着它是不可预测的。苹果的建议是不要尝试解释任意图像，而是绘制到您选择的位图，并对其进行解释。 - Ken Thomases

另一个问题：纹理的buffer属性只对最初从缓冲区创建的纹理有用。在这里不能依赖它。它很可能是nil。要将纹理数据复制到缓冲区，请使用MTLBlitCommandEncoder及其

copy(from:sourceSlice:sourceLevel:sourceOrigin:sourceSize:to:destinationOffset:destinationBytesPerRow:destinationBytesPerImage:)

方法。 - Ken Thomases

1个回答

网页内容由stack overflow 提供, 点击上面的

可以查看英文原文，
原文链接

- warrenm · Accepted Answer

传输数据是作为缓冲区还是纹理取决于在内核函数中要执行的操作。如果使用缓冲区，则无法获得纹理的一些优点：当采样越界、插值和从源像素格式自动转换为着色器中请求的组件类型时，无法定义行为。

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

为了讨论，假设我们希望将数据放在等同于.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;
}

注意：

我们使用 uchar4 作为缓冲区类型，因为每个4字节的序列代表一个像素。
我们使用带有 thread_position_in_grid 属性的参数来索引缓冲区，该参数表示我们使用计算命令编码器分派到网格中的全局索引。由于我们的“图像”是1D，因此该位置也是一维的。
通常，GPU上的整数运算操作非常昂贵。在这个函数中进行整数转换到浮点数的时间可能会主导使用包含浮点数的缓冲区的额外带宽，至少在某些处理器上是这样。

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