我的 iOS 金属计算内核是否存在编译器错误，或者我遗漏了什么？

Question

我需要一个 upper_bound 的实现，如我的金属计算内核的 STL 中所述。金属标准库中没有任何东西，我基本上将它从 <algorithm> 复制到我的着色器文件中，如下所示：

static device float* upper_bound( device float* first, device float* last, float val)
{
    ptrdiff_t count = last - first;
    while( count > 0){
        device float* it = first;
        ptrdiff_t step = count/2;
        it += step;
        if( !(val < *it)){
            first = ++it;
            count -= step + 1;
        }else count = step;
    }
    return first;
}

我创建了一个简单的内核来测试它：

kernel void upper_bound_test(
    device float* input [[buffer(0)]],
    device uint* output [[buffer(1)]]
)
{
    device float* where = upper_bound( input, input + 5, 3.1);
    output[0] = where - input;
}

此测试具有硬编码的输入大小和搜索值。我还在框架端硬编码了一个 5 元素输入缓冲区，如下所示。这个内核我期望return第一个输入的索引大于3.1

没用。事实上 output[0] 从来没有被写入——因为我用一个幻数预加载了缓冲区以查看它是否被覆盖了。它没有。事实上在 waitUntilCompleted 之后，commandBuffer.error 看起来像这样：

Error Domain = MTLCommandBufferErrorDomain
Code = 1
NSLocalizedDescription = "IOAcceleratorFamily returned error code 3"

错误代码 3 是什么意思？我的内核在有机会完成之前就被杀死了吗？

此外，我尝试了 upper_bound 的线性搜索版本，如下所示：

static device float* upper_bound2( device float* first, device float* last, float val)
{
    while( first < last && *first <= val)
        ++first;
    return first;
}

这个有效（有点）。我对 <algorithm> 中的二进制搜索 lower_bound 也有同样的问题——但一个天真的线性版本有效（有点）。顺便说一句，我测试了我从直接 C 代码复制的 STL 版本（显然删除了 device）并且它们在着色器域之外工作正常。请告诉我我做错了什么，这不是金属编译器错误。

关于上面的 "sort-of"：线性搜索版本适用于 5s 和 mini-2 (A7s)（上例中的 returns 索引 3），但适用于 6+ (A8) 它给出了正确答案 + 2^31。有没有搞错！完全相同的代码。请注意，在框架方面我使用 uint32_t，而在着色器方面我使用 uint——它们是同一回事。另请注意，每个指针减法（ptrdiff_t 是有符号的 8 字节的东西）都是小的非负值。为什么 6+ 设置为高位？当然，为什么我的真实二进制搜索版本不起作用？

这是框架方面的东西：

id<MTLFunction> upperBoundTestKernel = [_library newFunctionWithName: @"upper_bound_test"];
id <MTLComputePipelineState> upperBoundTestPipelineState = [_device
    newComputePipelineStateWithFunction: upperBoundTestKernel
    error: &err];


float sortedNumbers[] = {1., 2., 3., 4., 5.};
id<MTLBuffer> testInputBuffer = [_device
    newBufferWithBytes:(const void *)sortedNumbers
    length: sizeof(sortedNumbers)
    options: MTLResourceCPUCacheModeDefaultCache];

id<MTLBuffer> testOutputBuffer = [_device
    newBufferWithLength: sizeof(uint32_t)
    options: MTLResourceCPUCacheModeDefaultCache];

*(uint32_t*)testOutputBuffer.contents = 42;//magic number better get clobbered

id<MTLCommandBuffer> commandBuffer = [_commandQueue commandBuffer];
id<MTLComputeCommandEncoder> commandEncoder = [commandBuffer computeCommandEncoder];
[commandEncoder setComputePipelineState: upperBoundTestPipelineState];
[commandEncoder setBuffer: testInputBuffer offset: 0 atIndex: 0];
[commandEncoder setBuffer: testOutputBuffer offset: 0 atIndex: 1];
[commandEncoder
    dispatchThreadgroups: MTLSizeMake( 1, 1, 1)
    threadsPerThreadgroup: MTLSizeMake( 1, 1, 1)];
[commandEncoder endEncoding];
[commandBuffer commit];
[commandBuffer waitUntilCompleted];

uint32_t answer = *(uint32_t*)testOutputBuffer.contents;

Answer 1

嗯，我找到了 solution/work-around。我猜这是一个 pointer-aliasing 问题，因为 first 和 last 指向同一个缓冲区。所以我将它们更改为单个指针变量的偏移量。这是 re-written upper_bound2:

static uint upper_bound2( device float* input, uint first, uint last, float val)
{
    while( first < last && input[first] <= val)
        ++first;
    return first;
}

还有一个re-written测试内核：

kernel void upper_bound_test(
    device float* input [[buffer(0)]],
    device uint* output [[buffer(1)]]
)
{
    output[0] = upper_bound2( input, 0, 5, 3.1);
}

这完全奏效了。也就是说，它不仅解决了线性搜索的 "sort-of" 问题，而且类似的 re-written 二分搜索也起作用了。我不想相信这一点。金属着色器语言应该是 C++ 的子集，但标准指针语义不起作用？我真的不能比较或减去指针吗？

无论如何，我不记得看到任何文档说不能有指针别名或者什么声明咒语可以帮助我。还有其他帮助吗？

[更新]

郑重声明，正如 "slime" 在 Apple 开发论坛上指出的那样： https://developer.apple.com/library/ios/documentation/Metal/Reference/MetalShadingLanguageGuide/func-var-qual/func-var-qual.html#//apple_ref/doc/uid/TP40014364-CH4-SW3

"Buffers (device and constant) specified as argument values to a graphics or kernel function cannot be aliased—that is, a buffer passed as an argument value cannot overlap another buffer passed to a separate argument of the same graphics or kernel function."

但同样值得注意的是 upper_bound() 不是内核函数并且 upper_bound_test() 没有传递别名参数。 upper_bound_test() 所做的是创建一个本地临时对象，该临时对象指向与其参数之一相同的缓冲区。也许文档应该说明它的含义，例如："No pointer aliasing to device and constant memory in any function is allowed including rvalues." 我实际上不知道这是否太强了。

我的 iOS 金属计算内核是否存在编译器错误，或者我遗漏了什么？

Is there a compiler bug for my iOS metal compute kernel or am I missing something?

c++

stl

ios

metal