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

我需要一个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元素的输入缓冲区，如下所示。这个内核我期望返回的第一个输入的索引大于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>的二进制搜索下界也有同样的问题——然而一个简单的线性版本（有点）有效。顺便说一句，我测试了我从直接C代码中复制的STL版本（显然去掉了device），它们在着色器领域之外工作得很好。请告诉我我做错了什么，这不是一个金属编译器错误。

现在关于上面的"排序"：线性搜索版本适用于5s和mini-2（A7s）（在上例中返回索引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;