如何将纹理缓冲数据传递给金属

Question

我想将纹理数据作为计算着色器中的1D数组使用。我读到最好的方法是将其作为缓冲区而不是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"

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

对不起，如果这些是愚蠢的问题，我才刚刚开始使用金属，我找不到很多阅读。我购买并阅读了"以示例：iOS的高性能图形和数据并行编程"。附带问题，谁能推荐更多有关金属的书？

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. 我们将缓冲区作为uchar4，因为每个字节的每个序列代表一个像素。
2. 我们使用 thread_position_in_grid归因的参数将其索引到缓冲区，该参数指示我们使用Compute命令编码派遣的网格中的全局索引。由于我们的"图像"是1D，因此该位置也是一维。
3. 通常，整数算术操作在GPU上非常昂贵。在此功能中进行整数 -> float转换所花费的时间可能会在包含浮子的缓冲区上操作的额外带宽，至少在某些处理器上进行操作。

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