ios - Metal Compute 管道不适用于 MacOS，但适用于 iOS

Question

我正在尝试使用 Metal 进行一些 GPGPU 计算。我有一个基本的 Metal 管道:

• 创建所需的MTLComputePipelineState 管道和所有关联对象(MTLComputeCommandEncoder、命令队列等)；
• 创建用于写入的目标纹理(使用 desc.usage = MTLTextureUsageShaderWrite;)；
• 启动一个基本的着色器来用一些值填充这个纹理(在我的实验中，要么将其中一个颜色分量设置为 1，要么根据线程坐标创建一个灰度值渐变)；
• 从 GPU 读回此纹理的内容。

我在 2 个设置中测试此代码:

• 2013 年初在 OSX 10.11 上使用 MacBook Pro；
• 在 iOS 9 和 iPhone 6 上。

iOS 版本运行良好，我得到了我要求着色器执行的操作。但是在 OSX 上，我得到了一个有效的(非零，大小正确)输出纹理。但是，当取回数据时我得到的所有地方都是 0。

我是否遗漏了特定于 OS X 实现的步骤？这似乎发生在 NVIDIA GT650M 和 Intel HD4000 上，或者可能是运行时的错误？

由于我目前不知道如何进一步调查该问题，因此也将不胜感激这方面的任何帮助:-)

编辑 - 我当前的实现

这是我的实现的初始(失败)状态。它有点长，但主要是创建管道的样板代码:

id<MTLDevice> device = MTLCreateSystemDefaultDevice();
id<MTLLibrary> library = [device newDefaultLibrary];
id<MTLCommandQueue> commandQueue = [device newCommandQueue];

NSError *error = nil;
id<MTLComputePipelineState> pipeline = [device newComputePipelineStateWithFunction:[library
                                                                                    newFunctionWithName:@"dummy"]
                                                                             error:&error];
if (error)
{
    NSLog(@"%@", [error localizedDescription]);
}
MTLTextureDescriptor *desc = [MTLTextureDescriptor texture2DDescriptorWithPixelFormat:MTLPixelFormatRGBA8Unorm
                                                                                width:16
                                                                               height:1
                                                                            mipmapped:NO];
desc.usage = MTLTextureUsageShaderWrite;

id<MTLTexture> texture = [device newTextureWithDescriptor:desc];

MTLSize threadGroupCounts = MTLSizeMake(8, 1, 1);
MTLSize threadGroups = MTLSizeMake([texture width]  / threadGroupCounts.width,
                                   [texture height] / threadGroupCounts.height,
                                   1);

id<MTLCommandBuffer> commandBuffer = [commandQueue commandBuffer];

id<MTLComputeCommandEncoder> commandEncoder = [commandBuffer computeCommandEncoder];
[commandEncoder setComputePipelineState:pipeline];

[commandEncoder setTexture:texture atIndex:0];
[commandEncoder dispatchThreadgroups:threadGroups threadsPerThreadgroup:threadGroupCounts];
[commandEncoder endEncoding];

[commandBuffer commit];
[commandBuffer waitUntilCompleted];

用于获取数据的代码如下(我将文件分成两部分以获得更小的代码块):

// Get the data back
uint8_t* imageBytes = malloc([texture width] * [texture height] * 4);
assert(imageBytes);
MTLRegion region = MTLRegionMake2D(0, 0, [texture width], [texture height]);
[texture getBytes:imageBytes bytesPerRow:[texture width]*4 fromRegion:region mipmapLevel:0];
for (int i = 0; i < 16; ++i)
{
    NSLog(@"Pix = %d %d %d %d",
          *((uint8_t*)imageBytes + 4 * i),
          *((uint8_t*)imageBytes + 4 * i + 1),
          *((uint8_t*)imageBytes + 4 * i + 2),
          *((uint8_t*)imageBytes + 4 * i + 3));
}

这是着色器代码(将 1 写入红色和 alpha，在主机上读取时应该在输出缓冲区中变为 0xff):

#include <metal_stdlib>

using namespace metal;

kernel void dummy(texture2d<float, access::write> outTexture [[ texture(0) ]],
                  uint2 gid [[ thread_position_in_grid ]])
{
    outTexture.write(float4(1.0, 0.0, 0.0, 1.0), gid);
}

Answer 1

同步代码:

[commandEncoder setTexture:texture atIndex:0];
[commandEncoder dispatchThreadgroups:threadGroups threadsPerThreadgroup:threadGroupCounts];
[commandEncoder endEncoding];

//
// synchronize texture from gpu to host mem
//
id<MTLBlitCommandEncoder> blitEncoder = [commandBuffer blitCommandEncoder];
[blitEncoder synchronizeTexture:texture slice:0 level:0];
[blitEncoder endEncoding];

[commandBuffer commit];
[commandBuffer waitUntilCompleted];

这是在 2012 年年中使用与您拥有的相同 GPU 的 Mac Book 以及 2015 年年中使用 AMD Radeon R9 M370X 2048 МБ 的 Mac Book 进行测试的。

有时我使用以下技巧在不同步的情况下获取纹理数据:

id<MTLComputeCommandEncoder> commandEncoder = [commandBuffer computeCommandEncoder];
[commandEncoder setComputePipelineState:pipeline];

[commandEncoder setTexture:texture atIndex:0];
[commandEncoder dispatchThreadgroups:threadGroups threadsPerThreadgroup:threadGroupCounts];
[commandEncoder endEncoding];

//
// one trick: copy texture from GPU mem to shared
//
id<MTLBlitCommandEncoder> blitEncoder = [commandBuffer blitCommandEncoder];

[blitEncoder copyFromTexture:texture
                 sourceSlice: 0
                 sourceLevel: 0
                sourceOrigin: MTLOriginMake(0, 0, 0)
                  sourceSize: MTLSizeMake([texture width], [texture height], 1)
                    toBuffer: texturebuffer
           destinationOffset: 0
      destinationBytesPerRow: [texture width] * 4
    destinationBytesPerImage: 0];

[blitEncoder endEncoding];

[commandBuffer commit];
[commandBuffer waitUntilCompleted];


// Get the data back
uint8_t* imageBytes = [texturebuffer contents];

for (int i = 0; i < 16; ++i)
{
    NSLog(@"Pix = %d %d %d %d",
          *((uint8_t*)imageBytes + 4 * i),
          *((uint8_t*)imageBytes + 4 * i + 1),
          *((uint8_t*)imageBytes + 4 * i + 2),
          *((uint8_t*)imageBytes + 4 * i + 3));
}

这两种方法都能正常工作。